[msl] Add support for subgroupMatrixMultiplyAccumulate
Replace them with simdgroup_multiply_accumulate intrinsics in the
builtin polyfill transform. We need to declare a temporary variable to
receive the result.
Bug: 348702031
Change-Id: Ia2fe7b6b60f28600436b9d3eb1d5a61486d2eaa9
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/225138
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Commit-Queue: James Price <jrprice@google.com>
diff --git a/src/tint/lang/msl/builtin_fn.cc b/src/tint/lang/msl/builtin_fn.cc
index 5d6a433..df4d196 100644
--- a/src/tint/lang/msl/builtin_fn.cc
+++ b/src/tint/lang/msl/builtin_fn.cc
@@ -118,6 +118,8 @@
return "simdgroup_store";
case BuiltinFn::kSimdgroupMultiply:
return "simdgroup_multiply";
+ case BuiltinFn::kSimdgroupMultiplyAccumulate:
+ return "simdgroup_multiply_accumulate";
}
return "<unknown>";
}
@@ -170,6 +172,7 @@
case BuiltinFn::kNone:
case BuiltinFn::kConvert:
case BuiltinFn::kSimdgroupMultiply:
+ case BuiltinFn::kSimdgroupMultiplyAccumulate:
break;
}
return core::ir::Instruction::Accesses{};
diff --git a/src/tint/lang/msl/builtin_fn.cc.tmpl b/src/tint/lang/msl/builtin_fn.cc.tmpl
index 291b529..46780a4 100644
--- a/src/tint/lang/msl/builtin_fn.cc.tmpl
+++ b/src/tint/lang/msl/builtin_fn.cc.tmpl
@@ -75,6 +75,7 @@
case BuiltinFn::kNone:
case BuiltinFn::kConvert:
case BuiltinFn::kSimdgroupMultiply:
+ case BuiltinFn::kSimdgroupMultiplyAccumulate:
break;
}
return core::ir::Instruction::Accesses{};
diff --git a/src/tint/lang/msl/builtin_fn.h b/src/tint/lang/msl/builtin_fn.h
index dbbf03e..99a735e 100644
--- a/src/tint/lang/msl/builtin_fn.h
+++ b/src/tint/lang/msl/builtin_fn.h
@@ -86,6 +86,7 @@
kSimdgroupLoad,
kSimdgroupStore,
kSimdgroupMultiply,
+ kSimdgroupMultiplyAccumulate,
kNone,
};
diff --git a/src/tint/lang/msl/intrinsic/data.cc b/src/tint/lang/msl/intrinsic/data.cc
index 39ab930..9435456 100644
--- a/src/tint/lang/msl/intrinsic/data.cc
+++ b/src/tint/lang/msl/intrinsic/data.cc
@@ -2380,539 +2380,539 @@
{
/* [207] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(0),
+ /* matcher_indices */ MatcherIndicesIndex(10),
},
{
/* [208] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(3),
+ /* matcher_indices */ MatcherIndicesIndex(15),
},
{
/* [209] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(26),
+ /* matcher_indices */ MatcherIndicesIndex(20),
},
{
/* [210] */
- /* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(127),
- },
- {
- /* [211] */
- /* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(149),
- },
- {
- /* [212] */
- /* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(1),
- },
- {
- /* [213] */
- /* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(145),
- },
- {
- /* [214] */
- /* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(151),
- },
- {
- /* [215] */
- /* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(1),
- },
- {
- /* [216] */
- /* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(147),
- },
- {
- /* [217] */
- /* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(149),
- },
- {
- /* [218] */
- /* usage */ core::ParameterUsage::kSampleIndex,
- /* matcher_indices */ MatcherIndicesIndex(1),
- },
- {
- /* [219] */
- /* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(186),
- },
- {
- /* [220] */
- /* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(149),
- },
- {
- /* [221] */
- /* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(3),
- },
- {
- /* [222] */
- /* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(191),
- },
- {
- /* [223] */
- /* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(149),
- },
- {
- /* [224] */
- /* usage */ core::ParameterUsage::kSampleIndex,
- /* matcher_indices */ MatcherIndicesIndex(3),
- },
- {
- /* [225] */
- /* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(55),
- },
- {
- /* [226] */
- /* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(149),
- },
- {
- /* [227] */
- /* usage */ core::ParameterUsage::kArrayIndex,
- /* matcher_indices */ MatcherIndicesIndex(3),
- },
- {
- /* [228] */
- /* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(67),
- },
- {
- /* [229] */
- /* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(149),
- },
- {
- /* [230] */
- /* usage */ core::ParameterUsage::kArrayIndex,
- /* matcher_indices */ MatcherIndicesIndex(3),
- },
- {
- /* [231] */
- /* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(79),
- },
- {
- /* [232] */
- /* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(149),
- },
- {
- /* [233] */
- /* usage */ core::ParameterUsage::kArrayIndex,
- /* matcher_indices */ MatcherIndicesIndex(3),
- },
- {
- /* [234] */
- /* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(157),
- },
- {
- /* [235] */
- /* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(185),
- },
- {
- /* [236] */
- /* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(62),
- },
- {
- /* [237] */
- /* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(85),
- },
- {
- /* [238] */
- /* usage */ core::ParameterUsage::kValue,
- /* matcher_indices */ MatcherIndicesIndex(141),
- },
- {
- /* [239] */
- /* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(26),
- },
- {
- /* [240] */
- /* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(88),
- },
- {
- /* [241] */
- /* usage */ core::ParameterUsage::kValue,
- /* matcher_indices */ MatcherIndicesIndex(141),
- },
- {
- /* [242] */
- /* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(149),
- },
- {
- /* [243] */
- /* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(94),
- },
- {
- /* [244] */
- /* usage */ core::ParameterUsage::kValue,
- /* matcher_indices */ MatcherIndicesIndex(141),
- },
- {
- /* [245] */
- /* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(151),
- },
- {
- /* [246] */
- /* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(97),
- },
- {
- /* [247] */
- /* usage */ core::ParameterUsage::kValue,
- /* matcher_indices */ MatcherIndicesIndex(153),
- },
- {
- /* [248] */
- /* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(26),
- },
- {
- /* [249] */
- /* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(100),
- },
- {
- /* [250] */
- /* usage */ core::ParameterUsage::kValue,
- /* matcher_indices */ MatcherIndicesIndex(153),
- },
- {
- /* [251] */
- /* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(149),
- },
- {
- /* [252] */
- /* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(106),
- },
- {
- /* [253] */
- /* usage */ core::ParameterUsage::kValue,
- /* matcher_indices */ MatcherIndicesIndex(153),
- },
- {
- /* [254] */
- /* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(151),
- },
- {
- /* [255] */
- /* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(109),
- },
- {
- /* [256] */
- /* usage */ core::ParameterUsage::kValue,
- /* matcher_indices */ MatcherIndicesIndex(155),
- },
- {
- /* [257] */
- /* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(26),
- },
- {
- /* [258] */
- /* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(112),
- },
- {
- /* [259] */
- /* usage */ core::ParameterUsage::kValue,
- /* matcher_indices */ MatcherIndicesIndex(155),
- },
- {
- /* [260] */
- /* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(149),
- },
- {
- /* [261] */
- /* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(118),
- },
- {
- /* [262] */
- /* usage */ core::ParameterUsage::kValue,
- /* matcher_indices */ MatcherIndicesIndex(155),
- },
- {
- /* [263] */
- /* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(151),
- },
- {
- /* [264] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(10),
},
{
- /* [265] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(15),
- },
- {
- /* [266] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(20),
- },
- {
- /* [267] */
+ /* [211] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(0),
},
{
- /* [268] */
+ /* [212] */
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(3),
+ },
+ {
+ /* [213] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(26),
},
{
- /* [269] */
+ /* [214] */
/* usage */ core::ParameterUsage::kTexture,
/* matcher_indices */ MatcherIndicesIndex(127),
},
{
- /* [270] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(26),
+ /* [215] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(149),
},
{
- /* [271] */
- /* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(133),
+ /* [216] */
+ /* usage */ core::ParameterUsage::kLevel,
+ /* matcher_indices */ MatcherIndicesIndex(1),
},
{
- /* [272] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(26),
- },
- {
- /* [273] */
+ /* [217] */
/* usage */ core::ParameterUsage::kTexture,
/* matcher_indices */ MatcherIndicesIndex(145),
},
{
- /* [274] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(26),
+ /* [218] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(151),
},
{
- /* [275] */
+ /* [219] */
+ /* usage */ core::ParameterUsage::kLevel,
+ /* matcher_indices */ MatcherIndicesIndex(1),
+ },
+ {
+ /* [220] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(135),
+ /* matcher_indices */ MatcherIndicesIndex(147),
},
{
- /* [276] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(26),
+ /* [221] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(149),
},
{
- /* [277] */
- /* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(139),
+ /* [222] */
+ /* usage */ core::ParameterUsage::kSampleIndex,
+ /* matcher_indices */ MatcherIndicesIndex(1),
},
{
- /* [278] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(26),
- },
- {
- /* [279] */
+ /* [223] */
/* usage */ core::ParameterUsage::kTexture,
/* matcher_indices */ MatcherIndicesIndex(186),
},
{
- /* [280] */
+ /* [224] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(149),
+ },
+ {
+ /* [225] */
+ /* usage */ core::ParameterUsage::kLevel,
+ /* matcher_indices */ MatcherIndicesIndex(3),
+ },
+ {
+ /* [226] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(191),
+ },
+ {
+ /* [227] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(149),
+ },
+ {
+ /* [228] */
+ /* usage */ core::ParameterUsage::kSampleIndex,
+ /* matcher_indices */ MatcherIndicesIndex(3),
+ },
+ {
+ /* [229] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(55),
+ },
+ {
+ /* [230] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(149),
+ },
+ {
+ /* [231] */
+ /* usage */ core::ParameterUsage::kArrayIndex,
+ /* matcher_indices */ MatcherIndicesIndex(3),
+ },
+ {
+ /* [232] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(67),
+ },
+ {
+ /* [233] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(149),
+ },
+ {
+ /* [234] */
+ /* usage */ core::ParameterUsage::kArrayIndex,
+ /* matcher_indices */ MatcherIndicesIndex(3),
+ },
+ {
+ /* [235] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(79),
+ },
+ {
+ /* [236] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(149),
+ },
+ {
+ /* [237] */
+ /* usage */ core::ParameterUsage::kArrayIndex,
+ /* matcher_indices */ MatcherIndicesIndex(3),
+ },
+ {
+ /* [238] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(157),
+ },
+ {
+ /* [239] */
+ /* usage */ core::ParameterUsage::kSampler,
+ /* matcher_indices */ MatcherIndicesIndex(185),
+ },
+ {
+ /* [240] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(62),
+ },
+ {
+ /* [241] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(85),
+ },
+ {
+ /* [242] */
+ /* usage */ core::ParameterUsage::kValue,
+ /* matcher_indices */ MatcherIndicesIndex(141),
+ },
+ {
+ /* [243] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(26),
+ },
+ {
+ /* [244] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(88),
+ },
+ {
+ /* [245] */
+ /* usage */ core::ParameterUsage::kValue,
+ /* matcher_indices */ MatcherIndicesIndex(141),
+ },
+ {
+ /* [246] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(149),
+ },
+ {
+ /* [247] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(94),
+ },
+ {
+ /* [248] */
+ /* usage */ core::ParameterUsage::kValue,
+ /* matcher_indices */ MatcherIndicesIndex(141),
+ },
+ {
+ /* [249] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(151),
+ },
+ {
+ /* [250] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(97),
+ },
+ {
+ /* [251] */
+ /* usage */ core::ParameterUsage::kValue,
+ /* matcher_indices */ MatcherIndicesIndex(153),
+ },
+ {
+ /* [252] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(26),
+ },
+ {
+ /* [253] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(100),
+ },
+ {
+ /* [254] */
+ /* usage */ core::ParameterUsage::kValue,
+ /* matcher_indices */ MatcherIndicesIndex(153),
+ },
+ {
+ /* [255] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(149),
+ },
+ {
+ /* [256] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(106),
+ },
+ {
+ /* [257] */
+ /* usage */ core::ParameterUsage::kValue,
+ /* matcher_indices */ MatcherIndicesIndex(153),
+ },
+ {
+ /* [258] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(151),
+ },
+ {
+ /* [259] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(109),
+ },
+ {
+ /* [260] */
+ /* usage */ core::ParameterUsage::kValue,
+ /* matcher_indices */ MatcherIndicesIndex(155),
+ },
+ {
+ /* [261] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(26),
+ },
+ {
+ /* [262] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(112),
+ },
+ {
+ /* [263] */
+ /* usage */ core::ParameterUsage::kValue,
+ /* matcher_indices */ MatcherIndicesIndex(155),
+ },
+ {
+ /* [264] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(149),
+ },
+ {
+ /* [265] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(118),
+ },
+ {
+ /* [266] */
+ /* usage */ core::ParameterUsage::kValue,
+ /* matcher_indices */ MatcherIndicesIndex(155),
+ },
+ {
+ /* [267] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(151),
+ },
+ {
+ /* [268] */
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(0),
+ },
+ {
+ /* [269] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(26),
},
{
+ /* [270] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(127),
+ },
+ {
+ /* [271] */
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(26),
+ },
+ {
+ /* [272] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(133),
+ },
+ {
+ /* [273] */
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(26),
+ },
+ {
+ /* [274] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(145),
+ },
+ {
+ /* [275] */
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(26),
+ },
+ {
+ /* [276] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(135),
+ },
+ {
+ /* [277] */
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(26),
+ },
+ {
+ /* [278] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(139),
+ },
+ {
+ /* [279] */
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(26),
+ },
+ {
+ /* [280] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(186),
+ },
+ {
/* [281] */
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(26),
+ },
+ {
+ /* [282] */
/* usage */ core::ParameterUsage::kTexture,
/* matcher_indices */ MatcherIndicesIndex(187),
},
{
- /* [282] */
+ /* [283] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(26),
},
{
- /* [283] */
+ /* [284] */
/* usage */ core::ParameterUsage::kTexture,
/* matcher_indices */ MatcherIndicesIndex(188),
},
{
- /* [284] */
+ /* [285] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(26),
},
{
- /* [285] */
+ /* [286] */
/* usage */ core::ParameterUsage::kTexture,
/* matcher_indices */ MatcherIndicesIndex(189),
},
{
- /* [286] */
+ /* [287] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(26),
},
{
- /* [287] */
+ /* [288] */
/* usage */ core::ParameterUsage::kTexture,
/* matcher_indices */ MatcherIndicesIndex(40),
},
{
- /* [288] */
+ /* [289] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(26),
},
{
- /* [289] */
+ /* [290] */
/* usage */ core::ParameterUsage::kTexture,
/* matcher_indices */ MatcherIndicesIndex(43),
},
{
- /* [290] */
+ /* [291] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(26),
},
{
- /* [291] */
+ /* [292] */
/* usage */ core::ParameterUsage::kTexture,
/* matcher_indices */ MatcherIndicesIndex(46),
},
{
- /* [292] */
+ /* [293] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(26),
},
{
- /* [293] */
+ /* [294] */
/* usage */ core::ParameterUsage::kTexture,
/* matcher_indices */ MatcherIndicesIndex(143),
},
{
- /* [294] */
+ /* [295] */
/* usage */ core::ParameterUsage::kCoords,
/* matcher_indices */ MatcherIndicesIndex(26),
},
{
- /* [295] */
+ /* [296] */
/* usage */ core::ParameterUsage::kTexture,
/* matcher_indices */ MatcherIndicesIndex(49),
},
{
- /* [296] */
+ /* [297] */
/* usage */ core::ParameterUsage::kCoords,
/* matcher_indices */ MatcherIndicesIndex(26),
},
{
- /* [297] */
+ /* [298] */
/* usage */ core::ParameterUsage::kTexture,
/* matcher_indices */ MatcherIndicesIndex(52),
},
{
- /* [298] */
+ /* [299] */
/* usage */ core::ParameterUsage::kCoords,
/* matcher_indices */ MatcherIndicesIndex(149),
},
{
- /* [299] */
+ /* [300] */
/* usage */ core::ParameterUsage::kTexture,
/* matcher_indices */ MatcherIndicesIndex(58),
},
{
- /* [300] */
+ /* [301] */
/* usage */ core::ParameterUsage::kCoords,
/* matcher_indices */ MatcherIndicesIndex(151),
},
{
- /* [301] */
+ /* [302] */
/* usage */ core::ParameterUsage::kTexture,
/* matcher_indices */ MatcherIndicesIndex(61),
},
{
- /* [302] */
+ /* [303] */
/* usage */ core::ParameterUsage::kCoords,
/* matcher_indices */ MatcherIndicesIndex(26),
},
{
- /* [303] */
+ /* [304] */
/* usage */ core::ParameterUsage::kTexture,
/* matcher_indices */ MatcherIndicesIndex(64),
},
{
- /* [304] */
+ /* [305] */
/* usage */ core::ParameterUsage::kCoords,
/* matcher_indices */ MatcherIndicesIndex(149),
},
{
- /* [305] */
+ /* [306] */
/* usage */ core::ParameterUsage::kTexture,
/* matcher_indices */ MatcherIndicesIndex(70),
},
{
- /* [306] */
+ /* [307] */
/* usage */ core::ParameterUsage::kCoords,
/* matcher_indices */ MatcherIndicesIndex(151),
},
{
- /* [307] */
+ /* [308] */
/* usage */ core::ParameterUsage::kTexture,
/* matcher_indices */ MatcherIndicesIndex(73),
},
{
- /* [308] */
+ /* [309] */
/* usage */ core::ParameterUsage::kCoords,
/* matcher_indices */ MatcherIndicesIndex(26),
},
{
- /* [309] */
+ /* [310] */
/* usage */ core::ParameterUsage::kTexture,
/* matcher_indices */ MatcherIndicesIndex(76),
},
{
- /* [310] */
+ /* [311] */
/* usage */ core::ParameterUsage::kCoords,
/* matcher_indices */ MatcherIndicesIndex(149),
},
{
- /* [311] */
+ /* [312] */
/* usage */ core::ParameterUsage::kTexture,
/* matcher_indices */ MatcherIndicesIndex(82),
},
{
- /* [312] */
+ /* [313] */
/* usage */ core::ParameterUsage::kCoords,
/* matcher_indices */ MatcherIndicesIndex(151),
},
{
- /* [313] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(121),
- },
- {
/* [314] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(121),
@@ -2920,22 +2920,22 @@
{
/* [315] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(1),
+ /* matcher_indices */ MatcherIndicesIndex(121),
},
{
/* [316] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(121),
+ /* matcher_indices */ MatcherIndicesIndex(1),
},
{
/* [317] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(124),
+ /* matcher_indices */ MatcherIndicesIndex(121),
},
{
/* [318] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(3),
+ /* matcher_indices */ MatcherIndicesIndex(124),
},
{
/* [319] */
@@ -2945,60 +2945,65 @@
{
/* [320] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(32),
+ /* matcher_indices */ MatcherIndicesIndex(3),
},
{
/* [321] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(121),
+ /* matcher_indices */ MatcherIndicesIndex(32),
},
{
/* [322] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(26),
+ /* matcher_indices */ MatcherIndicesIndex(121),
},
{
/* [323] */
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(26),
+ },
+ {
+ /* [324] */
/* usage */ core::ParameterUsage::kTexture,
/* matcher_indices */ MatcherIndicesIndex(37),
},
{
- /* [324] */
+ /* [325] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(171),
},
{
- /* [325] */
+ /* [326] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(173),
},
{
- /* [326] */
+ /* [327] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(175),
},
{
- /* [327] */
+ /* [328] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(179),
},
{
- /* [328] */
+ /* [329] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(151),
},
{
- /* [329] */
+ /* [330] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(165),
},
{
- /* [330] */
+ /* [331] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(137),
},
{
- /* [331] */
+ /* [332] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(177),
},
@@ -3147,7 +3152,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(234),
+ /* parameters */ ParameterIndex(238),
/* return_matcher_indices */ MatcherIndicesIndex(141),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3642,7 +3647,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(293),
+ /* parameters */ ParameterIndex(294),
/* return_matcher_indices */ MatcherIndicesIndex(11),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3653,7 +3658,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(12),
- /* parameters */ ParameterIndex(210),
+ /* parameters */ ParameterIndex(214),
/* return_matcher_indices */ MatcherIndicesIndex(11),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3675,7 +3680,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(12),
- /* parameters */ ParameterIndex(213),
+ /* parameters */ ParameterIndex(217),
/* return_matcher_indices */ MatcherIndicesIndex(11),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3686,7 +3691,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(14),
- /* parameters */ ParameterIndex(216),
+ /* parameters */ ParameterIndex(220),
/* return_matcher_indices */ MatcherIndicesIndex(11),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3697,7 +3702,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(7),
- /* parameters */ ParameterIndex(219),
+ /* parameters */ ParameterIndex(223),
/* return_matcher_indices */ MatcherIndicesIndex(62),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3719,7 +3724,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(15),
- /* parameters */ ParameterIndex(222),
+ /* parameters */ ParameterIndex(226),
/* return_matcher_indices */ MatcherIndicesIndex(62),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3730,7 +3735,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(295),
+ /* parameters */ ParameterIndex(296),
/* return_matcher_indices */ MatcherIndicesIndex(141),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3741,7 +3746,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(297),
+ /* parameters */ ParameterIndex(298),
/* return_matcher_indices */ MatcherIndicesIndex(141),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3752,7 +3757,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(6),
- /* parameters */ ParameterIndex(225),
+ /* parameters */ ParameterIndex(229),
/* return_matcher_indices */ MatcherIndicesIndex(141),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3763,7 +3768,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(299),
+ /* parameters */ ParameterIndex(300),
/* return_matcher_indices */ MatcherIndicesIndex(141),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3774,7 +3779,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(301),
+ /* parameters */ ParameterIndex(302),
/* return_matcher_indices */ MatcherIndicesIndex(153),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3785,7 +3790,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(303),
+ /* parameters */ ParameterIndex(304),
/* return_matcher_indices */ MatcherIndicesIndex(153),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3796,7 +3801,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(6),
- /* parameters */ ParameterIndex(228),
+ /* parameters */ ParameterIndex(232),
/* return_matcher_indices */ MatcherIndicesIndex(153),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3807,7 +3812,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(305),
+ /* parameters */ ParameterIndex(306),
/* return_matcher_indices */ MatcherIndicesIndex(153),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3818,7 +3823,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(307),
+ /* parameters */ ParameterIndex(308),
/* return_matcher_indices */ MatcherIndicesIndex(155),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3829,7 +3834,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(309),
+ /* parameters */ ParameterIndex(310),
/* return_matcher_indices */ MatcherIndicesIndex(155),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3840,7 +3845,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(6),
- /* parameters */ ParameterIndex(231),
+ /* parameters */ ParameterIndex(235),
/* return_matcher_indices */ MatcherIndicesIndex(155),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3851,7 +3856,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(311),
+ /* parameters */ ParameterIndex(312),
/* return_matcher_indices */ MatcherIndicesIndex(155),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3862,7 +3867,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(293),
+ /* parameters */ ParameterIndex(294),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3873,7 +3878,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(269),
+ /* parameters */ ParameterIndex(270),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3884,7 +3889,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(271),
+ /* parameters */ ParameterIndex(272),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3895,7 +3900,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(273),
+ /* parameters */ ParameterIndex(274),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3906,7 +3911,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(275),
+ /* parameters */ ParameterIndex(276),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3917,7 +3922,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(277),
+ /* parameters */ ParameterIndex(278),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3928,7 +3933,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(216),
+ /* parameters */ ParameterIndex(220),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3939,7 +3944,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(279),
+ /* parameters */ ParameterIndex(280),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3950,7 +3955,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(281),
+ /* parameters */ ParameterIndex(282),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3961,7 +3966,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(283),
+ /* parameters */ ParameterIndex(284),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3972,7 +3977,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(285),
+ /* parameters */ ParameterIndex(286),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3983,7 +3988,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(222),
+ /* parameters */ ParameterIndex(226),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3994,7 +3999,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(10),
- /* parameters */ ParameterIndex(323),
+ /* parameters */ ParameterIndex(324),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4005,7 +4010,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(10),
- /* parameters */ ParameterIndex(287),
+ /* parameters */ ParameterIndex(288),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4016,7 +4021,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(10),
- /* parameters */ ParameterIndex(289),
+ /* parameters */ ParameterIndex(290),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4027,7 +4032,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(10),
- /* parameters */ ParameterIndex(291),
+ /* parameters */ ParameterIndex(292),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4038,7 +4043,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(269),
+ /* parameters */ ParameterIndex(270),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4049,7 +4054,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(271),
+ /* parameters */ ParameterIndex(272),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4060,7 +4065,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(273),
+ /* parameters */ ParameterIndex(274),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4071,7 +4076,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(275),
+ /* parameters */ ParameterIndex(276),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4082,7 +4087,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(277),
+ /* parameters */ ParameterIndex(278),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4093,7 +4098,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(216),
+ /* parameters */ ParameterIndex(220),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4104,7 +4109,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(279),
+ /* parameters */ ParameterIndex(280),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4115,7 +4120,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(281),
+ /* parameters */ ParameterIndex(282),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4126,7 +4131,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(283),
+ /* parameters */ ParameterIndex(284),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4137,7 +4142,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(285),
+ /* parameters */ ParameterIndex(286),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4148,7 +4153,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(222),
+ /* parameters */ ParameterIndex(226),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4159,7 +4164,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(10),
- /* parameters */ ParameterIndex(287),
+ /* parameters */ ParameterIndex(288),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4170,7 +4175,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(10),
- /* parameters */ ParameterIndex(289),
+ /* parameters */ ParameterIndex(290),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4181,7 +4186,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(10),
- /* parameters */ ParameterIndex(291),
+ /* parameters */ ParameterIndex(292),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4324,7 +4329,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(237),
+ /* parameters */ ParameterIndex(241),
/* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4335,7 +4340,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(240),
+ /* parameters */ ParameterIndex(244),
/* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4357,7 +4362,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(243),
+ /* parameters */ ParameterIndex(247),
/* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4368,7 +4373,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(246),
+ /* parameters */ ParameterIndex(250),
/* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4379,7 +4384,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(249),
+ /* parameters */ ParameterIndex(253),
/* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4401,7 +4406,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(252),
+ /* parameters */ ParameterIndex(256),
/* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4412,7 +4417,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(255),
+ /* parameters */ ParameterIndex(259),
/* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4423,7 +4428,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(258),
+ /* parameters */ ParameterIndex(262),
/* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4445,7 +4450,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(261),
+ /* parameters */ ParameterIndex(265),
/* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4456,7 +4461,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(293),
+ /* parameters */ ParameterIndex(294),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4489,7 +4494,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(213),
+ /* parameters */ ParameterIndex(217),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4566,7 +4571,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(324),
+ /* parameters */ ParameterIndex(325),
/* return_matcher_indices */ MatcherIndicesIndex(151),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4577,7 +4582,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(325),
+ /* parameters */ ParameterIndex(326),
/* return_matcher_indices */ MatcherIndicesIndex(165),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4588,7 +4593,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(326),
+ /* parameters */ ParameterIndex(327),
/* return_matcher_indices */ MatcherIndicesIndex(137),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4599,7 +4604,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(327),
+ /* parameters */ ParameterIndex(328),
/* return_matcher_indices */ MatcherIndicesIndex(177),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4610,7 +4615,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(328),
+ /* parameters */ ParameterIndex(329),
/* return_matcher_indices */ MatcherIndicesIndex(171),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4621,7 +4626,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(329),
+ /* parameters */ ParameterIndex(330),
/* return_matcher_indices */ MatcherIndicesIndex(173),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4632,7 +4637,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(330),
+ /* parameters */ ParameterIndex(331),
/* return_matcher_indices */ MatcherIndicesIndex(175),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4643,7 +4648,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(331),
+ /* parameters */ ParameterIndex(332),
/* return_matcher_indices */ MatcherIndicesIndex(179),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4863,7 +4868,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(10),
- /* parameters */ ParameterIndex(289),
+ /* parameters */ ParameterIndex(290),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4874,7 +4879,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(10),
- /* parameters */ ParameterIndex(323),
+ /* parameters */ ParameterIndex(324),
/* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4885,7 +4890,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(10),
- /* parameters */ ParameterIndex(287),
+ /* parameters */ ParameterIndex(288),
/* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4896,7 +4901,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(10),
- /* parameters */ ParameterIndex(289),
+ /* parameters */ ParameterIndex(290),
/* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4907,7 +4912,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(10),
- /* parameters */ ParameterIndex(291),
+ /* parameters */ ParameterIndex(292),
/* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4918,7 +4923,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(17),
- /* parameters */ ParameterIndex(318),
+ /* parameters */ ParameterIndex(319),
/* return_matcher_indices */ MatcherIndicesIndex(3),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4929,7 +4934,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(16),
- /* parameters */ ParameterIndex(313),
+ /* parameters */ ParameterIndex(314),
/* return_matcher_indices */ MatcherIndicesIndex(121),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4940,7 +4945,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(16),
- /* parameters */ ParameterIndex(315),
+ /* parameters */ ParameterIndex(316),
/* return_matcher_indices */ MatcherIndicesIndex(121),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4951,7 +4956,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(16),
- /* parameters */ ParameterIndex(314),
+ /* parameters */ ParameterIndex(315),
/* return_matcher_indices */ MatcherIndicesIndex(121),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4962,7 +4967,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(273),
+ /* parameters */ ParameterIndex(274),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4973,7 +4978,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(10),
- /* parameters */ ParameterIndex(291),
+ /* parameters */ ParameterIndex(292),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4984,7 +4989,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(216),
+ /* parameters */ ParameterIndex(220),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4995,7 +5000,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(222),
+ /* parameters */ ParameterIndex(226),
/* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -5006,7 +5011,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(17),
- /* parameters */ ParameterIndex(319),
+ /* parameters */ ParameterIndex(320),
/* return_matcher_indices */ MatcherIndicesIndex(3),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -5017,7 +5022,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(16),
- /* parameters */ ParameterIndex(316),
+ /* parameters */ ParameterIndex(317),
/* return_matcher_indices */ MatcherIndicesIndex(121),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -5039,7 +5044,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(16),
- /* parameters */ ParameterIndex(313),
+ /* parameters */ ParameterIndex(314),
/* return_matcher_indices */ MatcherIndicesIndex(121),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -5061,7 +5066,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(18),
- /* parameters */ ParameterIndex(321),
+ /* parameters */ ParameterIndex(322),
/* return_matcher_indices */ MatcherIndicesIndex(121),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -5083,7 +5088,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(8),
- /* parameters */ ParameterIndex(207),
+ /* parameters */ ParameterIndex(211),
/* return_matcher_indices */ MatcherIndicesIndex(3),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -5094,7 +5099,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(8),
- /* parameters */ ParameterIndex(267),
+ /* parameters */ ParameterIndex(268),
/* return_matcher_indices */ MatcherIndicesIndex(3),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -5105,7 +5110,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(8),
- /* parameters */ ParameterIndex(207),
+ /* parameters */ ParameterIndex(211),
/* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -5116,7 +5121,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(16),
- /* parameters */ ParameterIndex(313),
+ /* parameters */ ParameterIndex(314),
/* return_matcher_indices */ MatcherIndicesIndex(1),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -5127,7 +5132,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(16),
- /* parameters */ ParameterIndex(313),
+ /* parameters */ ParameterIndex(314),
/* return_matcher_indices */ MatcherIndicesIndex(1),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -5182,18 +5187,29 @@
/* num_explicit_templates */ 0,
/* num_templates */ 4,
/* templates */ TemplateIndex(1),
- /* parameters */ ParameterIndex(264),
+ /* parameters */ ParameterIndex(207),
/* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
/* [186] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 4,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 4,
+ /* templates */ TemplateIndex(1),
+ /* parameters */ ParameterIndex(207),
+ /* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [187] */
/* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(20),
- /* parameters */ ParameterIndex(318),
+ /* parameters */ ParameterIndex(319),
/* return_matcher_indices */ MatcherIndicesIndex(3),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -5589,6 +5605,12 @@
/* num overloads */ 1,
/* overloads */ OverloadIndex(185),
},
+ {
+ /* [38] */
+ /* fn simdgroup_multiply_accumulate[S : f32_f16, C : num, R : num, K : num](subgroup_matrix<subgroup_matrix_kind_result, S, C, R>, subgroup_matrix<subgroup_matrix_kind_left, S, K, R>, subgroup_matrix<subgroup_matrix_kind_right, S, C, K>, subgroup_matrix<subgroup_matrix_kind_result, S, C, R>) */
+ /* num overloads */ 1,
+ /* overloads */ OverloadIndex(186),
+ },
};
constexpr IntrinsicInfo kBinaryOperators[] = {
@@ -5596,13 +5618,13 @@
/* [0] */
/* op +[T : iu8](T, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(186),
+ /* overloads */ OverloadIndex(187),
},
{
/* [1] */
/* op *[T : iu8](T, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(186),
+ /* overloads */ OverloadIndex(187),
},
};
constexpr uint8_t kBinaryOperatorPlus = 0;
diff --git a/src/tint/lang/msl/msl.def b/src/tint/lang/msl/msl.def
index dc81272..5f5f8d8 100644
--- a/src/tint/lang/msl/msl.def
+++ b/src/tint/lang/msl/msl.def
@@ -371,6 +371,11 @@
subgroup_matrix<subgroup_matrix_kind_result, S, C, R>,
subgroup_matrix<subgroup_matrix_kind_left, S, K, R>,
subgroup_matrix<subgroup_matrix_kind_right, S, C, K>)
+@stage("compute") implicit(S: f32_f16, C: num, R: num, K: num) fn simdgroup_multiply_accumulate(
+ subgroup_matrix<subgroup_matrix_kind_result, S, C, R>,
+ subgroup_matrix<subgroup_matrix_kind_left, S, K, R>,
+ subgroup_matrix<subgroup_matrix_kind_right, S, C, K>,
+ subgroup_matrix<subgroup_matrix_kind_result, S, C, R>)
////////////////////////////////////////////////////////////////////////////////
// Binary Operators //
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
index 2581f69..a61b20a 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
@@ -111,6 +111,7 @@
case core::BuiltinFn::kSubgroupMatrixLoad:
case core::BuiltinFn::kSubgroupMatrixStore:
case core::BuiltinFn::kSubgroupMatrixMultiply:
+ case core::BuiltinFn::kSubgroupMatrixMultiplyAccumulate:
case core::BuiltinFn::kTextureDimensions:
case core::BuiltinFn::kTextureGather:
case core::BuiltinFn::kTextureGatherCompare:
@@ -282,6 +283,9 @@
case core::BuiltinFn::kSubgroupMatrixMultiply:
SubgroupMatrixMultiply(builtin);
break;
+ case core::BuiltinFn::kSubgroupMatrixMultiplyAccumulate:
+ SubgroupMatrixMultiplyAccumulate(builtin);
+ break;
default:
break;
@@ -1029,6 +1033,27 @@
});
builtin->Destroy();
}
+
+ /// Replace a subgroupMatrixMultiplyAccumulate builtin.
+ /// @param builtin the builtin call instruction
+ void SubgroupMatrixMultiplyAccumulate(core::ir::CoreBuiltinCall* builtin) {
+ b.InsertBefore(builtin, [&] {
+ auto* left = builtin->Args()[0];
+ auto* right = builtin->Args()[1];
+ auto* acc = builtin->Args()[2];
+
+ // Declare a local variable to receive the result.
+ auto* tmp = b.Var(ty.ptr<function>(builtin->Result(0)->Type()));
+ // 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::kSimdgroupMultiplyAccumulate,
+ b.Load(tmp->Result(0)), left, right, acc);
+ b.LoadWithResult(builtin->DetachResult(), tmp);
+ });
+ builtin->Destroy();
+ }
};
} // namespace
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
index 5d8eb2d..4e0bfe7 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
@@ -3546,5 +3546,85 @@
EXPECT_EQ(expect, str());
}
+TEST_F(MslWriter_BuiltinPolyfillTest, SubgroupMatrixMultiplyAccumulate_F32) {
+ auto* left = b.FunctionParam("left", ty.subgroup_matrix_left(ty.f32(), 4, 8));
+ auto* right = b.FunctionParam("right", ty.subgroup_matrix_right(ty.f32(), 8, 4));
+ auto* acc = b.FunctionParam("acc", ty.subgroup_matrix_result(ty.f32(), 8, 8));
+ auto* result = ty.subgroup_matrix_result(ty.f32(), 8, 8);
+ auto* func = b.Function("foo", result);
+ func->SetParams({left, right, acc});
+ b.Append(func->Block(), [&] {
+ auto* call =
+ b.Call(result, core::BuiltinFn::kSubgroupMatrixMultiplyAccumulate, left, right, acc);
+ b.Return(func, call);
+ });
+
+ auto* src = R"(
+%foo = func(%left:subgroup_matrix_left<f32, 4, 8>, %right:subgroup_matrix_right<f32, 8, 4>, %acc:subgroup_matrix_result<f32, 8, 8>):subgroup_matrix_result<f32, 8, 8> {
+ $B1: {
+ %5:subgroup_matrix_result<f32, 8, 8> = subgroupMatrixMultiplyAccumulate %left, %right, %acc
+ ret %5
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%left:subgroup_matrix_left<f32, 4, 8>, %right:subgroup_matrix_right<f32, 8, 4>, %acc:subgroup_matrix_result<f32, 8, 8>):subgroup_matrix_result<f32, 8, 8> {
+ $B1: {
+ %5:ptr<function, subgroup_matrix_result<f32, 8, 8>, read_write> = var
+ %6:subgroup_matrix_result<f32, 8, 8> = load %5
+ %7:void = msl.simdgroup_multiply_accumulate %6, %left, %right, %acc
+ %8:subgroup_matrix_result<f32, 8, 8> = load %5
+ ret %8
+ }
+}
+)";
+
+ Run(BuiltinPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_BuiltinPolyfillTest, SubgroupMatrixMultiplyAccumulate_F16) {
+ auto* left = b.FunctionParam("left", ty.subgroup_matrix_left(ty.f16(), 8, 4));
+ auto* right = b.FunctionParam("right", ty.subgroup_matrix_right(ty.f16(), 2, 8));
+ auto* acc = b.FunctionParam("acc", ty.subgroup_matrix_result(ty.f16(), 2, 4));
+ auto* result = ty.subgroup_matrix_result(ty.f16(), 2, 4);
+ auto* func = b.Function("foo", result);
+ func->SetParams({left, right, acc});
+ b.Append(func->Block(), [&] {
+ auto* call =
+ b.Call(result, core::BuiltinFn::kSubgroupMatrixMultiplyAccumulate, left, right, acc);
+ b.Return(func, call);
+ });
+
+ auto* src = R"(
+%foo = func(%left:subgroup_matrix_left<f16, 8, 4>, %right:subgroup_matrix_right<f16, 2, 8>, %acc:subgroup_matrix_result<f16, 2, 4>):subgroup_matrix_result<f16, 2, 4> {
+ $B1: {
+ %5:subgroup_matrix_result<f16, 2, 4> = subgroupMatrixMultiplyAccumulate %left, %right, %acc
+ ret %5
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%left:subgroup_matrix_left<f16, 8, 4>, %right:subgroup_matrix_right<f16, 2, 8>, %acc:subgroup_matrix_result<f16, 2, 4>):subgroup_matrix_result<f16, 2, 4> {
+ $B1: {
+ %5:ptr<function, subgroup_matrix_result<f16, 2, 4>, read_write> = var
+ %6:subgroup_matrix_result<f16, 2, 4> = load %5
+ %7:void = msl.simdgroup_multiply_accumulate %6, %left, %right, %acc
+ %8:subgroup_matrix_result<f16, 2, 4> = load %5
+ ret %8
+ }
+}
+)";
+
+ Run(BuiltinPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
} // namespace
} // namespace tint::msl::writer::raise
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixMultiplyAccumulate/8b907c.wgsl.expected.msl b/test/tint/builtins/gen/literal/subgroupMatrixMultiplyAccumulate/8b907c.wgsl.expected.msl
index 171ead9..bd5a0a2 100644
--- a/test/tint/builtins/gen/literal/subgroupMatrixMultiplyAccumulate/8b907c.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/subgroupMatrixMultiplyAccumulate/8b907c.wgsl.expected.msl
@@ -1,11 +1,33 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:1185 internal compiler error: TINT_UNREACHABLE unhandled: subgroupMatrixMultiplyAccumulate
-********************************************************************
-* 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. *
-********************************************************************
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-tint executable returned error: signal: trace/BPT trap
+struct tint_module_vars_struct {
+ device tint_array<half, 1024>* prevent_dce;
+};
+
+simdgroup_half8x8 subgroupMatrixMultiplyAccumulate_8b907c() {
+ simdgroup_half8x8 const v = simdgroup_half8x8();
+ simdgroup_half8x8 const v_1 = simdgroup_half8x8();
+ simdgroup_half8x8 const v_2 = simdgroup_half8x8();
+ simdgroup_half8x8 v_3 = make_filled_simdgroup_matrix<half, 8, 8>(0.0h);
+ simdgroup_multiply_accumulate(v_3, v, v_1, v_2);
+ simdgroup_half8x8 res = v_3;
+ 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(subgroupMatrixMultiplyAccumulate_8b907c(), (&(*tint_module_vars.prevent_dce)[0u]), ulong(64u), ulong2(0ul), false);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixMultiplyAccumulate/e7fb3c.wgsl.expected.msl b/test/tint/builtins/gen/literal/subgroupMatrixMultiplyAccumulate/e7fb3c.wgsl.expected.msl
index 171ead9..f5551a2 100644
--- a/test/tint/builtins/gen/literal/subgroupMatrixMultiplyAccumulate/e7fb3c.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/subgroupMatrixMultiplyAccumulate/e7fb3c.wgsl.expected.msl
@@ -1,11 +1,33 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:1185 internal compiler error: TINT_UNREACHABLE unhandled: subgroupMatrixMultiplyAccumulate
-********************************************************************
-* 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. *
-********************************************************************
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-tint executable returned error: signal: trace/BPT trap
+struct tint_module_vars_struct {
+ device tint_array<float, 1024>* prevent_dce;
+};
+
+simdgroup_float8x8 subgroupMatrixMultiplyAccumulate_e7fb3c() {
+ simdgroup_float8x8 const v = simdgroup_float8x8();
+ simdgroup_float8x8 const v_1 = simdgroup_float8x8();
+ simdgroup_float8x8 const v_2 = simdgroup_float8x8();
+ simdgroup_float8x8 v_3 = make_filled_simdgroup_matrix<float, 8, 8>(0.0f);
+ simdgroup_multiply_accumulate(v_3, v, v_1, v_2);
+ simdgroup_float8x8 res = v_3;
+ 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(subgroupMatrixMultiplyAccumulate_e7fb3c(), (&(*tint_module_vars.prevent_dce)[0u]), ulong(64u), ulong2(0ul), false);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixMultiplyAccumulate/8b907c.wgsl.expected.msl b/test/tint/builtins/gen/var/subgroupMatrixMultiplyAccumulate/8b907c.wgsl.expected.msl
index 171ead9..beca913 100644
--- a/test/tint/builtins/gen/var/subgroupMatrixMultiplyAccumulate/8b907c.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/subgroupMatrixMultiplyAccumulate/8b907c.wgsl.expected.msl
@@ -1,11 +1,33 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:1185 internal compiler error: TINT_UNREACHABLE unhandled: subgroupMatrixMultiplyAccumulate
-********************************************************************
-* 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. *
-********************************************************************
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-tint executable returned error: signal: trace/BPT trap
+struct tint_module_vars_struct {
+ device tint_array<half, 1024>* prevent_dce;
+};
+
+simdgroup_half8x8 subgroupMatrixMultiplyAccumulate_8b907c() {
+ simdgroup_half8x8 arg_0 = simdgroup_half8x8();
+ simdgroup_half8x8 arg_1 = simdgroup_half8x8();
+ simdgroup_half8x8 arg_2 = simdgroup_half8x8();
+ simdgroup_half8x8 v = make_filled_simdgroup_matrix<half, 8, 8>(0.0h);
+ simdgroup_multiply_accumulate(v, arg_0, arg_1, arg_2);
+ 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(subgroupMatrixMultiplyAccumulate_8b907c(), (&(*tint_module_vars.prevent_dce)[0u]), ulong(64u), ulong2(0ul), false);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixMultiplyAccumulate/e7fb3c.wgsl.expected.msl b/test/tint/builtins/gen/var/subgroupMatrixMultiplyAccumulate/e7fb3c.wgsl.expected.msl
index 171ead9..549a151 100644
--- a/test/tint/builtins/gen/var/subgroupMatrixMultiplyAccumulate/e7fb3c.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/subgroupMatrixMultiplyAccumulate/e7fb3c.wgsl.expected.msl
@@ -1,11 +1,33 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:1185 internal compiler error: TINT_UNREACHABLE unhandled: subgroupMatrixMultiplyAccumulate
-********************************************************************
-* 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. *
-********************************************************************
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-tint executable returned error: signal: trace/BPT trap
+struct tint_module_vars_struct {
+ device tint_array<float, 1024>* prevent_dce;
+};
+
+simdgroup_float8x8 subgroupMatrixMultiplyAccumulate_e7fb3c() {
+ simdgroup_float8x8 arg_0 = simdgroup_float8x8();
+ simdgroup_float8x8 arg_1 = simdgroup_float8x8();
+ simdgroup_float8x8 arg_2 = simdgroup_float8x8();
+ simdgroup_float8x8 v = make_filled_simdgroup_matrix<float, 8, 8>(0.0f);
+ simdgroup_multiply_accumulate(v, arg_0, arg_1, arg_2);
+ 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(subgroupMatrixMultiplyAccumulate_e7fb3c(), (&(*tint_module_vars.prevent_dce)[0u]), ulong(64u), ulong2(0ul), false);
+}
