tint/transform: Implement div / mod polyfill

Prevents UB for divide-by-zero and integer overflow when dividing

Fixed: tint:1349
Change-Id: Ieef66d27d7aec3011628ced076b2bccc7770a8af
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/108925
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>

test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.dxc.hlsl

@@ -1,3 +1,11 @@
+uint tint_div(uint lhs, uint rhs) {
+  return (lhs / ((rhs == 0u) ? 1u : rhs));
+}
+
+uint tint_mod(uint lhs, uint rhs) {
+  return (lhs % ((rhs == 0u) ? 1u : rhs));
+}
+
 static uint local_invocation_index_1 = 0u;
 groupshared uint wg[3][2][1];
 
@@ -12,8 +20,11 @@ void compute_main_inner(uint local_invocation_index) {
     const uint x_31 = idx;
     const uint x_33 = idx;
     const uint x_35 = idx;
+    const uint tint_symbol_2 = tint_div(x_31, 2u);
+    const uint tint_symbol_3 = tint_mod(x_33, 2u);
+    const uint tint_symbol_4 = tint_mod(x_35, 1u);
     uint atomic_result = 0u;
-    InterlockedExchange(wg[(x_31 / 2u)][(x_33 % 2u)][(x_35 % 1u)], 0u, atomic_result);
+    InterlockedExchange(wg[tint_symbol_2][tint_symbol_3][tint_symbol_4], 0u, atomic_result);
     {
       const uint x_42 = idx;
       idx = (x_42 + 1u);

test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.fxc.hlsl

@@ -1,3 +1,11 @@
+uint tint_div(uint lhs, uint rhs) {
+  return (lhs / ((rhs == 0u) ? 1u : rhs));
+}
+
+uint tint_mod(uint lhs, uint rhs) {
+  return (lhs % ((rhs == 0u) ? 1u : rhs));
+}
+
 static uint local_invocation_index_1 = 0u;
 groupshared uint wg[3][2][1];
 
@@ -12,8 +20,11 @@ void compute_main_inner(uint local_invocation_index) {
     const uint x_31 = idx;
     const uint x_33 = idx;
     const uint x_35 = idx;
+    const uint tint_symbol_2 = tint_div(x_31, 2u);
+    const uint tint_symbol_3 = tint_mod(x_33, 2u);
+    const uint tint_symbol_4 = tint_mod(x_35, 1u);
     uint atomic_result = 0u;
-    InterlockedExchange(wg[(x_31 / 2u)][(x_33 % 2u)][(x_35 % 1u)], 0u, atomic_result);
+    InterlockedExchange(wg[tint_symbol_2][tint_symbol_3][tint_symbol_4], 0u, atomic_result);
     {
       const uint x_42 = idx;
       idx = (x_42 + 1u);

test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.glsl

@@ -1,5 +1,13 @@
 #version 310 es
 
+uint tint_div(uint lhs, uint rhs) {
+  return (lhs / ((rhs == 0u) ? 1u : rhs));
+}
+
+uint tint_mod(uint lhs, uint rhs) {
+  return (lhs % ((rhs == 0u) ? 1u : rhs));
+}
+
 uint local_invocation_index_1 = 0u;
 shared uint wg[3][2][1];
 void compute_main_inner(uint local_invocation_index) {
@@ -13,7 +21,10 @@ void compute_main_inner(uint local_invocation_index) {
     uint x_31 = idx;
     uint x_33 = idx;
     uint x_35 = idx;
-    atomicExchange(wg[(x_31 / 2u)][(x_33 % 2u)][(x_35 % 1u)], 0u);
+    uint tint_symbol = tint_div(x_31, 2u);
+    uint tint_symbol_1 = tint_mod(x_33, 2u);
+    uint tint_symbol_2 = tint_mod(x_35, 1u);
+    atomicExchange(wg[tint_symbol][tint_symbol_1][tint_symbol_2], 0u);
     {
       uint x_42 = idx;
       idx = (x_42 + 1u);

test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.msl

@@ -14,7 +14,15 @@ struct tint_array {
     T elements[N];
 };
 
-void compute_main_inner(uint local_invocation_index, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol) {
+uint tint_div(uint lhs, uint rhs) {
+  return (lhs / select(rhs, 1u, (rhs == 0u)));
+}
+
+uint tint_mod(uint lhs, uint rhs) {
+  return (lhs % select(rhs, 1u, (rhs == 0u)));
+}
+
+void compute_main_inner(uint local_invocation_index, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_3) {
   uint idx = 0u;
   idx = local_invocation_index;
   while (true) {
@@ -25,39 +33,42 @@ void compute_main_inner(uint local_invocation_index, threadgroup tint_array<tint
     uint const x_31 = idx;
     uint const x_33 = idx;
     uint const x_35 = idx;
-    atomic_store_explicit(&((*(tint_symbol))[(x_31 / 2u)][(x_33 % 2u)][(x_35 % 1u)]), 0u, memory_order_relaxed);
+    uint const tint_symbol = tint_div(x_31, 2u);
+    uint const tint_symbol_1 = tint_mod(x_33, 2u);
+    uint const tint_symbol_2 = tint_mod(x_35, 1u);
+    atomic_store_explicit(&((*(tint_symbol_3))[tint_symbol][tint_symbol_1][tint_symbol_2]), 0u, memory_order_relaxed);
     {
       uint const x_42 = idx;
       idx = (x_42 + 1u);
     }
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  atomic_store_explicit(&((*(tint_symbol))[2][1][0]), 1u, memory_order_relaxed);
+  atomic_store_explicit(&((*(tint_symbol_3))[2][1][0]), 1u, memory_order_relaxed);
   return;
 }
 
-void compute_main_1(thread uint* const tint_symbol_1, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_2) {
-  uint const x_57 = *(tint_symbol_1);
-  compute_main_inner(x_57, tint_symbol_2);
+void compute_main_1(thread uint* const tint_symbol_4, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_5) {
+  uint const x_57 = *(tint_symbol_4);
+  compute_main_inner(x_57, tint_symbol_5);
   return;
 }
 
-void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_3, thread uint* const tint_symbol_4) {
+void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_6, thread uint* const tint_symbol_7) {
   for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 6u); idx_1 = (idx_1 + 1u)) {
     uint const i = (idx_1 / 2u);
     uint const i_1 = (idx_1 % 2u);
     uint const i_2 = (idx_1 % 1u);
-    atomic_store_explicit(&((*(tint_symbol_3))[i][i_1][i_2]), 0u, memory_order_relaxed);
+    atomic_store_explicit(&((*(tint_symbol_6))[i][i_1][i_2]), 0u, memory_order_relaxed);
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  *(tint_symbol_4) = local_invocation_index_1_param;
-  compute_main_1(tint_symbol_4, tint_symbol_3);
+  *(tint_symbol_7) = local_invocation_index_1_param;
+  compute_main_1(tint_symbol_7, tint_symbol_6);
 }
 
 kernel void compute_main(uint local_invocation_index_1_param [[thread_index_in_threadgroup]]) {
-  threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol_5;
-  thread uint tint_symbol_6 = 0u;
-  compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_5), &(tint_symbol_6));
+  threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol_8;
+  thread uint tint_symbol_9 = 0u;
+  compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_8), &(tint_symbol_9));
   return;
 }
 

test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 90
+; Bound: 105
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
@@ -10,6 +10,12 @@
                OpName %local_invocation_index_1_param_1 "local_invocation_index_1_param_1"
                OpName %local_invocation_index_1 "local_invocation_index_1"
                OpName %wg "wg"
+               OpName %tint_div "tint_div"
+               OpName %lhs "lhs"
+               OpName %rhs "rhs"
+               OpName %tint_mod "tint_mod"
+               OpName %lhs_0 "lhs"
+               OpName %rhs_0 "rhs"
                OpName %compute_main_inner "compute_main_inner"
                OpName %local_invocation_index "local_invocation_index"
                OpName %idx "idx"
@@ -36,105 +42,124 @@
 %_arr__arr__arr_uint_uint_1_uint_2_uint_3 = OpTypeArray %_arr__arr_uint_uint_1_uint_2 %uint_3
 %_ptr_Workgroup__arr__arr__arr_uint_uint_1_uint_2_uint_3 = OpTypePointer Workgroup %_arr__arr__arr_uint_uint_1_uint_2_uint_3
          %wg = OpVariable %_ptr_Workgroup__arr__arr__arr_uint_uint_1_uint_2_uint_3 Workgroup
+         %15 = OpTypeFunction %uint %uint %uint
+       %bool = OpTypeBool
        %void = OpTypeVoid
-         %15 = OpTypeFunction %void %uint
+         %31 = OpTypeFunction %void %uint
 %_ptr_Function_uint = OpTypePointer Function %uint
      %uint_6 = OpConstant %uint 6
-       %bool = OpTypeBool
      %uint_0 = OpConstant %uint 0
 %_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
    %uint_264 = OpConstant %uint 264
         %int = OpTypeInt 32 1
       %int_2 = OpConstant %int 2
       %int_1 = OpConstant %int 1
-         %53 = OpConstantNull %int
-         %55 = OpTypeFunction %void
-%compute_main_inner = OpFunction %void None %15
-%local_invocation_index = OpFunctionParameter %uint
+         %68 = OpConstantNull %int
+         %70 = OpTypeFunction %void
+   %tint_div = OpFunction %uint None %15
+        %lhs = OpFunctionParameter %uint
+        %rhs = OpFunctionParameter %uint
          %19 = OpLabel
+         %21 = OpIEqual %bool %rhs %6
+         %20 = OpSelect %uint %21 %uint_1 %rhs
+         %23 = OpUDiv %uint %lhs %20
+               OpReturnValue %23
+               OpFunctionEnd
+   %tint_mod = OpFunction %uint None %15
+      %lhs_0 = OpFunctionParameter %uint
+      %rhs_0 = OpFunctionParameter %uint
+         %27 = OpLabel
+         %29 = OpIEqual %bool %rhs_0 %6
+         %28 = OpSelect %uint %29 %uint_1 %rhs_0
+         %30 = OpUMod %uint %lhs_0 %28
+               OpReturnValue %30
+               OpFunctionEnd
+%compute_main_inner = OpFunction %void None %31
+%local_invocation_index = OpFunctionParameter %uint
+         %35 = OpLabel
         %idx = OpVariable %_ptr_Function_uint Function %6
                OpStore %idx %6
                OpStore %idx %local_invocation_index
-               OpBranch %22
-         %22 = OpLabel
-               OpLoopMerge %23 %24 None
-               OpBranch %25
-         %25 = OpLabel
-         %26 = OpLoad %uint %idx
-         %29 = OpULessThan %bool %26 %uint_6
-         %27 = OpLogicalNot %bool %29
-               OpSelectionMerge %31 None
-               OpBranchConditional %27 %32 %31
-         %32 = OpLabel
-               OpBranch %23
-         %31 = OpLabel
-         %33 = OpLoad %uint %idx
-         %34 = OpLoad %uint %idx
-         %35 = OpLoad %uint %idx
-         %39 = OpUDiv %uint %33 %uint_2
-         %40 = OpUMod %uint %34 %uint_2
-         %41 = OpUMod %uint %35 %uint_1
-         %43 = OpAccessChain %_ptr_Workgroup_uint %wg %39 %40 %41
-               OpAtomicStore %43 %uint_2 %uint_0 %6
-               OpBranch %24
-         %24 = OpLabel
-         %44 = OpLoad %uint %idx
-         %45 = OpIAdd %uint %44 %uint_1
-               OpStore %idx %45
-               OpBranch %22
-         %23 = OpLabel
+               OpBranch %38
+         %38 = OpLabel
+               OpLoopMerge %39 %40 None
+               OpBranch %41
+         %41 = OpLabel
+         %42 = OpLoad %uint %idx
+         %45 = OpULessThan %bool %42 %uint_6
+         %43 = OpLogicalNot %bool %45
+               OpSelectionMerge %46 None
+               OpBranchConditional %43 %47 %46
+         %47 = OpLabel
+               OpBranch %39
+         %46 = OpLabel
+         %48 = OpLoad %uint %idx
+         %49 = OpLoad %uint %idx
+         %50 = OpLoad %uint %idx
+         %51 = OpFunctionCall %uint %tint_div %48 %uint_2
+         %52 = OpFunctionCall %uint %tint_mod %49 %uint_2
+         %53 = OpFunctionCall %uint %tint_mod %50 %uint_1
+         %58 = OpAccessChain %_ptr_Workgroup_uint %wg %51 %52 %53
+               OpAtomicStore %58 %uint_2 %uint_0 %6
+               OpBranch %40
+         %40 = OpLabel
+         %59 = OpLoad %uint %idx
+         %60 = OpIAdd %uint %59 %uint_1
+               OpStore %idx %60
+               OpBranch %38
+         %39 = OpLabel
                OpControlBarrier %uint_2 %uint_2 %uint_264
-         %54 = OpAccessChain %_ptr_Workgroup_uint %wg %int_2 %int_1 %53
-               OpAtomicStore %54 %uint_2 %uint_0 %uint_1
+         %69 = OpAccessChain %_ptr_Workgroup_uint %wg %int_2 %int_1 %68
+               OpAtomicStore %69 %uint_2 %uint_0 %uint_1
                OpReturn
                OpFunctionEnd
-%compute_main_1 = OpFunction %void None %55
-         %57 = OpLabel
-         %58 = OpLoad %uint %local_invocation_index_1
-         %59 = OpFunctionCall %void %compute_main_inner %58
+%compute_main_1 = OpFunction %void None %70
+         %72 = OpLabel
+         %73 = OpLoad %uint %local_invocation_index_1
+         %74 = OpFunctionCall %void %compute_main_inner %73
                OpReturn
                OpFunctionEnd
-%compute_main_inner_1 = OpFunction %void None %15
+%compute_main_inner_1 = OpFunction %void None %31
 %local_invocation_index_1_param = OpFunctionParameter %uint
-         %62 = OpLabel
+         %77 = OpLabel
       %idx_1 = OpVariable %_ptr_Function_uint Function %6
                OpStore %idx_1 %local_invocation_index_1_param
-               OpBranch %64
-         %64 = OpLabel
-               OpLoopMerge %65 %66 None
-               OpBranch %67
-         %67 = OpLabel
-         %69 = OpLoad %uint %idx_1
-         %70 = OpULessThan %bool %69 %uint_6
-         %68 = OpLogicalNot %bool %70
-               OpSelectionMerge %71 None
-               OpBranchConditional %68 %72 %71
-         %72 = OpLabel
-               OpBranch %65
-         %71 = OpLabel
-         %73 = OpLoad %uint %idx_1
-         %74 = OpUDiv %uint %73 %uint_2
-         %75 = OpLoad %uint %idx_1
-         %76 = OpUMod %uint %75 %uint_2
-         %77 = OpLoad %uint %idx_1
-         %78 = OpUMod %uint %77 %uint_1
-         %81 = OpAccessChain %_ptr_Workgroup_uint %wg %74 %76 %78
-               OpAtomicStore %81 %uint_2 %uint_0 %6
-               OpBranch %66
-         %66 = OpLabel
-         %82 = OpLoad %uint %idx_1
-         %83 = OpIAdd %uint %82 %uint_1
-               OpStore %idx_1 %83
-               OpBranch %64
-         %65 = OpLabel
+               OpBranch %79
+         %79 = OpLabel
+               OpLoopMerge %80 %81 None
+               OpBranch %82
+         %82 = OpLabel
+         %84 = OpLoad %uint %idx_1
+         %85 = OpULessThan %bool %84 %uint_6
+         %83 = OpLogicalNot %bool %85
+               OpSelectionMerge %86 None
+               OpBranchConditional %83 %87 %86
+         %87 = OpLabel
+               OpBranch %80
+         %86 = OpLabel
+         %88 = OpLoad %uint %idx_1
+         %89 = OpUDiv %uint %88 %uint_2
+         %90 = OpLoad %uint %idx_1
+         %91 = OpUMod %uint %90 %uint_2
+         %92 = OpLoad %uint %idx_1
+         %93 = OpUMod %uint %92 %uint_1
+         %96 = OpAccessChain %_ptr_Workgroup_uint %wg %89 %91 %93
+               OpAtomicStore %96 %uint_2 %uint_0 %6
+               OpBranch %81
+         %81 = OpLabel
+         %97 = OpLoad %uint %idx_1
+         %98 = OpIAdd %uint %97 %uint_1
+               OpStore %idx_1 %98
+               OpBranch %79
+         %80 = OpLabel
                OpControlBarrier %uint_2 %uint_2 %uint_264
                OpStore %local_invocation_index_1 %local_invocation_index_1_param
-         %85 = OpFunctionCall %void %compute_main_1
+        %100 = OpFunctionCall %void %compute_main_1
                OpReturn
                OpFunctionEnd
-%compute_main = OpFunction %void None %55
-         %87 = OpLabel
-         %89 = OpLoad %uint %local_invocation_index_1_param_1
-         %88 = OpFunctionCall %void %compute_main_inner_1 %89
+%compute_main = OpFunction %void None %70
+        %102 = OpLabel
+        %104 = OpLoad %uint %local_invocation_index_1_param_1
+        %103 = OpFunctionCall %void %compute_main_inner_1 %104
                OpReturn
                OpFunctionEnd

test/tint/builtins/atomicStore/array/arrays.spvasm.expected.dxc.hlsl

@@ -1,3 +1,11 @@
+uint tint_div(uint lhs, uint rhs) {
+  return (lhs / ((rhs == 0u) ? 1u : rhs));
+}
+
+uint tint_mod(uint lhs, uint rhs) {
+  return (lhs % ((rhs == 0u) ? 1u : rhs));
+}
+
 static uint local_invocation_index_1 = 0u;
 groupshared uint wg[3][2][1];
 
@@ -12,8 +20,11 @@ void compute_main_inner(uint local_invocation_index) {
     const uint x_31 = idx;
     const uint x_33 = idx;
     const uint x_35 = idx;
+    const uint tint_symbol_2 = tint_div(x_31, 2u);
+    const uint tint_symbol_3 = tint_mod(x_33, 2u);
+    const uint tint_symbol_4 = tint_mod(x_35, 1u);
     uint atomic_result = 0u;
-    InterlockedExchange(wg[(x_31 / 2u)][(x_33 % 2u)][(x_35 % 1u)], 0u, atomic_result);
+    InterlockedExchange(wg[tint_symbol_2][tint_symbol_3][tint_symbol_4], 0u, atomic_result);
     {
       const uint x_42 = idx;
       idx = (x_42 + 1u);

test/tint/builtins/atomicStore/array/arrays.spvasm.expected.fxc.hlsl

@@ -1,3 +1,11 @@
+uint tint_div(uint lhs, uint rhs) {
+  return (lhs / ((rhs == 0u) ? 1u : rhs));
+}
+
+uint tint_mod(uint lhs, uint rhs) {
+  return (lhs % ((rhs == 0u) ? 1u : rhs));
+}
+
 static uint local_invocation_index_1 = 0u;
 groupshared uint wg[3][2][1];
 
@@ -12,8 +20,11 @@ void compute_main_inner(uint local_invocation_index) {
     const uint x_31 = idx;
     const uint x_33 = idx;
     const uint x_35 = idx;
+    const uint tint_symbol_2 = tint_div(x_31, 2u);
+    const uint tint_symbol_3 = tint_mod(x_33, 2u);
+    const uint tint_symbol_4 = tint_mod(x_35, 1u);
     uint atomic_result = 0u;
-    InterlockedExchange(wg[(x_31 / 2u)][(x_33 % 2u)][(x_35 % 1u)], 0u, atomic_result);
+    InterlockedExchange(wg[tint_symbol_2][tint_symbol_3][tint_symbol_4], 0u, atomic_result);
     {
       const uint x_42 = idx;
       idx = (x_42 + 1u);

test/tint/builtins/atomicStore/array/arrays.spvasm.expected.glsl

@@ -1,5 +1,13 @@
 #version 310 es
 
+uint tint_div(uint lhs, uint rhs) {
+  return (lhs / ((rhs == 0u) ? 1u : rhs));
+}
+
+uint tint_mod(uint lhs, uint rhs) {
+  return (lhs % ((rhs == 0u) ? 1u : rhs));
+}
+
 uint local_invocation_index_1 = 0u;
 shared uint wg[3][2][1];
 void compute_main_inner(uint local_invocation_index) {
@@ -13,7 +21,10 @@ void compute_main_inner(uint local_invocation_index) {
     uint x_31 = idx;
     uint x_33 = idx;
     uint x_35 = idx;
-    atomicExchange(wg[(x_31 / 2u)][(x_33 % 2u)][(x_35 % 1u)], 0u);
+    uint tint_symbol = tint_div(x_31, 2u);
+    uint tint_symbol_1 = tint_mod(x_33, 2u);
+    uint tint_symbol_2 = tint_mod(x_35, 1u);
+    atomicExchange(wg[tint_symbol][tint_symbol_1][tint_symbol_2], 0u);
     {
       uint x_42 = idx;
       idx = (x_42 + 1u);

test/tint/builtins/atomicStore/array/arrays.spvasm.expected.msl

@@ -14,7 +14,15 @@ struct tint_array {
     T elements[N];
 };
 
-void compute_main_inner(uint local_invocation_index, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol) {
+uint tint_div(uint lhs, uint rhs) {
+  return (lhs / select(rhs, 1u, (rhs == 0u)));
+}
+
+uint tint_mod(uint lhs, uint rhs) {
+  return (lhs % select(rhs, 1u, (rhs == 0u)));
+}
+
+void compute_main_inner(uint local_invocation_index, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_3) {
   uint idx = 0u;
   idx = local_invocation_index;
   while (true) {
@@ -25,39 +33,42 @@ void compute_main_inner(uint local_invocation_index, threadgroup tint_array<tint
     uint const x_31 = idx;
     uint const x_33 = idx;
     uint const x_35 = idx;
-    atomic_store_explicit(&((*(tint_symbol))[(x_31 / 2u)][(x_33 % 2u)][(x_35 % 1u)]), 0u, memory_order_relaxed);
+    uint const tint_symbol = tint_div(x_31, 2u);
+    uint const tint_symbol_1 = tint_mod(x_33, 2u);
+    uint const tint_symbol_2 = tint_mod(x_35, 1u);
+    atomic_store_explicit(&((*(tint_symbol_3))[tint_symbol][tint_symbol_1][tint_symbol_2]), 0u, memory_order_relaxed);
     {
       uint const x_42 = idx;
       idx = (x_42 + 1u);
     }
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  atomic_store_explicit(&((*(tint_symbol))[2][1][0]), 1u, memory_order_relaxed);
+  atomic_store_explicit(&((*(tint_symbol_3))[2][1][0]), 1u, memory_order_relaxed);
   return;
 }
 
-void compute_main_1(thread uint* const tint_symbol_1, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_2) {
-  uint const x_57 = *(tint_symbol_1);
-  compute_main_inner(x_57, tint_symbol_2);
+void compute_main_1(thread uint* const tint_symbol_4, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_5) {
+  uint const x_57 = *(tint_symbol_4);
+  compute_main_inner(x_57, tint_symbol_5);
   return;
 }
 
-void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_3, thread uint* const tint_symbol_4) {
+void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_6, thread uint* const tint_symbol_7) {
   for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 6u); idx_1 = (idx_1 + 1u)) {
     uint const i = (idx_1 / 2u);
     uint const i_1 = (idx_1 % 2u);
     uint const i_2 = (idx_1 % 1u);
-    atomic_store_explicit(&((*(tint_symbol_3))[i][i_1][i_2]), 0u, memory_order_relaxed);
+    atomic_store_explicit(&((*(tint_symbol_6))[i][i_1][i_2]), 0u, memory_order_relaxed);
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  *(tint_symbol_4) = local_invocation_index_1_param;
-  compute_main_1(tint_symbol_4, tint_symbol_3);
+  *(tint_symbol_7) = local_invocation_index_1_param;
+  compute_main_1(tint_symbol_7, tint_symbol_6);
 }
 
 kernel void compute_main(uint local_invocation_index_1_param [[thread_index_in_threadgroup]]) {
-  threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol_5;
-  thread uint tint_symbol_6 = 0u;
-  compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_5), &(tint_symbol_6));
+  threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol_8;
+  thread uint tint_symbol_9 = 0u;
+  compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_8), &(tint_symbol_9));
   return;
 }
 

test/tint/builtins/atomicStore/array/arrays.spvasm.expected.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 90
+; Bound: 105
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
@@ -10,6 +10,12 @@
                OpName %local_invocation_index_1_param_1 "local_invocation_index_1_param_1"
                OpName %local_invocation_index_1 "local_invocation_index_1"
                OpName %wg "wg"
+               OpName %tint_div "tint_div"
+               OpName %lhs "lhs"
+               OpName %rhs "rhs"
+               OpName %tint_mod "tint_mod"
+               OpName %lhs_0 "lhs"
+               OpName %rhs_0 "rhs"
                OpName %compute_main_inner "compute_main_inner"
                OpName %local_invocation_index "local_invocation_index"
                OpName %idx "idx"
@@ -36,105 +42,124 @@
 %_arr__arr__arr_uint_uint_1_uint_2_uint_3 = OpTypeArray %_arr__arr_uint_uint_1_uint_2 %uint_3
 %_ptr_Workgroup__arr__arr__arr_uint_uint_1_uint_2_uint_3 = OpTypePointer Workgroup %_arr__arr__arr_uint_uint_1_uint_2_uint_3
          %wg = OpVariable %_ptr_Workgroup__arr__arr__arr_uint_uint_1_uint_2_uint_3 Workgroup
+         %15 = OpTypeFunction %uint %uint %uint
+       %bool = OpTypeBool
        %void = OpTypeVoid
-         %15 = OpTypeFunction %void %uint
+         %31 = OpTypeFunction %void %uint
 %_ptr_Function_uint = OpTypePointer Function %uint
      %uint_6 = OpConstant %uint 6
-       %bool = OpTypeBool
      %uint_0 = OpConstant %uint 0
 %_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
    %uint_264 = OpConstant %uint 264
         %int = OpTypeInt 32 1
       %int_2 = OpConstant %int 2
       %int_1 = OpConstant %int 1
-         %53 = OpConstantNull %int
-         %55 = OpTypeFunction %void
-%compute_main_inner = OpFunction %void None %15
-%local_invocation_index = OpFunctionParameter %uint
+         %68 = OpConstantNull %int
+         %70 = OpTypeFunction %void
+   %tint_div = OpFunction %uint None %15
+        %lhs = OpFunctionParameter %uint
+        %rhs = OpFunctionParameter %uint
          %19 = OpLabel
+         %21 = OpIEqual %bool %rhs %6
+         %20 = OpSelect %uint %21 %uint_1 %rhs
+         %23 = OpUDiv %uint %lhs %20
+               OpReturnValue %23
+               OpFunctionEnd
+   %tint_mod = OpFunction %uint None %15
+      %lhs_0 = OpFunctionParameter %uint
+      %rhs_0 = OpFunctionParameter %uint
+         %27 = OpLabel
+         %29 = OpIEqual %bool %rhs_0 %6
+         %28 = OpSelect %uint %29 %uint_1 %rhs_0
+         %30 = OpUMod %uint %lhs_0 %28
+               OpReturnValue %30
+               OpFunctionEnd
+%compute_main_inner = OpFunction %void None %31
+%local_invocation_index = OpFunctionParameter %uint
+         %35 = OpLabel
         %idx = OpVariable %_ptr_Function_uint Function %6
                OpStore %idx %6
                OpStore %idx %local_invocation_index
-               OpBranch %22
-         %22 = OpLabel
-               OpLoopMerge %23 %24 None
-               OpBranch %25
-         %25 = OpLabel
-         %26 = OpLoad %uint %idx
-         %29 = OpULessThan %bool %26 %uint_6
-         %27 = OpLogicalNot %bool %29
-               OpSelectionMerge %31 None
-               OpBranchConditional %27 %32 %31
-         %32 = OpLabel
-               OpBranch %23
-         %31 = OpLabel
-         %33 = OpLoad %uint %idx
-         %34 = OpLoad %uint %idx
-         %35 = OpLoad %uint %idx
-         %39 = OpUDiv %uint %33 %uint_2
-         %40 = OpUMod %uint %34 %uint_2
-         %41 = OpUMod %uint %35 %uint_1
-         %43 = OpAccessChain %_ptr_Workgroup_uint %wg %39 %40 %41
-               OpAtomicStore %43 %uint_2 %uint_0 %6
-               OpBranch %24
-         %24 = OpLabel
-         %44 = OpLoad %uint %idx
-         %45 = OpIAdd %uint %44 %uint_1
-               OpStore %idx %45
-               OpBranch %22
-         %23 = OpLabel
+               OpBranch %38
+         %38 = OpLabel
+               OpLoopMerge %39 %40 None
+               OpBranch %41
+         %41 = OpLabel
+         %42 = OpLoad %uint %idx
+         %45 = OpULessThan %bool %42 %uint_6
+         %43 = OpLogicalNot %bool %45
+               OpSelectionMerge %46 None
+               OpBranchConditional %43 %47 %46
+         %47 = OpLabel
+               OpBranch %39
+         %46 = OpLabel
+         %48 = OpLoad %uint %idx
+         %49 = OpLoad %uint %idx
+         %50 = OpLoad %uint %idx
+         %51 = OpFunctionCall %uint %tint_div %48 %uint_2
+         %52 = OpFunctionCall %uint %tint_mod %49 %uint_2
+         %53 = OpFunctionCall %uint %tint_mod %50 %uint_1
+         %58 = OpAccessChain %_ptr_Workgroup_uint %wg %51 %52 %53
+               OpAtomicStore %58 %uint_2 %uint_0 %6
+               OpBranch %40
+         %40 = OpLabel
+         %59 = OpLoad %uint %idx
+         %60 = OpIAdd %uint %59 %uint_1
+               OpStore %idx %60
+               OpBranch %38
+         %39 = OpLabel
                OpControlBarrier %uint_2 %uint_2 %uint_264
-         %54 = OpAccessChain %_ptr_Workgroup_uint %wg %int_2 %int_1 %53
-               OpAtomicStore %54 %uint_2 %uint_0 %uint_1
+         %69 = OpAccessChain %_ptr_Workgroup_uint %wg %int_2 %int_1 %68
+               OpAtomicStore %69 %uint_2 %uint_0 %uint_1
                OpReturn
                OpFunctionEnd
-%compute_main_1 = OpFunction %void None %55
-         %57 = OpLabel
-         %58 = OpLoad %uint %local_invocation_index_1
-         %59 = OpFunctionCall %void %compute_main_inner %58
+%compute_main_1 = OpFunction %void None %70
+         %72 = OpLabel
+         %73 = OpLoad %uint %local_invocation_index_1
+         %74 = OpFunctionCall %void %compute_main_inner %73
                OpReturn
                OpFunctionEnd
-%compute_main_inner_1 = OpFunction %void None %15
+%compute_main_inner_1 = OpFunction %void None %31
 %local_invocation_index_1_param = OpFunctionParameter %uint
-         %62 = OpLabel
+         %77 = OpLabel
       %idx_1 = OpVariable %_ptr_Function_uint Function %6
                OpStore %idx_1 %local_invocation_index_1_param
-               OpBranch %64
-         %64 = OpLabel
-               OpLoopMerge %65 %66 None
-               OpBranch %67
-         %67 = OpLabel
-         %69 = OpLoad %uint %idx_1
-         %70 = OpULessThan %bool %69 %uint_6
-         %68 = OpLogicalNot %bool %70
-               OpSelectionMerge %71 None
-               OpBranchConditional %68 %72 %71
-         %72 = OpLabel
-               OpBranch %65
-         %71 = OpLabel
-         %73 = OpLoad %uint %idx_1
-         %74 = OpUDiv %uint %73 %uint_2
-         %75 = OpLoad %uint %idx_1
-         %76 = OpUMod %uint %75 %uint_2
-         %77 = OpLoad %uint %idx_1
-         %78 = OpUMod %uint %77 %uint_1
-         %81 = OpAccessChain %_ptr_Workgroup_uint %wg %74 %76 %78
-               OpAtomicStore %81 %uint_2 %uint_0 %6
-               OpBranch %66
-         %66 = OpLabel
-         %82 = OpLoad %uint %idx_1
-         %83 = OpIAdd %uint %82 %uint_1
-               OpStore %idx_1 %83
-               OpBranch %64
-         %65 = OpLabel
+               OpBranch %79
+         %79 = OpLabel
+               OpLoopMerge %80 %81 None
+               OpBranch %82
+         %82 = OpLabel
+         %84 = OpLoad %uint %idx_1
+         %85 = OpULessThan %bool %84 %uint_6
+         %83 = OpLogicalNot %bool %85
+               OpSelectionMerge %86 None
+               OpBranchConditional %83 %87 %86
+         %87 = OpLabel
+               OpBranch %80
+         %86 = OpLabel
+         %88 = OpLoad %uint %idx_1
+         %89 = OpUDiv %uint %88 %uint_2
+         %90 = OpLoad %uint %idx_1
+         %91 = OpUMod %uint %90 %uint_2
+         %92 = OpLoad %uint %idx_1
+         %93 = OpUMod %uint %92 %uint_1
+         %96 = OpAccessChain %_ptr_Workgroup_uint %wg %89 %91 %93
+               OpAtomicStore %96 %uint_2 %uint_0 %6
+               OpBranch %81
+         %81 = OpLabel
+         %97 = OpLoad %uint %idx_1
+         %98 = OpIAdd %uint %97 %uint_1
+               OpStore %idx_1 %98
+               OpBranch %79
+         %80 = OpLabel
                OpControlBarrier %uint_2 %uint_2 %uint_264
                OpStore %local_invocation_index_1 %local_invocation_index_1_param
-         %85 = OpFunctionCall %void %compute_main_1
+        %100 = OpFunctionCall %void %compute_main_1
                OpReturn
                OpFunctionEnd
-%compute_main = OpFunction %void None %55
-         %87 = OpLabel
-         %89 = OpLoad %uint %local_invocation_index_1_param_1
-         %88 = OpFunctionCall %void %compute_main_inner_1 %89
+%compute_main = OpFunction %void None %70
+        %102 = OpLabel
+        %104 = OpLoad %uint %local_invocation_index_1_param_1
+        %103 = OpFunctionCall %void %compute_main_inner_1 %104
                OpReturn
                OpFunctionEnd