tint: Resolve @builtin() args as expressions

This CL makes the builtin argument resolve as a shadowable enumerator expression. Bug: tint:1841 Bug: tint:1845 Change-Id: I2aaa2c63025752f25c113cf2fb38f4d91c2c16c2 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/120680 Kokoro: Ben Clayton <bclayton@google.com> Reviewed-by: Dan Sinclair <dsinclair@chromium.org> Commit-Queue: Ben Clayton <bclayton@google.com> Reviewed-by: James Price <jrprice@google.com>
2025-12-09 21:47:47 +00:00 · 2023-02-21 17:33:54 +00:00
parent 21571709a2
commit 4d3ff9711e
448 changed files with 2044 additions and 1536 deletions
--- a/test/tint/bug/tint/1088.spvasm.expected.dxc.hlsl
+++ b/test/tint/bug/tint/1088.spvasm.expected.dxc.hlsl
@@ -1,4 +1,4 @@
-static float3 position = float3(0.0f, 0.0f, 0.0f);
+static float3 position_1 = float3(0.0f, 0.0f, 0.0f);
 cbuffer cbuffer_x_14 : register(b2, space2) {
  uint4 x_14[17];
 };
@@ -18,13 +18,13 @@ float4x4 tint_symbol_4(uint4 buffer[17], uint offset) {
 void main_1() {
  float4 q = float4(0.0f, 0.0f, 0.0f, 0.0f);
  float3 p = float3(0.0f, 0.0f, 0.0f);
-  const float3 x_13 = position;
+  const float3 x_13 = position_1;
  q = float4(x_13.x, x_13.y, x_13.z, 1.0f);
  const float4 x_21 = q;
  p = float3(x_21.x, x_21.y, x_21.z);
  const float x_27 = p.x;
  const float x_41 = asfloat(x_14[13].x);
-  const float x_45 = position.y;
+  const float x_45 = position_1.y;
  const float x_49 = asfloat(x_14[4].x);
  p.x = (x_27 + sin(((x_41 * x_45) + x_49)));
  const float x_55 = p.y;
@@ -45,7 +45,7 @@ struct main_out {
  float2 vUV_1;
 };
 struct tint_symbol_1 {
-  float3 position_param : TEXCOORD0;
+  float3 position_1_param : TEXCOORD0;
  float3 normal_param : TEXCOORD1;
  float2 uv_param : TEXCOORD2;
 };
@@ -54,8 +54,8 @@ struct tint_symbol_2 {
  float4 gl_Position : SV_Position;
 };

-main_out main_inner(float3 position_param, float2 uv_param, float3 normal_param) {
-  position = position_param;
+main_out main_inner(float3 position_1_param, float2 uv_param, float3 normal_param) {
+  position_1 = position_1_param;
  uv = uv_param;
  normal = normal_param;
  main_1();
@@ -64,7 +64,7 @@ main_out main_inner(float3 position_param, float2 uv_param, float3 normal_param)
 }

 tint_symbol_2 main(tint_symbol_1 tint_symbol) {
-  const main_out inner_result = main_inner(tint_symbol.position_param, tint_symbol.uv_param, tint_symbol.normal_param);
+  const main_out inner_result = main_inner(tint_symbol.position_1_param, tint_symbol.uv_param, tint_symbol.normal_param);
  tint_symbol_2 wrapper_result = (tint_symbol_2)0;
  wrapper_result.gl_Position = inner_result.gl_Position;
  wrapper_result.vUV_1 = inner_result.vUV_1;
--- a/test/tint/bug/tint/1088.spvasm.expected.fxc.hlsl
+++ b/test/tint/bug/tint/1088.spvasm.expected.fxc.hlsl
@@ -1,4 +1,4 @@
-static float3 position = float3(0.0f, 0.0f, 0.0f);
+static float3 position_1 = float3(0.0f, 0.0f, 0.0f);
 cbuffer cbuffer_x_14 : register(b2, space2) {
  uint4 x_14[17];
 };
@@ -18,13 +18,13 @@ float4x4 tint_symbol_4(uint4 buffer[17], uint offset) {
 void main_1() {
  float4 q = float4(0.0f, 0.0f, 0.0f, 0.0f);
  float3 p = float3(0.0f, 0.0f, 0.0f);
-  const float3 x_13 = position;
+  const float3 x_13 = position_1;
  q = float4(x_13.x, x_13.y, x_13.z, 1.0f);
  const float4 x_21 = q;
  p = float3(x_21.x, x_21.y, x_21.z);
  const float x_27 = p.x;
  const float x_41 = asfloat(x_14[13].x);
-  const float x_45 = position.y;
+  const float x_45 = position_1.y;
  const float x_49 = asfloat(x_14[4].x);
  p.x = (x_27 + sin(((x_41 * x_45) + x_49)));
  const float x_55 = p.y;
@@ -45,7 +45,7 @@ struct main_out {
  float2 vUV_1;
 };
 struct tint_symbol_1 {
-  float3 position_param : TEXCOORD0;
+  float3 position_1_param : TEXCOORD0;
  float3 normal_param : TEXCOORD1;
  float2 uv_param : TEXCOORD2;
 };
@@ -54,8 +54,8 @@ struct tint_symbol_2 {
  float4 gl_Position : SV_Position;
 };

-main_out main_inner(float3 position_param, float2 uv_param, float3 normal_param) {
-  position = position_param;
+main_out main_inner(float3 position_1_param, float2 uv_param, float3 normal_param) {
+  position_1 = position_1_param;
  uv = uv_param;
  normal = normal_param;
  main_1();
@@ -64,7 +64,7 @@ main_out main_inner(float3 position_param, float2 uv_param, float3 normal_param)
 }

 tint_symbol_2 main(tint_symbol_1 tint_symbol) {
-  const main_out inner_result = main_inner(tint_symbol.position_param, tint_symbol.uv_param, tint_symbol.normal_param);
+  const main_out inner_result = main_inner(tint_symbol.position_1_param, tint_symbol.uv_param, tint_symbol.normal_param);
  tint_symbol_2 wrapper_result = (tint_symbol_2)0;
  wrapper_result.gl_Position = inner_result.gl_Position;
  wrapper_result.vUV_1 = inner_result.vUV_1;
--- a/test/tint/bug/tint/1088.spvasm.expected.glsl
+++ b/test/tint/bug/tint/1088.spvasm.expected.glsl
@@ -1,6 +1,6 @@
 #version 310 es

-layout(location = 0) in vec3 position_param_1;
+layout(location = 0) in vec3 position_1_param_1;
 layout(location = 2) in vec2 uv_param_1;
 layout(location = 1) in vec3 normal_param_1;
 layout(location = 0) out vec2 vUV_1_1;
@@ -21,7 +21,7 @@ struct LeftOver {
  strided_arr test[4];
 };

-vec3 position = vec3(0.0f, 0.0f, 0.0f);
+vec3 position_1 = vec3(0.0f, 0.0f, 0.0f);
 layout(binding = 2, std140) uniform x_14_block_ubo {
  LeftOver inner;
 } x_14;
@@ -33,13 +33,13 @@ vec4 tint_symbol = vec4(0.0f, 0.0f, 0.0f, 0.0f);
 void main_1() {
  vec4 q = vec4(0.0f, 0.0f, 0.0f, 0.0f);
  vec3 p = vec3(0.0f, 0.0f, 0.0f);
-  vec3 x_13 = position;
+  vec3 x_13 = position_1;
  q = vec4(x_13.x, x_13.y, x_13.z, 1.0f);
  vec4 x_21 = q;
  p = vec3(x_21.x, x_21.y, x_21.z);
  float x_27 = p.x;
  float x_41 = x_14.inner.test[0].el;
-  float x_45 = position.y;
+  float x_45 = position_1.y;
  float x_49 = x_14.inner.time;
  p.x = (x_27 + sin(((x_41 * x_45) + x_49)));
  float x_55 = p.y;
@@ -60,8 +60,8 @@ struct main_out {
  vec2 vUV_1;
 };

-main_out tint_symbol_1(vec3 position_param, vec2 uv_param, vec3 normal_param) {
-  position = position_param;
+main_out tint_symbol_1(vec3 position_1_param, vec2 uv_param, vec3 normal_param) {
+  position_1 = position_1_param;
  uv = uv_param;
  normal = normal_param;
  main_1();
@@ -71,7 +71,7 @@ main_out tint_symbol_1(vec3 position_param, vec2 uv_param, vec3 normal_param) {

 void main() {
  gl_PointSize = 1.0;
-  main_out inner_result = tint_symbol_1(position_param_1, uv_param_1, normal_param_1);
+  main_out inner_result = tint_symbol_1(position_1_param_1, uv_param_1, normal_param_1);
  gl_Position = inner_result.tint_symbol;
  vUV_1_1 = inner_result.vUV_1;
  gl_Position.y = -(gl_Position.y);
--- a/test/tint/bug/tint/1088.spvasm.expected.msl
+++ b/test/tint/bug/tint/1088.spvasm.expected.msl
@@ -58,7 +58,7 @@ struct main_out {
 };

 struct tint_symbol_2 {
-  float3 position_param [[attribute(0)]];
+  float3 position_1_param [[attribute(0)]];
  float3 normal_param [[attribute(1)]];
  float2 uv_param [[attribute(2)]];
 };
@@ -68,9 +68,9 @@ struct tint_symbol_3 {
  float4 gl_Position [[position]];
 };

-main_out tint_symbol_inner(float3 position_param, float2 uv_param, float3 normal_param, thread float3* const tint_symbol_10, thread float2* const tint_symbol_11, const constant LeftOver* const tint_symbol_13, thread float4* const tint_symbol_14, thread float2* const tint_symbol_15) {
+main_out tint_symbol_inner(float3 position_1_param, float2 uv_param, float3 normal_param, thread float3* const tint_symbol_10, thread float2* const tint_symbol_11, const constant LeftOver* const tint_symbol_13, thread float4* const tint_symbol_14, thread float2* const tint_symbol_15) {
  thread float3 tint_symbol_12 = 0.0f;
-  *(tint_symbol_10) = position_param;
+  *(tint_symbol_10) = position_1_param;
  *(tint_symbol_11) = uv_param;
  tint_symbol_12 = normal_param;
  main_1(tint_symbol_10, tint_symbol_13, tint_symbol_14, tint_symbol_11, tint_symbol_15);
@@ -83,7 +83,7 @@ vertex tint_symbol_3 tint_symbol(const constant LeftOver* tint_symbol_18 [[buffe
  thread float2 tint_symbol_17 = 0.0f;
  thread float4 tint_symbol_19 = 0.0f;
  thread float2 tint_symbol_20 = 0.0f;
-  main_out const inner_result = tint_symbol_inner(tint_symbol_1.position_param, tint_symbol_1.uv_param, tint_symbol_1.normal_param, &(tint_symbol_16), &(tint_symbol_17), tint_symbol_18, &(tint_symbol_19), &(tint_symbol_20));
+  main_out const inner_result = tint_symbol_inner(tint_symbol_1.position_1_param, tint_symbol_1.uv_param, tint_symbol_1.normal_param, &(tint_symbol_16), &(tint_symbol_17), tint_symbol_18, &(tint_symbol_19), &(tint_symbol_20));
  tint_symbol_3 wrapper_result = {};
  wrapper_result.gl_Position = inner_result.gl_Position;
  wrapper_result.vUV_1 = inner_result.vUV_1;
--- a/test/tint/bug/tint/1088.spvasm.expected.spvasm
+++ b/test/tint/bug/tint/1088.spvasm.expected.spvasm
@@ -6,14 +6,14 @@
               OpCapability Shader
         %76 = OpExtInstImport "GLSL.std.450"
               OpMemoryModel Logical GLSL450
-               OpEntryPoint Vertex %main "main" %position_param_1 %uv_param_1 %normal_param_1 %gl_Position_1 %vUV_1_1 %vertex_point_size
-               OpName %position_param_1 "position_param_1"
+               OpEntryPoint Vertex %main "main" %position_1_param_1 %uv_param_1 %normal_param_1 %gl_Position_1 %vUV_1_1 %vertex_point_size
+               OpName %position_1_param_1 "position_1_param_1"
               OpName %uv_param_1 "uv_param_1"
               OpName %normal_param_1 "normal_param_1"
               OpName %gl_Position_1 "gl_Position_1"
               OpName %vUV_1_1 "vUV_1_1"
               OpName %vertex_point_size "vertex_point_size"
-               OpName %position "position"
+               OpName %position_1 "position_1"
               OpName %x_14_block "x_14_block"
               OpMemberName %x_14_block 0 "inner"
               OpName %LeftOver "LeftOver"
@@ -35,11 +35,11 @@
               OpMemberName %main_out 0 "gl_Position"
               OpMemberName %main_out 1 "vUV_1"
               OpName %main_inner "main_inner"
-               OpName %position_param "position_param"
+               OpName %position_1_param "position_1_param"
               OpName %uv_param "uv_param"
               OpName %normal_param "normal_param"
               OpName %main "main"
-               OpDecorate %position_param_1 Location 0
+               OpDecorate %position_1_param_1 Location 0
               OpDecorate %uv_param_1 Location 2
               OpDecorate %normal_param_1 Location 1
               OpDecorate %gl_Position_1 BuiltIn Position
@@ -66,7 +66,7 @@
      %float = OpTypeFloat 32
    %v3float = OpTypeVector %float 3
 %_ptr_Input_v3float = OpTypePointer Input %v3float
-%position_param_1 = OpVariable %_ptr_Input_v3float Input
+%position_1_param_1 = OpVariable %_ptr_Input_v3float Input
    %v2float = OpTypeVector %float 2
 %_ptr_Input_v2float = OpTypePointer Input %v2float
 %uv_param_1 = OpVariable %_ptr_Input_v2float Input
@@ -83,7 +83,7 @@
 %vertex_point_size = OpVariable %_ptr_Output_float Output %18
 %_ptr_Private_v3float = OpTypePointer Private %v3float
         %21 = OpConstantNull %v3float
-   %position = OpVariable %_ptr_Private_v3float Private %21
+ %position_1 = OpVariable %_ptr_Private_v3float Private %21
 %mat4v4float = OpTypeMatrix %v4float 4
       %uint = OpTypeInt 32 0
     %uint_2 = OpConstant %uint 2
@@ -123,7 +123,7 @@
         %42 = OpLabel
          %q = OpVariable %_ptr_Function_v4float Function %12
          %p = OpVariable %_ptr_Function_v3float Function %21
-         %47 = OpLoad %v3float %position
+         %47 = OpLoad %v3float %position_1
         %48 = OpCompositeExtract %float %47 0
         %49 = OpCompositeExtract %float %47 1
         %50 = OpCompositeExtract %float %47 2
@@ -139,7 +139,7 @@
         %61 = OpLoad %float %60
         %66 = OpAccessChain %_ptr_Uniform_float %x_14 %uint_0 %uint_3 %64 %uint_0
         %67 = OpLoad %float %66
-         %70 = OpAccessChain %_ptr_Private_float %position %uint_1
+         %70 = OpAccessChain %_ptr_Private_float %position_1 %uint_1
         %71 = OpLoad %float %70
         %72 = OpAccessChain %_ptr_Uniform_float %x_14 %uint_0 %uint_1
         %73 = OpLoad %float %72
@@ -177,11 +177,11 @@
               OpReturn
               OpFunctionEnd
 %main_inner = OpFunction %main_out None %104
-%position_param = OpFunctionParameter %v3float
+%position_1_param = OpFunctionParameter %v3float
   %uv_param = OpFunctionParameter %v2float
 %normal_param = OpFunctionParameter %v3float
        %110 = OpLabel
-               OpStore %position %position_param
+               OpStore %position_1 %position_1_param
               OpStore %uv %uv_param
               OpStore %normal %normal_param
        %111 = OpFunctionCall %void %main_1
@@ -192,7 +192,7 @@
               OpFunctionEnd
       %main = OpFunction %void None %39
        %116 = OpLabel
-        %118 = OpLoad %v3float %position_param_1
+        %118 = OpLoad %v3float %position_1_param_1
        %119 = OpLoad %v2float %uv_param_1
        %120 = OpLoad %v3float %normal_param_1
        %117 = OpFunctionCall %main_out %main_inner %118 %119 %120
--- a/test/tint/bug/tint/1088.spvasm.expected.wgsl
+++ b/test/tint/bug/tint/1088.spvasm.expected.wgsl
@@ -18,7 +18,7 @@ struct LeftOver {
  test : Arr_1,
 }

-var<private> position : vec3<f32>;
+var<private> position_1 : vec3<f32>;

@group(2) @binding(2) var<uniform> x_14 : LeftOver;

@@ -33,13 +33,13 @@ var<private> gl_Position : vec4<f32>;
 fn main_1() {
  var q : vec4<f32>;
  var p : vec3<f32>;
-  let x_13 : vec3<f32> = position;
+  let x_13 : vec3<f32> = position_1;
  q = vec4<f32>(x_13.x, x_13.y, x_13.z, 1.0f);
  let x_21 : vec4<f32> = q;
  p = vec3<f32>(x_21.x, x_21.y, x_21.z);
  let x_27 : f32 = p.x;
  let x_41 : f32 = x_14.test[0i].el;
-  let x_45 : f32 = position.y;
+  let x_45 : f32 = position_1.y;
  let x_49 : f32 = x_14.time;
  p.x = (x_27 + sin(((x_41 * x_45) + x_49)));
  let x_55 : f32 = p.y;
@@ -63,8 +63,8 @@ struct main_out {
 }

@vertex
-fn main(@location(0) position_param : vec3<f32>, @location(2) uv_param : vec2<f32>, @location(1) normal_param : vec3<f32>) -> main_out {
-  position = position_param;
+fn main(@location(0) position_1_param : vec3<f32>, @location(2) uv_param : vec2<f32>, @location(1) normal_param : vec3<f32>) -> main_out {
+  position_1 = position_1_param;
  uv = uv_param;
  normal = normal_param;
  main_1();
--- a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.dxc.hlsl
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.dxc.hlsl
@@ -9,9 +9,9 @@ uint tint_mod(uint lhs, uint rhs) {
  return (lhs % ((rhs == 0u) ? 1u : rhs));
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  uint idx = 0u;
-  idx = local_invocation_index;
+  idx = local_invocation_index_2;
  while (true) {
    const uint x_25 = idx;
    if (!((x_25 < 6u))) {
--- a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.fxc.hlsl
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.fxc.hlsl
@@ -9,9 +9,9 @@ uint tint_mod(uint lhs, uint rhs) {
  return (lhs % ((rhs == 0u) ? 1u : rhs));
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  uint idx = 0u;
-  idx = local_invocation_index;
+  idx = local_invocation_index_2;
  while (true) {
    const uint x_25 = idx;
    if (!((x_25 < 6u))) {
--- a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.glsl
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.glsl
@@ -10,9 +10,9 @@ uint tint_mod(uint lhs, uint rhs) {
  return (lhs % ((rhs == 0u) ? 1u : rhs));
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  uint idx = 0u;
-  idx = local_invocation_index;
+  idx = local_invocation_index_2;
  while (true) {
    uint x_25 = idx;
    if (!((x_25 < 6u))) {
--- a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.msl
@@ -22,9 +22,9 @@ 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) {
+void compute_main_inner(uint local_invocation_index_2, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_3) {
  uint idx = 0u;
-  idx = local_invocation_index;
+  idx = local_invocation_index_2;
  while (true) {
    uint const x_25 = idx;
    if (!((x_25 < 6u))) {
--- a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.spvasm
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.spvasm
@@ -17,7 +17,7 @@
               OpName %lhs_0 "lhs"
               OpName %rhs_0 "rhs"
               OpName %compute_main_inner "compute_main_inner"
-               OpName %local_invocation_index "local_invocation_index"
+               OpName %local_invocation_index_2 "local_invocation_index_2"
               OpName %idx "idx"
               OpName %compute_main_1 "compute_main_1"
               OpName %compute_main_inner_1 "compute_main_inner_1"
@@ -75,11 +75,11 @@
               OpReturnValue %30
               OpFunctionEnd
 %compute_main_inner = OpFunction %void None %31
-%local_invocation_index = OpFunctionParameter %uint
+%local_invocation_index_2 = OpFunctionParameter %uint
         %35 = OpLabel
        %idx = OpVariable %_ptr_Function_uint Function %6
               OpStore %idx %6
-               OpStore %idx %local_invocation_index
+               OpStore %idx %local_invocation_index_2
               OpBranch %38
         %38 = OpLabel
               OpLoopMerge %39 %40 None
--- a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.wgsl
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.wgsl
@@ -8,9 +8,9 @@ var<private> local_invocation_index_1 : u32;

 var<workgroup> wg : array<array<array<atomic<u32>, 1u>, 2u>, 3u>;

-fn compute_main_inner(local_invocation_index : u32) {
+fn compute_main_inner(local_invocation_index_2 : u32) {
  var idx : u32 = 0u;
-  idx = local_invocation_index;
+  idx = local_invocation_index_2;
  loop {
    let x_25 : u32 = idx;
    if (!((x_25 < 6u))) {
--- a/test/tint/builtins/atomicStore/array/array.spvasm.expected.dxc.hlsl
+++ b/test/tint/builtins/atomicStore/array/array.spvasm.expected.dxc.hlsl
@@ -1,9 +1,9 @@
 static uint local_invocation_index_1 = 0u;
 groupshared uint wg[4];

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  uint idx = 0u;
-  idx = local_invocation_index;
+  idx = local_invocation_index_2;
  while (true) {
    const uint x_21 = idx;
    if (!((x_21 < 4u))) {
--- a/test/tint/builtins/atomicStore/array/array.spvasm.expected.fxc.hlsl
+++ b/test/tint/builtins/atomicStore/array/array.spvasm.expected.fxc.hlsl
@@ -1,9 +1,9 @@
 static uint local_invocation_index_1 = 0u;
 groupshared uint wg[4];

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  uint idx = 0u;
-  idx = local_invocation_index;
+  idx = local_invocation_index_2;
  while (true) {
    const uint x_21 = idx;
    if (!((x_21 < 4u))) {
--- a/test/tint/builtins/atomicStore/array/array.spvasm.expected.glsl
+++ b/test/tint/builtins/atomicStore/array/array.spvasm.expected.glsl
@@ -2,9 +2,9 @@

 uint local_invocation_index_1 = 0u;
 shared uint wg[4];
-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  uint idx = 0u;
-  idx = local_invocation_index;
+  idx = local_invocation_index_2;
  while (true) {
    uint x_21 = idx;
    if (!((x_21 < 4u))) {
--- a/test/tint/builtins/atomicStore/array/array.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/array/array.spvasm.expected.msl
@@ -14,9 +14,9 @@ struct tint_array {
    T elements[N];
 };

-void compute_main_inner(uint local_invocation_index, threadgroup tint_array<atomic_uint, 4>* const tint_symbol) {
+void compute_main_inner(uint local_invocation_index_2, threadgroup tint_array<atomic_uint, 4>* const tint_symbol) {
  uint idx = 0u;
-  idx = local_invocation_index;
+  idx = local_invocation_index_2;
  while (true) {
    uint const x_21 = idx;
    if (!((x_21 < 4u))) {
--- a/test/tint/builtins/atomicStore/array/array.spvasm.expected.spvasm
+++ b/test/tint/builtins/atomicStore/array/array.spvasm.expected.spvasm
@@ -11,7 +11,7 @@
               OpName %local_invocation_index_1 "local_invocation_index_1"
               OpName %wg "wg"
               OpName %compute_main_inner "compute_main_inner"
-               OpName %local_invocation_index "local_invocation_index"
+               OpName %local_invocation_index_2 "local_invocation_index_2"
               OpName %idx "idx"
               OpName %compute_main_1 "compute_main_1"
               OpName %compute_main_inner_1 "compute_main_inner_1"
@@ -43,11 +43,11 @@
      %int_1 = OpConstant %int 1
         %45 = OpTypeFunction %void
 %compute_main_inner = OpFunction %void None %11
-%local_invocation_index = OpFunctionParameter %uint
+%local_invocation_index_2 = OpFunctionParameter %uint
         %15 = OpLabel
        %idx = OpVariable %_ptr_Function_uint Function %6
               OpStore %idx %6
-               OpStore %idx %local_invocation_index
+               OpStore %idx %local_invocation_index_2
               OpBranch %18
         %18 = OpLabel
               OpLoopMerge %19 %20 None
--- a/test/tint/builtins/atomicStore/array/array.spvasm.expected.wgsl
+++ b/test/tint/builtins/atomicStore/array/array.spvasm.expected.wgsl
@@ -4,9 +4,9 @@ var<private> local_invocation_index_1 : u32;

 var<workgroup> wg : array<atomic<u32>, 4u>;

-fn compute_main_inner(local_invocation_index : u32) {
+fn compute_main_inner(local_invocation_index_2 : u32) {
  var idx : u32 = 0u;
-  idx = local_invocation_index;
+  idx = local_invocation_index_2;
  loop {
    let x_21 : u32 = idx;
    if (!((x_21 < 4u))) {
--- a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.dxc.hlsl
+++ b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.dxc.hlsl
@@ -9,9 +9,9 @@ uint tint_mod(uint lhs, uint rhs) {
  return (lhs % ((rhs == 0u) ? 1u : rhs));
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  uint idx = 0u;
-  idx = local_invocation_index;
+  idx = local_invocation_index_2;
  while (true) {
    const uint x_25 = idx;
    if (!((x_25 < 6u))) {
--- a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.fxc.hlsl
+++ b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.fxc.hlsl
@@ -9,9 +9,9 @@ uint tint_mod(uint lhs, uint rhs) {
  return (lhs % ((rhs == 0u) ? 1u : rhs));
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  uint idx = 0u;
-  idx = local_invocation_index;
+  idx = local_invocation_index_2;
  while (true) {
    const uint x_25 = idx;
    if (!((x_25 < 6u))) {
--- a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.glsl
+++ b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.glsl
@@ -10,9 +10,9 @@ uint tint_mod(uint lhs, uint rhs) {
  return (lhs % ((rhs == 0u) ? 1u : rhs));
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  uint idx = 0u;
-  idx = local_invocation_index;
+  idx = local_invocation_index_2;
  while (true) {
    uint x_25 = idx;
    if (!((x_25 < 6u))) {
--- a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.msl
@@ -22,9 +22,9 @@ 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) {
+void compute_main_inner(uint local_invocation_index_2, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_3) {
  uint idx = 0u;
-  idx = local_invocation_index;
+  idx = local_invocation_index_2;
  while (true) {
    uint const x_25 = idx;
    if (!((x_25 < 6u))) {
--- a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.spvasm
+++ b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.spvasm
@@ -17,7 +17,7 @@
               OpName %lhs_0 "lhs"
               OpName %rhs_0 "rhs"
               OpName %compute_main_inner "compute_main_inner"
-               OpName %local_invocation_index "local_invocation_index"
+               OpName %local_invocation_index_2 "local_invocation_index_2"
               OpName %idx "idx"
               OpName %compute_main_1 "compute_main_1"
               OpName %compute_main_inner_1 "compute_main_inner_1"
@@ -75,11 +75,11 @@
               OpReturnValue %30
               OpFunctionEnd
 %compute_main_inner = OpFunction %void None %31
-%local_invocation_index = OpFunctionParameter %uint
+%local_invocation_index_2 = OpFunctionParameter %uint
         %35 = OpLabel
        %idx = OpVariable %_ptr_Function_uint Function %6
               OpStore %idx %6
-               OpStore %idx %local_invocation_index
+               OpStore %idx %local_invocation_index_2
               OpBranch %38
         %38 = OpLabel
               OpLoopMerge %39 %40 None
--- a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.wgsl
+++ b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.wgsl
@@ -8,9 +8,9 @@ var<private> local_invocation_index_1 : u32;

 var<workgroup> wg : array<array<array<atomic<u32>, 1u>, 2u>, 3u>;

-fn compute_main_inner(local_invocation_index : u32) {
+fn compute_main_inner(local_invocation_index_2 : u32) {
  var idx : u32 = 0u;
-  idx = local_invocation_index;
+  idx = local_invocation_index_2;
  loop {
    let x_25 : u32 = idx;
    if (!((x_25 < 6u))) {
--- a/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.dxc.hlsl
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.dxc.hlsl
@@ -7,9 +7,9 @@ struct S_atomic {
 static uint local_invocation_index_1 = 0u;
 groupshared S_atomic wg[10];

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  uint idx = 0u;
-  idx = local_invocation_index;
+  idx = local_invocation_index_2;
  while (true) {
    const uint x_23 = idx;
    if (!((x_23 < 10u))) {
--- a/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.fxc.hlsl
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.fxc.hlsl
@@ -7,9 +7,9 @@ struct S_atomic {
 static uint local_invocation_index_1 = 0u;
 groupshared S_atomic wg[10];

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  uint idx = 0u;
-  idx = local_invocation_index;
+  idx = local_invocation_index_2;
  while (true) {
    const uint x_23 = idx;
    if (!((x_23 < 10u))) {
--- a/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.glsl
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.glsl
@@ -14,9 +14,9 @@ struct S {

 uint local_invocation_index_1 = 0u;
 shared S_atomic wg[10];
-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  uint idx = 0u;
-  idx = local_invocation_index;
+  idx = local_invocation_index_2;
  while (true) {
    uint x_23 = idx;
    if (!((x_23 < 10u))) {
--- a/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.msl
@@ -26,9 +26,9 @@ struct S {
  uint y;
 };

-void compute_main_inner(uint local_invocation_index, threadgroup tint_array<S_atomic, 10>* const tint_symbol) {
+void compute_main_inner(uint local_invocation_index_2, threadgroup tint_array<S_atomic, 10>* const tint_symbol) {
  uint idx = 0u;
-  idx = local_invocation_index;
+  idx = local_invocation_index_2;
  while (true) {
    uint const x_23 = idx;
    if (!((x_23 < 10u))) {
--- a/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.spvasm
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.spvasm
@@ -15,7 +15,7 @@
               OpMemberName %S_atomic 2 "y"
               OpName %wg "wg"
               OpName %compute_main_inner "compute_main_inner"
-               OpName %local_invocation_index "local_invocation_index"
+               OpName %local_invocation_index_2 "local_invocation_index_2"
               OpName %idx "idx"
               OpName %compute_main_1 "compute_main_1"
               OpName %compute_main_inner_1 "compute_main_inner_1"
@@ -54,11 +54,11 @@
      %int_4 = OpConstant %int 4
         %51 = OpTypeFunction %void
 %compute_main_inner = OpFunction %void None %13
-%local_invocation_index = OpFunctionParameter %uint
+%local_invocation_index_2 = OpFunctionParameter %uint
         %17 = OpLabel
        %idx = OpVariable %_ptr_Function_uint Function %6
               OpStore %idx %6
-               OpStore %idx %local_invocation_index
+               OpStore %idx %local_invocation_index_2
               OpBranch %20
         %20 = OpLabel
               OpLoopMerge %21 %22 None
--- a/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.wgsl
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.wgsl
@@ -22,9 +22,9 @@ var<private> local_invocation_index_1 : u32;

 var<workgroup> wg : array<S_atomic, 10u>;

-fn compute_main_inner(local_invocation_index : u32) {
+fn compute_main_inner(local_invocation_index_2 : u32) {
  var idx : u32 = 0u;
-  idx = local_invocation_index;
+  idx = local_invocation_index_2;
  loop {
    let x_23 : u32 = idx;
    if (!((x_23 < 10u))) {
--- a/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.spvasm.expected.dxc.hlsl
+++ b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.spvasm.expected.dxc.hlsl
@@ -7,7 +7,7 @@ struct S_atomic {
 static uint local_invocation_index_1 = 0u;
 groupshared S_atomic wg;

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  wg.x = 0;
  uint atomic_result = 0u;
  InterlockedExchange(wg.a, 0u, atomic_result);
--- a/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.spvasm.expected.fxc.hlsl
+++ b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.spvasm.expected.fxc.hlsl
@@ -7,7 +7,7 @@ struct S_atomic {
 static uint local_invocation_index_1 = 0u;
 groupshared S_atomic wg;

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  wg.x = 0;
  uint atomic_result = 0u;
  InterlockedExchange(wg.a, 0u, atomic_result);
--- a/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.spvasm.expected.glsl
+++ b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.spvasm.expected.glsl
@@ -14,7 +14,7 @@ struct S {

 uint local_invocation_index_1 = 0u;
 shared S_atomic wg;
-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  wg.x = 0;
  atomicExchange(wg.a, 0u);
  atomicExchange(wg.b, 0u);
--- a/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.spvasm.expected.msl
@@ -13,7 +13,7 @@ struct S {
  uint b;
 };

-void compute_main_inner(uint local_invocation_index, threadgroup S_atomic* const tint_symbol) {
+void compute_main_inner(uint local_invocation_index_2, threadgroup S_atomic* const tint_symbol) {
  (*(tint_symbol)).x = 0;
  atomic_store_explicit(&((*(tint_symbol)).a), 0u, memory_order_relaxed);
  atomic_store_explicit(&((*(tint_symbol)).b), 0u, memory_order_relaxed);
--- a/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.spvasm.expected.spvasm
+++ b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.spvasm.expected.spvasm
@@ -15,7 +15,7 @@
               OpMemberName %S_atomic 2 "b"
               OpName %wg "wg"
               OpName %compute_main_inner "compute_main_inner"
-               OpName %local_invocation_index "local_invocation_index"
+               OpName %local_invocation_index_2 "local_invocation_index_2"
               OpName %compute_main_1 "compute_main_1"
               OpName %compute_main_inner_1 "compute_main_inner_1"
               OpName %local_invocation_index_1_param "local_invocation_index_1_param"
@@ -45,7 +45,7 @@
   %uint_264 = OpConstant %uint 264
         %37 = OpTypeFunction %void
 %compute_main_inner = OpFunction %void None %11
-%local_invocation_index = OpFunctionParameter %uint
+%local_invocation_index_2 = OpFunctionParameter %uint
         %15 = OpLabel
         %18 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
               OpStore %18 %19
--- a/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.spvasm.expected.wgsl
+++ b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.spvasm.expected.wgsl
@@ -20,7 +20,7 @@ var<private> local_invocation_index_1 : u32;

 var<workgroup> wg : S_atomic;

-fn compute_main_inner(local_invocation_index : u32) {
+fn compute_main_inner(local_invocation_index_2 : u32) {
  wg.x = 0i;
  atomicStore(&(wg.a), 0u);
  atomicStore(&(wg.b), 0u);
--- a/test/tint/builtins/atomicStore/struct/flat_single_atomic.spvasm.expected.dxc.hlsl
+++ b/test/tint/builtins/atomicStore/struct/flat_single_atomic.spvasm.expected.dxc.hlsl
@@ -7,7 +7,7 @@ struct S_atomic {
 static uint local_invocation_index_1 = 0u;
 groupshared S_atomic wg;

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  wg.x = 0;
  uint atomic_result = 0u;
  InterlockedExchange(wg.a, 0u, atomic_result);
--- a/test/tint/builtins/atomicStore/struct/flat_single_atomic.spvasm.expected.fxc.hlsl
+++ b/test/tint/builtins/atomicStore/struct/flat_single_atomic.spvasm.expected.fxc.hlsl
@@ -7,7 +7,7 @@ struct S_atomic {
 static uint local_invocation_index_1 = 0u;
 groupshared S_atomic wg;

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  wg.x = 0;
  uint atomic_result = 0u;
  InterlockedExchange(wg.a, 0u, atomic_result);
--- a/test/tint/builtins/atomicStore/struct/flat_single_atomic.spvasm.expected.glsl
+++ b/test/tint/builtins/atomicStore/struct/flat_single_atomic.spvasm.expected.glsl
@@ -14,7 +14,7 @@ struct S {

 uint local_invocation_index_1 = 0u;
 shared S_atomic wg;
-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  wg.x = 0;
  atomicExchange(wg.a, 0u);
  wg.y = 0u;
--- a/test/tint/builtins/atomicStore/struct/flat_single_atomic.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/struct/flat_single_atomic.spvasm.expected.msl
@@ -13,7 +13,7 @@ struct S {
  uint y;
 };

-void compute_main_inner(uint local_invocation_index, threadgroup S_atomic* const tint_symbol) {
+void compute_main_inner(uint local_invocation_index_2, threadgroup S_atomic* const tint_symbol) {
  (*(tint_symbol)).x = 0;
  atomic_store_explicit(&((*(tint_symbol)).a), 0u, memory_order_relaxed);
  (*(tint_symbol)).y = 0u;
--- a/test/tint/builtins/atomicStore/struct/flat_single_atomic.spvasm.expected.spvasm
+++ b/test/tint/builtins/atomicStore/struct/flat_single_atomic.spvasm.expected.spvasm
@@ -15,7 +15,7 @@
               OpMemberName %S_atomic 2 "y"
               OpName %wg "wg"
               OpName %compute_main_inner "compute_main_inner"
-               OpName %local_invocation_index "local_invocation_index"
+               OpName %local_invocation_index_2 "local_invocation_index_2"
               OpName %compute_main_1 "compute_main_1"
               OpName %compute_main_inner_1 "compute_main_inner_1"
               OpName %local_invocation_index_1_param "local_invocation_index_1_param"
@@ -46,7 +46,7 @@
   %uint_264 = OpConstant %uint 264
         %33 = OpTypeFunction %void
 %compute_main_inner = OpFunction %void None %11
-%local_invocation_index = OpFunctionParameter %uint
+%local_invocation_index_2 = OpFunctionParameter %uint
         %15 = OpLabel
         %18 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
               OpStore %18 %19
--- a/test/tint/builtins/atomicStore/struct/flat_single_atomic.spvasm.expected.wgsl
+++ b/test/tint/builtins/atomicStore/struct/flat_single_atomic.spvasm.expected.wgsl
@@ -20,7 +20,7 @@ var<private> local_invocation_index_1 : u32;

 var<workgroup> wg : S_atomic;

-fn compute_main_inner(local_invocation_index : u32) {
+fn compute_main_inner(local_invocation_index_2 : u32) {
  wg.x = 0i;
  atomicStore(&(wg.a), 0u);
  wg.y = 0u;
--- a/test/tint/builtins/atomicStore/struct/nested.spvasm.expected.dxc.hlsl
+++ b/test/tint/builtins/atomicStore/struct/nested.spvasm.expected.dxc.hlsl
@@ -20,7 +20,7 @@ struct S2_atomic {
 static uint local_invocation_index_1 = 0u;
 groupshared S2_atomic wg;

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  wg.x = 0;
  wg.y = 0;
  wg.z = 0;
--- a/test/tint/builtins/atomicStore/struct/nested.spvasm.expected.fxc.hlsl
+++ b/test/tint/builtins/atomicStore/struct/nested.spvasm.expected.fxc.hlsl
@@ -20,7 +20,7 @@ struct S2_atomic {
 static uint local_invocation_index_1 = 0u;
 groupshared S2_atomic wg;

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  wg.x = 0;
  wg.y = 0;
  wg.z = 0;
--- a/test/tint/builtins/atomicStore/struct/nested.spvasm.expected.glsl
+++ b/test/tint/builtins/atomicStore/struct/nested.spvasm.expected.glsl
@@ -44,7 +44,7 @@ struct S2 {

 uint local_invocation_index_1 = 0u;
 shared S2_atomic wg;
-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  wg.x = 0;
  wg.y = 0;
  wg.z = 0;
--- a/test/tint/builtins/atomicStore/struct/nested.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/struct/nested.spvasm.expected.msl
@@ -43,7 +43,7 @@ struct S2 {
  S1 a;
 };

-void compute_main_inner(uint local_invocation_index, threadgroup S2_atomic* const tint_symbol) {
+void compute_main_inner(uint local_invocation_index_2, threadgroup S2_atomic* const tint_symbol) {
  (*(tint_symbol)).x = 0;
  (*(tint_symbol)).y = 0;
  (*(tint_symbol)).z = 0;
--- a/test/tint/builtins/atomicStore/struct/nested.spvasm.expected.spvasm
+++ b/test/tint/builtins/atomicStore/struct/nested.spvasm.expected.spvasm
@@ -26,7 +26,7 @@
               OpMemberName %S1_atomic 3 "z"
               OpName %wg "wg"
               OpName %compute_main_inner "compute_main_inner"
-               OpName %local_invocation_index "local_invocation_index"
+               OpName %local_invocation_index_2 "local_invocation_index_2"
               OpName %compute_main_1 "compute_main_1"
               OpName %compute_main_inner_1 "compute_main_inner_1"
               OpName %local_invocation_index_1_param "local_invocation_index_1_param"
@@ -68,7 +68,7 @@
   %uint_264 = OpConstant %uint 264
         %42 = OpTypeFunction %void
 %compute_main_inner = OpFunction %void None %13
-%local_invocation_index = OpFunctionParameter %uint
+%local_invocation_index_2 = OpFunctionParameter %uint
         %17 = OpLabel
         %20 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
               OpStore %20 %21
--- a/test/tint/builtins/atomicStore/struct/nested.spvasm.expected.wgsl
+++ b/test/tint/builtins/atomicStore/struct/nested.spvasm.expected.wgsl
@@ -68,7 +68,7 @@ var<private> local_invocation_index_1 : u32;

 var<workgroup> wg : S2_atomic;

-fn compute_main_inner(local_invocation_index : u32) {
+fn compute_main_inner(local_invocation_index_2 : u32) {
  wg.x = 0i;
  wg.y = 0i;
  wg.z = 0i;
--- a/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.dxc.hlsl
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.dxc.hlsl
@@ -7,11 +7,11 @@ struct S_atomic {
 static uint local_invocation_index_1 = 0u;
 groupshared S_atomic wg;

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  uint idx = 0u;
  wg.x = 0;
  wg.y = 0u;
-  idx = local_invocation_index;
+  idx = local_invocation_index_2;
  while (true) {
    const uint x_30 = idx;
    if (!((x_30 < 10u))) {
--- a/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.fxc.hlsl
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.fxc.hlsl
@@ -7,11 +7,11 @@ struct S_atomic {
 static uint local_invocation_index_1 = 0u;
 groupshared S_atomic wg;

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  uint idx = 0u;
  wg.x = 0;
  wg.y = 0u;
-  idx = local_invocation_index;
+  idx = local_invocation_index_2;
  while (true) {
    const uint x_30 = idx;
    if (!((x_30 < 10u))) {
--- a/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.glsl
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.glsl
@@ -14,11 +14,11 @@ struct S {

 uint local_invocation_index_1 = 0u;
 shared S_atomic wg;
-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  uint idx = 0u;
  wg.x = 0;
  wg.y = 0u;
-  idx = local_invocation_index;
+  idx = local_invocation_index_2;
  while (true) {
    uint x_30 = idx;
    if (!((x_30 < 10u))) {
--- a/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.msl
@@ -26,11 +26,11 @@ struct S {
  uint y;
 };

-void compute_main_inner(uint local_invocation_index, threadgroup S_atomic* const tint_symbol) {
+void compute_main_inner(uint local_invocation_index_2, threadgroup S_atomic* const tint_symbol) {
  uint idx = 0u;
  (*(tint_symbol)).x = 0;
  (*(tint_symbol)).y = 0u;
-  idx = local_invocation_index;
+  idx = local_invocation_index_2;
  while (true) {
    uint const x_30 = idx;
    if (!((x_30 < 10u))) {
--- a/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.spvasm
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.spvasm
@@ -15,7 +15,7 @@
               OpMemberName %S_atomic 2 "y"
               OpName %wg "wg"
               OpName %compute_main_inner "compute_main_inner"
-               OpName %local_invocation_index "local_invocation_index"
+               OpName %local_invocation_index_2 "local_invocation_index_2"
               OpName %idx "idx"
               OpName %compute_main_1 "compute_main_1"
               OpName %compute_main_inner_1 "compute_main_inner_1"
@@ -54,7 +54,7 @@
      %int_4 = OpConstant %int 4
         %51 = OpTypeFunction %void
 %compute_main_inner = OpFunction %void None %13
-%local_invocation_index = OpFunctionParameter %uint
+%local_invocation_index_2 = OpFunctionParameter %uint
         %17 = OpLabel
        %idx = OpVariable %_ptr_Function_uint Function %6
               OpStore %idx %6
@@ -62,7 +62,7 @@
               OpStore %22 %23
         %26 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_2
               OpStore %26 %6
-               OpStore %idx %local_invocation_index
+               OpStore %idx %local_invocation_index_2
               OpBranch %27
         %27 = OpLabel
               OpLoopMerge %28 %29 None
--- a/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.wgsl
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.wgsl
@@ -22,11 +22,11 @@ var<private> local_invocation_index_1 : u32;

 var<workgroup> wg : S_atomic;

-fn compute_main_inner(local_invocation_index : u32) {
+fn compute_main_inner(local_invocation_index_2 : u32) {
  var idx : u32 = 0u;
  wg.x = 0i;
  wg.y = 0u;
-  idx = local_invocation_index;
+  idx = local_invocation_index_2;
  loop {
    let x_30 : u32 = idx;
    if (!((x_30 < 10u))) {
--- a/test/tint/builtins/atomicStore/struct/via_ptr_let.spvasm.expected.dxc.hlsl
+++ b/test/tint/builtins/atomicStore/struct/via_ptr_let.spvasm.expected.dxc.hlsl
@@ -7,7 +7,7 @@ struct S_atomic {
 static uint local_invocation_index_1 = 0u;
 groupshared S_atomic wg;

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  wg.x = 0;
  uint atomic_result = 0u;
  InterlockedExchange(wg.a, 0u, atomic_result);
--- a/test/tint/builtins/atomicStore/struct/via_ptr_let.spvasm.expected.fxc.hlsl
+++ b/test/tint/builtins/atomicStore/struct/via_ptr_let.spvasm.expected.fxc.hlsl
@@ -7,7 +7,7 @@ struct S_atomic {
 static uint local_invocation_index_1 = 0u;
 groupshared S_atomic wg;

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  wg.x = 0;
  uint atomic_result = 0u;
  InterlockedExchange(wg.a, 0u, atomic_result);
--- a/test/tint/builtins/atomicStore/struct/via_ptr_let.spvasm.expected.glsl
+++ b/test/tint/builtins/atomicStore/struct/via_ptr_let.spvasm.expected.glsl
@@ -14,7 +14,7 @@ struct S {

 uint local_invocation_index_1 = 0u;
 shared S_atomic wg;
-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  wg.x = 0;
  atomicExchange(wg.a, 0u);
  wg.y = 0u;
--- a/test/tint/builtins/atomicStore/struct/via_ptr_let.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/struct/via_ptr_let.spvasm.expected.msl
@@ -13,7 +13,7 @@ struct S {
  uint y;
 };

-void compute_main_inner(uint local_invocation_index, threadgroup S_atomic* const tint_symbol) {
+void compute_main_inner(uint local_invocation_index_2, threadgroup S_atomic* const tint_symbol) {
  (*(tint_symbol)).x = 0;
  atomic_store_explicit(&((*(tint_symbol)).a), 0u, memory_order_relaxed);
  (*(tint_symbol)).y = 0u;
--- a/test/tint/builtins/atomicStore/struct/via_ptr_let.spvasm.expected.spvasm
+++ b/test/tint/builtins/atomicStore/struct/via_ptr_let.spvasm.expected.spvasm
@@ -15,7 +15,7 @@
               OpMemberName %S_atomic 2 "y"
               OpName %wg "wg"
               OpName %compute_main_inner "compute_main_inner"
-               OpName %local_invocation_index "local_invocation_index"
+               OpName %local_invocation_index_2 "local_invocation_index_2"
               OpName %compute_main_1 "compute_main_1"
               OpName %compute_main_inner_1 "compute_main_inner_1"
               OpName %local_invocation_index_1_param "local_invocation_index_1_param"
@@ -46,7 +46,7 @@
   %uint_264 = OpConstant %uint 264
         %33 = OpTypeFunction %void
 %compute_main_inner = OpFunction %void None %11
-%local_invocation_index = OpFunctionParameter %uint
+%local_invocation_index_2 = OpFunctionParameter %uint
         %15 = OpLabel
         %18 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
               OpStore %18 %19
--- a/test/tint/builtins/atomicStore/struct/via_ptr_let.spvasm.expected.wgsl
+++ b/test/tint/builtins/atomicStore/struct/via_ptr_let.spvasm.expected.wgsl
@@ -20,7 +20,7 @@ var<private> local_invocation_index_1 : u32;

 var<workgroup> wg : S_atomic;

-fn compute_main_inner(local_invocation_index : u32) {
+fn compute_main_inner(local_invocation_index_2 : u32) {
  wg.x = 0i;
  atomicStore(&(wg.a), 0u);
  wg.y = 0u;
--- a/test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_i32.spvasm.expected.dxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_i32.spvasm.expected.dxc.hlsl
@@ -10,7 +10,7 @@ void atomicAdd_794055() {
  return;
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  int atomic_result_1 = 0;
  InterlockedExchange(arg_0, 0, atomic_result_1);
  GroupMemoryBarrierWithGroupSync();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_i32.spvasm.expected.fxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_i32.spvasm.expected.fxc.hlsl
@@ -10,7 +10,7 @@ void atomicAdd_794055() {
  return;
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  int atomic_result_1 = 0;
  InterlockedExchange(arg_0, 0, atomic_result_1);
  GroupMemoryBarrierWithGroupSync();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_i32.spvasm.expected.glsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_i32.spvasm.expected.glsl
@@ -9,7 +9,7 @@ void atomicAdd_794055() {
  return;
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  atomicExchange(arg_0, 0);
  barrier();
  atomicAdd_794055();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_i32.spvasm.expected.msl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_i32.spvasm.expected.msl
@@ -8,7 +8,7 @@ void atomicAdd_794055(threadgroup atomic_int* const tint_symbol) {
  return;
 }

-void compute_main_inner(uint local_invocation_index, threadgroup atomic_int* const tint_symbol_1) {
+void compute_main_inner(uint local_invocation_index_2, threadgroup atomic_int* const tint_symbol_1) {
  atomic_store_explicit(tint_symbol_1, 0, memory_order_relaxed);
  threadgroup_barrier(mem_flags::mem_threadgroup);
  atomicAdd_794055(tint_symbol_1);
--- a/test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_i32.spvasm.expected.spvasm
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_i32.spvasm.expected.spvasm
@@ -13,7 +13,7 @@
               OpName %atomicAdd_794055 "atomicAdd_794055"
               OpName %res "res"
               OpName %compute_main_inner "compute_main_inner"
-               OpName %local_invocation_index "local_invocation_index"
+               OpName %local_invocation_index_2 "local_invocation_index_2"
               OpName %compute_main_1 "compute_main_1"
               OpName %compute_main_inner_1 "compute_main_inner_1"
               OpName %local_invocation_index_1_param "local_invocation_index_1_param"
@@ -46,7 +46,7 @@
               OpReturn
               OpFunctionEnd
 %compute_main_inner = OpFunction %void None %22
-%local_invocation_index = OpFunctionParameter %uint
+%local_invocation_index_2 = OpFunctionParameter %uint
         %25 = OpLabel
               OpAtomicStore %arg_0 %uint_2 %uint_0 %14
               OpControlBarrier %uint_2 %uint_2 %uint_264
--- a/test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_i32.spvasm.expected.wgsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_i32.spvasm.expected.wgsl
@@ -9,7 +9,7 @@ fn atomicAdd_794055() {
  return;
 }

-fn compute_main_inner(local_invocation_index : u32) {
+fn compute_main_inner(local_invocation_index_2 : u32) {
  atomicStore(&(arg_0), 0i);
  workgroupBarrier();
  atomicAdd_794055();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_u32.spvasm.expected.dxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_u32.spvasm.expected.dxc.hlsl
@@ -10,7 +10,7 @@ void atomicAdd_d5db1d() {
  return;
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  uint atomic_result_1 = 0u;
  InterlockedExchange(arg_0, 0u, atomic_result_1);
  GroupMemoryBarrierWithGroupSync();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_u32.spvasm.expected.fxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_u32.spvasm.expected.fxc.hlsl
@@ -10,7 +10,7 @@ void atomicAdd_d5db1d() {
  return;
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  uint atomic_result_1 = 0u;
  InterlockedExchange(arg_0, 0u, atomic_result_1);
  GroupMemoryBarrierWithGroupSync();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_u32.spvasm.expected.glsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_u32.spvasm.expected.glsl
@@ -9,7 +9,7 @@ void atomicAdd_d5db1d() {
  return;
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  atomicExchange(arg_0, 0u);
  barrier();
  atomicAdd_d5db1d();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_u32.spvasm.expected.msl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_u32.spvasm.expected.msl
@@ -8,7 +8,7 @@ void atomicAdd_d5db1d(threadgroup atomic_uint* const tint_symbol) {
  return;
 }

-void compute_main_inner(uint local_invocation_index, threadgroup atomic_uint* const tint_symbol_1) {
+void compute_main_inner(uint local_invocation_index_2, threadgroup atomic_uint* const tint_symbol_1) {
  atomic_store_explicit(tint_symbol_1, 0u, memory_order_relaxed);
  threadgroup_barrier(mem_flags::mem_threadgroup);
  atomicAdd_d5db1d(tint_symbol_1);
--- a/test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_u32.spvasm.expected.spvasm
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_u32.spvasm.expected.spvasm
@@ -13,7 +13,7 @@
               OpName %atomicAdd_d5db1d "atomicAdd_d5db1d"
               OpName %res "res"
               OpName %compute_main_inner "compute_main_inner"
-               OpName %local_invocation_index "local_invocation_index"
+               OpName %local_invocation_index_2 "local_invocation_index_2"
               OpName %compute_main_1 "compute_main_1"
               OpName %compute_main_inner_1 "compute_main_inner_1"
               OpName %local_invocation_index_1_param "local_invocation_index_1_param"
@@ -44,7 +44,7 @@
               OpReturn
               OpFunctionEnd
 %compute_main_inner = OpFunction %void None %20
-%local_invocation_index = OpFunctionParameter %uint
+%local_invocation_index_2 = OpFunctionParameter %uint
         %23 = OpLabel
               OpAtomicStore %arg_0 %uint_2 %uint_0 %6
               OpControlBarrier %uint_2 %uint_2 %uint_264
--- a/test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_u32.spvasm.expected.wgsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_u32.spvasm.expected.wgsl
@@ -9,7 +9,7 @@ fn atomicAdd_d5db1d() {
  return;
 }

-fn compute_main_inner(local_invocation_index : u32) {
+fn compute_main_inner(local_invocation_index_2 : u32) {
  atomicStore(&(arg_0), 0u);
  workgroupBarrier();
  atomicAdd_d5db1d();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicAnd/workgroup_i32.spvasm.expected.dxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicAnd/workgroup_i32.spvasm.expected.dxc.hlsl
@@ -10,7 +10,7 @@ void atomicAnd_45a819() {
  return;
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  int atomic_result_1 = 0;
  InterlockedExchange(arg_0, 0, atomic_result_1);
  GroupMemoryBarrierWithGroupSync();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicAnd/workgroup_i32.spvasm.expected.fxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicAnd/workgroup_i32.spvasm.expected.fxc.hlsl
@@ -10,7 +10,7 @@ void atomicAnd_45a819() {
  return;
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  int atomic_result_1 = 0;
  InterlockedExchange(arg_0, 0, atomic_result_1);
  GroupMemoryBarrierWithGroupSync();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicAnd/workgroup_i32.spvasm.expected.glsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicAnd/workgroup_i32.spvasm.expected.glsl
@@ -9,7 +9,7 @@ void atomicAnd_45a819() {
  return;
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  atomicExchange(arg_0, 0);
  barrier();
  atomicAnd_45a819();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicAnd/workgroup_i32.spvasm.expected.msl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicAnd/workgroup_i32.spvasm.expected.msl
@@ -8,7 +8,7 @@ void atomicAnd_45a819(threadgroup atomic_int* const tint_symbol) {
  return;
 }

-void compute_main_inner(uint local_invocation_index, threadgroup atomic_int* const tint_symbol_1) {
+void compute_main_inner(uint local_invocation_index_2, threadgroup atomic_int* const tint_symbol_1) {
  atomic_store_explicit(tint_symbol_1, 0, memory_order_relaxed);
  threadgroup_barrier(mem_flags::mem_threadgroup);
  atomicAnd_45a819(tint_symbol_1);
--- a/test/tint/builtins/atomics/from_gen/literal/atomicAnd/workgroup_i32.spvasm.expected.spvasm
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicAnd/workgroup_i32.spvasm.expected.spvasm
@@ -13,7 +13,7 @@
               OpName %atomicAnd_45a819 "atomicAnd_45a819"
               OpName %res "res"
               OpName %compute_main_inner "compute_main_inner"
-               OpName %local_invocation_index "local_invocation_index"
+               OpName %local_invocation_index_2 "local_invocation_index_2"
               OpName %compute_main_1 "compute_main_1"
               OpName %compute_main_inner_1 "compute_main_inner_1"
               OpName %local_invocation_index_1_param "local_invocation_index_1_param"
@@ -46,7 +46,7 @@
               OpReturn
               OpFunctionEnd
 %compute_main_inner = OpFunction %void None %22
-%local_invocation_index = OpFunctionParameter %uint
+%local_invocation_index_2 = OpFunctionParameter %uint
         %25 = OpLabel
               OpAtomicStore %arg_0 %uint_2 %uint_0 %14
               OpControlBarrier %uint_2 %uint_2 %uint_264
--- a/test/tint/builtins/atomics/from_gen/literal/atomicAnd/workgroup_i32.spvasm.expected.wgsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicAnd/workgroup_i32.spvasm.expected.wgsl
@@ -9,7 +9,7 @@ fn atomicAnd_45a819() {
  return;
 }

-fn compute_main_inner(local_invocation_index : u32) {
+fn compute_main_inner(local_invocation_index_2 : u32) {
  atomicStore(&(arg_0), 0i);
  workgroupBarrier();
  atomicAnd_45a819();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicAnd/workgroup_u32.spvasm.expected.dxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicAnd/workgroup_u32.spvasm.expected.dxc.hlsl
@@ -10,7 +10,7 @@ void atomicAnd_34edd3() {
  return;
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  uint atomic_result_1 = 0u;
  InterlockedExchange(arg_0, 0u, atomic_result_1);
  GroupMemoryBarrierWithGroupSync();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicAnd/workgroup_u32.spvasm.expected.fxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicAnd/workgroup_u32.spvasm.expected.fxc.hlsl
@@ -10,7 +10,7 @@ void atomicAnd_34edd3() {
  return;
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  uint atomic_result_1 = 0u;
  InterlockedExchange(arg_0, 0u, atomic_result_1);
  GroupMemoryBarrierWithGroupSync();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicAnd/workgroup_u32.spvasm.expected.glsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicAnd/workgroup_u32.spvasm.expected.glsl
@@ -9,7 +9,7 @@ void atomicAnd_34edd3() {
  return;
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  atomicExchange(arg_0, 0u);
  barrier();
  atomicAnd_34edd3();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicAnd/workgroup_u32.spvasm.expected.msl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicAnd/workgroup_u32.spvasm.expected.msl
@@ -8,7 +8,7 @@ void atomicAnd_34edd3(threadgroup atomic_uint* const tint_symbol) {
  return;
 }

-void compute_main_inner(uint local_invocation_index, threadgroup atomic_uint* const tint_symbol_1) {
+void compute_main_inner(uint local_invocation_index_2, threadgroup atomic_uint* const tint_symbol_1) {
  atomic_store_explicit(tint_symbol_1, 0u, memory_order_relaxed);
  threadgroup_barrier(mem_flags::mem_threadgroup);
  atomicAnd_34edd3(tint_symbol_1);
--- a/test/tint/builtins/atomics/from_gen/literal/atomicAnd/workgroup_u32.spvasm.expected.spvasm
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicAnd/workgroup_u32.spvasm.expected.spvasm
@@ -13,7 +13,7 @@
               OpName %atomicAnd_34edd3 "atomicAnd_34edd3"
               OpName %res "res"
               OpName %compute_main_inner "compute_main_inner"
-               OpName %local_invocation_index "local_invocation_index"
+               OpName %local_invocation_index_2 "local_invocation_index_2"
               OpName %compute_main_1 "compute_main_1"
               OpName %compute_main_inner_1 "compute_main_inner_1"
               OpName %local_invocation_index_1_param "local_invocation_index_1_param"
@@ -44,7 +44,7 @@
               OpReturn
               OpFunctionEnd
 %compute_main_inner = OpFunction %void None %20
-%local_invocation_index = OpFunctionParameter %uint
+%local_invocation_index_2 = OpFunctionParameter %uint
         %23 = OpLabel
               OpAtomicStore %arg_0 %uint_2 %uint_0 %6
               OpControlBarrier %uint_2 %uint_2 %uint_264
--- a/test/tint/builtins/atomics/from_gen/literal/atomicAnd/workgroup_u32.spvasm.expected.wgsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicAnd/workgroup_u32.spvasm.expected.wgsl
@@ -9,7 +9,7 @@ fn atomicAnd_34edd3() {
  return;
 }

-fn compute_main_inner(local_invocation_index : u32) {
+fn compute_main_inner(local_invocation_index_2 : u32) {
  atomicStore(&(arg_0), 0u);
  workgroupBarrier();
  atomicAnd_34edd3();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.dxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.dxc.hlsl
@@ -24,7 +24,7 @@ void atomicCompareExchangeWeak_e88938() {
  return;
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  int atomic_result_1 = 0;
  InterlockedExchange(arg_0, 0, atomic_result_1);
  GroupMemoryBarrierWithGroupSync();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.fxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.fxc.hlsl
@@ -24,7 +24,7 @@ void atomicCompareExchangeWeak_e88938() {
  return;
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  int atomic_result_1 = 0;
  InterlockedExchange(arg_0, 0, atomic_result_1);
  GroupMemoryBarrierWithGroupSync();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.glsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.glsl
@@ -26,7 +26,7 @@ void atomicCompareExchangeWeak_e88938() {
  return;
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  atomicExchange(arg_0, 0);
  barrier();
  atomicCompareExchangeWeak_e88938();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.msl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.msl
@@ -27,7 +27,7 @@ void atomicCompareExchangeWeak_e88938(threadgroup atomic_int* const tint_symbol_
  return;
 }

-void compute_main_inner(uint local_invocation_index, threadgroup atomic_int* const tint_symbol_3) {
+void compute_main_inner(uint local_invocation_index_2, threadgroup atomic_int* const tint_symbol_3) {
  atomic_store_explicit(tint_symbol_3, 0, memory_order_relaxed);
  threadgroup_barrier(mem_flags::mem_threadgroup);
  atomicCompareExchangeWeak_e88938(tint_symbol_3);
--- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.spvasm
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.spvasm
@@ -19,7 +19,7 @@
               OpMemberName %__atomic_compare_exchange_resulti32 0 "old_value"
               OpMemberName %__atomic_compare_exchange_resulti32 1 "exchanged"
               OpName %compute_main_inner "compute_main_inner"
-               OpName %local_invocation_index "local_invocation_index"
+               OpName %local_invocation_index_2 "local_invocation_index_2"
               OpName %compute_main_1 "compute_main_1"
               OpName %compute_main_inner_1 "compute_main_inner_1"
               OpName %local_invocation_index_1_param "local_invocation_index_1_param"
@@ -65,7 +65,7 @@
               OpReturn
               OpFunctionEnd
 %compute_main_inner = OpFunction %void None %30
-%local_invocation_index = OpFunctionParameter %uint
+%local_invocation_index_2 = OpFunctionParameter %uint
         %33 = OpLabel
               OpAtomicStore %arg_0 %uint_2 %uint_0 %36
               OpControlBarrier %uint_2 %uint_2 %uint_264
--- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.wgsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.wgsl
@@ -17,7 +17,7 @@ fn atomicCompareExchangeWeak_e88938() {
  return;
 }

-fn compute_main_inner(local_invocation_index : u32) {
+fn compute_main_inner(local_invocation_index_2 : u32) {
  atomicStore(&(arg_0), 0i);
  workgroupBarrier();
  atomicCompareExchangeWeak_e88938();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.dxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.dxc.hlsl
@@ -24,7 +24,7 @@ void atomicCompareExchangeWeak_83580d() {
  return;
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  uint atomic_result_1 = 0u;
  InterlockedExchange(arg_0, 0u, atomic_result_1);
  GroupMemoryBarrierWithGroupSync();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.fxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.fxc.hlsl
@@ -24,7 +24,7 @@ void atomicCompareExchangeWeak_83580d() {
  return;
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  uint atomic_result_1 = 0u;
  InterlockedExchange(arg_0, 0u, atomic_result_1);
  GroupMemoryBarrierWithGroupSync();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.glsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.glsl
@@ -26,7 +26,7 @@ void atomicCompareExchangeWeak_83580d() {
  return;
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  atomicExchange(arg_0, 0u);
  barrier();
  atomicCompareExchangeWeak_83580d();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.msl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.msl
@@ -27,7 +27,7 @@ void atomicCompareExchangeWeak_83580d(threadgroup atomic_uint* const tint_symbol
  return;
 }

-void compute_main_inner(uint local_invocation_index, threadgroup atomic_uint* const tint_symbol_3) {
+void compute_main_inner(uint local_invocation_index_2, threadgroup atomic_uint* const tint_symbol_3) {
  atomic_store_explicit(tint_symbol_3, 0u, memory_order_relaxed);
  threadgroup_barrier(mem_flags::mem_threadgroup);
  atomicCompareExchangeWeak_83580d(tint_symbol_3);
--- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.spvasm
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.spvasm
@@ -19,7 +19,7 @@
               OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value"
               OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged"
               OpName %compute_main_inner "compute_main_inner"
-               OpName %local_invocation_index "local_invocation_index"
+               OpName %local_invocation_index_2 "local_invocation_index_2"
               OpName %compute_main_1 "compute_main_1"
               OpName %compute_main_inner_1 "compute_main_inner_1"
               OpName %local_invocation_index_1_param "local_invocation_index_1_param"
@@ -63,7 +63,7 @@
               OpReturn
               OpFunctionEnd
 %compute_main_inner = OpFunction %void None %29
-%local_invocation_index = OpFunctionParameter %uint
+%local_invocation_index_2 = OpFunctionParameter %uint
         %32 = OpLabel
               OpAtomicStore %arg_0 %uint_2 %uint_0 %6
               OpControlBarrier %uint_2 %uint_2 %uint_264
--- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.wgsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.wgsl
@@ -17,7 +17,7 @@ fn atomicCompareExchangeWeak_83580d() {
  return;
 }

-fn compute_main_inner(local_invocation_index : u32) {
+fn compute_main_inner(local_invocation_index_2 : u32) {
  atomicStore(&(arg_0), 0u);
  workgroupBarrier();
  atomicCompareExchangeWeak_83580d();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicExchange/workgroup_i32.spvasm.expected.dxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicExchange/workgroup_i32.spvasm.expected.dxc.hlsl
@@ -10,7 +10,7 @@ void atomicExchange_e114ba() {
  return;
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  int atomic_result_1 = 0;
  InterlockedExchange(arg_0, 0, atomic_result_1);
  GroupMemoryBarrierWithGroupSync();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicExchange/workgroup_i32.spvasm.expected.fxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicExchange/workgroup_i32.spvasm.expected.fxc.hlsl
@@ -10,7 +10,7 @@ void atomicExchange_e114ba() {
  return;
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  int atomic_result_1 = 0;
  InterlockedExchange(arg_0, 0, atomic_result_1);
  GroupMemoryBarrierWithGroupSync();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicExchange/workgroup_i32.spvasm.expected.glsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicExchange/workgroup_i32.spvasm.expected.glsl
@@ -9,7 +9,7 @@ void atomicExchange_e114ba() {
  return;
 }

-void compute_main_inner(uint local_invocation_index) {
+void compute_main_inner(uint local_invocation_index_2) {
  atomicExchange(arg_0, 0);
  barrier();
  atomicExchange_e114ba();
--- a/test/tint/builtins/atomics/from_gen/literal/atomicExchange/workgroup_i32.spvasm.expected.msl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicExchange/workgroup_i32.spvasm.expected.msl
@@ -8,7 +8,7 @@ void atomicExchange_e114ba(threadgroup atomic_int* const tint_symbol) {
  return;
 }

-void compute_main_inner(uint local_invocation_index, threadgroup atomic_int* const tint_symbol_1) {
+void compute_main_inner(uint local_invocation_index_2, threadgroup atomic_int* const tint_symbol_1) {
  atomic_store_explicit(tint_symbol_1, 0, memory_order_relaxed);
  threadgroup_barrier(mem_flags::mem_threadgroup);
  atomicExchange_e114ba(tint_symbol_1);
--- a/Show More
+++ b/Show More