Fix regressions in logical operator handling in const exprs (gfx-rs#8721)

* Fix regressions in logical operator handling in const exprs

Fixes regressions introduced by gfx-rs#7339
Fixes gfx-rs#8711

Author: andyleiserson

cts_runner/test.lst

@@ -177,6 +177,10 @@ webgpu:shader,execution,flow_control,return:*
 // Many other vertex_buffer_access subtests also passing, but there are too many to enumerate.
 // Fails on Metal in CI only, not when running locally.
 fails-if(metal) webgpu:shader,execution,robust_access_vertex:vertex_buffer_access:indexed=true;indirect=false;drawCallTestParameter="baseVertex";type="float32x4";additionalBuffers=4;partialLastNumber=false;offsetVertexBuffer=true
+webgpu:shader,validation,const_assert,const_assert:*
+webgpu:shader,validation,expression,binary,short_circuiting_and_or:array_override:op="%26%26";a_val=1;b_val=1
+webgpu:shader,validation,expression,binary,short_circuiting_and_or:invalid_types:*
+webgpu:shader,validation,expression,binary,short_circuiting_and_or:scalar_vector:op="%26%26";lhs="bool";rhs="bool"
 webgpu:shader,validation,expression,call,builtin,all:arguments:test="ptr_deref"
 webgpu:shader,validation,expression,call,builtin,max:values:*
 webgpu:shader,validation,statement,statement_behavior:invalid_statements:body="break"

naga/src/front/wgsl/lower/mod.rs

@@ -560,6 +560,28 @@ impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> {
         }
     }
 
+    fn const_eval_expr_to_bool(&self, handle: Handle<ir::Expression>) -> Option<bool> {
+        match self.expr_type {
+            ExpressionContextType::Runtime(ref ctx) => {
+                if !ctx.local_expression_kind_tracker.is_const(handle) {
+                    return None;
+                }
+
+                self.module
+                    .to_ctx()
+                    .eval_expr_to_bool_from(handle, &ctx.function.expressions)
+            }
+            ExpressionContextType::Constant(Some(ref ctx)) => {
+                assert!(ctx.local_expression_kind_tracker.is_const(handle));
+                self.module
+                    .to_ctx()
+                    .eval_expr_to_bool_from(handle, &ctx.function.expressions)
+            }
+            ExpressionContextType::Constant(None) => self.module.to_ctx().eval_expr_to_bool(handle),
+            ExpressionContextType::Override => None,
+        }
+    }
+
     /// Return `true` if `handle` is a constant expression.
     fn is_const(&self, handle: Handle<ir::Expression>) -> bool {
         use ExpressionContextType as Ect;
@@ -2538,22 +2560,27 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
                 result_var,
             )))
         } else {
-            let left_expr = ctx.get(left);
-            // Constant or override context in either function or module scope
-            let &crate::Expression::Literal(crate::Literal::Bool(left_val)) = left_expr else {
-                return Err(Box::new(Error::NotBool(span)));
-            };
+            let left_val = ctx.const_eval_expr_to_bool(left);
+
+            if left_val.is_some_and(|left_val| {
+                op == crate::BinaryOperator::LogicalAnd && !left_val
+                    || op == crate::BinaryOperator::LogicalOr && left_val
+            }) {
+                // Short-circuit behavior: don't evaluate the RHS.
 
-            if op == crate::BinaryOperator::LogicalAnd && !left_val
-                || op == crate::BinaryOperator::LogicalOr && left_val
-            {
-                // Short-circuit behavior: don't evaluate the RHS. Ideally we
-                // would do _some_ validity checks of the RHS here, but that's
-                // tricky, because the RHS is allowed to have things that aren't
-                // legal in const contexts.
+                // TODO(https://github.com/gfx-rs/wgpu/issues/8440): We shouldn't ignore the
+                // RHS completely, it should still be type-checked. Preserving it for type
+                // checking is a bit tricky, because we're trying to produce an expression
+                // for a const context, but the RHS is allowed to have things that aren't
+                // const.
 
-                Ok(Typed::Plain(left_expr.clone()))
+                Ok(Typed::Plain(ctx.get(left).clone()))
             } else {
+                // Evaluate the RHS and construct the entire binary expression as we
+                // normally would. This case applies to well-formed constant logical
+                // expressions that don't short-circuit (handled by the constant evaluator
+                // shortly), to override expressions (handled when overrides are processed)
+                // and to non-well-formed expressions (rejected by type checking).
                 let right = self.expression_for_abstract(right, ctx)?;
                 ctx.grow_types(right)?;
 

naga/src/proc/mod.rs

@@ -440,7 +440,18 @@ pub struct GlobalCtx<'a> {
 
 impl GlobalCtx<'_> {
     /// Try to evaluate the expression in `self.global_expressions` using its `handle` and return it as a `u32`.
-    #[allow(dead_code)]
+    #[cfg_attr(
+        not(any(
+            feature = "glsl-in",
+            feature = "spv-in",
+            feature = "wgsl-in",
+            glsl_out,
+            hlsl_out,
+            msl_out,
+            wgsl_out
+        )),
+        allow(dead_code)
+    )]
     pub(super) fn eval_expr_to_u32(
         &self,
         handle: crate::Handle<crate::Expression>,
@@ -463,8 +474,17 @@ impl GlobalCtx<'_> {
         }
     }
 
