[msl] Add subgroupMatrixScalarMultiply
Add support for subgroupMatrixScalarMultiply by converting it into a
multiply of the matrix and an diagonal matrix created from the scalar
value
Fixed: 457426179
Change-Id: If711a92dedec8fbbcb65ed812dbd5c2064d22f40
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/271314
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
index 668cb4f..dcb1fd59 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
@@ -119,6 +119,7 @@
case core::BuiltinFn::kSubgroupMatrixMultiplyAccumulate:
case core::BuiltinFn::kSubgroupMatrixScalarAdd:
case core::BuiltinFn::kSubgroupMatrixScalarSubtract:
+ case core::BuiltinFn::kSubgroupMatrixScalarMultiply:
case core::BuiltinFn::kTextureDimensions:
case core::BuiltinFn::kTextureGather:
case core::BuiltinFn::kTextureGatherCompare:
@@ -315,6 +316,9 @@
case core::BuiltinFn::kSubgroupMatrixScalarSubtract:
SubgroupMatrixScalarSubtract(builtin);
break;
+ case core::BuiltinFn::kSubgroupMatrixScalarMultiply:
+ SubgroupMatrixScalarMultiply(builtin);
+ break;
default:
break;
@@ -1228,6 +1232,46 @@
});
builtin->Destroy();
}
+
+ /// Replace a subgroupMatrixScalarMultiply builtin.
+ /// @param builtin the builtin call instruction
+ void SubgroupMatrixScalarMultiply(core::ir::CoreBuiltinCall* builtin) {
+ b.InsertBefore(builtin, [&] {
+ auto* mat = builtin->Args()[0];
+ auto* scalar = builtin->Args()[1];
+
+ auto* sm_ty = mat->Type()->As<core::type::SubgroupMatrix>();
+ TINT_ASSERT(sm_ty);
+
+ auto* left_ty = ty.subgroup_matrix(core::SubgroupMatrixKind::kLeft, sm_ty->Type(),
+ sm_ty->Columns(), sm_ty->Rows());
+ auto* right_ty = ty.subgroup_matrix(core::SubgroupMatrixKind::kRight, sm_ty->Type(),
+ sm_ty->Columns(), sm_ty->Rows());
+ auto* result_ty = ty.subgroup_matrix(core::SubgroupMatrixKind::kResult, sm_ty->Type(),
+ sm_ty->Columns(), sm_ty->Rows());
+
+ // Declare a local variable to receive the result.
+ auto* tmp = b.Var(ty.ptr<function>(result_ty));
+
+ auto* val = b.CallExplicit<msl::ir::BuiltinCall>(left_ty, msl::BuiltinFn::kConvert,
+ Vector{left_ty}, mat);
+
+ auto* mul = b.CallExplicit<msl::ir::BuiltinCall>(
+ right_ty, msl::BuiltinFn::kMakeDiagonalSimdgroupMatrix, Vector{right_ty}, scalar);
+
+ // Note: We need to use a `load` instruction to pass the variable, as the intrinsic
+ // definition expects a value type (as we do not have reference types in the IR). The
+ // printer will just fold away the load, which achieves the pass-by-reference semantics
+ // that we want.
+ b.Call<msl::ir::BuiltinCall>(ty.void_(), msl::BuiltinFn::kSimdgroupMultiply,
+ b.Load(tmp->Result()), val, mul);
+ auto* ld = b.Load(tmp);
+
+ b.CallExplicitWithResult<msl::ir::BuiltinCall>(
+ builtin->DetachResult(), msl::BuiltinFn::kConvert, Vector{sm_ty}, ld);
+ });
+ builtin->Destroy();
+ }
};
} // namespace
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixScalarMultiply/0043a7.wgsl.expected.msl b/test/tint/builtins/gen/literal/subgroupMatrixScalarMultiply/0043a7.wgsl.expected.msl
index ee7074a..7f6f242 100644
--- a/test/tint/builtins/gen/literal/subgroupMatrixScalarMultiply/0043a7.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/subgroupMatrixScalarMultiply/0043a7.wgsl.expected.msl
@@ -1,12 +1,31 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-<dawn>/src/tint/lang/msl/writer/printer/printer.cc:1222 internal compiler error: TINT_UNREACHABLE unhandled: subgroupMatrixScalarMultiply
+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];
+};
-********************************************************************
-* 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. *
-********************************************************************
+struct tint_module_vars_struct {
+ device tint_array<float, 1024>* prevent_dce;
+};
-tint executable returned error: signal: trace/BPT trap
+simdgroup_float8x8 subgroupMatrixScalarMultiply_0043a7() {
+ simdgroup_float8x8 v = make_filled_simdgroup_matrix<float, 8, 8>(0.0f);
+ simdgroup_float8x8 const v_1 = simdgroup_float8x8();
+ simdgroup_multiply(v, v_1, simdgroup_float8x8(8.0f));
+ simdgroup_float8x8 res = v;
+ return res;
+}
+
+kernel void compute_main(device tint_array<float, 1024>* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ simdgroup_store(subgroupMatrixScalarMultiply_0043a7(), (&(*tint_module_vars.prevent_dce)[0u]), ulong(64u), ulong2(0ul), false);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixScalarMultiply/3ec99d.wgsl.expected.msl b/test/tint/builtins/gen/literal/subgroupMatrixScalarMultiply/3ec99d.wgsl.expected.msl
index ee7074a..d9f4e7d 100644
--- a/test/tint/builtins/gen/literal/subgroupMatrixScalarMultiply/3ec99d.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/subgroupMatrixScalarMultiply/3ec99d.wgsl.expected.msl
@@ -1,12 +1,31 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-<dawn>/src/tint/lang/msl/writer/printer/printer.cc:1222 internal compiler error: TINT_UNREACHABLE unhandled: subgroupMatrixScalarMultiply
+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];
+};
-********************************************************************
-* 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. *
-********************************************************************
+struct tint_module_vars_struct {
+ device tint_array<half, 1024>* prevent_dce;
+};
-tint executable returned error: signal: trace/BPT trap
+simdgroup_half8x8 subgroupMatrixScalarMultiply_3ec99d() {
+ simdgroup_half8x8 v = make_filled_simdgroup_matrix<half, 8, 8>(0.0h);
+ simdgroup_half8x8 const v_1 = simdgroup_half8x8();
+ simdgroup_multiply(v, v_1, simdgroup_half8x8(8.0h));
+ simdgroup_half8x8 res = v;
+ return res;
+}
+
+kernel void compute_main(device tint_array<half, 1024>* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ simdgroup_store(subgroupMatrixScalarMultiply_3ec99d(), (&(*tint_module_vars.prevent_dce)[0u]), ulong(64u), ulong2(0ul), false);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixScalarMultiply/4de795.wgsl.expected.msl b/test/tint/builtins/gen/literal/subgroupMatrixScalarMultiply/4de795.wgsl.expected.msl
index ee7074a..7dd2d11 100644
--- a/test/tint/builtins/gen/literal/subgroupMatrixScalarMultiply/4de795.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/subgroupMatrixScalarMultiply/4de795.wgsl.expected.msl
@@ -1,12 +1,31 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-<dawn>/src/tint/lang/msl/writer/printer/printer.cc:1222 internal compiler error: TINT_UNREACHABLE unhandled: subgroupMatrixScalarMultiply
+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];
+};
-********************************************************************
-* 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. *
-********************************************************************
+struct tint_module_vars_struct {
+ device tint_array<half, 1024>* prevent_dce;
+};
-tint executable returned error: signal: trace/BPT trap
+simdgroup_half8x8 subgroupMatrixScalarMultiply_4de795() {
+ simdgroup_half8x8 v = make_filled_simdgroup_matrix<half, 8, 8>(0.0h);
+ simdgroup_half8x8 const v_1 = simdgroup_half8x8();
+ simdgroup_multiply(v, v_1, simdgroup_half8x8(8.0h));
+ simdgroup_half8x8 res = v;
+ return res;
+}
+
+kernel void compute_main(device tint_array<half, 1024>* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ simdgroup_store(subgroupMatrixScalarMultiply_4de795(), (&(*tint_module_vars.prevent_dce)[0u]), ulong(64u), ulong2(0ul), false);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixScalarMultiply/9ec228.wgsl.expected.msl b/test/tint/builtins/gen/literal/subgroupMatrixScalarMultiply/9ec228.wgsl.expected.msl
index ee7074a..4a745c8 100644
--- a/test/tint/builtins/gen/literal/subgroupMatrixScalarMultiply/9ec228.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/subgroupMatrixScalarMultiply/9ec228.wgsl.expected.msl
@@ -1,12 +1,31 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-<dawn>/src/tint/lang/msl/writer/printer/printer.cc:1222 internal compiler error: TINT_UNREACHABLE unhandled: subgroupMatrixScalarMultiply
+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];
+};
-********************************************************************
-* 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. *
-********************************************************************
+struct tint_module_vars_struct {
+ device tint_array<float, 1024>* prevent_dce;
+};
-tint executable returned error: signal: trace/BPT trap
+simdgroup_float8x8 subgroupMatrixScalarMultiply_9ec228() {
+ simdgroup_float8x8 v = make_filled_simdgroup_matrix<float, 8, 8>(0.0f);
+ simdgroup_float8x8 const v_1 = simdgroup_float8x8();
+ simdgroup_multiply(v, v_1, simdgroup_float8x8(8.0f));
+ simdgroup_float8x8 res = v;
+ return res;
+}
+
+kernel void compute_main(device tint_array<float, 1024>* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ simdgroup_store(subgroupMatrixScalarMultiply_9ec228(), (&(*tint_module_vars.prevent_dce)[0u]), ulong(64u), ulong2(0ul), false);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixScalarMultiply/decf0b.wgsl.expected.msl b/test/tint/builtins/gen/literal/subgroupMatrixScalarMultiply/decf0b.wgsl.expected.msl
index ee7074a..36f2324 100644
--- a/test/tint/builtins/gen/literal/subgroupMatrixScalarMultiply/decf0b.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/subgroupMatrixScalarMultiply/decf0b.wgsl.expected.msl
@@ -1,12 +1,31 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-<dawn>/src/tint/lang/msl/writer/printer/printer.cc:1222 internal compiler error: TINT_UNREACHABLE unhandled: subgroupMatrixScalarMultiply
+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];
+};
-********************************************************************
-* 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. *
-********************************************************************
+struct tint_module_vars_struct {
+ device tint_array<float, 1024>* prevent_dce;
+};
-tint executable returned error: signal: trace/BPT trap
+simdgroup_float8x8 subgroupMatrixScalarMultiply_decf0b() {
+ simdgroup_float8x8 v = make_filled_simdgroup_matrix<float, 8, 8>(0.0f);
+ simdgroup_float8x8 const v_1 = simdgroup_float8x8();
+ simdgroup_multiply(v, v_1, simdgroup_float8x8(8.0f));
+ simdgroup_float8x8 res = v;
+ return res;
+}
+
+kernel void compute_main(device tint_array<float, 1024>* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ simdgroup_store(subgroupMatrixScalarMultiply_decf0b(), (&(*tint_module_vars.prevent_dce)[0u]), ulong(64u), ulong2(0ul), false);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixScalarMultiply/ec8d95.wgsl.expected.msl b/test/tint/builtins/gen/literal/subgroupMatrixScalarMultiply/ec8d95.wgsl.expected.msl
index ee7074a..152ec53 100644
--- a/test/tint/builtins/gen/literal/subgroupMatrixScalarMultiply/ec8d95.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/subgroupMatrixScalarMultiply/ec8d95.wgsl.expected.msl
@@ -1,12 +1,31 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-<dawn>/src/tint/lang/msl/writer/printer/printer.cc:1222 internal compiler error: TINT_UNREACHABLE unhandled: subgroupMatrixScalarMultiply
+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];
+};
-********************************************************************
-* 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. *
-********************************************************************
+struct tint_module_vars_struct {
+ device tint_array<half, 1024>* prevent_dce;
+};
-tint executable returned error: signal: trace/BPT trap
+simdgroup_half8x8 subgroupMatrixScalarMultiply_ec8d95() {
+ simdgroup_half8x8 v = make_filled_simdgroup_matrix<half, 8, 8>(0.0h);
+ simdgroup_half8x8 const v_1 = simdgroup_half8x8();
+ simdgroup_multiply(v, v_1, simdgroup_half8x8(8.0h));
+ simdgroup_half8x8 res = v;
+ return res;
+}
+
+kernel void compute_main(device tint_array<half, 1024>* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ simdgroup_store(subgroupMatrixScalarMultiply_ec8d95(), (&(*tint_module_vars.prevent_dce)[0u]), ulong(64u), ulong2(0ul), false);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixScalarMultiply/0043a7.wgsl.expected.msl b/test/tint/builtins/gen/var/subgroupMatrixScalarMultiply/0043a7.wgsl.expected.msl
index ee7074a..0614bcf 100644
--- a/test/tint/builtins/gen/var/subgroupMatrixScalarMultiply/0043a7.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/subgroupMatrixScalarMultiply/0043a7.wgsl.expected.msl
@@ -1,12 +1,32 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-<dawn>/src/tint/lang/msl/writer/printer/printer.cc:1222 internal compiler error: TINT_UNREACHABLE unhandled: subgroupMatrixScalarMultiply
+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];
+};
-********************************************************************
-* 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. *
-********************************************************************
+struct tint_module_vars_struct {
+ device tint_array<float, 1024>* prevent_dce;
+};
-tint executable returned error: signal: trace/BPT trap
+simdgroup_float8x8 subgroupMatrixScalarMultiply_0043a7() {
+ simdgroup_float8x8 arg_0 = simdgroup_float8x8();
+ float arg_1 = 8.0f;
+ simdgroup_float8x8 v = make_filled_simdgroup_matrix<float, 8, 8>(0.0f);
+ simdgroup_multiply(v, arg_0, simdgroup_float8x8(arg_1));
+ simdgroup_float8x8 res = v;
+ return res;
+}
+
+kernel void compute_main(device tint_array<float, 1024>* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ simdgroup_store(subgroupMatrixScalarMultiply_0043a7(), (&(*tint_module_vars.prevent_dce)[0u]), ulong(64u), ulong2(0ul), false);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixScalarMultiply/3ec99d.wgsl.expected.msl b/test/tint/builtins/gen/var/subgroupMatrixScalarMultiply/3ec99d.wgsl.expected.msl
index ee7074a..8207476 100644
--- a/test/tint/builtins/gen/var/subgroupMatrixScalarMultiply/3ec99d.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/subgroupMatrixScalarMultiply/3ec99d.wgsl.expected.msl
@@ -1,12 +1,32 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-<dawn>/src/tint/lang/msl/writer/printer/printer.cc:1222 internal compiler error: TINT_UNREACHABLE unhandled: subgroupMatrixScalarMultiply
+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];
+};
-********************************************************************
-* 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. *
-********************************************************************
+struct tint_module_vars_struct {
+ device tint_array<half, 1024>* prevent_dce;
+};
-tint executable returned error: signal: trace/BPT trap
+simdgroup_half8x8 subgroupMatrixScalarMultiply_3ec99d() {
+ simdgroup_half8x8 arg_0 = simdgroup_half8x8();
+ half arg_1 = 8.0h;
+ simdgroup_half8x8 v = make_filled_simdgroup_matrix<half, 8, 8>(0.0h);
+ simdgroup_multiply(v, arg_0, simdgroup_half8x8(arg_1));
+ simdgroup_half8x8 res = v;
+ return res;
+}
+
+kernel void compute_main(device tint_array<half, 1024>* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ simdgroup_store(subgroupMatrixScalarMultiply_3ec99d(), (&(*tint_module_vars.prevent_dce)[0u]), ulong(64u), ulong2(0ul), false);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixScalarMultiply/4de795.wgsl.expected.msl b/test/tint/builtins/gen/var/subgroupMatrixScalarMultiply/4de795.wgsl.expected.msl
index ee7074a..3ab7ad4 100644
--- a/test/tint/builtins/gen/var/subgroupMatrixScalarMultiply/4de795.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/subgroupMatrixScalarMultiply/4de795.wgsl.expected.msl
@@ -1,12 +1,32 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-<dawn>/src/tint/lang/msl/writer/printer/printer.cc:1222 internal compiler error: TINT_UNREACHABLE unhandled: subgroupMatrixScalarMultiply
+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];
+};
-********************************************************************
-* 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. *
-********************************************************************
+struct tint_module_vars_struct {
+ device tint_array<half, 1024>* prevent_dce;
+};
-tint executable returned error: signal: trace/BPT trap
+simdgroup_half8x8 subgroupMatrixScalarMultiply_4de795() {
+ simdgroup_half8x8 arg_0 = simdgroup_half8x8();
+ half arg_1 = 8.0h;
+ simdgroup_half8x8 v = make_filled_simdgroup_matrix<half, 8, 8>(0.0h);
+ simdgroup_multiply(v, arg_0, simdgroup_half8x8(arg_1));
+ simdgroup_half8x8 res = v;
+ return res;
+}
+
+kernel void compute_main(device tint_array<half, 1024>* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ simdgroup_store(subgroupMatrixScalarMultiply_4de795(), (&(*tint_module_vars.prevent_dce)[0u]), ulong(64u), ulong2(0ul), false);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixScalarMultiply/9ec228.wgsl.expected.msl b/test/tint/builtins/gen/var/subgroupMatrixScalarMultiply/9ec228.wgsl.expected.msl
index ee7074a..056e5c4 100644
--- a/test/tint/builtins/gen/var/subgroupMatrixScalarMultiply/9ec228.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/subgroupMatrixScalarMultiply/9ec228.wgsl.expected.msl
@@ -1,12 +1,32 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-<dawn>/src/tint/lang/msl/writer/printer/printer.cc:1222 internal compiler error: TINT_UNREACHABLE unhandled: subgroupMatrixScalarMultiply
+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];
+};
-********************************************************************
-* 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. *
-********************************************************************
+struct tint_module_vars_struct {
+ device tint_array<float, 1024>* prevent_dce;
+};
-tint executable returned error: signal: trace/BPT trap
+simdgroup_float8x8 subgroupMatrixScalarMultiply_9ec228() {
+ simdgroup_float8x8 arg_0 = simdgroup_float8x8();
+ float arg_1 = 8.0f;
+ simdgroup_float8x8 v = make_filled_simdgroup_matrix<float, 8, 8>(0.0f);
+ simdgroup_multiply(v, arg_0, simdgroup_float8x8(arg_1));
+ simdgroup_float8x8 res = v;
+ return res;
+}
+
+kernel void compute_main(device tint_array<float, 1024>* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ simdgroup_store(subgroupMatrixScalarMultiply_9ec228(), (&(*tint_module_vars.prevent_dce)[0u]), ulong(64u), ulong2(0ul), false);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixScalarMultiply/decf0b.wgsl.expected.msl b/test/tint/builtins/gen/var/subgroupMatrixScalarMultiply/decf0b.wgsl.expected.msl
index ee7074a..f771f82 100644
--- a/test/tint/builtins/gen/var/subgroupMatrixScalarMultiply/decf0b.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/subgroupMatrixScalarMultiply/decf0b.wgsl.expected.msl
@@ -1,12 +1,32 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-<dawn>/src/tint/lang/msl/writer/printer/printer.cc:1222 internal compiler error: TINT_UNREACHABLE unhandled: subgroupMatrixScalarMultiply
+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];
+};
-********************************************************************
-* 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. *
-********************************************************************
+struct tint_module_vars_struct {
+ device tint_array<float, 1024>* prevent_dce;
+};
-tint executable returned error: signal: trace/BPT trap
+simdgroup_float8x8 subgroupMatrixScalarMultiply_decf0b() {
+ simdgroup_float8x8 arg_0 = simdgroup_float8x8();
+ float arg_1 = 8.0f;
+ simdgroup_float8x8 v = make_filled_simdgroup_matrix<float, 8, 8>(0.0f);
+ simdgroup_multiply(v, arg_0, simdgroup_float8x8(arg_1));
+ simdgroup_float8x8 res = v;
+ return res;
+}
+
+kernel void compute_main(device tint_array<float, 1024>* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ simdgroup_store(subgroupMatrixScalarMultiply_decf0b(), (&(*tint_module_vars.prevent_dce)[0u]), ulong(64u), ulong2(0ul), false);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixScalarMultiply/ec8d95.wgsl.expected.msl b/test/tint/builtins/gen/var/subgroupMatrixScalarMultiply/ec8d95.wgsl.expected.msl
index ee7074a..ef0ff19 100644
--- a/test/tint/builtins/gen/var/subgroupMatrixScalarMultiply/ec8d95.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/subgroupMatrixScalarMultiply/ec8d95.wgsl.expected.msl
@@ -1,12 +1,32 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-<dawn>/src/tint/lang/msl/writer/printer/printer.cc:1222 internal compiler error: TINT_UNREACHABLE unhandled: subgroupMatrixScalarMultiply
+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];
+};
-********************************************************************
-* 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. *
-********************************************************************
+struct tint_module_vars_struct {
+ device tint_array<half, 1024>* prevent_dce;
+};
-tint executable returned error: signal: trace/BPT trap
+simdgroup_half8x8 subgroupMatrixScalarMultiply_ec8d95() {
+ simdgroup_half8x8 arg_0 = simdgroup_half8x8();
+ half arg_1 = 8.0h;
+ simdgroup_half8x8 v = make_filled_simdgroup_matrix<half, 8, 8>(0.0h);
+ simdgroup_multiply(v, arg_0, simdgroup_half8x8(arg_1));
+ simdgroup_half8x8 res = v;
+ return res;
+}
+
+kernel void compute_main(device tint_array<half, 1024>* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ simdgroup_store(subgroupMatrixScalarMultiply_ec8d95(), (&(*tint_module_vars.prevent_dce)[0u]), ulong(64u), ulong2(0ul), false);
+}
