diff --git a/naga/src/back/hlsl/writer.rs b/naga/src/back/hlsl/writer.rs index bc6086d539..459798c771 100644 --- 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 index f4b55bbbc5..dc444bddb5 100644 --- a/naga/src/back/msl/writer.rs +++ b/naga/src/back/msl/writer.rs @@ -2134,6 +2134,7 @@ impl Writer { } } fun @ (Mf::Unpack4xI8 | Mf::Unpack4xU8) => { + write!(self.out, "(")?; if matches!(fun, Mf::Unpack4xU8) { write!(self.out, "u")?; } @@ -2145,7 +2146,7 @@ impl Writer { 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 new file mode 100644 index 0000000000..72873dd667 --- /dev/null +++ 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 new file mode 100644 index 0000000000..ad2f8b8ffc --- /dev/null +++ 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 new file mode 100644 index 0000000000..e857de73fe --- /dev/null +++ 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 new file mode 100644 index 0000000000..e50f40c8fc --- /dev/null +++ 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 new file mode 100644 index 0000000000..a07b03300b --- /dev/null +++ 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 index c89eb19efe..1987e15514 100644 --- 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 new file mode 100644 index 0000000000..e00a1b4c54 --- /dev/null +++ b/naga/tests/out/msl/6772-unpack-expr-accesses.msl @@ -0,0 +1,12 @@ +// language: metal1.0 +#include +#include + +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 index 02613fcc04..a0138862d5 100644 --- 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(_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 new file mode 100644 index 0000000000..973557789e --- /dev/null +++ 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 \ No newline at end of file diff --git a/naga/tests/out/wgsl/6772-unpack-expr-accesses.wgsl b/naga/tests/out/wgsl/6772-unpack-expr-accesses.wgsl new file mode 100644 index 0000000000..4c13bc50f3 --- /dev/null +++ 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 index e277ae7456..6a234a0977 100644 --- 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() {