gfx-rs/wgpu

Accesses of expressions emitted by `unpack4x{I,U}8` in `{hlsl,msl}-out` cause compilation errors

Closed this issue · 2 comments

unpack4x{I,U}8 emits some bit-shifting expressions that aren't composable with further accesses of the members of the emitted uint4. For example, the following WGSL:

@compute @workgroup_size(1, 1)
fn main() {
	let idx = 2;
	_ = unpack4xI8(12u)[idx];
	_ = unpack4xU8(12u)[1];
}
// 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;
}

This needs to be properly parenthesized, i.e.:

@@ -7,6 +7,6 @@ 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;
+    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;
 }

This issue causes miscompilations on the DX12 and Metal backends. These manifest as internal errors emitted during pipeline compilation.

I first noticed this issue while running WebGPU CTS' webgpu:shader,execution,expression,call,builtin,textureLoad:multisampled:* in current Firefox Nightly.