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/local/ptr_workgroup.wgsl.expected.dxc.hlsl b/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.dxc.hlsl
index e1120a7..4c5de49 100644
--- a/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.dxc.hlsl
+++ b/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.dxc.hlsl
@@ -1,7 +1,7 @@
groupshared int i;
void tint_zero_workgroup_memory(uint local_idx) {
- {
+ if ((local_idx < 1u)) {
i = 0;
}
GroupMemoryBarrierWithGroupSync();
diff --git a/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.fxc.hlsl b/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.fxc.hlsl
index e1120a7..4c5de49 100644
--- a/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.fxc.hlsl
+++ b/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.fxc.hlsl
@@ -1,7 +1,7 @@
groupshared int i;
void tint_zero_workgroup_memory(uint local_idx) {
- {
+ if ((local_idx < 1u)) {
i = 0;
}
GroupMemoryBarrierWithGroupSync();
diff --git a/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.glsl b/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.glsl
index 1be2d8b..88dff26 100644
--- a/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.glsl
+++ b/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.glsl
@@ -2,7 +2,7 @@
shared int i;
void tint_zero_workgroup_memory(uint local_idx) {
- {
+ if ((local_idx < 1u)) {
i = 0;
}
barrier();
diff --git a/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.msl b/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.msl
index a09a9a6..4f4caf8 100644
--- a/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.msl
+++ b/test/tint/ptr_ref/load/local/ptr_workgroup.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);
diff --git a/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.spvasm b/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.spvasm
index 7b352f7..5db10ce 100644
--- a/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.spvasm
+++ b/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 29
+; Bound: 34
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -23,31 +23,39 @@
%i = OpVariable %_ptr_Workgroup_int Workgroup
%void = OpTypeVoid
%7 = OpTypeFunction %void %uint
- %12 = OpConstantNull %int
+ %uint_1 = OpConstant %uint 1
+ %bool = OpTypeBool
+ %17 = OpConstantNull %int
%uint_2 = OpConstant %uint 2
%uint_264 = OpConstant %uint 264
%int_123 = OpConstant %int 123
%int_1 = OpConstant %int 1
- %24 = OpTypeFunction %void
+ %29 = OpTypeFunction %void
%tint_zero_workgroup_memory = OpFunction %void None %7
%local_idx = OpFunctionParameter %uint
%11 = OpLabel
- OpStore %i %12
+ %13 = OpULessThan %bool %local_idx %uint_1
+ OpSelectionMerge %15 None
+ OpBranchConditional %13 %16 %15
+ %16 = OpLabel
+ OpStore %i %17
+ OpBranch %15
+ %15 = OpLabel
OpControlBarrier %uint_2 %uint_2 %uint_264
OpReturn
OpFunctionEnd
%main_inner = OpFunction %void None %7
%local_invocation_index = OpFunctionParameter %uint
- %18 = OpLabel
- %19 = OpFunctionCall %void %tint_zero_workgroup_memory %local_invocation_index
+ %23 = OpLabel
+ %24 = OpFunctionCall %void %tint_zero_workgroup_memory %local_invocation_index
OpStore %i %int_123
- %21 = OpLoad %int %i
- %23 = OpIAdd %int %21 %int_1
+ %26 = OpLoad %int %i
+ %28 = OpIAdd %int %26 %int_1
OpReturn
OpFunctionEnd
- %main = OpFunction %void None %24
- %26 = OpLabel
- %28 = OpLoad %uint %local_invocation_index_1
- %27 = OpFunctionCall %void %main_inner %28
+ %main = OpFunction %void None %29
+ %31 = OpLabel
+ %33 = OpLoad %uint %local_invocation_index_1
+ %32 = OpFunctionCall %void %main_inner %33
OpReturn
OpFunctionEnd
diff --git a/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.dxc.hlsl b/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.dxc.hlsl
index b581764..8ce73eb 100644
--- a/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.dxc.hlsl
+++ b/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.dxc.hlsl
@@ -1,7 +1,7 @@
groupshared int S;
void tint_zero_workgroup_memory(uint local_idx) {
- {
+ if ((local_idx < 1u)) {
S = 0;
}
GroupMemoryBarrierWithGroupSync();
diff --git a/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.fxc.hlsl b/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.fxc.hlsl
index b581764..8ce73eb 100644
--- a/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.fxc.hlsl
+++ b/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.fxc.hlsl
@@ -1,7 +1,7 @@
groupshared int S;
void tint_zero_workgroup_memory(uint local_idx) {
- {
+ if ((local_idx < 1u)) {
S = 0;
}
GroupMemoryBarrierWithGroupSync();
diff --git a/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.glsl b/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.glsl
index 66fd5ad..61f38d0 100644
--- a/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.glsl
+++ b/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.glsl
@@ -2,7 +2,7 @@
shared int S;
void tint_zero_workgroup_memory(uint local_idx) {
- {
+ if ((local_idx < 1u)) {
S = 0;
}
barrier();
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);
diff --git a/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.spvasm b/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.spvasm
index 25f57ba..b74dcdf 100644
--- a/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.spvasm
+++ b/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 30
+; Bound: 35
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -24,33 +24,41 @@
%S = OpVariable %_ptr_Workgroup_int Workgroup
%void = OpTypeVoid
%7 = OpTypeFunction %void %uint
- %12 = OpConstantNull %int
+ %uint_1 = OpConstant %uint 1
+ %bool = OpTypeBool
+ %17 = OpConstantNull %int
%uint_2 = OpConstant %uint 2
%uint_264 = OpConstant %uint 264
- %16 = OpTypeFunction %int
- %25 = OpTypeFunction %void
+ %21 = OpTypeFunction %int
+ %30 = OpTypeFunction %void
%tint_zero_workgroup_memory = OpFunction %void None %7
%local_idx = OpFunctionParameter %uint
%11 = OpLabel
- OpStore %S %12
+ %13 = OpULessThan %bool %local_idx %uint_1
+ OpSelectionMerge %15 None
+ OpBranchConditional %13 %16 %15
+ %16 = OpLabel
+ OpStore %S %17
+ OpBranch %15
+ %15 = OpLabel
OpControlBarrier %uint_2 %uint_2 %uint_264
OpReturn
OpFunctionEnd
- %func_S = OpFunction %int None %16
- %18 = OpLabel
- %19 = OpLoad %int %S
- OpReturnValue %19
+ %func_S = OpFunction %int None %21
+ %23 = OpLabel
+ %24 = OpLoad %int %S
+ OpReturnValue %24
OpFunctionEnd
%main_inner = OpFunction %void None %7
%local_invocation_index = OpFunctionParameter %uint
- %22 = OpLabel
- %23 = OpFunctionCall %void %tint_zero_workgroup_memory %local_invocation_index
- %24 = OpFunctionCall %int %func_S
+ %27 = OpLabel
+ %28 = OpFunctionCall %void %tint_zero_workgroup_memory %local_invocation_index
+ %29 = OpFunctionCall %int %func_S
OpReturn
OpFunctionEnd
- %main = OpFunction %void None %25
- %27 = OpLabel
- %29 = OpLoad %uint %local_invocation_index_1
- %28 = OpFunctionCall %void %main_inner %29
+ %main = OpFunction %void None %30
+ %32 = OpLabel
+ %34 = OpLoad %uint %local_invocation_index_1
+ %33 = OpFunctionCall %void %main_inner %34
OpReturn
OpFunctionEnd
diff --git a/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.dxc.hlsl b/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.dxc.hlsl
index b1faf09..798e60c 100644
--- a/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.dxc.hlsl
+++ b/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.dxc.hlsl
@@ -5,7 +5,7 @@
groupshared str S;
void tint_zero_workgroup_memory(uint local_idx) {
- {
+ if ((local_idx < 1u)) {
str tint_symbol_2 = (str)0;
S = tint_symbol_2;
}
diff --git a/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.fxc.hlsl b/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.fxc.hlsl
index b1faf09..798e60c 100644
--- a/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.fxc.hlsl
+++ b/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.fxc.hlsl
@@ -5,7 +5,7 @@
groupshared str S;
void tint_zero_workgroup_memory(uint local_idx) {
- {
+ if ((local_idx < 1u)) {
str tint_symbol_2 = (str)0;
S = tint_symbol_2;
}
diff --git a/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.glsl b/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.glsl
index fef574e..e87a8b9 100644
--- a/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.glsl
+++ b/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.glsl
@@ -6,7 +6,7 @@
shared str S;
void tint_zero_workgroup_memory(uint local_idx) {
- {
+ if ((local_idx < 1u)) {
str tint_symbol_1 = str(0);
S = tint_symbol_1;
}
diff --git a/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.msl b/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.msl
index c602af7..9f3be99 100644
--- a/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.msl
@@ -6,7 +6,7 @@
};
void tint_zero_workgroup_memory(uint local_idx, threadgroup str* const tint_symbol_2) {
- {
+ if ((local_idx < 1u)) {
str const tint_symbol_1 = str{};
*(tint_symbol_2) = tint_symbol_1;
}
diff --git a/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.spvasm b/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.spvasm
index f6cc01b..f08514e 100644
--- a/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.spvasm
+++ b/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 34
+; Bound: 39
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -28,36 +28,44 @@
%S = OpVariable %_ptr_Workgroup_str Workgroup
%void = OpTypeVoid
%8 = OpTypeFunction %void %uint
- %13 = OpConstantNull %str
+ %uint_1 = OpConstant %uint 1
+ %bool = OpTypeBool
+ %18 = OpConstantNull %str
%uint_2 = OpConstant %uint 2
%uint_264 = OpConstant %uint 264
- %17 = OpTypeFunction %int
+ %22 = OpTypeFunction %int
%uint_0 = OpConstant %uint 0
%_ptr_Workgroup_int = OpTypePointer Workgroup %int
- %29 = OpTypeFunction %void
+ %34 = OpTypeFunction %void
%tint_zero_workgroup_memory = OpFunction %void None %8
%local_idx = OpFunctionParameter %uint
%12 = OpLabel
- OpStore %S %13
+ %14 = OpULessThan %bool %local_idx %uint_1
+ OpSelectionMerge %16 None
+ OpBranchConditional %14 %17 %16
+ %17 = OpLabel
+ OpStore %S %18
+ OpBranch %16
+ %16 = OpLabel
OpControlBarrier %uint_2 %uint_2 %uint_264
OpReturn
OpFunctionEnd
- %func_S_i = OpFunction %int None %17
- %19 = OpLabel
- %22 = OpAccessChain %_ptr_Workgroup_int %S %uint_0
- %23 = OpLoad %int %22
- OpReturnValue %23
+ %func_S_i = OpFunction %int None %22
+ %24 = OpLabel
+ %27 = OpAccessChain %_ptr_Workgroup_int %S %uint_0
+ %28 = OpLoad %int %27
+ OpReturnValue %28
OpFunctionEnd
%main_inner = OpFunction %void None %8
%local_invocation_index = OpFunctionParameter %uint
- %26 = OpLabel
- %27 = OpFunctionCall %void %tint_zero_workgroup_memory %local_invocation_index
- %28 = OpFunctionCall %int %func_S_i
+ %31 = OpLabel
+ %32 = OpFunctionCall %void %tint_zero_workgroup_memory %local_invocation_index
+ %33 = OpFunctionCall %int %func_S_i
OpReturn
OpFunctionEnd
- %main = OpFunction %void None %29
- %31 = OpLabel
- %33 = OpLoad %uint %local_invocation_index_1
- %32 = OpFunctionCall %void %main_inner %33
+ %main = OpFunction %void None %34
+ %36 = OpLabel
+ %38 = OpLoad %uint %local_invocation_index_1
+ %37 = OpFunctionCall %void %main_inner %38
OpReturn
OpFunctionEnd
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.dxc.hlsl b/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.dxc.hlsl
index ab51c8863..f05c134 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.dxc.hlsl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.dxc.hlsl
@@ -1,7 +1,7 @@
groupshared float2x2 S;
void tint_zero_workgroup_memory(uint local_idx) {
- {
+ if ((local_idx < 1u)) {
S = float2x2((0.0f).xx, (0.0f).xx);
}
GroupMemoryBarrierWithGroupSync();
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.fxc.hlsl b/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.fxc.hlsl
index ab51c8863..f05c134 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.fxc.hlsl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.fxc.hlsl
@@ -1,7 +1,7 @@
groupshared float2x2 S;
void tint_zero_workgroup_memory(uint local_idx) {
- {
+ if ((local_idx < 1u)) {
S = float2x2((0.0f).xx, (0.0f).xx);
}
GroupMemoryBarrierWithGroupSync();
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.glsl b/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.glsl
index 9c0bd23..93f5bc8 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.glsl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.glsl
@@ -2,7 +2,7 @@
shared mat2 S;
void tint_zero_workgroup_memory(uint local_idx) {
- {
+ if ((local_idx < 1u)) {
S = mat2(vec2(0.0f), vec2(0.0f));
}
barrier();
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.msl b/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.msl
index 68817c2..7a56bbc 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.msl
@@ -6,7 +6,7 @@
};
void tint_zero_workgroup_memory(uint local_idx, threadgroup float2x2* const tint_symbol_1) {
- {
+ if ((local_idx < 1u)) {
*(tint_symbol_1) = float2x2(float2(0.0f), float2(0.0f));
}
threadgroup_barrier(mem_flags::mem_threadgroup);
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.spvasm b/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.spvasm
index 7125afd..719b833 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.spvasm
+++ b/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 41
+; Bound: 45
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -28,42 +28,49 @@
%S = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
%void = OpTypeVoid
%9 = OpTypeFunction %void %uint
- %14 = OpConstantNull %mat2v2float
+ %uint_1 = OpConstant %uint 1
+ %bool = OpTypeBool
+ %19 = OpConstantNull %mat2v2float
%uint_2 = OpConstant %uint 2
%uint_264 = OpConstant %uint 264
- %uint_1 = OpConstant %uint 1
%_arr_uint_uint_1 = OpTypeArray %uint %uint_1
- %18 = OpTypeFunction %v2float %_arr_uint_uint_1
+ %23 = OpTypeFunction %v2float %_arr_uint_uint_1
%int = OpTypeInt 32 1
- %25 = OpConstantNull %int
+ %29 = OpConstantNull %int
%_ptr_Workgroup_v2float = OpTypePointer Workgroup %v2float
- %35 = OpConstantComposite %_arr_uint_uint_1 %uint_1
- %36 = OpTypeFunction %void
+ %39 = OpConstantComposite %_arr_uint_uint_1 %uint_1
+ %40 = OpTypeFunction %void
%tint_zero_workgroup_memory = OpFunction %void None %9
%local_idx = OpFunctionParameter %uint
%13 = OpLabel
- OpStore %S %14
+ %15 = OpULessThan %bool %local_idx %uint_1
+ OpSelectionMerge %17 None
+ OpBranchConditional %15 %18 %17
+ %18 = OpLabel
+ OpStore %S %19
+ OpBranch %17
+ %17 = OpLabel
OpControlBarrier %uint_2 %uint_2 %uint_264
OpReturn
OpFunctionEnd
- %func_S_X = OpFunction %v2float None %18
+ %func_S_X = OpFunction %v2float None %23
%pointer = OpFunctionParameter %_arr_uint_uint_1
- %23 = OpLabel
- %26 = OpCompositeExtract %uint %pointer 0
- %28 = OpAccessChain %_ptr_Workgroup_v2float %S %26
- %29 = OpLoad %v2float %28
- OpReturnValue %29
+ %27 = OpLabel
+ %30 = OpCompositeExtract %uint %pointer 0
+ %32 = OpAccessChain %_ptr_Workgroup_v2float %S %30
+ %33 = OpLoad %v2float %32
+ OpReturnValue %33
OpFunctionEnd
%main_inner = OpFunction %void None %9
%local_invocation_index = OpFunctionParameter %uint
- %32 = OpLabel
- %33 = OpFunctionCall %void %tint_zero_workgroup_memory %local_invocation_index
- %34 = OpFunctionCall %v2float %func_S_X %35
+ %36 = OpLabel
+ %37 = OpFunctionCall %void %tint_zero_workgroup_memory %local_invocation_index
+ %38 = OpFunctionCall %v2float %func_S_X %39
OpReturn
OpFunctionEnd
- %main = OpFunction %void None %36
- %38 = OpLabel
- %40 = OpLoad %uint %local_invocation_index_1
- %39 = OpFunctionCall %void %main_inner %40
+ %main = OpFunction %void None %40
+ %42 = OpLabel
+ %44 = OpLoad %uint %local_invocation_index_1
+ %43 = OpFunctionCall %void %main_inner %44
OpReturn
OpFunctionEnd
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.dxc.hlsl b/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.dxc.hlsl
index 7214eab..5696a58 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.dxc.hlsl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.dxc.hlsl
@@ -1,7 +1,7 @@
groupshared float4 S;
void tint_zero_workgroup_memory(uint local_idx) {
- {
+ if ((local_idx < 1u)) {
S = (0.0f).xxxx;
}
GroupMemoryBarrierWithGroupSync();
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.fxc.hlsl b/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.fxc.hlsl
index 7214eab..5696a58 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.fxc.hlsl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.fxc.hlsl
@@ -1,7 +1,7 @@
groupshared float4 S;
void tint_zero_workgroup_memory(uint local_idx) {
- {
+ if ((local_idx < 1u)) {
S = (0.0f).xxxx;
}
GroupMemoryBarrierWithGroupSync();
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.glsl b/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.glsl
index 7dc1984..fa0f14c 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.glsl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.glsl
@@ -2,7 +2,7 @@
shared vec4 S;
void tint_zero_workgroup_memory(uint local_idx) {
- {
+ if ((local_idx < 1u)) {
S = vec4(0.0f);
}
barrier();
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.msl b/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.msl
index 989b8c1..64a47ae 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.msl
@@ -2,7 +2,7 @@
using namespace metal;
void tint_zero_workgroup_memory(uint local_idx, threadgroup float4* const tint_symbol_1) {
- {
+ if ((local_idx < 1u)) {
*(tint_symbol_1) = float4(0.0f);
}
threadgroup_barrier(mem_flags::mem_threadgroup);
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.spvasm b/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.spvasm
index 9458c04..337b173 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.spvasm
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 31
+; Bound: 36
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -25,33 +25,41 @@
%S = OpVariable %_ptr_Workgroup_v4float Workgroup
%void = OpTypeVoid
%8 = OpTypeFunction %void %uint
- %13 = OpConstantNull %v4float
+ %uint_1 = OpConstant %uint 1
+ %bool = OpTypeBool
+ %18 = OpConstantNull %v4float
%uint_2 = OpConstant %uint 2
%uint_264 = OpConstant %uint 264
- %17 = OpTypeFunction %v4float
- %26 = OpTypeFunction %void
+ %22 = OpTypeFunction %v4float
+ %31 = OpTypeFunction %void
%tint_zero_workgroup_memory = OpFunction %void None %8
%local_idx = OpFunctionParameter %uint
%12 = OpLabel
- OpStore %S %13
+ %14 = OpULessThan %bool %local_idx %uint_1
+ OpSelectionMerge %16 None
+ OpBranchConditional %14 %17 %16
+ %17 = OpLabel
+ OpStore %S %18
+ OpBranch %16
+ %16 = OpLabel
OpControlBarrier %uint_2 %uint_2 %uint_264
OpReturn
OpFunctionEnd
- %func_S = OpFunction %v4float None %17
- %19 = OpLabel
- %20 = OpLoad %v4float %S
- OpReturnValue %20
+ %func_S = OpFunction %v4float None %22
+ %24 = OpLabel
+ %25 = OpLoad %v4float %S
+ OpReturnValue %25
OpFunctionEnd
%main_inner = OpFunction %void None %8
%local_invocation_index = OpFunctionParameter %uint
- %23 = OpLabel
- %24 = OpFunctionCall %void %tint_zero_workgroup_memory %local_invocation_index
- %25 = OpFunctionCall %v4float %func_S
+ %28 = OpLabel
+ %29 = OpFunctionCall %void %tint_zero_workgroup_memory %local_invocation_index
+ %30 = OpFunctionCall %v4float %func_S
OpReturn
OpFunctionEnd
- %main = OpFunction %void None %26
- %28 = OpLabel
- %30 = OpLoad %uint %local_invocation_index_1
- %29 = OpFunctionCall %void %main_inner %30
+ %main = OpFunction %void None %31
+ %33 = OpLabel
+ %35 = OpLoad %uint %local_invocation_index_1
+ %34 = OpFunctionCall %void %main_inner %35
OpReturn
OpFunctionEnd
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.dxc.hlsl b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.dxc.hlsl
index c1280c2..7e48abb 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.dxc.hlsl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.dxc.hlsl
@@ -1,7 +1,7 @@
groupshared float2x4 S;
void tint_zero_workgroup_memory(uint local_idx) {
- {
+ if ((local_idx < 1u)) {
S = float2x4((0.0f).xxxx, (0.0f).xxxx);
}
GroupMemoryBarrierWithGroupSync();
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.fxc.hlsl b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.fxc.hlsl
index c1280c2..7e48abb 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.fxc.hlsl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.fxc.hlsl
@@ -1,7 +1,7 @@
groupshared float2x4 S;
void tint_zero_workgroup_memory(uint local_idx) {
- {
+ if ((local_idx < 1u)) {
S = float2x4((0.0f).xxxx, (0.0f).xxxx);
}
GroupMemoryBarrierWithGroupSync();
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.glsl b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.glsl
index 367dfff..d308c97 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.glsl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.glsl
@@ -2,7 +2,7 @@
shared mat2x4 S;
void tint_zero_workgroup_memory(uint local_idx) {
- {
+ if ((local_idx < 1u)) {
S = mat2x4(vec4(0.0f), vec4(0.0f));
}
barrier();
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.msl b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.msl
index f1c6242..1d5f314 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.msl
@@ -6,7 +6,7 @@
};
void tint_zero_workgroup_memory(uint local_idx, threadgroup float2x4* const tint_symbol_1) {
- {
+ if ((local_idx < 1u)) {
*(tint_symbol_1) = float2x4(float4(0.0f), float4(0.0f));
}
threadgroup_barrier(mem_flags::mem_threadgroup);
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.spvasm b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.spvasm
index 7185c1b..b59bfa3 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.spvasm
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 41
+; Bound: 45
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -28,42 +28,49 @@
%S = OpVariable %_ptr_Workgroup_mat2v4float Workgroup
%void = OpTypeVoid
%9 = OpTypeFunction %void %uint
- %14 = OpConstantNull %mat2v4float
+ %uint_1 = OpConstant %uint 1
+ %bool = OpTypeBool
+ %19 = OpConstantNull %mat2v4float
%uint_2 = OpConstant %uint 2
%uint_264 = OpConstant %uint 264
- %uint_1 = OpConstant %uint 1
%_arr_uint_uint_1 = OpTypeArray %uint %uint_1
- %18 = OpTypeFunction %v4float %_arr_uint_uint_1
+ %23 = OpTypeFunction %v4float %_arr_uint_uint_1
%int = OpTypeInt 32 1
- %25 = OpConstantNull %int
+ %29 = OpConstantNull %int
%_ptr_Workgroup_v4float = OpTypePointer Workgroup %v4float
- %35 = OpConstantComposite %_arr_uint_uint_1 %uint_1
- %36 = OpTypeFunction %void
+ %39 = OpConstantComposite %_arr_uint_uint_1 %uint_1
+ %40 = OpTypeFunction %void
%tint_zero_workgroup_memory = OpFunction %void None %9
%local_idx = OpFunctionParameter %uint
%13 = OpLabel
- OpStore %S %14
+ %15 = OpULessThan %bool %local_idx %uint_1
+ OpSelectionMerge %17 None
+ OpBranchConditional %15 %18 %17
+ %18 = OpLabel
+ OpStore %S %19
+ OpBranch %17
+ %17 = OpLabel
OpControlBarrier %uint_2 %uint_2 %uint_264
OpReturn
OpFunctionEnd
- %func_S_X = OpFunction %v4float None %18
+ %func_S_X = OpFunction %v4float None %23
%pointer = OpFunctionParameter %_arr_uint_uint_1
- %23 = OpLabel
- %26 = OpCompositeExtract %uint %pointer 0
- %28 = OpAccessChain %_ptr_Workgroup_v4float %S %26
- %29 = OpLoad %v4float %28
- OpReturnValue %29
+ %27 = OpLabel
+ %30 = OpCompositeExtract %uint %pointer 0
+ %32 = OpAccessChain %_ptr_Workgroup_v4float %S %30
+ %33 = OpLoad %v4float %32
+ OpReturnValue %33
OpFunctionEnd
%main_inner = OpFunction %void None %9
%local_invocation_index = OpFunctionParameter %uint
- %32 = OpLabel
- %33 = OpFunctionCall %void %tint_zero_workgroup_memory %local_invocation_index
- %34 = OpFunctionCall %v4float %func_S_X %35
+ %36 = OpLabel
+ %37 = OpFunctionCall %void %tint_zero_workgroup_memory %local_invocation_index
+ %38 = OpFunctionCall %v4float %func_S_X %39
OpReturn
OpFunctionEnd
- %main = OpFunction %void None %36
- %38 = OpLabel
- %40 = OpLoad %uint %local_invocation_index_1
- %39 = OpFunctionCall %void %main_inner %40
+ %main = OpFunction %void None %40
+ %42 = OpLabel
+ %44 = OpLoad %uint %local_invocation_index_1
+ %43 = OpFunctionCall %void %main_inner %44
OpReturn
OpFunctionEnd
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.dxc.hlsl b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.dxc.hlsl
index 92459e2..857ff46 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.dxc.hlsl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.dxc.hlsl
@@ -5,7 +5,7 @@
groupshared str S;
void tint_zero_workgroup_memory(uint local_idx) {
- {
+ if ((local_idx < 1u)) {
str tint_symbol_2 = (str)0;
S = tint_symbol_2;
}
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.fxc.hlsl b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.fxc.hlsl
index 92459e2..857ff46 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.fxc.hlsl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.fxc.hlsl
@@ -5,7 +5,7 @@
groupshared str S;
void tint_zero_workgroup_memory(uint local_idx) {
- {
+ if ((local_idx < 1u)) {
str tint_symbol_2 = (str)0;
S = tint_symbol_2;
}
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.glsl b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.glsl
index 2fc4367..8f0a5d4 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.glsl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.glsl
@@ -6,7 +6,7 @@
shared str S;
void tint_zero_workgroup_memory(uint local_idx) {
- {
+ if ((local_idx < 1u)) {
str tint_symbol_1 = str(vec4(0.0f));
S = tint_symbol_1;
}
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.msl b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.msl
index b20b54c..1ffdf08 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.msl
@@ -6,7 +6,7 @@
};
void tint_zero_workgroup_memory(uint local_idx, threadgroup str* const tint_symbol_2) {
- {
+ if ((local_idx < 1u)) {
str const tint_symbol_1 = str{};
*(tint_symbol_2) = tint_symbol_1;
}
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.spvasm b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.spvasm
index 516c63b..6b881a0 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.spvasm
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 35
+; Bound: 40
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -29,36 +29,44 @@
%S = OpVariable %_ptr_Workgroup_str Workgroup
%void = OpTypeVoid
%9 = OpTypeFunction %void %uint
- %14 = OpConstantNull %str
+ %uint_1 = OpConstant %uint 1
+ %bool = OpTypeBool
+ %19 = OpConstantNull %str
%uint_2 = OpConstant %uint 2
%uint_264 = OpConstant %uint 264
- %18 = OpTypeFunction %v4float
+ %23 = OpTypeFunction %v4float
%uint_0 = OpConstant %uint 0
%_ptr_Workgroup_v4float = OpTypePointer Workgroup %v4float
- %30 = OpTypeFunction %void
+ %35 = OpTypeFunction %void
%tint_zero_workgroup_memory = OpFunction %void None %9
%local_idx = OpFunctionParameter %uint
%13 = OpLabel
- OpStore %S %14
+ %15 = OpULessThan %bool %local_idx %uint_1
+ OpSelectionMerge %17 None
+ OpBranchConditional %15 %18 %17
+ %18 = OpLabel
+ OpStore %S %19
+ OpBranch %17
+ %17 = OpLabel
OpControlBarrier %uint_2 %uint_2 %uint_264
OpReturn
OpFunctionEnd
- %func_S_i = OpFunction %v4float None %18
- %20 = OpLabel
- %23 = OpAccessChain %_ptr_Workgroup_v4float %S %uint_0
- %24 = OpLoad %v4float %23
- OpReturnValue %24
+ %func_S_i = OpFunction %v4float None %23
+ %25 = OpLabel
+ %28 = OpAccessChain %_ptr_Workgroup_v4float %S %uint_0
+ %29 = OpLoad %v4float %28
+ OpReturnValue %29
OpFunctionEnd
%main_inner = OpFunction %void None %9
%local_invocation_index = OpFunctionParameter %uint
- %27 = OpLabel
- %28 = OpFunctionCall %void %tint_zero_workgroup_memory %local_invocation_index
- %29 = OpFunctionCall %v4float %func_S_i
+ %32 = OpLabel
+ %33 = OpFunctionCall %void %tint_zero_workgroup_memory %local_invocation_index
+ %34 = OpFunctionCall %v4float %func_S_i
OpReturn
OpFunctionEnd
- %main = OpFunction %void None %30
- %32 = OpLabel
- %34 = OpLoad %uint %local_invocation_index_1
- %33 = OpFunctionCall %void %main_inner %34
+ %main = OpFunction %void None %35
+ %37 = OpLabel
+ %39 = OpLoad %uint %local_invocation_index_1
+ %38 = OpFunctionCall %void %main_inner %39
OpReturn
OpFunctionEnd
