[naga] Handle phony statements properly by treating them as named expressions (#6328)

* [naga wgsl-in] phony assignments add named expressions

Signed-off-by: sagudev <16504129+sagudev@users.noreply.github.com>

* [naga wgsl-out] write out _naga_phony as phony

Signed-off-by: sagudev <16504129+sagudev@users.noreply.github.com>

* Add test

Signed-off-by: sagudev <16504129+sagudev@users.noreply.github.com>

* use statement span

Signed-off-by: sagudev <16504129+sagudev@users.noreply.github.com>

* every phony has same name

Signed-off-by: sagudev <16504129+sagudev@users.noreply.github.com>

---------

Signed-off-by: sagudev <16504129+sagudev@users.noreply.github.com>

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(());
             }
         };

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>>),
 }
 

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();

naga/tests/in/phony_assignment.ron

@@ -0,0 +1,2 @@
+(
+)

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;
+}

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;
+}
+

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;
+}

naga/tests/out/hlsl/phony_assignment.ron

@@ -0,0 +1,12 @@
+(
+    vertex:[
+    ],
+    fragment:[
+    ],
+    compute:[
+        (
+            entry_point:"main",
+            target_profile:"cs_5_1",
+        ),
+    ],
+)

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;
+}

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

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;
+}

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() {