Repository: gfx-rs/wgpu
Pull request #6333 (2024-09-27 00:39:15)

[naga] Permit only structs as binding array elements.
Require `T` to be a struct in `binding_array<T, ...>`; do not permit arrays.

In #5428, the validator was changed to accept binding array types that the SPIR-V backend couldn't properly emit. Specifically, the validator was changed to accept `binding_array<array<T>>`, but the SPIR-V backend wasn't changed to wrap the binding array elements in a SPIR-V struct type, as Vulkan requires. So the type would be accepted by the validator, and then rejected by the backend.


ErichDonGubler: CC @teoxoy as reviewer and @kvark as author of the regressor being fixed here.
JMS55: Wait, so can you no longer do `binding_array<array<T>>`? That's very unfortunate. I was using that with ray tracing (soon to be merged...) to bind multiple vertex data buffers easily. I _can_ switch to a single uberbuffer of all mesh data, but it's a lot more annoying to manage.

Patch:
diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs
index 7586f98dc..78e81350b 100644
--- a/naga/src/front/wgsl/lower/mod.rs
+++ b/naga/src/front/wgsl/lower/mod.rs
@@ -1778,12 +1778,14 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
 
                 return Ok(());
             }
-            ast::StatementKind::Ignore(expr) => {
+            ast::StatementKind::Phony(expr) => {
                 let mut emitter = Emitter::default();
                 emitter.start(&ctx.function.expressions);
 
-                let _ = self.expression(expr, &mut ctx.as_expression(block, &mut emitter))?;
+                let value = self.expression(expr, &mut ctx.as_expression(block, &mut emitter))?;
                 block.extend(emitter.finish(&ctx.function.expressions));
+                ctx.named_expressions
+                    .insert(value, ("phony".to_string(), stmt.span));
                 return Ok(());
             }
         };
