workgroupUniformLoad specialization for atomic<T> is missing. #8785
Description
Description
When given an atomic pointer, workgroupUniformLoad should return the inner T.
This works in Chrome (presumably using the tint / Dawn backend), but fails with naga / wgpu, which returns the atomic type instead.
Repro steps
Here is a simple shader:
var<workgroup> wg_scratch: atomic<u32>;
@compute @workgroup_size(4, 4, 4)
fn interval_tile_main(
@builtin(workgroup_id) workgroup_id: vec3u,
@builtin(local_invocation_id) local_id: vec3u
) {
let active_tile_index = workgroup_id.x + workgroup_id.y * 32768;
atomicOr(&wg_scratch, u32(active_tile_index >= 64));
workgroupBarrier();
if workgroupUniformLoad(&wg_scratch) == 0 {
return;
}
}It fails to parse in naga:
$ naga --version
28.0.0
$ naga shader.wgsl
Could not parse WGSL:
error: Incompatible operands: Equal(Atomic(Scalar { kind: Uint, width: 4 }), _)Expected vs observed behavior
This should be a valid WGSL program that naga can parse.
Platform
macOS 14.7.6 (23H626), naga 28.0.0