tint: work around zero-init bug on some GPUs for workgroup size(1)
This was seen specifically on Windows with an AMD GPU, when using DXC to
compile compute shaders with workgroup size of 1. Flaky, but
reproducible when running CTS test:
webgpu:shader,execution,zero_init:compute,zero_init:addressSpace="workgroup";workgroupSize=[1,1,1];*
We are generating valid code in Tint, but this workaround makes it so
that DXC opts to emit the zero-init instructions rather than tagging the
groupshared variable as 'zeroinitializer'. This seems to make the buggy
driver happy.
Also see: https://github.com/microsoft/DirectXShaderCompiler/issues/6352
Bug: tint:2143
Change-Id: If84b992c4961bafeb1a8ff328626d3f6cf1fb98d
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/176140
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.msl b/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.msl
index 6d50bbd..6f91f64 100644
--- a/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.msl
@@ -2,7 +2,7 @@
using namespace metal;
void tint_zero_workgroup_memory(uint local_idx, threadgroup int* const tint_symbol_1) {
- {
+ if ((local_idx < 1u)) {
*(tint_symbol_1) = 0;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