diff --git a/naga/src/front/wgsl/parse/ast.rs b/naga/src/front/wgsl/parse/ast.rs
index c4a798411..4307ca3d9 100644
--- a/naga/src/front/wgsl/parse/ast.rs
+++ b/naga/src/front/wgsl/parse/ast.rs
@@ -284,7 +284,7 @@ pub enum StatementKind<'a> {
     },
     Increment(Handle<Expression<'a>>),
     Decrement(Handle<Expression<'a>>),
-    Ignore(Handle<Expression<'a>>),
+    Phony(Handle<Expression<'a>>),
     ConstAssert(Handle<Expression<'a>>),
 }
 
diff --git a/naga/src/front/wgsl/parse/mod.rs b/naga/src/front/wgsl/parse/mod.rs
index 0157aa3a7..3b1d60620 100644
--- a/naga/src/front/wgsl/parse/mod.rs
+++ b/naga/src/front/wgsl/parse/mod.rs
@@ -1696,7 +1696,7 @@ impl Parser {
                         let expr = self.general_expression(lexer, ctx)?;
                         lexer.expect(Token::Separator(';'))?;
 
-                        ast::StatementKind::Ignore(expr)
+                        ast::StatementKind::Phony(expr)
                     }
                     "let" => {
                         let _ = lexer.next();
diff --git a/naga/src/valid/type.rs b/naga/src/valid/type.rs
index 0c4466077..c0c25dab7 100644
--- a/naga/src/valid/type.rs
+++ b/naga/src/valid/type.rs
@@ -677,7 +677,6 @@ impl super::Validator {
                     // Currently Naga only supports binding arrays of structs for non-handle types.
                     match gctx.types[base].inner {
                         crate::TypeInner::Struct { .. } => {}
-                        crate::TypeInner::Array { .. } => {}
                         _ => return Err(TypeError::BindingArrayBaseTypeNotStruct(base)),
                     };
                 }


Test patch:
diff --git a/naga/tests/in/phony_assignment.ron b/naga/tests/in/phony_assignment.ron
new file mode 100644
index 000000000..39679b87b
--- /dev/null
+++ b/naga/tests/in/phony_assignment.ron
@@ -0,0 +1,2 @@
+(
+)
\ No newline at end of file
diff --git a/naga/tests/in/phony_assignment.wgsl b/naga/tests/in/phony_assignment.wgsl
new file mode 100644
index 000000000..716a387b0
--- /dev/null
+++ b/naga/tests/in/phony_assignment.wgsl
@@ -0,0 +1,18 @@
+@group(0) @binding(0) var<uniform> binding: f32;
+
+fn five() -> i32 {
+  return 5;
+}
+
+@compute @workgroup_size(1) fn main(
+  @builtin(global_invocation_id) id: vec3<u32>
+) {
+    _ = binding;
+    _ = binding;
+    let a = 5;
+    _ = a;
+    _ = five();
+    let b = five();
+    // check for name collision
+    let phony = binding;
+}
\ No newline at end of file
diff --git a/naga/tests/out/glsl/phony_assignment.main.Compute.glsl b/naga/tests/out/glsl/phony_assignment.main.Compute.glsl
new file mode 100644
index 000000000..6118e9b4a
--- /dev/null
+++ b/naga/tests/out/glsl/phony_assignment.main.Compute.glsl
@@ -0,0 +1,23 @@
+#version 310 es
+
+precision highp float;
+precision highp int;
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+uniform type_block_0Compute { float _group_0_binding_0_cs; };
+
+
+int five() {
+    return 5;
+}
+
+void main() {
+    uvec3 id = gl_GlobalInvocationID;
+    float phony = _group_0_binding_0_cs;
+    float phony_1 = _group_0_binding_0_cs;
+    int _e6 = five();
+    int _e7 = five();
+    float phony_2 = _group_0_binding_0_cs;
+}
+
diff --git a/naga/tests/out/hlsl/phony_assignment.hlsl b/naga/tests/out/hlsl/phony_assignment.hlsl
new file mode 100644
index 000000000..f448290e0
--- /dev/null
+++ b/naga/tests/out/hlsl/phony_assignment.hlsl
@@ -0,0 +1,16 @@
+cbuffer binding : register(b0) { float binding; }
+
+int five()
+{
+    return 5;
+}
+
+[numthreads(1, 1, 1)]
+void main(uint3 id : SV_DispatchThreadID)
+{
+    float phony = binding;
+    float phony_1 = binding;
+    const int _e6 = five();
+    const int _e7 = five();
+    float phony_2 = binding;
+}
diff --git a/naga/tests/out/hlsl/phony_assignment.ron b/naga/tests/out/hlsl/phony_assignment.ron
new file mode 100644
index 000000000..a07b03300
--- /dev/null
+++ b/naga/tests/out/hlsl/phony_assignment.ron
@@ -0,0 +1,12 @@
+(
+    vertex:[
+    ],
+    fragment:[
+    ],
+    compute:[
+        (
+            entry_point:"main",
+            target_profile:"cs_5_1",
+        ),
+    ],
+)
diff --git a/naga/tests/out/msl/phony_assignment.msl b/naga/tests/out/msl/phony_assignment.msl
new file mode 100644
index 000000000..daad57190
--- /dev/null
+++ b/naga/tests/out/msl/phony_assignment.msl
@@ -0,0 +1,24 @@
+// language: metal1.0
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using metal::uint;
+
+
+int five(
+) {
+    return 5;
+}
+
+struct main_Input {
+};
+kernel void main_(
+  metal::uint3 id [[thread_position_in_grid]]
+, constant float& binding [[user(fake0)]]
+) {
+    float phony = binding;
+    float phony_1 = binding;
+    int _e6 = five();
+    int _e7 = five();
+    float phony_2 = binding;
+}
diff --git a/naga/tests/out/spv/phony_assignment.spvasm b/naga/tests/out/spv/phony_assignment.spvasm
new file mode 100644
index 000000000..842618390
--- /dev/null
+++ b/naga/tests/out/spv/phony_assignment.spvasm
@@ -0,0 +1,48 @@
+; SPIR-V
+; Version: 1.1
+; Generator: rspirv
+; Bound: 30
+OpCapability Shader
+%1 = OpExtInstImport "GLSL.std.450"
+OpMemoryModel Logical GLSL450
+OpEntryPoint GLCompute %19 "main" %16
+OpExecutionMode %19 LocalSize 1 1 1
+OpDecorate %7 DescriptorSet 0
+OpDecorate %7 Binding 0
+OpDecorate %8 Block
+OpMemberDecorate %8 0 Offset 0
+OpDecorate %16 BuiltIn GlobalInvocationId
+%2 = OpTypeVoid
+%3 = OpTypeFloat 32
+%4 = OpTypeInt 32 1
+%6 = OpTypeInt 32 0
+%5 = OpTypeVector %6 3
+%8 = OpTypeStruct %3
+%9 = OpTypePointer Uniform %8
+%7 = OpVariable  %9  Uniform
+%12 = OpTypeFunction %4
+%13 = OpConstant  %4  5
+%17 = OpTypePointer Input %5
+%16 = OpVariable  %17  Input
+%20 = OpTypeFunction %2
+%21 = OpTypePointer Uniform %3
+%22 = OpConstant  %6  0
+%11 = OpFunction  %4  None %12
+%10 = OpLabel
+OpBranch %14
+%14 = OpLabel
+OpReturnValue %13
+OpFunctionEnd
+%19 = OpFunction  %2  None %20
+%15 = OpLabel
+%18 = OpLoad  %5  %16
+%23 = OpAccessChain  %21  %7 %22
+OpBranch %24
+%24 = OpLabel
+%25 = OpLoad  %3  %23
+%26 = OpLoad  %3  %23
+%27 = OpFunctionCall  %4  %11
+%28 = OpFunctionCall  %4  %11
+%29 = OpLoad  %3  %23
+OpReturn
+OpFunctionEnd
\ No newline at end of file
diff --git a/naga/tests/out/wgsl/phony_assignment.wgsl b/naga/tests/out/wgsl/phony_assignment.wgsl
new file mode 100644
index 000000000..60987a15c
--- /dev/null
+++ b/naga/tests/out/wgsl/phony_assignment.wgsl
@@ -0,0 +1,15 @@
+@group(0) @binding(0) 
+var<uniform> binding: f32;
+
+fn five() -> i32 {
+    return 5i;
+}
+
+@compute @workgroup_size(1, 1, 1) 
+fn main(@builtin(global_invocation_id) id: vec3<u32>) {
+    let phony = binding;
+    let phony_1 = binding;
+    let _e6 = five();
+    let _e7 = five();
+    let phony_2 = binding;
+}
diff --git a/naga/tests/snapshots.rs b/naga/tests/snapshots.rs
index 936203986..596e4cea1 100644
--- a/naga/tests/snapshots.rs
+++ b/naga/tests/snapshots.rs
@@ -928,6 +928,10 @@ fn convert_wgsl() {
             "cross",
             Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL,
         ),
+        (
+            "phony_assignment",
+            Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL,
+        ),
     ];
 
     for &(name, targets) in inputs.iter() {
diff --git a/naga/tests/validation.rs b/naga/tests/validation.rs
index f20ae688b..b8a30c49b 100644
--- a/naga/tests/validation.rs
+++ b/naga/tests/validation.rs
@@ -496,3 +496,113 @@ fn main(input: VertexOutput) {{
         }
     }
 }
+
+#[allow(dead_code)]
+struct BindingArrayFixture {
+    module: naga::Module,
+    span: naga::Span,
+    ty_u32: naga::Handle<naga::Type>,
+    ty_array: naga::Handle<naga::Type>,
+    ty_struct: naga::Handle<naga::Type>,
+    validator: naga::valid::Validator,
+}
+
+impl BindingArrayFixture {
+    fn new() -> Self {
+        let mut module = naga::Module::default();
+        let span = naga::Span::default();
+        let ty_u32 = module.types.insert(
+            naga::Type {
+                name: Some("u32".into()),
+                inner: naga::TypeInner::Scalar(naga::Scalar::U32),
+            },
+            span,
+        );
+        let ty_array = module.types.insert(
+            naga::Type {
+                name: Some("array<u32, 10>".into()),
+                inner: naga::TypeInner::Array {
+                    base: ty_u32,
+                    size: naga::ArraySize::Constant(std::num::NonZeroU32::new(10).unwrap()),
+                    stride: 4,
+                },
+            },
+            span,
+        );
+        let ty_struct = module.types.insert(
+            naga::Type {
+                name: Some("S".into()),
+                inner: naga::TypeInner::Struct {
+                    members: vec![naga::StructMember {
+                        name: Some("m".into()),
+                        ty: ty_u32,
+                        binding: None,
+                        offset: 0,
+                    }],
+                    span: 4,
+                },
+            },
+            span,
+        );
+        let validator = naga::valid::Validator::new(Default::default(), Default::default());
+        BindingArrayFixture {
+            module,
+            span,
+            ty_u32,
+            ty_array,
+            ty_struct,
+            validator,
+        }
+    }
+}
+
+#[test]
+fn binding_arrays_hold_structs() {
+    let mut t = BindingArrayFixture::new();
+    let _binding_array = t.module.types.insert(
+        naga::Type {
+            name: Some("binding_array_of_struct".into()),
+            inner: naga::TypeInner::BindingArray {
+                base: t.ty_struct,
+                size: naga::ArraySize::Dynamic,
+            },
+        },
+        t.span,
+    );
+
+    assert!(t.validator.validate(&t.module).is_ok());
+}
+
+#[test]
+fn binding_arrays_cannot_hold_arrays() {
+    let mut t = BindingArrayFixture::new();
+    let _binding_array = t.module.types.insert(
+        naga::Type {
+            name: Some("binding_array_of_array".into()),
+            inner: naga::TypeInner::BindingArray {
+                base: t.ty_array,
+                size: naga::ArraySize::Dynamic,
+            },
+        },
+        t.span,
+    );
+
+    assert!(t.validator.validate(&t.module).is_err());
+}
+
+#[test]
+fn binding_arrays_cannot_hold_scalars() {
+    let mut t = BindingArrayFixture::new();
+    let _binding_array = t.module.types.insert(
+        naga::Type {
+            name: Some("binding_array_of_scalar".into()),
+            inner: naga::TypeInner::BindingArray {
+                base: t.ty_u32,
+                size: naga::ArraySize::Dynamic,
+            },
+        },
+        t.span,
+    );
+
+    assert!(t.validator.validate(&t.module).is_err());
+}