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);
}
