[msl] Support workgroup vars in ModuleScopeVars
Workgroup variables will be allocated by Dawn and passed as entry
point parameters (to workaround an MSL compiler bug with threadgroup
matrices), and we aggregate all of them into a single structure to
avoid hitting MSL's limit for threadgroup memory arguments.
Bug: 42251016
Change-Id: I5b4b793b2cebdccc8f88b130434610e3477f3b00
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/189801
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Commit-Queue: James Price <jrprice@google.com>
diff --git a/src/tint/lang/msl/writer/printer/var_test.cc b/src/tint/lang/msl/writer/printer/var_test.cc
index b91c2bd..4b3c2f9 100644
--- a/src/tint/lang/msl/writer/printer/var_test.cc
+++ b/src/tint/lang/msl/writer/printer/var_test.cc
@@ -282,8 +282,7 @@
)");
}
-// TODO(jrprice): Requires ModuleScopeVarToEntryPointParam transform
-TEST_F(MslPrinterTest, DISABLED_VarGlobalWorkgroup) {
+TEST_F(MslPrinterTest, VarGlobalWorkgroup) {
core::ir::Var* v = nullptr;
b.Append(mod.root_block,
[&] { v = b.Var("v", ty.ptr<core::AddressSpace::kWorkgroup, f32>()); });
@@ -297,10 +296,12 @@
});
ASSERT_TRUE(Generate()) << err_ << output_;
- EXPECT_EQ(output_, MetalHeader() + R"(
-threadgroup float v;
-void foo() {
- float a = v;
+ EXPECT_EQ(output_, MetalHeader() + R"(struct tint_module_vars_struct {
+ threadgroup float* v;
+};
+
+void foo(tint_module_vars_struct tint_module_vars) {
+ float a = (*tint_module_vars.v);
}
)");
}
diff --git a/src/tint/lang/msl/writer/raise/module_scope_vars.cc b/src/tint/lang/msl/writer/raise/module_scope_vars.cc
index 00a569f..6a4ecb2 100644
--- a/src/tint/lang/msl/writer/raise/module_scope_vars.cc
+++ b/src/tint/lang/msl/writer/raise/module_scope_vars.cc
@@ -174,6 +174,9 @@
core::ir::Function* func,
const core::ir::ReferencedModuleVars::VarSet& referenced_vars) {
core::ir::Value* module_var_struct = nullptr;
+ core::ir::FunctionParam* workgroup_allocation_param = nullptr;
+ Vector<core::type::Manager::StructMemberDesc, 4> workgroup_struct_members;
+
// Add parameters and insert instruction at the top of the entry point to set up the
// module-scope variables structure.
b.InsertBefore(func->Block()->Front(), [&] { //
@@ -206,6 +209,24 @@
decl = param;
break;
}
+ case core::AddressSpace::kWorkgroup: {
+ // Workgroup variables are received as a function parameter (to workaround
+ // an MSL compiler bug with threadgroup matrices), and we aggregate all
+ // workgroup variables into a structure to avoid hitting MSL's limit for
+ // threadgroup memory arguments.
+ if (!workgroup_allocation_param) {
+ workgroup_allocation_param = b.FunctionParam(nullptr);
+ func->AppendParam(workgroup_allocation_param);
+ }
+ decl = b.Access(ptr, workgroup_allocation_param,
+ u32(workgroup_struct_members.Length()))
+ ->Result(0);
+ workgroup_struct_members.Push(core::type::Manager::StructMemberDesc{
+ ir.symbols.New(),
+ ptr->StoreType(),
+ });
+ break;
+ }
case core::AddressSpace::kHandle: {
// Handle types become function parameters and drop the pointer.
auto* param = b.FunctionParam(ptr->UnwrapPtr());
@@ -230,6 +251,14 @@
auto* construct = b.Construct(struct_type, std::move(construct_args));
module_var_struct = b.Let(kModuleVarsName, construct)->Result(0);
});
+
+ // Create the workgroup variable structure if needed.
+ if (!workgroup_struct_members.IsEmpty()) {
+ auto* workgroup_struct =
+ ty.Struct(ir.symbols.New(), std::move(workgroup_struct_members));
+ workgroup_allocation_param->SetType(ty.ptr<workgroup>(workgroup_struct));
+ }
+
return module_var_struct;
}
diff --git a/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc b/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc
index 3bef228..cf3fe76 100644
--- a/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc
+++ b/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc
@@ -364,6 +364,73 @@
EXPECT_EQ(expect, str());
}
+TEST_F(MslWriter_ModuleScopeVarsTest, Workgroup) {
+ 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* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kCompute,
+ std::array<uint32_t, 3>{1u, 1u, 1u});
+ b.Append(func->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(func);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %a:ptr<workgroup, i32, read_write> = var
+ %b:ptr<workgroup, i32, read_write> = var
+}
+
+%foo = @compute @workgroup_size(1, 1, 1) func():void {
+ $B2: {
+ %4:i32 = load %a
+ %5:i32 = load %b
+ %6:i32 = add %4, %5
+ store %a, %6
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+tint_module_vars_struct = struct @align(1) {
+ a:ptr<workgroup, i32, read_write> @offset(0)
+ b:ptr<workgroup, i32, read_write> @offset(0)
+}
+
+tint_symbol_2 = struct @align(4) {
+ tint_symbol:i32 @offset(0)
+ tint_symbol_1:i32 @offset(4)
+}
+
+%foo = @compute @workgroup_size(1, 1, 1) func(%2:ptr<workgroup, tint_symbol_2, read_write>):void {
+ $B1: {
+ %a:ptr<workgroup, i32, read_write> = access %2, 0u
+ %b:ptr<workgroup, i32, read_write> = access %2, 1u
+ %5:tint_module_vars_struct = construct %a, %b
+ %tint_module_vars:tint_module_vars_struct = let %5
+ %7:ptr<workgroup, i32, read_write> = access %tint_module_vars, 0u
+ %8:i32 = load %7
+ %9:ptr<workgroup, i32, read_write> = access %tint_module_vars, 1u
+ %10:i32 = load %9
+ %11:i32 = add %8, %10
+ %12:ptr<workgroup, i32, read_write> = access %tint_module_vars, 0u
+ store %12, %11
+ ret
+ }
+}
+)";
+
+ Run(ModuleScopeVars);
+
+ EXPECT_EQ(expect, str());
+}
+
TEST_F(MslWriter_ModuleScopeVarsTest, MultipleAddressSpaces) {
auto* var_a = b.Var("a", ty.ptr<uniform, i32, core::Access::kRead>());
auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>());
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 37f0e9b..db158da 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
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+ half2x2 tint_symbol;
+};
+struct tint_module_vars_struct {
+ const constant half2x2* u;
+ threadgroup half2x2* w;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat2x2<f16>, read> = var @binding_point(0, 0)
- %w:ptr<workgroup, mat2x2<f16>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
- $B2: {
- %5:bool = eq %tint_local_index, 0u
- if %5 [t: $B3] { # if_1
- $B3: { # true
- store %w, mat2x2<f16>(vec2<f16>(0.0h))
- exit_if # if_1
- }
- }
- %6:void = msl.threadgroup_barrier 4u
- %7:mat2x2<f16> = load %u
- store %w, %7
- %8:ptr<workgroup, vec2<f16>, read_write> = access %w, 1i
- %9:ptr<uniform, vec2<f16>, read> = access %u, 0i
- %10:vec2<f16> = load %9
- store %8, %10
- %11:ptr<workgroup, vec2<f16>, read_write> = access %w, 1i
- %12:ptr<uniform, vec2<f16>, read> = access %u, 0i
- %13:vec2<f16> = load %12
- %14:vec2<f16> = swizzle %13, yx
- store %11, %14
- %15:ptr<workgroup, vec2<f16>, read_write> = access %w, 0i
- %16:ptr<uniform, vec2<f16>, read> = access %u, 1i
- %17:f16 = load_vector_element %16, 0i
- store_vector_element %15, 1i, %17
- ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half2x2* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+ 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));
}
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ (*tint_module_vars.w) = (*tint_module_vars.u);
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].yx;
+ (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
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 363c971..938d418 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
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+ float2x2 tint_symbol;
+};
+struct tint_module_vars_struct {
+ const constant float2x2* u;
+ threadgroup float2x2* w;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat2x2<f32>, read> = var @binding_point(0, 0)
- %w:ptr<workgroup, mat2x2<f32>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
- $B2: {
- %5:bool = eq %tint_local_index, 0u
- if %5 [t: $B3] { # if_1
- $B3: { # true
- store %w, mat2x2<f32>(vec2<f32>(0.0f))
- exit_if # if_1
- }
- }
- %6:void = msl.threadgroup_barrier 4u
- %7:mat2x2<f32> = load %u
- store %w, %7
- %8:ptr<workgroup, vec2<f32>, read_write> = access %w, 1i
- %9:ptr<uniform, vec2<f32>, read> = access %u, 0i
- %10:vec2<f32> = load %9
- store %8, %10
- %11:ptr<workgroup, vec2<f32>, read_write> = access %w, 1i
- %12:ptr<uniform, vec2<f32>, read> = access %u, 0i
- %13:vec2<f32> = load %12
- %14:vec2<f32> = swizzle %13, yx
- store %11, %14
- %15:ptr<workgroup, vec2<f32>, read_write> = access %w, 0i
- %16:ptr<uniform, vec2<f32>, read> = access %u, 1i
- %17:f32 = load_vector_element %16, 0i
- store_vector_element %15, 1i, %17
- ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float2x2* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+ 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));
}
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ (*tint_module_vars.w) = (*tint_module_vars.u);
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].yx;
+ (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
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 2a286d3..f0e615f 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
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+ half2x3 tint_symbol;
+};
+struct tint_module_vars_struct {
+ const constant half2x3* u;
+ threadgroup half2x3* w;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat2x3<f16>, read> = var @binding_point(0, 0)
- %w:ptr<workgroup, mat2x3<f16>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
- $B2: {
- %5:bool = eq %tint_local_index, 0u
- if %5 [t: $B3] { # if_1
- $B3: { # true
- store %w, mat2x3<f16>(vec3<f16>(0.0h))
- exit_if # if_1
- }
- }
- %6:void = msl.threadgroup_barrier 4u
- %7:mat2x3<f16> = load %u
- store %w, %7
- %8:ptr<workgroup, vec3<f16>, read_write> = access %w, 1i
- %9:ptr<uniform, vec3<f16>, read> = access %u, 0i
- %10:vec3<f16> = load %9
- store %8, %10
- %11:ptr<workgroup, vec3<f16>, read_write> = access %w, 1i
- %12:ptr<uniform, vec3<f16>, read> = access %u, 0i
- %13:vec3<f16> = load %12
- %14:vec3<f16> = swizzle %13, zxy
- store %11, %14
- %15:ptr<workgroup, vec3<f16>, read_write> = access %w, 0i
- %16:ptr<uniform, vec3<f16>, read> = access %u, 1i
- %17:f16 = load_vector_element %16, 0i
- store_vector_element %15, 1i, %17
- ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half2x3* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+ 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));
}
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ (*tint_module_vars.w) = (*tint_module_vars.u);
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].zxy;
+ (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
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 93c51a4..bb2ee0c 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
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+ float2x3 tint_symbol;
+};
+struct tint_module_vars_struct {
+ const constant float2x3* u;
+ threadgroup float2x3* w;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat2x3<f32>, read> = var @binding_point(0, 0)
- %w:ptr<workgroup, mat2x3<f32>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
- $B2: {
- %5:bool = eq %tint_local_index, 0u
- if %5 [t: $B3] { # if_1
- $B3: { # true
- store %w, mat2x3<f32>(vec3<f32>(0.0f))
- exit_if # if_1
- }
- }
- %6:void = msl.threadgroup_barrier 4u
- %7:mat2x3<f32> = load %u
- store %w, %7
- %8:ptr<workgroup, vec3<f32>, read_write> = access %w, 1i
- %9:ptr<uniform, vec3<f32>, read> = access %u, 0i
- %10:vec3<f32> = load %9
- store %8, %10
- %11:ptr<workgroup, vec3<f32>, read_write> = access %w, 1i
- %12:ptr<uniform, vec3<f32>, read> = access %u, 0i
- %13:vec3<f32> = load %12
- %14:vec3<f32> = swizzle %13, zxy
- store %11, %14
- %15:ptr<workgroup, vec3<f32>, read_write> = access %w, 0i
- %16:ptr<uniform, vec3<f32>, read> = access %u, 1i
- %17:f32 = load_vector_element %16, 0i
- store_vector_element %15, 1i, %17
- ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float2x3* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+ 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));
}
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ (*tint_module_vars.w) = (*tint_module_vars.u);
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].zxy;
+ (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
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 9a8c6ea..85389c0 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
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+ half2x4 tint_symbol;
+};
+struct tint_module_vars_struct {
+ const constant half2x4* u;
+ threadgroup half2x4* w;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat2x4<f16>, read> = var @binding_point(0, 0)
- %w:ptr<workgroup, mat2x4<f16>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
- $B2: {
- %5:bool = eq %tint_local_index, 0u
- if %5 [t: $B3] { # if_1
- $B3: { # true
- store %w, mat2x4<f16>(vec4<f16>(0.0h))
- exit_if # if_1
- }
- }
- %6:void = msl.threadgroup_barrier 4u
- %7:mat2x4<f16> = load %u
- store %w, %7
- %8:ptr<workgroup, vec4<f16>, read_write> = access %w, 1i
- %9:ptr<uniform, vec4<f16>, read> = access %u, 0i
- %10:vec4<f16> = load %9
- store %8, %10
- %11:ptr<workgroup, vec4<f16>, read_write> = access %w, 1i
- %12:ptr<uniform, vec4<f16>, read> = access %u, 0i
- %13:vec4<f16> = load %12
- %14:vec4<f16> = swizzle %13, ywxz
- store %11, %14
- %15:ptr<workgroup, vec4<f16>, read_write> = access %w, 0i
- %16:ptr<uniform, vec4<f16>, read> = access %u, 1i
- %17:f16 = load_vector_element %16, 0i
- store_vector_element %15, 1i, %17
- ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half2x4* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+ 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));
}
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ (*tint_module_vars.w) = (*tint_module_vars.u);
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].ywxz;
+ (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
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 0a303a5..52da3cb 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
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+ float2x4 tint_symbol;
+};
+struct tint_module_vars_struct {
+ const constant float2x4* u;
+ threadgroup float2x4* w;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat2x4<f32>, read> = var @binding_point(0, 0)
- %w:ptr<workgroup, mat2x4<f32>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
- $B2: {
- %5:bool = eq %tint_local_index, 0u
- if %5 [t: $B3] { # if_1
- $B3: { # true
- store %w, mat2x4<f32>(vec4<f32>(0.0f))
- exit_if # if_1
- }
- }
- %6:void = msl.threadgroup_barrier 4u
- %7:mat2x4<f32> = load %u
- store %w, %7
- %8:ptr<workgroup, vec4<f32>, read_write> = access %w, 1i
- %9:ptr<uniform, vec4<f32>, read> = access %u, 0i
- %10:vec4<f32> = load %9
- store %8, %10
- %11:ptr<workgroup, vec4<f32>, read_write> = access %w, 1i
- %12:ptr<uniform, vec4<f32>, read> = access %u, 0i
- %13:vec4<f32> = load %12
- %14:vec4<f32> = swizzle %13, ywxz
- store %11, %14
- %15:ptr<workgroup, vec4<f32>, read_write> = access %w, 0i
- %16:ptr<uniform, vec4<f32>, read> = access %u, 1i
- %17:f32 = load_vector_element %16, 0i
- store_vector_element %15, 1i, %17
- ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float2x4* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+ 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));
}
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ (*tint_module_vars.w) = (*tint_module_vars.u);
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].ywxz;
+ (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
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 d5fae9d..5893867 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
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+ half3x2 tint_symbol;
+};
+struct tint_module_vars_struct {
+ const constant half3x2* u;
+ threadgroup half3x2* w;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat3x2<f16>, read> = var @binding_point(0, 0)
- %w:ptr<workgroup, mat3x2<f16>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
- $B2: {
- %5:bool = eq %tint_local_index, 0u
- if %5 [t: $B3] { # if_1
- $B3: { # true
- store %w, mat3x2<f16>(vec2<f16>(0.0h))
- exit_if # if_1
- }
- }
- %6:void = msl.threadgroup_barrier 4u
- %7:mat3x2<f16> = load %u
- store %w, %7
- %8:ptr<workgroup, vec2<f16>, read_write> = access %w, 1i
- %9:ptr<uniform, vec2<f16>, read> = access %u, 0i
- %10:vec2<f16> = load %9
- store %8, %10
- %11:ptr<workgroup, vec2<f16>, read_write> = access %w, 1i
- %12:ptr<uniform, vec2<f16>, read> = access %u, 0i
- %13:vec2<f16> = load %12
- %14:vec2<f16> = swizzle %13, yx
- store %11, %14
- %15:ptr<workgroup, vec2<f16>, read_write> = access %w, 0i
- %16:ptr<uniform, vec2<f16>, read> = access %u, 1i
- %17:f16 = load_vector_element %16, 0i
- store_vector_element %15, 1i, %17
- ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half3x2* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+ 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));
}
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ (*tint_module_vars.w) = (*tint_module_vars.u);
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].yx;
+ (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
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 45c60a8..da3e2fd 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
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+ float3x2 tint_symbol;
+};
+struct tint_module_vars_struct {
+ const constant float3x2* u;
+ threadgroup float3x2* w;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat3x2<f32>, read> = var @binding_point(0, 0)
- %w:ptr<workgroup, mat3x2<f32>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
- $B2: {
- %5:bool = eq %tint_local_index, 0u
- if %5 [t: $B3] { # if_1
- $B3: { # true
- store %w, mat3x2<f32>(vec2<f32>(0.0f))
- exit_if # if_1
- }
- }
- %6:void = msl.threadgroup_barrier 4u
- %7:mat3x2<f32> = load %u
- store %w, %7
- %8:ptr<workgroup, vec2<f32>, read_write> = access %w, 1i
- %9:ptr<uniform, vec2<f32>, read> = access %u, 0i
- %10:vec2<f32> = load %9
- store %8, %10
- %11:ptr<workgroup, vec2<f32>, read_write> = access %w, 1i
- %12:ptr<uniform, vec2<f32>, read> = access %u, 0i
- %13:vec2<f32> = load %12
- %14:vec2<f32> = swizzle %13, yx
- store %11, %14
- %15:ptr<workgroup, vec2<f32>, read_write> = access %w, 0i
- %16:ptr<uniform, vec2<f32>, read> = access %u, 1i
- %17:f32 = load_vector_element %16, 0i
- store_vector_element %15, 1i, %17
- ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float3x2* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+ 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));
}
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ (*tint_module_vars.w) = (*tint_module_vars.u);
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].yx;
+ (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
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 aef9a32..3276f46 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
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+ half3x3 tint_symbol;
+};
+struct tint_module_vars_struct {
+ const constant half3x3* u;
+ threadgroup half3x3* w;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat3x3<f16>, read> = var @binding_point(0, 0)
- %w:ptr<workgroup, mat3x3<f16>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
- $B2: {
- %5:bool = eq %tint_local_index, 0u
- if %5 [t: $B3] { # if_1
- $B3: { # true
- store %w, mat3x3<f16>(vec3<f16>(0.0h))
- exit_if # if_1
- }
- }
- %6:void = msl.threadgroup_barrier 4u
- %7:mat3x3<f16> = load %u
- store %w, %7
- %8:ptr<workgroup, vec3<f16>, read_write> = access %w, 1i
- %9:ptr<uniform, vec3<f16>, read> = access %u, 0i
- %10:vec3<f16> = load %9
- store %8, %10
- %11:ptr<workgroup, vec3<f16>, read_write> = access %w, 1i
- %12:ptr<uniform, vec3<f16>, read> = access %u, 0i
- %13:vec3<f16> = load %12
- %14:vec3<f16> = swizzle %13, zxy
- store %11, %14
- %15:ptr<workgroup, vec3<f16>, read_write> = access %w, 0i
- %16:ptr<uniform, vec3<f16>, read> = access %u, 1i
- %17:f16 = load_vector_element %16, 0i
- store_vector_element %15, 1i, %17
- ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half3x3* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+ 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));
}
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ (*tint_module_vars.w) = (*tint_module_vars.u);
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].zxy;
+ (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
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 db86b8c..105e86d 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
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+ float3x3 tint_symbol;
+};
+struct tint_module_vars_struct {
+ const constant float3x3* u;
+ threadgroup float3x3* w;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat3x3<f32>, read> = var @binding_point(0, 0)
- %w:ptr<workgroup, mat3x3<f32>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
- $B2: {
- %5:bool = eq %tint_local_index, 0u
- if %5 [t: $B3] { # if_1
- $B3: { # true
- store %w, mat3x3<f32>(vec3<f32>(0.0f))
- exit_if # if_1
- }
- }
- %6:void = msl.threadgroup_barrier 4u
- %7:mat3x3<f32> = load %u
- store %w, %7
- %8:ptr<workgroup, vec3<f32>, read_write> = access %w, 1i
- %9:ptr<uniform, vec3<f32>, read> = access %u, 0i
- %10:vec3<f32> = load %9
- store %8, %10
- %11:ptr<workgroup, vec3<f32>, read_write> = access %w, 1i
- %12:ptr<uniform, vec3<f32>, read> = access %u, 0i
- %13:vec3<f32> = load %12
- %14:vec3<f32> = swizzle %13, zxy
- store %11, %14
- %15:ptr<workgroup, vec3<f32>, read_write> = access %w, 0i
- %16:ptr<uniform, vec3<f32>, read> = access %u, 1i
- %17:f32 = load_vector_element %16, 0i
- store_vector_element %15, 1i, %17
- ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float3x3* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+ 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));
}
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ (*tint_module_vars.w) = (*tint_module_vars.u);
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].zxy;
+ (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
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 6f09501..fce2da8 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
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+ half3x4 tint_symbol;
+};
+struct tint_module_vars_struct {
+ const constant half3x4* u;
+ threadgroup half3x4* w;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat3x4<f16>, read> = var @binding_point(0, 0)
- %w:ptr<workgroup, mat3x4<f16>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
- $B2: {
- %5:bool = eq %tint_local_index, 0u
- if %5 [t: $B3] { # if_1
- $B3: { # true
- store %w, mat3x4<f16>(vec4<f16>(0.0h))
- exit_if # if_1
- }
- }
- %6:void = msl.threadgroup_barrier 4u
- %7:mat3x4<f16> = load %u
- store %w, %7
- %8:ptr<workgroup, vec4<f16>, read_write> = access %w, 1i
- %9:ptr<uniform, vec4<f16>, read> = access %u, 0i
- %10:vec4<f16> = load %9
- store %8, %10
- %11:ptr<workgroup, vec4<f16>, read_write> = access %w, 1i
- %12:ptr<uniform, vec4<f16>, read> = access %u, 0i
- %13:vec4<f16> = load %12
- %14:vec4<f16> = swizzle %13, ywxz
- store %11, %14
- %15:ptr<workgroup, vec4<f16>, read_write> = access %w, 0i
- %16:ptr<uniform, vec4<f16>, read> = access %u, 1i
- %17:f16 = load_vector_element %16, 0i
- store_vector_element %15, 1i, %17
- ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half3x4* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+ 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));
}
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ (*tint_module_vars.w) = (*tint_module_vars.u);
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].ywxz;
+ (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
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 675d4d3..bc3018c 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
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+ float3x4 tint_symbol;
+};
+struct tint_module_vars_struct {
+ const constant float3x4* u;
+ threadgroup float3x4* w;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat3x4<f32>, read> = var @binding_point(0, 0)
- %w:ptr<workgroup, mat3x4<f32>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
- $B2: {
- %5:bool = eq %tint_local_index, 0u
- if %5 [t: $B3] { # if_1
- $B3: { # true
- store %w, mat3x4<f32>(vec4<f32>(0.0f))
- exit_if # if_1
- }
- }
- %6:void = msl.threadgroup_barrier 4u
- %7:mat3x4<f32> = load %u
- store %w, %7
- %8:ptr<workgroup, vec4<f32>, read_write> = access %w, 1i
- %9:ptr<uniform, vec4<f32>, read> = access %u, 0i
- %10:vec4<f32> = load %9
- store %8, %10
- %11:ptr<workgroup, vec4<f32>, read_write> = access %w, 1i
- %12:ptr<uniform, vec4<f32>, read> = access %u, 0i
- %13:vec4<f32> = load %12
- %14:vec4<f32> = swizzle %13, ywxz
- store %11, %14
- %15:ptr<workgroup, vec4<f32>, read_write> = access %w, 0i
- %16:ptr<uniform, vec4<f32>, read> = access %u, 1i
- %17:f32 = load_vector_element %16, 0i
- store_vector_element %15, 1i, %17
- ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float3x4* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+ 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));
}
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ (*tint_module_vars.w) = (*tint_module_vars.u);
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].ywxz;
+ (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
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 ced9ab1..2a21f89 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
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+ half4x2 tint_symbol;
+};
+struct tint_module_vars_struct {
+ const constant half4x2* u;
+ threadgroup half4x2* w;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat4x2<f16>, read> = var @binding_point(0, 0)
- %w:ptr<workgroup, mat4x2<f16>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
- $B2: {
- %5:bool = eq %tint_local_index, 0u
- if %5 [t: $B3] { # if_1
- $B3: { # true
- store %w, mat4x2<f16>(vec2<f16>(0.0h))
- exit_if # if_1
- }
- }
- %6:void = msl.threadgroup_barrier 4u
- %7:mat4x2<f16> = load %u
- store %w, %7
- %8:ptr<workgroup, vec2<f16>, read_write> = access %w, 1i
- %9:ptr<uniform, vec2<f16>, read> = access %u, 0i
- %10:vec2<f16> = load %9
- store %8, %10
- %11:ptr<workgroup, vec2<f16>, read_write> = access %w, 1i
- %12:ptr<uniform, vec2<f16>, read> = access %u, 0i
- %13:vec2<f16> = load %12
- %14:vec2<f16> = swizzle %13, yx
- store %11, %14
- %15:ptr<workgroup, vec2<f16>, read_write> = access %w, 0i
- %16:ptr<uniform, vec2<f16>, read> = access %u, 1i
- %17:f16 = load_vector_element %16, 0i
- store_vector_element %15, 1i, %17
- ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half4x2* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+ 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));
}
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ (*tint_module_vars.w) = (*tint_module_vars.u);
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].yx;
+ (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
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 ff6fa47..5ffe4bd 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
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+ float4x2 tint_symbol;
+};
+struct tint_module_vars_struct {
+ const constant float4x2* u;
+ threadgroup float4x2* w;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat4x2<f32>, read> = var @binding_point(0, 0)
- %w:ptr<workgroup, mat4x2<f32>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
- $B2: {
- %5:bool = eq %tint_local_index, 0u
- if %5 [t: $B3] { # if_1
- $B3: { # true
- store %w, mat4x2<f32>(vec2<f32>(0.0f))
- exit_if # if_1
- }
- }
- %6:void = msl.threadgroup_barrier 4u
- %7:mat4x2<f32> = load %u
- store %w, %7
- %8:ptr<workgroup, vec2<f32>, read_write> = access %w, 1i
- %9:ptr<uniform, vec2<f32>, read> = access %u, 0i
- %10:vec2<f32> = load %9
- store %8, %10
- %11:ptr<workgroup, vec2<f32>, read_write> = access %w, 1i
- %12:ptr<uniform, vec2<f32>, read> = access %u, 0i
- %13:vec2<f32> = load %12
- %14:vec2<f32> = swizzle %13, yx
- store %11, %14
- %15:ptr<workgroup, vec2<f32>, read_write> = access %w, 0i
- %16:ptr<uniform, vec2<f32>, read> = access %u, 1i
- %17:f32 = load_vector_element %16, 0i
- store_vector_element %15, 1i, %17
- ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float4x2* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+ 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));
}
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ (*tint_module_vars.w) = (*tint_module_vars.u);
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].yx;
+ (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
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 4ec8615..2a836c1 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
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+ half4x3 tint_symbol;
+};
+struct tint_module_vars_struct {
+ const constant half4x3* u;
+ threadgroup half4x3* w;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat4x3<f16>, read> = var @binding_point(0, 0)
- %w:ptr<workgroup, mat4x3<f16>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
- $B2: {
- %5:bool = eq %tint_local_index, 0u
- if %5 [t: $B3] { # if_1
- $B3: { # true
- store %w, mat4x3<f16>(vec3<f16>(0.0h))
- exit_if # if_1
- }
- }
- %6:void = msl.threadgroup_barrier 4u
- %7:mat4x3<f16> = load %u
- store %w, %7
- %8:ptr<workgroup, vec3<f16>, read_write> = access %w, 1i
- %9:ptr<uniform, vec3<f16>, read> = access %u, 0i
- %10:vec3<f16> = load %9
- store %8, %10
- %11:ptr<workgroup, vec3<f16>, read_write> = access %w, 1i
- %12:ptr<uniform, vec3<f16>, read> = access %u, 0i
- %13:vec3<f16> = load %12
- %14:vec3<f16> = swizzle %13, zxy
- store %11, %14
- %15:ptr<workgroup, vec3<f16>, read_write> = access %w, 0i
- %16:ptr<uniform, vec3<f16>, read> = access %u, 1i
- %17:f16 = load_vector_element %16, 0i
- store_vector_element %15, 1i, %17
- ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half4x3* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+ 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));
}
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ (*tint_module_vars.w) = (*tint_module_vars.u);
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].zxy;
+ (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
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 ded9db0..1d930a4 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
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+ float4x3 tint_symbol;
+};
+struct tint_module_vars_struct {
+ const constant float4x3* u;
+ threadgroup float4x3* w;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat4x3<f32>, read> = var @binding_point(0, 0)
- %w:ptr<workgroup, mat4x3<f32>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
- $B2: {
- %5:bool = eq %tint_local_index, 0u
- if %5 [t: $B3] { # if_1
- $B3: { # true
- store %w, mat4x3<f32>(vec3<f32>(0.0f))
- exit_if # if_1
- }
- }
- %6:void = msl.threadgroup_barrier 4u
- %7:mat4x3<f32> = load %u
- store %w, %7
- %8:ptr<workgroup, vec3<f32>, read_write> = access %w, 1i
- %9:ptr<uniform, vec3<f32>, read> = access %u, 0i
- %10:vec3<f32> = load %9
- store %8, %10
- %11:ptr<workgroup, vec3<f32>, read_write> = access %w, 1i
- %12:ptr<uniform, vec3<f32>, read> = access %u, 0i
- %13:vec3<f32> = load %12
- %14:vec3<f32> = swizzle %13, zxy
- store %11, %14
- %15:ptr<workgroup, vec3<f32>, read_write> = access %w, 0i
- %16:ptr<uniform, vec3<f32>, read> = access %u, 1i
- %17:f32 = load_vector_element %16, 0i
- store_vector_element %15, 1i, %17
- ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float4x3* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+ 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));
}
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ (*tint_module_vars.w) = (*tint_module_vars.u);
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].zxy;
+ (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
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 68cc6fd..67b0ff7 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
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+ half4x4 tint_symbol;
+};
+struct tint_module_vars_struct {
+ const constant half4x4* u;
+ threadgroup half4x4* w;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat4x4<f16>, read> = var @binding_point(0, 0)
- %w:ptr<workgroup, mat4x4<f16>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
- $B2: {
- %5:bool = eq %tint_local_index, 0u
- if %5 [t: $B3] { # if_1
- $B3: { # true
- store %w, mat4x4<f16>(vec4<f16>(0.0h))
- exit_if # if_1
- }
- }
- %6:void = msl.threadgroup_barrier 4u
- %7:mat4x4<f16> = load %u
- store %w, %7
- %8:ptr<workgroup, vec4<f16>, read_write> = access %w, 1i
- %9:ptr<uniform, vec4<f16>, read> = access %u, 0i
- %10:vec4<f16> = load %9
- store %8, %10
- %11:ptr<workgroup, vec4<f16>, read_write> = access %w, 1i
- %12:ptr<uniform, vec4<f16>, read> = access %u, 0i
- %13:vec4<f16> = load %12
- %14:vec4<f16> = swizzle %13, ywxz
- store %11, %14
- %15:ptr<workgroup, vec4<f16>, read_write> = access %w, 0i
- %16:ptr<uniform, vec4<f16>, read> = access %u, 1i
- %17:f16 = load_vector_element %16, 0i
- store_vector_element %15, 1i, %17
- ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half4x4* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+ 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));
}
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ (*tint_module_vars.w) = (*tint_module_vars.u);
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].ywxz;
+ (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
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 9818ee4..1143198 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
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+ float4x4 tint_symbol;
+};
+struct tint_module_vars_struct {
+ const constant float4x4* u;
+ threadgroup float4x4* w;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat4x4<f32>, read> = var @binding_point(0, 0)
- %w:ptr<workgroup, mat4x4<f32>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
- $B2: {
- %5:bool = eq %tint_local_index, 0u
- if %5 [t: $B3] { # if_1
- $B3: { # true
- store %w, mat4x4<f32>(vec4<f32>(0.0f))
- exit_if # if_1
- }
- }
- %6:void = msl.threadgroup_barrier 4u
- %7:mat4x4<f32> = load %u
- store %w, %7
- %8:ptr<workgroup, vec4<f32>, read_write> = access %w, 1i
- %9:ptr<uniform, vec4<f32>, read> = access %u, 0i
- %10:vec4<f32> = load %9
- store %8, %10
- %11:ptr<workgroup, vec4<f32>, read_write> = access %w, 1i
- %12:ptr<uniform, vec4<f32>, read> = access %u, 0i
- %13:vec4<f32> = load %12
- %14:vec4<f32> = swizzle %13, ywxz
- store %11, %14
- %15:ptr<workgroup, vec4<f32>, read_write> = access %w, 0i
- %16:ptr<uniform, vec4<f32>, read> = access %u, 1i
- %17:f32 = load_vector_element %16, 0i
- store_vector_element %15, 1i, %17
- ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float4x4* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+ 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));
}
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ (*tint_module_vars.w) = (*tint_module_vars.u);
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+ (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].ywxz;
+ (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/bug/chromium/40943165.wgsl.expected.ir.msl b/test/tint/bug/chromium/40943165.wgsl.expected.ir.msl
index fbd61a0..a7b5929 100644
--- a/test/tint/bug/chromium/40943165.wgsl.expected.ir.msl
+++ b/test/tint/bug/chromium/40943165.wgsl.expected.ir.msl
@@ -1,27 +1,17 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
+struct tint_symbol_1 {
+ float2x2 tint_symbol;
+};
+struct tint_module_vars_struct {
+ threadgroup float2x2* W;
+};
-threadgroup float2x2 W;
-kernel void F(uint mat2x2 [[thread_index_in_threadgroup]]) {
+kernel void F(uint mat2x2 [[thread_index_in_threadgroup]], threadgroup tint_symbol_1* v) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.W=(&(*v).tint_symbol)};
if ((mat2x2 == 0u)) {
- W = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.W) = float2x2(float2(0.0f), float2(0.0f));
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- W[0] = (W[0] + 0.0f);
+ (*tint_module_vars.W)[0] = ((*tint_module_vars.W)[0] + 0.0f);
}
-program_source:4:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 W;
- ^
-program_source:6:15: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
- if ((mat2x2 == 0u)) {
- ~~~~~~~^~~~~
-program_source:6:15: note: remove extraneous parentheses around the comparison to silence this warning
- if ((mat2x2 == 0u)) {
- ~ ^ ~
-program_source:6:15: note: use '=' to turn this equality comparison into an assignment
- if ((mat2x2 == 0u)) {
- ^~
- =
-
diff --git a/test/tint/bug/tint/1926.wgsl.expected.ir.msl b/test/tint/bug/tint/1926.wgsl.expected.ir.msl
index c450d39..812684c 100644
--- a/test/tint/bug/tint/1926.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1926.wgsl.expected.ir.msl
@@ -1,41 +1,24 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_2 {
+ uint tint_symbol_1;
+};
+struct tint_module_vars_struct {
+ threadgroup uint* sh_atomic_failed;
+ device uint* output;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %sh_atomic_failed:ptr<workgroup, u32, read_write> = var
- %output:ptr<storage, u32, read_write> = var @binding_point(0, 4)
-}
-
-%tint_symbol = @compute @workgroup_size(256, 1, 1) func(%global_id:vec3<u32> [@global_invocation_id], %local_id:vec3<u32> [@local_invocation_id], %tint_local_index:u32 [@local_invocation_index]):void {
- $B2: {
- %7:bool = eq %tint_local_index, 0u
- if %7 [t: $B3] { # if_1
- $B3: { # true
- store %sh_atomic_failed, 0u
- exit_if # if_1
- }
- }
- %8:void = msl.threadgroup_barrier 4u
- %9:void = msl.threadgroup_barrier 4u
- %10:u32 = load %sh_atomic_failed
- %11:u32 = let %10
- %12:void = msl.threadgroup_barrier 4u
- %failed:u32 = let %11
- %14:u32 = access %local_id, 0u
- %15:bool = eq %14, 0u
- if %15 [t: $B4] { # if_2
- $B4: { # true
- store %output, %failed
- exit_if # if_2
- }
- }
- ret
+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)]]) {
+ 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;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ uint const v_1 = (*tint_module_vars.sh_atomic_failed);
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ uint const failed = v_1;
+ if ((local_id[0u] == 0u)) {
+ (*tint_module_vars.output) = failed;
}
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
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 5a6429d..5f96aa4 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
@@ -1,32 +1,19 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
+struct tint_symbol_2 {
+ int tint_symbol_1;
+};
+struct tint_module_vars_struct {
+ threadgroup int* i;
+};
-threadgroup int i;
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.i=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
- i = 0;
+ (*tint_module_vars.i) = 0;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- i = 123;
- threadgroup int* const p = i;
- int const u = (p + 1);
+ (*tint_module_vars.i) = 123;
+ threadgroup int* const p = tint_module_vars.i;
+ int const u = ((*p) + 1);
}
-program_source:4:17: error: program scope variable must reside in constant address space
-threadgroup int i;
- ^
-program_source:6:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
- if ((tint_local_index == 0u)) {
- ~~~~~~~~~~~~~~~~~^~~~~
-program_source:6:25: note: remove extraneous parentheses around the comparison to silence this warning
- if ((tint_local_index == 0u)) {
- ~ ^ ~
-program_source:6:25: note: use '=' to turn this equality comparison into an assignment
- if ((tint_local_index == 0u)) {
- ^~
- =
-program_source:12:13: error: cannot initialize a variable of type 'const int' with an rvalue of type 'threadgroup int *'
- int const u = (p + 1);
- ^ ~~~~~~~
-
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 25d79e7..fbd0e7d 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
@@ -1,36 +1,20 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
+struct tint_symbol_2 {
+ int tint_symbol_1;
+};
+struct tint_module_vars_struct {
+ threadgroup int* S;
+};
-threadgroup int S;
int func(threadgroup int* const pointer) {
- return pointer;
+ return (*pointer);
}
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
- S = 0;
+ (*tint_module_vars.S) = 0;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- int const r = func(S);
+ int const r = func(tint_module_vars.S);
}
-program_source:4:17: error: program scope variable must reside in constant address space
-threadgroup int S;
- ^
-program_source:6:10: error: cannot initialize return object of type 'int' with an lvalue of type 'threadgroup int *const'
- return pointer;
- ^~~~~~~
-program_source:9:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
- if ((tint_local_index == 0u)) {
- ~~~~~~~~~~~~~~~~~^~~~~
-program_source:9:25: note: remove extraneous parentheses around the comparison to silence this warning
- if ((tint_local_index == 0u)) {
- ~ ^ ~
-program_source:9:25: note: use '=' to turn this equality comparison into an assignment
- if ((tint_local_index == 0u)) {
- ^~
- =
-program_source:13:13: warning: unused variable 'r' [-Wunused-variable]
- int const r = func(S);
- ^
-
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 4be387c..0e25c93 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
@@ -1,39 +1,23 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
struct str {
int i;
};
+struct tint_symbol_2 {
+ str tint_symbol_1;
+};
+struct tint_module_vars_struct {
+ threadgroup str* S;
+};
-threadgroup str S;
int func(threadgroup int* const pointer) {
- return pointer;
+ return (*pointer);
}
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
- S = str{};
+ (*tint_module_vars.S) = str{};
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- int const r = func(S.i);
+ int const r = func((&(*tint_module_vars.S).i));
}
-program_source:7:17: error: program scope variable must reside in constant address space
-threadgroup str S;
- ^
-program_source:9:10: error: cannot initialize return object of type 'int' with an lvalue of type 'threadgroup int *const'
- return pointer;
- ^~~~~~~
-program_source:12:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
- if ((tint_local_index == 0u)) {
- ~~~~~~~~~~~~~~~~~^~~~~
-program_source:12:25: note: remove extraneous parentheses around the comparison to silence this warning
- if ((tint_local_index == 0u)) {
- ~ ^ ~
-program_source:12:25: note: use '=' to turn this equality comparison into an assignment
- if ((tint_local_index == 0u)) {
- ^~
- =
-program_source:16:13: warning: unused variable 'r' [-Wunused-variable]
- int const r = func(S.i);
- ^
-
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 6866579..9e76b8b 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
@@ -1,36 +1,20 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
+struct tint_symbol_2 {
+ float2x2 tint_symbol_1;
+};
+struct tint_module_vars_struct {
+ threadgroup float2x2* S;
+};
-threadgroup float2x2 S;
float2 func(threadgroup float2* const pointer) {
- return pointer;
+ return (*pointer);
}
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
- S = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.S) = float2x2(float2(0.0f), float2(0.0f));
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- float2 const r = func(S[1]);
+ float2 const r = func((&(*tint_module_vars.S)[1]));
}
-program_source:4:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 S;
- ^
-program_source:6:10: error: cannot initialize return object of type 'float2' (vector of 2 'float' values) with an lvalue of type 'threadgroup float2 *const'
- return pointer;
- ^~~~~~~
-program_source:9:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
- if ((tint_local_index == 0u)) {
- ~~~~~~~~~~~~~~~~~^~~~~
-program_source:9:25: note: remove extraneous parentheses around the comparison to silence this warning
- if ((tint_local_index == 0u)) {
- ~ ^ ~
-program_source:9:25: note: use '=' to turn this equality comparison into an assignment
- if ((tint_local_index == 0u)) {
- ^~
- =
-program_source:13:16: warning: unused variable 'r' [-Wunused-variable]
- float2 const r = func(S[1]);
- ^
-
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 1e11dc4..3764937 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
@@ -1,36 +1,20 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
+struct tint_symbol_2 {
+ float4 tint_symbol_1;
+};
+struct tint_module_vars_struct {
+ threadgroup float4* S;
+};
-threadgroup float4 S;
float4 func(threadgroup float4* const pointer) {
- return pointer;
+ return (*pointer);
}
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
- S = float4(0.0f);
+ (*tint_module_vars.S) = float4(0.0f);
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- float4 const r = func(S);
+ float4 const r = func(tint_module_vars.S);
}
-program_source:4:20: error: program scope variable must reside in constant address space
-threadgroup float4 S;
- ^
-program_source:6:10: error: cannot initialize return object of type 'float4' (vector of 4 'float' values) with an lvalue of type 'threadgroup float4 *const'
- return pointer;
- ^~~~~~~
-program_source:9:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
- if ((tint_local_index == 0u)) {
- ~~~~~~~~~~~~~~~~~^~~~~
-program_source:9:25: note: remove extraneous parentheses around the comparison to silence this warning
- if ((tint_local_index == 0u)) {
- ~ ^ ~
-program_source:9:25: note: use '=' to turn this equality comparison into an assignment
- if ((tint_local_index == 0u)) {
- ^~
- =
-program_source:13:16: warning: unused variable 'r' [-Wunused-variable]
- float4 const r = func(S);
- ^
-
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 a6f5b85..fb1ab2d 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
@@ -1,36 +1,20 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
+struct tint_symbol_2 {
+ float2x4 tint_symbol_1;
+};
+struct tint_module_vars_struct {
+ threadgroup float2x4* S;
+};
-threadgroup float2x4 S;
float4 func(threadgroup float4* const pointer) {
- return pointer;
+ return (*pointer);
}
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
- S = float2x4(float4(0.0f), float4(0.0f));
+ (*tint_module_vars.S) = float2x4(float4(0.0f), float4(0.0f));
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- float4 const r = func(S[1]);
+ float4 const r = func((&(*tint_module_vars.S)[1]));
}
-program_source:4:22: error: program scope variable must reside in constant address space
-threadgroup float2x4 S;
- ^
-program_source:6:10: error: cannot initialize return object of type 'float4' (vector of 4 'float' values) with an lvalue of type 'threadgroup float4 *const'
- return pointer;
- ^~~~~~~
-program_source:9:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
- if ((tint_local_index == 0u)) {
- ~~~~~~~~~~~~~~~~~^~~~~
-program_source:9:25: note: remove extraneous parentheses around the comparison to silence this warning
- if ((tint_local_index == 0u)) {
- ~ ^ ~
-program_source:9:25: note: use '=' to turn this equality comparison into an assignment
- if ((tint_local_index == 0u)) {
- ^~
- =
-program_source:13:16: warning: unused variable 'r' [-Wunused-variable]
- float4 const r = func(S[1]);
- ^
-
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 62142f9..96d683f 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
@@ -1,39 +1,23 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
struct str {
float4 i;
};
+struct tint_symbol_2 {
+ str tint_symbol_1;
+};
+struct tint_module_vars_struct {
+ threadgroup str* S;
+};
-threadgroup str S;
float4 func(threadgroup float4* const pointer) {
- return pointer;
+ return (*pointer);
}
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
- S = str{};
+ (*tint_module_vars.S) = str{};
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- float4 const r = func(S.i);
+ float4 const r = func((&(*tint_module_vars.S).i));
}
-program_source:7:17: error: program scope variable must reside in constant address space
-threadgroup str S;
- ^
-program_source:9:10: error: cannot initialize return object of type 'float4' (vector of 4 'float' values) with an lvalue of type 'threadgroup float4 *const'
- return pointer;
- ^~~~~~~
-program_source:12:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
- if ((tint_local_index == 0u)) {
- ~~~~~~~~~~~~~~~~~^~~~~
-program_source:12:25: note: remove extraneous parentheses around the comparison to silence this warning
- if ((tint_local_index == 0u)) {
- ~ ^ ~
-program_source:12:25: note: use '=' to turn this equality comparison into an assignment
- if ((tint_local_index == 0u)) {
- ^~
- =
-program_source:16:16: warning: unused variable 'r' [-Wunused-variable]
- float4 const r = func(S.i);
- ^
-
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 b5cca68..57dbb78 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
@@ -1,36 +1,20 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
+struct tint_symbol_2 {
+ int tint_symbol_1;
+};
+struct tint_module_vars_struct {
+ threadgroup int* S;
+};
-threadgroup int S;
void func(threadgroup int* const pointer) {
- pointer = 42;
+ (*pointer) = 42;
}
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
- S = 0;
+ (*tint_module_vars.S) = 0;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- func(S);
+ func(tint_module_vars.S);
}
-program_source:4:17: error: program scope variable must reside in constant address space
-threadgroup int S;
- ^
-program_source:6:11: error: cannot assign to variable 'pointer' with const-qualified type 'threadgroup int *const'
- pointer = 42;
- ~~~~~~~ ^
-program_source:5:34: note: variable 'pointer' declared const here
-void func(threadgroup int* const pointer) {
- ~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~
-program_source:9:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
- if ((tint_local_index == 0u)) {
- ~~~~~~~~~~~~~~~~~^~~~~
-program_source:9:25: note: remove extraneous parentheses around the comparison to silence this warning
- if ((tint_local_index == 0u)) {
- ~ ^ ~
-program_source:9:25: note: use '=' to turn this equality comparison into an assignment
- if ((tint_local_index == 0u)) {
- ^~
- =
-
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 cc3cff5..1defcd2 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
@@ -1,39 +1,23 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
struct str {
int i;
};
+struct tint_symbol_2 {
+ str tint_symbol_1;
+};
+struct tint_module_vars_struct {
+ threadgroup str* S;
+};
-threadgroup str S;
void func(threadgroup int* const pointer) {
- pointer = 42;
+ (*pointer) = 42;
}
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
- S = str{};
+ (*tint_module_vars.S) = str{};
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- func(S.i);
+ func((&(*tint_module_vars.S).i));
}
-program_source:7:17: error: program scope variable must reside in constant address space
-threadgroup str S;
- ^
-program_source:9:11: error: cannot assign to variable 'pointer' with const-qualified type 'threadgroup int *const'
- pointer = 42;
- ~~~~~~~ ^
-program_source:8:34: note: variable 'pointer' declared const here
-void func(threadgroup int* const pointer) {
- ~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~
-program_source:12:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
- if ((tint_local_index == 0u)) {
- ~~~~~~~~~~~~~~~~~^~~~~
-program_source:12:25: note: remove extraneous parentheses around the comparison to silence this warning
- if ((tint_local_index == 0u)) {
- ~ ^ ~
-program_source:12:25: note: use '=' to turn this equality comparison into an assignment
- if ((tint_local_index == 0u)) {
- ^~
- =
-
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 ef57e5a..f089435 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
@@ -1,36 +1,20 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
+struct tint_symbol_2 {
+ float2x2 tint_symbol_1;
+};
+struct tint_module_vars_struct {
+ threadgroup float2x2* S;
+};
-threadgroup float2x2 S;
void func(threadgroup float2* const pointer) {
- pointer = float2(0.0f);
+ (*pointer) = float2(0.0f);
}
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
- S = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.S) = float2x2(float2(0.0f), float2(0.0f));
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- func(S[1]);
+ func((&(*tint_module_vars.S)[1]));
}
-program_source:4:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 S;
- ^
-program_source:6:11: error: cannot assign to variable 'pointer' with const-qualified type 'threadgroup float2 *const'
- pointer = float2(0.0f);
- ~~~~~~~ ^
-program_source:5:37: note: variable 'pointer' declared const here
-void func(threadgroup float2* const pointer) {
- ~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~
-program_source:9:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
- if ((tint_local_index == 0u)) {
- ~~~~~~~~~~~~~~~~~^~~~~
-program_source:9:25: note: remove extraneous parentheses around the comparison to silence this warning
- if ((tint_local_index == 0u)) {
- ~ ^ ~
-program_source:9:25: note: use '=' to turn this equality comparison into an assignment
- if ((tint_local_index == 0u)) {
- ^~
- =
-
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 2de34c8..416f51b 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
@@ -1,36 +1,20 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
+struct tint_symbol_2 {
+ float4 tint_symbol_1;
+};
+struct tint_module_vars_struct {
+ threadgroup float4* S;
+};
-threadgroup float4 S;
void func(threadgroup float4* const pointer) {
- pointer = float4(0.0f);
+ (*pointer) = float4(0.0f);
}
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
- S = float4(0.0f);
+ (*tint_module_vars.S) = float4(0.0f);
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- func(S);
+ func(tint_module_vars.S);
}
-program_source:4:20: error: program scope variable must reside in constant address space
-threadgroup float4 S;
- ^
-program_source:6:11: error: cannot assign to variable 'pointer' with const-qualified type 'threadgroup float4 *const'
- pointer = float4(0.0f);
- ~~~~~~~ ^
-program_source:5:37: note: variable 'pointer' declared const here
-void func(threadgroup float4* const pointer) {
- ~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~
-program_source:9:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
- if ((tint_local_index == 0u)) {
- ~~~~~~~~~~~~~~~~~^~~~~
-program_source:9:25: note: remove extraneous parentheses around the comparison to silence this warning
- if ((tint_local_index == 0u)) {
- ~ ^ ~
-program_source:9:25: note: use '=' to turn this equality comparison into an assignment
- if ((tint_local_index == 0u)) {
- ^~
- =
-
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 0f5c1cc..c4d6342 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
@@ -1,36 +1,20 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
+struct tint_symbol_2 {
+ float2x4 tint_symbol_1;
+};
+struct tint_module_vars_struct {
+ threadgroup float2x4* S;
+};
-threadgroup float2x4 S;
void func(threadgroup float4* const pointer) {
- pointer = float4(0.0f);
+ (*pointer) = float4(0.0f);
}
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
- S = float2x4(float4(0.0f), float4(0.0f));
+ (*tint_module_vars.S) = float2x4(float4(0.0f), float4(0.0f));
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- func(S[1]);
+ func((&(*tint_module_vars.S)[1]));
}
-program_source:4:22: error: program scope variable must reside in constant address space
-threadgroup float2x4 S;
- ^
-program_source:6:11: error: cannot assign to variable 'pointer' with const-qualified type 'threadgroup float4 *const'
- pointer = float4(0.0f);
- ~~~~~~~ ^
-program_source:5:37: note: variable 'pointer' declared const here
-void func(threadgroup float4* const pointer) {
- ~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~
-program_source:9:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
- if ((tint_local_index == 0u)) {
- ~~~~~~~~~~~~~~~~~^~~~~
-program_source:9:25: note: remove extraneous parentheses around the comparison to silence this warning
- if ((tint_local_index == 0u)) {
- ~ ^ ~
-program_source:9:25: note: use '=' to turn this equality comparison into an assignment
- if ((tint_local_index == 0u)) {
- ^~
- =
-
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 9017799..e943602 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
@@ -1,39 +1,23 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
struct str {
float4 i;
};
+struct tint_symbol_2 {
+ str tint_symbol_1;
+};
+struct tint_module_vars_struct {
+ threadgroup str* S;
+};
-threadgroup str S;
void func(threadgroup float4* const pointer) {
- pointer = float4(0.0f);
+ (*pointer) = float4(0.0f);
}
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
- S = str{};
+ (*tint_module_vars.S) = str{};
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- func(S.i);
+ func((&(*tint_module_vars.S).i));
}
-program_source:7:17: error: program scope variable must reside in constant address space
-threadgroup str S;
- ^
-program_source:9:11: error: cannot assign to variable 'pointer' with const-qualified type 'threadgroup float4 *const'
- pointer = float4(0.0f);
- ~~~~~~~ ^
-program_source:8:37: note: variable 'pointer' declared const here
-void func(threadgroup float4* const pointer) {
- ~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~
-program_source:12:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
- if ((tint_local_index == 0u)) {
- ~~~~~~~~~~~~~~~~~^~~~~
-program_source:12:25: note: remove extraneous parentheses around the comparison to silence this warning
- if ((tint_local_index == 0u)) {
- ~ ^ ~
-program_source:12:25: note: use '=' to turn this equality comparison into an assignment
- if ((tint_local_index == 0u)) {
- ^~
- =
-
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 85e0261..2b6c649 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
@@ -1,63 +1,49 @@
-SKIP: FAILED
+#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];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %p:ptr<private, f32, read_write> = var
- %w:ptr<workgroup, f32, read_write> = var
- %uniforms:ptr<storage, vec2<f32>, read> = var @binding_point(0, 1)
- %storages:ptr<storage, array<f32>, read_write> = var @binding_point(0, 0)
-}
+struct tint_module_vars_struct {
+ thread float* p;
+ threadgroup float* w;
+ const device float2* uniforms;
+ device tint_array<float, 1>* storages;
+};
+struct tint_symbol_2 {
+ float tint_symbol_1;
+};
-%no_uses = func():void {
- $B2: {
- ret
- }
+void no_uses() {
}
-%zoo = func():void {
- $B3: {
- %7:f32 = load %p
- %8:f32 = mul %7, 2.0f
- store %p, %8
- ret
- }
+void zoo(tint_module_vars_struct tint_module_vars) {
+ (*tint_module_vars.p) = ((*tint_module_vars.p) * 2.0f);
}
-%bar = func(%a:f32, %b:f32):void {
- $B4: {
- store %p, %a
- store %w, %b
- %12:ptr<storage, f32, read_write> = access %storages, 0i
- %13:f32 = load_vector_element %uniforms, 0u
- store %12, %13
- %14:void = call %zoo
- ret
- }
+void bar(float a, float b, tint_module_vars_struct tint_module_vars) {
+ (*tint_module_vars.p) = a;
+ (*tint_module_vars.w) = b;
+ (*tint_module_vars.storages)[0] = (*tint_module_vars.uniforms)[0u];
+ zoo(tint_module_vars);
}
-%foo = func(%a_1:f32):void { # %a_1: 'a'
- $B5: {
- %b_1:f32 = let 2.0f # %b_1: 'b'
- %18:void = call %bar, %a_1, %b_1
- %19:void = call %no_uses
- ret
- }
+void foo(float a, tint_module_vars_struct tint_module_vars) {
+ float const b = 2.0f;
+ bar(a, b, tint_module_vars);
+ no_uses();
}
-%tint_symbol = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
- $B6: {
- %22:bool = eq %tint_local_index, 0u
- if %22 [t: $B7] { # if_1
- $B7: { # true
- store %w, 0.0f
- exit_if # if_1
- }
- }
- %23:void = msl.threadgroup_barrier 4u
- %24:void = call %foo, 1.0f
- ret
+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)]]) {
+ 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)) {
+ (*tint_module_vars.w) = 0.0f;
}
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ foo(1.0f, tint_module_vars);
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
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 e3cf1cd..b44443a 100644
--- a/test/tint/types/module_scope_var.wgsl.expected.ir.msl
+++ b/test/tint/types/module_scope_var.wgsl.expected.ir.msl
@@ -1,7 +1,8 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
+struct tint_symbol_2 {
+ float tint_symbol_1;
+};
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
@@ -17,76 +18,45 @@
struct S {
float a;
};
+struct tint_module_vars_struct {
+ thread bool* bool_var;
+ thread int* i32_var;
+ thread uint* u32_var;
+ thread float* f32_var;
+ thread int2* v2i32_var;
+ thread uint3* v3u32_var;
+ thread float4* v4f32_var;
+ thread float2x3* m2x3_var;
+ thread tint_array<float, 4>* arr_var;
+ thread S* struct_var;
+ threadgroup float* wg_var;
+};
-thread bool bool_var = false;
-thread int i32_var = 0;
-thread uint u32_var = 0u;
-thread float f32_var = 0.0f;
-thread int2 v2i32_var = 0;
-thread uint3 v3u32_var = 0u;
-thread float4 v4f32_var = 0.0f;
-thread float2x3 m2x3_var = float2x3(0.0f);
-thread tint_array<float, 4> arr_var = {};
-thread S struct_var = {};
-threadgroup float wg_var;
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+ thread bool bool_var = false;
+ thread int i32_var = 0;
+ thread uint u32_var = 0u;
+ thread float f32_var = 0.0f;
+ thread int2 v2i32_var = 0;
+ thread uint3 v3u32_var = 0u;
+ thread float4 v4f32_var = 0.0f;
+ thread float2x3 m2x3_var = float2x3(0.0f);
+ thread tint_array<float, 4> arr_var = {};
+ thread S struct_var = {};
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.bool_var=(&bool_var), .i32_var=(&i32_var), .u32_var=(&u32_var), .f32_var=(&f32_var), .v2i32_var=(&v2i32_var), .v3u32_var=(&v3u32_var), .v4f32_var=(&v4f32_var), .m2x3_var=(&m2x3_var), .arr_var=(&arr_var), .struct_var=(&struct_var), .wg_var=(&(*v).tint_symbol_1)};
if ((tint_local_index == 0u)) {
- wg_var = 0.0f;
+ (*tint_module_vars.wg_var) = 0.0f;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- bool_var = false;
- i32_var = 0;
- u32_var = 0u;
- f32_var = 0.0f;
- v2i32_var = int2(0);
- v3u32_var = uint3(0u);
- v4f32_var = float4(0.0f);
- m2x3_var = float2x3(float3(0.0f), float3(0.0f));
- arr_var = tint_array<float, 4>{};
- struct_var = S{};
- wg_var = 42.0f;
+ (*tint_module_vars.bool_var) = false;
+ (*tint_module_vars.i32_var) = 0;
+ (*tint_module_vars.u32_var) = 0u;
+ (*tint_module_vars.f32_var) = 0.0f;
+ (*tint_module_vars.v2i32_var) = int2(0);
+ (*tint_module_vars.v3u32_var) = uint3(0u);
+ (*tint_module_vars.v4f32_var) = float4(0.0f);
+ (*tint_module_vars.m2x3_var) = float2x3(float3(0.0f), float3(0.0f));
+ (*tint_module_vars.arr_var) = tint_array<float, 4>{};
+ (*tint_module_vars.struct_var) = S{};
+ (*tint_module_vars.wg_var) = 42.0f;
}
-program_source:19:13: error: program scope variable must reside in constant address space
-thread bool bool_var = false;
- ^
-program_source:20:12: error: program scope variable must reside in constant address space
-thread int i32_var = 0;
- ^
-program_source:21:13: error: program scope variable must reside in constant address space
-thread uint u32_var = 0u;
- ^
-program_source:22:14: error: program scope variable must reside in constant address space
-thread float f32_var = 0.0f;
- ^
-program_source:23:13: error: program scope variable must reside in constant address space
-thread int2 v2i32_var = 0;
- ^
-program_source:24:14: error: program scope variable must reside in constant address space
-thread uint3 v3u32_var = 0u;
- ^
-program_source:25:15: error: program scope variable must reside in constant address space
-thread float4 v4f32_var = 0.0f;
- ^
-program_source:26:17: error: program scope variable must reside in constant address space
-thread float2x3 m2x3_var = float2x3(0.0f);
- ^
-program_source:27:29: error: program scope variable must reside in constant address space
-thread tint_array<float, 4> arr_var = {};
- ^
-program_source:28:10: error: program scope variable must reside in constant address space
-thread S struct_var = {};
- ^
-program_source:29:19: error: program scope variable must reside in constant address space
-threadgroup float wg_var;
- ^
-program_source:31:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
- if ((tint_local_index == 0u)) {
- ~~~~~~~~~~~~~~~~~^~~~~
-program_source:31:25: note: remove extraneous parentheses around the comparison to silence this warning
- if ((tint_local_index == 0u)) {
- ~ ^ ~
-program_source:31:25: note: use '=' to turn this equality comparison into an assignment
- if ((tint_local_index == 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 68d0820..bc857e6 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
@@ -1,43 +1,22 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
+struct tint_symbol_2 {
+ float tint_symbol_1;
+};
+struct tint_module_vars_struct {
+ thread float* p;
+ threadgroup float* w;
+};
-thread float p = 0.0f;
-threadgroup float w;
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+ 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)) {
- w = 0.0f;
+ (*tint_module_vars.w) = 0.0f;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- thread float* const p_ptr = p;
- threadgroup float* const w_ptr = w;
- float const x = (p_ptr + w_ptr);
- p_ptr = x;
+ thread float* const p_ptr = tint_module_vars.p;
+ threadgroup float* const w_ptr = tint_module_vars.w;
+ float const x = ((*p_ptr) + (*w_ptr));
+ (*p_ptr) = x;
}
-program_source:4:14: error: program scope variable must reside in constant address space
-thread float p = 0.0f;
- ^
-program_source:5:19: error: program scope variable must reside in constant address space
-threadgroup float w;
- ^
-program_source:7:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
- if ((tint_local_index == 0u)) {
- ~~~~~~~~~~~~~~~~~^~~~~
-program_source:7:25: note: remove extraneous parentheses around the comparison to silence this warning
- if ((tint_local_index == 0u)) {
- ~ ^ ~
-program_source:7:25: note: use '=' to turn this equality comparison into an assignment
- if ((tint_local_index == 0u)) {
- ^~
- =
-program_source:13:26: error: invalid operands to binary expression ('float *const' and 'threadgroup float *const')
- float const x = (p_ptr + w_ptr);
- ~~~~~ ^ ~~~~~
-program_source:14:9: error: cannot assign to variable 'p_ptr' with const-qualified type 'float *const'
- p_ptr = x;
- ~~~~~ ^
-program_source:11:23: note: variable 'p_ptr' declared const here
- thread float* const p_ptr = p;
- ~~~~~~~~~~~~~~~~~~~~^~~~~~~~~
-
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 a1c80e0..44e6f0a 100644
--- a/test/tint/var/initialization/workgroup/matrix.wgsl.expected.ir.msl
+++ b/test/tint/var/initialization/workgroup/matrix.wgsl.expected.ir.msl
@@ -1,26 +1,16 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
+struct tint_symbol_2 {
+ float2x3 tint_symbol_1;
+};
+struct tint_module_vars_struct {
+ threadgroup float2x3* v;
+};
-threadgroup float2x3 v;
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v_1) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.v=(&(*v_1).tint_symbol_1)};
if ((tint_local_index == 0u)) {
- v = float2x3(float3(0.0f), float3(0.0f));
+ (*tint_module_vars.v) = float2x3(float3(0.0f), float3(0.0f));
}
threadgroup_barrier(mem_flags::mem_threadgroup);
}
-program_source:4:22: error: program scope variable must reside in constant address space
-threadgroup float2x3 v;
- ^
-program_source:6:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
- if ((tint_local_index == 0u)) {
- ~~~~~~~~~~~~~~~~~^~~~~
-program_source:6:25: note: remove extraneous parentheses around the comparison to silence this warning
- if ((tint_local_index == 0u)) {
- ~ ^ ~
-program_source:6:25: note: use '=' to turn this equality comparison into an assignment
- if ((tint_local_index == 0u)) {
- ^~
- =
-
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 e1f46e9..97b64c8 100644
--- a/test/tint/var/initialization/workgroup/scalar.wgsl.expected.ir.msl
+++ b/test/tint/var/initialization/workgroup/scalar.wgsl.expected.ir.msl
@@ -1,30 +1,17 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
+struct tint_symbol_2 {
+ int tint_symbol_1;
+};
+struct tint_module_vars_struct {
+ threadgroup int* v;
+};
-threadgroup int v;
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v_1) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.v=(&(*v_1).tint_symbol_1)};
if ((tint_local_index == 0u)) {
- v = 0;
+ (*tint_module_vars.v) = 0;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- int const i = v;
+ int const i = (*tint_module_vars.v);
}
-program_source:4:17: error: program scope variable must reside in constant address space
-threadgroup int v;
- ^
-program_source:6:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
- if ((tint_local_index == 0u)) {
- ~~~~~~~~~~~~~~~~~^~~~~
-program_source:6:25: note: remove extraneous parentheses around the comparison to silence this warning
- if ((tint_local_index == 0u)) {
- ~ ^ ~
-program_source:6:25: note: use '=' to turn this equality comparison into an assignment
- if ((tint_local_index == 0u)) {
- ^~
- =
-program_source:10:13: warning: unused variable 'i' [-Wunused-variable]
- int const i = v;
- ^
-
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 6655af9..e974d43 100644
--- a/test/tint/var/initialization/workgroup/struct.wgsl.expected.ir.msl
+++ b/test/tint/var/initialization/workgroup/struct.wgsl.expected.ir.msl
@@ -1,30 +1,20 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
struct S {
int a;
float b;
};
+struct tint_symbol_2 {
+ S tint_symbol_1;
+};
+struct tint_module_vars_struct {
+ threadgroup S* v;
+};
-threadgroup S v;
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v_1) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.v=(&(*v_1).tint_symbol_1)};
if ((tint_local_index == 0u)) {
- v = S{};
+ (*tint_module_vars.v) = S{};
}
threadgroup_barrier(mem_flags::mem_threadgroup);
}
-program_source:8:15: error: program scope variable must reside in constant address space
-threadgroup S v;
- ^
-program_source:10:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
- if ((tint_local_index == 0u)) {
- ~~~~~~~~~~~~~~~~~^~~~~
-program_source:10:25: note: remove extraneous parentheses around the comparison to silence this warning
- if ((tint_local_index == 0u)) {
- ~ ^ ~
-program_source:10:25: note: use '=' to turn this equality comparison into an assignment
- if ((tint_local_index == 0u)) {
- ^~
- =
-
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 a8308b3..d24feeb 100644
--- a/test/tint/var/initialization/workgroup/vector.wgsl.expected.ir.msl
+++ b/test/tint/var/initialization/workgroup/vector.wgsl.expected.ir.msl
@@ -1,26 +1,16 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
+struct tint_symbol_2 {
+ int3 tint_symbol_1;
+};
+struct tint_module_vars_struct {
+ threadgroup int3* v;
+};
-threadgroup int3 v;
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v_1) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.v=(&(*v_1).tint_symbol_1)};
if ((tint_local_index == 0u)) {
- v = int3(0);
+ (*tint_module_vars.v) = int3(0);
}
threadgroup_barrier(mem_flags::mem_threadgroup);
}
-program_source:4:18: error: program scope variable must reside in constant address space
-threadgroup int3 v;
- ^
-program_source:6:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
- if ((tint_local_index == 0u)) {
- ~~~~~~~~~~~~~~~~~^~~~~
-program_source:6:25: note: remove extraneous parentheses around the comparison to silence this warning
- if ((tint_local_index == 0u)) {
- ~ ^ ~
-program_source:6:25: note: use '=' to turn this equality comparison into an assignment
- if ((tint_local_index == 0u)) {
- ^~
- =
-
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 81f00cc..0a83ab7 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
@@ -1,621 +1,413 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
+struct tint_symbol_101 {
+ float2x2 tint_symbol_1;
+ float2x2 tint_symbol_2;
+ float2x2 tint_symbol_3;
+ float2x2 tint_symbol_4;
+ float2x2 tint_symbol_5;
+ float2x2 tint_symbol_6;
+ float2x2 tint_symbol_7;
+ float2x2 tint_symbol_8;
+ float2x2 tint_symbol_9;
+ float2x2 tint_symbol_10;
+ float2x2 tint_symbol_11;
+ float2x2 tint_symbol_12;
+ float2x2 tint_symbol_13;
+ float2x2 tint_symbol_14;
+ float2x2 tint_symbol_15;
+ float2x2 tint_symbol_16;
+ float2x2 tint_symbol_17;
+ float2x2 tint_symbol_18;
+ float2x2 tint_symbol_19;
+ float2x2 tint_symbol_20;
+ float2x2 tint_symbol_21;
+ float2x2 tint_symbol_22;
+ float2x2 tint_symbol_23;
+ float2x2 tint_symbol_24;
+ float2x2 tint_symbol_25;
+ float2x2 tint_symbol_26;
+ float2x2 tint_symbol_27;
+ float2x2 tint_symbol_28;
+ float2x2 tint_symbol_29;
+ float2x2 tint_symbol_30;
+ float2x2 tint_symbol_31;
+ float2x2 tint_symbol_32;
+ float2x2 tint_symbol_33;
+ float2x2 tint_symbol_34;
+ float2x2 tint_symbol_35;
+ float2x2 tint_symbol_36;
+ float2x2 tint_symbol_37;
+ float2x2 tint_symbol_38;
+ float2x2 tint_symbol_39;
+ float2x2 tint_symbol_40;
+ float2x2 tint_symbol_41;
+ float2x2 tint_symbol_42;
+ float2x2 tint_symbol_43;
+ float2x2 tint_symbol_44;
+ float2x2 tint_symbol_45;
+ float2x2 tint_symbol_46;
+ float2x2 tint_symbol_47;
+ float2x2 tint_symbol_48;
+ float2x2 tint_symbol_49;
+ float2x2 tint_symbol_50;
+ float2x2 tint_symbol_51;
+ float2x2 tint_symbol_52;
+ float2x2 tint_symbol_53;
+ float2x2 tint_symbol_54;
+ float2x2 tint_symbol_55;
+ float2x2 tint_symbol_56;
+ float2x2 tint_symbol_57;
+ float2x2 tint_symbol_58;
+ float2x2 tint_symbol_59;
+ float2x2 tint_symbol_60;
+ float2x2 tint_symbol_61;
+ float2x2 tint_symbol_62;
+ float2x2 tint_symbol_63;
+ float2x2 tint_symbol_64;
+ float2x2 tint_symbol_65;
+ float2x2 tint_symbol_66;
+ float2x2 tint_symbol_67;
+ float2x2 tint_symbol_68;
+ float2x2 tint_symbol_69;
+ float2x2 tint_symbol_70;
+ float2x2 tint_symbol_71;
+ float2x2 tint_symbol_72;
+ float2x2 tint_symbol_73;
+ float2x2 tint_symbol_74;
+ float2x2 tint_symbol_75;
+ float2x2 tint_symbol_76;
+ float2x2 tint_symbol_77;
+ float2x2 tint_symbol_78;
+ float2x2 tint_symbol_79;
+ float2x2 tint_symbol_80;
+ float2x2 tint_symbol_81;
+ float2x2 tint_symbol_82;
+ float2x2 tint_symbol_83;
+ float2x2 tint_symbol_84;
+ float2x2 tint_symbol_85;
+ float2x2 tint_symbol_86;
+ float2x2 tint_symbol_87;
+ float2x2 tint_symbol_88;
+ float2x2 tint_symbol_89;
+ float2x2 tint_symbol_90;
+ float2x2 tint_symbol_91;
+ float2x2 tint_symbol_92;
+ float2x2 tint_symbol_93;
+ float2x2 tint_symbol_94;
+ float2x2 tint_symbol_95;
+ float2x2 tint_symbol_96;
+ float2x2 tint_symbol_97;
+ float2x2 tint_symbol_98;
+ float2x2 tint_symbol_99;
+ float2x2 tint_symbol_100;
+};
+struct tint_module_vars_struct {
+ threadgroup float2x2* m00;
+ threadgroup float2x2* m01;
+ threadgroup float2x2* m02;
+ threadgroup float2x2* m03;
+ threadgroup float2x2* m04;
+ threadgroup float2x2* m05;
+ threadgroup float2x2* m06;
+ threadgroup float2x2* m07;
+ threadgroup float2x2* m08;
+ threadgroup float2x2* m09;
+ threadgroup float2x2* m10;
+ threadgroup float2x2* m11;
+ threadgroup float2x2* m12;
+ threadgroup float2x2* m13;
+ threadgroup float2x2* m14;
+ threadgroup float2x2* m15;
+ threadgroup float2x2* m16;
+ threadgroup float2x2* m17;
+ threadgroup float2x2* m18;
+ threadgroup float2x2* m19;
+ threadgroup float2x2* m20;
+ threadgroup float2x2* m21;
+ threadgroup float2x2* m22;
+ threadgroup float2x2* m23;
+ threadgroup float2x2* m24;
+ threadgroup float2x2* m25;
+ threadgroup float2x2* m26;
+ threadgroup float2x2* m27;
+ threadgroup float2x2* m28;
+ threadgroup float2x2* m29;
+ threadgroup float2x2* m30;
+ threadgroup float2x2* m31;
+ threadgroup float2x2* m32;
+ threadgroup float2x2* m33;
+ threadgroup float2x2* m34;
+ threadgroup float2x2* m35;
+ threadgroup float2x2* m36;
+ threadgroup float2x2* m37;
+ threadgroup float2x2* m38;
+ threadgroup float2x2* m39;
+ threadgroup float2x2* m40;
+ threadgroup float2x2* m41;
+ threadgroup float2x2* m42;
+ threadgroup float2x2* m43;
+ threadgroup float2x2* m44;
+ threadgroup float2x2* m45;
+ threadgroup float2x2* m46;
+ threadgroup float2x2* m47;
+ threadgroup float2x2* m48;
+ threadgroup float2x2* m49;
+ threadgroup float2x2* m50;
+ threadgroup float2x2* m51;
+ threadgroup float2x2* m52;
+ threadgroup float2x2* m53;
+ threadgroup float2x2* m54;
+ threadgroup float2x2* m55;
+ threadgroup float2x2* m56;
+ threadgroup float2x2* m57;
+ threadgroup float2x2* m58;
+ threadgroup float2x2* m59;
+ threadgroup float2x2* m60;
+ threadgroup float2x2* m61;
+ threadgroup float2x2* m62;
+ threadgroup float2x2* m63;
+ threadgroup float2x2* m64;
+ threadgroup float2x2* m65;
+ threadgroup float2x2* m66;
+ threadgroup float2x2* m67;
+ threadgroup float2x2* m68;
+ threadgroup float2x2* m69;
+ threadgroup float2x2* m70;
+ threadgroup float2x2* m71;
+ threadgroup float2x2* m72;
+ threadgroup float2x2* m73;
+ threadgroup float2x2* m74;
+ threadgroup float2x2* m75;
+ threadgroup float2x2* m76;
+ threadgroup float2x2* m77;
+ threadgroup float2x2* m78;
+ threadgroup float2x2* m79;
+ threadgroup float2x2* m80;
+ threadgroup float2x2* m81;
+ threadgroup float2x2* m82;
+ threadgroup float2x2* m83;
+ threadgroup float2x2* m84;
+ threadgroup float2x2* m85;
+ threadgroup float2x2* m86;
+ threadgroup float2x2* m87;
+ threadgroup float2x2* m88;
+ threadgroup float2x2* m89;
+ threadgroup float2x2* m90;
+ threadgroup float2x2* m91;
+ threadgroup float2x2* m92;
+ threadgroup float2x2* m93;
+ threadgroup float2x2* m94;
+ threadgroup float2x2* m95;
+ threadgroup float2x2* m96;
+ threadgroup float2x2* m97;
+ threadgroup float2x2* m98;
+ threadgroup float2x2* m99;
+};
-threadgroup float2x2 m00;
-threadgroup float2x2 m01;
-threadgroup float2x2 m02;
-threadgroup float2x2 m03;
-threadgroup float2x2 m04;
-threadgroup float2x2 m05;
-threadgroup float2x2 m06;
-threadgroup float2x2 m07;
-threadgroup float2x2 m08;
-threadgroup float2x2 m09;
-threadgroup float2x2 m10;
-threadgroup float2x2 m11;
-threadgroup float2x2 m12;
-threadgroup float2x2 m13;
-threadgroup float2x2 m14;
-threadgroup float2x2 m15;
-threadgroup float2x2 m16;
-threadgroup float2x2 m17;
-threadgroup float2x2 m18;
-threadgroup float2x2 m19;
-threadgroup float2x2 m20;
-threadgroup float2x2 m21;
-threadgroup float2x2 m22;
-threadgroup float2x2 m23;
-threadgroup float2x2 m24;
-threadgroup float2x2 m25;
-threadgroup float2x2 m26;
-threadgroup float2x2 m27;
-threadgroup float2x2 m28;
-threadgroup float2x2 m29;
-threadgroup float2x2 m30;
-threadgroup float2x2 m31;
-threadgroup float2x2 m32;
-threadgroup float2x2 m33;
-threadgroup float2x2 m34;
-threadgroup float2x2 m35;
-threadgroup float2x2 m36;
-threadgroup float2x2 m37;
-threadgroup float2x2 m38;
-threadgroup float2x2 m39;
-threadgroup float2x2 m40;
-threadgroup float2x2 m41;
-threadgroup float2x2 m42;
-threadgroup float2x2 m43;
-threadgroup float2x2 m44;
-threadgroup float2x2 m45;
-threadgroup float2x2 m46;
-threadgroup float2x2 m47;
-threadgroup float2x2 m48;
-threadgroup float2x2 m49;
-threadgroup float2x2 m50;
-threadgroup float2x2 m51;
-threadgroup float2x2 m52;
-threadgroup float2x2 m53;
-threadgroup float2x2 m54;
-threadgroup float2x2 m55;
-threadgroup float2x2 m56;
-threadgroup float2x2 m57;
-threadgroup float2x2 m58;
-threadgroup float2x2 m59;
-threadgroup float2x2 m60;
-threadgroup float2x2 m61;
-threadgroup float2x2 m62;
-threadgroup float2x2 m63;
-threadgroup float2x2 m64;
-threadgroup float2x2 m65;
-threadgroup float2x2 m66;
-threadgroup float2x2 m67;
-threadgroup float2x2 m68;
-threadgroup float2x2 m69;
-threadgroup float2x2 m70;
-threadgroup float2x2 m71;
-threadgroup float2x2 m72;
-threadgroup float2x2 m73;
-threadgroup float2x2 m74;
-threadgroup float2x2 m75;
-threadgroup float2x2 m76;
-threadgroup float2x2 m77;
-threadgroup float2x2 m78;
-threadgroup float2x2 m79;
-threadgroup float2x2 m80;
-threadgroup float2x2 m81;
-threadgroup float2x2 m82;
-threadgroup float2x2 m83;
-threadgroup float2x2 m84;
-threadgroup float2x2 m85;
-threadgroup float2x2 m86;
-threadgroup float2x2 m87;
-threadgroup float2x2 m88;
-threadgroup float2x2 m89;
-threadgroup float2x2 m90;
-threadgroup float2x2 m91;
-threadgroup float2x2 m92;
-threadgroup float2x2 m93;
-threadgroup float2x2 m94;
-threadgroup float2x2 m95;
-threadgroup float2x2 m96;
-threadgroup float2x2 m97;
-threadgroup float2x2 m98;
-threadgroup float2x2 m99;
-kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], threadgroup tint_symbol_101* v) {
+ 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)) {
- m00 = float2x2(float2(0.0f), float2(0.0f));
- m01 = float2x2(float2(0.0f), float2(0.0f));
- m02 = float2x2(float2(0.0f), float2(0.0f));
- m03 = float2x2(float2(0.0f), float2(0.0f));
- m04 = float2x2(float2(0.0f), float2(0.0f));
- m05 = float2x2(float2(0.0f), float2(0.0f));
- m06 = float2x2(float2(0.0f), float2(0.0f));
- m07 = float2x2(float2(0.0f), float2(0.0f));
- m08 = float2x2(float2(0.0f), float2(0.0f));
- m09 = float2x2(float2(0.0f), float2(0.0f));
- m10 = float2x2(float2(0.0f), float2(0.0f));
- m11 = float2x2(float2(0.0f), float2(0.0f));
- m12 = float2x2(float2(0.0f), float2(0.0f));
- m13 = float2x2(float2(0.0f), float2(0.0f));
- m14 = float2x2(float2(0.0f), float2(0.0f));
- m15 = float2x2(float2(0.0f), float2(0.0f));
- m16 = float2x2(float2(0.0f), float2(0.0f));
- m17 = float2x2(float2(0.0f), float2(0.0f));
- m18 = float2x2(float2(0.0f), float2(0.0f));
- m19 = float2x2(float2(0.0f), float2(0.0f));
- m20 = float2x2(float2(0.0f), float2(0.0f));
- m21 = float2x2(float2(0.0f), float2(0.0f));
- m22 = float2x2(float2(0.0f), float2(0.0f));
- m23 = float2x2(float2(0.0f), float2(0.0f));
- m24 = float2x2(float2(0.0f), float2(0.0f));
- m25 = float2x2(float2(0.0f), float2(0.0f));
- m26 = float2x2(float2(0.0f), float2(0.0f));
- m27 = float2x2(float2(0.0f), float2(0.0f));
- m28 = float2x2(float2(0.0f), float2(0.0f));
- m29 = float2x2(float2(0.0f), float2(0.0f));
- m30 = float2x2(float2(0.0f), float2(0.0f));
- m31 = float2x2(float2(0.0f), float2(0.0f));
- m32 = float2x2(float2(0.0f), float2(0.0f));
- m33 = float2x2(float2(0.0f), float2(0.0f));
- m34 = float2x2(float2(0.0f), float2(0.0f));
- m35 = float2x2(float2(0.0f), float2(0.0f));
- m36 = float2x2(float2(0.0f), float2(0.0f));
- m37 = float2x2(float2(0.0f), float2(0.0f));
- m38 = float2x2(float2(0.0f), float2(0.0f));
- m39 = float2x2(float2(0.0f), float2(0.0f));
- m40 = float2x2(float2(0.0f), float2(0.0f));
- m41 = float2x2(float2(0.0f), float2(0.0f));
- m42 = float2x2(float2(0.0f), float2(0.0f));
- m43 = float2x2(float2(0.0f), float2(0.0f));
- m44 = float2x2(float2(0.0f), float2(0.0f));
- m45 = float2x2(float2(0.0f), float2(0.0f));
- m46 = float2x2(float2(0.0f), float2(0.0f));
- m47 = float2x2(float2(0.0f), float2(0.0f));
- m48 = float2x2(float2(0.0f), float2(0.0f));
- m49 = float2x2(float2(0.0f), float2(0.0f));
- m50 = float2x2(float2(0.0f), float2(0.0f));
- m51 = float2x2(float2(0.0f), float2(0.0f));
- m52 = float2x2(float2(0.0f), float2(0.0f));
- m53 = float2x2(float2(0.0f), float2(0.0f));
- m54 = float2x2(float2(0.0f), float2(0.0f));
- m55 = float2x2(float2(0.0f), float2(0.0f));
- m56 = float2x2(float2(0.0f), float2(0.0f));
- m57 = float2x2(float2(0.0f), float2(0.0f));
- m58 = float2x2(float2(0.0f), float2(0.0f));
- m59 = float2x2(float2(0.0f), float2(0.0f));
- m60 = float2x2(float2(0.0f), float2(0.0f));
- m61 = float2x2(float2(0.0f), float2(0.0f));
- m62 = float2x2(float2(0.0f), float2(0.0f));
- m63 = float2x2(float2(0.0f), float2(0.0f));
- m64 = float2x2(float2(0.0f), float2(0.0f));
- m65 = float2x2(float2(0.0f), float2(0.0f));
- m66 = float2x2(float2(0.0f), float2(0.0f));
- m67 = float2x2(float2(0.0f), float2(0.0f));
- m68 = float2x2(float2(0.0f), float2(0.0f));
- m69 = float2x2(float2(0.0f), float2(0.0f));
- m70 = float2x2(float2(0.0f), float2(0.0f));
- m71 = float2x2(float2(0.0f), float2(0.0f));
- m72 = float2x2(float2(0.0f), float2(0.0f));
- m73 = float2x2(float2(0.0f), float2(0.0f));
- m74 = float2x2(float2(0.0f), float2(0.0f));
- m75 = float2x2(float2(0.0f), float2(0.0f));
- m76 = float2x2(float2(0.0f), float2(0.0f));
- m77 = float2x2(float2(0.0f), float2(0.0f));
- m78 = float2x2(float2(0.0f), float2(0.0f));
- m79 = float2x2(float2(0.0f), float2(0.0f));
- m80 = float2x2(float2(0.0f), float2(0.0f));
- m81 = float2x2(float2(0.0f), float2(0.0f));
- m82 = float2x2(float2(0.0f), float2(0.0f));
- m83 = float2x2(float2(0.0f), float2(0.0f));
- m84 = float2x2(float2(0.0f), float2(0.0f));
- m85 = float2x2(float2(0.0f), float2(0.0f));
- m86 = float2x2(float2(0.0f), float2(0.0f));
- m87 = float2x2(float2(0.0f), float2(0.0f));
- m88 = float2x2(float2(0.0f), float2(0.0f));
- m89 = float2x2(float2(0.0f), float2(0.0f));
- m90 = float2x2(float2(0.0f), float2(0.0f));
- m91 = float2x2(float2(0.0f), float2(0.0f));
- m92 = float2x2(float2(0.0f), float2(0.0f));
- m93 = float2x2(float2(0.0f), float2(0.0f));
- m94 = float2x2(float2(0.0f), float2(0.0f));
- m95 = float2x2(float2(0.0f), float2(0.0f));
- m96 = float2x2(float2(0.0f), float2(0.0f));
- m97 = float2x2(float2(0.0f), float2(0.0f));
- m98 = float2x2(float2(0.0f), float2(0.0f));
- m99 = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m00) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m01) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m02) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m03) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m04) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m05) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m06) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m07) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m08) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m09) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m10) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m11) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m12) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m13) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m14) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m15) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m16) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m17) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m18) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m19) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m20) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m21) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m22) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m23) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m24) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m25) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m26) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m27) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m28) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m29) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m30) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m31) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m32) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m33) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m34) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m35) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m36) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m37) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m38) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m39) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m40) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m41) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m42) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m43) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m44) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m45) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m46) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m47) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m48) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m49) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m50) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m51) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m52) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m53) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m54) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m55) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m56) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m57) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m58) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m59) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m60) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m61) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m62) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m63) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m64) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m65) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m66) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m67) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m68) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m69) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m70) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m71) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m72) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m73) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m74) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m75) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m76) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m77) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m78) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m79) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m80) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m81) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m82) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m83) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m84) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m85) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m86) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m87) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m88) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m89) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m90) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m91) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m92) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m93) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m94) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m95) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m96) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m97) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m98) = float2x2(float2(0.0f), float2(0.0f));
+ (*tint_module_vars.m99) = float2x2(float2(0.0f), float2(0.0f));
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- m00[0][0] = 1.0f;
- m01[0][0] = 1.0f;
- m02[0][0] = 1.0f;
- m03[0][0] = 1.0f;
- m04[0][0] = 1.0f;
- m05[0][0] = 1.0f;
- m06[0][0] = 1.0f;
- m07[0][0] = 1.0f;
- m08[0][0] = 1.0f;
- m09[0][0] = 1.0f;
- m10[0][0] = 1.0f;
- m11[0][0] = 1.0f;
- m12[0][0] = 1.0f;
- m13[0][0] = 1.0f;
- m14[0][0] = 1.0f;
- m15[0][0] = 1.0f;
- m16[0][0] = 1.0f;
- m17[0][0] = 1.0f;
- m18[0][0] = 1.0f;
- m19[0][0] = 1.0f;
- m20[0][0] = 1.0f;
- m21[0][0] = 1.0f;
- m22[0][0] = 1.0f;
- m23[0][0] = 1.0f;
- m24[0][0] = 1.0f;
- m25[0][0] = 1.0f;
- m26[0][0] = 1.0f;
- m27[0][0] = 1.0f;
- m28[0][0] = 1.0f;
- m29[0][0] = 1.0f;
- m30[0][0] = 1.0f;
- m31[0][0] = 1.0f;
- m32[0][0] = 1.0f;
- m33[0][0] = 1.0f;
- m34[0][0] = 1.0f;
- m35[0][0] = 1.0f;
- m36[0][0] = 1.0f;
- m37[0][0] = 1.0f;
- m38[0][0] = 1.0f;
- m39[0][0] = 1.0f;
- m40[0][0] = 1.0f;
- m41[0][0] = 1.0f;
- m42[0][0] = 1.0f;
- m43[0][0] = 1.0f;
- m44[0][0] = 1.0f;
- m45[0][0] = 1.0f;
- m46[0][0] = 1.0f;
- m47[0][0] = 1.0f;
- m48[0][0] = 1.0f;
- m49[0][0] = 1.0f;
- m50[0][0] = 1.0f;
- m51[0][0] = 1.0f;
- m52[0][0] = 1.0f;
- m53[0][0] = 1.0f;
- m54[0][0] = 1.0f;
- m55[0][0] = 1.0f;
- m56[0][0] = 1.0f;
- m57[0][0] = 1.0f;
- m58[0][0] = 1.0f;
- m59[0][0] = 1.0f;
- m60[0][0] = 1.0f;
- m61[0][0] = 1.0f;
- m62[0][0] = 1.0f;
- m63[0][0] = 1.0f;
- m64[0][0] = 1.0f;
- m65[0][0] = 1.0f;
- m66[0][0] = 1.0f;
- m67[0][0] = 1.0f;
- m68[0][0] = 1.0f;
- m69[0][0] = 1.0f;
- m70[0][0] = 1.0f;
- m71[0][0] = 1.0f;
- m72[0][0] = 1.0f;
- m73[0][0] = 1.0f;
- m74[0][0] = 1.0f;
- m75[0][0] = 1.0f;
- m76[0][0] = 1.0f;
- m77[0][0] = 1.0f;
- m78[0][0] = 1.0f;
- m79[0][0] = 1.0f;
- m80[0][0] = 1.0f;
- m81[0][0] = 1.0f;
- m82[0][0] = 1.0f;
- m83[0][0] = 1.0f;
- m84[0][0] = 1.0f;
- m85[0][0] = 1.0f;
- m86[0][0] = 1.0f;
- m87[0][0] = 1.0f;
- m88[0][0] = 1.0f;
- m89[0][0] = 1.0f;
- m90[0][0] = 1.0f;
- m91[0][0] = 1.0f;
- m92[0][0] = 1.0f;
- m93[0][0] = 1.0f;
- m94[0][0] = 1.0f;
- m95[0][0] = 1.0f;
- m96[0][0] = 1.0f;
- m97[0][0] = 1.0f;
- m98[0][0] = 1.0f;
- m99[0][0] = 1.0f;
+ (*tint_module_vars.m00)[0][0] = 1.0f;
+ (*tint_module_vars.m01)[0][0] = 1.0f;
+ (*tint_module_vars.m02)[0][0] = 1.0f;
+ (*tint_module_vars.m03)[0][0] = 1.0f;
+ (*tint_module_vars.m04)[0][0] = 1.0f;
+ (*tint_module_vars.m05)[0][0] = 1.0f;
+ (*tint_module_vars.m06)[0][0] = 1.0f;
+ (*tint_module_vars.m07)[0][0] = 1.0f;
+ (*tint_module_vars.m08)[0][0] = 1.0f;
+ (*tint_module_vars.m09)[0][0] = 1.0f;
+ (*tint_module_vars.m10)[0][0] = 1.0f;
+ (*tint_module_vars.m11)[0][0] = 1.0f;
+ (*tint_module_vars.m12)[0][0] = 1.0f;
+ (*tint_module_vars.m13)[0][0] = 1.0f;
+ (*tint_module_vars.m14)[0][0] = 1.0f;
+ (*tint_module_vars.m15)[0][0] = 1.0f;
+ (*tint_module_vars.m16)[0][0] = 1.0f;
+ (*tint_module_vars.m17)[0][0] = 1.0f;
+ (*tint_module_vars.m18)[0][0] = 1.0f;
+ (*tint_module_vars.m19)[0][0] = 1.0f;
+ (*tint_module_vars.m20)[0][0] = 1.0f;
+ (*tint_module_vars.m21)[0][0] = 1.0f;
+ (*tint_module_vars.m22)[0][0] = 1.0f;
+ (*tint_module_vars.m23)[0][0] = 1.0f;
+ (*tint_module_vars.m24)[0][0] = 1.0f;
+ (*tint_module_vars.m25)[0][0] = 1.0f;
+ (*tint_module_vars.m26)[0][0] = 1.0f;
+ (*tint_module_vars.m27)[0][0] = 1.0f;
+ (*tint_module_vars.m28)[0][0] = 1.0f;
+ (*tint_module_vars.m29)[0][0] = 1.0f;
+ (*tint_module_vars.m30)[0][0] = 1.0f;
+ (*tint_module_vars.m31)[0][0] = 1.0f;
+ (*tint_module_vars.m32)[0][0] = 1.0f;
+ (*tint_module_vars.m33)[0][0] = 1.0f;
+ (*tint_module_vars.m34)[0][0] = 1.0f;
+ (*tint_module_vars.m35)[0][0] = 1.0f;
+ (*tint_module_vars.m36)[0][0] = 1.0f;
+ (*tint_module_vars.m37)[0][0] = 1.0f;
+ (*tint_module_vars.m38)[0][0] = 1.0f;
+ (*tint_module_vars.m39)[0][0] = 1.0f;
+ (*tint_module_vars.m40)[0][0] = 1.0f;
+ (*tint_module_vars.m41)[0][0] = 1.0f;
+ (*tint_module_vars.m42)[0][0] = 1.0f;
+ (*tint_module_vars.m43)[0][0] = 1.0f;
+ (*tint_module_vars.m44)[0][0] = 1.0f;
+ (*tint_module_vars.m45)[0][0] = 1.0f;
+ (*tint_module_vars.m46)[0][0] = 1.0f;
+ (*tint_module_vars.m47)[0][0] = 1.0f;
+ (*tint_module_vars.m48)[0][0] = 1.0f;
+ (*tint_module_vars.m49)[0][0] = 1.0f;
+ (*tint_module_vars.m50)[0][0] = 1.0f;
+ (*tint_module_vars.m51)[0][0] = 1.0f;
+ (*tint_module_vars.m52)[0][0] = 1.0f;
+ (*tint_module_vars.m53)[0][0] = 1.0f;
+ (*tint_module_vars.m54)[0][0] = 1.0f;
+ (*tint_module_vars.m55)[0][0] = 1.0f;
+ (*tint_module_vars.m56)[0][0] = 1.0f;
+ (*tint_module_vars.m57)[0][0] = 1.0f;
+ (*tint_module_vars.m58)[0][0] = 1.0f;
+ (*tint_module_vars.m59)[0][0] = 1.0f;
+ (*tint_module_vars.m60)[0][0] = 1.0f;
+ (*tint_module_vars.m61)[0][0] = 1.0f;
+ (*tint_module_vars.m62)[0][0] = 1.0f;
+ (*tint_module_vars.m63)[0][0] = 1.0f;
+ (*tint_module_vars.m64)[0][0] = 1.0f;
+ (*tint_module_vars.m65)[0][0] = 1.0f;
+ (*tint_module_vars.m66)[0][0] = 1.0f;
+ (*tint_module_vars.m67)[0][0] = 1.0f;
+ (*tint_module_vars.m68)[0][0] = 1.0f;
+ (*tint_module_vars.m69)[0][0] = 1.0f;
+ (*tint_module_vars.m70)[0][0] = 1.0f;
+ (*tint_module_vars.m71)[0][0] = 1.0f;
+ (*tint_module_vars.m72)[0][0] = 1.0f;
+ (*tint_module_vars.m73)[0][0] = 1.0f;
+ (*tint_module_vars.m74)[0][0] = 1.0f;
+ (*tint_module_vars.m75)[0][0] = 1.0f;
+ (*tint_module_vars.m76)[0][0] = 1.0f;
+ (*tint_module_vars.m77)[0][0] = 1.0f;
+ (*tint_module_vars.m78)[0][0] = 1.0f;
+ (*tint_module_vars.m79)[0][0] = 1.0f;
+ (*tint_module_vars.m80)[0][0] = 1.0f;
+ (*tint_module_vars.m81)[0][0] = 1.0f;
+ (*tint_module_vars.m82)[0][0] = 1.0f;
+ (*tint_module_vars.m83)[0][0] = 1.0f;
+ (*tint_module_vars.m84)[0][0] = 1.0f;
+ (*tint_module_vars.m85)[0][0] = 1.0f;
+ (*tint_module_vars.m86)[0][0] = 1.0f;
+ (*tint_module_vars.m87)[0][0] = 1.0f;
+ (*tint_module_vars.m88)[0][0] = 1.0f;
+ (*tint_module_vars.m89)[0][0] = 1.0f;
+ (*tint_module_vars.m90)[0][0] = 1.0f;
+ (*tint_module_vars.m91)[0][0] = 1.0f;
+ (*tint_module_vars.m92)[0][0] = 1.0f;
+ (*tint_module_vars.m93)[0][0] = 1.0f;
+ (*tint_module_vars.m94)[0][0] = 1.0f;
+ (*tint_module_vars.m95)[0][0] = 1.0f;
+ (*tint_module_vars.m96)[0][0] = 1.0f;
+ (*tint_module_vars.m97)[0][0] = 1.0f;
+ (*tint_module_vars.m98)[0][0] = 1.0f;
+ (*tint_module_vars.m99)[0][0] = 1.0f;
}
-program_source:4:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m00;
- ^
-program_source:5:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m01;
- ^
-program_source:6:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m02;
- ^
-program_source:7:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m03;
- ^
-program_source:8:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m04;
- ^
-program_source:9:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m05;
- ^
-program_source:10:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m06;
- ^
-program_source:11:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m07;
- ^
-program_source:12:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m08;
- ^
-program_source:13:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m09;
- ^
-program_source:14:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m10;
- ^
-program_source:15:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m11;
- ^
-program_source:16:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m12;
- ^
-program_source:17:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m13;
- ^
-program_source:18:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m14;
- ^
-program_source:19:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m15;
- ^
-program_source:20:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m16;
- ^
-program_source:21:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m17;
- ^
-program_source:22:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m18;
- ^
-program_source:23:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m19;
- ^
-program_source:24:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m20;
- ^
-program_source:25:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m21;
- ^
-program_source:26:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m22;
- ^
-program_source:27:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m23;
- ^
-program_source:28:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m24;
- ^
-program_source:29:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m25;
- ^
-program_source:30:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m26;
- ^
-program_source:31:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m27;
- ^
-program_source:32:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m28;
- ^
-program_source:33:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m29;
- ^
-program_source:34:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m30;
- ^
-program_source:35:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m31;
- ^
-program_source:36:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m32;
- ^
-program_source:37:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m33;
- ^
-program_source:38:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m34;
- ^
-program_source:39:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m35;
- ^
-program_source:40:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m36;
- ^
-program_source:41:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m37;
- ^
-program_source:42:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m38;
- ^
-program_source:43:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m39;
- ^
-program_source:44:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m40;
- ^
-program_source:45:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m41;
- ^
-program_source:46:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m42;
- ^
-program_source:47:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m43;
- ^
-program_source:48:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m44;
- ^
-program_source:49:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m45;
- ^
-program_source:50:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m46;
- ^
-program_source:51:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m47;
- ^
-program_source:52:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m48;
- ^
-program_source:53:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m49;
- ^
-program_source:54:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m50;
- ^
-program_source:55:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m51;
- ^
-program_source:56:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m52;
- ^
-program_source:57:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m53;
- ^
-program_source:58:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m54;
- ^
-program_source:59:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m55;
- ^
-program_source:60:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m56;
- ^
-program_source:61:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m57;
- ^
-program_source:62:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m58;
- ^
-program_source:63:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m59;
- ^
-program_source:64:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m60;
- ^
-program_source:65:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m61;
- ^
-program_source:66:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m62;
- ^
-program_source:67:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m63;
- ^
-program_source:68:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m64;
- ^
-program_source:69:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m65;
- ^
-program_source:70:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m66;
- ^
-program_source:71:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m67;
- ^
-program_source:72:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m68;
- ^
-program_source:73:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m69;
- ^
-program_source:74:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m70;
- ^
-program_source:75:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m71;
- ^
-program_source:76:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m72;
- ^
-program_source:77:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m73;
- ^
-program_source:78:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m74;
- ^
-program_source:79:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m75;
- ^
-program_source:80:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m76;
- ^
-program_source:81:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m77;
- ^
-program_source:82:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m78;
- ^
-program_source:83:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m79;
- ^
-program_source:84:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m80;
- ^
-program_source:85:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m81;
- ^
-program_source:86:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m82;
- ^
-program_source:87:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m83;
- ^
-program_source:88:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m84;
- ^
-program_source:89:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m85;
- ^
-program_source:90:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m86;
- ^
-program_source:91:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m87;
- ^
-program_source:92:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m88;
- ^
-program_source:93:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m89;
- ^
-program_source:94:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m90;
- ^
-program_source:95:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m91;
- ^
-program_source:96:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m92;
- ^
-program_source:97:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m93;
- ^
-program_source:98:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m94;
- ^
-program_source:99:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m95;
- ^
-program_source:100:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m96;
- ^
-program_source:101:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m97;
- ^
-program_source:102:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m98;
- ^
-program_source:103:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m99;
- ^
-program_source:105:12: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
- if ((idx == 0u)) {
- ~~~~^~~~~
-program_source:105:12: note: remove extraneous parentheses around the comparison to silence this warning
- if ((idx == 0u)) {
- ~ ^ ~
-program_source:105:12: note: use '=' to turn this equality comparison into an assignment
- if ((idx == 0u)) {
- ^~
- =
-
