Google Git
Sign in
dawn/dawn/fa00fe9d417fbbbfcdae4390a85faedc9dc6dfb0^!/.
commit
fa00fe9d417fbbbfcdae4390a85faedc9dc6dfb0
[log]
author
Antonio Maiorano <amaiorano@google.com>
Wed May 03 15:30:54 2023 +0000
committer
Dawn LUCI CQ <dawn-scoped@luci-project-accounts.iam.gserviceaccount.com>
Wed May 03 15:30:54 2023 +0000
tree
d757a7389b1c6d0cafb9d57f7bb414045677572b
parent
e903396ff2c5d4c205bffc5c6b91cb48ebefd5f7 [diff]
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>

diff --git a/src/tint/transform/builtin_polyfill_test.cc b/src/tint/transform/builtin_polyfill_test.cc
index 65c1f29..ad47af6 100644
--- a/src/tint/transform/builtin_polyfill_test.cc
+++ b/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 @@
 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 @@
     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
 ////////////////////////////////////////////////////////////////////////////////

diff --git a/src/tint/writer/glsl/generator_impl.cc b/src/tint/writer/glsl/generator_impl.cc
index 8fcdd3d..181bce8 100644
--- a/src/tint/writer/glsl/generator_impl.cc
+++ b/src/tint/writer/glsl/generator_impl.cc

@@ -170,7 +170,6 @@
     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 @@
         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.

diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc
index ed86e9e..fefd424 100644
--- a/src/tint/writer/hlsl/generator_impl.cc
+++ b/src/tint/writer/hlsl/generator_impl.cc

@@ -176,8 +176,6 @@
 
     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 @@
         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.

diff --git a/src/tint/writer/spirv/generator_impl.cc b/src/tint/writer/spirv/generator_impl.cc
index 80485c9..0a3f99f 100644
--- a/src/tint/writer/spirv/generator_impl.cc
+++ b/src/tint/writer/spirv/generator_impl.cc

@@ -119,7 +119,7 @@
         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 =

diff --git a/test/tint/bug/tint/1926.wgsl b/test/tint/bug/tint/1926.wgsl
new file mode 100644
index 0000000..5e86097
--- /dev/null
+++ b/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;
+    }
+}

diff --git a/test/tint/bug/tint/1926.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/1926.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..9367ce7
--- /dev/null
+++ b/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;
+}

diff --git a/test/tint/bug/tint/1926.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/1926.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..9367ce7
--- /dev/null
+++ b/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;
+}

diff --git a/test/tint/bug/tint/1926.wgsl.expected.glsl b/test/tint/bug/tint/1926.wgsl.expected.glsl
new file mode 100644
index 0000000..5e3cd06
--- /dev/null
+++ b/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;
+}

diff --git a/test/tint/bug/tint/1926.wgsl.expected.msl b/test/tint/bug/tint/1926.wgsl.expected.msl
new file mode 100644
index 0000000..1a9d82f
--- /dev/null
+++ b/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;
+}
+

diff --git a/test/tint/bug/tint/1926.wgsl.expected.spvasm b/test/tint/bug/tint/1926.wgsl.expected.spvasm
new file mode 100644
index 0000000..48ba716
--- /dev/null
+++ b/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

diff --git a/test/tint/bug/tint/1926.wgsl.expected.wgsl b/test/tint/bug/tint/1926.wgsl.expected.wgsl
new file mode 100644
index 0000000..c485055
--- /dev/null
+++ b/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;
+  }
+}

diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.dxc.hlsl
index 7db1f5c..33ea9c0 100644
--- a/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.dxc.hlsl
+++ b/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));
 }
 

diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.fxc.hlsl
index 7db1f5c..33ea9c0 100644
--- a/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.fxc.hlsl
+++ b/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));
 }
 

diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.glsl b/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.glsl
index 08650d7..fd29d9c 100644
--- a/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.glsl
+++ b/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;
 }
 

diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.dxc.hlsl
index 31bcd63..05e4c85 100644
--- a/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.dxc.hlsl
+++ b/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));
 }
 

diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.fxc.hlsl
index 31bcd63..05e4c85 100644
--- a/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.fxc.hlsl
+++ b/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));
 }
 

diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.glsl b/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.glsl
index b0d7db3..1dc5546 100644
--- a/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.glsl
+++ b/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;
 }
 

diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.dxc.hlsl
index c495ebc..6d8b261 100644
--- a/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.dxc.hlsl
+++ b/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));
 }
 

diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.fxc.hlsl
index c495ebc..6d8b261 100644
--- a/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.fxc.hlsl
+++ b/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));
 }
 

diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.glsl b/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.glsl
index 8d3a9e6..5919644 100644
--- a/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.glsl
+++ b/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;
 }
 

diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.dxc.hlsl
index ed848f2..c8e40ee 100644
--- a/test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.dxc.hlsl
+++ b/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);
 }
 

diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.glsl b/test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.glsl
index 9d804d4..fe2cfd0 100644
--- a/test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.glsl
+++ b/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;
 }
 

diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.dxc.hlsl
index 7db1f5c..33ea9c0 100644
--- a/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.dxc.hlsl
+++ b/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));
 }
 

diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.fxc.hlsl
index 7db1f5c..33ea9c0 100644
--- a/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.fxc.hlsl
+++ b/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));
 }
 

diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.glsl b/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.glsl
index 08650d7..fd29d9c 100644
--- a/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.glsl
+++ b/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;
 }
 

diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.dxc.hlsl
index 31bcd63..05e4c85 100644
--- a/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.dxc.hlsl
+++ b/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));
 }
 

diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.fxc.hlsl
index 31bcd63..05e4c85 100644
--- a/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.fxc.hlsl
+++ b/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));
 }
 

diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.glsl b/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.glsl
index b0d7db3..1dc5546 100644
--- a/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.glsl
+++ b/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;
 }
 

diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.dxc.hlsl
index c495ebc..6d8b261 100644
--- a/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.dxc.hlsl
+++ b/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));
 }
 

diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.fxc.hlsl
index c495ebc..6d8b261 100644
--- a/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.fxc.hlsl
+++ b/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));
 }
 

diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.glsl b/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.glsl
index 8d3a9e6..5919644 100644
--- a/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.glsl
+++ b/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;
 }
 

diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl.expected.dxc.hlsl
index ed848f2..c8e40ee 100644
--- a/test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl.expected.dxc.hlsl
+++ b/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);
 }
 

diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl.expected.glsl b/test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl.expected.glsl
index 9d804d4..fe2cfd0 100644
--- a/test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl.expected.glsl
+++ b/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;
 }
 

diff --git a/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.dxc.hlsl b/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.dxc.hlsl
index d8a47b9..a9b6ea5 100644
--- a/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.dxc.hlsl

@@ -3,17 +3,17 @@
   return;
 }
 
-typedef int tint_workgroupUniformLoad_ret[4];
-tint_workgroupUniformLoad_ret tint_workgroupUniformLoad(inout int p[4]) {
+groupshared int v[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();
 }

diff --git a/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.fxc.hlsl b/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.fxc.hlsl
index d8a47b9..a9b6ea5 100644
--- a/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.fxc.hlsl

@@ -3,17 +3,17 @@
   return;
 }
 
-typedef int tint_workgroupUniformLoad_ret[4];
-tint_workgroupUniformLoad_ret tint_workgroupUniformLoad(inout int p[4]) {
+groupshared int v[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();
 }

diff --git a/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.glsl b/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.glsl
index 661c13b..cda9b79 100644
--- a/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.glsl
+++ b/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.glsl

@@ -4,15 +4,15 @@
 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();
 }
 

diff --git a/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.dxc.hlsl b/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.dxc.hlsl
index 1958f4a..a42cbe5 100644
--- a/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.dxc.hlsl

@@ -3,17 +3,17 @@
   return;
 }
 
-typedef int tint_workgroupUniformLoad_ret[128];
-tint_workgroupUniformLoad_ret tint_workgroupUniformLoad(inout int p[128]) {
+groupshared int v[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];
 }

diff --git a/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.fxc.hlsl b/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.fxc.hlsl
index 1958f4a..a42cbe5 100644
--- a/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.fxc.hlsl

@@ -3,17 +3,17 @@
   return;
 }
 
-typedef int tint_workgroupUniformLoad_ret[128];
-tint_workgroupUniformLoad_ret tint_workgroupUniformLoad(inout int p[128]) {
+groupshared int v[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];
 }

diff --git a/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.glsl b/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.glsl
index fa13387..6e551e8 100644
--- a/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.glsl
+++ b/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.glsl

@@ -4,16 +4,16 @@
 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];
 }
 

diff --git a/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.dxc.hlsl b/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.dxc.hlsl
index 1958f4a..a42cbe5 100644
--- a/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.dxc.hlsl

@@ -3,17 +3,17 @@
   return;
 }
 
-typedef int tint_workgroupUniformLoad_ret[128];
-tint_workgroupUniformLoad_ret tint_workgroupUniformLoad(inout int p[128]) {
+groupshared int v[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];
 }

diff --git a/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.fxc.hlsl b/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.fxc.hlsl
index 1958f4a..a42cbe5 100644
--- a/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.fxc.hlsl

@@ -3,17 +3,17 @@
   return;
 }
 
-typedef int tint_workgroupUniformLoad_ret[128];
-tint_workgroupUniformLoad_ret tint_workgroupUniformLoad(inout int p[128]) {
+groupshared int v[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];
 }

diff --git a/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.glsl b/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.glsl
index fa13387..6e551e8 100644
--- a/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.glsl
+++ b/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.glsl

@@ -4,16 +4,16 @@
 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];
 }
 

