fix: make unpack4x{I,U}8 output exprs. composable in `{hlsl,msl}-ou…

…t` (#6773)
gfx-rs · Dec 23, 2024 · 9d9b99a · 9d9b99a
1 parent ee3ae0e
commit 9d9b99a
Show file tree

Hide file tree

Showing 13 changed files with 108 additions and 6 deletions.
diff --git a/naga/src/back/hlsl/writer.rs b/naga/src/back/hlsl/writer.rs
@@ -3294,6 +3294,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                         write!(self.out, " >> 24) / {scale}.0)")?;
                     }
                     fun @ (Function::Unpack4xI8 | Function::Unpack4xU8) => {
+                        write!(self.out, "(")?;
                         if matches!(fun, Function::Unpack4xU8) {
                             write!(self.out, "u")?;
                         }
@@ -3305,7 +3306,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                         self.write_expr(module, arg, func_ctx)?;
                         write!(self.out, " >> 16, ")?;
                         self.write_expr(module, arg, func_ctx)?;
-                        write!(self.out, " >> 24) << 24 >> 24")?;
+                        write!(self.out, " >> 24) << 24 >> 24)")?;
                     }
                     Function::QuantizeToF16 => {
                         write!(self.out, "f16tof32(f32tof16(")?;

diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs
@@ -2134,6 +2134,7 @@ impl<W: Write> Writer<W> {
                         }
                     }
                     fun @ (Mf::Unpack4xI8 | Mf::Unpack4xU8) => {
+                        write!(self.out, "(")?;
                         if matches!(fun, Mf::Unpack4xU8) {
                             write!(self.out, "u")?;
                         }
@@ -2145,7 +2146,7 @@ impl<W: Write> Writer<W> {
                         self.put_expression(arg, context, true)?;
                         write!(self.out, " >> 16, ")?;
                         self.put_expression(arg, context, true)?;
-                        write!(self.out, " >> 24) << 24 >> 24")?;
+                        write!(self.out, " >> 24) << 24 >> 24)")?;
                     }
                     Mf::QuantizeToF16 => {
                         match *context.resolve_type(arg) {

diff --git a/naga/tests/in/6772-unpack-expr-accesses.ron b/naga/tests/in/6772-unpack-expr-accesses.ron
@@ -0,0 +1,2 @@
+(
+)
diff --git a/naga/tests/in/6772-unpack-expr-accesses.wgsl b/naga/tests/in/6772-unpack-expr-accesses.wgsl
@@ -0,0 +1,6 @@
+@compute @workgroup_size(1, 1)
+fn main() {
+    let idx = 2;
+    _ = unpack4xI8(12u)[idx];
+    _ = unpack4xU8(12u)[1];
+}
diff --git a/naga/tests/out/glsl/6772-unpack-expr-accesses.main.Compute.glsl b/naga/tests/out/glsl/6772-unpack-expr-accesses.main.Compute.glsl
@@ -0,0 +1,13 @@
+#version 310 es
+
+precision highp float;
+precision highp int;
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+
+void main() {
+    int phony = ivec4(bitfieldExtract(int(12u), 0, 8), bitfieldExtract(int(12u), 8, 8), bitfieldExtract(int(12u), 16, 8), bitfieldExtract(int(12u), 24, 8))[2];
+    uint phony_1 = uvec4(bitfieldExtract(12u, 0, 8), bitfieldExtract(12u, 8, 8), bitfieldExtract(12u, 16, 8), bitfieldExtract(12u, 24, 8)).y;
+}
+
diff --git a/naga/tests/out/hlsl/6772-unpack-expr-accesses.hlsl b/naga/tests/out/hlsl/6772-unpack-expr-accesses.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void main()
+{
+    int phony = (int4(12u, 12u >> 8, 12u >> 16, 12u >> 24) << 24 >> 24)[2];
+    uint phony_1 = (uint4(12u, 12u >> 8, 12u >> 16, 12u >> 24) << 24 >> 24).y;
+}
diff --git a/naga/tests/out/hlsl/6772-unpack-expr-accesses.ron b/naga/tests/out/hlsl/6772-unpack-expr-accesses.ron
@@ -0,0 +1,12 @@
+(
+    vertex:[
+    ],
+    fragment:[
+    ],
+    compute:[
+        (
+            entry_point:"main",
+            target_profile:"cs_5_1",
+        ),
+    ],
+)
diff --git a/naga/tests/out/hlsl/bits.hlsl b/naga/tests/out/hlsl/bits.hlsl
@@ -213,9 +213,9 @@ void main()
     uint _e50 = u;
     f2_ = float2(f16tof32(_e50), f16tof32((_e50) >> 16));
     uint _e52 = u;
-    i4_ = int4(_e52, _e52 >> 8, _e52 >> 16, _e52 >> 24) << 24 >> 24;
+    i4_ = (int4(_e52, _e52 >> 8, _e52 >> 16, _e52 >> 24) << 24 >> 24);
     uint _e54 = u;
-    u4_ = uint4(_e54, _e54 >> 8, _e54 >> 16, _e54 >> 24) << 24 >> 24;
+    u4_ = (uint4(_e54, _e54 >> 8, _e54 >> 16, _e54 >> 24) << 24 >> 24);
     int _e56 = i;
     int _e57 = i;
     i = naga_insertBits(_e56, _e57, 5u, 10u);

diff --git a/naga/tests/out/msl/6772-unpack-expr-accesses.msl b/naga/tests/out/msl/6772-unpack-expr-accesses.msl
@@ -0,0 +1,12 @@
+// language: metal1.0
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using metal::uint;
+
+
+kernel void main_(
+) {
+    int phony = (int4(12u, 12u >> 8, 12u >> 16, 12u >> 24) << 24 >> 24)[2];
+    uint phony_1 = (uint4(12u, 12u >> 8, 12u >> 16, 12u >> 24) << 24 >> 24).y;
+}
diff --git a/naga/tests/out/msl/bits.msl b/naga/tests/out/msl/bits.msl
@@ -42,9 +42,9 @@ kernel void main_(
     uint _e50 = u;
     f2_ = float2(as_type<half2>(_e50));
     uint _e52 = u;
-    i4_ = int4(_e52, _e52 >> 8, _e52 >> 16, _e52 >> 24) << 24 >> 24;
+    i4_ = (int4(_e52, _e52 >> 8, _e52 >> 16, _e52 >> 24) << 24 >> 24);
     uint _e54 = u;
-    u4_ = uint4(_e54, _e54 >> 8, _e54 >> 16, _e54 >> 24) << 24 >> 24;
+    u4_ = (uint4(_e54, _e54 >> 8, _e54 >> 16, _e54 >> 24) << 24 >> 24);
     int _e56 = i;
     int _e57 = i;
     i = metal::insert_bits(_e56, _e57, metal::min(5u, 32u), metal::min(10u, 32u - metal::min(5u, 32u)));

diff --git a/naga/tests/out/spv/6772-unpack-expr-accesses.spvasm b/naga/tests/out/spv/6772-unpack-expr-accesses.spvasm
@@ -0,0 +1,40 @@
+; SPIR-V
+; Version: 1.1
+; Generator: rspirv
+; Bound: 30
+OpCapability Shader
+%1 = OpExtInstImport "GLSL.std.450"
+OpMemoryModel Logical GLSL450
+OpEntryPoint GLCompute %4 "main"
+OpExecutionMode %4 LocalSize 1 1 1
+%2 = OpTypeVoid
+%5 = OpTypeFunction %2
+%6 = OpTypeInt 32 1
+%7 = OpConstant  %6  2
+%8 = OpTypeInt 32 0
+%9 = OpConstant  %8  12
+%11 = OpTypeVector %6 4
+%13 = OpConstant  %8  8
+%19 = OpConstant  %8  0
+%20 = OpConstant  %8  16
+%21 = OpConstant  %8  24
+%23 = OpTypeVector %8 4
+%4 = OpFunction  %2  None %5
+%3 = OpLabel
+OpBranch %10
+%10 = OpLabel
+%14 = OpBitcast  %6  %9
+%15 = OpBitFieldSExtract  %6  %14 %19 %13
+%16 = OpBitFieldSExtract  %6  %14 %13 %13
+%17 = OpBitFieldSExtract  %6  %14 %20 %13
+%18 = OpBitFieldSExtract  %6  %14 %21 %13
+%12 = OpCompositeConstruct  %11  %15 %16 %17 %18
+%22 = OpCompositeExtract  %6  %12 2
+%25 = OpBitFieldUExtract  %8  %9 %19 %13
+%26 = OpBitFieldUExtract  %8  %9 %13 %13
+%27 = OpBitFieldUExtract  %8  %9 %20 %13
+%28 = OpBitFieldUExtract  %8  %9 %21 %13
+%24 = OpCompositeConstruct  %23  %25 %26 %27 %28
+%29 = OpCompositeExtract  %8  %24 1
+OpReturn
+OpFunctionEnd
diff --git a/naga/tests/out/wgsl/6772-unpack-expr-accesses.wgsl b/naga/tests/out/wgsl/6772-unpack-expr-accesses.wgsl
@@ -0,0 +1,5 @@
+@compute @workgroup_size(1, 1, 1) 
+fn main() {
+    let phony = unpack4xI8(12u)[2i];
+    let phony_1 = unpack4xU8(12u).y;
+}
diff --git a/naga/tests/snapshots.rs b/naga/tests/snapshots.rs
@@ -945,6 +945,10 @@ fn convert_wgsl() {
             Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL,
         ),
         ("diagnostic-filter", Targets::IR),
+        (
+            "6772-unpack-expr-accesses",
+            Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL,
+        ),
     ];
 
     for &(name, targets) in inputs.iter() {