tint/hlsl+glsl: fix workgroupUniformLoad polyfills

The BuiltinPolyfill transform expects the DirectVariableAccess transform
to run after it, but this regressed as part of
https://dawn-review.googlesource.com/c/dawn/+/122203

Add unit test along with e2e 1926.wgsl test.

Bug: tint:1926
Change-Id: I5107453ce152b12e6f2f36930846e1fffa775708
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/131020
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>

src/tint/transform/builtin_polyfill_test.cc

@@ -16,6 +16,7 @@
 
 #include <utility>
 
+#include "src/tint/transform/direct_variable_access.h"
 #include "src/tint/transform/test_helper.h"
 
 namespace tint::transform {
@@ -3673,8 +3674,23 @@ fn f() {
 DataMap polyfillWorkgroupUniformLoad() {
     BuiltinPolyfill::Builtins builtins;
     builtins.workgroup_uniform_load = true;
+
     DataMap data;
     data.Add<BuiltinPolyfill::Config>(builtins);
+
+    return data;
+}
+
+DataMap polyfillWorkgroupUniformLoadWithDirectVariableAccess() {
+    DataMap data;
+
+    BuiltinPolyfill::Builtins builtins;
+    builtins.workgroup_uniform_load = true;
+    data.Add<BuiltinPolyfill::Config>(builtins);
+
+    DirectVariableAccess::Options options;
+    data.Add<DirectVariableAccess::Config>(options);
+
     return data;
 }
 
@@ -3830,6 +3846,50 @@ fn f() {
     EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(BuiltinPolyfillTest, WorkgroupUniformLoad_DirectVariableAccess) {
+    auto* src = R"(
+var<workgroup> v : i32;
+var<workgroup> v2 : i32;
+
+fn f() {
+  let r = workgroupUniformLoad(&v);
+  let s = workgroupUniformLoad(&v2);
+}
+)";
+
+    auto* expect = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+fn tint_workgroupUniformLoad_v() -> i32 {
+  workgroupBarrier();
+  let result = v;
+  workgroupBarrier();
+  return result;
+}
+
+fn tint_workgroupUniformLoad_v2() -> i32 {
+  workgroupBarrier();
+  let result = v2;
+  workgroupBarrier();
+  return result;
+}
+
+var<workgroup> v : i32;
+
+var<workgroup> v2 : i32;
+
+fn f() {
+  let r = tint_workgroupUniformLoad_v();
+  let s = tint_workgroupUniformLoad_v2();
+}
+)";
+
+    auto got = Run<BuiltinPolyfill, DirectVariableAccess>(
+        src, polyfillWorkgroupUniformLoadWithDirectVariableAccess());
+
+    EXPECT_EQ(expect, str(got));
+}
+
 ////////////////////////////////////////////////////////////////////////////////
 // quantizeToF16
 ////////////////////////////////////////////////////////////////////////////////

src/tint/writer/glsl/generator_impl.cc

@@ -170,7 +170,6 @@ SanitizedResult Sanitize(const Program* in,
     manager.Add<transform::PreservePadding>();  // Must come before DirectVariableAccess
 
     manager.Add<transform::Unshadow>();  // Must come before DirectVariableAccess
-    manager.Add<transform::DirectVariableAccess>();
 
     manager.Add<transform::PromoteSideEffectsToDecl>();
 
@@ -203,9 +202,11 @@ SanitizedResult Sanitize(const Program* in,
         polyfills.texture_sample_base_clamp_to_edge_2d_f32 = true;
         polyfills.workgroup_uniform_load = true;
         data.Add<transform::BuiltinPolyfill::Config>(polyfills);
-        manager.Add<transform::BuiltinPolyfill>();
+        manager.Add<transform::BuiltinPolyfill>();  // Must come before DirectVariableAccess
     }
 
+    manager.Add<transform::DirectVariableAccess>();
+
     if (!options.disable_workgroup_init) {
         // ZeroInitWorkgroupMemory must come before CanonicalizeEntryPointIO as
         // ZeroInitWorkgroupMemory may inject new builtin parameters.

src/tint/writer/hlsl/generator_impl.cc

@@ -176,8 +176,6 @@ SanitizedResult Sanitize(const Program* in, const Options& options) {
 
     manager.Add<transform::Unshadow>();  // Must come before DirectVariableAccess
 
-    manager.Add<transform::DirectVariableAccess>();
-
     // LocalizeStructArrayAssignment must come after:
     // * SimplifyPointers, because it assumes assignment to arrays in structs are
     // done directly, not indirectly.
@@ -229,9 +227,11 @@ SanitizedResult Sanitize(const Program* in, const Options& options) {
         polyfills.texture_sample_base_clamp_to_edge_2d_f32 = true;
         polyfills.workgroup_uniform_load = true;
         data.Add<transform::BuiltinPolyfill::Config>(polyfills);
-        manager.Add<transform::BuiltinPolyfill>();
+        manager.Add<transform::BuiltinPolyfill>();  // Must come before DirectVariableAccess
     }
 
+    manager.Add<transform::DirectVariableAccess>();
+
     if (!options.disable_workgroup_init) {
         // ZeroInitWorkgroupMemory must come before CanonicalizeEntryPointIO as
         // ZeroInitWorkgroupMemory may inject new builtin parameters.

src/tint/writer/spirv/generator_impl.cc

@@ -119,7 +119,7 @@ SanitizedResult Sanitize(const Program* in, const Options& options) {
         polyfills.quantize_to_vec_f16 = true;  // crbug.com/tint/1741
         polyfills.workgroup_uniform_load = true;
         data.Add<transform::BuiltinPolyfill::Config>(polyfills);
-        manager.Add<transform::BuiltinPolyfill>();
+        manager.Add<transform::BuiltinPolyfill>();  // Must come before DirectVariableAccess
     }
 
     bool disable_workgroup_init_in_sanitizer =

test/tint/bug/tint/1926.wgsl

@@ -0,0 +1,16 @@
+var<workgroup> sh_atomic_failed: u32;
+
+@group(0) @binding(4)
+var<storage, read_write> output: u32;
+
+@compute @workgroup_size(256)
+fn main(
+    @builtin(global_invocation_id) global_id: vec3<u32>,
+    @builtin(local_invocation_id) local_id: vec3<u32>,
+) {
+    let failed = workgroupUniformLoad(&sh_atomic_failed);
+    
+    if (local_id.x == 0) {
+        output = failed;
+    }
+}

test/tint/bug/tint/1926.wgsl.expected.dxc.hlsl

@@ -0,0 +1,33 @@
+groupshared uint sh_atomic_failed;
+
+uint tint_workgroupUniformLoad_sh_atomic_failed() {
+  GroupMemoryBarrierWithGroupSync();
+  const uint result = sh_atomic_failed;
+  GroupMemoryBarrierWithGroupSync();
+  return result;
+}
+
+RWByteAddressBuffer output : register(u4);
+
+struct tint_symbol_1 {
+  uint3 local_id : SV_GroupThreadID;
+  uint local_invocation_index : SV_GroupIndex;
+  uint3 global_id : SV_DispatchThreadID;
+};
+
+void main_inner(uint3 global_id, uint3 local_id, uint local_invocation_index) {
+  if ((local_invocation_index < 1u)) {
+    sh_atomic_failed = 0u;
+  }
+  GroupMemoryBarrierWithGroupSync();
+  const uint failed = tint_workgroupUniformLoad_sh_atomic_failed();
+  if ((local_id.x == 0u)) {
+    output.Store(0u, asuint(failed));
+  }
+}
+
+[numthreads(256, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.global_id, tint_symbol.local_id, tint_symbol.local_invocation_index);
+  return;
+}

test/tint/bug/tint/1926.wgsl.expected.fxc.hlsl

@@ -0,0 +1,33 @@
+groupshared uint sh_atomic_failed;
+
+uint tint_workgroupUniformLoad_sh_atomic_failed() {
+  GroupMemoryBarrierWithGroupSync();
+  const uint result = sh_atomic_failed;
+  GroupMemoryBarrierWithGroupSync();
+  return result;
+}
+
+RWByteAddressBuffer output : register(u4);
+
+struct tint_symbol_1 {
+  uint3 local_id : SV_GroupThreadID;
+  uint local_invocation_index : SV_GroupIndex;
+  uint3 global_id : SV_DispatchThreadID;
+};
+
+void main_inner(uint3 global_id, uint3 local_id, uint local_invocation_index) {
+  if ((local_invocation_index < 1u)) {
+    sh_atomic_failed = 0u;
+  }
+  GroupMemoryBarrierWithGroupSync();
+  const uint failed = tint_workgroupUniformLoad_sh_atomic_failed();
+  if ((local_id.x == 0u)) {
+    output.Store(0u, asuint(failed));
+  }
+}
+
+[numthreads(256, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.global_id, tint_symbol.local_id, tint_symbol.local_invocation_index);
+  return;
+}

test/tint/bug/tint/1926.wgsl.expected.glsl

@@ -0,0 +1,30 @@
+#version 310 es
+
+shared uint sh_atomic_failed;
+uint tint_workgroupUniformLoad_sh_atomic_failed() {
+  barrier();
+  uint result = sh_atomic_failed;
+  barrier();
+  return result;
+}
+
+layout(binding = 4, std430) buffer tint_symbol_block_ssbo {
+  uint inner;
+} tint_symbol;
+
+void tint_symbol_1(uvec3 global_id, uvec3 local_id, uint local_invocation_index) {
+  if ((local_invocation_index < 1u)) {
+    sh_atomic_failed = 0u;
+  }
+  barrier();
+  uint failed = tint_workgroupUniformLoad_sh_atomic_failed();
+  if ((local_id.x == 0u)) {
+    tint_symbol.inner = failed;
+  }
+}
+
+layout(local_size_x = 256, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol_1(gl_GlobalInvocationID, gl_LocalInvocationID, gl_LocalInvocationIndex);
+  return;
+}

test/tint/bug/tint/1926.wgsl.expected.msl

@@ -0,0 +1,27 @@
+#include <metal_stdlib>
+
+using namespace metal;
+uint tint_workgroupUniformLoad(threadgroup uint* const p) {
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  uint const result = *(p);
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  return result;
+}
+
+void tint_symbol_inner(uint3 global_id, uint3 local_id, uint local_invocation_index, threadgroup uint* const tint_symbol_1, device uint* const tint_symbol_2) {
+  if ((local_invocation_index < 1u)) {
+    *(tint_symbol_1) = 0u;
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  uint const failed = tint_workgroupUniformLoad(tint_symbol_1);
+  if ((local_id[0] == 0u)) {
+    *(tint_symbol_2) = failed;
+  }
+}
+
+kernel void tint_symbol(device uint* tint_symbol_4 [[buffer(0)]], uint3 global_id [[thread_position_in_grid]], uint3 local_id [[thread_position_in_threadgroup]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup uint tint_symbol_3;
+  tint_symbol_inner(global_id, local_id, local_invocation_index, &(tint_symbol_3), tint_symbol_4);
+  return;
+}
+

test/tint/bug/tint/1926.wgsl.expected.spvasm

@@ -0,0 +1,92 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 50
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main" %global_id_1 %local_id_1 %local_invocation_index_1
+               OpExecutionMode %main LocalSize 256 1 1
+               OpName %global_id_1 "global_id_1"
+               OpName %local_id_1 "local_id_1"
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %sh_atomic_failed "sh_atomic_failed"
+               OpName %output_block "output_block"
+               OpMemberName %output_block 0 "inner"
+               OpName %output "output"
+               OpName %tint_workgroupUniformLoad_sh_atomic_failed "tint_workgroupUniformLoad_sh_atomic_failed"
+               OpName %main_inner "main_inner"
+               OpName %global_id "global_id"
+               OpName %local_id "local_id"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %main "main"
+               OpDecorate %global_id_1 BuiltIn GlobalInvocationId
+               OpDecorate %local_id_1 BuiltIn LocalInvocationId
+               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+               OpDecorate %output_block Block
+               OpMemberDecorate %output_block 0 Offset 0
+               OpDecorate %output DescriptorSet 0
+               OpDecorate %output Binding 4
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%global_id_1 = OpVariable %_ptr_Input_v3uint Input
+ %local_id_1 = OpVariable %_ptr_Input_v3uint Input
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+%sh_atomic_failed = OpVariable %_ptr_Workgroup_uint Workgroup
+%output_block = OpTypeStruct %uint
+%_ptr_StorageBuffer_output_block = OpTypePointer StorageBuffer %output_block
+     %output = OpVariable %_ptr_StorageBuffer_output_block StorageBuffer
+         %13 = OpTypeFunction %uint
+       %void = OpTypeVoid
+     %uint_2 = OpConstant %uint 2
+   %uint_264 = OpConstant %uint 264
+         %22 = OpTypeFunction %void %v3uint %v3uint %uint
+     %uint_1 = OpConstant %uint 1
+       %bool = OpTypeBool
+         %33 = OpConstantNull %uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+         %43 = OpTypeFunction %void
+%tint_workgroupUniformLoad_sh_atomic_failed = OpFunction %uint None %13
+         %15 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %20 = OpLoad %uint %sh_atomic_failed
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+               OpReturnValue %20
+               OpFunctionEnd
+ %main_inner = OpFunction %void None %22
+  %global_id = OpFunctionParameter %v3uint
+   %local_id = OpFunctionParameter %v3uint
+%local_invocation_index = OpFunctionParameter %uint
+         %27 = OpLabel
+         %29 = OpULessThan %bool %local_invocation_index %uint_1
+               OpSelectionMerge %31 None
+               OpBranchConditional %29 %32 %31
+         %32 = OpLabel
+               OpStore %sh_atomic_failed %33
+               OpBranch %31
+         %31 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %35 = OpFunctionCall %uint %tint_workgroupUniformLoad_sh_atomic_failed
+         %36 = OpCompositeExtract %uint %local_id 0
+         %37 = OpIEqual %bool %36 %33
+               OpSelectionMerge %38 None
+               OpBranchConditional %37 %39 %38
+         %39 = OpLabel
+         %42 = OpAccessChain %_ptr_StorageBuffer_uint %output %uint_0
+               OpStore %42 %35
+               OpBranch %38
+         %38 = OpLabel
+               OpReturn
+               OpFunctionEnd
+       %main = OpFunction %void None %43
+         %45 = OpLabel
+         %47 = OpLoad %v3uint %global_id_1
+         %48 = OpLoad %v3uint %local_id_1
+         %49 = OpLoad %uint %local_invocation_index_1
+         %46 = OpFunctionCall %void %main_inner %47 %48 %49
+               OpReturn
+               OpFunctionEnd

test/tint/bug/tint/1926.wgsl.expected.wgsl

@@ -0,0 +1,11 @@
+var<workgroup> sh_atomic_failed : u32;
+
+@group(0) @binding(4) var<storage, read_write> output : u32;
+
+@compute @workgroup_size(256)
+fn main(@builtin(global_invocation_id) global_id : vec3<u32>, @builtin(local_invocation_id) local_id : vec3<u32>) {
+  let failed = workgroupUniformLoad(&(sh_atomic_failed));
+  if ((local_id.x == 0)) {
+    output = failed;
+  }
+}

test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.dxc.hlsl

@@ -1,15 +1,16 @@
-uint tint_workgroupUniformLoad(inout uint p) {
+groupshared uint arg_0;
+
+uint tint_workgroupUniformLoad_arg_0() {
   GroupMemoryBarrierWithGroupSync();
-  const uint result = p;
+  const uint result = arg_0;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared uint arg_0;
 RWByteAddressBuffer prevent_dce : register(u0, space2);
 
 void workgroupUniformLoad_37307c() {
-  uint res = tint_workgroupUniformLoad(arg_0);
+  uint res = tint_workgroupUniformLoad_arg_0();
   prevent_dce.Store(0u, asuint(res));
 }
 

test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.fxc.hlsl

@@ -1,15 +1,16 @@
-uint tint_workgroupUniformLoad(inout uint p) {
+groupshared uint arg_0;
+
+uint tint_workgroupUniformLoad_arg_0() {
   GroupMemoryBarrierWithGroupSync();
-  const uint result = p;
+  const uint result = arg_0;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared uint arg_0;
 RWByteAddressBuffer prevent_dce : register(u0, space2);
 
 void workgroupUniformLoad_37307c() {
-  uint res = tint_workgroupUniformLoad(arg_0);
+  uint res = tint_workgroupUniformLoad_arg_0();
   prevent_dce.Store(0u, asuint(res));
 }
 

test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.glsl

@@ -1,19 +1,19 @@
 #version 310 es
 
-uint tint_workgroupUniformLoad(inout uint p) {
+shared uint arg_0;
+uint tint_workgroupUniformLoad_arg_0() {
   barrier();
-  uint result = p;
+  uint result = arg_0;
   barrier();
   return result;
 }
 
-shared uint arg_0;
 layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
   uint inner;
 } prevent_dce;
 
 void workgroupUniformLoad_37307c() {
-  uint res = tint_workgroupUniformLoad(arg_0);
+  uint res = tint_workgroupUniformLoad_arg_0();
   prevent_dce.inner = res;
 }
 

test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.dxc.hlsl

@@ -1,15 +1,16 @@
-float tint_workgroupUniformLoad(inout float p) {
+groupshared float arg_0;
+
+float tint_workgroupUniformLoad_arg_0() {
   GroupMemoryBarrierWithGroupSync();
-  const float result = p;
+  const float result = arg_0;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared float arg_0;
 RWByteAddressBuffer prevent_dce : register(u0, space2);
 
 void workgroupUniformLoad_7a857c() {
-  float res = tint_workgroupUniformLoad(arg_0);
+  float res = tint_workgroupUniformLoad_arg_0();
   prevent_dce.Store(0u, asuint(res));
 }
 

test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.fxc.hlsl

@@ -1,15 +1,16 @@
-float tint_workgroupUniformLoad(inout float p) {
+groupshared float arg_0;
+
+float tint_workgroupUniformLoad_arg_0() {
   GroupMemoryBarrierWithGroupSync();
-  const float result = p;
+  const float result = arg_0;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared float arg_0;
 RWByteAddressBuffer prevent_dce : register(u0, space2);
 
 void workgroupUniformLoad_7a857c() {
-  float res = tint_workgroupUniformLoad(arg_0);
+  float res = tint_workgroupUniformLoad_arg_0();
   prevent_dce.Store(0u, asuint(res));
 }
 

test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.glsl

@@ -1,19 +1,19 @@
 #version 310 es
 
-float tint_workgroupUniformLoad(inout float p) {
+shared float arg_0;
+float tint_workgroupUniformLoad_arg_0() {
   barrier();
-  float result = p;
+  float result = arg_0;
   barrier();
   return result;
 }
 
-shared float arg_0;
 layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
   float inner;
 } prevent_dce;
 
 void workgroupUniformLoad_7a857c() {
-  float res = tint_workgroupUniformLoad(arg_0);
+  float res = tint_workgroupUniformLoad_arg_0();
   prevent_dce.inner = res;
 }
 

test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.dxc.hlsl

@@ -1,15 +1,16 @@
-int tint_workgroupUniformLoad(inout int p) {
+groupshared int arg_0;
+
+int tint_workgroupUniformLoad_arg_0() {
   GroupMemoryBarrierWithGroupSync();
-  const int result = p;
+  const int result = arg_0;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared int arg_0;
 RWByteAddressBuffer prevent_dce : register(u0, space2);
 
 void workgroupUniformLoad_9d33de() {
-  int res = tint_workgroupUniformLoad(arg_0);
+  int res = tint_workgroupUniformLoad_arg_0();
   prevent_dce.Store(0u, asuint(res));
 }
 

test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.fxc.hlsl

@@ -1,15 +1,16 @@
-int tint_workgroupUniformLoad(inout int p) {
+groupshared int arg_0;
+
+int tint_workgroupUniformLoad_arg_0() {
   GroupMemoryBarrierWithGroupSync();
-  const int result = p;
+  const int result = arg_0;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared int arg_0;
 RWByteAddressBuffer prevent_dce : register(u0, space2);
 
 void workgroupUniformLoad_9d33de() {
-  int res = tint_workgroupUniformLoad(arg_0);
+  int res = tint_workgroupUniformLoad_arg_0();
   prevent_dce.Store(0u, asuint(res));
 }
 

test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.glsl

@@ -1,19 +1,19 @@
 #version 310 es
 
-int tint_workgroupUniformLoad(inout int p) {
+shared int arg_0;
+int tint_workgroupUniformLoad_arg_0() {
   barrier();
-  int result = p;
+  int result = arg_0;
   barrier();
   return result;
 }
 
-shared int arg_0;
 layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
   int inner;
 } prevent_dce;
 
 void workgroupUniformLoad_9d33de() {
-  int res = tint_workgroupUniformLoad(arg_0);
+  int res = tint_workgroupUniformLoad_arg_0();
   prevent_dce.inner = res;
 }
 

test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.dxc.hlsl

@@ -1,15 +1,16 @@
-float16_t tint_workgroupUniformLoad(inout float16_t p) {
+groupshared float16_t arg_0;
+
+float16_t tint_workgroupUniformLoad_arg_0() {
   GroupMemoryBarrierWithGroupSync();
-  const float16_t result = p;
+  const float16_t result = arg_0;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared float16_t arg_0;
 RWByteAddressBuffer prevent_dce : register(u0, space2);
 
 void workgroupUniformLoad_e07d08() {
-  float16_t res = tint_workgroupUniformLoad(arg_0);
+  float16_t res = tint_workgroupUniformLoad_arg_0();
   prevent_dce.Store<float16_t>(0u, res);
 }
 

test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.glsl

@@ -1,20 +1,20 @@
 #version 310 es
 #extension GL_AMD_gpu_shader_half_float : require
 
-float16_t tint_workgroupUniformLoad(inout float16_t p) {
+shared float16_t arg_0;
+float16_t tint_workgroupUniformLoad_arg_0() {
   barrier();
-  float16_t result = p;
+  float16_t result = arg_0;
   barrier();
   return result;
 }
 
-shared float16_t arg_0;
 layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
   float16_t inner;
 } prevent_dce;
 
 void workgroupUniformLoad_e07d08() {
-  float16_t res = tint_workgroupUniformLoad(arg_0);
+  float16_t res = tint_workgroupUniformLoad_arg_0();
   prevent_dce.inner = res;
 }
 

test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.dxc.hlsl

@@ -1,15 +1,16 @@
-uint tint_workgroupUniformLoad(inout uint p) {
+groupshared uint arg_0;
+
+uint tint_workgroupUniformLoad_arg_0() {
   GroupMemoryBarrierWithGroupSync();
-  const uint result = p;
+  const uint result = arg_0;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared uint arg_0;
 RWByteAddressBuffer prevent_dce : register(u0, space2);
 
 void workgroupUniformLoad_37307c() {
-  uint res = tint_workgroupUniformLoad(arg_0);
+  uint res = tint_workgroupUniformLoad_arg_0();
   prevent_dce.Store(0u, asuint(res));
 }
 

test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.fxc.hlsl

@@ -1,15 +1,16 @@
-uint tint_workgroupUniformLoad(inout uint p) {
+groupshared uint arg_0;
+
+uint tint_workgroupUniformLoad_arg_0() {
   GroupMemoryBarrierWithGroupSync();
-  const uint result = p;
+  const uint result = arg_0;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared uint arg_0;
 RWByteAddressBuffer prevent_dce : register(u0, space2);
 
 void workgroupUniformLoad_37307c() {
-  uint res = tint_workgroupUniformLoad(arg_0);
+  uint res = tint_workgroupUniformLoad_arg_0();
   prevent_dce.Store(0u, asuint(res));
 }
 

test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.glsl

@@ -1,19 +1,19 @@
 #version 310 es
 
-uint tint_workgroupUniformLoad(inout uint p) {
+shared uint arg_0;
+uint tint_workgroupUniformLoad_arg_0() {
   barrier();
-  uint result = p;
+  uint result = arg_0;
   barrier();
   return result;
 }
 
-shared uint arg_0;
 layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
   uint inner;
 } prevent_dce;
 
 void workgroupUniformLoad_37307c() {
-  uint res = tint_workgroupUniformLoad(arg_0);
+  uint res = tint_workgroupUniformLoad_arg_0();
   prevent_dce.inner = res;
 }
 

test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.dxc.hlsl

@@ -1,15 +1,16 @@
-float tint_workgroupUniformLoad(inout float p) {
+groupshared float arg_0;
+
+float tint_workgroupUniformLoad_arg_0() {
   GroupMemoryBarrierWithGroupSync();
-  const float result = p;
+  const float result = arg_0;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared float arg_0;
 RWByteAddressBuffer prevent_dce : register(u0, space2);
 
 void workgroupUniformLoad_7a857c() {
-  float res = tint_workgroupUniformLoad(arg_0);
+  float res = tint_workgroupUniformLoad_arg_0();
   prevent_dce.Store(0u, asuint(res));
 }
 

test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.fxc.hlsl

@@ -1,15 +1,16 @@
-float tint_workgroupUniformLoad(inout float p) {
+groupshared float arg_0;
+
+float tint_workgroupUniformLoad_arg_0() {
   GroupMemoryBarrierWithGroupSync();
-  const float result = p;
+  const float result = arg_0;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared float arg_0;
 RWByteAddressBuffer prevent_dce : register(u0, space2);
 
 void workgroupUniformLoad_7a857c() {
-  float res = tint_workgroupUniformLoad(arg_0);
+  float res = tint_workgroupUniformLoad_arg_0();
   prevent_dce.Store(0u, asuint(res));
 }
 

test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.glsl

@@ -1,19 +1,19 @@
 #version 310 es
 
-float tint_workgroupUniformLoad(inout float p) {
+shared float arg_0;
+float tint_workgroupUniformLoad_arg_0() {
   barrier();
-  float result = p;
+  float result = arg_0;
   barrier();
   return result;
 }
 
-shared float arg_0;
 layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
   float inner;
 } prevent_dce;
 
 void workgroupUniformLoad_7a857c() {
-  float res = tint_workgroupUniformLoad(arg_0);
+  float res = tint_workgroupUniformLoad_arg_0();
   prevent_dce.inner = res;
 }
 

test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.dxc.hlsl

@@ -1,15 +1,16 @@
-int tint_workgroupUniformLoad(inout int p) {
+groupshared int arg_0;
+
+int tint_workgroupUniformLoad_arg_0() {
   GroupMemoryBarrierWithGroupSync();
-  const int result = p;
+  const int result = arg_0;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared int arg_0;
 RWByteAddressBuffer prevent_dce : register(u0, space2);
 
 void workgroupUniformLoad_9d33de() {
-  int res = tint_workgroupUniformLoad(arg_0);
+  int res = tint_workgroupUniformLoad_arg_0();
   prevent_dce.Store(0u, asuint(res));
 }
 

test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.fxc.hlsl

@@ -1,15 +1,16 @@
-int tint_workgroupUniformLoad(inout int p) {
+groupshared int arg_0;
+
+int tint_workgroupUniformLoad_arg_0() {
   GroupMemoryBarrierWithGroupSync();
-  const int result = p;
+  const int result = arg_0;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared int arg_0;
 RWByteAddressBuffer prevent_dce : register(u0, space2);
 
 void workgroupUniformLoad_9d33de() {
-  int res = tint_workgroupUniformLoad(arg_0);
+  int res = tint_workgroupUniformLoad_arg_0();
   prevent_dce.Store(0u, asuint(res));
 }
 

test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.glsl

@@ -1,19 +1,19 @@
 #version 310 es
 
-int tint_workgroupUniformLoad(inout int p) {
+shared int arg_0;
+int tint_workgroupUniformLoad_arg_0() {
   barrier();
-  int result = p;
+  int result = arg_0;
   barrier();
   return result;
 }
 
-shared int arg_0;
 layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
   int inner;
 } prevent_dce;
 
 void workgroupUniformLoad_9d33de() {
-  int res = tint_workgroupUniformLoad(arg_0);
+  int res = tint_workgroupUniformLoad_arg_0();
   prevent_dce.inner = res;
 }
 

test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl.expected.dxc.hlsl

@@ -1,15 +1,16 @@
-float16_t tint_workgroupUniformLoad(inout float16_t p) {
+groupshared float16_t arg_0;
+
+float16_t tint_workgroupUniformLoad_arg_0() {
   GroupMemoryBarrierWithGroupSync();
-  const float16_t result = p;
+  const float16_t result = arg_0;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared float16_t arg_0;
 RWByteAddressBuffer prevent_dce : register(u0, space2);
 
 void workgroupUniformLoad_e07d08() {
-  float16_t res = tint_workgroupUniformLoad(arg_0);
+  float16_t res = tint_workgroupUniformLoad_arg_0();
   prevent_dce.Store<float16_t>(0u, res);
 }
 

test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl.expected.glsl

@@ -1,20 +1,20 @@
 #version 310 es
 #extension GL_AMD_gpu_shader_half_float : require
 
-float16_t tint_workgroupUniformLoad(inout float16_t p) {
+shared float16_t arg_0;
+float16_t tint_workgroupUniformLoad_arg_0() {
   barrier();
-  float16_t result = p;
+  float16_t result = arg_0;
   barrier();
   return result;
 }
 
-shared float16_t arg_0;
 layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
   float16_t inner;
 } prevent_dce;
 
 void workgroupUniformLoad_e07d08() {
-  float16_t res = tint_workgroupUniformLoad(arg_0);
+  float16_t res = tint_workgroupUniformLoad_arg_0();
   prevent_dce.inner = res;
 }
 

test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.dxc.hlsl

@@ -3,17 +3,17 @@ void unused_entry_point() {
   return;
 }
 
-typedef int tint_workgroupUniformLoad_ret[4];
+groupshared int v[4];
-tint_workgroupUniformLoad_ret tint_workgroupUniformLoad(inout int p[4]) {
+
+typedef int tint_workgroupUniformLoad_v_ret[4];
+tint_workgroupUniformLoad_v_ret tint_workgroupUniformLoad_v() {
   GroupMemoryBarrierWithGroupSync();
-  const int result[4] = p;
+  const int result[4] = v;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared int v[4];
-
 typedef int foo_ret[4];
 foo_ret foo() {
-  return tint_workgroupUniformLoad(v);
+  return tint_workgroupUniformLoad_v();
 }

test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.fxc.hlsl

@@ -3,17 +3,17 @@ void unused_entry_point() {
   return;
 }
 
-typedef int tint_workgroupUniformLoad_ret[4];
+groupshared int v[4];
-tint_workgroupUniformLoad_ret tint_workgroupUniformLoad(inout int p[4]) {
+
+typedef int tint_workgroupUniformLoad_v_ret[4];
+tint_workgroupUniformLoad_v_ret tint_workgroupUniformLoad_v() {
   GroupMemoryBarrierWithGroupSync();
-  const int result[4] = p;
+  const int result[4] = v;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared int v[4];
-
 typedef int foo_ret[4];
 foo_ret foo() {
-  return tint_workgroupUniformLoad(v);
+  return tint_workgroupUniformLoad_v();
 }

test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.glsl

@@ -4,15 +4,15 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 void unused_entry_point() {
   return;
 }
-int[4] tint_workgroupUniformLoad(inout int p[4]) {
+shared int v[4];
+int[4] tint_workgroupUniformLoad_v() {
   barrier();
-  int result[4] = p;
+  int result[4] = v;
   barrier();
   return result;
 }
 
-shared int v[4];
 int[4] foo() {
-  return tint_workgroupUniformLoad(v);
+  return tint_workgroupUniformLoad_v();
 }
 

test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.dxc.hlsl

@@ -3,17 +3,17 @@ void unused_entry_point() {
   return;
 }
 
-typedef int tint_workgroupUniformLoad_ret[128];
+groupshared int v[128];
-tint_workgroupUniformLoad_ret tint_workgroupUniformLoad(inout int p[128]) {
+
+typedef int tint_workgroupUniformLoad_v_ret[128];
+tint_workgroupUniformLoad_v_ret tint_workgroupUniformLoad_v() {
   GroupMemoryBarrierWithGroupSync();
-  const int result[128] = p;
+  const int result[128] = v;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared int v[128];
-
 int foo() {
-  const int tint_symbol[128] = tint_workgroupUniformLoad(v);
+  const int tint_symbol[128] = tint_workgroupUniformLoad_v();
   return tint_symbol[0];
 }

test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.fxc.hlsl

@@ -3,17 +3,17 @@ void unused_entry_point() {
   return;
 }
 
-typedef int tint_workgroupUniformLoad_ret[128];
+groupshared int v[128];
-tint_workgroupUniformLoad_ret tint_workgroupUniformLoad(inout int p[128]) {
+
+typedef int tint_workgroupUniformLoad_v_ret[128];
+tint_workgroupUniformLoad_v_ret tint_workgroupUniformLoad_v() {
   GroupMemoryBarrierWithGroupSync();
-  const int result[128] = p;
+  const int result[128] = v;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared int v[128];
-
 int foo() {
-  const int tint_symbol[128] = tint_workgroupUniformLoad(v);
+  const int tint_symbol[128] = tint_workgroupUniformLoad_v();
   return tint_symbol[0];
 }

test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.glsl

@@ -4,16 +4,16 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 void unused_entry_point() {
   return;
 }
-int[128] tint_workgroupUniformLoad(inout int p[128]) {
+shared int v[128];
+int[128] tint_workgroupUniformLoad_v() {
   barrier();
-  int result[128] = p;
+  int result[128] = v;
   barrier();
   return result;
 }
 
-shared int v[128];
 int foo() {
-  int tint_symbol[128] = tint_workgroupUniformLoad(v);
+  int tint_symbol[128] = tint_workgroupUniformLoad_v();
   return tint_symbol[0];
 }
 

test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.dxc.hlsl

@@ -3,17 +3,17 @@ void unused_entry_point() {
   return;
 }
 
-typedef int tint_workgroupUniformLoad_ret[128];
+groupshared int v[128];
-tint_workgroupUniformLoad_ret tint_workgroupUniformLoad(inout int p[128]) {
+
+typedef int tint_workgroupUniformLoad_v_ret[128];
+tint_workgroupUniformLoad_v_ret tint_workgroupUniformLoad_v() {
   GroupMemoryBarrierWithGroupSync();
-  const int result[128] = p;
+  const int result[128] = v;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared int v[128];
-
 int foo() {
-  const int tint_symbol[128] = tint_workgroupUniformLoad(v);
+  const int tint_symbol[128] = tint_workgroupUniformLoad_v();
   return tint_symbol[0];
 }

test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.fxc.hlsl

@@ -3,17 +3,17 @@ void unused_entry_point() {
   return;
 }
 
-typedef int tint_workgroupUniformLoad_ret[128];
+groupshared int v[128];
-tint_workgroupUniformLoad_ret tint_workgroupUniformLoad(inout int p[128]) {
+
+typedef int tint_workgroupUniformLoad_v_ret[128];
+tint_workgroupUniformLoad_v_ret tint_workgroupUniformLoad_v() {
   GroupMemoryBarrierWithGroupSync();
-  const int result[128] = p;
+  const int result[128] = v;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared int v[128];
-
 int foo() {
-  const int tint_symbol[128] = tint_workgroupUniformLoad(v);
+  const int tint_symbol[128] = tint_workgroupUniformLoad_v();
   return tint_symbol[0];
 }

test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.glsl

@@ -4,16 +4,16 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 void unused_entry_point() {
   return;
 }
-int[128] tint_workgroupUniformLoad(inout int p[128]) {
+shared int v[128];
+int[128] tint_workgroupUniformLoad_v() {
   barrier();
-  int result[128] = p;
+  int result[128] = v;
   barrier();
   return result;
 }
 
-shared int v[128];
 int foo() {
-  int tint_symbol[128] = tint_workgroupUniformLoad(v);
+  int tint_symbol[128] = tint_workgroupUniformLoad_v();
   return tint_symbol[0];
 }
 

test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.dxc.hlsl

@@ -3,15 +3,15 @@ void unused_entry_point() {
   return;
 }
 
-bool tint_workgroupUniformLoad(inout bool p) {
+groupshared bool v;
+
+bool tint_workgroupUniformLoad_v() {
   GroupMemoryBarrierWithGroupSync();
-  const bool result = p;
+  const bool result = v;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared bool v;
-
 bool foo() {
-  return tint_workgroupUniformLoad(v);
+  return tint_workgroupUniformLoad_v();
 }

test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.fxc.hlsl

@@ -3,15 +3,15 @@ void unused_entry_point() {
   return;
 }
 
-bool tint_workgroupUniformLoad(inout bool p) {
+groupshared bool v;
+
+bool tint_workgroupUniformLoad_v() {
   GroupMemoryBarrierWithGroupSync();
-  const bool result = p;
+  const bool result = v;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared bool v;
-
 bool foo() {
-  return tint_workgroupUniformLoad(v);
+  return tint_workgroupUniformLoad_v();
 }

test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.glsl

@@ -4,15 +4,15 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 void unused_entry_point() {
   return;
 }
-bool tint_workgroupUniformLoad(inout bool p) {
+shared bool v;
+bool tint_workgroupUniformLoad_v() {
   barrier();
-  bool result = p;
+  bool result = v;
   barrier();
   return result;
 }
 
-shared bool v;
 bool foo() {
-  return tint_workgroupUniformLoad(v);
+  return tint_workgroupUniformLoad_v();
 }
 

test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.dxc.hlsl

@@ -3,22 +3,30 @@ void unused_entry_point() {
   return;
 }
 
-int tint_workgroupUniformLoad(inout int p) {
+groupshared int a;
+
+int tint_workgroupUniformLoad_a() {
   GroupMemoryBarrierWithGroupSync();
-  const int result = p;
+  const int result = a;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared int a;
 groupshared int b;
 
+int tint_workgroupUniformLoad_b() {
+  GroupMemoryBarrierWithGroupSync();
+  const int result = b;
+  GroupMemoryBarrierWithGroupSync();
+  return result;
+}
+
 void foo() {
   {
     int i = 0;
     while (true) {
       const int tint_symbol = i;
-      const int tint_symbol_1 = tint_workgroupUniformLoad(a);
+      const int tint_symbol_1 = tint_workgroupUniformLoad_a();
       if (!((tint_symbol < tint_symbol_1))) {
         break;
       }
@@ -26,7 +34,7 @@ void foo() {
       }
       {
         const int tint_symbol_2 = i;
-        const int tint_symbol_3 = tint_workgroupUniformLoad(b);
+        const int tint_symbol_3 = tint_workgroupUniformLoad_b();
         i = (tint_symbol_2 + tint_symbol_3);
       }
     }

test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.fxc.hlsl

@@ -3,22 +3,30 @@ void unused_entry_point() {
   return;
 }
 
-int tint_workgroupUniformLoad(inout int p) {
+groupshared int a;
+
+int tint_workgroupUniformLoad_a() {
   GroupMemoryBarrierWithGroupSync();
-  const int result = p;
+  const int result = a;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared int a;
 groupshared int b;
 
+int tint_workgroupUniformLoad_b() {
+  GroupMemoryBarrierWithGroupSync();
+  const int result = b;
+  GroupMemoryBarrierWithGroupSync();
+  return result;
+}
+
 void foo() {
   {
     int i = 0;
     while (true) {
       const int tint_symbol = i;
-      const int tint_symbol_1 = tint_workgroupUniformLoad(a);
+      const int tint_symbol_1 = tint_workgroupUniformLoad_a();
       if (!((tint_symbol < tint_symbol_1))) {
         break;
       }
@@ -26,7 +34,7 @@ void foo() {
       }
       {
         const int tint_symbol_2 = i;
-        const int tint_symbol_3 = tint_workgroupUniformLoad(b);
+        const int tint_symbol_3 = tint_workgroupUniformLoad_b();
         i = (tint_symbol_2 + tint_symbol_3);
       }
     }

test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.glsl

@@ -4,21 +4,28 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 void unused_entry_point() {
   return;
 }
-int tint_workgroupUniformLoad(inout int p) {
+shared int a;
+int tint_workgroupUniformLoad_a() {
   barrier();
-  int result = p;
+  int result = a;
   barrier();
   return result;
 }
 
-shared int a;
 shared int b;
+int tint_workgroupUniformLoad_b() {
+  barrier();
+  int result = b;
+  barrier();
+  return result;
+}
+
 void foo() {
   {
     int i = 0;
     while (true) {
       int tint_symbol = i;
-      int tint_symbol_1 = tint_workgroupUniformLoad(a);
+      int tint_symbol_1 = tint_workgroupUniformLoad_a();
       if (!((tint_symbol < tint_symbol_1))) {
         break;
       }
@@ -26,7 +33,7 @@ void foo() {
       }
       {
         int tint_symbol_2 = i;
-        int tint_symbol_3 = tint_workgroupUniformLoad(b);
+        int tint_symbol_3 = tint_workgroupUniformLoad_b();
         i = (tint_symbol_2 + tint_symbol_3);
       }
     }

test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.dxc.hlsl

@@ -3,17 +3,17 @@ void unused_entry_point() {
   return;
 }
 
-bool tint_workgroupUniformLoad(inout bool p) {
+groupshared bool v;
+
+bool tint_workgroupUniformLoad_v() {
   GroupMemoryBarrierWithGroupSync();
-  const bool result = p;
+  const bool result = v;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared bool v;
-
 int foo() {
-  if (tint_workgroupUniformLoad(v)) {
+  if (tint_workgroupUniformLoad_v()) {
     return 42;
   }
   return 0;

test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.fxc.hlsl

@@ -3,17 +3,17 @@ void unused_entry_point() {
   return;
 }
 
-bool tint_workgroupUniformLoad(inout bool p) {
+groupshared bool v;
+
+bool tint_workgroupUniformLoad_v() {
   GroupMemoryBarrierWithGroupSync();
-  const bool result = p;
+  const bool result = v;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared bool v;
-
 int foo() {
-  if (tint_workgroupUniformLoad(v)) {
+  if (tint_workgroupUniformLoad_v()) {
     return 42;
   }
   return 0;

test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.glsl

@@ -4,16 +4,16 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 void unused_entry_point() {
   return;
 }
-bool tint_workgroupUniformLoad(inout bool p) {
+shared bool v;
+bool tint_workgroupUniformLoad_v() {
   barrier();
-  bool result = p;
+  bool result = v;
   barrier();
   return result;
 }
 
-shared bool v;
 int foo() {
-  if (tint_workgroupUniformLoad(v)) {
+  if (tint_workgroupUniformLoad_v()) {
     return 42;
   }
   return 0;

test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.dxc.hlsl

@@ -3,15 +3,15 @@ void unused_entry_point() {
   return;
 }
 
-float3x3 tint_workgroupUniformLoad(inout float3x3 p) {
+groupshared float3x3 v;
+
+float3x3 tint_workgroupUniformLoad_v() {
   GroupMemoryBarrierWithGroupSync();
-  const float3x3 result = p;
+  const float3x3 result = v;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared float3x3 v;
-
 float3x3 foo() {
-  return tint_workgroupUniformLoad(v);
+  return tint_workgroupUniformLoad_v();
 }

test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.fxc.hlsl

@@ -3,15 +3,15 @@ void unused_entry_point() {
   return;
 }
 
-float3x3 tint_workgroupUniformLoad(inout float3x3 p) {
+groupshared float3x3 v;
+
+float3x3 tint_workgroupUniformLoad_v() {
   GroupMemoryBarrierWithGroupSync();
-  const float3x3 result = p;
+  const float3x3 result = v;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared float3x3 v;
-
 float3x3 foo() {
-  return tint_workgroupUniformLoad(v);
+  return tint_workgroupUniformLoad_v();
 }

test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.glsl

@@ -4,15 +4,15 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 void unused_entry_point() {
   return;
 }
-mat3 tint_workgroupUniformLoad(inout mat3 p) {
+shared mat3 v;
+mat3 tint_workgroupUniformLoad_v() {
   barrier();
-  mat3 result = p;
+  mat3 result = v;
   barrier();
   return result;
 }
 
-shared mat3 v;
 mat3 foo() {
-  return tint_workgroupUniformLoad(v);
+  return tint_workgroupUniformLoad_v();
 }
 

test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.dxc.hlsl

@@ -12,15 +12,15 @@ struct Outer {
   Inner a[4];
 };
 
-Outer tint_workgroupUniformLoad(inout Outer p) {
+groupshared Outer v;
+
+Outer tint_workgroupUniformLoad_v() {
   GroupMemoryBarrierWithGroupSync();
-  const Outer result = p;
+  const Outer result = v;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared Outer v;
-
 Outer foo() {
-  return tint_workgroupUniformLoad(v);
+  return tint_workgroupUniformLoad_v();
 }

test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.fxc.hlsl

@@ -12,15 +12,15 @@ struct Outer {
   Inner a[4];
 };
 
-Outer tint_workgroupUniformLoad(inout Outer p) {
+groupshared Outer v;
+
+Outer tint_workgroupUniformLoad_v() {
   GroupMemoryBarrierWithGroupSync();
-  const Outer result = p;
+  const Outer result = v;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared Outer v;
-
 Outer foo() {
-  return tint_workgroupUniformLoad(v);
+  return tint_workgroupUniformLoad_v();
 }

test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.glsl

@@ -14,15 +14,15 @@ struct Outer {
   Inner a[4];
 };
 
-Outer tint_workgroupUniformLoad(inout Outer p) {
+shared Outer v;
+Outer tint_workgroupUniformLoad_v() {
   barrier();
-  Outer result = p;
+  Outer result = v;
   barrier();
   return result;
 }
 
-shared Outer v;
 Outer foo() {
-  return tint_workgroupUniformLoad(v);
+  return tint_workgroupUniformLoad_v();
 }
 

test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.dxc.hlsl

@@ -3,15 +3,15 @@ void unused_entry_point() {
   return;
 }
 
-float4 tint_workgroupUniformLoad(inout float4 p) {
+groupshared float4 v;
+
+float4 tint_workgroupUniformLoad_v() {
   GroupMemoryBarrierWithGroupSync();
-  const float4 result = p;
+  const float4 result = v;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared float4 v;
-
 float4 foo() {
-  return tint_workgroupUniformLoad(v);
+  return tint_workgroupUniformLoad_v();
 }

test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.fxc.hlsl

@@ -3,15 +3,15 @@ void unused_entry_point() {
   return;
 }
 
-float4 tint_workgroupUniformLoad(inout float4 p) {
+groupshared float4 v;
+
+float4 tint_workgroupUniformLoad_v() {
   GroupMemoryBarrierWithGroupSync();
-  const float4 result = p;
+  const float4 result = v;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared float4 v;
-
 float4 foo() {
-  return tint_workgroupUniformLoad(v);
+  return tint_workgroupUniformLoad_v();
 }

test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.glsl

@@ -4,15 +4,15 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 void unused_entry_point() {
   return;
 }
-vec4 tint_workgroupUniformLoad(inout vec4 p) {
+shared vec4 v;
+vec4 tint_workgroupUniformLoad_v() {
   barrier();
-  vec4 result = p;
+  vec4 result = v;
   barrier();
   return result;
 }
 
-shared vec4 v;
 vec4 foo() {
-  return tint_workgroupUniformLoad(v);
+  return tint_workgroupUniformLoad_v();
 }
 

test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.dxc.hlsl

@@ -3,20 +3,21 @@ void unused_entry_point() {
   return;
 }
 
-int tint_workgroupUniformLoad(inout int p) {
+groupshared int v[4];
+
+int tint_workgroupUniformLoad_v_X(uint p[1]) {
   GroupMemoryBarrierWithGroupSync();
-  const int result = p;
+  const int result = v[p[0]];
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared int v[4];
-
 int foo_v_X(uint p[1]) {
-  return tint_workgroupUniformLoad(v[p[0]]);
+  const uint tint_symbol[1] = {p[0u]};
+  return tint_workgroupUniformLoad_v_X(tint_symbol);
 }
 
 int bar() {
-  const uint tint_symbol[1] = (uint[1])0;
+  const uint tint_symbol_1[1] = (uint[1])0;
-  return foo_v_X(tint_symbol);
+  return foo_v_X(tint_symbol_1);
 }

test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.fxc.hlsl

@@ -3,20 +3,21 @@ void unused_entry_point() {
   return;
 }
 
-int tint_workgroupUniformLoad(inout int p) {
+groupshared int v[4];
+
+int tint_workgroupUniformLoad_v_X(uint p[1]) {
   GroupMemoryBarrierWithGroupSync();
-  const int result = p;
+  const int result = v[p[0]];
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
-groupshared int v[4];
-
 int foo_v_X(uint p[1]) {
-  return tint_workgroupUniformLoad(v[p[0]]);
+  const uint tint_symbol[1] = {p[0u]};
+  return tint_workgroupUniformLoad_v_X(tint_symbol);
 }
 
 int bar() {
-  const uint tint_symbol[1] = (uint[1])0;
+  const uint tint_symbol_1[1] = (uint[1])0;
-  return foo_v_X(tint_symbol);
+  return foo_v_X(tint_symbol_1);
 }

test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.glsl

@@ -4,20 +4,21 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 void unused_entry_point() {
   return;
 }
-int tint_workgroupUniformLoad(inout int p) {
+shared int v[4];
+int tint_workgroupUniformLoad_v_X(uint p[1]) {
   barrier();
-  int result = p;
+  int result = v[p[0]];
   barrier();
   return result;
 }
 
-shared int v[4];
 int foo_v_X(uint p[1]) {
-  return tint_workgroupUniformLoad(v[p[0]]);
+  uint tint_symbol[1] = uint[1](p[0u]);
+  return tint_workgroupUniformLoad_v_X(tint_symbol);
 }
 
 int bar() {
-  uint tint_symbol[1] = uint[1](0u);
+  uint tint_symbol_1[1] = uint[1](0u);
-  return foo_v_X(tint_symbol);
+  return foo_v_X(tint_symbol_1);
 }