diff --git a/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.dxc.hlsl b/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.dxc.hlsl
index 28f1e19..afb3e79 100644
--- a/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.dxc.hlsl

@@ -3,15 +3,15 @@
   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();
 }

diff --git a/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.fxc.hlsl b/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.fxc.hlsl
index 28f1e19..afb3e79 100644
--- a/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.fxc.hlsl

@@ -3,15 +3,15 @@
   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();
 }

diff --git a/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.glsl b/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.glsl
index 9793dba..4962454 100644
--- a/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.glsl
+++ b/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.glsl

@@ -4,15 +4,15 @@
 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();
 }
 

diff --git a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.dxc.hlsl b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.dxc.hlsl
index 98df2c2..ec23355 100644
--- a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.dxc.hlsl

@@ -3,22 +3,30 @@
   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 @@
       }
       {
         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);
       }
     }

diff --git a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.fxc.hlsl b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.fxc.hlsl
index 98df2c2..ec23355 100644
--- a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.fxc.hlsl

@@ -3,22 +3,30 @@
   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 @@
       }
       {
         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);
       }
     }

diff --git a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.glsl b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.glsl
index 97926d0..09c4abe 100644
--- a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.glsl
+++ b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.glsl

@@ -4,21 +4,28 @@
 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 @@
       }
       {
         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);
       }
     }

diff --git a/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.dxc.hlsl b/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.dxc.hlsl
index bfc2a25..48a589f 100644
--- a/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.dxc.hlsl

@@ -3,17 +3,17 @@
   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;

diff --git a/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.fxc.hlsl b/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.fxc.hlsl
index bfc2a25..48a589f 100644
--- a/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.fxc.hlsl

@@ -3,17 +3,17 @@
   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;

diff --git a/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.glsl b/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.glsl
index ef53245..ad15d25 100644
--- a/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.glsl
+++ b/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.glsl

@@ -4,16 +4,16 @@
 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;

diff --git a/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.dxc.hlsl b/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.dxc.hlsl
index 053093e..1b22b46 100644
--- a/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.dxc.hlsl

@@ -3,15 +3,15 @@
   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();
 }

diff --git a/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.fxc.hlsl b/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.fxc.hlsl
index 053093e..1b22b46 100644
--- a/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.fxc.hlsl

@@ -3,15 +3,15 @@
   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();
 }

diff --git a/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.glsl b/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.glsl
index 86ea506..ac70b34a 100644
--- a/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.glsl
+++ b/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.glsl

@@ -4,15 +4,15 @@
 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();
 }
 

diff --git a/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.dxc.hlsl b/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.dxc.hlsl
index 201beac..eb3a52c 100644
--- a/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.dxc.hlsl

@@ -12,15 +12,15 @@
   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();
 }

diff --git a/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.fxc.hlsl b/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.fxc.hlsl
index 201beac..eb3a52c 100644
--- a/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.fxc.hlsl

@@ -12,15 +12,15 @@
   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();
 }

diff --git a/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.glsl b/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.glsl
index 203e90d..7f9766a 100644
--- a/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.glsl
+++ b/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.glsl

@@ -14,15 +14,15 @@
   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();
 }
 

diff --git a/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.dxc.hlsl b/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.dxc.hlsl
index 77473c206..80cd6df 100644
--- a/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.dxc.hlsl

@@ -3,15 +3,15 @@
   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();
 }

diff --git a/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.fxc.hlsl b/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.fxc.hlsl
index 77473c206..80cd6df 100644
--- a/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.fxc.hlsl

@@ -3,15 +3,15 @@
   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();
 }

diff --git a/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.glsl b/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.glsl
index 76e29ed..f4dc643 100644
--- a/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.glsl
+++ b/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.glsl

@@ -4,15 +4,15 @@
 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();
 }
 

diff --git a/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.dxc.hlsl b/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.dxc.hlsl
index 2d0fe18..de20d9c 100644
--- a/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.dxc.hlsl

@@ -3,20 +3,21 @@
   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;
-  return foo_v_X(tint_symbol);
+  const uint tint_symbol_1[1] = (uint[1])0;
+  return foo_v_X(tint_symbol_1);
 }

diff --git a/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.fxc.hlsl b/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.fxc.hlsl
index 2d0fe18..de20d9c 100644
--- a/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.fxc.hlsl

@@ -3,20 +3,21 @@
   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;
-  return foo_v_X(tint_symbol);
+  const uint tint_symbol_1[1] = (uint[1])0;
+  return foo_v_X(tint_symbol_1);
 }

diff --git a/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.glsl b/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.glsl
index aa2e549..7555463 100644
--- a/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.glsl
+++ b/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.glsl

@@ -4,20 +4,21 @@
 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);
-  return foo_v_X(tint_symbol);
+  uint tint_symbol_1[1] = uint[1](0u);
+  return foo_v_X(tint_symbol_1);
 }