[tint] Add f16 overload of subgroupBroadcast()
Bug: tint:2041
Change-Id: Ib3b8b832318e97f2e130c50565030d7133027bb3
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/172901
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: James Price <jrprice@google.com>
diff --git a/src/tint/lang/core/core.def b/src/tint/lang/core/core.def
index 4a5ed76..1e57b92 100644
--- a/src/tint/lang/core/core.def
+++ b/src/tint/lang/core/core.def
@@ -851,7 +851,7 @@
@stage("fragment", "compute") fn atomicCompareExchangeWeak<T: iu32, S: workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T, T) -> __atomic_compare_exchange_result<T>
@must_use @stage("compute") fn subgroupBallot() -> vec4<u32>
-@must_use @stage("compute") fn subgroupBroadcast<T: fiu32>(value: T, @const sourceLaneIndex: u32) -> T
+@must_use @stage("compute") fn subgroupBroadcast<T: fiu32_f16>(value: T, @const sourceLaneIndex: u32) -> T
////////////////////////////////////////////////////////////////////////////////
// Value constructors //
diff --git a/src/tint/lang/core/intrinsic/data.cc b/src/tint/lang/core/intrinsic/data.cc
index 97482d5..a127690 100644
--- a/src/tint/lang/core/intrinsic/data.cc
+++ b/src/tint/lang/core/intrinsic/data.cc
@@ -4686,41 +4686,46 @@
{
/* [29] */
/* name */ "T",
- /* matcher_index */ TypeMatcherIndex(65),
+ /* matcher_index */ TypeMatcherIndex(67),
},
{
/* [30] */
/* name */ "T",
- /* matcher_index */ TypeMatcherIndex(57),
+ /* matcher_index */ TypeMatcherIndex(65),
},
{
/* [31] */
/* name */ "T",
- /* matcher_index */ TypeMatcherIndex(58),
+ /* matcher_index */ TypeMatcherIndex(57),
},
{
/* [32] */
/* name */ "T",
- /* matcher_index */ TypeMatcherIndex(55),
+ /* matcher_index */ TypeMatcherIndex(58),
},
{
/* [33] */
/* name */ "T",
- /* matcher_index */ TypeMatcherIndex(56),
+ /* matcher_index */ TypeMatcherIndex(55),
},
{
/* [34] */
/* name */ "T",
- /* matcher_index */ TypeMatcherIndex(59),
+ /* matcher_index */ TypeMatcherIndex(56),
},
{
/* [35] */
/* name */ "T",
- /* matcher_index */ TypeMatcherIndex(54),
+ /* matcher_index */ TypeMatcherIndex(59),
},
{
/* [36] */
/* name */ "T",
+ /* matcher_index */ TypeMatcherIndex(54),
+ },
+ {
+ /* [37] */
+ /* name */ "T",
/* matcher_index */ TypeMatcherIndex(71),
},
};
@@ -5540,7 +5545,7 @@
/* num_parameters */ 0,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(35),
+ /* template_types */ TemplateTypeIndex(36),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(52),
@@ -6437,7 +6442,7 @@
/* num_parameters */ 0,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(35),
+ /* template_types */ TemplateTypeIndex(36),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(10),
@@ -6723,7 +6728,7 @@
/* num_parameters */ 0,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(35),
+ /* template_types */ TemplateTypeIndex(36),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(72),
@@ -7399,7 +7404,7 @@
/* num_parameters */ 0,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(102),
@@ -7412,7 +7417,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(385),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(102),
@@ -7477,7 +7482,7 @@
/* num_parameters */ 0,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(108),
@@ -7490,7 +7495,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(388),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(108),
@@ -7555,7 +7560,7 @@
/* num_parameters */ 0,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(114),
@@ -7568,7 +7573,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(391),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(114),
@@ -7633,7 +7638,7 @@
/* num_parameters */ 0,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(120),
@@ -7646,7 +7651,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(394),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(120),
@@ -7711,7 +7716,7 @@
/* num_parameters */ 0,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(126),
@@ -7724,7 +7729,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(397),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(126),
@@ -7789,7 +7794,7 @@
/* num_parameters */ 0,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(132),
@@ -7802,7 +7807,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(400),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(132),
@@ -7867,7 +7872,7 @@
/* num_parameters */ 0,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(138),
@@ -7880,7 +7885,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(403),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(138),
@@ -7945,7 +7950,7 @@
/* num_parameters */ 0,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(144),
@@ -7958,7 +7963,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(406),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(144),
@@ -8023,7 +8028,7 @@
/* num_parameters */ 0,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(150),
@@ -8036,7 +8041,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(409),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(150),
@@ -8426,7 +8431,7 @@
/* num_parameters */ 2,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(29),
+ /* template_types */ TemplateTypeIndex(30),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(1),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(2),
@@ -8439,7 +8444,7 @@
/* num_parameters */ 2,
/* num_template_types */ 1,
/* num_template_numbers */ 1,
- /* template_types */ TemplateTypeIndex(29),
+ /* template_types */ TemplateTypeIndex(30),
/* template_numbers */ TemplateNumberIndex(4),
/* parameters */ ParameterIndex(149),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(6),
@@ -8478,7 +8483,7 @@
/* num_parameters */ 2,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(29),
+ /* template_types */ TemplateTypeIndex(30),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(1),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(2),
@@ -8491,7 +8496,7 @@
/* num_parameters */ 2,
/* num_template_types */ 1,
/* num_template_numbers */ 1,
- /* template_types */ TemplateTypeIndex(29),
+ /* template_types */ TemplateTypeIndex(30),
/* template_numbers */ TemplateNumberIndex(4),
/* parameters */ ParameterIndex(149),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(6),
@@ -8608,7 +8613,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(30),
+ /* template_types */ TemplateTypeIndex(31),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(1),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(31),
@@ -8647,7 +8652,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(31),
+ /* template_types */ TemplateTypeIndex(32),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(1),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(33),
@@ -8686,7 +8691,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(32),
+ /* template_types */ TemplateTypeIndex(33),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(1),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(15),
@@ -8725,7 +8730,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(33),
+ /* template_types */ TemplateTypeIndex(34),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(1),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(85),
@@ -8764,7 +8769,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(34),
+ /* template_types */ TemplateTypeIndex(35),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(1),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(9),
@@ -10233,7 +10238,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(29),
+ /* template_types */ TemplateTypeIndex(30),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(1),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(2),
@@ -10246,7 +10251,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 1,
- /* template_types */ TemplateTypeIndex(29),
+ /* template_types */ TemplateTypeIndex(30),
/* template_numbers */ TemplateNumberIndex(4),
/* parameters */ ParameterIndex(149),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(6),
@@ -10285,7 +10290,7 @@
/* num_parameters */ 2,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(29),
+ /* template_types */ TemplateTypeIndex(30),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(1),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(2),
@@ -10298,7 +10303,7 @@
/* num_parameters */ 2,
/* num_template_types */ 1,
/* num_template_numbers */ 1,
- /* template_types */ TemplateTypeIndex(29),
+ /* template_types */ TemplateTypeIndex(30),
/* template_numbers */ TemplateNumberIndex(4),
/* parameters */ ParameterIndex(149),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(6),
@@ -10467,7 +10472,7 @@
/* num_parameters */ 2,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(29),
+ /* template_types */ TemplateTypeIndex(30),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(16),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(2),
@@ -10480,7 +10485,7 @@
/* num_parameters */ 2,
/* num_template_types */ 1,
/* num_template_numbers */ 1,
- /* template_types */ TemplateTypeIndex(29),
+ /* template_types */ TemplateTypeIndex(30),
/* template_numbers */ TemplateNumberIndex(4),
/* parameters */ ParameterIndex(351),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(6),
@@ -10493,7 +10498,7 @@
/* num_parameters */ 2,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(29),
+ /* template_types */ TemplateTypeIndex(30),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(16),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(2),
@@ -10506,7 +10511,7 @@
/* num_parameters */ 2,
/* num_template_types */ 1,
/* num_template_numbers */ 1,
- /* template_types */ TemplateTypeIndex(29),
+ /* template_types */ TemplateTypeIndex(30),
/* template_numbers */ TemplateNumberIndex(4),
/* parameters */ ParameterIndex(351),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(6),
@@ -10948,7 +10953,7 @@
/* num_parameters */ 2,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(0),
+ /* template_types */ TemplateTypeIndex(29),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(348),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(2),
@@ -10987,7 +10992,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(35),
+ /* template_types */ TemplateTypeIndex(36),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(213),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(156),
@@ -11926,7 +11931,7 @@
},
{
/* [120] */
- /* fn subgroupBroadcast<T : fiu32>(value: T, @const sourceLaneIndex: u32) -> T */
+ /* fn subgroupBroadcast<T : fiu32_f16>(value: T, @const sourceLaneIndex: u32) -> T */
/* num overloads */ 1,
/* overloads */ OverloadIndex(465),
},
diff --git a/src/tint/lang/wgsl/intrinsic/data.cc b/src/tint/lang/wgsl/intrinsic/data.cc
index 73b6fa4..8266687 100644
--- a/src/tint/lang/wgsl/intrinsic/data.cc
+++ b/src/tint/lang/wgsl/intrinsic/data.cc
@@ -4698,41 +4698,46 @@
{
/* [29] */
/* name */ "T",
- /* matcher_index */ TypeMatcherIndex(65),
+ /* matcher_index */ TypeMatcherIndex(67),
},
{
/* [30] */
/* name */ "T",
- /* matcher_index */ TypeMatcherIndex(57),
+ /* matcher_index */ TypeMatcherIndex(65),
},
{
/* [31] */
/* name */ "T",
- /* matcher_index */ TypeMatcherIndex(58),
+ /* matcher_index */ TypeMatcherIndex(57),
},
{
/* [32] */
/* name */ "T",
- /* matcher_index */ TypeMatcherIndex(55),
+ /* matcher_index */ TypeMatcherIndex(58),
},
{
/* [33] */
/* name */ "T",
- /* matcher_index */ TypeMatcherIndex(56),
+ /* matcher_index */ TypeMatcherIndex(55),
},
{
/* [34] */
/* name */ "T",
- /* matcher_index */ TypeMatcherIndex(59),
+ /* matcher_index */ TypeMatcherIndex(56),
},
{
/* [35] */
/* name */ "T",
- /* matcher_index */ TypeMatcherIndex(54),
+ /* matcher_index */ TypeMatcherIndex(59),
},
{
/* [36] */
/* name */ "T",
+ /* matcher_index */ TypeMatcherIndex(54),
+ },
+ {
+ /* [37] */
+ /* name */ "T",
/* matcher_index */ TypeMatcherIndex(71),
},
};
@@ -5554,7 +5559,7 @@
/* num_parameters */ 0,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(35),
+ /* template_types */ TemplateTypeIndex(36),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(54),
@@ -6451,7 +6456,7 @@
/* num_parameters */ 0,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(35),
+ /* template_types */ TemplateTypeIndex(36),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(10),
@@ -6737,7 +6742,7 @@
/* num_parameters */ 0,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(35),
+ /* template_types */ TemplateTypeIndex(36),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(74),
@@ -7413,7 +7418,7 @@
/* num_parameters */ 0,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(104),
@@ -7426,7 +7431,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(386),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(104),
@@ -7491,7 +7496,7 @@
/* num_parameters */ 0,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(110),
@@ -7504,7 +7509,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(389),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(110),
@@ -7569,7 +7574,7 @@
/* num_parameters */ 0,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(116),
@@ -7582,7 +7587,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(392),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(116),
@@ -7647,7 +7652,7 @@
/* num_parameters */ 0,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(122),
@@ -7660,7 +7665,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(395),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(122),
@@ -7725,7 +7730,7 @@
/* num_parameters */ 0,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(128),
@@ -7738,7 +7743,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(398),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(128),
@@ -7803,7 +7808,7 @@
/* num_parameters */ 0,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(134),
@@ -7816,7 +7821,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(401),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(134),
@@ -7881,7 +7886,7 @@
/* num_parameters */ 0,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(140),
@@ -7894,7 +7899,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(404),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(140),
@@ -7959,7 +7964,7 @@
/* num_parameters */ 0,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(146),
@@ -7972,7 +7977,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(407),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(146),
@@ -8037,7 +8042,7 @@
/* num_parameters */ 0,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(152),
@@ -8050,7 +8055,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(36),
+ /* template_types */ TemplateTypeIndex(37),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(410),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(152),
@@ -8440,7 +8445,7 @@
/* num_parameters */ 2,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(29),
+ /* template_types */ TemplateTypeIndex(30),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(1),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(2),
@@ -8453,7 +8458,7 @@
/* num_parameters */ 2,
/* num_template_types */ 1,
/* num_template_numbers */ 1,
- /* template_types */ TemplateTypeIndex(29),
+ /* template_types */ TemplateTypeIndex(30),
/* template_numbers */ TemplateNumberIndex(4),
/* parameters */ ParameterIndex(149),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(6),
@@ -8492,7 +8497,7 @@
/* num_parameters */ 2,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(29),
+ /* template_types */ TemplateTypeIndex(30),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(1),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(2),
@@ -8505,7 +8510,7 @@
/* num_parameters */ 2,
/* num_template_types */ 1,
/* num_template_numbers */ 1,
- /* template_types */ TemplateTypeIndex(29),
+ /* template_types */ TemplateTypeIndex(30),
/* template_numbers */ TemplateNumberIndex(4),
/* parameters */ ParameterIndex(149),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(6),
@@ -8622,7 +8627,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(30),
+ /* template_types */ TemplateTypeIndex(31),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(1),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(31),
@@ -8661,7 +8666,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(31),
+ /* template_types */ TemplateTypeIndex(32),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(1),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(33),
@@ -8700,7 +8705,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(32),
+ /* template_types */ TemplateTypeIndex(33),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(1),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(15),
@@ -8739,7 +8744,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(33),
+ /* template_types */ TemplateTypeIndex(34),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(1),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(87),
@@ -8778,7 +8783,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(34),
+ /* template_types */ TemplateTypeIndex(35),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(1),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(9),
@@ -10247,7 +10252,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(29),
+ /* template_types */ TemplateTypeIndex(30),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(1),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(2),
@@ -10260,7 +10265,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 1,
- /* template_types */ TemplateTypeIndex(29),
+ /* template_types */ TemplateTypeIndex(30),
/* template_numbers */ TemplateNumberIndex(4),
/* parameters */ ParameterIndex(149),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(6),
@@ -10299,7 +10304,7 @@
/* num_parameters */ 2,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(29),
+ /* template_types */ TemplateTypeIndex(30),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(1),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(2),
@@ -10312,7 +10317,7 @@
/* num_parameters */ 2,
/* num_template_types */ 1,
/* num_template_numbers */ 1,
- /* template_types */ TemplateTypeIndex(29),
+ /* template_types */ TemplateTypeIndex(30),
/* template_numbers */ TemplateNumberIndex(4),
/* parameters */ ParameterIndex(149),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(6),
@@ -10481,7 +10486,7 @@
/* num_parameters */ 2,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(29),
+ /* template_types */ TemplateTypeIndex(30),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(16),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(2),
@@ -10494,7 +10499,7 @@
/* num_parameters */ 2,
/* num_template_types */ 1,
/* num_template_numbers */ 1,
- /* template_types */ TemplateTypeIndex(29),
+ /* template_types */ TemplateTypeIndex(30),
/* template_numbers */ TemplateNumberIndex(4),
/* parameters */ ParameterIndex(351),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(6),
@@ -10507,7 +10512,7 @@
/* num_parameters */ 2,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(29),
+ /* template_types */ TemplateTypeIndex(30),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(16),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(2),
@@ -10520,7 +10525,7 @@
/* num_parameters */ 2,
/* num_template_types */ 1,
/* num_template_numbers */ 1,
- /* template_types */ TemplateTypeIndex(29),
+ /* template_types */ TemplateTypeIndex(30),
/* template_numbers */ TemplateNumberIndex(4),
/* parameters */ ParameterIndex(351),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(6),
@@ -10975,7 +10980,7 @@
/* num_parameters */ 2,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(0),
+ /* template_types */ TemplateTypeIndex(29),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(348),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(2),
@@ -11027,7 +11032,7 @@
/* num_parameters */ 1,
/* num_template_types */ 1,
/* num_template_numbers */ 0,
- /* template_types */ TemplateTypeIndex(35),
+ /* template_types */ TemplateTypeIndex(36),
/* template_numbers */ TemplateNumberIndex(/* invalid */),
/* parameters */ ParameterIndex(213),
/* return_type_matcher_indices */ TypeMatcherIndicesIndex(158),
@@ -11972,7 +11977,7 @@
},
{
/* [121] */
- /* fn subgroupBroadcast<T : fiu32>(value: T, @const sourceLaneIndex: u32) -> T */
+ /* fn subgroupBroadcast<T : fiu32_f16>(value: T, @const sourceLaneIndex: u32) -> T */
/* num overloads */ 1,
/* overloads */ OverloadIndex(466),
},
diff --git a/src/tint/lang/wgsl/wgsl.def b/src/tint/lang/wgsl/wgsl.def
index e23de8a..3e3282c 100644
--- a/src/tint/lang/wgsl/wgsl.def
+++ b/src/tint/lang/wgsl/wgsl.def
@@ -736,7 +736,7 @@
@stage("fragment", "compute") fn atomicCompareExchangeWeak<T: iu32, S: workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T, T) -> __atomic_compare_exchange_result<T>
@must_use @stage("compute") fn subgroupBallot() -> vec4<u32>
-@must_use @stage("compute") fn subgroupBroadcast<T: fiu32>(value: T, @const sourceLaneIndex: u32) -> T
+@must_use @stage("compute") fn subgroupBroadcast<T: fiu32_f16>(value: T, @const sourceLaneIndex: u32) -> T
////////////////////////////////////////////////////////////////////////////////
// Value constructors //
diff --git a/test/tint/builtins/gen/literal/subgroupBroadcast/07e2d8.wgsl b/test/tint/builtins/gen/literal/subgroupBroadcast/07e2d8.wgsl
new file mode 100644
index 0000000..1b0957f
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupBroadcast/07e2d8.wgsl
@@ -0,0 +1,55 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by 'tools/src/cmd/gen' using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// To regenerate run: './tools/run gen'
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// flags: --hlsl_shader_model 62
+
+
+enable chromium_experimental_subgroups;
+
+enable f16;
+
+// fn subgroupBroadcast(value: f16, @const sourceLaneIndex: u32) -> f16
+fn subgroupBroadcast_07e2d8() {
+ var res: f16 = subgroupBroadcast(1.h, 1u);
+ prevent_dce = res;
+}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : f16;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupBroadcast_07e2d8();
+}
diff --git a/test/tint/builtins/gen/literal/subgroupBroadcast/07e2d8.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/subgroupBroadcast/07e2d8.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..b0e0844
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupBroadcast/07e2d8.wgsl.expected.dxc.hlsl
@@ -0,0 +1,12 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void subgroupBroadcast_07e2d8() {
+ float16_t res = WaveReadLaneAt(float16_t(1.0h), 1u);
+ prevent_dce.Store<float16_t>(0u, res);
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ subgroupBroadcast_07e2d8();
+ return;
+}
diff --git a/test/tint/builtins/gen/literal/subgroupBroadcast/07e2d8.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupBroadcast/07e2d8.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..157e3dd
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupBroadcast/07e2d8.wgsl.expected.fxc.hlsl
@@ -0,0 +1,14 @@
+SKIP: FAILED
+
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void subgroupBroadcast_07e2d8() {
+ float16_t res = WaveReadLaneAt(float16_t(1.0h), 1u);
+ prevent_dce.Store<float16_t>(0u, res);
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ subgroupBroadcast_07e2d8();
+ return;
+}
diff --git a/test/tint/builtins/gen/literal/subgroupBroadcast/07e2d8.wgsl.expected.glsl b/test/tint/builtins/gen/literal/subgroupBroadcast/07e2d8.wgsl.expected.glsl
new file mode 100644
index 0000000..2d78f62
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupBroadcast/07e2d8.wgsl.expected.glsl
@@ -0,0 +1,22 @@
+SKIP: FAILED
+
+
+enable chromium_experimental_subgroups;
+enable f16;
+
+fn subgroupBroadcast_07e2d8() {
+ var res : f16 = subgroupBroadcast(1.0h, 1u);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : f16;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupBroadcast_07e2d8();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/literal/subgroupBroadcast/07e2d8.wgsl:41:8 error: GLSL backend does not support extension 'chromium_experimental_subgroups'
+enable chromium_experimental_subgroups;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/literal/subgroupBroadcast/07e2d8.wgsl.expected.msl b/test/tint/builtins/gen/literal/subgroupBroadcast/07e2d8.wgsl.expected.msl
new file mode 100644
index 0000000..b6a3aa7
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupBroadcast/07e2d8.wgsl.expected.msl
@@ -0,0 +1,13 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void subgroupBroadcast_07e2d8(device half* const tint_symbol) {
+ half res = simd_broadcast(1.0h,ushort(1u));
+ *(tint_symbol) = res;
+}
+
+kernel void compute_main(device half* tint_symbol_1 [[buffer(0)]]) {
+ subgroupBroadcast_07e2d8(tint_symbol_1);
+ return;
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupBroadcast/07e2d8.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/subgroupBroadcast/07e2d8.wgsl.expected.spvasm
new file mode 100644
index 0000000..f5ff160
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupBroadcast/07e2d8.wgsl.expected.spvasm
@@ -0,0 +1,53 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 24
+; Schema: 0
+ OpCapability Shader
+ OpCapability Float16
+ OpCapability UniformAndStorageBuffer16BitAccess
+ OpCapability StorageBuffer16BitAccess
+ OpCapability StorageInputOutput16
+ OpCapability GroupNonUniformBallot
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %compute_main "compute_main"
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %prevent_dce_block "prevent_dce_block"
+ OpMemberName %prevent_dce_block 0 "inner"
+ OpName %prevent_dce "prevent_dce"
+ OpName %subgroupBroadcast_07e2d8 "subgroupBroadcast_07e2d8"
+ OpName %res "res"
+ OpName %compute_main "compute_main"
+ OpDecorate %prevent_dce_block Block
+ OpMemberDecorate %prevent_dce_block 0 Offset 0
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
+ %half = OpTypeFloat 16
+%prevent_dce_block = OpTypeStruct %half
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+ %void = OpTypeVoid
+ %5 = OpTypeFunction %void
+%half_0x1p_0 = OpConstant %half 0x1p+0
+ %uint = OpTypeInt 32 0
+ %uint_1 = OpConstant %uint 1
+ %uint_3 = OpConstant %uint 3
+%_ptr_Function_half = OpTypePointer Function %half
+ %16 = OpConstantNull %half
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_half = OpTypePointer StorageBuffer %half
+%subgroupBroadcast_07e2d8 = OpFunction %void None %5
+ %8 = OpLabel
+ %res = OpVariable %_ptr_Function_half Function %16
+ %9 = OpGroupNonUniformBroadcast %half %uint_3 %half_0x1p_0 %uint_1
+ OpStore %res %9
+ %19 = OpAccessChain %_ptr_StorageBuffer_half %prevent_dce %uint_0
+ %20 = OpLoad %half %res
+ OpStore %19 %20
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %5
+ %22 = OpLabel
+ %23 = OpFunctionCall %void %subgroupBroadcast_07e2d8
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/subgroupBroadcast/07e2d8.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupBroadcast/07e2d8.wgsl.expected.wgsl
new file mode 100644
index 0000000..763d168
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupBroadcast/07e2d8.wgsl.expected.wgsl
@@ -0,0 +1,14 @@
+enable chromium_experimental_subgroups;
+enable f16;
+
+fn subgroupBroadcast_07e2d8() {
+ var res : f16 = subgroupBroadcast(1.0h, 1u);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : f16;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupBroadcast_07e2d8();
+}
diff --git a/test/tint/builtins/gen/var/subgroupBroadcast/07e2d8.wgsl b/test/tint/builtins/gen/var/subgroupBroadcast/07e2d8.wgsl
new file mode 100644
index 0000000..b7f1c10
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupBroadcast/07e2d8.wgsl
@@ -0,0 +1,57 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by 'tools/src/cmd/gen' using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// To regenerate run: './tools/run gen'
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// flags: --hlsl_shader_model 62
+
+
+enable chromium_experimental_subgroups;
+
+enable f16;
+
+// fn subgroupBroadcast(value: f16, @const sourceLaneIndex: u32) -> f16
+fn subgroupBroadcast_07e2d8() {
+ var arg_0 = 1.h;
+ const arg_1 = 1u;
+ var res: f16 = subgroupBroadcast(arg_0, arg_1);
+ prevent_dce = res;
+}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : f16;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupBroadcast_07e2d8();
+}
diff --git a/test/tint/builtins/gen/var/subgroupBroadcast/07e2d8.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/subgroupBroadcast/07e2d8.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..281c550c
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupBroadcast/07e2d8.wgsl.expected.dxc.hlsl
@@ -0,0 +1,13 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void subgroupBroadcast_07e2d8() {
+ float16_t arg_0 = float16_t(1.0h);
+ float16_t res = WaveReadLaneAt(arg_0, 1u);
+ prevent_dce.Store<float16_t>(0u, res);
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ subgroupBroadcast_07e2d8();
+ return;
+}
diff --git a/test/tint/builtins/gen/var/subgroupBroadcast/07e2d8.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/subgroupBroadcast/07e2d8.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..6980fb0
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupBroadcast/07e2d8.wgsl.expected.fxc.hlsl
@@ -0,0 +1,15 @@
+SKIP: FAILED
+
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void subgroupBroadcast_07e2d8() {
+ float16_t arg_0 = float16_t(1.0h);
+ float16_t res = WaveReadLaneAt(arg_0, 1u);
+ prevent_dce.Store<float16_t>(0u, res);
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ subgroupBroadcast_07e2d8();
+ return;
+}
diff --git a/test/tint/builtins/gen/var/subgroupBroadcast/07e2d8.wgsl.expected.glsl b/test/tint/builtins/gen/var/subgroupBroadcast/07e2d8.wgsl.expected.glsl
new file mode 100644
index 0000000..6ab776a
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupBroadcast/07e2d8.wgsl.expected.glsl
@@ -0,0 +1,24 @@
+SKIP: FAILED
+
+
+enable chromium_experimental_subgroups;
+enable f16;
+
+fn subgroupBroadcast_07e2d8() {
+ var arg_0 = 1.0h;
+ const arg_1 = 1u;
+ var res : f16 = subgroupBroadcast(arg_0, arg_1);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : f16;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupBroadcast_07e2d8();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/var/subgroupBroadcast/07e2d8.wgsl:41:8 error: GLSL backend does not support extension 'chromium_experimental_subgroups'
+enable chromium_experimental_subgroups;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/var/subgroupBroadcast/07e2d8.wgsl.expected.msl b/test/tint/builtins/gen/var/subgroupBroadcast/07e2d8.wgsl.expected.msl
new file mode 100644
index 0000000..0d11a88
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupBroadcast/07e2d8.wgsl.expected.msl
@@ -0,0 +1,14 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void subgroupBroadcast_07e2d8(device half* const tint_symbol) {
+ half arg_0 = 1.0h;
+ half res = simd_broadcast(arg_0,ushort(1u));
+ *(tint_symbol) = res;
+}
+
+kernel void compute_main(device half* tint_symbol_1 [[buffer(0)]]) {
+ subgroupBroadcast_07e2d8(tint_symbol_1);
+ return;
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupBroadcast/07e2d8.wgsl.expected.spvasm b/test/tint/builtins/gen/var/subgroupBroadcast/07e2d8.wgsl.expected.spvasm
new file mode 100644
index 0000000..a59bcd4
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupBroadcast/07e2d8.wgsl.expected.spvasm
@@ -0,0 +1,57 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 26
+; Schema: 0
+ OpCapability Shader
+ OpCapability Float16
+ OpCapability UniformAndStorageBuffer16BitAccess
+ OpCapability StorageBuffer16BitAccess
+ OpCapability StorageInputOutput16
+ OpCapability GroupNonUniformBallot
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %compute_main "compute_main"
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %prevent_dce_block "prevent_dce_block"
+ OpMemberName %prevent_dce_block 0 "inner"
+ OpName %prevent_dce "prevent_dce"
+ OpName %subgroupBroadcast_07e2d8 "subgroupBroadcast_07e2d8"
+ OpName %arg_0 "arg_0"
+ OpName %res "res"
+ OpName %compute_main "compute_main"
+ OpDecorate %prevent_dce_block Block
+ OpMemberDecorate %prevent_dce_block 0 Offset 0
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
+ %half = OpTypeFloat 16
+%prevent_dce_block = OpTypeStruct %half
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+ %void = OpTypeVoid
+ %5 = OpTypeFunction %void
+%half_0x1p_0 = OpConstant %half 0x1p+0
+%_ptr_Function_half = OpTypePointer Function %half
+ %12 = OpConstantNull %half
+ %uint = OpTypeInt 32 0
+ %uint_1 = OpConstant %uint 1
+ %uint_3 = OpConstant %uint 3
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_half = OpTypePointer StorageBuffer %half
+%subgroupBroadcast_07e2d8 = OpFunction %void None %5
+ %8 = OpLabel
+ %arg_0 = OpVariable %_ptr_Function_half Function %12
+ %res = OpVariable %_ptr_Function_half Function %12
+ OpStore %arg_0 %half_0x1p_0
+ %14 = OpLoad %half %arg_0
+ %13 = OpGroupNonUniformBroadcast %half %uint_3 %14 %uint_1
+ OpStore %res %13
+ %21 = OpAccessChain %_ptr_StorageBuffer_half %prevent_dce %uint_0
+ %22 = OpLoad %half %res
+ OpStore %21 %22
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %5
+ %24 = OpLabel
+ %25 = OpFunctionCall %void %subgroupBroadcast_07e2d8
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/subgroupBroadcast/07e2d8.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupBroadcast/07e2d8.wgsl.expected.wgsl
new file mode 100644
index 0000000..7ef010d
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupBroadcast/07e2d8.wgsl.expected.wgsl
@@ -0,0 +1,16 @@
+enable chromium_experimental_subgroups;
+enable f16;
+
+fn subgroupBroadcast_07e2d8() {
+ var arg_0 = 1.0h;
+ const arg_1 = 1u;
+ var res : f16 = subgroupBroadcast(arg_0, arg_1);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : f16;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupBroadcast_07e2d8();
+}
