tint: Add support for workgroupUniformLoad
Accept any type in the intrinsics definition, and then manually
validate that there are no atomics in the type. Add manual E2E tests
for composite types.
Use the BuiltinPolyfill transform to implement it for all backends.
Update the uniformity analysis with special-case tags for the builtin.
Fixed: tint:1780
Change-Id: I95786dff4df70a0b16ed1c53b853b5d0ec6bc501
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/114862
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: James Price <jrprice@google.com>
Kokoro: James Price <jrprice@google.com>
diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl b/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl
new file mode 100644
index 0000000..ff17877
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl
@@ -0,0 +1,33 @@
+// Copyright 2023 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/src/cmd/gen
+// using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+var<workgroup> arg_0: u32;
+
+// fn workgroupUniformLoad(ptr<workgroup, u32, read_write>) -> u32
+fn workgroupUniformLoad_37307c() {
+ var res: u32 = workgroupUniformLoad(&arg_0);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ workgroupUniformLoad_37307c();
+}
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
new file mode 100644
index 0000000..26bb61f
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+groupshared uint arg_0;
+
+uint tint_workgroupUniformLoad_arg_0() {
+ GroupMemoryBarrierWithGroupSync();
+ const uint result = arg_0;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+void workgroupUniformLoad_37307c() {
+ uint res = tint_workgroupUniformLoad_arg_0();
+}
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ arg_0 = 0u;
+ }
+ GroupMemoryBarrierWithGroupSync();
+ workgroupUniformLoad_37307c();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
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
new file mode 100644
index 0000000..26bb61f
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.fxc.hlsl
@@ -0,0 +1,30 @@
+groupshared uint arg_0;
+
+uint tint_workgroupUniformLoad_arg_0() {
+ GroupMemoryBarrierWithGroupSync();
+ const uint result = arg_0;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+void workgroupUniformLoad_37307c() {
+ uint res = tint_workgroupUniformLoad_arg_0();
+}
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ arg_0 = 0u;
+ }
+ GroupMemoryBarrierWithGroupSync();
+ workgroupUniformLoad_37307c();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.glsl b/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.glsl
new file mode 100644
index 0000000..5b51b2d
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.glsl
@@ -0,0 +1,27 @@
+#version 310 es
+
+shared uint arg_0;
+uint tint_workgroupUniformLoad_arg_0() {
+ barrier();
+ uint result = arg_0;
+ barrier();
+ return result;
+}
+
+void workgroupUniformLoad_37307c() {
+ uint res = tint_workgroupUniformLoad_arg_0();
+}
+
+void compute_main(uint local_invocation_index) {
+ {
+ arg_0 = 0u;
+ }
+ barrier();
+ workgroupUniformLoad_37307c();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main(gl_LocalInvocationIndex);
+ return;
+}
diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.msl b/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.msl
new file mode 100644
index 0000000..3553390
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.msl
@@ -0,0 +1,28 @@
+#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 workgroupUniformLoad_37307c(threadgroup uint* const tint_symbol) {
+ uint res = tint_workgroupUniformLoad(tint_symbol);
+}
+
+void compute_main_inner(uint local_invocation_index, threadgroup uint* const tint_symbol_1) {
+ {
+ *(tint_symbol_1) = 0u;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ workgroupUniformLoad_37307c(tint_symbol_1);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup uint tint_symbol_2;
+ compute_main_inner(local_invocation_index, &(tint_symbol_2));
+ return;
+}
+
diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.spvasm
new file mode 100644
index 0000000..b3755dc
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.spvasm
@@ -0,0 +1,59 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 32
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %local_invocation_index_1 "local_invocation_index_1"
+ OpName %arg_0 "arg_0"
+ OpName %tint_workgroupUniformLoad_arg_0 "tint_workgroupUniformLoad_arg_0"
+ OpName %workgroupUniformLoad_37307c "workgroupUniformLoad_37307c"
+ OpName %res "res"
+ OpName %compute_main_inner "compute_main_inner"
+ OpName %local_invocation_index "local_invocation_index"
+ OpName %compute_main "compute_main"
+ OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+ %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+ %arg_0 = OpVariable %_ptr_Workgroup_uint Workgroup
+ %6 = OpTypeFunction %uint
+ %void = OpTypeVoid
+ %uint_2 = OpConstant %uint 2
+ %uint_264 = OpConstant %uint 264
+ %15 = OpTypeFunction %void
+%_ptr_Function_uint = OpTypePointer Function %uint
+ %21 = OpConstantNull %uint
+ %22 = OpTypeFunction %void %uint
+%tint_workgroupUniformLoad_arg_0 = OpFunction %uint None %6
+ %8 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %13 = OpLoad %uint %arg_0
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ OpReturnValue %13
+ OpFunctionEnd
+%workgroupUniformLoad_37307c = OpFunction %void None %15
+ %17 = OpLabel
+ %res = OpVariable %_ptr_Function_uint Function %21
+ %18 = OpFunctionCall %uint %tint_workgroupUniformLoad_arg_0
+ OpStore %res %18
+ OpReturn
+ OpFunctionEnd
+%compute_main_inner = OpFunction %void None %22
+%local_invocation_index = OpFunctionParameter %uint
+ %25 = OpLabel
+ OpStore %arg_0 %21
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %27 = OpFunctionCall %void %workgroupUniformLoad_37307c
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %15
+ %29 = OpLabel
+ %31 = OpLoad %uint %local_invocation_index_1
+ %30 = OpFunctionCall %void %compute_main_inner %31
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.wgsl
new file mode 100644
index 0000000..f4d48fa
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.wgsl
@@ -0,0 +1,10 @@
+var<workgroup> arg_0 : u32;
+
+fn workgroupUniformLoad_37307c() {
+ var res : u32 = workgroupUniformLoad(&(arg_0));
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ workgroupUniformLoad_37307c();
+}
diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl b/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl
new file mode 100644
index 0000000..c5cf95e
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl
@@ -0,0 +1,33 @@
+// Copyright 2023 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/src/cmd/gen
+// using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+var<workgroup> arg_0: f32;
+
+// fn workgroupUniformLoad(ptr<workgroup, f32, read_write>) -> f32
+fn workgroupUniformLoad_7a857c() {
+ var res: f32 = workgroupUniformLoad(&arg_0);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ workgroupUniformLoad_7a857c();
+}
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
new file mode 100644
index 0000000..fe6d27e
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+groupshared float arg_0;
+
+float tint_workgroupUniformLoad_arg_0() {
+ GroupMemoryBarrierWithGroupSync();
+ const float result = arg_0;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+void workgroupUniformLoad_7a857c() {
+ float res = tint_workgroupUniformLoad_arg_0();
+}
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ arg_0 = 0.0f;
+ }
+ GroupMemoryBarrierWithGroupSync();
+ workgroupUniformLoad_7a857c();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
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
new file mode 100644
index 0000000..fe6d27e
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.fxc.hlsl
@@ -0,0 +1,30 @@
+groupshared float arg_0;
+
+float tint_workgroupUniformLoad_arg_0() {
+ GroupMemoryBarrierWithGroupSync();
+ const float result = arg_0;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+void workgroupUniformLoad_7a857c() {
+ float res = tint_workgroupUniformLoad_arg_0();
+}
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ arg_0 = 0.0f;
+ }
+ GroupMemoryBarrierWithGroupSync();
+ workgroupUniformLoad_7a857c();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.glsl b/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.glsl
new file mode 100644
index 0000000..bfdeb87
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.glsl
@@ -0,0 +1,27 @@
+#version 310 es
+
+shared float arg_0;
+float tint_workgroupUniformLoad_arg_0() {
+ barrier();
+ float result = arg_0;
+ barrier();
+ return result;
+}
+
+void workgroupUniformLoad_7a857c() {
+ float res = tint_workgroupUniformLoad_arg_0();
+}
+
+void compute_main(uint local_invocation_index) {
+ {
+ arg_0 = 0.0f;
+ }
+ barrier();
+ workgroupUniformLoad_7a857c();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main(gl_LocalInvocationIndex);
+ return;
+}
diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.msl b/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.msl
new file mode 100644
index 0000000..7637172
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.msl
@@ -0,0 +1,28 @@
+#include <metal_stdlib>
+
+using namespace metal;
+float tint_workgroupUniformLoad(threadgroup float* const p) {
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ float const result = *(p);
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ return result;
+}
+
+void workgroupUniformLoad_7a857c(threadgroup float* const tint_symbol) {
+ float res = tint_workgroupUniformLoad(tint_symbol);
+}
+
+void compute_main_inner(uint local_invocation_index, threadgroup float* const tint_symbol_1) {
+ {
+ *(tint_symbol_1) = 0.0f;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ workgroupUniformLoad_7a857c(tint_symbol_1);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup float tint_symbol_2;
+ compute_main_inner(local_invocation_index, &(tint_symbol_2));
+ return;
+}
+
diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.spvasm
new file mode 100644
index 0000000..abd0f6a
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.spvasm
@@ -0,0 +1,60 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 33
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %local_invocation_index_1 "local_invocation_index_1"
+ OpName %arg_0 "arg_0"
+ OpName %tint_workgroupUniformLoad_arg_0 "tint_workgroupUniformLoad_arg_0"
+ OpName %workgroupUniformLoad_7a857c "workgroupUniformLoad_7a857c"
+ OpName %res "res"
+ OpName %compute_main_inner "compute_main_inner"
+ OpName %local_invocation_index "local_invocation_index"
+ OpName %compute_main "compute_main"
+ OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+ %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+ %float = OpTypeFloat 32
+%_ptr_Workgroup_float = OpTypePointer Workgroup %float
+ %arg_0 = OpVariable %_ptr_Workgroup_float Workgroup
+ %7 = OpTypeFunction %float
+ %void = OpTypeVoid
+ %uint_2 = OpConstant %uint 2
+ %uint_264 = OpConstant %uint 264
+ %16 = OpTypeFunction %void
+%_ptr_Function_float = OpTypePointer Function %float
+ %22 = OpConstantNull %float
+ %23 = OpTypeFunction %void %uint
+%tint_workgroupUniformLoad_arg_0 = OpFunction %float None %7
+ %9 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %14 = OpLoad %float %arg_0
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ OpReturnValue %14
+ OpFunctionEnd
+%workgroupUniformLoad_7a857c = OpFunction %void None %16
+ %18 = OpLabel
+ %res = OpVariable %_ptr_Function_float Function %22
+ %19 = OpFunctionCall %float %tint_workgroupUniformLoad_arg_0
+ OpStore %res %19
+ OpReturn
+ OpFunctionEnd
+%compute_main_inner = OpFunction %void None %23
+%local_invocation_index = OpFunctionParameter %uint
+ %26 = OpLabel
+ OpStore %arg_0 %22
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %28 = OpFunctionCall %void %workgroupUniformLoad_7a857c
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %16
+ %30 = OpLabel
+ %32 = OpLoad %uint %local_invocation_index_1
+ %31 = OpFunctionCall %void %compute_main_inner %32
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.wgsl
new file mode 100644
index 0000000..2fd46aa
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.wgsl
@@ -0,0 +1,10 @@
+var<workgroup> arg_0 : f32;
+
+fn workgroupUniformLoad_7a857c() {
+ var res : f32 = workgroupUniformLoad(&(arg_0));
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ workgroupUniformLoad_7a857c();
+}
diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl b/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl
new file mode 100644
index 0000000..6ebb067
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl
@@ -0,0 +1,33 @@
+// Copyright 2023 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/src/cmd/gen
+// using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+var<workgroup> arg_0: i32;
+
+// fn workgroupUniformLoad(ptr<workgroup, i32, read_write>) -> i32
+fn workgroupUniformLoad_9d33de() {
+ var res: i32 = workgroupUniformLoad(&arg_0);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ workgroupUniformLoad_9d33de();
+}
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
new file mode 100644
index 0000000..87cf14c
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+groupshared int arg_0;
+
+int tint_workgroupUniformLoad_arg_0() {
+ GroupMemoryBarrierWithGroupSync();
+ const int result = arg_0;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+void workgroupUniformLoad_9d33de() {
+ int res = tint_workgroupUniformLoad_arg_0();
+}
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ arg_0 = 0;
+ }
+ GroupMemoryBarrierWithGroupSync();
+ workgroupUniformLoad_9d33de();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
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
new file mode 100644
index 0000000..87cf14c
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.fxc.hlsl
@@ -0,0 +1,30 @@
+groupshared int arg_0;
+
+int tint_workgroupUniformLoad_arg_0() {
+ GroupMemoryBarrierWithGroupSync();
+ const int result = arg_0;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+void workgroupUniformLoad_9d33de() {
+ int res = tint_workgroupUniformLoad_arg_0();
+}
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ arg_0 = 0;
+ }
+ GroupMemoryBarrierWithGroupSync();
+ workgroupUniformLoad_9d33de();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.glsl b/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.glsl
new file mode 100644
index 0000000..31ae3f9
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.glsl
@@ -0,0 +1,27 @@
+#version 310 es
+
+shared int arg_0;
+int tint_workgroupUniformLoad_arg_0() {
+ barrier();
+ int result = arg_0;
+ barrier();
+ return result;
+}
+
+void workgroupUniformLoad_9d33de() {
+ int res = tint_workgroupUniformLoad_arg_0();
+}
+
+void compute_main(uint local_invocation_index) {
+ {
+ arg_0 = 0;
+ }
+ barrier();
+ workgroupUniformLoad_9d33de();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main(gl_LocalInvocationIndex);
+ return;
+}
diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.msl b/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.msl
new file mode 100644
index 0000000..4ef9721
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.msl
@@ -0,0 +1,28 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int tint_workgroupUniformLoad(threadgroup int* const p) {
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ int const result = *(p);
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ return result;
+}
+
+void workgroupUniformLoad_9d33de(threadgroup int* const tint_symbol) {
+ int res = tint_workgroupUniformLoad(tint_symbol);
+}
+
+void compute_main_inner(uint local_invocation_index, threadgroup int* const tint_symbol_1) {
+ {
+ *(tint_symbol_1) = 0;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ workgroupUniformLoad_9d33de(tint_symbol_1);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup int tint_symbol_2;
+ compute_main_inner(local_invocation_index, &(tint_symbol_2));
+ return;
+}
+
diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.spvasm
new file mode 100644
index 0000000..50b0e3f
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.spvasm
@@ -0,0 +1,60 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 33
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %local_invocation_index_1 "local_invocation_index_1"
+ OpName %arg_0 "arg_0"
+ OpName %tint_workgroupUniformLoad_arg_0 "tint_workgroupUniformLoad_arg_0"
+ OpName %workgroupUniformLoad_9d33de "workgroupUniformLoad_9d33de"
+ OpName %res "res"
+ OpName %compute_main_inner "compute_main_inner"
+ OpName %local_invocation_index "local_invocation_index"
+ OpName %compute_main "compute_main"
+ OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+ %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+ %int = OpTypeInt 32 1
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+ %arg_0 = OpVariable %_ptr_Workgroup_int Workgroup
+ %7 = OpTypeFunction %int
+ %void = OpTypeVoid
+ %uint_2 = OpConstant %uint 2
+ %uint_264 = OpConstant %uint 264
+ %16 = OpTypeFunction %void
+%_ptr_Function_int = OpTypePointer Function %int
+ %22 = OpConstantNull %int
+ %23 = OpTypeFunction %void %uint
+%tint_workgroupUniformLoad_arg_0 = OpFunction %int None %7
+ %9 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %14 = OpLoad %int %arg_0
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ OpReturnValue %14
+ OpFunctionEnd
+%workgroupUniformLoad_9d33de = OpFunction %void None %16
+ %18 = OpLabel
+ %res = OpVariable %_ptr_Function_int Function %22
+ %19 = OpFunctionCall %int %tint_workgroupUniformLoad_arg_0
+ OpStore %res %19
+ OpReturn
+ OpFunctionEnd
+%compute_main_inner = OpFunction %void None %23
+%local_invocation_index = OpFunctionParameter %uint
+ %26 = OpLabel
+ OpStore %arg_0 %22
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %28 = OpFunctionCall %void %workgroupUniformLoad_9d33de
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %16
+ %30 = OpLabel
+ %32 = OpLoad %uint %local_invocation_index_1
+ %31 = OpFunctionCall %void %compute_main_inner %32
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.wgsl
new file mode 100644
index 0000000..5228ac7
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.wgsl
@@ -0,0 +1,10 @@
+var<workgroup> arg_0 : i32;
+
+fn workgroupUniformLoad_9d33de() {
+ var res : i32 = workgroupUniformLoad(&(arg_0));
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ workgroupUniformLoad_9d33de();
+}
diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl b/test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl
new file mode 100644
index 0000000..a0820bb
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl
@@ -0,0 +1,35 @@
+// Copyright 2023 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/src/cmd/gen
+// using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+enable f16;
+var<workgroup> arg_0: f16;
+
+// fn workgroupUniformLoad(ptr<workgroup, f16, read_write>) -> f16
+fn workgroupUniformLoad_e07d08() {
+ var res: f16 = workgroupUniformLoad(&arg_0);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ workgroupUniformLoad_e07d08();
+}
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
new file mode 100644
index 0000000..d34a7c4
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+groupshared float16_t arg_0;
+
+float16_t tint_workgroupUniformLoad_arg_0() {
+ GroupMemoryBarrierWithGroupSync();
+ const float16_t result = arg_0;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+void workgroupUniformLoad_e07d08() {
+ float16_t res = tint_workgroupUniformLoad_arg_0();
+}
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ arg_0 = float16_t(0.0h);
+ }
+ GroupMemoryBarrierWithGroupSync();
+ workgroupUniformLoad_e07d08();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..bd76914
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.fxc.hlsl
@@ -0,0 +1,35 @@
+SKIP: FAILED
+
+groupshared float16_t arg_0;
+
+float16_t tint_workgroupUniformLoad_arg_0() {
+ GroupMemoryBarrierWithGroupSync();
+ const float16_t result = arg_0;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+void workgroupUniformLoad_e07d08() {
+ float16_t res = tint_workgroupUniformLoad_arg_0();
+}
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ arg_0 = float16_t(0.0h);
+ }
+ GroupMemoryBarrierWithGroupSync();
+ workgroupUniformLoad_e07d08();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
+
+FXC validation failure:
+T:\tmp\dawn-temp\dawn-src\test\tint\Shader@0x000001D7B7FE0DC0(1,13-21): error X3000: unrecognized identifier 'float16_t'
diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.glsl b/test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.glsl
new file mode 100644
index 0000000..0b882bd
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.glsl
@@ -0,0 +1,28 @@
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+shared float16_t arg_0;
+float16_t tint_workgroupUniformLoad_arg_0() {
+ barrier();
+ float16_t result = arg_0;
+ barrier();
+ return result;
+}
+
+void workgroupUniformLoad_e07d08() {
+ float16_t res = tint_workgroupUniformLoad_arg_0();
+}
+
+void compute_main(uint local_invocation_index) {
+ {
+ arg_0 = 0.0hf;
+ }
+ barrier();
+ workgroupUniformLoad_e07d08();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main(gl_LocalInvocationIndex);
+ return;
+}
diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.msl b/test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.msl
new file mode 100644
index 0000000..c4c40c9
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.msl
@@ -0,0 +1,28 @@
+#include <metal_stdlib>
+
+using namespace metal;
+half tint_workgroupUniformLoad(threadgroup half* const p) {
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ half const result = *(p);
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ return result;
+}
+
+void workgroupUniformLoad_e07d08(threadgroup half* const tint_symbol) {
+ half res = tint_workgroupUniformLoad(tint_symbol);
+}
+
+void compute_main_inner(uint local_invocation_index, threadgroup half* const tint_symbol_1) {
+ {
+ *(tint_symbol_1) = 0.0h;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ workgroupUniformLoad_e07d08(tint_symbol_1);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup half tint_symbol_2;
+ compute_main_inner(local_invocation_index, &(tint_symbol_2));
+ return;
+}
+
diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.spvasm
new file mode 100644
index 0000000..9c5cf13
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.spvasm
@@ -0,0 +1,64 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 33
+; Schema: 0
+ OpCapability Shader
+ OpCapability Float16
+ OpCapability UniformAndStorageBuffer16BitAccess
+ OpCapability StorageBuffer16BitAccess
+ OpCapability StorageInputOutput16
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %local_invocation_index_1 "local_invocation_index_1"
+ OpName %arg_0 "arg_0"
+ OpName %tint_workgroupUniformLoad_arg_0 "tint_workgroupUniformLoad_arg_0"
+ OpName %workgroupUniformLoad_e07d08 "workgroupUniformLoad_e07d08"
+ OpName %res "res"
+ OpName %compute_main_inner "compute_main_inner"
+ OpName %local_invocation_index "local_invocation_index"
+ OpName %compute_main "compute_main"
+ OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+ %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+ %half = OpTypeFloat 16
+%_ptr_Workgroup_half = OpTypePointer Workgroup %half
+ %arg_0 = OpVariable %_ptr_Workgroup_half Workgroup
+ %7 = OpTypeFunction %half
+ %void = OpTypeVoid
+ %uint_2 = OpConstant %uint 2
+ %uint_264 = OpConstant %uint 264
+ %16 = OpTypeFunction %void
+%_ptr_Function_half = OpTypePointer Function %half
+ %22 = OpConstantNull %half
+ %23 = OpTypeFunction %void %uint
+%tint_workgroupUniformLoad_arg_0 = OpFunction %half None %7
+ %9 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %14 = OpLoad %half %arg_0
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ OpReturnValue %14
+ OpFunctionEnd
+%workgroupUniformLoad_e07d08 = OpFunction %void None %16
+ %18 = OpLabel
+ %res = OpVariable %_ptr_Function_half Function %22
+ %19 = OpFunctionCall %half %tint_workgroupUniformLoad_arg_0
+ OpStore %res %19
+ OpReturn
+ OpFunctionEnd
+%compute_main_inner = OpFunction %void None %23
+%local_invocation_index = OpFunctionParameter %uint
+ %26 = OpLabel
+ OpStore %arg_0 %22
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %28 = OpFunctionCall %void %workgroupUniformLoad_e07d08
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %16
+ %30 = OpLabel
+ %32 = OpLoad %uint %local_invocation_index_1
+ %31 = OpFunctionCall %void %compute_main_inner %32
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.wgsl
new file mode 100644
index 0000000..79b6a47
--- /dev/null
+++ b/test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+enable f16;
+
+var<workgroup> arg_0 : f16;
+
+fn workgroupUniformLoad_e07d08() {
+ var res : f16 = workgroupUniformLoad(&(arg_0));
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ workgroupUniformLoad_e07d08();
+}
diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl b/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl
new file mode 100644
index 0000000..ff17877
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl
@@ -0,0 +1,33 @@
+// Copyright 2023 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/src/cmd/gen
+// using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+var<workgroup> arg_0: u32;
+
+// fn workgroupUniformLoad(ptr<workgroup, u32, read_write>) -> u32
+fn workgroupUniformLoad_37307c() {
+ var res: u32 = workgroupUniformLoad(&arg_0);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ workgroupUniformLoad_37307c();
+}
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
new file mode 100644
index 0000000..26bb61f
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+groupshared uint arg_0;
+
+uint tint_workgroupUniformLoad_arg_0() {
+ GroupMemoryBarrierWithGroupSync();
+ const uint result = arg_0;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+void workgroupUniformLoad_37307c() {
+ uint res = tint_workgroupUniformLoad_arg_0();
+}
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ arg_0 = 0u;
+ }
+ GroupMemoryBarrierWithGroupSync();
+ workgroupUniformLoad_37307c();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
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
new file mode 100644
index 0000000..26bb61f
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.fxc.hlsl
@@ -0,0 +1,30 @@
+groupshared uint arg_0;
+
+uint tint_workgroupUniformLoad_arg_0() {
+ GroupMemoryBarrierWithGroupSync();
+ const uint result = arg_0;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+void workgroupUniformLoad_37307c() {
+ uint res = tint_workgroupUniformLoad_arg_0();
+}
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ arg_0 = 0u;
+ }
+ GroupMemoryBarrierWithGroupSync();
+ workgroupUniformLoad_37307c();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.glsl b/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.glsl
new file mode 100644
index 0000000..5b51b2d
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.glsl
@@ -0,0 +1,27 @@
+#version 310 es
+
+shared uint arg_0;
+uint tint_workgroupUniformLoad_arg_0() {
+ barrier();
+ uint result = arg_0;
+ barrier();
+ return result;
+}
+
+void workgroupUniformLoad_37307c() {
+ uint res = tint_workgroupUniformLoad_arg_0();
+}
+
+void compute_main(uint local_invocation_index) {
+ {
+ arg_0 = 0u;
+ }
+ barrier();
+ workgroupUniformLoad_37307c();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main(gl_LocalInvocationIndex);
+ return;
+}
diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.msl b/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.msl
new file mode 100644
index 0000000..3553390
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.msl
@@ -0,0 +1,28 @@
+#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 workgroupUniformLoad_37307c(threadgroup uint* const tint_symbol) {
+ uint res = tint_workgroupUniformLoad(tint_symbol);
+}
+
+void compute_main_inner(uint local_invocation_index, threadgroup uint* const tint_symbol_1) {
+ {
+ *(tint_symbol_1) = 0u;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ workgroupUniformLoad_37307c(tint_symbol_1);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup uint tint_symbol_2;
+ compute_main_inner(local_invocation_index, &(tint_symbol_2));
+ return;
+}
+
diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.spvasm b/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.spvasm
new file mode 100644
index 0000000..b3755dc
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.spvasm
@@ -0,0 +1,59 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 32
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %local_invocation_index_1 "local_invocation_index_1"
+ OpName %arg_0 "arg_0"
+ OpName %tint_workgroupUniformLoad_arg_0 "tint_workgroupUniformLoad_arg_0"
+ OpName %workgroupUniformLoad_37307c "workgroupUniformLoad_37307c"
+ OpName %res "res"
+ OpName %compute_main_inner "compute_main_inner"
+ OpName %local_invocation_index "local_invocation_index"
+ OpName %compute_main "compute_main"
+ OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+ %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+ %arg_0 = OpVariable %_ptr_Workgroup_uint Workgroup
+ %6 = OpTypeFunction %uint
+ %void = OpTypeVoid
+ %uint_2 = OpConstant %uint 2
+ %uint_264 = OpConstant %uint 264
+ %15 = OpTypeFunction %void
+%_ptr_Function_uint = OpTypePointer Function %uint
+ %21 = OpConstantNull %uint
+ %22 = OpTypeFunction %void %uint
+%tint_workgroupUniformLoad_arg_0 = OpFunction %uint None %6
+ %8 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %13 = OpLoad %uint %arg_0
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ OpReturnValue %13
+ OpFunctionEnd
+%workgroupUniformLoad_37307c = OpFunction %void None %15
+ %17 = OpLabel
+ %res = OpVariable %_ptr_Function_uint Function %21
+ %18 = OpFunctionCall %uint %tint_workgroupUniformLoad_arg_0
+ OpStore %res %18
+ OpReturn
+ OpFunctionEnd
+%compute_main_inner = OpFunction %void None %22
+%local_invocation_index = OpFunctionParameter %uint
+ %25 = OpLabel
+ OpStore %arg_0 %21
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %27 = OpFunctionCall %void %workgroupUniformLoad_37307c
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %15
+ %29 = OpLabel
+ %31 = OpLoad %uint %local_invocation_index_1
+ %30 = OpFunctionCall %void %compute_main_inner %31
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.wgsl b/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.wgsl
new file mode 100644
index 0000000..f4d48fa
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.wgsl
@@ -0,0 +1,10 @@
+var<workgroup> arg_0 : u32;
+
+fn workgroupUniformLoad_37307c() {
+ var res : u32 = workgroupUniformLoad(&(arg_0));
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ workgroupUniformLoad_37307c();
+}
diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl b/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl
new file mode 100644
index 0000000..c5cf95e
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl
@@ -0,0 +1,33 @@
+// Copyright 2023 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/src/cmd/gen
+// using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+var<workgroup> arg_0: f32;
+
+// fn workgroupUniformLoad(ptr<workgroup, f32, read_write>) -> f32
+fn workgroupUniformLoad_7a857c() {
+ var res: f32 = workgroupUniformLoad(&arg_0);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ workgroupUniformLoad_7a857c();
+}
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
new file mode 100644
index 0000000..fe6d27e
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+groupshared float arg_0;
+
+float tint_workgroupUniformLoad_arg_0() {
+ GroupMemoryBarrierWithGroupSync();
+ const float result = arg_0;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+void workgroupUniformLoad_7a857c() {
+ float res = tint_workgroupUniformLoad_arg_0();
+}
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ arg_0 = 0.0f;
+ }
+ GroupMemoryBarrierWithGroupSync();
+ workgroupUniformLoad_7a857c();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
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
new file mode 100644
index 0000000..fe6d27e
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.fxc.hlsl
@@ -0,0 +1,30 @@
+groupshared float arg_0;
+
+float tint_workgroupUniformLoad_arg_0() {
+ GroupMemoryBarrierWithGroupSync();
+ const float result = arg_0;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+void workgroupUniformLoad_7a857c() {
+ float res = tint_workgroupUniformLoad_arg_0();
+}
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ arg_0 = 0.0f;
+ }
+ GroupMemoryBarrierWithGroupSync();
+ workgroupUniformLoad_7a857c();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.glsl b/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.glsl
new file mode 100644
index 0000000..bfdeb87
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.glsl
@@ -0,0 +1,27 @@
+#version 310 es
+
+shared float arg_0;
+float tint_workgroupUniformLoad_arg_0() {
+ barrier();
+ float result = arg_0;
+ barrier();
+ return result;
+}
+
+void workgroupUniformLoad_7a857c() {
+ float res = tint_workgroupUniformLoad_arg_0();
+}
+
+void compute_main(uint local_invocation_index) {
+ {
+ arg_0 = 0.0f;
+ }
+ barrier();
+ workgroupUniformLoad_7a857c();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main(gl_LocalInvocationIndex);
+ return;
+}
diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.msl b/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.msl
new file mode 100644
index 0000000..7637172
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.msl
@@ -0,0 +1,28 @@
+#include <metal_stdlib>
+
+using namespace metal;
+float tint_workgroupUniformLoad(threadgroup float* const p) {
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ float const result = *(p);
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ return result;
+}
+
+void workgroupUniformLoad_7a857c(threadgroup float* const tint_symbol) {
+ float res = tint_workgroupUniformLoad(tint_symbol);
+}
+
+void compute_main_inner(uint local_invocation_index, threadgroup float* const tint_symbol_1) {
+ {
+ *(tint_symbol_1) = 0.0f;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ workgroupUniformLoad_7a857c(tint_symbol_1);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup float tint_symbol_2;
+ compute_main_inner(local_invocation_index, &(tint_symbol_2));
+ return;
+}
+
diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.spvasm b/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.spvasm
new file mode 100644
index 0000000..abd0f6a
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.spvasm
@@ -0,0 +1,60 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 33
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %local_invocation_index_1 "local_invocation_index_1"
+ OpName %arg_0 "arg_0"
+ OpName %tint_workgroupUniformLoad_arg_0 "tint_workgroupUniformLoad_arg_0"
+ OpName %workgroupUniformLoad_7a857c "workgroupUniformLoad_7a857c"
+ OpName %res "res"
+ OpName %compute_main_inner "compute_main_inner"
+ OpName %local_invocation_index "local_invocation_index"
+ OpName %compute_main "compute_main"
+ OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+ %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+ %float = OpTypeFloat 32
+%_ptr_Workgroup_float = OpTypePointer Workgroup %float
+ %arg_0 = OpVariable %_ptr_Workgroup_float Workgroup
+ %7 = OpTypeFunction %float
+ %void = OpTypeVoid
+ %uint_2 = OpConstant %uint 2
+ %uint_264 = OpConstant %uint 264
+ %16 = OpTypeFunction %void
+%_ptr_Function_float = OpTypePointer Function %float
+ %22 = OpConstantNull %float
+ %23 = OpTypeFunction %void %uint
+%tint_workgroupUniformLoad_arg_0 = OpFunction %float None %7
+ %9 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %14 = OpLoad %float %arg_0
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ OpReturnValue %14
+ OpFunctionEnd
+%workgroupUniformLoad_7a857c = OpFunction %void None %16
+ %18 = OpLabel
+ %res = OpVariable %_ptr_Function_float Function %22
+ %19 = OpFunctionCall %float %tint_workgroupUniformLoad_arg_0
+ OpStore %res %19
+ OpReturn
+ OpFunctionEnd
+%compute_main_inner = OpFunction %void None %23
+%local_invocation_index = OpFunctionParameter %uint
+ %26 = OpLabel
+ OpStore %arg_0 %22
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %28 = OpFunctionCall %void %workgroupUniformLoad_7a857c
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %16
+ %30 = OpLabel
+ %32 = OpLoad %uint %local_invocation_index_1
+ %31 = OpFunctionCall %void %compute_main_inner %32
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.wgsl b/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.wgsl
new file mode 100644
index 0000000..2fd46aa
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/7a857c.wgsl.expected.wgsl
@@ -0,0 +1,10 @@
+var<workgroup> arg_0 : f32;
+
+fn workgroupUniformLoad_7a857c() {
+ var res : f32 = workgroupUniformLoad(&(arg_0));
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ workgroupUniformLoad_7a857c();
+}
diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl b/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl
new file mode 100644
index 0000000..6ebb067
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl
@@ -0,0 +1,33 @@
+// Copyright 2023 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/src/cmd/gen
+// using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+var<workgroup> arg_0: i32;
+
+// fn workgroupUniformLoad(ptr<workgroup, i32, read_write>) -> i32
+fn workgroupUniformLoad_9d33de() {
+ var res: i32 = workgroupUniformLoad(&arg_0);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ workgroupUniformLoad_9d33de();
+}
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
new file mode 100644
index 0000000..87cf14c
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+groupshared int arg_0;
+
+int tint_workgroupUniformLoad_arg_0() {
+ GroupMemoryBarrierWithGroupSync();
+ const int result = arg_0;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+void workgroupUniformLoad_9d33de() {
+ int res = tint_workgroupUniformLoad_arg_0();
+}
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ arg_0 = 0;
+ }
+ GroupMemoryBarrierWithGroupSync();
+ workgroupUniformLoad_9d33de();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
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
new file mode 100644
index 0000000..87cf14c
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.fxc.hlsl
@@ -0,0 +1,30 @@
+groupshared int arg_0;
+
+int tint_workgroupUniformLoad_arg_0() {
+ GroupMemoryBarrierWithGroupSync();
+ const int result = arg_0;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+void workgroupUniformLoad_9d33de() {
+ int res = tint_workgroupUniformLoad_arg_0();
+}
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ arg_0 = 0;
+ }
+ GroupMemoryBarrierWithGroupSync();
+ workgroupUniformLoad_9d33de();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.glsl b/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.glsl
new file mode 100644
index 0000000..31ae3f9
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.glsl
@@ -0,0 +1,27 @@
+#version 310 es
+
+shared int arg_0;
+int tint_workgroupUniformLoad_arg_0() {
+ barrier();
+ int result = arg_0;
+ barrier();
+ return result;
+}
+
+void workgroupUniformLoad_9d33de() {
+ int res = tint_workgroupUniformLoad_arg_0();
+}
+
+void compute_main(uint local_invocation_index) {
+ {
+ arg_0 = 0;
+ }
+ barrier();
+ workgroupUniformLoad_9d33de();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main(gl_LocalInvocationIndex);
+ return;
+}
diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.msl b/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.msl
new file mode 100644
index 0000000..4ef9721
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.msl
@@ -0,0 +1,28 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int tint_workgroupUniformLoad(threadgroup int* const p) {
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ int const result = *(p);
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ return result;
+}
+
+void workgroupUniformLoad_9d33de(threadgroup int* const tint_symbol) {
+ int res = tint_workgroupUniformLoad(tint_symbol);
+}
+
+void compute_main_inner(uint local_invocation_index, threadgroup int* const tint_symbol_1) {
+ {
+ *(tint_symbol_1) = 0;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ workgroupUniformLoad_9d33de(tint_symbol_1);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup int tint_symbol_2;
+ compute_main_inner(local_invocation_index, &(tint_symbol_2));
+ return;
+}
+
diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.spvasm b/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.spvasm
new file mode 100644
index 0000000..50b0e3f
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.spvasm
@@ -0,0 +1,60 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 33
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %local_invocation_index_1 "local_invocation_index_1"
+ OpName %arg_0 "arg_0"
+ OpName %tint_workgroupUniformLoad_arg_0 "tint_workgroupUniformLoad_arg_0"
+ OpName %workgroupUniformLoad_9d33de "workgroupUniformLoad_9d33de"
+ OpName %res "res"
+ OpName %compute_main_inner "compute_main_inner"
+ OpName %local_invocation_index "local_invocation_index"
+ OpName %compute_main "compute_main"
+ OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+ %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+ %int = OpTypeInt 32 1
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+ %arg_0 = OpVariable %_ptr_Workgroup_int Workgroup
+ %7 = OpTypeFunction %int
+ %void = OpTypeVoid
+ %uint_2 = OpConstant %uint 2
+ %uint_264 = OpConstant %uint 264
+ %16 = OpTypeFunction %void
+%_ptr_Function_int = OpTypePointer Function %int
+ %22 = OpConstantNull %int
+ %23 = OpTypeFunction %void %uint
+%tint_workgroupUniformLoad_arg_0 = OpFunction %int None %7
+ %9 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %14 = OpLoad %int %arg_0
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ OpReturnValue %14
+ OpFunctionEnd
+%workgroupUniformLoad_9d33de = OpFunction %void None %16
+ %18 = OpLabel
+ %res = OpVariable %_ptr_Function_int Function %22
+ %19 = OpFunctionCall %int %tint_workgroupUniformLoad_arg_0
+ OpStore %res %19
+ OpReturn
+ OpFunctionEnd
+%compute_main_inner = OpFunction %void None %23
+%local_invocation_index = OpFunctionParameter %uint
+ %26 = OpLabel
+ OpStore %arg_0 %22
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %28 = OpFunctionCall %void %workgroupUniformLoad_9d33de
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %16
+ %30 = OpLabel
+ %32 = OpLoad %uint %local_invocation_index_1
+ %31 = OpFunctionCall %void %compute_main_inner %32
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.wgsl b/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.wgsl
new file mode 100644
index 0000000..5228ac7
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/9d33de.wgsl.expected.wgsl
@@ -0,0 +1,10 @@
+var<workgroup> arg_0 : i32;
+
+fn workgroupUniformLoad_9d33de() {
+ var res : i32 = workgroupUniformLoad(&(arg_0));
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ workgroupUniformLoad_9d33de();
+}
diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl b/test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl
new file mode 100644
index 0000000..a0820bb
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl
@@ -0,0 +1,35 @@
+// Copyright 2023 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/src/cmd/gen
+// using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+enable f16;
+var<workgroup> arg_0: f16;
+
+// fn workgroupUniformLoad(ptr<workgroup, f16, read_write>) -> f16
+fn workgroupUniformLoad_e07d08() {
+ var res: f16 = workgroupUniformLoad(&arg_0);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ workgroupUniformLoad_e07d08();
+}
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
new file mode 100644
index 0000000..d34a7c4
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+groupshared float16_t arg_0;
+
+float16_t tint_workgroupUniformLoad_arg_0() {
+ GroupMemoryBarrierWithGroupSync();
+ const float16_t result = arg_0;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+void workgroupUniformLoad_e07d08() {
+ float16_t res = tint_workgroupUniformLoad_arg_0();
+}
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ arg_0 = float16_t(0.0h);
+ }
+ GroupMemoryBarrierWithGroupSync();
+ workgroupUniformLoad_e07d08();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..8b6f290
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl.expected.fxc.hlsl
@@ -0,0 +1,35 @@
+SKIP: FAILED
+
+groupshared float16_t arg_0;
+
+float16_t tint_workgroupUniformLoad_arg_0() {
+ GroupMemoryBarrierWithGroupSync();
+ const float16_t result = arg_0;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+void workgroupUniformLoad_e07d08() {
+ float16_t res = tint_workgroupUniformLoad_arg_0();
+}
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ arg_0 = float16_t(0.0h);
+ }
+ GroupMemoryBarrierWithGroupSync();
+ workgroupUniformLoad_e07d08();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
+
+FXC validation failure:
+T:\tmp\dawn-temp\dawn-src\test\tint\Shader@0x000001D6A00D1F90(1,13-21): error X3000: unrecognized identifier 'float16_t'
diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl.expected.glsl b/test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl.expected.glsl
new file mode 100644
index 0000000..0b882bd
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl.expected.glsl
@@ -0,0 +1,28 @@
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+shared float16_t arg_0;
+float16_t tint_workgroupUniformLoad_arg_0() {
+ barrier();
+ float16_t result = arg_0;
+ barrier();
+ return result;
+}
+
+void workgroupUniformLoad_e07d08() {
+ float16_t res = tint_workgroupUniformLoad_arg_0();
+}
+
+void compute_main(uint local_invocation_index) {
+ {
+ arg_0 = 0.0hf;
+ }
+ barrier();
+ workgroupUniformLoad_e07d08();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main(gl_LocalInvocationIndex);
+ return;
+}
diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl.expected.msl b/test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl.expected.msl
new file mode 100644
index 0000000..c4c40c9
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl.expected.msl
@@ -0,0 +1,28 @@
+#include <metal_stdlib>
+
+using namespace metal;
+half tint_workgroupUniformLoad(threadgroup half* const p) {
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ half const result = *(p);
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ return result;
+}
+
+void workgroupUniformLoad_e07d08(threadgroup half* const tint_symbol) {
+ half res = tint_workgroupUniformLoad(tint_symbol);
+}
+
+void compute_main_inner(uint local_invocation_index, threadgroup half* const tint_symbol_1) {
+ {
+ *(tint_symbol_1) = 0.0h;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ workgroupUniformLoad_e07d08(tint_symbol_1);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup half tint_symbol_2;
+ compute_main_inner(local_invocation_index, &(tint_symbol_2));
+ return;
+}
+
diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl.expected.spvasm b/test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl.expected.spvasm
new file mode 100644
index 0000000..9c5cf13
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl.expected.spvasm
@@ -0,0 +1,64 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 33
+; Schema: 0
+ OpCapability Shader
+ OpCapability Float16
+ OpCapability UniformAndStorageBuffer16BitAccess
+ OpCapability StorageBuffer16BitAccess
+ OpCapability StorageInputOutput16
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %local_invocation_index_1 "local_invocation_index_1"
+ OpName %arg_0 "arg_0"
+ OpName %tint_workgroupUniformLoad_arg_0 "tint_workgroupUniformLoad_arg_0"
+ OpName %workgroupUniformLoad_e07d08 "workgroupUniformLoad_e07d08"
+ OpName %res "res"
+ OpName %compute_main_inner "compute_main_inner"
+ OpName %local_invocation_index "local_invocation_index"
+ OpName %compute_main "compute_main"
+ OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+ %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+ %half = OpTypeFloat 16
+%_ptr_Workgroup_half = OpTypePointer Workgroup %half
+ %arg_0 = OpVariable %_ptr_Workgroup_half Workgroup
+ %7 = OpTypeFunction %half
+ %void = OpTypeVoid
+ %uint_2 = OpConstant %uint 2
+ %uint_264 = OpConstant %uint 264
+ %16 = OpTypeFunction %void
+%_ptr_Function_half = OpTypePointer Function %half
+ %22 = OpConstantNull %half
+ %23 = OpTypeFunction %void %uint
+%tint_workgroupUniformLoad_arg_0 = OpFunction %half None %7
+ %9 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %14 = OpLoad %half %arg_0
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ OpReturnValue %14
+ OpFunctionEnd
+%workgroupUniformLoad_e07d08 = OpFunction %void None %16
+ %18 = OpLabel
+ %res = OpVariable %_ptr_Function_half Function %22
+ %19 = OpFunctionCall %half %tint_workgroupUniformLoad_arg_0
+ OpStore %res %19
+ OpReturn
+ OpFunctionEnd
+%compute_main_inner = OpFunction %void None %23
+%local_invocation_index = OpFunctionParameter %uint
+ %26 = OpLabel
+ OpStore %arg_0 %22
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %28 = OpFunctionCall %void %workgroupUniformLoad_e07d08
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %16
+ %30 = OpLabel
+ %32 = OpLoad %uint %local_invocation_index_1
+ %31 = OpFunctionCall %void %compute_main_inner %32
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl.expected.wgsl b/test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl.expected.wgsl
new file mode 100644
index 0000000..79b6a47
--- /dev/null
+++ b/test/tint/builtins/gen/var/workgroupUniformLoad/e07d08.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+enable f16;
+
+var<workgroup> arg_0 : f16;
+
+fn workgroupUniformLoad_e07d08() {
+ var res : f16 = workgroupUniformLoad(&(arg_0));
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ workgroupUniformLoad_e07d08();
+}
diff --git a/test/tint/builtins/workgroupUniformLoad/array.wgsl b/test/tint/builtins/workgroupUniformLoad/array.wgsl
new file mode 100644
index 0000000..bdc07fc
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/array.wgsl
@@ -0,0 +1,5 @@
+var<workgroup> v : array<i32, 4>;
+
+fn foo() -> array<i32, 4> {
+ return workgroupUniformLoad(&v);
+}
diff --git a/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.dxc.hlsl b/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..a9b6ea5
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.dxc.hlsl
@@ -0,0 +1,19 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+ return;
+}
+
+groupshared int v[4];
+
+typedef int tint_workgroupUniformLoad_v_ret[4];
+tint_workgroupUniformLoad_v_ret tint_workgroupUniformLoad_v() {
+ GroupMemoryBarrierWithGroupSync();
+ const int result[4] = v;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+typedef int foo_ret[4];
+foo_ret foo() {
+ 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
new file mode 100644
index 0000000..a9b6ea5
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.fxc.hlsl
@@ -0,0 +1,19 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+ return;
+}
+
+groupshared int v[4];
+
+typedef int tint_workgroupUniformLoad_v_ret[4];
+tint_workgroupUniformLoad_v_ret tint_workgroupUniformLoad_v() {
+ GroupMemoryBarrierWithGroupSync();
+ const int result[4] = v;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+typedef int foo_ret[4];
+foo_ret foo() {
+ return tint_workgroupUniformLoad_v();
+}
diff --git a/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.glsl b/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.glsl
new file mode 100644
index 0000000..cda9b79
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.glsl
@@ -0,0 +1,18 @@
+#version 310 es
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void unused_entry_point() {
+ return;
+}
+shared int v[4];
+int[4] tint_workgroupUniformLoad_v() {
+ barrier();
+ int result[4] = v;
+ barrier();
+ return result;
+}
+
+int[4] foo() {
+ return tint_workgroupUniformLoad_v();
+}
+
diff --git a/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.msl b/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.msl
new file mode 100644
index 0000000..a349a06
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.msl
@@ -0,0 +1,27 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+tint_array<int, 4> tint_workgroupUniformLoad(threadgroup tint_array<int, 4>* const p) {
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ tint_array<int, 4> const result = *(p);
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ return result;
+}
+
+tint_array<int, 4> foo(threadgroup tint_array<int, 4>* const tint_symbol) {
+ return tint_workgroupUniformLoad(tint_symbol);
+}
+
diff --git a/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.spvasm b/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.spvasm
new file mode 100644
index 0000000..9e7b9d1
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.spvasm
@@ -0,0 +1,41 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 22
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+ OpExecutionMode %unused_entry_point LocalSize 1 1 1
+ OpName %v "v"
+ OpName %unused_entry_point "unused_entry_point"
+ OpName %tint_workgroupUniformLoad_v "tint_workgroupUniformLoad_v"
+ OpName %foo "foo"
+ OpDecorate %_arr_int_uint_4 ArrayStride 4
+ %int = OpTypeInt 32 1
+ %uint = OpTypeInt 32 0
+ %uint_4 = OpConstant %uint 4
+%_arr_int_uint_4 = OpTypeArray %int %uint_4
+%_ptr_Workgroup__arr_int_uint_4 = OpTypePointer Workgroup %_arr_int_uint_4
+ %v = OpVariable %_ptr_Workgroup__arr_int_uint_4 Workgroup
+ %void = OpTypeVoid
+ %7 = OpTypeFunction %void
+ %11 = OpTypeFunction %_arr_int_uint_4
+ %uint_2 = OpConstant %uint 2
+ %uint_264 = OpConstant %uint 264
+%unused_entry_point = OpFunction %void None %7
+ %10 = OpLabel
+ OpReturn
+ OpFunctionEnd
+%tint_workgroupUniformLoad_v = OpFunction %_arr_int_uint_4 None %11
+ %13 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %17 = OpLoad %_arr_int_uint_4 %v
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ OpReturnValue %17
+ OpFunctionEnd
+ %foo = OpFunction %_arr_int_uint_4 None %11
+ %20 = OpLabel
+ %21 = OpFunctionCall %_arr_int_uint_4 %tint_workgroupUniformLoad_v
+ OpReturnValue %21
+ OpFunctionEnd
diff --git a/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.wgsl b/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.wgsl
new file mode 100644
index 0000000..15fde7d
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/array.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+var<workgroup> v : array<i32, 4>;
+
+fn foo() -> array<i32, 4> {
+ return workgroupUniformLoad(&(v));
+}
diff --git a/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl b/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl
new file mode 100644
index 0000000..d1f5bda
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl
@@ -0,0 +1,7 @@
+// flags: --overrides wgsize=64
+override wgsize : i32;
+var<workgroup> v : array<i32, wgsize * 2>;
+
+fn foo() -> i32 {
+ return workgroupUniformLoad(&v)[0];
+}
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
new file mode 100644
index 0000000..a42cbe5
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.dxc.hlsl
@@ -0,0 +1,19 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+ return;
+}
+
+groupshared int v[128];
+
+typedef int tint_workgroupUniformLoad_v_ret[128];
+tint_workgroupUniformLoad_v_ret tint_workgroupUniformLoad_v() {
+ GroupMemoryBarrierWithGroupSync();
+ const int result[128] = v;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+int foo() {
+ 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
new file mode 100644
index 0000000..a42cbe5
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.fxc.hlsl
@@ -0,0 +1,19 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+ return;
+}
+
+groupshared int v[128];
+
+typedef int tint_workgroupUniformLoad_v_ret[128];
+tint_workgroupUniformLoad_v_ret tint_workgroupUniformLoad_v() {
+ GroupMemoryBarrierWithGroupSync();
+ const int result[128] = v;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+int foo() {
+ 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
new file mode 100644
index 0000000..6e551e8
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.glsl
@@ -0,0 +1,19 @@
+#version 310 es
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void unused_entry_point() {
+ return;
+}
+shared int v[128];
+int[128] tint_workgroupUniformLoad_v() {
+ barrier();
+ int result[128] = v;
+ barrier();
+ return result;
+}
+
+int foo() {
+ int tint_symbol[128] = tint_workgroupUniformLoad_v();
+ return tint_symbol[0];
+}
+
diff --git a/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.msl b/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.msl
new file mode 100644
index 0000000..aeee60c
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.msl
@@ -0,0 +1,28 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+tint_array<int, 128> tint_workgroupUniformLoad(threadgroup tint_array<int, 128>* const p) {
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ tint_array<int, 128> const result = *(p);
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ return result;
+}
+
+int foo(threadgroup tint_array<int, 128>* const tint_symbol_1) {
+ tint_array<int, 128> const tint_symbol = tint_workgroupUniformLoad(tint_symbol_1);
+ return tint_symbol[0];
+}
+
diff --git a/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.spvasm b/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.spvasm
new file mode 100644
index 0000000..8789073
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.spvasm
@@ -0,0 +1,44 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 25
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+ OpExecutionMode %unused_entry_point LocalSize 1 1 1
+ OpName %v "v"
+ OpName %unused_entry_point "unused_entry_point"
+ OpName %tint_workgroupUniformLoad_v "tint_workgroupUniformLoad_v"
+ OpName %foo "foo"
+ OpDecorate %_arr_int_uint_128 ArrayStride 4
+ %int = OpTypeInt 32 1
+ %uint = OpTypeInt 32 0
+ %uint_128 = OpConstant %uint 128
+%_arr_int_uint_128 = OpTypeArray %int %uint_128
+%_ptr_Workgroup__arr_int_uint_128 = OpTypePointer Workgroup %_arr_int_uint_128
+ %v = OpVariable %_ptr_Workgroup__arr_int_uint_128 Workgroup
+ %void = OpTypeVoid
+ %7 = OpTypeFunction %void
+ %11 = OpTypeFunction %_arr_int_uint_128
+ %uint_2 = OpConstant %uint 2
+ %uint_264 = OpConstant %uint 264
+ %19 = OpTypeFunction %int
+ %23 = OpConstantNull %int
+%unused_entry_point = OpFunction %void None %7
+ %10 = OpLabel
+ OpReturn
+ OpFunctionEnd
+%tint_workgroupUniformLoad_v = OpFunction %_arr_int_uint_128 None %11
+ %13 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %17 = OpLoad %_arr_int_uint_128 %v
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ OpReturnValue %17
+ OpFunctionEnd
+ %foo = OpFunction %int None %19
+ %21 = OpLabel
+ %22 = OpFunctionCall %_arr_int_uint_128 %tint_workgroupUniformLoad_v
+ %24 = OpCompositeExtract %int %22 0
+ OpReturnValue %24
+ OpFunctionEnd
diff --git a/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.wgsl b/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.wgsl
new file mode 100644
index 0000000..3afe732
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/array_overridable_count.wgsl.expected.wgsl
@@ -0,0 +1,7 @@
+const wgsize : i32 = 64i;
+
+var<workgroup> v : array<i32, (wgsize * 2)>;
+
+fn foo() -> i32 {
+ return workgroupUniformLoad(&(v))[0];
+}
diff --git a/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl b/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl
new file mode 100644
index 0000000..f9d42d8
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl
@@ -0,0 +1,8 @@
+// flags: --overrides wgsize=64
+override wgsize : i32;
+type Array = array<i32, wgsize * 2>;
+var<workgroup> v : Array;
+
+fn foo() -> i32 {
+ return workgroupUniformLoad(&v)[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
new file mode 100644
index 0000000..a42cbe5
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.dxc.hlsl
@@ -0,0 +1,19 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+ return;
+}
+
+groupshared int v[128];
+
+typedef int tint_workgroupUniformLoad_v_ret[128];
+tint_workgroupUniformLoad_v_ret tint_workgroupUniformLoad_v() {
+ GroupMemoryBarrierWithGroupSync();
+ const int result[128] = v;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+int foo() {
+ 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
new file mode 100644
index 0000000..a42cbe5
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.fxc.hlsl
@@ -0,0 +1,19 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+ return;
+}
+
+groupshared int v[128];
+
+typedef int tint_workgroupUniformLoad_v_ret[128];
+tint_workgroupUniformLoad_v_ret tint_workgroupUniformLoad_v() {
+ GroupMemoryBarrierWithGroupSync();
+ const int result[128] = v;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+int foo() {
+ 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
new file mode 100644
index 0000000..6e551e8
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.glsl
@@ -0,0 +1,19 @@
+#version 310 es
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void unused_entry_point() {
+ return;
+}
+shared int v[128];
+int[128] tint_workgroupUniformLoad_v() {
+ barrier();
+ int result[128] = v;
+ barrier();
+ return result;
+}
+
+int foo() {
+ int tint_symbol[128] = tint_workgroupUniformLoad_v();
+ return tint_symbol[0];
+}
+
diff --git a/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.msl b/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.msl
new file mode 100644
index 0000000..aeee60c
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.msl
@@ -0,0 +1,28 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+tint_array<int, 128> tint_workgroupUniformLoad(threadgroup tint_array<int, 128>* const p) {
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ tint_array<int, 128> const result = *(p);
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ return result;
+}
+
+int foo(threadgroup tint_array<int, 128>* const tint_symbol_1) {
+ tint_array<int, 128> const tint_symbol = tint_workgroupUniformLoad(tint_symbol_1);
+ return tint_symbol[0];
+}
+
diff --git a/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.spvasm b/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.spvasm
new file mode 100644
index 0000000..8789073
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.spvasm
@@ -0,0 +1,44 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 25
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+ OpExecutionMode %unused_entry_point LocalSize 1 1 1
+ OpName %v "v"
+ OpName %unused_entry_point "unused_entry_point"
+ OpName %tint_workgroupUniformLoad_v "tint_workgroupUniformLoad_v"
+ OpName %foo "foo"
+ OpDecorate %_arr_int_uint_128 ArrayStride 4
+ %int = OpTypeInt 32 1
+ %uint = OpTypeInt 32 0
+ %uint_128 = OpConstant %uint 128
+%_arr_int_uint_128 = OpTypeArray %int %uint_128
+%_ptr_Workgroup__arr_int_uint_128 = OpTypePointer Workgroup %_arr_int_uint_128
+ %v = OpVariable %_ptr_Workgroup__arr_int_uint_128 Workgroup
+ %void = OpTypeVoid
+ %7 = OpTypeFunction %void
+ %11 = OpTypeFunction %_arr_int_uint_128
+ %uint_2 = OpConstant %uint 2
+ %uint_264 = OpConstant %uint 264
+ %19 = OpTypeFunction %int
+ %23 = OpConstantNull %int
+%unused_entry_point = OpFunction %void None %7
+ %10 = OpLabel
+ OpReturn
+ OpFunctionEnd
+%tint_workgroupUniformLoad_v = OpFunction %_arr_int_uint_128 None %11
+ %13 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %17 = OpLoad %_arr_int_uint_128 %v
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ OpReturnValue %17
+ OpFunctionEnd
+ %foo = OpFunction %int None %19
+ %21 = OpLabel
+ %22 = OpFunctionCall %_arr_int_uint_128 %tint_workgroupUniformLoad_v
+ %24 = OpCompositeExtract %int %22 0
+ OpReturnValue %24
+ OpFunctionEnd
diff --git a/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.wgsl b/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.wgsl
new file mode 100644
index 0000000..df6498e
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/array_overridable_count_aliased.wgsl.expected.wgsl
@@ -0,0 +1,9 @@
+const wgsize : i32 = 64i;
+
+type Array = array<i32, (wgsize * 2)>;
+
+var<workgroup> v : Array;
+
+fn foo() -> i32 {
+ return workgroupUniformLoad(&(v))[0];
+}
diff --git a/test/tint/builtins/workgroupUniformLoad/bool.wgsl b/test/tint/builtins/workgroupUniformLoad/bool.wgsl
new file mode 100644
index 0000000..1fb9c02
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/bool.wgsl
@@ -0,0 +1,5 @@
+var<workgroup> v : bool;
+
+fn foo() -> bool {
+ return workgroupUniformLoad(&v);
+}
diff --git a/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.dxc.hlsl b/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..afb3e79
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.dxc.hlsl
@@ -0,0 +1,17 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+ return;
+}
+
+groupshared bool v;
+
+bool tint_workgroupUniformLoad_v() {
+ GroupMemoryBarrierWithGroupSync();
+ const bool result = v;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+bool foo() {
+ 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
new file mode 100644
index 0000000..afb3e79
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.fxc.hlsl
@@ -0,0 +1,17 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+ return;
+}
+
+groupshared bool v;
+
+bool tint_workgroupUniformLoad_v() {
+ GroupMemoryBarrierWithGroupSync();
+ const bool result = v;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+bool foo() {
+ return tint_workgroupUniformLoad_v();
+}
diff --git a/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.glsl b/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.glsl
new file mode 100644
index 0000000..4962454
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.glsl
@@ -0,0 +1,18 @@
+#version 310 es
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void unused_entry_point() {
+ return;
+}
+shared bool v;
+bool tint_workgroupUniformLoad_v() {
+ barrier();
+ bool result = v;
+ barrier();
+ return result;
+}
+
+bool foo() {
+ return tint_workgroupUniformLoad_v();
+}
+
diff --git a/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.msl b/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.msl
new file mode 100644
index 0000000..f593e45
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.msl
@@ -0,0 +1,14 @@
+#include <metal_stdlib>
+
+using namespace metal;
+bool tint_workgroupUniformLoad(threadgroup bool* const p) {
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ bool const result = *(p);
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ return result;
+}
+
+bool foo(threadgroup bool* const tint_symbol) {
+ return tint_workgroupUniformLoad(tint_symbol);
+}
+
diff --git a/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.spvasm b/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.spvasm
new file mode 100644
index 0000000..df9773b
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.spvasm
@@ -0,0 +1,38 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 20
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+ OpExecutionMode %unused_entry_point LocalSize 1 1 1
+ OpName %v "v"
+ OpName %unused_entry_point "unused_entry_point"
+ OpName %tint_workgroupUniformLoad_v "tint_workgroupUniformLoad_v"
+ OpName %foo "foo"
+ %bool = OpTypeBool
+%_ptr_Workgroup_bool = OpTypePointer Workgroup %bool
+ %v = OpVariable %_ptr_Workgroup_bool Workgroup
+ %void = OpTypeVoid
+ %4 = OpTypeFunction %void
+ %8 = OpTypeFunction %bool
+ %uint = OpTypeInt 32 0
+ %uint_2 = OpConstant %uint 2
+ %uint_264 = OpConstant %uint 264
+%unused_entry_point = OpFunction %void None %4
+ %7 = OpLabel
+ OpReturn
+ OpFunctionEnd
+%tint_workgroupUniformLoad_v = OpFunction %bool None %8
+ %10 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %15 = OpLoad %bool %v
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ OpReturnValue %15
+ OpFunctionEnd
+ %foo = OpFunction %bool None %8
+ %18 = OpLabel
+ %19 = OpFunctionCall %bool %tint_workgroupUniformLoad_v
+ OpReturnValue %19
+ OpFunctionEnd
diff --git a/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.wgsl b/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.wgsl
new file mode 100644
index 0000000..f5a2a83
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/bool.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+var<workgroup> v : bool;
+
+fn foo() -> bool {
+ return workgroupUniformLoad(&(v));
+}
diff --git a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl
new file mode 100644
index 0000000..5de56d5
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl
@@ -0,0 +1,7 @@
+var<workgroup> a : i32;
+var<workgroup> b : i32;
+
+fn foo() {
+ for (var i = 0; i < workgroupUniformLoad(&a); i += workgroupUniformLoad(&b)) {
+ }
+}
diff --git a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.dxc.hlsl b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..ec23355
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.dxc.hlsl
@@ -0,0 +1,42 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+ return;
+}
+
+groupshared int a;
+
+int tint_workgroupUniformLoad_a() {
+ GroupMemoryBarrierWithGroupSync();
+ const int result = a;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+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();
+ if (!((tint_symbol < tint_symbol_1))) {
+ break;
+ }
+ {
+ }
+ {
+ const int tint_symbol_2 = i;
+ 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
new file mode 100644
index 0000000..ec23355
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.fxc.hlsl
@@ -0,0 +1,42 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+ return;
+}
+
+groupshared int a;
+
+int tint_workgroupUniformLoad_a() {
+ GroupMemoryBarrierWithGroupSync();
+ const int result = a;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+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();
+ if (!((tint_symbol < tint_symbol_1))) {
+ break;
+ }
+ {
+ }
+ {
+ const int tint_symbol_2 = i;
+ 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
new file mode 100644
index 0000000..09c4abe
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.glsl
@@ -0,0 +1,42 @@
+#version 310 es
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void unused_entry_point() {
+ return;
+}
+shared int a;
+int tint_workgroupUniformLoad_a() {
+ barrier();
+ int result = a;
+ barrier();
+ return result;
+}
+
+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();
+ if (!((tint_symbol < tint_symbol_1))) {
+ break;
+ }
+ {
+ }
+ {
+ int tint_symbol_2 = i;
+ 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.msl b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.msl
new file mode 100644
index 0000000..57923b6
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.msl
@@ -0,0 +1,30 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int tint_workgroupUniformLoad(threadgroup int* const p) {
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ int const result = *(p);
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ return result;
+}
+
+void foo(threadgroup int* const tint_symbol_4, threadgroup int* const tint_symbol_5) {
+ {
+ int i = 0;
+ while (true) {
+ int const tint_symbol = i;
+ int const tint_symbol_1 = tint_workgroupUniformLoad(tint_symbol_4);
+ if (!((tint_symbol < tint_symbol_1))) {
+ break;
+ }
+ {
+ }
+ {
+ int const tint_symbol_2 = i;
+ int const tint_symbol_3 = tint_workgroupUniformLoad(tint_symbol_5);
+ i = as_type<int>((as_type<uint>(tint_symbol_2) + as_type<uint>(tint_symbol_3)));
+ }
+ }
+ }
+}
+
diff --git a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.spvasm b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.spvasm
new file mode 100644
index 0000000..55f1f12
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.spvasm
@@ -0,0 +1,75 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 42
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+ OpExecutionMode %unused_entry_point LocalSize 1 1 1
+ OpName %a "a"
+ OpName %b "b"
+ OpName %unused_entry_point "unused_entry_point"
+ OpName %tint_workgroupUniformLoad_a "tint_workgroupUniformLoad_a"
+ OpName %tint_workgroupUniformLoad_b "tint_workgroupUniformLoad_b"
+ OpName %foo "foo"
+ OpName %i "i"
+ %int = OpTypeInt 32 1
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+ %a = OpVariable %_ptr_Workgroup_int Workgroup
+ %b = OpVariable %_ptr_Workgroup_int Workgroup
+ %void = OpTypeVoid
+ %5 = OpTypeFunction %void
+ %9 = OpTypeFunction %int
+ %uint = OpTypeInt 32 0
+ %uint_2 = OpConstant %uint 2
+ %uint_264 = OpConstant %uint 264
+ %25 = OpConstantNull %int
+%_ptr_Function_int = OpTypePointer Function %int
+ %bool = OpTypeBool
+%unused_entry_point = OpFunction %void None %5
+ %8 = OpLabel
+ OpReturn
+ OpFunctionEnd
+%tint_workgroupUniformLoad_a = OpFunction %int None %9
+ %11 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %16 = OpLoad %int %a
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ OpReturnValue %16
+ OpFunctionEnd
+%tint_workgroupUniformLoad_b = OpFunction %int None %9
+ %19 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %21 = OpLoad %int %b
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ OpReturnValue %21
+ OpFunctionEnd
+ %foo = OpFunction %void None %5
+ %24 = OpLabel
+ %i = OpVariable %_ptr_Function_int Function %25
+ OpStore %i %25
+ OpBranch %28
+ %28 = OpLabel
+ OpLoopMerge %29 %30 None
+ OpBranch %31
+ %31 = OpLabel
+ %32 = OpLoad %int %i
+ %33 = OpFunctionCall %int %tint_workgroupUniformLoad_a
+ %35 = OpSLessThan %bool %32 %33
+ %34 = OpLogicalNot %bool %35
+ OpSelectionMerge %37 None
+ OpBranchConditional %34 %38 %37
+ %38 = OpLabel
+ OpBranch %29
+ %37 = OpLabel
+ OpBranch %30
+ %30 = OpLabel
+ %39 = OpLoad %int %i
+ %40 = OpFunctionCall %int %tint_workgroupUniformLoad_b
+ %41 = OpIAdd %int %39 %40
+ OpStore %i %41
+ OpBranch %28
+ %29 = OpLabel
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.wgsl b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.wgsl
new file mode 100644
index 0000000..1de7b54
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+var<workgroup> a : i32;
+
+var<workgroup> b : i32;
+
+fn foo() {
+ for(var i = 0; (i < workgroupUniformLoad(&(a))); i += workgroupUniformLoad(&(b))) {
+ }
+}
diff --git a/test/tint/builtins/workgroupUniformLoad/for_loops.wgsl.expected.msl b/test/tint/builtins/workgroupUniformLoad/for_loops.wgsl.expected.msl
new file mode 100644
index 0000000..57923b6
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/for_loops.wgsl.expected.msl
@@ -0,0 +1,30 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int tint_workgroupUniformLoad(threadgroup int* const p) {
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ int const result = *(p);
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ return result;
+}
+
+void foo(threadgroup int* const tint_symbol_4, threadgroup int* const tint_symbol_5) {
+ {
+ int i = 0;
+ while (true) {
+ int const tint_symbol = i;
+ int const tint_symbol_1 = tint_workgroupUniformLoad(tint_symbol_4);
+ if (!((tint_symbol < tint_symbol_1))) {
+ break;
+ }
+ {
+ }
+ {
+ int const tint_symbol_2 = i;
+ int const tint_symbol_3 = tint_workgroupUniformLoad(tint_symbol_5);
+ i = as_type<int>((as_type<uint>(tint_symbol_2) + as_type<uint>(tint_symbol_3)));
+ }
+ }
+ }
+}
+
diff --git a/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl b/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl
new file mode 100644
index 0000000..8592270
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl
@@ -0,0 +1,8 @@
+var<workgroup> v : bool;
+
+fn foo() -> i32 {
+ if (workgroupUniformLoad(&v)) {
+ return 42;
+ }
+ return 0;
+}
diff --git a/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.dxc.hlsl b/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..48a589f
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.dxc.hlsl
@@ -0,0 +1,20 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+ return;
+}
+
+groupshared bool v;
+
+bool tint_workgroupUniformLoad_v() {
+ GroupMemoryBarrierWithGroupSync();
+ const bool result = v;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+int foo() {
+ 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
new file mode 100644
index 0000000..48a589f
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.fxc.hlsl
@@ -0,0 +1,20 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+ return;
+}
+
+groupshared bool v;
+
+bool tint_workgroupUniformLoad_v() {
+ GroupMemoryBarrierWithGroupSync();
+ const bool result = v;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+int foo() {
+ 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
new file mode 100644
index 0000000..ad15d25
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.glsl
@@ -0,0 +1,21 @@
+#version 310 es
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void unused_entry_point() {
+ return;
+}
+shared bool v;
+bool tint_workgroupUniformLoad_v() {
+ barrier();
+ bool result = v;
+ barrier();
+ return result;
+}
+
+int foo() {
+ if (tint_workgroupUniformLoad_v()) {
+ return 42;
+ }
+ return 0;
+}
+
diff --git a/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.msl b/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.msl
new file mode 100644
index 0000000..a55470c
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.msl
@@ -0,0 +1,17 @@
+#include <metal_stdlib>
+
+using namespace metal;
+bool tint_workgroupUniformLoad(threadgroup bool* const p) {
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ bool const result = *(p);
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ return result;
+}
+
+int foo(threadgroup bool* const tint_symbol) {
+ if (tint_workgroupUniformLoad(tint_symbol)) {
+ return 42;
+ }
+ return 0;
+}
+
diff --git a/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.spvasm b/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.spvasm
new file mode 100644
index 0000000..083d997
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.spvasm
@@ -0,0 +1,67 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 37
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+ OpExecutionMode %unused_entry_point LocalSize 1 1 1
+ OpName %v "v"
+ OpName %unused_entry_point "unused_entry_point"
+ OpName %tint_workgroupUniformLoad_v "tint_workgroupUniformLoad_v"
+ OpName %foo "foo"
+ OpName %tint_return_flag "tint_return_flag"
+ OpName %tint_return_value "tint_return_value"
+ %bool = OpTypeBool
+%_ptr_Workgroup_bool = OpTypePointer Workgroup %bool
+ %v = OpVariable %_ptr_Workgroup_bool Workgroup
+ %void = OpTypeVoid
+ %4 = OpTypeFunction %void
+ %8 = OpTypeFunction %bool
+ %uint = OpTypeInt 32 0
+ %uint_2 = OpConstant %uint 2
+ %uint_264 = OpConstant %uint 264
+ %int = OpTypeInt 32 1
+ %17 = OpTypeFunction %int
+%_ptr_Function_bool = OpTypePointer Function %bool
+ %23 = OpConstantNull %bool
+%_ptr_Function_int = OpTypePointer Function %int
+ %26 = OpConstantNull %int
+ %true = OpConstantTrue %bool
+ %int_42 = OpConstant %int 42
+%unused_entry_point = OpFunction %void None %4
+ %7 = OpLabel
+ OpReturn
+ OpFunctionEnd
+%tint_workgroupUniformLoad_v = OpFunction %bool None %8
+ %10 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %15 = OpLoad %bool %v
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ OpReturnValue %15
+ OpFunctionEnd
+ %foo = OpFunction %int None %17
+ %20 = OpLabel
+%tint_return_flag = OpVariable %_ptr_Function_bool Function %23
+%tint_return_value = OpVariable %_ptr_Function_int Function %26
+ %27 = OpFunctionCall %bool %tint_workgroupUniformLoad_v
+ OpSelectionMerge %28 None
+ OpBranchConditional %27 %29 %28
+ %29 = OpLabel
+ OpStore %tint_return_flag %true
+ OpStore %tint_return_value %int_42
+ OpBranch %28
+ %28 = OpLabel
+ %33 = OpLoad %bool %tint_return_flag
+ %32 = OpLogicalNot %bool %33
+ OpSelectionMerge %34 None
+ OpBranchConditional %32 %35 %34
+ %35 = OpLabel
+ OpStore %tint_return_flag %true
+ OpStore %tint_return_value %26
+ OpBranch %34
+ %34 = OpLabel
+ %36 = OpLoad %int %tint_return_value
+ OpReturnValue %36
+ OpFunctionEnd
diff --git a/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.wgsl b/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.wgsl
new file mode 100644
index 0000000..8e5a28b
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/if_condition.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+var<workgroup> v : bool;
+
+fn foo() -> i32 {
+ if (workgroupUniformLoad(&(v))) {
+ return 42;
+ }
+ return 0;
+}
diff --git a/test/tint/builtins/workgroupUniformLoad/matrix.wgsl b/test/tint/builtins/workgroupUniformLoad/matrix.wgsl
new file mode 100644
index 0000000..97cf5ac
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/matrix.wgsl
@@ -0,0 +1,5 @@
+var<workgroup> v : mat3x3<f32>;
+
+fn foo() -> mat3x3<f32> {
+ return workgroupUniformLoad(&v);
+}
diff --git a/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.dxc.hlsl b/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..1b22b46
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.dxc.hlsl
@@ -0,0 +1,17 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+ return;
+}
+
+groupshared float3x3 v;
+
+float3x3 tint_workgroupUniformLoad_v() {
+ GroupMemoryBarrierWithGroupSync();
+ const float3x3 result = v;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+float3x3 foo() {
+ 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
new file mode 100644
index 0000000..1b22b46
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.fxc.hlsl
@@ -0,0 +1,17 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+ return;
+}
+
+groupshared float3x3 v;
+
+float3x3 tint_workgroupUniformLoad_v() {
+ GroupMemoryBarrierWithGroupSync();
+ const float3x3 result = v;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+float3x3 foo() {
+ return tint_workgroupUniformLoad_v();
+}
diff --git a/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.glsl b/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.glsl
new file mode 100644
index 0000000..ac70b34a
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.glsl
@@ -0,0 +1,18 @@
+#version 310 es
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void unused_entry_point() {
+ return;
+}
+shared mat3 v;
+mat3 tint_workgroupUniformLoad_v() {
+ barrier();
+ mat3 result = v;
+ barrier();
+ return result;
+}
+
+mat3 foo() {
+ return tint_workgroupUniformLoad_v();
+}
+
diff --git a/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.msl b/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.msl
new file mode 100644
index 0000000..c37c301
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.msl
@@ -0,0 +1,14 @@
+#include <metal_stdlib>
+
+using namespace metal;
+float3x3 tint_workgroupUniformLoad(threadgroup float3x3* const p) {
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ float3x3 const result = *(p);
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ return result;
+}
+
+float3x3 foo(threadgroup float3x3* const tint_symbol) {
+ return tint_workgroupUniformLoad(tint_symbol);
+}
+
diff --git a/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.spvasm b/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.spvasm
new file mode 100644
index 0000000..e77b58f
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.spvasm
@@ -0,0 +1,40 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 22
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+ OpExecutionMode %unused_entry_point LocalSize 1 1 1
+ OpName %v "v"
+ OpName %unused_entry_point "unused_entry_point"
+ OpName %tint_workgroupUniformLoad_v "tint_workgroupUniformLoad_v"
+ OpName %foo "foo"
+ %float = OpTypeFloat 32
+ %v3float = OpTypeVector %float 3
+%mat3v3float = OpTypeMatrix %v3float 3
+%_ptr_Workgroup_mat3v3float = OpTypePointer Workgroup %mat3v3float
+ %v = OpVariable %_ptr_Workgroup_mat3v3float Workgroup
+ %void = OpTypeVoid
+ %6 = OpTypeFunction %void
+ %10 = OpTypeFunction %mat3v3float
+ %uint = OpTypeInt 32 0
+ %uint_2 = OpConstant %uint 2
+ %uint_264 = OpConstant %uint 264
+%unused_entry_point = OpFunction %void None %6
+ %9 = OpLabel
+ OpReturn
+ OpFunctionEnd
+%tint_workgroupUniformLoad_v = OpFunction %mat3v3float None %10
+ %12 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %17 = OpLoad %mat3v3float %v
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ OpReturnValue %17
+ OpFunctionEnd
+ %foo = OpFunction %mat3v3float None %10
+ %20 = OpLabel
+ %21 = OpFunctionCall %mat3v3float %tint_workgroupUniformLoad_v
+ OpReturnValue %21
+ OpFunctionEnd
diff --git a/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.wgsl b/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.wgsl
new file mode 100644
index 0000000..4f4dbb0
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/matrix.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+var<workgroup> v : mat3x3<f32>;
+
+fn foo() -> mat3x3<f32> {
+ return workgroupUniformLoad(&(v));
+}
diff --git a/test/tint/builtins/workgroupUniformLoad/structures.wgsl b/test/tint/builtins/workgroupUniformLoad/structures.wgsl
new file mode 100644
index 0000000..d8ae634
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/structures.wgsl
@@ -0,0 +1,15 @@
+struct Inner {
+ b : bool,
+ v : vec4<i32>,
+ m : mat3x3<f32>,
+}
+
+struct Outer {
+ a : array<Inner, 4>,
+}
+
+var<workgroup> v : Outer;
+
+fn foo() -> Outer {
+ return workgroupUniformLoad(&v);
+}
diff --git a/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.dxc.hlsl b/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..eb3a52c
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.dxc.hlsl
@@ -0,0 +1,26 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+ return;
+}
+
+struct Inner {
+ bool b;
+ int4 v;
+ float3x3 m;
+};
+struct Outer {
+ Inner a[4];
+};
+
+groupshared Outer v;
+
+Outer tint_workgroupUniformLoad_v() {
+ GroupMemoryBarrierWithGroupSync();
+ const Outer result = v;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+Outer foo() {
+ 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
new file mode 100644
index 0000000..eb3a52c
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.fxc.hlsl
@@ -0,0 +1,26 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+ return;
+}
+
+struct Inner {
+ bool b;
+ int4 v;
+ float3x3 m;
+};
+struct Outer {
+ Inner a[4];
+};
+
+groupshared Outer v;
+
+Outer tint_workgroupUniformLoad_v() {
+ GroupMemoryBarrierWithGroupSync();
+ const Outer result = v;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+Outer foo() {
+ return tint_workgroupUniformLoad_v();
+}
diff --git a/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.glsl b/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.glsl
new file mode 100644
index 0000000..7f9766a
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.glsl
@@ -0,0 +1,28 @@
+#version 310 es
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void unused_entry_point() {
+ return;
+}
+struct Inner {
+ bool b;
+ ivec4 v;
+ mat3 m;
+};
+
+struct Outer {
+ Inner a[4];
+};
+
+shared Outer v;
+Outer tint_workgroupUniformLoad_v() {
+ barrier();
+ Outer result = v;
+ barrier();
+ return result;
+}
+
+Outer foo() {
+ return tint_workgroupUniformLoad_v();
+}
+
diff --git a/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.msl b/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.msl
new file mode 100644
index 0000000..09b35ef
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.msl
@@ -0,0 +1,37 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+struct Inner {
+ bool b;
+ int4 v;
+ float3x3 m;
+};
+
+struct Outer {
+ tint_array<Inner, 4> a;
+};
+
+Outer tint_workgroupUniformLoad(threadgroup Outer* const p) {
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ Outer const result = *(p);
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ return result;
+}
+
+Outer foo(threadgroup Outer* const tint_symbol) {
+ return tint_workgroupUniformLoad(tint_symbol);
+}
+
diff --git a/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.spvasm b/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.spvasm
new file mode 100644
index 0000000..1d5be36
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.spvasm
@@ -0,0 +1,60 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 29
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+ OpExecutionMode %unused_entry_point LocalSize 1 1 1
+ OpName %Outer "Outer"
+ OpMemberName %Outer 0 "a"
+ OpName %Inner "Inner"
+ OpMemberName %Inner 0 "b"
+ OpMemberName %Inner 1 "v"
+ OpMemberName %Inner 2 "m"
+ OpName %v "v"
+ OpName %unused_entry_point "unused_entry_point"
+ OpName %tint_workgroupUniformLoad_v "tint_workgroupUniformLoad_v"
+ OpName %foo "foo"
+ OpMemberDecorate %Outer 0 Offset 0
+ OpMemberDecorate %Inner 0 Offset 0
+ OpMemberDecorate %Inner 1 Offset 16
+ OpMemberDecorate %Inner 2 Offset 32
+ OpMemberDecorate %Inner 2 ColMajor
+ OpMemberDecorate %Inner 2 MatrixStride 16
+ OpDecorate %_arr_Inner_uint_4 ArrayStride 80
+ %bool = OpTypeBool
+ %int = OpTypeInt 32 1
+ %v4int = OpTypeVector %int 4
+ %float = OpTypeFloat 32
+ %v3float = OpTypeVector %float 3
+%mat3v3float = OpTypeMatrix %v3float 3
+ %Inner = OpTypeStruct %bool %v4int %mat3v3float
+ %uint = OpTypeInt 32 0
+ %uint_4 = OpConstant %uint 4
+%_arr_Inner_uint_4 = OpTypeArray %Inner %uint_4
+ %Outer = OpTypeStruct %_arr_Inner_uint_4
+%_ptr_Workgroup_Outer = OpTypePointer Workgroup %Outer
+ %v = OpVariable %_ptr_Workgroup_Outer Workgroup
+ %void = OpTypeVoid
+ %14 = OpTypeFunction %void
+ %18 = OpTypeFunction %Outer
+ %uint_2 = OpConstant %uint 2
+ %uint_264 = OpConstant %uint 264
+%unused_entry_point = OpFunction %void None %14
+ %17 = OpLabel
+ OpReturn
+ OpFunctionEnd
+%tint_workgroupUniformLoad_v = OpFunction %Outer None %18
+ %20 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %24 = OpLoad %Outer %v
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ OpReturnValue %24
+ OpFunctionEnd
+ %foo = OpFunction %Outer None %18
+ %27 = OpLabel
+ %28 = OpFunctionCall %Outer %tint_workgroupUniformLoad_v
+ OpReturnValue %28
+ OpFunctionEnd
diff --git a/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.wgsl b/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.wgsl
new file mode 100644
index 0000000..ece72ad
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.wgsl
@@ -0,0 +1,15 @@
+struct Inner {
+ b : bool,
+ v : vec4<i32>,
+ m : mat3x3<f32>,
+}
+
+struct Outer {
+ a : array<Inner, 4>,
+}
+
+var<workgroup> v : Outer;
+
+fn foo() -> Outer {
+ return workgroupUniformLoad(&(v));
+}
diff --git a/test/tint/builtins/workgroupUniformLoad/vec.wgsl b/test/tint/builtins/workgroupUniformLoad/vec.wgsl
new file mode 100644
index 0000000..ca04431
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/vec.wgsl
@@ -0,0 +1,5 @@
+var<workgroup> v : vec4<f32>;
+
+fn foo() -> vec4<f32> {
+ return workgroupUniformLoad(&v);
+}
diff --git a/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.dxc.hlsl b/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..80cd6df
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.dxc.hlsl
@@ -0,0 +1,17 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+ return;
+}
+
+groupshared float4 v;
+
+float4 tint_workgroupUniformLoad_v() {
+ GroupMemoryBarrierWithGroupSync();
+ const float4 result = v;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+float4 foo() {
+ 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
new file mode 100644
index 0000000..80cd6df
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.fxc.hlsl
@@ -0,0 +1,17 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+ return;
+}
+
+groupshared float4 v;
+
+float4 tint_workgroupUniformLoad_v() {
+ GroupMemoryBarrierWithGroupSync();
+ const float4 result = v;
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+float4 foo() {
+ return tint_workgroupUniformLoad_v();
+}
diff --git a/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.glsl b/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.glsl
new file mode 100644
index 0000000..f4dc643
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.glsl
@@ -0,0 +1,18 @@
+#version 310 es
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void unused_entry_point() {
+ return;
+}
+shared vec4 v;
+vec4 tint_workgroupUniformLoad_v() {
+ barrier();
+ vec4 result = v;
+ barrier();
+ return result;
+}
+
+vec4 foo() {
+ return tint_workgroupUniformLoad_v();
+}
+
diff --git a/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.msl b/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.msl
new file mode 100644
index 0000000..29f6e72
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.msl
@@ -0,0 +1,14 @@
+#include <metal_stdlib>
+
+using namespace metal;
+float4 tint_workgroupUniformLoad(threadgroup float4* const p) {
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ float4 const result = *(p);
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ return result;
+}
+
+float4 foo(threadgroup float4* const tint_symbol) {
+ return tint_workgroupUniformLoad(tint_symbol);
+}
+
diff --git a/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.spvasm b/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.spvasm
new file mode 100644
index 0000000..8b2ea4e
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.spvasm
@@ -0,0 +1,39 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 21
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+ OpExecutionMode %unused_entry_point LocalSize 1 1 1
+ OpName %v "v"
+ OpName %unused_entry_point "unused_entry_point"
+ OpName %tint_workgroupUniformLoad_v "tint_workgroupUniformLoad_v"
+ OpName %foo "foo"
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+%_ptr_Workgroup_v4float = OpTypePointer Workgroup %v4float
+ %v = OpVariable %_ptr_Workgroup_v4float Workgroup
+ %void = OpTypeVoid
+ %5 = OpTypeFunction %void
+ %9 = OpTypeFunction %v4float
+ %uint = OpTypeInt 32 0
+ %uint_2 = OpConstant %uint 2
+ %uint_264 = OpConstant %uint 264
+%unused_entry_point = OpFunction %void None %5
+ %8 = OpLabel
+ OpReturn
+ OpFunctionEnd
+%tint_workgroupUniformLoad_v = OpFunction %v4float None %9
+ %11 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %16 = OpLoad %v4float %v
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ OpReturnValue %16
+ OpFunctionEnd
+ %foo = OpFunction %v4float None %9
+ %19 = OpLabel
+ %20 = OpFunctionCall %v4float %tint_workgroupUniformLoad_v
+ OpReturnValue %20
+ OpFunctionEnd
diff --git a/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.wgsl b/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.wgsl
new file mode 100644
index 0000000..e3eea9b
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/vec.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+var<workgroup> v : vec4<f32>;
+
+fn foo() -> vec4<f32> {
+ return workgroupUniformLoad(&(v));
+}
diff --git a/test/tint/builtins/workgroupUniformLoad/via_param.wgsl b/test/tint/builtins/workgroupUniformLoad/via_param.wgsl
new file mode 100644
index 0000000..59f5e8b
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/via_param.wgsl
@@ -0,0 +1,11 @@
+enable chromium_experimental_full_ptr_parameters;
+
+var<workgroup> v : array<i32, 4>;
+
+fn foo(p : ptr<workgroup, i32>) -> i32 {
+ return workgroupUniformLoad(p);
+}
+
+fn bar() -> i32 {
+ return foo(&(v[0]));
+}
diff --git a/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.dxc.hlsl b/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..de20d9c
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.dxc.hlsl
@@ -0,0 +1,23 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+ return;
+}
+
+groupshared int v[4];
+
+int tint_workgroupUniformLoad_v_X(uint p[1]) {
+ GroupMemoryBarrierWithGroupSync();
+ const int result = v[p[0]];
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+int foo_v_X(uint p[1]) {
+ const uint tint_symbol[1] = {p[0u]};
+ return tint_workgroupUniformLoad_v_X(tint_symbol);
+}
+
+int bar() {
+ 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
new file mode 100644
index 0000000..de20d9c
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.fxc.hlsl
@@ -0,0 +1,23 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+ return;
+}
+
+groupshared int v[4];
+
+int tint_workgroupUniformLoad_v_X(uint p[1]) {
+ GroupMemoryBarrierWithGroupSync();
+ const int result = v[p[0]];
+ GroupMemoryBarrierWithGroupSync();
+ return result;
+}
+
+int foo_v_X(uint p[1]) {
+ const uint tint_symbol[1] = {p[0u]};
+ return tint_workgroupUniformLoad_v_X(tint_symbol);
+}
+
+int bar() {
+ 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
new file mode 100644
index 0000000..7555463
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.glsl
@@ -0,0 +1,24 @@
+#version 310 es
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void unused_entry_point() {
+ return;
+}
+shared int v[4];
+int tint_workgroupUniformLoad_v_X(uint p[1]) {
+ barrier();
+ int result = v[p[0]];
+ barrier();
+ return result;
+}
+
+int foo_v_X(uint p[1]) {
+ uint tint_symbol[1] = uint[1](p[0u]);
+ return tint_workgroupUniformLoad_v_X(tint_symbol);
+}
+
+int bar() {
+ uint tint_symbol_1[1] = uint[1](0u);
+ return foo_v_X(tint_symbol_1);
+}
+
diff --git a/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.msl b/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.msl
new file mode 100644
index 0000000..f666133
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.msl
@@ -0,0 +1,31 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+int tint_workgroupUniformLoad(threadgroup int* const p) {
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ int const result = *(p);
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ return result;
+}
+
+int foo(threadgroup int* const p) {
+ return tint_workgroupUniformLoad(p);
+}
+
+int bar(threadgroup tint_array<int, 4>* const tint_symbol) {
+ return foo(&((*(tint_symbol))[0]));
+}
+
diff --git a/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.spvasm b/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.spvasm
new file mode 100644
index 0000000..439d8d2
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.spvasm
@@ -0,0 +1,63 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 38
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+ OpExecutionMode %unused_entry_point LocalSize 1 1 1
+ OpName %v "v"
+ OpName %unused_entry_point "unused_entry_point"
+ OpName %tint_workgroupUniformLoad_v_X "tint_workgroupUniformLoad_v_X"
+ OpName %p "p"
+ OpName %foo_v_X "foo_v_X"
+ OpName %p_0 "p"
+ OpName %bar "bar"
+ OpDecorate %_arr_int_uint_4 ArrayStride 4
+ OpDecorate %_arr_uint_uint_1 ArrayStride 4
+ %int = OpTypeInt 32 1
+ %uint = OpTypeInt 32 0
+ %uint_4 = OpConstant %uint 4
+%_arr_int_uint_4 = OpTypeArray %int %uint_4
+%_ptr_Workgroup__arr_int_uint_4 = OpTypePointer Workgroup %_arr_int_uint_4
+ %v = OpVariable %_ptr_Workgroup__arr_int_uint_4 Workgroup
+ %void = OpTypeVoid
+ %7 = OpTypeFunction %void
+ %uint_1 = OpConstant %uint 1
+%_arr_uint_uint_1 = OpTypeArray %uint %uint_1
+ %11 = OpTypeFunction %int %_arr_uint_uint_1
+ %uint_2 = OpConstant %uint 2
+ %uint_264 = OpConstant %uint 264
+ %20 = OpConstantNull %int
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+ %30 = OpConstantNull %uint
+ %33 = OpTypeFunction %int
+ %37 = OpConstantNull %_arr_uint_uint_1
+%unused_entry_point = OpFunction %void None %7
+ %10 = OpLabel
+ OpReturn
+ OpFunctionEnd
+%tint_workgroupUniformLoad_v_X = OpFunction %int None %11
+ %p = OpFunctionParameter %_arr_uint_uint_1
+ %16 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %21 = OpCompositeExtract %uint %p 0
+ %23 = OpAccessChain %_ptr_Workgroup_int %v %21
+ %24 = OpLoad %int %23
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ OpReturnValue %24
+ OpFunctionEnd
+ %foo_v_X = OpFunction %int None %11
+ %p_0 = OpFunctionParameter %_arr_uint_uint_1
+ %28 = OpLabel
+ %31 = OpCompositeExtract %uint %p_0 0
+ %32 = OpCompositeConstruct %_arr_uint_uint_1 %31
+ %29 = OpFunctionCall %int %tint_workgroupUniformLoad_v_X %32
+ OpReturnValue %29
+ OpFunctionEnd
+ %bar = OpFunction %int None %33
+ %35 = OpLabel
+ %36 = OpFunctionCall %int %foo_v_X %37
+ OpReturnValue %36
+ OpFunctionEnd
diff --git a/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.wgsl b/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.wgsl
new file mode 100644
index 0000000..59f5e8b
--- /dev/null
+++ b/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.wgsl
@@ -0,0 +1,11 @@
+enable chromium_experimental_full_ptr_parameters;
+
+var<workgroup> v : array<i32, 4>;
+
+fn foo(p : ptr<workgroup, i32>) -> i32 {
+ return workgroupUniformLoad(p);
+}
+
+fn bar() -> i32 {
+ return foo(&(v[0]));
+}
