[msl] Handle subgroupMatrixStore builtins

Replace them with simdgroup_store intrinsics in the builtin polyfill
transform. The MSL intrinsics take a pointer to the first element, and
have a `matrix_origin` parameter that must be a `ulong2`.

Bug: 348702031
Change-Id: If45047dc291174c8d605ad1f29b449a6f8bde884
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/224254
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: dan sinclair <dsinclair@chromium.org>

src/tint/lang/core/intrinsic/type_matchers.h

@@ -54,6 +54,7 @@
 #include "src/tint/lang/core/type/storage_texture.h"
 #include "src/tint/lang/core/type/texture_dimension.h"
 #include "src/tint/lang/core/type/u32.h"
+#include "src/tint/lang/core/type/u64.h"
 #include "src/tint/lang/core/type/u8.h"
 #include "src/tint/lang/core/type/vector.h"
 
@@ -109,6 +110,14 @@ inline bool MatchU32(intrinsic::MatchState&, const type::Type* ty) {
     return ty->IsAnyOf<intrinsic::Any, type::U32, type::AbstractInt>();
 }
 
+inline const type::U64* BuildU64(intrinsic::MatchState& state, const type::Type*) {
+    return state.types.u64();
+}
+
+inline bool MatchU64(intrinsic::MatchState&, const type::Type* ty) {
+    return ty->IsAnyOf<intrinsic::Any, type::U64, type::AbstractInt>();
+}
+
 inline const type::U8* BuildU8(intrinsic::MatchState& state, const type::Type*) {
     return state.types.u8();
 }

src/tint/lang/core/ir/transform/remove_terminator_args.h

@@ -41,6 +41,7 @@ namespace tint::core::ir::transform {
 /// The capabilities that the transform can support.
 const core::ir::Capabilities kRemoveTerminatorArgsCapabilities{
     core::ir::Capability::kAllow8BitIntegers,
+    core::ir::Capability::kAllow64BitIntegers,
     core::ir::Capability::kAllowPointersAndHandlesInStructures,
     core::ir::Capability::kAllowVectorElementPointer,
     core::ir::Capability::kAllowHandleVarsWithoutBindings,

src/tint/lang/core/ir/transform/rename_conflicts.h

@@ -41,6 +41,7 @@ namespace tint::core::ir::transform {
 /// The capabilities that the transform can support.
 const core::ir::Capabilities kRenameConflictsCapabilities{
     core::ir::Capability::kAllow8BitIntegers,
+    core::ir::Capability::kAllow64BitIntegers,
     core::ir::Capability::kAllowPointersAndHandlesInStructures,
     core::ir::Capability::kAllowVectorElementPointer,
     core::ir::Capability::kAllowHandleVarsWithoutBindings,

src/tint/lang/core/ir/transform/value_to_let.h

@@ -42,6 +42,7 @@ namespace tint::core::ir::transform {
 /// The capabilities that the transform can support.
 const core::ir::Capabilities kValueToLetCapabilities{
     core::ir::Capability::kAllow8BitIntegers,
+    core::ir::Capability::kAllow64BitIntegers,
     core::ir::Capability::kAllowPointersAndHandlesInStructures,
     core::ir::Capability::kAllowVectorElementPointer,
     core::ir::Capability::kAllowHandleVarsWithoutBindings,

src/tint/lang/msl/builtin_fn.cc

@@ -112,6 +112,8 @@ const char* str(BuiltinFn i) {
             return "quad_shuffle_xor";
         case BuiltinFn::kConvert:
             return "convert";
+        case BuiltinFn::kSimdgroupStore:
+            return "simdgroup_store";
     }
     return "<unknown>";
 }
@@ -144,6 +146,7 @@ tint::core::ir::Instruction::Accesses GetSideEffects(BuiltinFn fn) {
             return core::ir::Instruction::Accesses{core::ir::Instruction::Access::kLoad};
 
         case BuiltinFn::kWrite:
+        case BuiltinFn::kSimdgroupStore:
             return core::ir::Instruction::Accesses{core::ir::Instruction::Access::kStore};
 
         case BuiltinFn::kDistance:

src/tint/lang/msl/builtin_fn.cc.tmpl

@@ -55,6 +55,7 @@ tint::core::ir::Instruction::Accesses GetSideEffects(BuiltinFn fn) {
             return core::ir::Instruction::Accesses{core::ir::Instruction::Access::kLoad};
 
         case BuiltinFn::kWrite:
+        case BuiltinFn::kSimdgroupStore:
             return core::ir::Instruction::Accesses{core::ir::Instruction::Access::kStore};
 
         case BuiltinFn::kDistance:

src/tint/lang/msl/builtin_fn.h

@@ -83,6 +83,7 @@ enum class BuiltinFn : uint8_t {
     kSimdBallot,
     kQuadShuffleXor,
     kConvert,
+    kSimdgroupStore,
     kNone,
 };