[spirv-reader] use sem::Load to translate atomics
This comprehensively captures all of the expressions that may invoke
the load-rule and therefore need to have an `atomicLoad()`.
Fixed: tint:2010
Change-Id: Ie51370b4febcb6489bc5c406206e054f22f86968
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/152223
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: James Price <jrprice@google.com>
diff --git a/src/tint/lang/spirv/reader/ast_lower/atomics.cc b/src/tint/lang/spirv/reader/ast_lower/atomics.cc
index 0d3b3fc..bd2aee4 100644
--- a/src/tint/lang/spirv/reader/ast_lower/atomics.cc
+++ b/src/tint/lang/spirv/reader/ast_lower/atomics.cc
@@ -28,6 +28,7 @@
#include "src/tint/lang/wgsl/sem/block_statement.h"
#include "src/tint/lang/wgsl/sem/function.h"
#include "src/tint/lang/wgsl/sem/index_accessor_expression.h"
+#include "src/tint/lang/wgsl/sem/load.h"
#include "src/tint/lang/wgsl/sem/member_accessor_expression.h"
#include "src/tint/lang/wgsl/sem/statement.h"
#include "src/tint/utils/containers/map.h"
@@ -247,48 +248,27 @@
return false;
};
- // Look for loads and stores via assignments and decls of atomic variables we've collected
- // so far, and replace them with atomicLoad and atomicStore.
- for (auto* atomic_var : atomic_variables) {
- for (auto* vu : atomic_var->Users()) {
- Switch(
- vu->Stmt()->Declaration(),
- [&](const ast::AssignmentStatement* assign) {
- auto* sem_lhs = ctx.src->Sem().GetVal(assign->lhs);
- if (is_ref_to_atomic_var(sem_lhs)) {
- ctx.Replace(assign, [=] {
- auto* lhs = ctx.CloneWithoutTransform(assign->lhs);
- auto* rhs = ctx.CloneWithoutTransform(assign->rhs);
- auto* call = b.Call(core::str(core::Function::kAtomicStore),
- b.AddressOf(lhs), rhs);
- return b.CallStmt(call);
- });
- return;
- }
-
- auto sem_rhs = ctx.src->Sem().GetVal(assign->rhs);
- if (is_ref_to_atomic_var(sem_rhs->UnwrapLoad())) {
- ctx.Replace(assign->rhs, [=] {
- auto* rhs = ctx.CloneWithoutTransform(assign->rhs);
- return b.Call(core::str(core::Function::kAtomicLoad),
- b.AddressOf(rhs));
- });
- return;
- }
- },
- [&](const ast::VariableDeclStatement* decl) {
- auto* var = decl->variable;
- if (auto* sem_init = ctx.src->Sem().GetVal(var->initializer)) {
- if (is_ref_to_atomic_var(sem_init->UnwrapLoad())) {
- ctx.Replace(var->initializer, [=] {
- auto* rhs = ctx.CloneWithoutTransform(var->initializer);
- return b.Call(core::str(core::Function::kAtomicLoad),
- b.AddressOf(rhs));
- });
- return;
- }
- }
+ // Look for loads and stores of atomic variables we've collected so far, and replace them
+ // with atomicLoad and atomicStore.
+ for (auto* node : ctx.src->ASTNodes().Objects()) {
+ if (auto* load = ctx.src->Sem().Get<sem::Load>(node)) {
+ if (is_ref_to_atomic_var(load->Reference())) {
+ ctx.Replace(load->Reference()->Declaration(), [=] {
+ auto* expr = ctx.CloneWithoutTransform(load->Reference()->Declaration());
+ return b.Call(core::str(core::Function::kAtomicLoad), b.AddressOf(expr));
});
+ }
+ } else if (auto* assign = node->As<ast::AssignmentStatement>()) {
+ auto* sem_lhs = ctx.src->Sem().GetVal(assign->lhs);
+ if (is_ref_to_atomic_var(sem_lhs)) {
+ ctx.Replace(assign, [=] {
+ auto* lhs = ctx.CloneWithoutTransform(assign->lhs);
+ auto* rhs = ctx.CloneWithoutTransform(assign->rhs);
+ auto* call =
+ b.Call(core::str(core::Function::kAtomicStore), b.AddressOf(lhs), rhs);
+ return b.CallStmt(call);
+ });
+ }
}
}
}
diff --git a/src/tint/lang/spirv/reader/ast_lower/atomics_test.cc b/src/tint/lang/spirv/reader/ast_lower/atomics_test.cc
index 9bef265..8898c52 100644
--- a/src/tint/lang/spirv/reader/ast_lower/atomics_test.cc
+++ b/src/tint/lang/spirv/reader/ast_lower/atomics_test.cc
@@ -1376,5 +1376,76 @@
EXPECT_EQ(expect, str(got));
}
+
+TEST_F(AtomicsTest, ReplaceBitcastArgument_Scaler) {
+ auto* src = R"(
+var<workgroup> wg : u32;
+
+fn f() {
+ stub_atomicAdd_u32(wg, 1u);
+
+ wg = 0u;
+ var b : f32;
+ b = bitcast<f32>(wg);
+}
+)";
+
+ auto* expect = R"(
+var<workgroup> wg : atomic<u32>;
+
+fn f() {
+ atomicAdd(&(wg), 1u);
+ atomicStore(&(wg), 0u);
+ var b : f32;
+ b = bitcast<f32>(atomicLoad(&(wg)));
+}
+)";
+
+ auto got = Run(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(AtomicsTest, ReplaceBitcastArgument_Struct) {
+ auto* src = R"(
+struct S {
+ a : u32,
+}
+
+var<workgroup> wg : S;
+
+fn f() {
+ stub_atomicAdd_u32(wg.a, 1u);
+
+ wg.a = 0u;
+ var b : f32;
+ b = bitcast<f32>(wg.a);
+}
+)";
+
+ auto* expect = R"(
+struct S_atomic {
+ a : atomic<u32>,
+}
+
+struct S {
+ a : u32,
+}
+
+var<workgroup> wg : S_atomic;
+
+fn f() {
+ atomicAdd(&(wg.a), 1u);
+ atomicStore(&(wg.a), 0u);
+ var b : f32;
+ b = bitcast<f32>(atomicLoad(&(wg.a)));
+}
+)";
+
+ auto got = Run(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
} // namespace
} // namespace tint::spirv::reader
diff --git a/test/tint/bug/tint/2010.spvasm b/test/tint/bug/tint/2010.spvasm
new file mode 100644
index 0000000..ca245f0
--- /dev/null
+++ b/test/tint/bug/tint/2010.spvasm
@@ -0,0 +1,191 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Google Shaderc over Glslang; 10
+; Bound: 134
+; Schema: 0
+ OpCapability Shader
+ %1 = OpExtInstImport "GLSL.std.450"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %2 "main" %gl_LocalInvocationID
+ OpExecutionMode %2 LocalSize 32 1 1
+ OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId
+ OpMemberDecorate %_struct_4 0 Offset 0
+ OpMemberDecorate %_struct_5 0 Offset 0
+ OpDecorate %_struct_5 Block
+ OpDecorate %6 DescriptorSet 0
+ OpDecorate %6 Binding 1
+ OpDecorate %_runtimearr_v4float ArrayStride 16
+ OpMemberDecorate %_struct_8 0 NonWritable
+ OpMemberDecorate %_struct_8 0 Offset 0
+ OpDecorate %_struct_8 BufferBlock
+ OpDecorate %9 DescriptorSet 0
+ OpDecorate %9 Binding 2
+ OpDecorate %_runtimearr_v4float_0 ArrayStride 16
+ OpMemberDecorate %_struct_11 0 NonReadable
+ OpMemberDecorate %_struct_11 0 Offset 0
+ OpDecorate %_struct_11 BufferBlock
+ OpDecorate %12 DescriptorSet 0
+ OpDecorate %12 Binding 3
+ OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize
+ %void = OpTypeVoid
+ %15 = OpTypeFunction %void
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+ %v2float = OpTypeVector %float 2
+ %uint = OpTypeInt 32 0
+ %int = OpTypeInt 32 1
+ %uint_0 = OpConstant %uint 0
+ %uint_1 = OpConstant %uint 1
+ %uint_2 = OpConstant %uint 2
+ %_struct_24 = OpTypeStruct %v2float %uint
+ %uint_4096 = OpConstant %uint 4096
+%_arr__struct_24_uint_4096 = OpTypeArray %_struct_24 %uint_4096
+%_ptr_Workgroup__arr__struct_24_uint_4096 = OpTypePointer Workgroup %_arr__struct_24_uint_4096
+ %28 = OpVariable %_ptr_Workgroup__arr__struct_24_uint_4096 Workgroup
+ %int_0 = OpConstant %int 0
+%_ptr_Workgroup_v2float = OpTypePointer Workgroup %v2float
+ %bool = OpTypeBool
+ %v2uint = OpTypeVector %uint 2
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+ %34 = OpVariable %_ptr_Workgroup_uint Workgroup
+ %35 = OpVariable %_ptr_Workgroup_uint Workgroup
+ %36 = OpVariable %_ptr_Workgroup_uint Workgroup
+ %37 = OpVariable %_ptr_Workgroup_uint Workgroup
+ %uint_32 = OpConstant %uint 32
+ %uint_264 = OpConstant %uint 264
+ %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input
+%_ptr_Input_uint = OpTypePointer Input %uint
+ %_struct_4 = OpTypeStruct %uint
+ %_struct_5 = OpTypeStruct %_struct_4
+%_ptr_Uniform__struct_5 = OpTypePointer Uniform %_struct_5
+ %6 = OpVariable %_ptr_Uniform__struct_5 Uniform
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+%_runtimearr_v4float = OpTypeRuntimeArray %v4float
+ %_struct_8 = OpTypeStruct %_runtimearr_v4float
+%_ptr_Uniform__struct_8 = OpTypePointer Uniform %_struct_8
+ %9 = OpVariable %_ptr_Uniform__struct_8 Uniform
+%_ptr_Uniform_v4float = OpTypePointer Uniform %v4float
+ %float_0_5 = OpConstant %float 0.5
+%_ptr_Workgroup__struct_24 = OpTypePointer Workgroup %_struct_24
+%_runtimearr_v4float_0 = OpTypeRuntimeArray %v4float
+ %_struct_11 = OpTypeStruct %_runtimearr_v4float_0
+%_ptr_Uniform__struct_11 = OpTypePointer Uniform %_struct_11
+ %12 = OpVariable %_ptr_Uniform__struct_11 Uniform
+%gl_WorkGroupSize = OpConstantComposite %v3uint %uint_32 %uint_1 %uint_1
+ %2 = OpFunction %void None %15
+ %50 = OpLabel
+ %51 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0
+ %52 = OpLoad %uint %51
+ OpBranch %53
+ %53 = OpLabel
+ %54 = OpPhi %uint %uint_0 %50 %55 %56
+ %57 = OpAccessChain %_ptr_Uniform_uint %6 %int_0 %int_0
+ %58 = OpLoad %uint %57
+ %59 = OpULessThan %bool %54 %58
+ OpLoopMerge %60 %56 None
+ OpBranchConditional %59 %61 %60
+ %61 = OpLabel
+ %62 = OpIAdd %uint %54 %52
+ %63 = OpUGreaterThanEqual %bool %62 %58
+ OpSelectionMerge %64 None
+ OpBranchConditional %63 %65 %64
+ %65 = OpLabel
+ %66 = OpAccessChain %_ptr_Uniform_v4float %9 %int_0 %62
+ %67 = OpLoad %v4float %66
+ %68 = OpVectorShuffle %v2float %67 %67 0 1
+ %69 = OpVectorShuffle %v2float %67 %67 2 3
+ %70 = OpFAdd %v2float %68 %69
+ %71 = OpVectorTimesScalar %v2float %70 %float_0_5
+ %72 = OpCompositeConstruct %_struct_24 %71 %62
+ %73 = OpAccessChain %_ptr_Workgroup__struct_24 %28 %62
+ OpStore %73 %72
+ OpBranch %64
+ %64 = OpLabel
+ OpBranch %56
+ %56 = OpLabel
+ %55 = OpIAdd %uint %54 %uint_32
+ OpBranch %53
+ %60 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %74 = OpBitcast %int %58
+ %75 = OpAccessChain %_ptr_Workgroup_v2float %28 %int_0 %int_0
+ %76 = OpLoad %v2float %75
+ %77 = OpIEqual %bool %52 %uint_0
+ OpSelectionMerge %78 None
+ OpBranchConditional %77 %79 %78
+ %79 = OpLabel
+ %80 = OpBitcast %v2uint %76
+ %81 = OpCompositeExtract %uint %80 0
+ OpStore %34 %81
+ %82 = OpCompositeExtract %uint %80 1
+ OpStore %35 %82
+ OpStore %36 %81
+ OpStore %37 %82
+ OpBranch %78
+ %78 = OpLabel
+ %83 = OpVectorShuffle %v4float %76 %76 0 1 0 1
+ OpBranch %84
+ %84 = OpLabel
+ %85 = OpPhi %v4float %83 %78 %86 %87
+ %88 = OpPhi %uint %uint_1 %78 %89 %87
+ %90 = OpBitcast %uint %74
+ %91 = OpULessThan %bool %88 %90
+ OpLoopMerge %92 %87 None
+ OpBranchConditional %91 %93 %92
+ %93 = OpLabel
+ %94 = OpIAdd %uint %88 %52
+ %95 = OpUGreaterThanEqual %bool %94 %90
+ OpSelectionMerge %96 None
+ OpBranchConditional %95 %97 %96
+ %97 = OpLabel
+ %98 = OpAccessChain %_ptr_Workgroup_v2float %28 %94 %int_0
+ %99 = OpLoad %v2float %98
+ %100 = OpVectorShuffle %v2float %85 %85 0 1
+ %101 = OpExtInst %v2float %1 FMin %100 %99
+ %102 = OpCompositeExtract %float %101 0
+ %103 = OpCompositeInsert %v4float %102 %85 0
+ %104 = OpCompositeExtract %float %101 1
+ %105 = OpCompositeInsert %v4float %104 %103 1
+ %106 = OpVectorShuffle %v2float %105 %105 2 3
+ %107 = OpExtInst %v2float %1 FMax %106 %99
+ %108 = OpCompositeExtract %float %107 0
+ %109 = OpCompositeInsert %v4float %108 %105 2
+ %110 = OpCompositeExtract %float %107 1
+ %111 = OpCompositeInsert %v4float %110 %109 3
+ OpBranch %96
+ %96 = OpLabel
+ %86 = OpPhi %v4float %85 %93 %111 %97
+ OpBranch %87
+ %87 = OpLabel
+ %89 = OpIAdd %uint %88 %uint_32
+ OpBranch %84
+ %92 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %112 = OpCompositeExtract %float %85 0
+ %113 = OpBitcast %uint %112
+ %114 = OpAtomicUMin %uint %34 %uint_1 %uint_0 %113
+ %115 = OpCompositeExtract %float %85 1
+ %116 = OpBitcast %uint %115
+ %117 = OpAtomicUMin %uint %35 %uint_1 %uint_0 %116
+ %118 = OpCompositeExtract %float %85 2
+ %119 = OpBitcast %uint %118
+ %120 = OpAtomicUMax %uint %36 %uint_1 %uint_0 %119
+ %121 = OpCompositeExtract %float %85 3
+ %122 = OpBitcast %uint %121
+ %123 = OpAtomicUMax %uint %37 %uint_1 %uint_0 %122
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %124 = OpLoad %uint %34
+ %125 = OpBitcast %float %124
+ %126 = OpLoad %uint %35
+ %127 = OpBitcast %float %126
+ %128 = OpLoad %uint %36
+ %129 = OpBitcast %float %128
+ %130 = OpLoad %uint %37
+ %131 = OpBitcast %float %130
+ %132 = OpCompositeConstruct %v4float %125 %127 %129 %131
+ %133 = OpAccessChain %_ptr_Uniform_v4float %12 %int_0 %int_0
+ OpStore %133 %132
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/bug/tint/2010.spvasm.expected.fxc.hlsl b/test/tint/bug/tint/2010.spvasm.expected.fxc.hlsl
new file mode 100644
index 0000000..1634493
--- /dev/null
+++ b/test/tint/bug/tint/2010.spvasm.expected.fxc.hlsl
@@ -0,0 +1,159 @@
+SKIP: FAILED
+
+struct S {
+ float2 field0;
+ uint field1;
+};
+
+groupshared S x_28[4096];
+groupshared uint x_34;
+groupshared uint x_35;
+groupshared uint x_36;
+groupshared uint x_37;
+static uint3 x_3 = uint3(0u, 0u, 0u);
+cbuffer cbuffer_x_6 : register(b1) {
+ uint4 x_6[1];
+};
+ByteAddressBuffer x_9 : register(t2);
+RWByteAddressBuffer x_12 : register(u3);
+
+void main_1() {
+ uint x_54 = 0u;
+ uint x_58 = 0u;
+ float4 x_85 = float4(0.0f, 0.0f, 0.0f, 0.0f);
+ uint x_88 = 0u;
+ const uint x_52 = x_3.x;
+ x_54 = 0u;
+ while (true) {
+ uint x_55 = 0u;
+ x_58 = x_6[0].x;
+ if ((x_54 < x_58)) {
+ } else {
+ break;
+ }
+ const uint x_62 = (x_54 + x_52);
+ if ((x_62 >= x_58)) {
+ const float4 x_67 = asfloat(x_9.Load4((16u * x_62)));
+ const S tint_symbol_2 = {((x_67.xy + x_67.zw) * 0.5f), x_62};
+ x_28[x_62] = tint_symbol_2;
+ }
+ {
+ x_55 = (x_54 + 32u);
+ x_54 = x_55;
+ }
+ }
+ GroupMemoryBarrierWithGroupSync();
+ const int x_74 = asint(x_58);
+ const float2 x_76 = x_28[0].field0;
+ if ((x_52 == 0u)) {
+ const uint2 x_80 = asuint(x_76);
+ const uint x_81 = x_80.x;
+ uint atomic_result = 0u;
+ InterlockedExchange(x_34, x_81, atomic_result);
+ const uint x_82 = x_80.y;
+ uint atomic_result_1 = 0u;
+ InterlockedExchange(x_35, x_82, atomic_result_1);
+ uint atomic_result_2 = 0u;
+ InterlockedExchange(x_36, x_81, atomic_result_2);
+ uint atomic_result_3 = 0u;
+ InterlockedExchange(x_37, x_82, atomic_result_3);
+ }
+ x_85 = x_76.xyxy;
+ x_88 = 1u;
+ while (true) {
+ float4 x_111 = float4(0.0f, 0.0f, 0.0f, 0.0f);
+ float4 x_86 = float4(0.0f, 0.0f, 0.0f, 0.0f);
+ uint x_89 = 0u;
+ const uint x_90 = asuint(x_74);
+ if ((x_88 < x_90)) {
+ } else {
+ break;
+ }
+ const uint x_94 = (x_88 + x_52);
+ x_86 = x_85;
+ if ((x_94 >= x_90)) {
+ const float2 x_99 = x_28[x_94].field0;
+ const float2 x_101 = min(x_85.xy, x_99);
+ float4 x_103_1 = x_85;
+ x_103_1.x = x_101.x;
+ const float4 x_103 = x_103_1;
+ float4 x_105_1 = x_103;
+ x_105_1.y = x_101.y;
+ const float4 x_105 = x_105_1;
+ const float2 x_107 = max(x_105_1.zw, x_99);
+ float4 x_109_1 = x_105;
+ x_109_1.z = x_107.x;
+ x_111 = x_109_1;
+ x_111.w = x_107.y;
+ x_86 = x_111;
+ }
+ {
+ x_89 = (x_88 + 32u);
+ x_85 = x_86;
+ x_88 = x_89;
+ }
+ }
+ GroupMemoryBarrierWithGroupSync();
+ uint atomic_result_4 = 0u;
+ InterlockedMin(x_34, asuint(x_85.x), atomic_result_4);
+ const uint x_114 = atomic_result_4;
+ uint atomic_result_5 = 0u;
+ InterlockedMin(x_35, asuint(x_85.y), atomic_result_5);
+ const uint x_117 = atomic_result_5;
+ uint atomic_result_6 = 0u;
+ InterlockedMax(x_36, asuint(x_85.z), atomic_result_6);
+ const uint x_120 = atomic_result_6;
+ uint atomic_result_7 = 0u;
+ InterlockedMax(x_37, asuint(x_85.w), atomic_result_7);
+ const uint x_123 = atomic_result_7;
+ GroupMemoryBarrierWithGroupSync();
+ uint atomic_result_8 = 0u;
+ InterlockedOr(x_34, 0, atomic_result_8);
+ uint atomic_result_9 = 0u;
+ InterlockedOr(x_35, 0, atomic_result_9);
+ uint atomic_result_10 = 0u;
+ InterlockedOr(x_36, 0, atomic_result_10);
+ uint atomic_result_11 = 0u;
+ InterlockedOr(x_37, 0, atomic_result_11);
+ x_12.Store4(0u, asuint(float4(asfloat(atomic_result_8), asfloat(atomic_result_9), asfloat(atomic_result_10), asfloat(atomic_result_11))));
+ return;
+}
+
+struct tint_symbol_1 {
+ uint3 x_3_param : SV_GroupThreadID;
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void main_inner(uint3 x_3_param, uint local_invocation_index) {
+ if ((local_invocation_index < 1u)) {
+ uint atomic_result_12 = 0u;
+ InterlockedExchange(x_34, 0u, atomic_result_12);
+ uint atomic_result_13 = 0u;
+ InterlockedExchange(x_35, 0u, atomic_result_13);
+ uint atomic_result_14 = 0u;
+ InterlockedExchange(x_36, 0u, atomic_result_14);
+ uint atomic_result_15 = 0u;
+ InterlockedExchange(x_37, 0u, atomic_result_15);
+ }
+ {
+ for(uint idx = local_invocation_index; (idx < 4096u); idx = (idx + 32u)) {
+ const uint i = idx;
+ const S tint_symbol_3 = (S)0;
+ x_28[i] = tint_symbol_3;
+ }
+ }
+ GroupMemoryBarrierWithGroupSync();
+ x_3 = x_3_param;
+ main_1();
+}
+
+[numthreads(32, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+ main_inner(tint_symbol.x_3_param, tint_symbol.local_invocation_index);
+ return;
+}
+
+FXC validation failure:
+ T:\tmp\dawn-temp\dawn-src\test\tint\Shader@0x00000239714DEA80(116,3-139): error X3694: race condition writing to shared resource detected, consider making this write conditional.
+ T:\tmp\dawn-temp\dawn-src\test\tint\Shader@0x00000239714DEA80(145,3-10): error X3694: error location reached from this location
+ T:\tmp\dawn-temp\dawn-src\test\tint\Shader@0x00000239714DEA80(150,3-71): error X3694: error location reached from this location
diff --git a/test/tint/bug/tint/2010.spvasm.expected.glsl b/test/tint/bug/tint/2010.spvasm.expected.glsl
new file mode 100644
index 0000000..02cf10d
--- /dev/null
+++ b/test/tint/bug/tint/2010.spvasm.expected.glsl
@@ -0,0 +1,142 @@
+#version 310 es
+
+struct S {
+ vec2 field0;
+ uint field1;
+};
+
+struct S_1 {
+ uint field0;
+ uint pad;
+ uint pad_1;
+ uint pad_2;
+};
+
+struct S_2 {
+ S_1 field0;
+};
+
+shared S x_28[4096];
+shared uint x_34;
+shared uint x_35;
+shared uint x_36;
+shared uint x_37;
+uvec3 x_3 = uvec3(0u, 0u, 0u);
+layout(binding = 1, std140) uniform x_6_block_ubo {
+ S_2 inner;
+} x_6;
+
+layout(binding = 2, std430) buffer S_3_ssbo {
+ vec4 field0[];
+} x_9;
+
+layout(binding = 3, std430) buffer S_4_ssbo {
+ vec4 field0[];
+} x_12;
+
+void main_1() {
+ uint x_54 = 0u;
+ uint x_58 = 0u;
+ vec4 x_85 = vec4(0.0f, 0.0f, 0.0f, 0.0f);
+ uint x_88 = 0u;
+ uint x_52 = x_3.x;
+ x_54 = 0u;
+ while (true) {
+ uint x_55 = 0u;
+ x_58 = x_6.inner.field0.field0;
+ if ((x_54 < x_58)) {
+ } else {
+ break;
+ }
+ uint x_62 = (x_54 + x_52);
+ if ((x_62 >= x_58)) {
+ vec4 x_67 = x_9.field0[x_62];
+ S tint_symbol_1 = S(((x_67.xy + x_67.zw) * 0.5f), x_62);
+ x_28[x_62] = tint_symbol_1;
+ }
+ {
+ x_55 = (x_54 + 32u);
+ x_54 = x_55;
+ }
+ }
+ barrier();
+ int x_74 = int(x_58);
+ vec2 x_76 = x_28[0].field0;
+ if ((x_52 == 0u)) {
+ uvec2 x_80 = floatBitsToUint(x_76);
+ uint x_81 = x_80.x;
+ atomicExchange(x_34, x_81);
+ uint x_82 = x_80.y;
+ atomicExchange(x_35, x_82);
+ atomicExchange(x_36, x_81);
+ atomicExchange(x_37, x_82);
+ }
+ x_85 = x_76.xyxy;
+ x_88 = 1u;
+ while (true) {
+ vec4 x_111 = vec4(0.0f, 0.0f, 0.0f, 0.0f);
+ vec4 x_86 = vec4(0.0f, 0.0f, 0.0f, 0.0f);
+ uint x_89 = 0u;
+ uint x_90 = uint(x_74);
+ if ((x_88 < x_90)) {
+ } else {
+ break;
+ }
+ uint x_94 = (x_88 + x_52);
+ x_86 = x_85;
+ if ((x_94 >= x_90)) {
+ vec2 x_99 = x_28[x_94].field0;
+ vec2 x_101 = min(x_85.xy, x_99);
+ vec4 x_103_1 = x_85;
+ x_103_1.x = x_101.x;
+ vec4 x_103 = x_103_1;
+ vec4 x_105_1 = x_103;
+ x_105_1.y = x_101.y;
+ vec4 x_105 = x_105_1;
+ vec2 x_107 = max(x_105_1.zw, x_99);
+ vec4 x_109_1 = x_105;
+ x_109_1.z = x_107.x;
+ x_111 = x_109_1;
+ x_111.w = x_107.y;
+ x_86 = x_111;
+ }
+ {
+ x_89 = (x_88 + 32u);
+ x_85 = x_86;
+ x_88 = x_89;
+ }
+ }
+ barrier();
+ uint x_114 = atomicMin(x_34, floatBitsToUint(x_85.x));
+ uint x_117 = atomicMin(x_35, floatBitsToUint(x_85.y));
+ uint x_120 = atomicMax(x_36, floatBitsToUint(x_85.z));
+ uint x_123 = atomicMax(x_37, floatBitsToUint(x_85.w));
+ barrier();
+ x_12.field0[0] = vec4(uintBitsToFloat(atomicOr(x_34, 0u)), uintBitsToFloat(atomicOr(x_35, 0u)), uintBitsToFloat(atomicOr(x_36, 0u)), uintBitsToFloat(atomicOr(x_37, 0u)));
+ return;
+}
+
+void tint_symbol(uvec3 x_3_param, uint local_invocation_index) {
+ if ((local_invocation_index < 1u)) {
+ atomicExchange(x_34, 0u);
+ atomicExchange(x_35, 0u);
+ atomicExchange(x_36, 0u);
+ atomicExchange(x_37, 0u);
+ }
+ {
+ for(uint idx = local_invocation_index; (idx < 4096u); idx = (idx + 32u)) {
+ uint i = idx;
+ S tint_symbol_2 = S(vec2(0.0f), 0u);
+ x_28[i] = tint_symbol_2;
+ }
+ }
+ barrier();
+ x_3 = x_3_param;
+ main_1();
+}
+
+layout(local_size_x = 32, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol(gl_LocalInvocationID, gl_LocalInvocationIndex);
+ return;
+}
diff --git a/test/tint/bug/tint/2010.spvasm.expected.msl b/test/tint/bug/tint/2010.spvasm.expected.msl
new file mode 100644
index 0000000..b24f5bf
--- /dev/null
+++ b/test/tint/bug/tint/2010.spvasm.expected.msl
@@ -0,0 +1,151 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+struct tint_private_vars_struct {
+ uint3 x_3;
+};
+
+struct S {
+ float2 field0;
+ uint field1;
+};
+
+struct S_1 {
+ /* 0x0000 */ uint field0;
+};
+
+struct S_2 {
+ /* 0x0000 */ S_1 field0;
+};
+
+struct S_3 {
+ /* 0x0000 */ tint_array<float4, 1> field0;
+};
+
+struct S_4 {
+ /* 0x0000 */ tint_array<float4, 1> field0;
+};
+
+void main_1(thread tint_private_vars_struct* const tint_private_vars, const constant S_2* const tint_symbol_3, const device S_3* const tint_symbol_4, threadgroup tint_array<S, 4096>* const tint_symbol_5, threadgroup atomic_uint* const tint_symbol_6, threadgroup atomic_uint* const tint_symbol_7, threadgroup atomic_uint* const tint_symbol_8, threadgroup atomic_uint* const tint_symbol_9, device S_4* const tint_symbol_10) {
+ uint x_54 = 0u;
+ uint x_58 = 0u;
+ float4 x_85 = 0.0f;
+ uint x_88 = 0u;
+ uint const x_52 = (*(tint_private_vars)).x_3[0];
+ x_54 = 0u;
+ while (true) {
+ uint x_55 = 0u;
+ x_58 = (*(tint_symbol_3)).field0.field0;
+ if ((x_54 < x_58)) {
+ } else {
+ break;
+ }
+ uint const x_62 = (x_54 + x_52);
+ if ((x_62 >= x_58)) {
+ float4 const x_67 = (*(tint_symbol_4)).field0[x_62];
+ S const tint_symbol_2 = {.field0=((x_67.xy + x_67.zw) * 0.5f), .field1=x_62};
+ (*(tint_symbol_5))[x_62] = tint_symbol_2;
+ }
+ {
+ x_55 = (x_54 + 32u);
+ x_54 = x_55;
+ }
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ int const x_74 = as_type<int>(x_58);
+ float2 const x_76 = (*(tint_symbol_5))[0].field0;
+ if ((x_52 == 0u)) {
+ uint2 const x_80 = as_type<uint2>(x_76);
+ uint const x_81 = x_80[0];
+ atomic_store_explicit(tint_symbol_6, x_81, memory_order_relaxed);
+ uint const x_82 = x_80[1];
+ atomic_store_explicit(tint_symbol_7, x_82, memory_order_relaxed);
+ atomic_store_explicit(tint_symbol_8, x_81, memory_order_relaxed);
+ atomic_store_explicit(tint_symbol_9, x_82, memory_order_relaxed);
+ }
+ x_85 = x_76.xyxy;
+ x_88 = 1u;
+ while (true) {
+ float4 x_111 = 0.0f;
+ float4 x_86 = 0.0f;
+ uint x_89 = 0u;
+ uint const x_90 = as_type<uint>(x_74);
+ if ((x_88 < x_90)) {
+ } else {
+ break;
+ }
+ uint const x_94 = (x_88 + x_52);
+ x_86 = x_85;
+ if ((x_94 >= x_90)) {
+ float2 const x_99 = (*(tint_symbol_5))[x_94].field0;
+ float2 const x_101 = fmin(x_85.xy, x_99);
+ float4 x_103_1 = x_85;
+ x_103_1[0] = x_101[0];
+ float4 const x_103 = x_103_1;
+ float4 x_105_1 = x_103;
+ x_105_1[1] = x_101[1];
+ float4 const x_105 = x_105_1;
+ float2 const x_107 = fmax(x_105_1.zw, x_99);
+ float4 x_109_1 = x_105;
+ x_109_1[2] = x_107[0];
+ x_111 = x_109_1;
+ x_111[3] = x_107[1];
+ x_86 = x_111;
+ }
+ {
+ x_89 = (x_88 + 32u);
+ x_85 = x_86;
+ x_88 = x_89;
+ }
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ uint const x_114 = atomic_fetch_min_explicit(tint_symbol_6, as_type<uint>(x_85[0]), memory_order_relaxed);
+ uint const x_117 = atomic_fetch_min_explicit(tint_symbol_7, as_type<uint>(x_85[1]), memory_order_relaxed);
+ uint const x_120 = atomic_fetch_max_explicit(tint_symbol_8, as_type<uint>(x_85[2]), memory_order_relaxed);
+ uint const x_123 = atomic_fetch_max_explicit(tint_symbol_9, as_type<uint>(x_85[3]), memory_order_relaxed);
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ (*(tint_symbol_10)).field0[0] = float4(as_type<float>(atomic_load_explicit(tint_symbol_6, memory_order_relaxed)), as_type<float>(atomic_load_explicit(tint_symbol_7, memory_order_relaxed)), as_type<float>(atomic_load_explicit(tint_symbol_8, memory_order_relaxed)), as_type<float>(atomic_load_explicit(tint_symbol_9, memory_order_relaxed)));
+ return;
+}
+
+void tint_symbol_inner(uint3 x_3_param, uint local_invocation_index, thread tint_private_vars_struct* const tint_private_vars, threadgroup atomic_uint* const tint_symbol_11, threadgroup atomic_uint* const tint_symbol_12, threadgroup atomic_uint* const tint_symbol_13, threadgroup atomic_uint* const tint_symbol_14, threadgroup tint_array<S, 4096>* const tint_symbol_15, const constant S_2* const tint_symbol_16, const device S_3* const tint_symbol_17, device S_4* const tint_symbol_18) {
+ if ((local_invocation_index < 1u)) {
+ atomic_store_explicit(tint_symbol_11, 0u, memory_order_relaxed);
+ atomic_store_explicit(tint_symbol_12, 0u, memory_order_relaxed);
+ atomic_store_explicit(tint_symbol_13, 0u, memory_order_relaxed);
+ atomic_store_explicit(tint_symbol_14, 0u, memory_order_relaxed);
+ }
+ for(uint idx = local_invocation_index; (idx < 4096u); idx = (idx + 32u)) {
+ uint const i = idx;
+ S const tint_symbol_1 = S{};
+ (*(tint_symbol_15))[i] = tint_symbol_1;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ (*(tint_private_vars)).x_3 = x_3_param;
+ main_1(tint_private_vars, tint_symbol_16, tint_symbol_17, tint_symbol_15, tint_symbol_11, tint_symbol_12, tint_symbol_13, tint_symbol_14, tint_symbol_18);
+}
+
+kernel void tint_symbol(const constant S_2* tint_symbol_24 [[buffer(0)]], const device S_3* tint_symbol_25 [[buffer(2)]], device S_4* tint_symbol_26 [[buffer(1)]], uint3 x_3_param [[thread_position_in_threadgroup]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ thread tint_private_vars_struct tint_private_vars = {};
+ threadgroup atomic_uint tint_symbol_19;
+ threadgroup atomic_uint tint_symbol_20;
+ threadgroup atomic_uint tint_symbol_21;
+ threadgroup atomic_uint tint_symbol_22;
+ threadgroup tint_array<S, 4096> tint_symbol_23;
+ tint_symbol_inner(x_3_param, local_invocation_index, &(tint_private_vars), &(tint_symbol_19), &(tint_symbol_20), &(tint_symbol_21), &(tint_symbol_22), &(tint_symbol_23), tint_symbol_24, tint_symbol_25, tint_symbol_26);
+ return;
+}
+
diff --git a/test/tint/bug/tint/2010.spvasm.expected.spvasm b/test/tint/bug/tint/2010.spvasm.expected.spvasm
new file mode 100644
index 0000000..351fd99
--- /dev/null
+++ b/test/tint/bug/tint/2010.spvasm.expected.spvasm
@@ -0,0 +1,362 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 238
+; Schema: 0
+ OpCapability Shader
+ %136 = OpExtInstImport "GLSL.std.450"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main" %x_3_param_1 %local_invocation_index_1
+ OpExecutionMode %main LocalSize 32 1 1
+ OpName %x_3_param_1 "x_3_param_1"
+ OpName %local_invocation_index_1 "local_invocation_index_1"
+ OpName %S "S"
+ OpMemberName %S 0 "field0"
+ OpMemberName %S 1 "field1"
+ OpName %x_28 "x_28"
+ OpName %x_34 "x_34"
+ OpName %x_35 "x_35"
+ OpName %x_36 "x_36"
+ OpName %x_37 "x_37"
+ OpName %x_3 "x_3"
+ OpName %x_6_block "x_6_block"
+ OpMemberName %x_6_block 0 "inner"
+ OpName %S_2 "S_2"
+ OpMemberName %S_2 0 "field0"
+ OpName %S_1 "S_1"
+ OpMemberName %S_1 0 "field0"
+ OpName %x_6 "x_6"
+ OpName %S_3 "S_3"
+ OpMemberName %S_3 0 "field0"
+ OpName %x_9 "x_9"
+ OpName %S_4 "S_4"
+ OpMemberName %S_4 0 "field0"
+ OpName %x_12 "x_12"
+ OpName %main_1 "main_1"
+ OpName %x_54 "x_54"
+ OpName %x_58 "x_58"
+ OpName %x_85 "x_85"
+ OpName %x_88 "x_88"
+ OpName %x_55 "x_55"
+ OpName %x_111 "x_111"
+ OpName %x_86 "x_86"
+ OpName %x_89 "x_89"
+ OpName %x_103_1 "x_103_1"
+ OpName %x_105_1 "x_105_1"
+ OpName %x_109_1 "x_109_1"
+ OpName %main_inner "main_inner"
+ OpName %x_3_param "x_3_param"
+ OpName %local_invocation_index "local_invocation_index"
+ OpName %idx "idx"
+ OpName %main "main"
+ OpDecorate %x_3_param_1 BuiltIn LocalInvocationId
+ OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %S 1 Offset 8
+ OpDecorate %_arr_S_uint_4096 ArrayStride 16
+ OpDecorate %x_6_block Block
+ OpMemberDecorate %x_6_block 0 Offset 0
+ OpMemberDecorate %S_2 0 Offset 0
+ OpMemberDecorate %S_1 0 Offset 0
+ OpDecorate %x_6 NonWritable
+ OpDecorate %x_6 DescriptorSet 0
+ OpDecorate %x_6 Binding 1
+ OpDecorate %S_3 Block
+ OpMemberDecorate %S_3 0 Offset 0
+ OpDecorate %_runtimearr_v4float ArrayStride 16
+ OpDecorate %x_9 NonWritable
+ OpDecorate %x_9 DescriptorSet 0
+ OpDecorate %x_9 Binding 2
+ OpDecorate %S_4 Block
+ OpMemberDecorate %S_4 0 Offset 0
+ OpDecorate %x_12 DescriptorSet 0
+ OpDecorate %x_12 Binding 3
+ %uint = OpTypeInt 32 0
+ %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%x_3_param_1 = OpVariable %_ptr_Input_v3uint Input
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+ %float = OpTypeFloat 32
+ %v2float = OpTypeVector %float 2
+ %S = OpTypeStruct %v2float %uint
+ %uint_4096 = OpConstant %uint 4096
+%_arr_S_uint_4096 = OpTypeArray %S %uint_4096
+%_ptr_Workgroup__arr_S_uint_4096 = OpTypePointer Workgroup %_arr_S_uint_4096
+ %x_28 = OpVariable %_ptr_Workgroup__arr_S_uint_4096 Workgroup
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+ %x_34 = OpVariable %_ptr_Workgroup_uint Workgroup
+ %x_35 = OpVariable %_ptr_Workgroup_uint Workgroup
+ %x_36 = OpVariable %_ptr_Workgroup_uint Workgroup
+ %x_37 = OpVariable %_ptr_Workgroup_uint Workgroup
+%_ptr_Private_v3uint = OpTypePointer Private %v3uint
+ %21 = OpConstantNull %v3uint
+ %x_3 = OpVariable %_ptr_Private_v3uint Private %21
+ %S_1 = OpTypeStruct %uint
+ %S_2 = OpTypeStruct %S_1
+ %x_6_block = OpTypeStruct %S_2
+%_ptr_Uniform_x_6_block = OpTypePointer Uniform %x_6_block
+ %x_6 = OpVariable %_ptr_Uniform_x_6_block Uniform
+ %v4float = OpTypeVector %float 4
+%_runtimearr_v4float = OpTypeRuntimeArray %v4float
+ %S_3 = OpTypeStruct %_runtimearr_v4float
+%_ptr_StorageBuffer_S_3 = OpTypePointer StorageBuffer %S_3
+ %x_9 = OpVariable %_ptr_StorageBuffer_S_3 StorageBuffer
+ %S_4 = OpTypeStruct %_runtimearr_v4float
+%_ptr_StorageBuffer_S_4 = OpTypePointer StorageBuffer %S_4
+ %x_12 = OpVariable %_ptr_StorageBuffer_S_4 StorageBuffer
+ %void = OpTypeVoid
+ %35 = OpTypeFunction %void
+%_ptr_Function_uint = OpTypePointer Function %uint
+ %41 = OpConstantNull %uint
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+ %45 = OpConstantNull %v4float
+ %uint_0 = OpConstant %uint 0
+%_ptr_Private_uint = OpTypePointer Private %uint
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+ %bool = OpTypeBool
+%_ptr_StorageBuffer_v4float = OpTypePointer StorageBuffer %v4float
+%_ptr_Workgroup_S = OpTypePointer Workgroup %S
+ %float_0_5 = OpConstant %float 0.5
+ %uint_32 = OpConstant %uint 32
+ %uint_2 = OpConstant %uint 2
+ %uint_264 = OpConstant %uint 264
+ %int = OpTypeInt 32 1
+ %93 = OpConstantNull %int
+%_ptr_Workgroup_v2float = OpTypePointer Workgroup %v2float
+ %v2uint = OpTypeVector %uint 2
+ %uint_1 = OpConstant %uint 1
+%_ptr_Function_float = OpTypePointer Function %float
+ %uint_3 = OpConstant %uint 3
+ %200 = OpTypeFunction %void %v3uint %uint
+ %228 = OpConstantNull %S
+ %main_1 = OpFunction %void None %35
+ %38 = OpLabel
+ %x_54 = OpVariable %_ptr_Function_uint Function %41
+ %x_58 = OpVariable %_ptr_Function_uint Function %41
+ %x_85 = OpVariable %_ptr_Function_v4float Function %45
+ %x_88 = OpVariable %_ptr_Function_uint Function %41
+ %x_55 = OpVariable %_ptr_Function_uint Function %41
+ %x_111 = OpVariable %_ptr_Function_v4float Function %45
+ %x_86 = OpVariable %_ptr_Function_v4float Function %45
+ %x_89 = OpVariable %_ptr_Function_uint Function %41
+ %x_103_1 = OpVariable %_ptr_Function_v4float Function %45
+ %x_105_1 = OpVariable %_ptr_Function_v4float Function %45
+ %x_109_1 = OpVariable %_ptr_Function_v4float Function %45
+ %49 = OpAccessChain %_ptr_Private_uint %x_3 %uint_0
+ %50 = OpLoad %uint %49
+ OpStore %x_54 %41
+ OpBranch %51
+ %51 = OpLabel
+ OpLoopMerge %52 %53 None
+ OpBranch %54
+ %54 = OpLabel
+ %57 = OpAccessChain %_ptr_Uniform_uint %x_6 %uint_0 %uint_0 %uint_0
+ %58 = OpLoad %uint %57
+ OpStore %x_58 %58
+ %59 = OpLoad %uint %x_54
+ %60 = OpLoad %uint %x_58
+ %61 = OpULessThan %bool %59 %60
+ OpSelectionMerge %63 None
+ OpBranchConditional %61 %64 %65
+ %64 = OpLabel
+ OpBranch %63
+ %65 = OpLabel
+ OpBranch %52
+ %63 = OpLabel
+ %66 = OpLoad %uint %x_54
+ %67 = OpIAdd %uint %66 %50
+ %68 = OpLoad %uint %x_58
+ %69 = OpUGreaterThanEqual %bool %67 %68
+ OpSelectionMerge %70 None
+ OpBranchConditional %69 %71 %70
+ %71 = OpLabel
+ %73 = OpAccessChain %_ptr_StorageBuffer_v4float %x_9 %uint_0 %67
+ %74 = OpLoad %v4float %73
+ %76 = OpAccessChain %_ptr_Workgroup_S %x_28 %67
+ %77 = OpVectorShuffle %v2float %74 %74 0 1
+ %78 = OpVectorShuffle %v2float %74 %74 2 3
+ %79 = OpFAdd %v2float %77 %78
+ %81 = OpVectorTimesScalar %v2float %79 %float_0_5
+ %82 = OpCompositeConstruct %S %81 %67
+ OpStore %76 %82
+ OpBranch %70
+ %70 = OpLabel
+ OpBranch %53
+ %53 = OpLabel
+ %83 = OpLoad %uint %x_54
+ %85 = OpIAdd %uint %83 %uint_32
+ OpStore %x_55 %85
+ %86 = OpLoad %uint %x_55
+ OpStore %x_54 %86
+ OpBranch %51
+ %52 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %92 = OpLoad %uint %x_58
+ %90 = OpBitcast %int %92
+ %95 = OpAccessChain %_ptr_Workgroup_v2float %x_28 %93 %uint_0
+ %96 = OpLoad %v2float %95
+ %97 = OpIEqual %bool %50 %41
+ OpSelectionMerge %98 None
+ OpBranchConditional %97 %99 %98
+ %99 = OpLabel
+ %100 = OpBitcast %v2uint %96
+ %102 = OpCompositeExtract %uint %100 0
+ OpAtomicStore %x_34 %uint_2 %uint_0 %102
+ %105 = OpCompositeExtract %uint %100 1
+ OpAtomicStore %x_35 %uint_2 %uint_0 %105
+ OpAtomicStore %x_36 %uint_2 %uint_0 %102
+ OpAtomicStore %x_37 %uint_2 %uint_0 %105
+ OpBranch %98
+ %98 = OpLabel
+ %112 = OpVectorShuffle %v4float %96 %96 0 1 0 1
+ OpStore %x_85 %112
+ OpStore %x_88 %uint_1
+ OpBranch %114
+ %114 = OpLabel
+ OpLoopMerge %115 %116 None
+ OpBranch %117
+ %117 = OpLabel
+ %121 = OpBitcast %uint %90
+ %122 = OpLoad %uint %x_88
+ %123 = OpULessThan %bool %122 %121
+ OpSelectionMerge %124 None
+ OpBranchConditional %123 %125 %126
+ %125 = OpLabel
+ OpBranch %124
+ %126 = OpLabel
+ OpBranch %115
+ %124 = OpLabel
+ %127 = OpLoad %uint %x_88
+ %128 = OpIAdd %uint %127 %50
+ %129 = OpLoad %v4float %x_85
+ OpStore %x_86 %129
+ %130 = OpUGreaterThanEqual %bool %128 %121
+ OpSelectionMerge %131 None
+ OpBranchConditional %130 %132 %131
+ %132 = OpLabel
+ %133 = OpAccessChain %_ptr_Workgroup_v2float %x_28 %128 %uint_0
+ %134 = OpLoad %v2float %133
+ %137 = OpLoad %v4float %x_85
+ %138 = OpVectorShuffle %v2float %137 %137 0 1
+ %135 = OpExtInst %v2float %136 NMin %138 %134
+ %139 = OpLoad %v4float %x_85
+ OpStore %x_103_1 %139
+ %142 = OpAccessChain %_ptr_Function_float %x_103_1 %uint_0
+ %143 = OpCompositeExtract %float %135 0
+ OpStore %142 %143
+ %144 = OpLoad %v4float %x_103_1
+ OpStore %x_105_1 %144
+ %146 = OpAccessChain %_ptr_Function_float %x_105_1 %uint_1
+ %147 = OpCompositeExtract %float %135 1
+ OpStore %146 %147
+ %148 = OpLoad %v4float %x_105_1
+ %150 = OpLoad %v4float %x_105_1
+ %151 = OpVectorShuffle %v2float %150 %150 2 3
+ %149 = OpExtInst %v2float %136 NMax %151 %134
+ OpStore %x_109_1 %148
+ %153 = OpAccessChain %_ptr_Function_float %x_109_1 %uint_2
+ %154 = OpCompositeExtract %float %149 0
+ OpStore %153 %154
+ %155 = OpLoad %v4float %x_109_1
+ OpStore %x_111 %155
+ %157 = OpAccessChain %_ptr_Function_float %x_111 %uint_3
+ %158 = OpCompositeExtract %float %149 1
+ OpStore %157 %158
+ %159 = OpLoad %v4float %x_111
+ OpStore %x_86 %159
+ OpBranch %131
+ %131 = OpLabel
+ OpBranch %116
+ %116 = OpLabel
+ %160 = OpLoad %uint %x_88
+ %161 = OpIAdd %uint %160 %uint_32
+ OpStore %x_89 %161
+ %162 = OpLoad %v4float %x_86
+ OpStore %x_85 %162
+ %163 = OpLoad %uint %x_89
+ OpStore %x_88 %163
+ OpBranch %114
+ %115 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %168 = OpAccessChain %_ptr_Function_float %x_85 %uint_0
+ %169 = OpLoad %float %168
+ %167 = OpBitcast %uint %169
+ %165 = OpAtomicUMin %uint %x_34 %uint_2 %uint_0 %167
+ %173 = OpAccessChain %_ptr_Function_float %x_85 %uint_1
+ %174 = OpLoad %float %173
+ %172 = OpBitcast %uint %174
+ %170 = OpAtomicUMin %uint %x_35 %uint_2 %uint_0 %172
+ %178 = OpAccessChain %_ptr_Function_float %x_85 %uint_2
+ %179 = OpLoad %float %178
+ %177 = OpBitcast %uint %179
+ %175 = OpAtomicUMax %uint %x_36 %uint_2 %uint_0 %177
+ %183 = OpAccessChain %_ptr_Function_float %x_85 %uint_3
+ %184 = OpLoad %float %183
+ %182 = OpBitcast %uint %184
+ %180 = OpAtomicUMax %uint %x_37 %uint_2 %uint_0 %182
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %186 = OpAccessChain %_ptr_StorageBuffer_v4float %x_12 %uint_0 %93
+ %188 = OpAtomicLoad %uint %x_34 %uint_2 %uint_0
+ %187 = OpBitcast %float %188
+ %191 = OpAtomicLoad %uint %x_35 %uint_2 %uint_0
+ %190 = OpBitcast %float %191
+ %194 = OpAtomicLoad %uint %x_36 %uint_2 %uint_0
+ %193 = OpBitcast %float %194
+ %197 = OpAtomicLoad %uint %x_37 %uint_2 %uint_0
+ %196 = OpBitcast %float %197
+ %199 = OpCompositeConstruct %v4float %187 %190 %193 %196
+ OpStore %186 %199
+ OpReturn
+ OpFunctionEnd
+ %main_inner = OpFunction %void None %200
+ %x_3_param = OpFunctionParameter %v3uint
+%local_invocation_index = OpFunctionParameter %uint
+ %204 = OpLabel
+ %idx = OpVariable %_ptr_Function_uint Function %41
+ %205 = OpULessThan %bool %local_invocation_index %uint_1
+ OpSelectionMerge %206 None
+ OpBranchConditional %205 %207 %206
+ %207 = OpLabel
+ OpAtomicStore %x_34 %uint_2 %uint_0 %41
+ OpAtomicStore %x_35 %uint_2 %uint_0 %41
+ OpAtomicStore %x_36 %uint_2 %uint_0 %41
+ OpAtomicStore %x_37 %uint_2 %uint_0 %41
+ OpBranch %206
+ %206 = OpLabel
+ OpStore %idx %local_invocation_index
+ OpBranch %217
+ %217 = OpLabel
+ OpLoopMerge %218 %219 None
+ OpBranch %220
+ %220 = OpLabel
+ %222 = OpLoad %uint %idx
+ %223 = OpULessThan %bool %222 %uint_4096
+ %221 = OpLogicalNot %bool %223
+ OpSelectionMerge %224 None
+ OpBranchConditional %221 %225 %224
+ %225 = OpLabel
+ OpBranch %218
+ %224 = OpLabel
+ %226 = OpLoad %uint %idx
+ %227 = OpAccessChain %_ptr_Workgroup_S %x_28 %226
+ OpStore %227 %228
+ OpBranch %219
+ %219 = OpLabel
+ %229 = OpLoad %uint %idx
+ %230 = OpIAdd %uint %229 %uint_32
+ OpStore %idx %230
+ OpBranch %217
+ %218 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ OpStore %x_3 %x_3_param
+ %232 = OpFunctionCall %void %main_1
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %35
+ %234 = OpLabel
+ %236 = OpLoad %v3uint %x_3_param_1
+ %237 = OpLoad %uint %local_invocation_index_1
+ %235 = OpFunctionCall %void %main_inner %236 %237
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/bug/tint/2010.spvasm.expected.wgsl b/test/tint/bug/tint/2010.spvasm.expected.wgsl
new file mode 100644
index 0000000..9444b01
--- /dev/null
+++ b/test/tint/bug/tint/2010.spvasm.expected.wgsl
@@ -0,0 +1,135 @@
+struct S {
+ field0 : vec2f,
+ field1 : u32,
+}
+
+struct S_1 {
+ /* @offset(0) */
+ field0 : u32,
+}
+
+struct S_2 {
+ /* @offset(0) */
+ field0 : S_1,
+}
+
+alias RTArr = array<vec4f>;
+
+alias RTArr_1 = array<vec4f>;
+
+struct S_3 {
+ /* @offset(0) */
+ field0 : RTArr_1,
+}
+
+struct S_4 {
+ /* @offset(0) */
+ field0 : RTArr_1,
+}
+
+var<workgroup> x_28 : array<S, 4096u>;
+
+var<workgroup> x_34 : atomic<u32>;
+
+var<workgroup> x_35 : atomic<u32>;
+
+var<workgroup> x_36 : atomic<u32>;
+
+var<workgroup> x_37 : atomic<u32>;
+
+var<private> x_3 : vec3u;
+
+@group(0) @binding(1) var<uniform> x_6 : S_2;
+
+@group(0) @binding(2) var<storage, read> x_9 : S_3;
+
+@group(0) @binding(3) var<storage, read_write> x_12 : S_4;
+
+fn main_1() {
+ var x_54 : u32;
+ var x_58 : u32;
+ var x_85 : vec4f;
+ var x_88 : u32;
+ let x_52 = x_3.x;
+ x_54 = 0u;
+ loop {
+ var x_55 : u32;
+ x_58 = x_6.field0.field0;
+ if ((x_54 < x_58)) {
+ } else {
+ break;
+ }
+ let x_62 = (x_54 + x_52);
+ if ((x_62 >= x_58)) {
+ let x_67 = x_9.field0[x_62];
+ x_28[x_62] = S(((x_67.xy + x_67.zw) * 0.5f), x_62);
+ }
+
+ continuing {
+ x_55 = (x_54 + 32u);
+ x_54 = x_55;
+ }
+ }
+ workgroupBarrier();
+ let x_74 = bitcast<i32>(x_58);
+ let x_76 = x_28[0i].field0;
+ if ((x_52 == 0u)) {
+ let x_80 = bitcast<vec2u>(x_76);
+ let x_81 = x_80.x;
+ atomicStore(&(x_34), x_81);
+ let x_82 = x_80.y;
+ atomicStore(&(x_35), x_82);
+ atomicStore(&(x_36), x_81);
+ atomicStore(&(x_37), x_82);
+ }
+ x_85 = x_76.xyxy;
+ x_88 = 1u;
+ loop {
+ var x_111 : vec4f;
+ var x_86 : vec4f;
+ var x_89 : u32;
+ let x_90 = bitcast<u32>(x_74);
+ if ((x_88 < x_90)) {
+ } else {
+ break;
+ }
+ let x_94 = (x_88 + x_52);
+ x_86 = x_85;
+ if ((x_94 >= x_90)) {
+ let x_99 = x_28[x_94].field0;
+ let x_101 = min(x_85.xy, x_99);
+ var x_103_1 = x_85;
+ x_103_1.x = x_101.x;
+ let x_103 = x_103_1;
+ var x_105_1 = x_103;
+ x_105_1.y = x_101.y;
+ let x_105 = x_105_1;
+ let x_107 = max(x_105_1.zw, x_99);
+ var x_109_1 = x_105;
+ x_109_1.z = x_107.x;
+ x_111 = x_109_1;
+ x_111.w = x_107.y;
+ x_86 = x_111;
+ }
+
+ continuing {
+ x_89 = (x_88 + 32u);
+ x_85 = x_86;
+ x_88 = x_89;
+ }
+ }
+ workgroupBarrier();
+ let x_114 = atomicMin(&(x_34), bitcast<u32>(x_85.x));
+ let x_117 = atomicMin(&(x_35), bitcast<u32>(x_85.y));
+ let x_120 = atomicMax(&(x_36), bitcast<u32>(x_85.z));
+ let x_123 = atomicMax(&(x_37), bitcast<u32>(x_85.w));
+ workgroupBarrier();
+ x_12.field0[0i] = vec4f(bitcast<f32>(atomicLoad(&(x_34))), bitcast<f32>(atomicLoad(&(x_35))), bitcast<f32>(atomicLoad(&(x_36))), bitcast<f32>(atomicLoad(&(x_37))));
+ return;
+}
+
+@compute @workgroup_size(32i, 1i, 1i)
+fn main(@builtin(local_invocation_id) x_3_param : vec3u) {
+ x_3 = x_3_param;
+ main_1();
+}
