[spirv-reader][ir] Add `OpGroupNonUniformShuffle` support.
Add support to convert `OpGroupNonUniformShuffle` into a
`subgroupShuffle` instruction.
Fixed: 431032569
Change-Id: I9d3f312f3ad32d4c068306e1ecc03e96ef883e76
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/252514
Reviewed-by: David Neto <dneto@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/spirv/builtin_fn.cc b/src/tint/lang/spirv/builtin_fn.cc
index 590f37f..cfa588c 100644
--- a/src/tint/lang/spirv/builtin_fn.cc
+++ b/src/tint/lang/spirv/builtin_fn.cc
@@ -246,6 +246,8 @@
return "group_non_uniform_quad_broadcast";
case BuiltinFn::kGroupNonUniformQuadSwap:
return "group_non_uniform_quad_swap";
+ case BuiltinFn::kGroupNonUniformShuffle:
+ return "group_non_uniform_shuffle";
}
return "<unknown>";
}
@@ -360,6 +362,7 @@
case BuiltinFn::kOuterProduct:
case BuiltinFn::kGroupNonUniformBroadcast:
case BuiltinFn::kGroupNonUniformBroadcastFirst:
+ case BuiltinFn::kGroupNonUniformShuffle:
case BuiltinFn::kGroupNonUniformQuadBroadcast:
case BuiltinFn::kGroupNonUniformQuadSwap:
break;
diff --git a/src/tint/lang/spirv/builtin_fn.cc.tmpl b/src/tint/lang/spirv/builtin_fn.cc.tmpl
index 69a02e6..4c614f3 100644
--- a/src/tint/lang/spirv/builtin_fn.cc.tmpl
+++ b/src/tint/lang/spirv/builtin_fn.cc.tmpl
@@ -137,6 +137,7 @@
case BuiltinFn::kOuterProduct:
case BuiltinFn::kGroupNonUniformBroadcast:
case BuiltinFn::kGroupNonUniformBroadcastFirst:
+ case BuiltinFn::kGroupNonUniformShuffle:
case BuiltinFn::kGroupNonUniformQuadBroadcast:
case BuiltinFn::kGroupNonUniformQuadSwap:
break;
diff --git a/src/tint/lang/spirv/builtin_fn.h b/src/tint/lang/spirv/builtin_fn.h
index dc690b2..11760fe 100644
--- a/src/tint/lang/spirv/builtin_fn.h
+++ b/src/tint/lang/spirv/builtin_fn.h
@@ -150,6 +150,7 @@
kGroupNonUniformBroadcastFirst,
kGroupNonUniformQuadBroadcast,
kGroupNonUniformQuadSwap,
+ kGroupNonUniformShuffle,
kNone,
};
diff --git a/src/tint/lang/spirv/intrinsic/data.cc b/src/tint/lang/spirv/intrinsic/data.cc
index 50962ea..f8b6299 100644
--- a/src/tint/lang/spirv/intrinsic/data.cc
+++ b/src/tint/lang/spirv/intrinsic/data.cc
@@ -6187,7 +6187,7 @@
{
/* [746] */
/* usage */ core::ParameterUsage::kId,
- /* matcher_indices */ MatcherIndicesIndex(1068),
+ /* matcher_indices */ MatcherIndicesIndex(592),
},
{
/* [747] */
@@ -6202,7 +6202,7 @@
{
/* [749] */
/* usage */ core::ParameterUsage::kId,
- /* matcher_indices */ MatcherIndicesIndex(1068),
+ /* matcher_indices */ MatcherIndicesIndex(592),
},
{
/* [750] */
@@ -6236,48 +6236,48 @@
},
{
/* [756] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(943),
+ /* usage */ core::ParameterUsage::kScope,
+ /* matcher_indices */ MatcherIndicesIndex(592),
},
{
/* [757] */
- /* usage */ core::ParameterUsage::kNone,
+ /* usage */ core::ParameterUsage::kValue,
/* matcher_indices */ MatcherIndicesIndex(8),
},
{
/* [758] */
- /* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(465),
+ /* usage */ core::ParameterUsage::kId,
+ /* matcher_indices */ MatcherIndicesIndex(592),
},
{
/* [759] */
- /* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(21),
+ /* usage */ core::ParameterUsage::kScope,
+ /* matcher_indices */ MatcherIndicesIndex(592),
},
{
/* [760] */
- /* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(474),
+ /* usage */ core::ParameterUsage::kValue,
+ /* matcher_indices */ MatcherIndicesIndex(1003),
},
{
/* [761] */
- /* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(21),
+ /* usage */ core::ParameterUsage::kId,
+ /* matcher_indices */ MatcherIndicesIndex(592),
},
{
/* [762] */
- /* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(483),
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(943),
},
{
/* [763] */
- /* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(21),
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(8),
},
{
/* [764] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(492),
+ /* matcher_indices */ MatcherIndicesIndex(465),
},
{
/* [765] */
@@ -6287,7 +6287,7 @@
{
/* [766] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(501),
+ /* matcher_indices */ MatcherIndicesIndex(474),
},
{
/* [767] */
@@ -6297,7 +6297,7 @@
{
/* [768] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(510),
+ /* matcher_indices */ MatcherIndicesIndex(483),
},
{
/* [769] */
@@ -6307,37 +6307,37 @@
{
/* [770] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(519),
+ /* matcher_indices */ MatcherIndicesIndex(492),
},
{
/* [771] */
/* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(32),
+ /* matcher_indices */ MatcherIndicesIndex(21),
},
{
/* [772] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(528),
+ /* matcher_indices */ MatcherIndicesIndex(501),
},
{
/* [773] */
/* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(32),
+ /* matcher_indices */ MatcherIndicesIndex(21),
},
{
/* [774] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(537),
+ /* matcher_indices */ MatcherIndicesIndex(510),
},
{
/* [775] */
/* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(32),
+ /* matcher_indices */ MatcherIndicesIndex(21),
},
{
/* [776] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(546),
+ /* matcher_indices */ MatcherIndicesIndex(519),
},
{
/* [777] */
@@ -6347,221 +6347,231 @@
{
/* [778] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(681),
+ /* matcher_indices */ MatcherIndicesIndex(528),
},
{
/* [779] */
+ /* usage */ core::ParameterUsage::kLevel,
+ /* matcher_indices */ MatcherIndicesIndex(32),
+ },
+ {
+ /* [780] */
+ /* usage */ core::ParameterUsage::kImage,
+ /* matcher_indices */ MatcherIndicesIndex(537),
+ },
+ {
+ /* [781] */
+ /* usage */ core::ParameterUsage::kLevel,
+ /* matcher_indices */ MatcherIndicesIndex(32),
+ },
+ {
+ /* [782] */
+ /* usage */ core::ParameterUsage::kImage,
+ /* matcher_indices */ MatcherIndicesIndex(546),
+ },
+ {
+ /* [783] */
+ /* usage */ core::ParameterUsage::kLevel,
+ /* matcher_indices */ MatcherIndicesIndex(32),
+ },
+ {
+ /* [784] */
+ /* usage */ core::ParameterUsage::kImage,
+ /* matcher_indices */ MatcherIndicesIndex(681),
+ },
+ {
+ /* [785] */
/* usage */ core::ParameterUsage::kCoords,
/* matcher_indices */ MatcherIndicesIndex(20),
},
{
- /* [780] */
+ /* [786] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(959),
},
{
- /* [781] */
+ /* [787] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(963),
},
{
- /* [782] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(967),
- },
- {
- /* [783] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(8),
- },
- {
- /* [784] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(967),
- },
- {
- /* [785] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(903),
- },
- {
- /* [786] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(1015),
- },
- {
- /* [787] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(909),
- },
- {
/* [788] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(1024),
+ /* matcher_indices */ MatcherIndicesIndex(967),
},
{
/* [789] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(1027),
+ /* matcher_indices */ MatcherIndicesIndex(8),
},
{
/* [790] */
- /* usage */ core::ParameterUsage::kX,
- /* matcher_indices */ MatcherIndicesIndex(8),
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(967),
},
{
/* [791] */
- /* usage */ core::ParameterUsage::kI,
- /* matcher_indices */ MatcherIndicesIndex(971),
- },
- {
- /* [792] */
- /* usage */ core::ParameterUsage::kX,
+ /* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(903),
},
{
+ /* [792] */
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(1015),
+ },
+ {
/* [793] */
- /* usage */ core::ParameterUsage::kI,
- /* matcher_indices */ MatcherIndicesIndex(901),
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(909),
},
{
/* [794] */
- /* usage */ core::ParameterUsage::kX,
- /* matcher_indices */ MatcherIndicesIndex(8),
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(1024),
},
{
/* [795] */
- /* usage */ core::ParameterUsage::kI,
- /* matcher_indices */ MatcherIndicesIndex(975),
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(1027),
},
{
/* [796] */
/* usage */ core::ParameterUsage::kX,
- /* matcher_indices */ MatcherIndicesIndex(1015),
+ /* matcher_indices */ MatcherIndicesIndex(8),
},
{
/* [797] */
/* usage */ core::ParameterUsage::kI,
- /* matcher_indices */ MatcherIndicesIndex(907),
+ /* matcher_indices */ MatcherIndicesIndex(971),
},
{
/* [798] */
- /* usage */ core::ParameterUsage::kScope,
- /* matcher_indices */ MatcherIndicesIndex(592),
+ /* usage */ core::ParameterUsage::kX,
+ /* matcher_indices */ MatcherIndicesIndex(903),
},
{
/* [799] */
- /* usage */ core::ParameterUsage::kValue,
- /* matcher_indices */ MatcherIndicesIndex(8),
+ /* usage */ core::ParameterUsage::kI,
+ /* matcher_indices */ MatcherIndicesIndex(901),
},
{
/* [800] */
- /* usage */ core::ParameterUsage::kScope,
- /* matcher_indices */ MatcherIndicesIndex(592),
+ /* usage */ core::ParameterUsage::kX,
+ /* matcher_indices */ MatcherIndicesIndex(8),
},
{
/* [801] */
- /* usage */ core::ParameterUsage::kValue,
- /* matcher_indices */ MatcherIndicesIndex(1003),
+ /* usage */ core::ParameterUsage::kI,
+ /* matcher_indices */ MatcherIndicesIndex(975),
},
{
/* [802] */
- /* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(357),
+ /* usage */ core::ParameterUsage::kX,
+ /* matcher_indices */ MatcherIndicesIndex(1015),
},
{
/* [803] */
- /* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(366),
+ /* usage */ core::ParameterUsage::kI,
+ /* matcher_indices */ MatcherIndicesIndex(907),
},
{
/* [804] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(375),
+ /* matcher_indices */ MatcherIndicesIndex(357),
},
{
/* [805] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(384),
+ /* matcher_indices */ MatcherIndicesIndex(366),
},
{
/* [806] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(393),
+ /* matcher_indices */ MatcherIndicesIndex(375),
},
{
/* [807] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(402),
+ /* matcher_indices */ MatcherIndicesIndex(384),
},
{
/* [808] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(411),
+ /* matcher_indices */ MatcherIndicesIndex(393),
},
{
/* [809] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(420),
+ /* matcher_indices */ MatcherIndicesIndex(402),
},
{
/* [810] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(429),
+ /* matcher_indices */ MatcherIndicesIndex(411),
},
{
/* [811] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(438),
+ /* matcher_indices */ MatcherIndicesIndex(420),
},
{
/* [812] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(447),
+ /* matcher_indices */ MatcherIndicesIndex(429),
},
{
/* [813] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(456),
+ /* matcher_indices */ MatcherIndicesIndex(438),
},
{
/* [814] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(555),
+ /* matcher_indices */ MatcherIndicesIndex(447),
},
{
/* [815] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(564),
+ /* matcher_indices */ MatcherIndicesIndex(456),
},
{
/* [816] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(1058),
+ /* usage */ core::ParameterUsage::kImage,
+ /* matcher_indices */ MatcherIndicesIndex(555),
},
{
/* [817] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(1060),
+ /* usage */ core::ParameterUsage::kImage,
+ /* matcher_indices */ MatcherIndicesIndex(564),
},
{
/* [818] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(929),
+ /* matcher_indices */ MatcherIndicesIndex(1058),
},
{
/* [819] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(1062),
+ /* matcher_indices */ MatcherIndicesIndex(1060),
},
{
/* [820] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(1064),
+ /* matcher_indices */ MatcherIndicesIndex(929),
},
{
/* [821] */
/* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(1062),
+ },
+ {
+ /* [822] */
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(1064),
+ },
+ {
+ /* [823] */
+ /* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(1066),
},
};
@@ -9166,7 +9176,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 4,
/* templates */ TemplateIndex(72),
- /* parameters */ ParameterIndex(778),
+ /* parameters */ ParameterIndex(784),
/* return_matcher_indices */ MatcherIndicesIndex(1056),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9309,7 +9319,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(76),
- /* parameters */ ParameterIndex(802),
+ /* parameters */ ParameterIndex(804),
/* return_matcher_indices */ MatcherIndicesIndex(8),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9320,7 +9330,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(76),
- /* parameters */ ParameterIndex(803),
+ /* parameters */ ParameterIndex(805),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9331,7 +9341,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(76),
- /* parameters */ ParameterIndex(804),
+ /* parameters */ ParameterIndex(806),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9342,7 +9352,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(76),
- /* parameters */ ParameterIndex(805),
+ /* parameters */ ParameterIndex(807),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9353,7 +9363,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(76),
- /* parameters */ ParameterIndex(806),
+ /* parameters */ ParameterIndex(808),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9364,7 +9374,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(76),
- /* parameters */ ParameterIndex(807),
+ /* parameters */ ParameterIndex(809),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9375,7 +9385,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 5,
/* templates */ TemplateIndex(18),
- /* parameters */ ParameterIndex(808),
+ /* parameters */ ParameterIndex(810),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9386,7 +9396,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(129),
- /* parameters */ ParameterIndex(809),
+ /* parameters */ ParameterIndex(811),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9397,7 +9407,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(129),
- /* parameters */ ParameterIndex(810),
+ /* parameters */ ParameterIndex(812),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9408,7 +9418,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(129),
- /* parameters */ ParameterIndex(811),
+ /* parameters */ ParameterIndex(813),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9419,7 +9429,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(129),
- /* parameters */ ParameterIndex(812),
+ /* parameters */ ParameterIndex(814),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9430,7 +9440,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(80),
- /* parameters */ ParameterIndex(813),
+ /* parameters */ ParameterIndex(815),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9573,7 +9583,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 6,
/* templates */ TemplateIndex(12),
- /* parameters */ ParameterIndex(758),
+ /* parameters */ ParameterIndex(764),
/* return_matcher_indices */ MatcherIndicesIndex(8),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9584,7 +9594,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 6,
/* templates */ TemplateIndex(12),
- /* parameters */ ParameterIndex(760),
+ /* parameters */ ParameterIndex(766),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9595,7 +9605,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 6,
/* templates */ TemplateIndex(12),
- /* parameters */ ParameterIndex(762),
+ /* parameters */ ParameterIndex(768),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9606,7 +9616,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 6,
/* templates */ TemplateIndex(12),
- /* parameters */ ParameterIndex(764),
+ /* parameters */ ParameterIndex(770),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9617,7 +9627,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 6,
/* templates */ TemplateIndex(12),
- /* parameters */ ParameterIndex(766),
+ /* parameters */ ParameterIndex(772),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9628,7 +9638,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 6,
/* templates */ TemplateIndex(12),
- /* parameters */ ParameterIndex(768),
+ /* parameters */ ParameterIndex(774),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9639,7 +9649,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 5,
/* templates */ TemplateIndex(23),
- /* parameters */ ParameterIndex(770),
+ /* parameters */ ParameterIndex(776),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9650,7 +9660,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 5,
/* templates */ TemplateIndex(23),
- /* parameters */ ParameterIndex(772),
+ /* parameters */ ParameterIndex(778),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9661,7 +9671,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 5,
/* templates */ TemplateIndex(23),
- /* parameters */ ParameterIndex(774),
+ /* parameters */ ParameterIndex(780),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9672,7 +9682,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 5,
/* templates */ TemplateIndex(23),
- /* parameters */ ParameterIndex(776),
+ /* parameters */ ParameterIndex(782),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9947,7 +9957,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(144),
- /* parameters */ ParameterIndex(816),
+ /* parameters */ ParameterIndex(818),
/* return_matcher_indices */ MatcherIndicesIndex(8),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9958,7 +9968,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(148),
- /* parameters */ ParameterIndex(816),
+ /* parameters */ ParameterIndex(818),
/* return_matcher_indices */ MatcherIndicesIndex(8),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9969,7 +9979,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 3,
/* templates */ TemplateIndex(145),
- /* parameters */ ParameterIndex(817),
+ /* parameters */ ParameterIndex(819),
/* return_matcher_indices */ MatcherIndicesIndex(21),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9980,7 +9990,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 3,
/* templates */ TemplateIndex(149),
- /* parameters */ ParameterIndex(817),
+ /* parameters */ ParameterIndex(819),
/* return_matcher_indices */ MatcherIndicesIndex(21),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9991,7 +10001,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(91),
- /* parameters */ ParameterIndex(818),
+ /* parameters */ ParameterIndex(820),
/* return_matcher_indices */ MatcherIndicesIndex(32),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10002,7 +10012,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(96),
- /* parameters */ ParameterIndex(818),
+ /* parameters */ ParameterIndex(820),
/* return_matcher_indices */ MatcherIndicesIndex(32),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10244,7 +10254,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(10),
- /* parameters */ ParameterIndex(819),
+ /* parameters */ ParameterIndex(821),
/* return_matcher_indices */ MatcherIndicesIndex(1062),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10255,7 +10265,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(10),
- /* parameters */ ParameterIndex(820),
+ /* parameters */ ParameterIndex(822),
/* return_matcher_indices */ MatcherIndicesIndex(1064),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10266,7 +10276,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(10),
- /* parameters */ ParameterIndex(821),
+ /* parameters */ ParameterIndex(823),
/* return_matcher_indices */ MatcherIndicesIndex(1066),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10464,7 +10474,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(161),
- /* parameters */ ParameterIndex(787),
+ /* parameters */ ParameterIndex(793),
/* return_matcher_indices */ MatcherIndicesIndex(1015),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10486,7 +10496,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(110),
- /* parameters */ ParameterIndex(788),
+ /* parameters */ ParameterIndex(794),
/* return_matcher_indices */ MatcherIndicesIndex(1021),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10607,7 +10617,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(181),
- /* parameters */ ParameterIndex(790),
+ /* parameters */ ParameterIndex(796),
/* return_matcher_indices */ MatcherIndicesIndex(8),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10618,7 +10628,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 3,
/* templates */ TemplateIndex(164),
- /* parameters */ ParameterIndex(792),
+ /* parameters */ ParameterIndex(798),
/* return_matcher_indices */ MatcherIndicesIndex(903),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10629,7 +10639,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 3,
/* templates */ TemplateIndex(167),
- /* parameters */ ParameterIndex(794),
+ /* parameters */ ParameterIndex(800),
/* return_matcher_indices */ MatcherIndicesIndex(8),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10640,7 +10650,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 4,
/* templates */ TemplateIndex(114),
- /* parameters */ ParameterIndex(796),
+ /* parameters */ ParameterIndex(802),
/* return_matcher_indices */ MatcherIndicesIndex(1015),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10706,7 +10716,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(118),
- /* parameters */ ParameterIndex(788),
+ /* parameters */ ParameterIndex(794),
/* return_matcher_indices */ MatcherIndicesIndex(1021),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10728,7 +10738,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 3,
/* templates */ TemplateIndex(119),
- /* parameters */ ParameterIndex(786),
+ /* parameters */ ParameterIndex(792),
/* return_matcher_indices */ MatcherIndicesIndex(1045),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10750,7 +10760,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(170),
- /* parameters */ ParameterIndex(787),
+ /* parameters */ ParameterIndex(793),
/* return_matcher_indices */ MatcherIndicesIndex(1015),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10772,7 +10782,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(173),
- /* parameters */ ParameterIndex(787),
+ /* parameters */ ParameterIndex(793),
/* return_matcher_indices */ MatcherIndicesIndex(1015),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10794,7 +10804,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(176),
- /* parameters */ ParameterIndex(787),
+ /* parameters */ ParameterIndex(793),
/* return_matcher_indices */ MatcherIndicesIndex(1015),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10827,7 +10837,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(33),
- /* parameters */ ParameterIndex(798),
+ /* parameters */ ParameterIndex(756),
/* return_matcher_indices */ MatcherIndicesIndex(8),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10838,7 +10848,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(32),
- /* parameters */ ParameterIndex(800),
+ /* parameters */ ParameterIndex(759),
/* return_matcher_indices */ MatcherIndicesIndex(1003),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10866,17 +10876,39 @@
},
{
/* [290] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMustUse),
+ /* num_parameters */ 3,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 1,
+ /* templates */ TemplateIndex(33),
+ /* parameters */ ParameterIndex(756),
+ /* return_matcher_indices */ MatcherIndicesIndex(8),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [291] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMustUse),
+ /* num_parameters */ 3,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 2,
+ /* templates */ TemplateIndex(32),
+ /* parameters */ ParameterIndex(759),
+ /* return_matcher_indices */ MatcherIndicesIndex(1003),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [292] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(179),
- /* parameters */ ParameterIndex(756),
+ /* parameters */ ParameterIndex(762),
/* return_matcher_indices */ MatcherIndicesIndex(592),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [291] */
+ /* [293] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -10887,7 +10919,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [292] */
+ /* [294] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -10898,47 +10930,25 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [293] */
- /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 1,
- /* num_explicit_templates */ 1,
- /* num_templates */ 7,
- /* templates */ TemplateIndex(0),
- /* parameters */ ParameterIndex(814),
- /* return_matcher_indices */ MatcherIndicesIndex(8),
- /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
- },
- {
- /* [294] */
- /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 1,
- /* num_explicit_templates */ 1,
- /* num_templates */ 7,
- /* templates */ TemplateIndex(0),
- /* parameters */ ParameterIndex(815),
- /* return_matcher_indices */ MatcherIndicesIndex(8),
- /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
- },
- {
/* [295] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 2,
- /* num_explicit_templates */ 0,
- /* num_templates */ 4,
- /* templates */ TemplateIndex(106),
- /* parameters */ ParameterIndex(780),
- /* return_matcher_indices */ MatcherIndicesIndex(955),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 1,
+ /* num_templates */ 7,
+ /* templates */ TemplateIndex(0),
+ /* parameters */ ParameterIndex(816),
+ /* return_matcher_indices */ MatcherIndicesIndex(8),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
/* [296] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 2,
- /* num_explicit_templates */ 0,
- /* num_templates */ 3,
- /* templates */ TemplateIndex(158),
- /* parameters */ ParameterIndex(782),
- /* return_matcher_indices */ MatcherIndicesIndex(967),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 1,
+ /* num_templates */ 7,
+ /* templates */ TemplateIndex(0),
+ /* parameters */ ParameterIndex(817),
+ /* return_matcher_indices */ MatcherIndicesIndex(8),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -10946,10 +10956,10 @@
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
- /* num_templates */ 3,
- /* templates */ TemplateIndex(158),
- /* parameters */ ParameterIndex(784),
- /* return_matcher_indices */ MatcherIndicesIndex(1015),
+ /* num_templates */ 4,
+ /* templates */ TemplateIndex(106),
+ /* parameters */ ParameterIndex(786),
+ /* return_matcher_indices */ MatcherIndicesIndex(955),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -10959,12 +10969,34 @@
/* num_explicit_templates */ 0,
/* num_templates */ 3,
/* templates */ TemplateIndex(158),
+ /* parameters */ ParameterIndex(788),
+ /* return_matcher_indices */ MatcherIndicesIndex(967),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [299] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 2,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 3,
+ /* templates */ TemplateIndex(158),
+ /* parameters */ ParameterIndex(790),
+ /* return_matcher_indices */ MatcherIndicesIndex(1015),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [300] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 2,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 3,
+ /* templates */ TemplateIndex(158),
/* parameters */ ParameterIndex(736),
/* return_matcher_indices */ MatcherIndicesIndex(1015),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [299] */
+ /* [301] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -10975,18 +11007,18 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [300] */
+ /* [302] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 3,
/* templates */ TemplateIndex(158),
- /* parameters */ ParameterIndex(785),
+ /* parameters */ ParameterIndex(791),
/* return_matcher_indices */ MatcherIndicesIndex(963),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [301] */
+ /* [303] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -10997,7 +11029,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [302] */
+ /* [304] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -11017,7 +11049,7 @@
/* [0] */
/* fn array_length[I : u32, A : access](ptr<storage, struct_with_runtime_array, A>, I) -> u32 */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(290),
+ /* overloads */ OverloadIndex(292),
},
{
/* [1] */
@@ -11114,19 +11146,19 @@
/* [14] */
/* fn atomic_i_increment[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, T, writable>, U, U) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(291),
+ /* overloads */ OverloadIndex(293),
},
{
/* [15] */
/* fn atomic_i_decrement[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, T, writable>, U, U) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(291),
+ /* overloads */ OverloadIndex(293),
},
{
/* [16] */
/* fn dot[N : num, T : f32_f16](vec<N, T>, vec<N, T>) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(292),
+ /* overloads */ OverloadIndex(294),
},
{
/* [17] */
@@ -11213,13 +11245,13 @@
/* [22] */
/* fn image_query_levels<Z : iu32>[T : fiu32, D : depth, R : arrayed, S : sampled, F : texel_format, A : access](image: image<T, dim_1d_2d_3d_or_cube, D, R, single_sampled, S, F, A>) -> Z */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(293),
+ /* overloads */ OverloadIndex(295),
},
{
/* [23] */
/* fn image_query_samples<Z : iu32>[T : fiu32, D : depth, R : arrayed, S : sampled, F : texel_format, A : access](image: image<T, dim_2d, D, R, multi_sampled, S, F, A>) -> Z */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(294),
+ /* overloads */ OverloadIndex(296),
},
{
/* [24] */
@@ -11455,19 +11487,19 @@
/* [36] */
/* fn matrix_times_matrix[T : f32_f16, K : num, C : num, R : num](mat<K, R, T>, mat<C, K, T>) -> mat<C, R, T> */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(295),
+ /* overloads */ OverloadIndex(297),
},
{
/* [37] */
/* fn matrix_times_scalar[T : f32_f16, N : num, M : num](mat<N, M, T>, T) -> mat<N, M, T> */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(296),
+ /* overloads */ OverloadIndex(298),
},
{
/* [38] */
/* fn matrix_times_vector[T : f32_f16, N : num, M : num](mat<N, M, T>, vec<N, T>) -> vec<M, T> */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(297),
+ /* overloads */ OverloadIndex(299),
},
{
/* [39] */
@@ -11481,13 +11513,13 @@
/* [40] */
/* fn vector_times_matrix[T : f32_f16, N : num, M : num](vec<N, T>, mat<M, N, T>) -> vec<M, T> */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(298),
+ /* overloads */ OverloadIndex(300),
},
{
/* [41] */
/* fn vector_times_scalar[T : f32_f16, N : num](vec<N, T>, T) -> vec<N, T> */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(299),
+ /* overloads */ OverloadIndex(301),
},
{
/* [42] */
@@ -11844,19 +11876,19 @@
/* [92] */
/* fn outer_product[T : f32_f16, N : num, M : num](vec<N, T>, vec<M, T>) -> mat<M, N, T> */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(300),
+ /* overloads */ OverloadIndex(302),
},
{
/* [93] */
/* fn s_dot(u32, u32, u32) -> i32 */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(301),
+ /* overloads */ OverloadIndex(303),
},
{
/* [94] */
/* fn u_dot(u32, u32, u32) -> u32 */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(302),
+ /* overloads */ OverloadIndex(304),
},
{
/* [95] */
@@ -11885,8 +11917,8 @@
},
{
/* [98] */
- /* fn group_non_uniform_broadcast[T : scalar](scope: u32, e: T, @const id: iu32) -> T */
- /* fn group_non_uniform_broadcast[N : num, T : scalar](scope: u32, e: vec<N, T>, @const id: iu32) -> vec<N, T> */
+ /* fn group_non_uniform_broadcast[T : scalar](scope: u32, e: T, @const id: u32) -> T */
+ /* fn group_non_uniform_broadcast[N : num, T : scalar](scope: u32, e: vec<N, T>, @const id: u32) -> vec<N, T> */
/* num overloads */ 2,
/* overloads */ OverloadIndex(284),
},
@@ -11899,8 +11931,8 @@
},
{
/* [100] */
- /* fn group_non_uniform_quad_broadcast[T : scalar](scope: u32, e: T, @const id: iu32) -> T */
- /* fn group_non_uniform_quad_broadcast[N : num, T : scalar](scope: u32, e: vec<N, T>, @const id: iu32) -> vec<N, T> */
+ /* fn group_non_uniform_quad_broadcast[T : scalar](scope: u32, e: T, @const id: u32) -> T */
+ /* fn group_non_uniform_quad_broadcast[N : num, T : scalar](scope: u32, e: vec<N, T>, @const id: u32) -> vec<N, T> */
/* num overloads */ 2,
/* overloads */ OverloadIndex(284),
},
@@ -11911,6 +11943,13 @@
/* num overloads */ 2,
/* overloads */ OverloadIndex(288),
},
+ {
+ /* [102] */
+ /* fn group_non_uniform_shuffle[T : scalar](scope: u32, value: T, id: u32) -> T */
+ /* fn group_non_uniform_shuffle[N : num, T : scalar](scope: u32, value: vec<N, T>, id: u32) -> vec<N, T> */
+ /* num overloads */ 2,
+ /* overloads */ OverloadIndex(290),
+ },
};
// clang-format on
diff --git a/src/tint/lang/spirv/reader/lower/builtins.cc b/src/tint/lang/spirv/reader/lower/builtins.cc
index c5e95a9..14f87f4 100644
--- a/src/tint/lang/spirv/reader/lower/builtins.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins.cc
@@ -219,13 +219,16 @@
OuterProduct(builtin);
break;
case spirv::BuiltinFn::kGroupNonUniformBroadcast:
- GroupNonUniformBroadcast(builtin, core::BuiltinFn::kSubgroupBroadcast);
+ GroupNonUniformBuiltin(builtin, core::BuiltinFn::kSubgroupBroadcast);
+ break;
+ case spirv::BuiltinFn::kGroupNonUniformShuffle:
+ GroupNonUniformBuiltin(builtin, core::BuiltinFn::kSubgroupShuffle);
break;
case spirv::BuiltinFn::kGroupNonUniformBroadcastFirst:
GroupNonUniformBroadcastFirst(builtin);
break;
case spirv::BuiltinFn::kGroupNonUniformQuadBroadcast:
- GroupNonUniformBroadcast(builtin, core::BuiltinFn::kQuadBroadcast);
+ GroupNonUniformBuiltin(builtin, core::BuiltinFn::kQuadBroadcast);
break;
case spirv::BuiltinFn::kGroupNonUniformQuadSwap:
GroupNonUniformQuadSwap(builtin);
@@ -316,7 +319,7 @@
call->Destroy();
}
- void GroupNonUniformBroadcast(spirv::ir::BuiltinCall* call, core::BuiltinFn fn) {
+ void GroupNonUniformBuiltin(spirv::ir::BuiltinCall* call, core::BuiltinFn fn) {
auto* value = call->Args()[1];
auto* id = call->Args()[2];
diff --git a/src/tint/lang/spirv/reader/lower/builtins_test.cc b/src/tint/lang/spirv/reader/lower/builtins_test.cc
index 658b493..60ebe0e 100644
--- a/src/tint/lang/spirv/reader/lower/builtins_test.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins_test.cc
@@ -9758,5 +9758,138 @@
EXPECT_EQ(expect, str());
}
+TEST_F(SpirvReader_BuiltinsTest, NonUniformShuffle_Constant_BoolScalar) {
+ auto* ep = b.ComputeFunction("main");
+
+ b.Append(ep->Block(), [&] { //
+ b.Call<spirv::ir::BuiltinCall>(ty.bool_(), spirv::BuiltinFn::kGroupNonUniformShuffle, 3_u,
+ true, 1_u);
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:bool = spirv.group_non_uniform_shuffle 3u, true, 1u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ Run(Builtins);
+
+ auto expect = R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = convert true
+ %3:u32 = subgroupShuffle %2, 1u
+ %4:bool = convert %3
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, NonUniformShuffle_Constant_BoolVector) {
+ auto* ep = b.ComputeFunction("main");
+
+ b.Append(ep->Block(), [&] { //
+ b.Call<spirv::ir::BuiltinCall>(ty.vec3(ty.bool_()),
+ spirv::BuiltinFn::kGroupNonUniformShuffle, 3_u,
+ b.Composite(ty.vec3(ty.bool_()), true, false, true), 1_u);
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec3<bool> = spirv.group_non_uniform_shuffle 3u, vec3<bool>(true, false, true), 1u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ Run(Builtins);
+
+ auto expect = R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec3<u32> = convert vec3<bool>(true, false, true)
+ %3:vec3<u32> = subgroupShuffle %2, 1u
+ %4:vec3<bool> = convert %3
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, NonUniformShuffle_Constant_NumericScalar) {
+ auto* ep = b.ComputeFunction("main");
+
+ b.Append(ep->Block(), [&] { //
+ b.Call<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kGroupNonUniformShuffle, 3_u,
+ 2_u, 1_u);
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = spirv.group_non_uniform_shuffle 3u, 2u, 1u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ Run(Builtins);
+
+ auto expect = R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = subgroupShuffle 2u, 1u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, NonUniformShuffle_Constant_NumericVector) {
+ auto* ep = b.ComputeFunction("main");
+
+ b.Append(ep->Block(), [&] { //
+ b.Call<spirv::ir::BuiltinCall>(ty.vec3<u32>(), spirv::BuiltinFn::kGroupNonUniformShuffle,
+ 3_u, b.Composite(ty.vec3<u32>(), 2_u, 3_u, 2_u), 1_u);
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec3<u32> = spirv.group_non_uniform_shuffle 3u, vec3<u32>(2u, 3u, 2u), 1u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ Run(Builtins);
+
+ auto expect = R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec3<u32> = subgroupShuffle vec3<u32>(2u, 3u, 2u), 1u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
} // namespace
} // namespace tint::spirv::reader::lower
diff --git a/src/tint/lang/spirv/reader/parser/builtin_test.cc b/src/tint/lang/spirv/reader/parser/builtin_test.cc
index 72e8e31..bb4fd5b 100644
--- a/src/tint/lang/spirv/reader/parser/builtin_test.cc
+++ b/src/tint/lang/spirv/reader/parser/builtin_test.cc
@@ -2357,5 +2357,139 @@
SPV_ENV_VULKAN_1_1);
}
+TEST_F(SpirvParserTest, NonUniformShuffle_Constant_BoolScalar) {
+ EXPECT_IR_SPV(R"(
+ OpCapability Shader
+ OpCapability GroupNonUniformShuffle
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ %uint = OpTypeInt 32 0
+ %uint_1 = OpConstant %uint 1
+ %uint_3 = OpConstant %uint 3
+ %bool = OpTypeBool
+ %true = OpConstantTrue %bool
+ %void = OpTypeVoid
+ %23 = OpTypeFunction %void
+ %main = OpFunction %void None %23
+ %24 = OpLabel
+ %8 = OpGroupNonUniformShuffle %bool %uint_3 %true %uint_1
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:bool = spirv.group_non_uniform_shuffle 3u, true, 1u
+ ret
+ }
+}
+)",
+ SPV_ENV_VULKAN_1_1);
+}
+
+TEST_F(SpirvParserTest, NonUniformShuffle_Constant_BoolVector) {
+ EXPECT_IR_SPV(R"(
+ OpCapability Shader
+ OpCapability GroupNonUniformShuffle
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ %bool = OpTypeBool
+ %true = OpConstantTrue %bool
+ %false = OpConstantFalse %bool
+ %uint = OpTypeInt 32 0
+ %uint_1 = OpConstant %uint 1
+ %uint_3 = OpConstant %uint 3
+ %v3bool = OpTypeVector %bool 3
+ %12 = OpConstantComposite %v3bool %true %false %true
+ %void = OpTypeVoid
+ %23 = OpTypeFunction %void
+ %main = OpFunction %void None %23
+ %24 = OpLabel
+ %8 = OpGroupNonUniformShuffle %v3bool %uint_3 %12 %uint_1
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec3<bool> = spirv.group_non_uniform_shuffle 3u, vec3<bool>(true, false, true), 1u
+ ret
+ }
+}
+)",
+ SPV_ENV_VULKAN_1_1);
+}
+
+TEST_F(SpirvParserTest, NonUniformShuffle_Constant_NumericScalar) {
+ EXPECT_IR_SPV(R"(
+ OpCapability Shader
+ OpCapability GroupNonUniformShuffle
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ %uint = OpTypeInt 32 0
+ %uint_1 = OpConstant %uint 1
+ %uint_3 = OpConstant %uint 3
+ %v3uint = OpTypeVector %uint 3
+ %bool = OpTypeBool
+ %true = OpConstantTrue %bool
+ %void = OpTypeVoid
+ %23 = OpTypeFunction %void
+ %main = OpFunction %void None %23
+ %24 = OpLabel
+ %8 = OpGroupNonUniformShuffle %uint %uint_3 %uint_3 %uint_1
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = spirv.group_non_uniform_shuffle 3u, 3u, 1u
+ ret
+ }
+}
+)",
+ SPV_ENV_VULKAN_1_1);
+}
+
+TEST_F(SpirvParserTest, NonUniformShuffle_Constant_NumericVector) {
+ EXPECT_IR_SPV(R"(
+ OpCapability Shader
+ OpCapability GroupNonUniformShuffle
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ %uint = OpTypeInt 32 0
+ %uint_1 = OpConstant %uint 1
+ %uint_3 = OpConstant %uint 3
+ %v3uint = OpTypeVector %uint 3
+ %12 = OpConstantComposite %v3uint %uint_1 %uint_3 %uint_1
+ %bool = OpTypeBool
+ %true = OpConstantTrue %bool
+ %void = OpTypeVoid
+ %23 = OpTypeFunction %void
+ %main = OpFunction %void None %23
+ %24 = OpLabel
+ %8 = OpGroupNonUniformShuffle %v3uint %uint_3 %12 %uint_1
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec3<u32> = spirv.group_non_uniform_shuffle 3u, vec3<u32>(1u, 3u, 1u), 1u
+ ret
+ }
+}
+)",
+ SPV_ENV_VULKAN_1_1);
+}
+
} // namespace
} // namespace tint::spirv::reader
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index 7abce8a..36ad945 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -2228,6 +2228,9 @@
case spv::Op::OpGroupNonUniformBroadcast:
EmitSubgroupBuiltin(inst, spirv::BuiltinFn::kGroupNonUniformBroadcast);
break;
+ case spv::Op::OpGroupNonUniformShuffle:
+ EmitSubgroupBuiltin(inst, spirv::BuiltinFn::kGroupNonUniformShuffle);
+ break;
case spv::Op::OpGroupNonUniformQuadBroadcast:
EmitSubgroupBuiltin(inst, spirv::BuiltinFn::kGroupNonUniformQuadBroadcast);
break;
diff --git a/src/tint/lang/spirv/spirv.def b/src/tint/lang/spirv/spirv.def
index 2f3b2f4..9a37634 100644
--- a/src/tint/lang/spirv/spirv.def
+++ b/src/tint/lang/spirv/spirv.def
@@ -1817,9 +1817,9 @@
// Subgroups
////////////////////////////////////////////////////////////////////////////////
@must_use @stage("fragment", "compute") implicit(T: scalar)
-fn group_non_uniform_broadcast(scope: u32, e: T, @const id: iu32) -> T
+fn group_non_uniform_broadcast(scope: u32, e: T, @const id: u32) -> T
@must_use @stage("fragment", "compute") implicit(N: num, T: scalar)
-fn group_non_uniform_broadcast(scope: u32, e: vec<N, T>, @const id: iu32) -> vec<N, T>
+fn group_non_uniform_broadcast(scope: u32, e: vec<N, T>, @const id: u32) -> vec<N, T>
@must_use @stage("fragment", "compute") implicit(T: scalar)
fn group_non_uniform_broadcast_first(scope: u32, value: T) -> T
@@ -1827,12 +1827,17 @@
fn group_non_uniform_broadcast_first(scope: u32, value: vec<N, T>) -> vec<N, T>
@must_use @stage("fragment", "compute") implicit(T: scalar)
-fn group_non_uniform_quad_broadcast(scope: u32, e: T, @const id: iu32) -> T
+fn group_non_uniform_quad_broadcast(scope: u32, e: T, @const id: u32) -> T
@must_use @stage("fragment", "compute") implicit(N: num, T: scalar)
-fn group_non_uniform_quad_broadcast(scope: u32, e: vec<N, T>, @const id: iu32) -> vec<N, T>
+fn group_non_uniform_quad_broadcast(scope: u32, e: vec<N, T>, @const id: u32) -> vec<N, T>
@must_use @stage("fragment", "compute") implicit(T: scalar)
fn group_non_uniform_quad_swap(scope: u32, e: T, @const dir: u32) -> T
@must_use @stage("fragment", "compute") implicit(N: num, T: scalar)
fn group_non_uniform_quad_swap(scope:u32, e: vec<N, T>, @const dir: u32) -> vec<N, T>
+@must_use @stage("fragment", "compute") implicit(T: scalar)
+fn group_non_uniform_shuffle(scope: u32, value: T, id: u32) -> T
+@must_use @stage("fragment", "compute") implicit(N: num, T: scalar)
+fn group_non_uniform_shuffle(scope: u32, value: vec<N, T>, id: u32) -> vec<N, T>
+
diff --git a/src/tint/lang/spirv/writer/printer/printer.cc b/src/tint/lang/spirv/writer/printer/printer.cc
index 31917b1..782d4c5 100644
--- a/src/tint/lang/spirv/writer/printer/printer.cc
+++ b/src/tint/lang/spirv/writer/printer/printer.cc
@@ -1697,6 +1697,9 @@
case BuiltinFn::kGroupNonUniformQuadSwap:
op = spv::Op::OpGroupNonUniformQuadSwap;
break;
+ case BuiltinFn::kGroupNonUniformShuffle:
+ op = spv::Op::OpGroupNonUniformShuffle;
+ break;
case spirv::BuiltinFn::kNone:
TINT_ICE() << "undefined spirv ir function";
}
