[spirv-reader][ir] Add support for `OpGroupNonUniformQuadSwap`
Add support to convert the `OpGroupNonUniformQuadSwap` instruction into
the `quadSwapX`, `quadSwapY`, or `quadSwapDiagonal` instruction as
needed.
Fixed: 431031432
Change-Id: Ie7d095ae0c8ffa3ddf8fb55b2da6c0cccbc86bd4
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/252455
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: David Neto <dneto@google.com>
diff --git a/src/tint/cmd/fuzz/wgsl/dictionary.txt b/src/tint/cmd/fuzz/wgsl/dictionary.txt
index 4daf792..ef18175 100644
--- a/src/tint/cmd/fuzz/wgsl/dictionary.txt
+++ b/src/tint/cmd/fuzz/wgsl/dictionary.txt
@@ -195,6 +195,7 @@
"derivative_uniformity"
"determinant"
"diagnostic"
+"dir"
"discard"
"distance"
"dot"
diff --git a/src/tint/lang/core/core.def b/src/tint/lang/core/core.def
index f854ba1..4d2fafe 100644
--- a/src/tint/lang/core/core.def
+++ b/src/tint/lang/core/core.def
@@ -248,6 +248,7 @@
coords
count
depth
+ dir
dref
compare_value
elements
diff --git a/src/tint/lang/core/enums.cc b/src/tint/lang/core/enums.cc
index ba8e361..5ccba9f 100644
--- a/src/tint/lang/core/enums.cc
+++ b/src/tint/lang/core/enums.cc
@@ -1197,6 +1197,8 @@
return "depth";
case ParameterUsage::kDepthRef:
return "depth_ref";
+ case ParameterUsage::kDir:
+ return "dir";
case ParameterUsage::kDref:
return "dref";
case ParameterUsage::kE:
diff --git a/src/tint/lang/core/enums.h b/src/tint/lang/core/enums.h
index 8b87249..5def098 100644
--- a/src/tint/lang/core/enums.h
+++ b/src/tint/lang/core/enums.h
@@ -652,6 +652,7 @@
kDelta,
kDepth,
kDepthRef,
+ kDir,
kDref,
kE,
kElements,
diff --git a/src/tint/lang/spirv/builtin_fn.cc b/src/tint/lang/spirv/builtin_fn.cc
index 0f8eb25..590f37f 100644
--- a/src/tint/lang/spirv/builtin_fn.cc
+++ b/src/tint/lang/spirv/builtin_fn.cc
@@ -244,6 +244,8 @@
return "group_non_uniform_broadcast_first";
case BuiltinFn::kGroupNonUniformQuadBroadcast:
return "group_non_uniform_quad_broadcast";
+ case BuiltinFn::kGroupNonUniformQuadSwap:
+ return "group_non_uniform_quad_swap";
}
return "<unknown>";
}
@@ -359,6 +361,7 @@
case BuiltinFn::kGroupNonUniformBroadcast:
case BuiltinFn::kGroupNonUniformBroadcastFirst:
case BuiltinFn::kGroupNonUniformQuadBroadcast:
+ case BuiltinFn::kGroupNonUniformQuadSwap:
break;
}
return core::ir::Instruction::Accesses{};
diff --git a/src/tint/lang/spirv/builtin_fn.cc.tmpl b/src/tint/lang/spirv/builtin_fn.cc.tmpl
index 79d0356..69a02e6 100644
--- a/src/tint/lang/spirv/builtin_fn.cc.tmpl
+++ b/src/tint/lang/spirv/builtin_fn.cc.tmpl
@@ -138,6 +138,7 @@
case BuiltinFn::kGroupNonUniformBroadcast:
case BuiltinFn::kGroupNonUniformBroadcastFirst:
case BuiltinFn::kGroupNonUniformQuadBroadcast:
+ case BuiltinFn::kGroupNonUniformQuadSwap:
break;
}
return core::ir::Instruction::Accesses{};
diff --git a/src/tint/lang/spirv/builtin_fn.h b/src/tint/lang/spirv/builtin_fn.h
index bbef717..dc690b2 100644
--- a/src/tint/lang/spirv/builtin_fn.h
+++ b/src/tint/lang/spirv/builtin_fn.h
@@ -149,6 +149,7 @@
kGroupNonUniformBroadcast,
kGroupNonUniformBroadcastFirst,
kGroupNonUniformQuadBroadcast,
+ kGroupNonUniformQuadSwap,
kNone,
};
diff --git a/src/tint/lang/spirv/intrinsic/data.cc b/src/tint/lang/spirv/intrinsic/data.cc
index 95334d5..50962ea 100644
--- a/src/tint/lang/spirv/intrinsic/data.cc
+++ b/src/tint/lang/spirv/intrinsic/data.cc
@@ -6206,48 +6206,48 @@
},
{
/* [750] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(943),
+ /* usage */ core::ParameterUsage::kScope,
+ /* matcher_indices */ MatcherIndicesIndex(592),
},
{
/* [751] */
- /* usage */ core::ParameterUsage::kNone,
+ /* usage */ core::ParameterUsage::kE,
/* matcher_indices */ MatcherIndicesIndex(8),
},
{
/* [752] */
- /* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(465),
+ /* usage */ core::ParameterUsage::kDir,
+ /* matcher_indices */ MatcherIndicesIndex(592),
},
{
/* [753] */
- /* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(21),
+ /* usage */ core::ParameterUsage::kScope,
+ /* matcher_indices */ MatcherIndicesIndex(592),
},
{
/* [754] */
- /* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(474),
+ /* usage */ core::ParameterUsage::kE,
+ /* matcher_indices */ MatcherIndicesIndex(1003),
},
{
/* [755] */
- /* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(21),
+ /* usage */ core::ParameterUsage::kDir,
+ /* matcher_indices */ MatcherIndicesIndex(592),
},
{
/* [756] */
- /* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(483),
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(943),
},
{
/* [757] */
- /* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(21),
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(8),
},
{
/* [758] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(492),
+ /* matcher_indices */ MatcherIndicesIndex(465),
},
{
/* [759] */
@@ -6257,7 +6257,7 @@
{
/* [760] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(501),
+ /* matcher_indices */ MatcherIndicesIndex(474),
},
{
/* [761] */
@@ -6267,7 +6267,7 @@
{
/* [762] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(510),
+ /* matcher_indices */ MatcherIndicesIndex(483),
},
{
/* [763] */
@@ -6277,37 +6277,37 @@
{
/* [764] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(519),
+ /* matcher_indices */ MatcherIndicesIndex(492),
},
{
/* [765] */
/* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(32),
+ /* matcher_indices */ MatcherIndicesIndex(21),
},
{
/* [766] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(528),
+ /* matcher_indices */ MatcherIndicesIndex(501),
},
{
/* [767] */
/* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(32),
+ /* matcher_indices */ MatcherIndicesIndex(21),
},
{
/* [768] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(537),
+ /* matcher_indices */ MatcherIndicesIndex(510),
},
{
/* [769] */
/* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(32),
+ /* matcher_indices */ MatcherIndicesIndex(21),
},
{
/* [770] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(546),
+ /* matcher_indices */ MatcherIndicesIndex(519),
},
{
/* [771] */
@@ -6317,220 +6317,250 @@
{
/* [772] */
/* usage */ core::ParameterUsage::kImage,
- /* matcher_indices */ MatcherIndicesIndex(681),
+ /* matcher_indices */ MatcherIndicesIndex(528),
},
{
/* [773] */
+ /* usage */ core::ParameterUsage::kLevel,
+ /* matcher_indices */ MatcherIndicesIndex(32),
+ },
+ {
+ /* [774] */
+ /* usage */ core::ParameterUsage::kImage,
+ /* matcher_indices */ MatcherIndicesIndex(537),
+ },
+ {
+ /* [775] */
+ /* usage */ core::ParameterUsage::kLevel,
+ /* matcher_indices */ MatcherIndicesIndex(32),
+ },
+ {
+ /* [776] */
+ /* usage */ core::ParameterUsage::kImage,
+ /* matcher_indices */ MatcherIndicesIndex(546),
+ },
+ {
+ /* [777] */
+ /* usage */ core::ParameterUsage::kLevel,
+ /* matcher_indices */ MatcherIndicesIndex(32),
+ },
+ {
+ /* [778] */
+ /* usage */ core::ParameterUsage::kImage,
+ /* matcher_indices */ MatcherIndicesIndex(681),
+ },
+ {
+ /* [779] */
/* usage */ core::ParameterUsage::kCoords,
/* matcher_indices */ MatcherIndicesIndex(20),
},
{
- /* [774] */
+ /* [780] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(959),
},
{
- /* [775] */
+ /* [781] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(963),
},
{
- /* [776] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(967),
- },
- {
- /* [777] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(8),
- },
- {
- /* [778] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(967),
- },
- {
- /* [779] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(903),
- },
- {
- /* [780] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(1015),
- },
- {
- /* [781] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(909),
- },
- {
/* [782] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(1024),
+ /* matcher_indices */ MatcherIndicesIndex(967),
},
{
/* [783] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(1027),
+ /* matcher_indices */ MatcherIndicesIndex(8),
},
{
/* [784] */
- /* usage */ core::ParameterUsage::kX,
- /* matcher_indices */ MatcherIndicesIndex(8),
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(967),
},
{
/* [785] */
- /* usage */ core::ParameterUsage::kI,
- /* matcher_indices */ MatcherIndicesIndex(971),
- },
- {
- /* [786] */
- /* usage */ core::ParameterUsage::kX,
+ /* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(903),
},
{
+ /* [786] */
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(1015),
+ },
+ {
/* [787] */
- /* usage */ core::ParameterUsage::kI,
- /* matcher_indices */ MatcherIndicesIndex(901),
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(909),
},
{
/* [788] */
- /* usage */ core::ParameterUsage::kX,
- /* matcher_indices */ MatcherIndicesIndex(8),
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(1024),
},
{
/* [789] */
- /* usage */ core::ParameterUsage::kI,
- /* matcher_indices */ MatcherIndicesIndex(975),
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(1027),
},
{
/* [790] */
/* usage */ core::ParameterUsage::kX,
- /* matcher_indices */ MatcherIndicesIndex(1015),
+ /* matcher_indices */ MatcherIndicesIndex(8),
},
{
/* [791] */
/* usage */ core::ParameterUsage::kI,
- /* matcher_indices */ MatcherIndicesIndex(907),
+ /* matcher_indices */ MatcherIndicesIndex(971),
},
{
/* [792] */
+ /* usage */ core::ParameterUsage::kX,
+ /* matcher_indices */ MatcherIndicesIndex(903),
+ },
+ {
+ /* [793] */
+ /* usage */ core::ParameterUsage::kI,
+ /* matcher_indices */ MatcherIndicesIndex(901),
+ },
+ {
+ /* [794] */
+ /* usage */ core::ParameterUsage::kX,
+ /* matcher_indices */ MatcherIndicesIndex(8),
+ },
+ {
+ /* [795] */
+ /* usage */ core::ParameterUsage::kI,
+ /* matcher_indices */ MatcherIndicesIndex(975),
+ },
+ {
+ /* [796] */
+ /* usage */ core::ParameterUsage::kX,
+ /* matcher_indices */ MatcherIndicesIndex(1015),
+ },
+ {
+ /* [797] */
+ /* usage */ core::ParameterUsage::kI,
+ /* matcher_indices */ MatcherIndicesIndex(907),
+ },
+ {
+ /* [798] */
/* usage */ core::ParameterUsage::kScope,
/* matcher_indices */ MatcherIndicesIndex(592),
},
{
- /* [793] */
+ /* [799] */
/* usage */ core::ParameterUsage::kValue,
/* matcher_indices */ MatcherIndicesIndex(8),
},
{
- /* [794] */
+ /* [800] */
/* usage */ core::ParameterUsage::kScope,
/* matcher_indices */ MatcherIndicesIndex(592),
},
{
- /* [795] */
+ /* [801] */
/* usage */ core::ParameterUsage::kValue,
/* matcher_indices */ MatcherIndicesIndex(1003),
},
{
- /* [796] */
+ /* [802] */
/* usage */ core::ParameterUsage::kImage,
/* matcher_indices */ MatcherIndicesIndex(357),
},
{
- /* [797] */
+ /* [803] */
/* usage */ core::ParameterUsage::kImage,
/* matcher_indices */ MatcherIndicesIndex(366),
},
{
- /* [798] */
+ /* [804] */
/* usage */ core::ParameterUsage::kImage,
/* matcher_indices */ MatcherIndicesIndex(375),
},
{
- /* [799] */
+ /* [805] */
/* usage */ core::ParameterUsage::kImage,
/* matcher_indices */ MatcherIndicesIndex(384),
},
{
- /* [800] */
+ /* [806] */
/* usage */ core::ParameterUsage::kImage,
/* matcher_indices */ MatcherIndicesIndex(393),
},
{
- /* [801] */
+ /* [807] */
/* usage */ core::ParameterUsage::kImage,
/* matcher_indices */ MatcherIndicesIndex(402),
},
{
- /* [802] */
+ /* [808] */
/* usage */ core::ParameterUsage::kImage,
/* matcher_indices */ MatcherIndicesIndex(411),
},
{
- /* [803] */
+ /* [809] */
/* usage */ core::ParameterUsage::kImage,
/* matcher_indices */ MatcherIndicesIndex(420),
},
{
- /* [804] */
+ /* [810] */
/* usage */ core::ParameterUsage::kImage,
/* matcher_indices */ MatcherIndicesIndex(429),
},
{
- /* [805] */
+ /* [811] */
/* usage */ core::ParameterUsage::kImage,
/* matcher_indices */ MatcherIndicesIndex(438),
},
{
- /* [806] */
+ /* [812] */
/* usage */ core::ParameterUsage::kImage,
/* matcher_indices */ MatcherIndicesIndex(447),
},
{
- /* [807] */
+ /* [813] */
/* usage */ core::ParameterUsage::kImage,
/* matcher_indices */ MatcherIndicesIndex(456),
},
{
- /* [808] */
+ /* [814] */
/* usage */ core::ParameterUsage::kImage,
/* matcher_indices */ MatcherIndicesIndex(555),
},
{
- /* [809] */
+ /* [815] */
/* usage */ core::ParameterUsage::kImage,
/* matcher_indices */ MatcherIndicesIndex(564),
},
{
- /* [810] */
+ /* [816] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(1058),
},
{
- /* [811] */
+ /* [817] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(1060),
},
{
- /* [812] */
+ /* [818] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(929),
},
{
- /* [813] */
+ /* [819] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(1062),
},
{
- /* [814] */
+ /* [820] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(1064),
},
{
- /* [815] */
+ /* [821] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(1066),
},
@@ -9136,7 +9166,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 4,
/* templates */ TemplateIndex(72),
- /* parameters */ ParameterIndex(772),
+ /* parameters */ ParameterIndex(778),
/* return_matcher_indices */ MatcherIndicesIndex(1056),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9279,7 +9309,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(76),
- /* parameters */ ParameterIndex(796),
+ /* parameters */ ParameterIndex(802),
/* return_matcher_indices */ MatcherIndicesIndex(8),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9290,7 +9320,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(76),
- /* parameters */ ParameterIndex(797),
+ /* parameters */ ParameterIndex(803),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9301,7 +9331,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(76),
- /* parameters */ ParameterIndex(798),
+ /* parameters */ ParameterIndex(804),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9312,7 +9342,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(76),
- /* parameters */ ParameterIndex(799),
+ /* parameters */ ParameterIndex(805),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9323,7 +9353,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(76),
- /* parameters */ ParameterIndex(800),
+ /* parameters */ ParameterIndex(806),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9334,7 +9364,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(76),
- /* parameters */ ParameterIndex(801),
+ /* parameters */ ParameterIndex(807),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9345,7 +9375,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 5,
/* templates */ TemplateIndex(18),
- /* parameters */ ParameterIndex(802),
+ /* parameters */ ParameterIndex(808),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9356,7 +9386,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(129),
- /* parameters */ ParameterIndex(803),
+ /* parameters */ ParameterIndex(809),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9367,7 +9397,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(129),
- /* parameters */ ParameterIndex(804),
+ /* parameters */ ParameterIndex(810),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9378,7 +9408,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(129),
- /* parameters */ ParameterIndex(805),
+ /* parameters */ ParameterIndex(811),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9389,7 +9419,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(129),
- /* parameters */ ParameterIndex(806),
+ /* parameters */ ParameterIndex(812),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9400,7 +9430,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(80),
- /* parameters */ ParameterIndex(807),
+ /* parameters */ ParameterIndex(813),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9543,7 +9573,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 6,
/* templates */ TemplateIndex(12),
- /* parameters */ ParameterIndex(752),
+ /* parameters */ ParameterIndex(758),
/* return_matcher_indices */ MatcherIndicesIndex(8),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9554,7 +9584,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 6,
/* templates */ TemplateIndex(12),
- /* parameters */ ParameterIndex(754),
+ /* parameters */ ParameterIndex(760),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9565,7 +9595,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 6,
/* templates */ TemplateIndex(12),
- /* parameters */ ParameterIndex(756),
+ /* parameters */ ParameterIndex(762),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9576,7 +9606,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 6,
/* templates */ TemplateIndex(12),
- /* parameters */ ParameterIndex(758),
+ /* parameters */ ParameterIndex(764),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9587,7 +9617,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 */),
},
@@ -9598,7 +9628,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 */),
},
@@ -9609,7 +9639,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 5,
/* templates */ TemplateIndex(23),
- /* parameters */ ParameterIndex(764),
+ /* parameters */ ParameterIndex(770),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9620,7 +9650,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 5,
/* templates */ TemplateIndex(23),
- /* parameters */ ParameterIndex(766),
+ /* parameters */ ParameterIndex(772),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9631,7 +9661,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 5,
/* templates */ TemplateIndex(23),
- /* parameters */ ParameterIndex(768),
+ /* parameters */ ParameterIndex(774),
/* return_matcher_indices */ MatcherIndicesIndex(42),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9642,7 +9672,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 5,
/* templates */ TemplateIndex(23),
- /* parameters */ ParameterIndex(770),
+ /* parameters */ ParameterIndex(776),
/* return_matcher_indices */ MatcherIndicesIndex(718),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9917,7 +9947,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(144),
- /* parameters */ ParameterIndex(810),
+ /* parameters */ ParameterIndex(816),
/* return_matcher_indices */ MatcherIndicesIndex(8),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9928,7 +9958,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(148),
- /* parameters */ ParameterIndex(810),
+ /* parameters */ ParameterIndex(816),
/* return_matcher_indices */ MatcherIndicesIndex(8),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9939,7 +9969,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 3,
/* templates */ TemplateIndex(145),
- /* parameters */ ParameterIndex(811),
+ /* parameters */ ParameterIndex(817),
/* return_matcher_indices */ MatcherIndicesIndex(21),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9950,7 +9980,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 3,
/* templates */ TemplateIndex(149),
- /* parameters */ ParameterIndex(811),
+ /* parameters */ ParameterIndex(817),
/* return_matcher_indices */ MatcherIndicesIndex(21),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9961,7 +9991,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(91),
- /* parameters */ ParameterIndex(812),
+ /* parameters */ ParameterIndex(818),
/* return_matcher_indices */ MatcherIndicesIndex(32),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -9972,7 +10002,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(96),
- /* parameters */ ParameterIndex(812),
+ /* parameters */ ParameterIndex(818),
/* return_matcher_indices */ MatcherIndicesIndex(32),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10214,7 +10244,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(10),
- /* parameters */ ParameterIndex(813),
+ /* parameters */ ParameterIndex(819),
/* return_matcher_indices */ MatcherIndicesIndex(1062),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10225,7 +10255,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(10),
- /* parameters */ ParameterIndex(814),
+ /* parameters */ ParameterIndex(820),
/* return_matcher_indices */ MatcherIndicesIndex(1064),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10236,7 +10266,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(10),
- /* parameters */ ParameterIndex(815),
+ /* parameters */ ParameterIndex(821),
/* return_matcher_indices */ MatcherIndicesIndex(1066),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10434,7 +10464,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(161),
- /* parameters */ ParameterIndex(781),
+ /* parameters */ ParameterIndex(787),
/* return_matcher_indices */ MatcherIndicesIndex(1015),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10456,7 +10486,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(110),
- /* parameters */ ParameterIndex(782),
+ /* parameters */ ParameterIndex(788),
/* return_matcher_indices */ MatcherIndicesIndex(1021),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10577,7 +10607,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(181),
- /* parameters */ ParameterIndex(784),
+ /* parameters */ ParameterIndex(790),
/* return_matcher_indices */ MatcherIndicesIndex(8),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10588,7 +10618,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 3,
/* templates */ TemplateIndex(164),
- /* parameters */ ParameterIndex(786),
+ /* parameters */ ParameterIndex(792),
/* return_matcher_indices */ MatcherIndicesIndex(903),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10599,7 +10629,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 3,
/* templates */ TemplateIndex(167),
- /* parameters */ ParameterIndex(788),
+ /* parameters */ ParameterIndex(794),
/* return_matcher_indices */ MatcherIndicesIndex(8),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10610,7 +10640,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 4,
/* templates */ TemplateIndex(114),
- /* parameters */ ParameterIndex(790),
+ /* parameters */ ParameterIndex(796),
/* return_matcher_indices */ MatcherIndicesIndex(1015),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10676,7 +10706,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(118),
- /* parameters */ ParameterIndex(782),
+ /* parameters */ ParameterIndex(788),
/* return_matcher_indices */ MatcherIndicesIndex(1021),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10698,7 +10728,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 3,
/* templates */ TemplateIndex(119),
- /* parameters */ ParameterIndex(780),
+ /* parameters */ ParameterIndex(786),
/* return_matcher_indices */ MatcherIndicesIndex(1045),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10720,7 +10750,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(170),
- /* parameters */ ParameterIndex(781),
+ /* parameters */ ParameterIndex(787),
/* return_matcher_indices */ MatcherIndicesIndex(1015),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10742,7 +10772,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(173),
- /* parameters */ ParameterIndex(781),
+ /* parameters */ ParameterIndex(787),
/* return_matcher_indices */ MatcherIndicesIndex(1015),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10764,7 +10794,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(176),
- /* parameters */ ParameterIndex(781),
+ /* parameters */ ParameterIndex(787),
/* return_matcher_indices */ MatcherIndicesIndex(1015),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10797,7 +10827,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(33),
- /* parameters */ ParameterIndex(792),
+ /* parameters */ ParameterIndex(798),
/* return_matcher_indices */ MatcherIndicesIndex(8),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -10808,23 +10838,45 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(32),
- /* parameters */ ParameterIndex(794),
+ /* parameters */ ParameterIndex(800),
/* return_matcher_indices */ MatcherIndicesIndex(1003),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
/* [288] */
+ /* 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(750),
+ /* return_matcher_indices */ MatcherIndicesIndex(8),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [289] */
+ /* 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(753),
+ /* return_matcher_indices */ MatcherIndicesIndex(1003),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [290] */
/* 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(750),
+ /* parameters */ ParameterIndex(756),
/* return_matcher_indices */ MatcherIndicesIndex(592),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [289] */
+ /* [291] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -10835,7 +10887,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [290] */
+ /* [292] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -10846,47 +10898,25 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [291] */
- /* 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(808),
- /* return_matcher_indices */ MatcherIndicesIndex(8),
- /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
- },
- {
- /* [292] */
- /* 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(809),
- /* return_matcher_indices */ MatcherIndicesIndex(8),
- /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
- },
- {
/* [293] */
/* 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(774),
- /* return_matcher_indices */ MatcherIndicesIndex(955),
+ /* 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 */ 2,
- /* num_explicit_templates */ 0,
- /* num_templates */ 3,
- /* templates */ TemplateIndex(158),
- /* parameters */ ParameterIndex(776),
- /* return_matcher_indices */ MatcherIndicesIndex(967),
+ /* 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 */),
},
{
@@ -10894,10 +10924,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(778),
- /* return_matcher_indices */ MatcherIndicesIndex(1015),
+ /* num_templates */ 4,
+ /* templates */ TemplateIndex(106),
+ /* parameters */ ParameterIndex(780),
+ /* return_matcher_indices */ MatcherIndicesIndex(955),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -10907,12 +10937,34 @@
/* num_explicit_templates */ 0,
/* num_templates */ 3,
/* templates */ TemplateIndex(158),
+ /* parameters */ ParameterIndex(782),
+ /* return_matcher_indices */ MatcherIndicesIndex(967),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [297] */
+ /* 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),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [298] */
+ /* 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 */),
},
{
- /* [297] */
+ /* [299] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -10923,18 +10975,18 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [298] */
+ /* [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(779),
+ /* parameters */ ParameterIndex(785),
/* return_matcher_indices */ MatcherIndicesIndex(963),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [299] */
+ /* [301] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -10945,7 +10997,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [300] */
+ /* [302] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -10965,7 +11017,7 @@
/* [0] */
/* fn array_length[I : u32, A : access](ptr<storage, struct_with_runtime_array, A>, I) -> u32 */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(288),
+ /* overloads */ OverloadIndex(290),
},
{
/* [1] */
@@ -11062,19 +11114,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(289),
+ /* overloads */ OverloadIndex(291),
},
{
/* [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(289),
+ /* overloads */ OverloadIndex(291),
},
{
/* [16] */
/* fn dot[N : num, T : f32_f16](vec<N, T>, vec<N, T>) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(290),
+ /* overloads */ OverloadIndex(292),
},
{
/* [17] */
@@ -11161,13 +11213,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(291),
+ /* overloads */ OverloadIndex(293),
},
{
/* [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(292),
+ /* overloads */ OverloadIndex(294),
},
{
/* [24] */
@@ -11403,19 +11455,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(293),
+ /* overloads */ OverloadIndex(295),
},
{
/* [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(294),
+ /* overloads */ OverloadIndex(296),
},
{
/* [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(295),
+ /* overloads */ OverloadIndex(297),
},
{
/* [39] */
@@ -11429,13 +11481,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(296),
+ /* overloads */ OverloadIndex(298),
},
{
/* [41] */
/* fn vector_times_scalar[T : f32_f16, N : num](vec<N, T>, T) -> vec<N, T> */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(297),
+ /* overloads */ OverloadIndex(299),
},
{
/* [42] */
@@ -11792,19 +11844,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(298),
+ /* overloads */ OverloadIndex(300),
},
{
/* [93] */
/* fn s_dot(u32, u32, u32) -> i32 */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(299),
+ /* overloads */ OverloadIndex(301),
},
{
/* [94] */
/* fn u_dot(u32, u32, u32) -> u32 */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(300),
+ /* overloads */ OverloadIndex(302),
},
{
/* [95] */
@@ -11852,6 +11904,13 @@
/* num overloads */ 2,
/* overloads */ OverloadIndex(284),
},
+ {
+ /* [101] */
+ /* fn group_non_uniform_quad_swap[T : scalar](scope: u32, e: T, @const dir: u32) -> T */
+ /* fn group_non_uniform_quad_swap[N : num, T : scalar](scope: u32, e: vec<N, T>, @const dir: u32) -> vec<N, T> */
+ /* num overloads */ 2,
+ /* overloads */ OverloadIndex(288),
+ },
};
// clang-format on
diff --git a/src/tint/lang/spirv/reader/lower/builtins.cc b/src/tint/lang/spirv/reader/lower/builtins.cc
index 74d29f9..c5e95a9 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::kGroupNonUniformQuadBroadcast:
GroupNonUniformBroadcast(builtin, core::BuiltinFn::kQuadBroadcast);
break;
+ case spirv::BuiltinFn::kGroupNonUniformQuadSwap:
+ GroupNonUniformQuadSwap(builtin);
+ break;
case spirv::BuiltinFn::kAtomicLoad:
case spirv::BuiltinFn::kAtomicStore:
case spirv::BuiltinFn::kAtomicExchange:
@@ -271,6 +274,48 @@
}
}
+ void GroupNonUniformQuadSwap(spirv::ir::BuiltinCall* call) {
+ auto* value = call->Args()[1];
+ auto* dir_val = call->Args()[2];
+
+ TINT_ASSERT(dir_val->Is<core::ir::Constant>());
+ auto* cnst = dir_val->As<core::ir::Constant>();
+ TINT_ASSERT(cnst);
+
+ uint32_t dir = cnst->Value()->ValueAs<uint32_t>();
+ core::BuiltinFn fn = core::BuiltinFn::kNone;
+ switch (dir) {
+ case 0:
+ fn = core::BuiltinFn::kQuadSwapX;
+ break;
+ case 1:
+ fn = core::BuiltinFn::kQuadSwapY;
+ break;
+ case 2:
+ fn = core::BuiltinFn::kQuadSwapDiagonal;
+ break;
+ default:
+ TINT_UNREACHABLE();
+ }
+
+ auto* type = call->Result()->Type();
+ b.InsertBefore(call, [&] {
+ if (type->DeepestElement()->Is<core::type::Bool>()) {
+ type = ty.MatchWidth(ty.u32(), type);
+ value = b.Convert(type, value)->Result();
+ }
+
+ core::ir::Value* c = b.Call(type, fn, Vector{value})->Result();
+
+ if (type != call->Result()->Type()) {
+ c = b.Convert(call->Result()->Type(), c)->Result();
+ }
+
+ call->Result()->ReplaceAllUsesWith(c);
+ });
+ call->Destroy();
+ }
+
void GroupNonUniformBroadcast(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 4feb75b..658b493 100644
--- a/src/tint/lang/spirv/reader/lower/builtins_test.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins_test.cc
@@ -9625,5 +9625,138 @@
EXPECT_EQ(expect, str());
}
+TEST_F(SpirvReader_BuiltinsTest, NonUniformQuadSwap_Constant_BoolScalar) {
+ auto* ep = b.ComputeFunction("main");
+
+ b.Append(ep->Block(), [&] { //
+ b.Call<spirv::ir::BuiltinCall>(ty.bool_(), spirv::BuiltinFn::kGroupNonUniformQuadSwap, 3_u,
+ true, 0_u);
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:bool = spirv.group_non_uniform_quad_swap 3u, true, 0u
+ 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 = quadSwapX %2
+ %4:bool = convert %3
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, NonUniformQuadSwap_Constant_BoolVector) {
+ auto* ep = b.ComputeFunction("main");
+
+ b.Append(ep->Block(), [&] { //
+ b.Call<spirv::ir::BuiltinCall>(ty.vec3(ty.bool_()),
+ spirv::BuiltinFn::kGroupNonUniformQuadSwap, 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_quad_swap 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> = quadSwapY %2
+ %4:vec3<bool> = convert %3
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, NonUniformQuadSwap_Constant_NumericScalar) {
+ auto* ep = b.ComputeFunction("main");
+
+ b.Append(ep->Block(), [&] { //
+ b.Call<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kGroupNonUniformQuadSwap, 3_u,
+ 2_u, 2_u);
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = spirv.group_non_uniform_quad_swap 3u, 2u, 2u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ Run(Builtins);
+
+ auto expect = R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = quadSwapDiagonal 2u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, NonUniformQuadSwap_Constant_NumericVector) {
+ auto* ep = b.ComputeFunction("main");
+
+ b.Append(ep->Block(), [&] { //
+ b.Call<spirv::ir::BuiltinCall>(ty.vec3<u32>(), spirv::BuiltinFn::kGroupNonUniformQuadSwap,
+ 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_quad_swap 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> = quadSwapY vec3<u32>(2u, 3u, 2u)
+ 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 5ca4005..72e8e31 100644
--- a/src/tint/lang/spirv/reader/parser/builtin_test.cc
+++ b/src/tint/lang/spirv/reader/parser/builtin_test.cc
@@ -2221,5 +2221,141 @@
SPV_ENV_VULKAN_1_1);
}
+TEST_F(SpirvParserTest, NonUniformQuadSwap_Constant_BoolScalar) {
+ EXPECT_IR_SPV(R"(
+ OpCapability Shader
+ OpCapability GroupNonUniformQuad
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 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 = OpGroupNonUniformQuadSwap %bool %uint_3 %true %uint_0
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:bool = spirv.group_non_uniform_quad_swap 3u, true, 0u
+ ret
+ }
+}
+)",
+ SPV_ENV_VULKAN_1_1);
+}
+
+TEST_F(SpirvParserTest, NonUniformQuadSwap_Constant_BoolVector) {
+ EXPECT_IR_SPV(R"(
+ OpCapability Shader
+ OpCapability GroupNonUniformQuad
+ 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 = OpGroupNonUniformQuadSwap %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_quad_swap 3u, vec3<bool>(true, false, true), 1u
+ ret
+ }
+}
+)",
+ SPV_ENV_VULKAN_1_1);
+}
+
+TEST_F(SpirvParserTest, NonUniformQuadSwap_Constant_NumericScalar) {
+ EXPECT_IR_SPV(R"(
+ OpCapability Shader
+ OpCapability GroupNonUniformQuad
+ 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_2 = OpConstant %uint 2
+ %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 = OpGroupNonUniformQuadSwap %uint %uint_3 %uint_1 %uint_2
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = spirv.group_non_uniform_quad_swap 3u, 1u, 2u
+ ret
+ }
+}
+)",
+ SPV_ENV_VULKAN_1_1);
+}
+
+TEST_F(SpirvParserTest, NonUniformQuadSwap_Constant_NumericVector) {
+ EXPECT_IR_SPV(R"(
+ OpCapability Shader
+ OpCapability GroupNonUniformQuad
+ 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 = OpGroupNonUniformQuadSwap %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_quad_swap 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 4cc4bd9..7abce8a 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -2223,13 +2223,16 @@
EmitSubgroupBuiltin(inst, core::BuiltinFn::kSubgroupBallot);
break;
case spv::Op::OpGroupNonUniformBroadcastFirst:
- EmitSubgroupBroadcast(inst, spirv::BuiltinFn::kGroupNonUniformBroadcastFirst);
+ EmitSubgroupBuiltin(inst, spirv::BuiltinFn::kGroupNonUniformBroadcastFirst);
break;
case spv::Op::OpGroupNonUniformBroadcast:
- EmitSubgroupBroadcast(inst, spirv::BuiltinFn::kGroupNonUniformBroadcast);
+ EmitSubgroupBuiltin(inst, spirv::BuiltinFn::kGroupNonUniformBroadcast);
break;
case spv::Op::OpGroupNonUniformQuadBroadcast:
- EmitSubgroupBroadcast(inst, spirv::BuiltinFn::kGroupNonUniformQuadBroadcast);
+ EmitSubgroupBuiltin(inst, spirv::BuiltinFn::kGroupNonUniformQuadBroadcast);
+ break;
+ case spv::Op::OpGroupNonUniformQuadSwap:
+ EmitSubgroupBuiltin(inst, spirv::BuiltinFn::kGroupNonUniformQuadSwap);
break;
default:
TINT_UNIMPLEMENTED()
@@ -2249,15 +2252,16 @@
}
}
- void EmitSubgroupBroadcast(spvtools::opt::Instruction& inst, spirv::BuiltinFn fn) {
+ void EmitSubgroupBuiltin(spvtools::opt::Instruction& inst, spirv::BuiltinFn fn) {
auto val = Value(inst.GetSingleWordInOperand(1));
// TODO(431054356): Convert core::BuiltinFn::kSubgroupBroadcast non-constant values into a
// `subgroupShuffle` when we support SPIR-V >= 1.5 source.
//
// For QuadBroadcast this will remain an error as there is no WGSL equivalent.
+ // For QuadSwap this will remain an error as there is no WGSL equivalent.
if (!val->Is<core::ir::Constant>()) {
- TINT_ICE() << "non-constant Broadcast values not supported";
+ TINT_ICE() << "non-constant GroupNonUniform `value` not supported";
}
ValidateScope(inst);
diff --git a/src/tint/lang/spirv/spirv.def b/src/tint/lang/spirv/spirv.def
index 331354f..2f3b2f4 100644
--- a/src/tint/lang/spirv/spirv.def
+++ b/src/tint/lang/spirv/spirv.def
@@ -1831,3 +1831,8 @@
@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>
+@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>
+
diff --git a/src/tint/lang/spirv/writer/printer/printer.cc b/src/tint/lang/spirv/writer/printer/printer.cc
index 27d95d3..31917b1 100644
--- a/src/tint/lang/spirv/writer/printer/printer.cc
+++ b/src/tint/lang/spirv/writer/printer/printer.cc
@@ -1694,6 +1694,9 @@
case BuiltinFn::kGroupNonUniformQuadBroadcast:
op = spv::Op::OpGroupNonUniformQuadBroadcast;
break;
+ case BuiltinFn::kGroupNonUniformQuadSwap:
+ op = spv::Op::OpGroupNonUniformQuadSwap;
+ break;
case spirv::BuiltinFn::kNone:
TINT_ICE() << "undefined spirv ir function";
}
