[msl] Populate Output::workgroup_allocations
This is used by Dawn to allocate threadgroup memory.
Add a `PrintResult` struct to produce extra metadata from the Printer.
Emit the `[[threadgroup]]` attribute on threadgroup memory arguments.
Bug: 42251016
Change-Id: I7ce4e45f44a146442b00bd1bf62a3d911e3bba6f
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/189803
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: dan sinclair <dsinclair@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/lang/msl/writer/BUILD.bazel b/src/tint/lang/msl/writer/BUILD.bazel
index f7ccfaf..c578d9d 100644
--- a/src/tint/lang/msl/writer/BUILD.bazel
+++ b/src/tint/lang/msl/writer/BUILD.bazel
@@ -104,6 +104,7 @@
"return_test.cc",
"type_test.cc",
"var_test.cc",
+ "writer_test.cc",
],
deps = [
"//src/tint/api/common",
diff --git a/src/tint/lang/msl/writer/BUILD.cmake b/src/tint/lang/msl/writer/BUILD.cmake
index 0236187..4aeb2ed 100644
--- a/src/tint/lang/msl/writer/BUILD.cmake
+++ b/src/tint/lang/msl/writer/BUILD.cmake
@@ -116,6 +116,7 @@
lang/msl/writer/return_test.cc
lang/msl/writer/type_test.cc
lang/msl/writer/var_test.cc
+ lang/msl/writer/writer_test.cc
)
tint_target_add_dependencies(tint_lang_msl_writer_test test
diff --git a/src/tint/lang/msl/writer/BUILD.gn b/src/tint/lang/msl/writer/BUILD.gn
index 4fdf454..5635792 100644
--- a/src/tint/lang/msl/writer/BUILD.gn
+++ b/src/tint/lang/msl/writer/BUILD.gn
@@ -106,6 +106,7 @@
"return_test.cc",
"type_test.cc",
"var_test.cc",
+ "writer_test.cc",
]
deps = [
"${tint_src_dir}:gmock_and_gtest",
diff --git a/src/tint/lang/msl/writer/printer/printer.cc b/src/tint/lang/msl/writer/printer/printer.cc
index 0ae4046..b2de07b 100644
--- a/src/tint/lang/msl/writer/printer/printer.cc
+++ b/src/tint/lang/msl/writer/printer/printer.cc
@@ -106,7 +106,7 @@
explicit Printer(core::ir::Module& module) : ir_(module) {}
/// @returns the generated MSL shader
- tint::Result<std::string> Generate() {
+ tint::Result<PrintResult> Generate() {
auto valid = core::ir::ValidateAndDumpIfNeeded(ir_, "MSL writer");
if (valid != Success) {
return std::move(valid.Failure());
@@ -128,10 +128,15 @@
StringStream ss;
ss << preamble_buffer_.String() << std::endl << main_buffer_.String();
- return ss.str();
+ result_.msl = ss.str();
+
+ return std::move(result_);
}
private:
+ /// The result of printing the module.
+ PrintResult result_;
+
/// Map of builtin structure to unique generated name
std::unordered_map<const core::type::Struct*, std::string> builtin_struct_names_;
@@ -253,6 +258,8 @@
{
auto out = Line();
+ auto func_name = NameOf(func);
+
switch (func->Stage()) {
case core::ir::Function::PipelineStage::kCompute:
out << "kernel ";
@@ -266,11 +273,14 @@
case core::ir::Function::PipelineStage::kUndefined:
break;
}
+ if (func->Stage() != core::ir::Function::PipelineStage::kUndefined) {
+ result_.workgroup_allocations.insert({func_name, {}});
+ }
// TODO(dsinclair): Handle return type attributes
EmitType(out, func->ReturnType());
- out << " " << NameOf(func) << "(";
+ out << " " << func_name << "(";
size_t i = 0;
for (auto* param : func->Params()) {
@@ -328,9 +338,10 @@
out << "]]";
}
+ auto ptr = param->Type()->As<core::type::Pointer>();
if (auto binding_point = param->BindingPoint()) {
TINT_ASSERT(binding_point->group == 0);
- if (auto ptr = param->Type()->As<core::type::Pointer>()) {
+ if (ptr) {
switch (ptr->AddressSpace()) {
case core::AddressSpace::kStorage:
case core::AddressSpace::kUniform:
@@ -353,6 +364,12 @@
TINT_ICE_ON_NO_MATCH);
}
}
+ if (ptr && ptr->AddressSpace() == core::AddressSpace::kWorkgroup &&
+ func->Stage() == core::ir::Function::PipelineStage::kCompute) {
+ auto& allocations = result_.workgroup_allocations.at(func_name);
+ out << " [[threadgroup(" << allocations.size() << ")]]";
+ allocations.push_back(ptr->StoreType()->Size());
+ }
}
out << ") {";
@@ -1502,8 +1519,16 @@
} // namespace
-Result<std::string> Print(core::ir::Module& module) {
+Result<PrintResult> Print(core::ir::Module& module) {
return Printer{module}.Generate();
}
+PrintResult::PrintResult() = default;
+
+PrintResult::~PrintResult() = default;
+
+PrintResult::PrintResult(const PrintResult&) = default;
+
+PrintResult& PrintResult::operator=(const PrintResult&) = default;
+
} // namespace tint::msl::writer
diff --git a/src/tint/lang/msl/writer/printer/printer.h b/src/tint/lang/msl/writer/printer/printer.h
index a5218c3..bba5c8c 100644
--- a/src/tint/lang/msl/writer/printer/printer.h
+++ b/src/tint/lang/msl/writer/printer/printer.h
@@ -29,6 +29,8 @@
#define SRC_TINT_LANG_MSL_WRITER_PRINTER_PRINTER_H_
#include <string>
+#include <unordered_map>
+#include <vector>
#include "src/tint/utils/result/result.h"
@@ -39,9 +41,33 @@
namespace tint::msl::writer {
-/// @returns the generated MSL shader on success, or failure
+/// The output produced when printing MSL.
+struct PrintResult {
+ /// Constructor
+ PrintResult();
+
+ /// Destructor
+ ~PrintResult();
+
+ /// Copy constructor
+ PrintResult(const PrintResult&);
+
+ /// Copy assignment
+ /// @returns this
+ PrintResult& operator=(const PrintResult&);
+
+ /// The generated MSL.
+ std::string msl = "";
+
+ /// A map from entry point name to a list of dynamic workgroup allocations.
+ /// Each element of the vector is the size of the workgroup allocation that should be created
+ /// for that index.
+ std::unordered_map<std::string, std::vector<uint32_t>> workgroup_allocations;
+};
+
/// @param module the Tint IR module to generate
-Result<std::string> Print(core::ir::Module& module);
+/// @returns the result of printing the MSL shader on success, or failure
+Result<PrintResult> Print(core::ir::Module& module);
} // namespace tint::msl::writer
diff --git a/src/tint/lang/msl/writer/writer.cc b/src/tint/lang/msl/writer/writer.cc
index c0f9d41..af54c6f 100644
--- a/src/tint/lang/msl/writer/writer.cc
+++ b/src/tint/lang/msl/writer/writer.cc
@@ -62,7 +62,11 @@
if (result != Success) {
return result.Failure();
}
- output.msl = result.Get();
+ output.msl = result->msl;
+ output.workgroup_allocations = std::move(result->workgroup_allocations);
+ // TODO(crbug.com/42251016): Set has_invariant.
+ // TODO(crbug.com/42251016): Set needs_storage_buffer_sizes.
+ // TODO(crbug.com/42251016): Set used_array_length_from_uniform_indices.
return output;
}
diff --git a/src/tint/lang/msl/writer/writer_test.cc b/src/tint/lang/msl/writer/writer_test.cc
new file mode 100644
index 0000000..52896f7
--- /dev/null
+++ b/src/tint/lang/msl/writer/writer_test.cc
@@ -0,0 +1,90 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/msl/writer/helper_test.h"
+
+#include "gmock/gmock.h"
+
+namespace tint::msl::writer {
+namespace {
+
+using namespace tint::core::fluent_types; // NOLINT
+using namespace tint::core::number_suffixes; // NOLINT
+
+TEST_F(MslWriterTest, WorkgroupAllocations) {
+ auto* var_a = b.Var("a", ty.ptr<workgroup, i32>());
+ auto* var_b = b.Var("b", ty.ptr<workgroup, i32>());
+ mod.root_block->Append(var_a);
+ mod.root_block->Append(var_b);
+
+ auto* foo = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kCompute,
+ std::array<uint32_t, 3>{1u, 1u, 1u});
+ b.Append(foo->Block(), [&] {
+ auto* load_a = b.Load(var_a);
+ auto* load_b = b.Load(var_b);
+ b.Store(var_a, b.Add<i32>(load_a, load_b));
+ b.Return(foo);
+ });
+
+ // No allocations, but still needs an entry in the map.
+ auto* bar = b.Function("bar", ty.void_(), core::ir::Function::PipelineStage::kCompute,
+ std::array<uint32_t, 3>{1u, 1u, 1u});
+ b.Append(bar->Block(), [&] { b.Return(bar); });
+
+ ASSERT_TRUE(Generate()) << err_ << output_.msl;
+ EXPECT_EQ(output_.msl, R"(#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_2 {
+ int tint_symbol;
+ int tint_symbol_1;
+};
+struct tint_module_vars_struct {
+ threadgroup int* a;
+ threadgroup int* b;
+};
+
+kernel void foo(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v [[threadgroup(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.a=(&(*v).tint_symbol), .b=(&(*v).tint_symbol_1)};
+ if ((tint_local_index == 0u)) {
+ (*tint_module_vars.a) = 0;
+ (*tint_module_vars.b) = 0;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ (*tint_module_vars.a) = ((*tint_module_vars.a) + (*tint_module_vars.b));
+}
+kernel void bar() {
+}
+)");
+ ASSERT_EQ(output_.workgroup_allocations.size(), 2u);
+ ASSERT_EQ(output_.workgroup_allocations.count("foo"), 1u);
+ ASSERT_EQ(output_.workgroup_allocations.count("bar"), 1u);
+ EXPECT_THAT(output_.workgroup_allocations.at("foo"), testing::ElementsAre(8u));
+ EXPECT_THAT(output_.workgroup_allocations.at("bar"), testing::ElementsAre());
+}
+
+} // namespace
+} // namespace tint::msl::writer
diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x2_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x2_f16/to_workgroup.wgsl.expected.ir.msl
index db158da..03a9fb8 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x2_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x2_f16/to_workgroup.wgsl.expected.ir.msl
@@ -8,7 +8,7 @@
threadgroup half2x2* w;
};
-kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half2x2* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half2x2* u [[buffer(0)]], threadgroup tint_symbol_1* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.w) = half2x2(half2(0.0h), half2(0.0h));
diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
index 938d418..2a0c9bb 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -8,7 +8,7 @@
threadgroup float2x2* w;
};
-kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float2x2* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float2x2* u [[buffer(0)]], threadgroup tint_symbol_1* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.w) = float2x2(float2(0.0f), float2(0.0f));
diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
index f0e615f..0ad0481 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -8,7 +8,7 @@
threadgroup half2x3* w;
};
-kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half2x3* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half2x3* u [[buffer(0)]], threadgroup tint_symbol_1* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.w) = half2x3(half3(0.0h), half3(0.0h));
diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
index bb2ee0c..326af71 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -8,7 +8,7 @@
threadgroup float2x3* w;
};
-kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float2x3* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float2x3* u [[buffer(0)]], threadgroup tint_symbol_1* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.w) = float2x3(float3(0.0f), float3(0.0f));
diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
index 85389c0..52ed137 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -8,7 +8,7 @@
threadgroup half2x4* w;
};
-kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half2x4* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half2x4* u [[buffer(0)]], threadgroup tint_symbol_1* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.w) = half2x4(half4(0.0h), half4(0.0h));
diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
index 52da3cb..33c1232 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -8,7 +8,7 @@
threadgroup float2x4* w;
};
-kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float2x4* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float2x4* u [[buffer(0)]], threadgroup tint_symbol_1* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.w) = float2x4(float4(0.0f), float4(0.0f));
diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x2_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x2_f16/to_workgroup.wgsl.expected.ir.msl
index 5893867..f7d1040 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x2_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x2_f16/to_workgroup.wgsl.expected.ir.msl
@@ -8,7 +8,7 @@
threadgroup half3x2* w;
};
-kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half3x2* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half3x2* u [[buffer(0)]], threadgroup tint_symbol_1* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.w) = half3x2(half2(0.0h), half2(0.0h), half2(0.0h));
diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x2_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x2_f32/to_workgroup.wgsl.expected.ir.msl
index da3e2fd..a90a83d 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -8,7 +8,7 @@
threadgroup float3x2* w;
};
-kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float3x2* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float3x2* u [[buffer(0)]], threadgroup tint_symbol_1* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.w) = float3x2(float2(0.0f), float2(0.0f), float2(0.0f));
diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x3_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x3_f16/to_workgroup.wgsl.expected.ir.msl
index 3276f46..e267930 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -8,7 +8,7 @@
threadgroup half3x3* w;
};
-kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half3x3* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half3x3* u [[buffer(0)]], threadgroup tint_symbol_1* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.w) = half3x3(half3(0.0h), half3(0.0h), half3(0.0h));
diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
index 105e86d..9095e62 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -8,7 +8,7 @@
threadgroup float3x3* w;
};
-kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float3x3* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float3x3* u [[buffer(0)]], threadgroup tint_symbol_1* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.w) = float3x3(float3(0.0f), float3(0.0f), float3(0.0f));
diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x4_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x4_f16/to_workgroup.wgsl.expected.ir.msl
index fce2da8..572e381 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -8,7 +8,7 @@
threadgroup half3x4* w;
};
-kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half3x4* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half3x4* u [[buffer(0)]], threadgroup tint_symbol_1* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.w) = half3x4(half4(0.0h), half4(0.0h), half4(0.0h));
diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
index bc3018c..f3ab07e 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -8,7 +8,7 @@
threadgroup float3x4* w;
};
-kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float3x4* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float3x4* u [[buffer(0)]], threadgroup tint_symbol_1* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.w) = float3x4(float4(0.0f), float4(0.0f), float4(0.0f));
diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
index 2a21f89..1cabc56 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
@@ -8,7 +8,7 @@
threadgroup half4x2* w;
};
-kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half4x2* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half4x2* u [[buffer(0)]], threadgroup tint_symbol_1* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.w) = half4x2(half2(0.0h), half2(0.0h), half2(0.0h), half2(0.0h));
diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
index 5ffe4bd..53be8e0 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -8,7 +8,7 @@
threadgroup float4x2* w;
};
-kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float4x2* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float4x2* u [[buffer(0)]], threadgroup tint_symbol_1* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.w) = float4x2(float2(0.0f), float2(0.0f), float2(0.0f), float2(0.0f));
diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
index 2a836c1..14b97e2 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -8,7 +8,7 @@
threadgroup half4x3* w;
};
-kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half4x3* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half4x3* u [[buffer(0)]], threadgroup tint_symbol_1* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.w) = half4x3(half3(0.0h), half3(0.0h), half3(0.0h), half3(0.0h));
diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
index 1d930a4..7b9f2fc 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -8,7 +8,7 @@
threadgroup float4x3* w;
};
-kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float4x3* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float4x3* u [[buffer(0)]], threadgroup tint_symbol_1* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.w) = float4x3(float3(0.0f), float3(0.0f), float3(0.0f), float3(0.0f));
diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
index 67b0ff7..3fd1741 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -8,7 +8,7 @@
threadgroup half4x4* w;
};
-kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half4x4* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half4x4* u [[buffer(0)]], threadgroup tint_symbol_1* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.w) = half4x4(half4(0.0h), half4(0.0h), half4(0.0h), half4(0.0h));
diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
index 1143198..8f9c882 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -8,7 +8,7 @@
threadgroup float4x4* w;
};
-kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float4x4* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float4x4* u [[buffer(0)]], threadgroup tint_symbol_1* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.w) = float4x4(float4(0.0f), float4(0.0f), float4(0.0f), float4(0.0f));
diff --git a/test/tint/bug/chromium/40943165.wgsl.expected.ir.msl b/test/tint/bug/chromium/40943165.wgsl.expected.ir.msl
index a7b5929..7c814a2 100644
--- a/test/tint/bug/chromium/40943165.wgsl.expected.ir.msl
+++ b/test/tint/bug/chromium/40943165.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@
threadgroup float2x2* W;
};
-kernel void F(uint mat2x2 [[thread_index_in_threadgroup]], threadgroup tint_symbol_1* v) {
+kernel void F(uint mat2x2 [[thread_index_in_threadgroup]], threadgroup tint_symbol_1* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.W=(&(*v).tint_symbol)};
if ((mat2x2 == 0u)) {
(*tint_module_vars.W) = float2x2(float2(0.0f), float2(0.0f));
diff --git a/test/tint/bug/tint/1926.wgsl.expected.ir.msl b/test/tint/bug/tint/1926.wgsl.expected.ir.msl
index 812684c..a6e9dba 100644
--- a/test/tint/bug/tint/1926.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1926.wgsl.expected.ir.msl
@@ -8,7 +8,7 @@
device uint* output;
};
-kernel void tint_symbol(uint3 global_id [[thread_position_in_grid]], uint3 local_id [[thread_position_in_threadgroup]], uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v, device uint* output [[buffer(4)]]) {
+kernel void tint_symbol(uint3 global_id [[thread_position_in_grid]], uint3 local_id [[thread_position_in_threadgroup]], uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v [[threadgroup(0)]], device uint* output [[buffer(4)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.sh_atomic_failed=(&(*v).tint_symbol_1), .output=output};
if ((tint_local_index == 0u)) {
(*tint_module_vars.sh_atomic_failed) = 0u;
diff --git a/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.ir.msl
index 5f96aa4..7e06d54 100644
--- a/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@
threadgroup int* i;
};
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.i=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.i) = 0;
diff --git a/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.ir.msl
index fbd0e7d..53a535a 100644
--- a/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.ir.msl
@@ -10,7 +10,7 @@
int func(threadgroup int* const pointer) {
return (*pointer);
}
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.S) = 0;
diff --git a/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.ir.msl
index 0e25c93..500d34a 100644
--- a/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.ir.msl
@@ -13,7 +13,7 @@
int func(threadgroup int* const pointer) {
return (*pointer);
}
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.S) = str{};
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.ir.msl
index 9e76b8b..b5b4302 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.ir.msl
@@ -10,7 +10,7 @@
float2 func(threadgroup float2* const pointer) {
return (*pointer);
}
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.S) = float2x2(float2(0.0f), float2(0.0f));
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.ir.msl
index 3764937..9d0a19c 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.ir.msl
@@ -10,7 +10,7 @@
float4 func(threadgroup float4* const pointer) {
return (*pointer);
}
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.S) = float4(0.0f);
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.ir.msl
index fb1ab2d..5ba2e4b 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.ir.msl
@@ -10,7 +10,7 @@
float4 func(threadgroup float4* const pointer) {
return (*pointer);
}
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.S) = float2x4(float4(0.0f), float4(0.0f));
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.ir.msl
index 96d683f..e7e11ac 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.ir.msl
@@ -13,7 +13,7 @@
float4 func(threadgroup float4* const pointer) {
return (*pointer);
}
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.S) = str{};
diff --git a/test/tint/ptr_ref/store/param/workgroup/i32.wgsl.expected.ir.msl b/test/tint/ptr_ref/store/param/workgroup/i32.wgsl.expected.ir.msl
index 57dbb78..59eefc1 100644
--- a/test/tint/ptr_ref/store/param/workgroup/i32.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/i32.wgsl.expected.ir.msl
@@ -10,7 +10,7 @@
void func(threadgroup int* const pointer) {
(*pointer) = 42;
}
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.S) = 0;
diff --git a/test/tint/ptr_ref/store/param/workgroup/i32_in_struct.wgsl.expected.ir.msl b/test/tint/ptr_ref/store/param/workgroup/i32_in_struct.wgsl.expected.ir.msl
index 1defcd2..172d809 100644
--- a/test/tint/ptr_ref/store/param/workgroup/i32_in_struct.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/i32_in_struct.wgsl.expected.ir.msl
@@ -13,7 +13,7 @@
void func(threadgroup int* const pointer) {
(*pointer) = 42;
}
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.S) = str{};
diff --git a/test/tint/ptr_ref/store/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.ir.msl b/test/tint/ptr_ref/store/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.ir.msl
index f089435..f5b5fe3 100644
--- a/test/tint/ptr_ref/store/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.ir.msl
@@ -10,7 +10,7 @@
void func(threadgroup float2* const pointer) {
(*pointer) = float2(0.0f);
}
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.S) = float2x2(float2(0.0f), float2(0.0f));
diff --git a/test/tint/ptr_ref/store/param/workgroup/vec4_f32.wgsl.expected.ir.msl b/test/tint/ptr_ref/store/param/workgroup/vec4_f32.wgsl.expected.ir.msl
index 416f51b..74aff82 100644
--- a/test/tint/ptr_ref/store/param/workgroup/vec4_f32.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/vec4_f32.wgsl.expected.ir.msl
@@ -10,7 +10,7 @@
void func(threadgroup float4* const pointer) {
(*pointer) = float4(0.0f);
}
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.S) = float4(0.0f);
diff --git a/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.ir.msl b/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.ir.msl
index c4d6342..1a4fde6 100644
--- a/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.ir.msl
@@ -10,7 +10,7 @@
void func(threadgroup float4* const pointer) {
(*pointer) = float4(0.0f);
}
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.S) = float2x4(float4(0.0f), float4(0.0f));
diff --git a/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_struct.wgsl.expected.ir.msl b/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_struct.wgsl.expected.ir.msl
index e943602..fe7c279 100644
--- a/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_struct.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_struct.wgsl.expected.ir.msl
@@ -13,7 +13,7 @@
void func(threadgroup float4* const pointer) {
(*pointer) = float4(0.0f);
}
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.S) = str{};
diff --git a/test/tint/types/module_scope_used_in_functions.wgsl.expected.ir.msl b/test/tint/types/module_scope_used_in_functions.wgsl.expected.ir.msl
index 2b6c649..0b9f5f3 100644
--- a/test/tint/types/module_scope_used_in_functions.wgsl.expected.ir.msl
+++ b/test/tint/types/module_scope_used_in_functions.wgsl.expected.ir.msl
@@ -38,7 +38,7 @@
bar(a, b, tint_module_vars);
no_uses();
}
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v, const device float2* uniforms [[buffer(1)]], device tint_array<float, 1>* storages [[buffer(0)]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v [[threadgroup(0)]], const device float2* uniforms [[buffer(1)]], device tint_array<float, 1>* storages [[buffer(0)]]) {
thread float p = 0.0f;
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.p=(&p), .w=(&(*v).tint_symbol_1), .uniforms=uniforms, .storages=storages};
if ((tint_local_index == 0u)) {
diff --git a/test/tint/types/module_scope_var.wgsl.expected.ir.msl b/test/tint/types/module_scope_var.wgsl.expected.ir.msl
index b44443a..35f373a 100644
--- a/test/tint/types/module_scope_var.wgsl.expected.ir.msl
+++ b/test/tint/types/module_scope_var.wgsl.expected.ir.msl
@@ -32,7 +32,7 @@
threadgroup float* wg_var;
};
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v [[threadgroup(0)]]) {
thread bool bool_var = false;
thread int i32_var = 0;
thread uint u32_var = 0u;
diff --git a/test/tint/types/module_scope_vars_pointers.wgsl.expected.ir.msl b/test/tint/types/module_scope_vars_pointers.wgsl.expected.ir.msl
index bc857e6..e643a4b 100644
--- a/test/tint/types/module_scope_vars_pointers.wgsl.expected.ir.msl
+++ b/test/tint/types/module_scope_vars_pointers.wgsl.expected.ir.msl
@@ -8,7 +8,7 @@
threadgroup float* w;
};
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v [[threadgroup(0)]]) {
thread float p = 0.0f;
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.p=(&p), .w=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
diff --git a/test/tint/var/initialization/workgroup/matrix.wgsl.expected.ir.msl b/test/tint/var/initialization/workgroup/matrix.wgsl.expected.ir.msl
index 44e6f0a..058a371 100644
--- a/test/tint/var/initialization/workgroup/matrix.wgsl.expected.ir.msl
+++ b/test/tint/var/initialization/workgroup/matrix.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@
threadgroup float2x3* v;
};
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v_1) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v_1 [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.v=(&(*v_1).tint_symbol_1)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.v) = float2x3(float3(0.0f), float3(0.0f));
diff --git a/test/tint/var/initialization/workgroup/scalar.wgsl.expected.ir.msl b/test/tint/var/initialization/workgroup/scalar.wgsl.expected.ir.msl
index 97b64c8..a039a36 100644
--- a/test/tint/var/initialization/workgroup/scalar.wgsl.expected.ir.msl
+++ b/test/tint/var/initialization/workgroup/scalar.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@
threadgroup int* v;
};
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v_1) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v_1 [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.v=(&(*v_1).tint_symbol_1)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.v) = 0;
diff --git a/test/tint/var/initialization/workgroup/struct.wgsl.expected.ir.msl b/test/tint/var/initialization/workgroup/struct.wgsl.expected.ir.msl
index e974d43..98f5004 100644
--- a/test/tint/var/initialization/workgroup/struct.wgsl.expected.ir.msl
+++ b/test/tint/var/initialization/workgroup/struct.wgsl.expected.ir.msl
@@ -11,7 +11,7 @@
threadgroup S* v;
};
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v_1) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v_1 [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.v=(&(*v_1).tint_symbol_1)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.v) = S{};
diff --git a/test/tint/var/initialization/workgroup/vector.wgsl.expected.ir.msl b/test/tint/var/initialization/workgroup/vector.wgsl.expected.ir.msl
index d24feeb..bc51209 100644
--- a/test/tint/var/initialization/workgroup/vector.wgsl.expected.ir.msl
+++ b/test/tint/var/initialization/workgroup/vector.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@
threadgroup int3* v;
};
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v_1) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v_1 [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.v=(&(*v_1).tint_symbol_1)};
if ((tint_local_index == 0u)) {
(*tint_module_vars.v) = int3(0);
diff --git a/test/tint/var/uses/many_workgroup_vars.wgsl.expected.ir.msl b/test/tint/var/uses/many_workgroup_vars.wgsl.expected.ir.msl
index 0a83ab7..6ae0add 100644
--- a/test/tint/var/uses/many_workgroup_vars.wgsl.expected.ir.msl
+++ b/test/tint/var/uses/many_workgroup_vars.wgsl.expected.ir.msl
@@ -205,7 +205,7 @@
threadgroup float2x2* m99;
};
-kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], threadgroup tint_symbol_101* v) {
+kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], threadgroup tint_symbol_101* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.m00=(&(*v).tint_symbol_1), .m01=(&(*v).tint_symbol_2), .m02=(&(*v).tint_symbol_3), .m03=(&(*v).tint_symbol_4), .m04=(&(*v).tint_symbol_5), .m05=(&(*v).tint_symbol_6), .m06=(&(*v).tint_symbol_7), .m07=(&(*v).tint_symbol_8), .m08=(&(*v).tint_symbol_9), .m09=(&(*v).tint_symbol_10), .m10=(&(*v).tint_symbol_11), .m11=(&(*v).tint_symbol_12), .m12=(&(*v).tint_symbol_13), .m13=(&(*v).tint_symbol_14), .m14=(&(*v).tint_symbol_15), .m15=(&(*v).tint_symbol_16), .m16=(&(*v).tint_symbol_17), .m17=(&(*v).tint_symbol_18), .m18=(&(*v).tint_symbol_19), .m19=(&(*v).tint_symbol_20), .m20=(&(*v).tint_symbol_21), .m21=(&(*v).tint_symbol_22), .m22=(&(*v).tint_symbol_23), .m23=(&(*v).tint_symbol_24), .m24=(&(*v).tint_symbol_25), .m25=(&(*v).tint_symbol_26), .m26=(&(*v).tint_symbol_27), .m27=(&(*v).tint_symbol_28), .m28=(&(*v).tint_symbol_29), .m29=(&(*v).tint_symbol_30), .m30=(&(*v).tint_symbol_31), .m31=(&(*v).tint_symbol_32), .m32=(&(*v).tint_symbol_33), .m33=(&(*v).tint_symbol_34), .m34=(&(*v).tint_symbol_35), .m35=(&(*v).tint_symbol_36), .m36=(&(*v).tint_symbol_37), .m37=(&(*v).tint_symbol_38), .m38=(&(*v).tint_symbol_39), .m39=(&(*v).tint_symbol_40), .m40=(&(*v).tint_symbol_41), .m41=(&(*v).tint_symbol_42), .m42=(&(*v).tint_symbol_43), .m43=(&(*v).tint_symbol_44), .m44=(&(*v).tint_symbol_45), .m45=(&(*v).tint_symbol_46), .m46=(&(*v).tint_symbol_47), .m47=(&(*v).tint_symbol_48), .m48=(&(*v).tint_symbol_49), .m49=(&(*v).tint_symbol_50), .m50=(&(*v).tint_symbol_51), .m51=(&(*v).tint_symbol_52), .m52=(&(*v).tint_symbol_53), .m53=(&(*v).tint_symbol_54), .m54=(&(*v).tint_symbol_55), .m55=(&(*v).tint_symbol_56), .m56=(&(*v).tint_symbol_57), .m57=(&(*v).tint_symbol_58), .m58=(&(*v).tint_symbol_59), .m59=(&(*v).tint_symbol_60), .m60=(&(*v).tint_symbol_61), .m61=(&(*v).tint_symbol_62), .m62=(&(*v).tint_symbol_63), .m63=(&(*v).tint_symbol_64), .m64=(&(*v).tint_symbol_65), .m65=(&(*v).tint_symbol_66), .m66=(&(*v).tint_symbol_67), .m67=(&(*v).tint_symbol_68), .m68=(&(*v).tint_symbol_69), .m69=(&(*v).tint_symbol_70), .m70=(&(*v).tint_symbol_71), .m71=(&(*v).tint_symbol_72), .m72=(&(*v).tint_symbol_73), .m73=(&(*v).tint_symbol_74), .m74=(&(*v).tint_symbol_75), .m75=(&(*v).tint_symbol_76), .m76=(&(*v).tint_symbol_77), .m77=(&(*v).tint_symbol_78), .m78=(&(*v).tint_symbol_79), .m79=(&(*v).tint_symbol_80), .m80=(&(*v).tint_symbol_81), .m81=(&(*v).tint_symbol_82), .m82=(&(*v).tint_symbol_83), .m83=(&(*v).tint_symbol_84), .m84=(&(*v).tint_symbol_85), .m85=(&(*v).tint_symbol_86), .m86=(&(*v).tint_symbol_87), .m87=(&(*v).tint_symbol_88), .m88=(&(*v).tint_symbol_89), .m89=(&(*v).tint_symbol_90), .m90=(&(*v).tint_symbol_91), .m91=(&(*v).tint_symbol_92), .m92=(&(*v).tint_symbol_93), .m93=(&(*v).tint_symbol_94), .m94=(&(*v).tint_symbol_95), .m95=(&(*v).tint_symbol_96), .m96=(&(*v).tint_symbol_97), .m97=(&(*v).tint_symbol_98), .m98=(&(*v).tint_symbol_99), .m99=(&(*v).tint_symbol_100)};
if ((idx == 0u)) {
(*tint_module_vars.m00) = float2x2(float2(0.0f), float2(0.0f));