+    /// Try to evaluate the expression in `self.global_expressions` using its `handle` and return it as a `bool`.
+    #[cfg_attr(not(feature = "wgsl-in"), allow(dead_code))]
+    pub(super) fn eval_expr_to_bool(
+        &self,
+        handle: crate::Handle<crate::Expression>,
+    ) -> Option<bool> {
+        self.eval_expr_to_bool_from(handle, self.global_expressions)
+    }
+
     /// Try to evaluate the expression in the `arena` using its `handle` and return it as a `bool`.
-    #[allow(dead_code)]
+    #[cfg_attr(not(feature = "wgsl-in"), allow(dead_code))]
     pub(super) fn eval_expr_to_bool_from(
         &self,
         handle: crate::Handle<crate::Expression>,
@@ -476,7 +496,7 @@ impl GlobalCtx<'_> {
         }
     }
 
-    #[allow(dead_code)]
+    #[expect(dead_code)]
     pub(crate) fn eval_expr_to_literal(
         &self,
         handle: crate::Handle<crate::Expression>,

naga/tests/in/wgsl/const_assert.wgsl

@@ -1,4 +1,4 @@
-// Sourced from https://www.w3.org/TR/WGSL/#const-assert-statement
+// Sourced from https://www.w3.org/TR/WGSL/#const-assert-statement and the CTS.
 const x = 1;
 const y = 2;
 const_assert x < y; // valid at module-scope.
@@ -9,13 +9,20 @@ const_assert x == 1i;
 const_assert x > 0u;
 const_assert x < 2.0f;
 
+const g_false = false;
+const_assert (!((g_false) || (any(vec3(false, false, false)))));
+
 @compute @workgroup_size(1)
 fn foo() {
   const z = x + y - 2;
+  const l_false = false;
   const_assert z > 0; // valid in functions.
   const_assert(z > 0);
 
   const_assert z == 1i;
   const_assert z > 0u;
   const_assert z < 2.0f;
+
+  const_assert (!((g_false) || (any(vec3(false, false, false)))));
+  const_assert (!((l_false) || (any(vec3(false, false, false)))));
 }

naga/tests/in/wgsl/operators.wgsl

@@ -3,6 +3,9 @@ const v_f32_zero = vec4<f32>(0.0, 0.0, 0.0, 0.0);
 const v_f32_half = vec4<f32>(0.5, 0.5, 0.5, 0.5);
 const v_i32_one = vec4<i32>(1, 1, 1, 1);
 
+const b_false = false;
+const b_true = true;
+
 fn builtins() -> vec4<f32> {
     // select()
     let condition = true;
@@ -45,6 +48,14 @@ fn q() -> bool { return false; }
 fn r() -> bool { return true; }
 fn s() -> bool { return false; }
 
+const short_circuit_1_invalid_rhs = false && sqrt(-1) != 0;
+// TODO(https://github.com/gfx-rs/wgpu/issues/8440):
+// The following should not be accepted, but it currently is.
+// When fixed, move this to a wgsl_errors test.
+const short_circuit_2_invalid_rhs = false && (0u + 1.0f > 0);
+const short_circuit_3 = b_false || b_true;
+const short_circuit_4 = !((b_false) || (any(vec3(false, false, false))));
+
 fn logical() {
     let t = true;
     let f = false;
@@ -60,7 +71,18 @@ fn logical() {
     let bitwise_or1 = vec3(t) | vec3(f);
     let bitwise_and0 = t & f;
     let bitwise_and1 = vec4(t) & vec4(f);
-    let short_circuit = (p() || q()) && (r() || s());
+
+    const short_circuit_1_invalid = false && sqrt(-1) != 0;
+    // TODO(https://github.com/gfx-rs/wgpu/issues/8440):
+    // The following should not be accepted, but it currently is.
+    // When fixed, move this to a wgsl_errors test.
+    const short_circuit_2_invalid_rhs = false && (0u + 1.0f > 0);
+    const short_circuit_3 = b_false || b_true;
+    const short_circuit_4 = !((b_false) || (any(vec3(false, false, false))));
+
+    let short_circuit_5 = !((f) || (any(vec3(false, false, false))));
+    let short_circuit_6 = (p() || q()) && (r() || s());
+    let short_circuit_7 = true || q();
 }
 
 fn arithmetic() {
@@ -330,4 +352,3 @@ fn main(@builtin(workgroup_id) id: vec3<u32>) {
 
     negation_avoids_prefix_decrement();
 }
-

naga/tests/out/glsl/wgsl-operators.main.Compute.glsl

@@ -9,6 +9,12 @@ const vec4 v_f32_one = vec4(1.0, 1.0, 1.0, 1.0);
 const vec4 v_f32_zero = vec4(0.0, 0.0, 0.0, 0.0);
 const vec4 v_f32_half = vec4(0.5, 0.5, 0.5, 0.5);
 const ivec4 v_i32_one = ivec4(1, 1, 1, 1);
+const bool b_false = false;
+const bool b_true = true;
+const bool short_circuit_1_invalid_rhs = false;
+const bool short_circuit_2_invalid_rhs = false;
+const bool short_circuit_3_ = true;
+const bool short_circuit_4_ = true;
 
 
 vec4 builtins() {
@@ -68,6 +74,8 @@ void logical() {
     bool local_2 = false;
     bool local_3 = false;
     bool local_4 = false;
+    bool local_5 = false;
+    bool local_6 = false;
     bool neg0_ = !(true);
     bvec2 neg1_ = not(bvec2(true));
     if (!(true)) {
@@ -86,28 +94,42 @@ void logical() {
     bvec3 bitwise_or1_ = bvec3(bvec3(true).x || bvec3(false).x, bvec3(true).y || bvec3(false).y, bvec3(true).z || bvec3(false).z);
     bool bitwise_and0_ = (true && false);
     bvec4 bitwise_and1_ = bvec4(bvec4(true).x && bvec4(false).x, bvec4(true).y && bvec4(false).y, bvec4(true).z && bvec4(false).z, bvec4(true).w && bvec4(false).w);
-    bool _e22 = p();
-    if (!(_e22)) {
-        bool _e26 = q();
-        local_2 = _e26;
+    if (!(false)) {
+        local_2 = false;
     } else {
         local_2 = true;
     }
-    bool _e28 = local_2;
-    if (_e28) {
-        bool _e31 = r();
-        if (!(_e31)) {
-            bool _e35 = s();
-            local_4 = _e35;
+    bool _e27 = local_2;
+    bool short_circuit_5_ = !(_e27);
+    bool _e29 = p();
+    if (!(_e29)) {
+        bool _e33 = q();
+        local_3 = _e33;
+    } else {
+        local_3 = true;
+    }
+    bool _e35 = local_3;
+    if (_e35) {
+        bool _e38 = r();
+        if (!(_e38)) {
+            bool _e42 = s();
+            local_5 = _e42;
         } else {
-            local_4 = true;
+            local_5 = true;
         }
-        bool _e37 = local_4;
-        local_3 = _e37;
+        bool _e44 = local_5;
+        local_4 = _e44;
+    } else {
+        local_4 = false;
+    }
+    bool short_circuit_6_ = local_4;
+    if (false) {
+        bool _e50 = q();
+        local_6 = _e50;
     } else {
-        local_3 = false;
+        local_6 = true;
     }
-    bool short_circuit = local_3;
+    bool short_circuit_7_ = local_6;
     return;
 }
 

naga/tests/out/hlsl/wgsl-operators.hlsl

@@ -2,6 +2,12 @@ static const float4 v_f32_one = float4(1.0, 1.0, 1.0, 1.0);
 static const float4 v_f32_zero = float4(0.0, 0.0, 0.0, 0.0);
 static const float4 v_f32_half = float4(0.5, 0.5, 0.5, 0.5);
 static const int4 v_i32_one = int4(int(1), int(1), int(1), int(1));
+static const bool b_false = false;
+static const bool b_true = true;
+static const bool short_circuit_1_invalid_rhs = false;
+static const bool short_circuit_2_invalid_rhs = false;
+static const bool short_circuit_3_ = true;
+static const bool short_circuit_4_ = true;
 
 float4 builtins()
 {
@@ -75,6 +81,8 @@ void logical()
     bool local_2 = (bool)0;
     bool local_3 = (bool)0;
     bool local_4 = (bool)0;
+    bool local_5 = (bool)0;
+    bool local_6 = (bool)0;
 
     bool neg0_ = !(true);
     bool2 neg1_ = !((true).xx);
@@ -94,28 +102,42 @@ void logical()
     bool3 bitwise_or1_ = ((true).xxx | (false).xxx);
     bool bitwise_and0_ = (true & false);
     bool4 bitwise_and1_ = ((true).xxxx & (false).xxxx);
-    const bool _e22 = p();
-    if (!(_e22)) {
-        const bool _e26 = q();
-        local_2 = _e26;
+    if (!(false)) {
+        local_2 = false;
     } else {
         local_2 = true;
     }
-    bool _e28 = local_2;
-    if (_e28) {
-        const bool _e31 = r();
-        if (!(_e31)) {
-            const bool _e35 = s();
-            local_4 = _e35;
+    bool _e27 = local_2;
+    bool short_circuit_5_ = !(_e27);
+    const bool _e29 = p();
+    if (!(_e29)) {
+        const bool _e33 = q();
+        local_3 = _e33;
+    } else {
+        local_3 = true;
+    }
+    bool _e35 = local_3;
+    if (_e35) {
+        const bool _e38 = r();
+        if (!(_e38)) {
+            const bool _e42 = s();
+            local_5 = _e42;
         } else {
-            local_4 = true;
+            local_5 = true;
         }
-        bool _e37 = local_4;
-        local_3 = _e37;
+        bool _e44 = local_5;
+        local_4 = _e44;
+    } else {
+        local_4 = false;
+    }
+    bool short_circuit_6_ = local_4;
+    if (false) {
+        const bool _e50 = q();
+        local_6 = _e50;
     } else {
-        local_3 = false;
+        local_6 = true;
     }
-    bool short_circuit = local_3;
+    bool short_circuit_7_ = local_6;
     return;
 }