[spirv-reader][ir] Support `OpGroupNonUniformShuffleDown
Add capability to convert a `OpGroupNonUniformShuffleDown into a
`subgroupShuffleDown instruction.
Fixed: 431031975
Change-Id: I8f2d9e8dbd93f1661e19318371d9e07799396dc6
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/252634
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 ef0216c..938cec2 100644
--- a/src/tint/lang/spirv/builtin_fn.cc
+++ b/src/tint/lang/spirv/builtin_fn.cc
@@ -250,6 +250,8 @@
return "group_non_uniform_shuffle";
case BuiltinFn::kGroupNonUniformShuffleXor:
return "group_non_uniform_shuffle_xor";
+ case BuiltinFn::kGroupNonUniformShuffleDown:
+ return "group_non_uniform_shuffle_down";
}
return "<unknown>";
}
@@ -366,6 +368,7 @@
case BuiltinFn::kGroupNonUniformBroadcastFirst:
case BuiltinFn::kGroupNonUniformShuffle:
case BuiltinFn::kGroupNonUniformShuffleXor:
+ case BuiltinFn::kGroupNonUniformShuffleDown:
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 fd5c64c..da2d8b9 100644
--- a/src/tint/lang/spirv/builtin_fn.cc.tmpl
+++ b/src/tint/lang/spirv/builtin_fn.cc.tmpl
@@ -139,6 +139,7 @@
case BuiltinFn::kGroupNonUniformBroadcastFirst:
case BuiltinFn::kGroupNonUniformShuffle:
case BuiltinFn::kGroupNonUniformShuffleXor:
+ case BuiltinFn::kGroupNonUniformShuffleDown:
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 4d6d656..078e996 100644
--- a/src/tint/lang/spirv/builtin_fn.h
+++ b/src/tint/lang/spirv/builtin_fn.h
@@ -152,6 +152,7 @@
kGroupNonUniformQuadSwap,
kGroupNonUniformShuffle,
kGroupNonUniformShuffleXor,
+ kGroupNonUniformShuffleDown,
kNone,
};
diff --git a/src/tint/lang/spirv/intrinsic/data.cc b/src/tint/lang/spirv/intrinsic/data.cc
index d38660c..c606332 100644
--- a/src/tint/lang/spirv/intrinsic/data.cc
+++ b/src/tint/lang/spirv/intrinsic/data.cc
@@ -6296,48 +6296,48 @@
},
{
/* [768] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(943),
+ /* usage */ core::ParameterUsage::kScope,
+ /* matcher_indices */ MatcherIndicesIndex(592),
},
{
/* [769] */
- /* usage */ core::ParameterUsage::kNone,
+ /* usage */ core::ParameterUsage::kValue,
/* matcher_indices */ MatcherIndicesIndex(8),
},
{
/* [770] */
- /* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(465),
+ /* usage */ core::ParameterUsage::kDelta,
+ /* matcher_indices */ MatcherIndicesIndex(592),
},
{
/* [771] */
- /* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(21),
+ /* usage */ core::ParameterUsage::kScope,
+ /* matcher_indices */ MatcherIndicesIndex(592),
},
{
/* [772] */
- /* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(474),
+ /* usage */ core::ParameterUsage::kValue,
+ /* matcher_indices */ MatcherIndicesIndex(1003),
},
{
/* [773] */
- /* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(21),
+ /* usage */ core::ParameterUsage::kDelta,
+ /* matcher_indices */ MatcherIndicesIndex(592),
},
{
/* [774] */
- /* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(483),
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(943),
},
{
/* [775] */
- /* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(21),
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(8),
},
{
/* [776] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(492),
+ /* matcher_indices */ MatcherIndicesIndex(465),
},
{
/* [777] */
@@ -6347,7 +6347,7 @@
{
/* [778] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(501),
+ /* matcher_indices */ MatcherIndicesIndex(474),
},
{
/* [779] */
@@ -6357,7 +6357,7 @@
{
/* [780] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(510),
+ /* matcher_indices */ MatcherIndicesIndex(483),
},
{
/* [781] */
@@ -6367,37 +6367,37 @@
{
/* [782] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(519),
+ /* matcher_indices */ MatcherIndicesIndex(492),
},
{
/* [783] */
/* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(32),
+ /* matcher_indices */ MatcherIndicesIndex(21),
},
{
/* [784] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(528),
+ /* matcher_indices */ MatcherIndicesIndex(501),
},
{
/* [785] */
/* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(32),
+ /* matcher_indices */ MatcherIndicesIndex(21),
},
{
/* [786] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(537),
+ /* matcher_indices */ MatcherIndicesIndex(510),
},
{
/* [787] */
/* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(32),
+ /* matcher_indices */ MatcherIndicesIndex(21),
},
{
/* [788] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(546),
+ /* matcher_indices */ MatcherIndicesIndex(519),
},
{
/* [789] */
@@ -6407,200 +6407,230 @@
{
/* [790] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(681),
+ /* matcher_indices */ MatcherIndicesIndex(528),
},
{
/* [791] */
+ /* usage */ core::ParameterUsage::kLevel,
+ /* matcher_indices */ MatcherIndicesIndex(32),
+ },
+ {
+ /* [792] */
+ /* usage */ core::ParameterUsage::kImage,
+ /* matcher_indices */ MatcherIndicesIndex(537),
+ },
+ {
+ /* [793] */
+ /* usage */ core::ParameterUsage::kLevel,
+ /* matcher_indices */ MatcherIndicesIndex(32),
+ },
+ {
+ /* [794] */
+ /* usage */ core::ParameterUsage::kImage,
+ /* matcher_indices */ MatcherIndicesIndex(546),
+ },
+ {
+ /* [795] */
+ /* usage */ core::ParameterUsage::kLevel,
+ /* matcher_indices */ MatcherIndicesIndex(32),
+ },
+ {
+ /* [796] */
+ /* usage */ core::ParameterUsage::kImage,
+ /* matcher_indices */ MatcherIndicesIndex(681),
+ },
+ {
+ /* [797] */
/* usage */ core::ParameterUsage::kCoords,
/* matcher_indices */ MatcherIndicesIndex(20),
},
{
- /* [792] */
+ /* [798] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(959),
},
{
- /* [793] */
+ /* [799] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(963),
},
{
- /* [794] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(967),
- },
- {
- /* [795] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(8),
- },
- {
- /* [796] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(967),
- },
- {
- /* [797] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(903),
- },
- {
- /* [798] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(1015),
- },
- {
- /* [799] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(909),
- },
- {
/* [800] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(1024),
+ /* matcher_indices */ MatcherIndicesIndex(967),
},
{
/* [801] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(1027),
+ /* matcher_indices */ MatcherIndicesIndex(8),
},
{
/* [802] */
- /* usage */ core::ParameterUsage::kX,
- /* matcher_indices */ MatcherIndicesIndex(8),
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(967),
},
{
/* [803] */
- /* usage */ core::ParameterUsage::kI,
- /* matcher_indices */ MatcherIndicesIndex(971),
- },
- {
- /* [804] */
- /* usage */ core::ParameterUsage::kX,
+ /* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(903),
},
{
+ /* [804] */
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(1015),
+ },
+ {
/* [805] */
- /* usage */ core::ParameterUsage::kI,
- /* matcher_indices */ MatcherIndicesIndex(901),
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(909),
},
{
/* [806] */
- /* usage */ core::ParameterUsage::kX,
- /* matcher_indices */ MatcherIndicesIndex(8),
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(1024),
},
{
/* [807] */
- /* usage */ core::ParameterUsage::kI,
- /* matcher_indices */ MatcherIndicesIndex(975),
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(1027),
},
{
/* [808] */
/* usage */ core::ParameterUsage::kX,
- /* matcher_indices */ MatcherIndicesIndex(1015),
+ /* matcher_indices */ MatcherIndicesIndex(8),
},
{
/* [809] */
/* usage */ core::ParameterUsage::kI,
- /* matcher_indices */ MatcherIndicesIndex(907),
+ /* matcher_indices */ MatcherIndicesIndex(971),
},
{
/* [810] */
- /* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(357),
+ /* usage */ core::ParameterUsage::kX,
+ /* matcher_indices */ MatcherIndicesIndex(903),
},
{
/* [811] */
- /* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(366),
+ /* usage */ core::ParameterUsage::kI,
+ /* matcher_indices */ MatcherIndicesIndex(901),
},
{
/* [812] */
- /* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(375),
+ /* usage */ core::ParameterUsage::kX,
+ /* matcher_indices */ MatcherIndicesIndex(8),
},
{
/* [813] */
- /* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(384),
+ /* usage */ core::ParameterUsage::kI,
+ /* matcher_indices */ MatcherIndicesIndex(975),
},
{
/* [814] */
- /* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(393),
+ /* usage */ core::ParameterUsage::kX,
+ /* matcher_indices */ MatcherIndicesIndex(1015),
},
{
/* [815] */
- /* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(402),
+ /* usage */ core::ParameterUsage::kI,
+ /* matcher_indices */ MatcherIndicesIndex(907),
},
{
/* [816] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(411),
+ /* matcher_indices */ MatcherIndicesIndex(357),
},
{
/* [817] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(420),
+ /* matcher_indices */ MatcherIndicesIndex(366),
},
{
/* [818] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(429),
+ /* matcher_indices */ MatcherIndicesIndex(375),
},
{
/* [819] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(438),
+ /* matcher_indices */ MatcherIndicesIndex(384),
},
{
/* [820] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(447),
+ /* matcher_indices */ MatcherIndicesIndex(393),
},
{
/* [821] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(456),
+ /* matcher_indices */ MatcherIndicesIndex(402),
},
{
/* [822] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(555),
+ /* matcher_indices */ MatcherIndicesIndex(411),
},
{
/* [823] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(564),
+ /* matcher_indices */ MatcherIndicesIndex(420),
},
{
/* [824] */
+ /* usage */ core::ParameterUsage::kImage,
+ /* matcher_indices */ MatcherIndicesIndex(429),
+ },
+ {
+ /* [825] */
+ /* usage */ core::ParameterUsage::kImage,
+ /* matcher_indices */ MatcherIndicesIndex(438),
+ },
+ {
+ /* [826] */
+ /* usage */ core::ParameterUsage::kImage,
+ /* matcher_indices */ MatcherIndicesIndex(447),
+ },
+ {
+ /* [827] */
+ /* usage */ core::ParameterUsage::kImage,
+ /* matcher_indices */ MatcherIndicesIndex(456),
+ },
+ {
+ /* [828] */
+ /* usage */ core::ParameterUsage::kImage,
+ /* matcher_indices */ MatcherIndicesIndex(555),
+ },
+ {
+ /* [829] */
+ /* usage */ core::ParameterUsage::kImage,
+ /* matcher_indices */ MatcherIndicesIndex(564),
+ },
+ {
+ /* [830] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(1058),
},
{
- /* [825] */
+ /* [831] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(1060),
},
{
- /* [826] */
+ /* [832] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(929),
},
{
- /* [827] */
+ /* [833] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(1062),
},
{
- /* [828] */
+ /* [834] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(1064),
},
{
- /* [829] */
+ /* [835] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(1066),
},
@@ -9206,7 +9236,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 4,
/* templates */ TemplateIndex(72),
- /* parameters */ ParameterIndex(790),
+ /* parameters */ ParameterIndex(796),
/* return_matcher_indices */ MatcherIndicesIndex(1056),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9349,7 +9379,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(76),
- /* parameters */ ParameterIndex(810),
+ /* parameters */ ParameterIndex(816),
/* return_matcher_indices */ MatcherIndicesIndex(8),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9360,7 +9390,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(76),
- /* parameters */ ParameterIndex(811),
+ /* parameters */ ParameterIndex(817),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9371,7 +9401,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(76),
- /* parameters */ ParameterIndex(812),
+ /* parameters */ ParameterIndex(818),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9382,7 +9412,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(76),
- /* parameters */ ParameterIndex(813),
+ /* parameters */ ParameterIndex(819),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9393,7 +9423,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(76),
- /* parameters */ ParameterIndex(814),
+ /* parameters */ ParameterIndex(820),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9404,7 +9434,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(76),
- /* parameters */ ParameterIndex(815),
+ /* parameters */ ParameterIndex(821),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9415,7 +9445,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 5,
/* templates */ TemplateIndex(18),
- /* parameters */ ParameterIndex(816),
+ /* parameters */ ParameterIndex(822),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9426,7 +9456,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(129),
- /* parameters */ ParameterIndex(817),
+ /* parameters */ ParameterIndex(823),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9437,7 +9467,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(129),
- /* parameters */ ParameterIndex(818),
+ /* parameters */ ParameterIndex(824),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9448,7 +9478,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(129),
- /* parameters */ ParameterIndex(819),
+ /* parameters */ ParameterIndex(825),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9459,7 +9489,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(129),
- /* parameters */ ParameterIndex(820),
+ /* parameters */ ParameterIndex(826),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9470,7 +9500,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(80),
- /* parameters */ ParameterIndex(821),
+ /* parameters */ ParameterIndex(827),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9613,7 +9643,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 6,
/* templates */ TemplateIndex(12),
- /* parameters */ ParameterIndex(770),
+ /* parameters */ ParameterIndex(776),
/* return_matcher_indices */ MatcherIndicesIndex(8),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9624,7 +9654,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 6,
/* templates */ TemplateIndex(12),
- /* parameters */ ParameterIndex(772),
+ /* parameters */ ParameterIndex(778),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9635,7 +9665,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 6,
/* templates */ TemplateIndex(12),
- /* parameters */ ParameterIndex(774),
+ /* parameters */ ParameterIndex(780),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9646,7 +9676,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 6,
/* templates */ TemplateIndex(12),
- /* parameters */ ParameterIndex(776),
+ /* parameters */ ParameterIndex(782),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9657,7 +9687,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 6,
/* templates */ TemplateIndex(12),
- /* parameters */ ParameterIndex(778),
+ /* parameters */ ParameterIndex(784),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9668,7 +9698,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 6,
/* templates */ TemplateIndex(12),
- /* parameters */ ParameterIndex(780),
+ /* parameters */ ParameterIndex(786),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9679,7 +9709,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 5,
/* templates */ TemplateIndex(23),
- /* parameters */ ParameterIndex(782),
+ /* parameters */ ParameterIndex(788),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9690,7 +9720,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 5,
/* templates */ TemplateIndex(23),
- /* parameters */ ParameterIndex(784),
+ /* parameters */ ParameterIndex(790),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9701,7 +9731,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 5,
/* templates */ TemplateIndex(23),
- /* parameters */ ParameterIndex(786),
+ /* parameters */ ParameterIndex(792),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9712,7 +9742,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 5,
/* templates */ TemplateIndex(23),
- /* parameters */ ParameterIndex(788),
+ /* parameters */ ParameterIndex(794),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9987,7 +10017,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(144),
- /* parameters */ ParameterIndex(824),
+ /* parameters */ ParameterIndex(830),
/* return_matcher_indices */ MatcherIndicesIndex(8),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9998,7 +10028,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(148),
- /* parameters */ ParameterIndex(824),
+ /* parameters */ ParameterIndex(830),
/* return_matcher_indices */ MatcherIndicesIndex(8),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10009,7 +10039,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 3,
/* templates */ TemplateIndex(145),
- /* parameters */ ParameterIndex(825),
+ /* parameters */ ParameterIndex(831),
/* return_matcher_indices */ MatcherIndicesIndex(21),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10020,7 +10050,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 3,
/* templates */ TemplateIndex(149),
- /* parameters */ ParameterIndex(825),
+ /* parameters */ ParameterIndex(831),
/* return_matcher_indices */ MatcherIndicesIndex(21),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10031,7 +10061,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(91),
- /* parameters */ ParameterIndex(826),
+ /* parameters */ ParameterIndex(832),
/* return_matcher_indices */ MatcherIndicesIndex(32),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10042,7 +10072,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(96),
- /* parameters */ ParameterIndex(826),
+ /* parameters */ ParameterIndex(832),
/* return_matcher_indices */ MatcherIndicesIndex(32),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10284,7 +10314,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(10),
- /* parameters */ ParameterIndex(827),
+ /* parameters */ ParameterIndex(833),
/* return_matcher_indices */ MatcherIndicesIndex(1062),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10295,7 +10325,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(10),
- /* parameters */ ParameterIndex(828),
+ /* parameters */ ParameterIndex(834),
/* return_matcher_indices */ MatcherIndicesIndex(1064),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10306,7 +10336,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(10),
- /* parameters */ ParameterIndex(829),
+ /* parameters */ ParameterIndex(835),
/* return_matcher_indices */ MatcherIndicesIndex(1066),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10504,7 +10534,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(161),
- /* parameters */ ParameterIndex(799),
+ /* parameters */ ParameterIndex(805),
/* return_matcher_indices */ MatcherIndicesIndex(1015),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10526,7 +10556,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(110),
- /* parameters */ ParameterIndex(800),
+ /* parameters */ ParameterIndex(806),
/* return_matcher_indices */ MatcherIndicesIndex(1021),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10647,7 +10677,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(181),
- /* parameters */ ParameterIndex(802),
+ /* parameters */ ParameterIndex(808),
/* return_matcher_indices */ MatcherIndicesIndex(8),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10658,7 +10688,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 3,
/* templates */ TemplateIndex(164),
- /* parameters */ ParameterIndex(804),
+ /* parameters */ ParameterIndex(810),
/* return_matcher_indices */ MatcherIndicesIndex(903),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10669,7 +10699,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 3,
/* templates */ TemplateIndex(167),
- /* parameters */ ParameterIndex(806),
+ /* parameters */ ParameterIndex(812),
/* return_matcher_indices */ MatcherIndicesIndex(8),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10680,7 +10710,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 4,
/* templates */ TemplateIndex(114),
- /* parameters */ ParameterIndex(808),
+ /* parameters */ ParameterIndex(814),
/* return_matcher_indices */ MatcherIndicesIndex(1015),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10746,7 +10776,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(118),
- /* parameters */ ParameterIndex(800),
+ /* parameters */ ParameterIndex(806),
/* return_matcher_indices */ MatcherIndicesIndex(1021),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10768,7 +10798,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 3,
/* templates */ TemplateIndex(119),
- /* parameters */ ParameterIndex(798),
+ /* parameters */ ParameterIndex(804),
/* return_matcher_indices */ MatcherIndicesIndex(1045),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10790,7 +10820,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(170),
- /* parameters */ ParameterIndex(799),
+ /* parameters */ ParameterIndex(805),
/* return_matcher_indices */ MatcherIndicesIndex(1015),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10812,7 +10842,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(173),
- /* parameters */ ParameterIndex(799),
+ /* parameters */ ParameterIndex(805),
/* return_matcher_indices */ MatcherIndicesIndex(1015),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10834,7 +10864,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(176),
- /* parameters */ ParameterIndex(799),
+ /* parameters */ ParameterIndex(805),
/* return_matcher_indices */ MatcherIndicesIndex(1015),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10950,17 +10980,39 @@
},
{
/* [294] */
+ /* 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(768),
+ /* return_matcher_indices */ MatcherIndicesIndex(8),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [295] */
+ /* 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(771),
+ /* return_matcher_indices */ MatcherIndicesIndex(1003),
+ /* 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 */ 2,
/* templates */ TemplateIndex(179),
- /* parameters */ ParameterIndex(768),
+ /* parameters */ ParameterIndex(774),
/* return_matcher_indices */ MatcherIndicesIndex(592),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [295] */
+ /* [297] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -10971,7 +11023,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [296] */
+ /* [298] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -10982,47 +11034,25 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [297] */
- /* 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(822),
- /* return_matcher_indices */ MatcherIndicesIndex(8),
- /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
- },
- {
- /* [298] */
- /* 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(823),
- /* return_matcher_indices */ MatcherIndicesIndex(8),
- /* 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 */ 4,
- /* templates */ TemplateIndex(106),
- /* parameters */ ParameterIndex(792),
- /* return_matcher_indices */ MatcherIndicesIndex(955),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 1,
+ /* num_templates */ 7,
+ /* templates */ TemplateIndex(0),
+ /* parameters */ ParameterIndex(828),
+ /* return_matcher_indices */ MatcherIndicesIndex(8),
/* 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(794),
- /* return_matcher_indices */ MatcherIndicesIndex(967),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 1,
+ /* num_templates */ 7,
+ /* templates */ TemplateIndex(0),
+ /* parameters */ ParameterIndex(829),
+ /* return_matcher_indices */ MatcherIndicesIndex(8),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -11030,10 +11060,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(796),
- /* return_matcher_indices */ MatcherIndicesIndex(1015),
+ /* num_templates */ 4,
+ /* templates */ TemplateIndex(106),
+ /* parameters */ ParameterIndex(798),
+ /* return_matcher_indices */ MatcherIndicesIndex(955),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -11043,12 +11073,34 @@
/* num_explicit_templates */ 0,
/* num_templates */ 3,
/* templates */ TemplateIndex(158),
+ /* parameters */ ParameterIndex(800),
+ /* return_matcher_indices */ MatcherIndicesIndex(967),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [303] */
+ /* 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(802),
+ /* return_matcher_indices */ MatcherIndicesIndex(1015),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [304] */
+ /* 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 */),
},
{
- /* [303] */
+ /* [305] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -11059,18 +11111,18 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [304] */
+ /* [306] */
/* 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(797),
+ /* parameters */ ParameterIndex(803),
/* return_matcher_indices */ MatcherIndicesIndex(963),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [305] */
+ /* [307] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -11081,7 +11133,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [306] */
+ /* [308] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -11101,7 +11153,7 @@
/* [0] */
/* fn array_length[I : u32, A : access](ptr<storage, struct_with_runtime_array, A>, I) -> u32 */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(294),
+ /* overloads */ OverloadIndex(296),
},
{
/* [1] */
@@ -11198,19 +11250,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(295),
+ /* overloads */ OverloadIndex(297),
},
{
/* [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(295),
+ /* overloads */ OverloadIndex(297),
},
{
/* [16] */
/* fn dot[N : num, T : f32_f16](vec<N, T>, vec<N, T>) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(296),
+ /* overloads */ OverloadIndex(298),
},
{
/* [17] */
@@ -11297,13 +11349,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(297),
+ /* overloads */ OverloadIndex(299),
},
{
/* [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(298),
+ /* overloads */ OverloadIndex(300),
},
{
/* [24] */
@@ -11539,19 +11591,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(299),
+ /* overloads */ OverloadIndex(301),
},
{
/* [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(300),
+ /* overloads */ OverloadIndex(302),
},
{
/* [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(301),
+ /* overloads */ OverloadIndex(303),
},
{
/* [39] */
@@ -11565,13 +11617,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(302),
+ /* overloads */ OverloadIndex(304),
},
{
/* [41] */
/* fn vector_times_scalar[T : f32_f16, N : num](vec<N, T>, T) -> vec<N, T> */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(303),
+ /* overloads */ OverloadIndex(305),
},
{
/* [42] */
@@ -11928,19 +11980,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(304),
+ /* overloads */ OverloadIndex(306),
},
{
/* [93] */
/* fn s_dot(u32, u32, u32) -> i32 */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(305),
+ /* overloads */ OverloadIndex(307),
},
{
/* [94] */
/* fn u_dot(u32, u32, u32) -> u32 */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(306),
+ /* overloads */ OverloadIndex(308),
},
{
/* [95] */
@@ -12009,6 +12061,13 @@
/* num overloads */ 2,
/* overloads */ OverloadIndex(292),
},
+ {
+ /* [104] */
+ /* fn group_non_uniform_shuffle_down[T : scalar](scope: u32, value: T, delta: u32) -> T */
+ /* fn group_non_uniform_shuffle_down[N : num, T : scalar](scope: u32, value: vec<N, T>, delta: u32) -> vec<N, T> */
+ /* num overloads */ 2,
+ /* overloads */ OverloadIndex(294),
+ },
};
// clang-format on
diff --git a/src/tint/lang/spirv/reader/lower/builtins.cc b/src/tint/lang/spirv/reader/lower/builtins.cc
index 04cc6c1..f3f2b90 100644
--- a/src/tint/lang/spirv/reader/lower/builtins.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins.cc
@@ -227,6 +227,9 @@
case spirv::BuiltinFn::kGroupNonUniformShuffleXor:
GroupNonUniformBuiltin(builtin, core::BuiltinFn::kSubgroupShuffleXor);
break;
+ case spirv::BuiltinFn::kGroupNonUniformShuffleDown:
+ GroupNonUniformBuiltin(builtin, core::BuiltinFn::kSubgroupShuffleDown);
+ break;
case spirv::BuiltinFn::kGroupNonUniformBroadcastFirst:
GroupNonUniformBroadcastFirst(builtin);
break;
diff --git a/src/tint/lang/spirv/reader/lower/builtins_test.cc b/src/tint/lang/spirv/reader/lower/builtins_test.cc
index fe27e70..ae08b9e 100644
--- a/src/tint/lang/spirv/reader/lower/builtins_test.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins_test.cc
@@ -10024,5 +10024,139 @@
EXPECT_EQ(expect, str());
}
+TEST_F(SpirvReader_BuiltinsTest, NonUniformShuffleDown_Constant_BoolScalar) {
+ auto* ep = b.ComputeFunction("main");
+
+ b.Append(ep->Block(), [&] { //
+ b.Call<spirv::ir::BuiltinCall>(ty.bool_(), spirv::BuiltinFn::kGroupNonUniformShuffleDown,
+ 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_down 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 = subgroupShuffleDown %2, 1u
+ %4:bool = convert %3
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, NonUniformShuffleDown_Constant_BoolVector) {
+ auto* ep = b.ComputeFunction("main");
+
+ b.Append(ep->Block(), [&] { //
+ b.Call<spirv::ir::BuiltinCall>(ty.vec3(ty.bool_()),
+ spirv::BuiltinFn::kGroupNonUniformShuffleDown, 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_down 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> = subgroupShuffleDown %2, 1u
+ %4:vec3<bool> = convert %3
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, NonUniformShuffleDown_Constant_NumericScalar) {
+ auto* ep = b.ComputeFunction("main");
+
+ b.Append(ep->Block(), [&] { //
+ b.Call<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kGroupNonUniformShuffleDown, 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_down 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 = subgroupShuffleDown 2u, 1u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, NonUniformShuffleDown_Constant_NumericVector) {
+ auto* ep = b.ComputeFunction("main");
+
+ b.Append(ep->Block(), [&] { //
+ b.Call<spirv::ir::BuiltinCall>(ty.vec3<u32>(),
+ spirv::BuiltinFn::kGroupNonUniformShuffleDown, 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_down 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> = subgroupShuffleDown 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 dc05570..639f19d 100644
--- a/src/tint/lang/spirv/reader/parser/builtin_test.cc
+++ b/src/tint/lang/spirv/reader/parser/builtin_test.cc
@@ -2625,5 +2625,139 @@
SPV_ENV_VULKAN_1_1);
}
+TEST_F(SpirvParserTest, NonUniformShuffleDown_Constant_BoolScalar) {
+ EXPECT_IR_SPV(R"(
+ OpCapability Shader
+ OpCapability GroupNonUniformShuffleRelative
+ 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 = OpGroupNonUniformShuffleDown %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_down 3u, true, 1u
+ ret
+ }
+}
+)",
+ SPV_ENV_VULKAN_1_1);
+}
+
+TEST_F(SpirvParserTest, NonUniformShuffleDown_Constant_BoolVector) {
+ EXPECT_IR_SPV(R"(
+ OpCapability Shader
+ OpCapability GroupNonUniformShuffleRelative
+ 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 = OpGroupNonUniformShuffleDown %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_down 3u, vec3<bool>(true, false, true), 1u
+ ret
+ }
+}
+)",
+ SPV_ENV_VULKAN_1_1);
+}
+
+TEST_F(SpirvParserTest, NonUniformShuffleDown_Constant_NumericScalar) {
+ EXPECT_IR_SPV(R"(
+ OpCapability Shader
+ OpCapability GroupNonUniformShuffleRelative
+ 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 = OpGroupNonUniformShuffleDown %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_down 3u, 3u, 1u
+ ret
+ }
+}
+)",
+ SPV_ENV_VULKAN_1_1);
+}
+
+TEST_F(SpirvParserTest, NonUniformShuffleDown_Constant_NumericVector) {
+ EXPECT_IR_SPV(R"(
+ OpCapability Shader
+ OpCapability GroupNonUniformShuffleRelative
+ 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 = OpGroupNonUniformShuffleDown %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_down 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 970d0bf..0021d1c 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -2234,6 +2234,9 @@
case spv::Op::OpGroupNonUniformShuffleXor:
EmitSubgroupBuiltin(inst, spirv::BuiltinFn::kGroupNonUniformShuffleXor);
break;
+ case spv::Op::OpGroupNonUniformShuffleDown:
+ EmitSubgroupBuiltin(inst, spirv::BuiltinFn::kGroupNonUniformShuffleDown);
+ 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 4b69db5..e24b251 100644
--- a/src/tint/lang/spirv/spirv.def
+++ b/src/tint/lang/spirv/spirv.def
@@ -1846,3 +1846,8 @@
@must_use @stage("fragment", "compute") implicit(N: num, T: scalar)
fn group_non_uniform_shuffle_xor(scope: u32, value: vec<N, T>, mask: u32) -> vec<N, T>
+@must_use @stage("fragment", "compute") implicit(T: scalar)
+fn group_non_uniform_shuffle_down(scope: u32, value: T, delta: u32) -> T
+@must_use @stage("fragment", "compute") implicit(N: num, T: scalar)
+fn group_non_uniform_shuffle_down(scope: u32, value: vec<N, T>, delta: 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 364066f..8f29286 100644
--- a/src/tint/lang/spirv/writer/printer/printer.cc
+++ b/src/tint/lang/spirv/writer/printer/printer.cc
@@ -1703,6 +1703,9 @@
case BuiltinFn::kGroupNonUniformShuffleXor:
op = spv::Op::OpGroupNonUniformShuffleXor;
break;
+ case BuiltinFn::kGroupNonUniformShuffleDown:
+ op = spv::Op::OpGroupNonUniformShuffleDown;
+ break;
case spirv::BuiltinFn::kNone:
TINT_ICE() << "undefined spirv ir function";
}
