Accesses of expressions emitted by `unpack4x{I,U}8` in `{hlsl,msl}-out` cause compilation errors
Closed this issue · 2 comments
ErichDonGubler commented
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;
}
ErichDonGubler commented
This issue causes miscompilations on the DX12 and Metal backends. These manifest as internal errors emitted during pipeline compilation.
ErichDonGubler commented
I first noticed this issue while running WebGPU CTS' webgpu:shader,execution,expression,call,builtin,textureLoad:multisampled:*
in current Firefox Nightly.