[msl] Avoid generating packed_bool3 types
These are reserved in MSL and are causing issues on Intel drivers.
Bug: 424772881
Change-Id: I3a0819f0928cc544c56202411df951e9908e2485
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/247856
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Commit-Queue: James Price <jrprice@google.com>
diff --git a/src/tint/lang/msl/intrinsic/data.cc b/src/tint/lang/msl/intrinsic/data.cc
index bda5b13..9435456 100644
--- a/src/tint/lang/msl/intrinsic/data.cc
+++ b/src/tint/lang/msl/intrinsic/data.cc
@@ -1307,39 +1307,35 @@
/* [168] */ MatcherIndex(10),
/* [169] */ MatcherIndex(25),
/* [170] */ MatcherIndex(10),
- /* [171] */ MatcherIndex(13),
- /* [172] */ MatcherIndex(4),
+ /* [171] */ MatcherIndex(42),
+ /* [172] */ MatcherIndex(7),
/* [173] */ MatcherIndex(42),
- /* [174] */ MatcherIndex(4),
+ /* [174] */ MatcherIndex(5),
/* [175] */ MatcherIndex(42),
- /* [176] */ MatcherIndex(7),
- /* [177] */ MatcherIndex(42),
- /* [178] */ MatcherIndex(5),
+ /* [176] */ MatcherIndex(10),
+ /* [177] */ MatcherIndex(13),
+ /* [178] */ MatcherIndex(11),
/* [179] */ MatcherIndex(42),
- /* [180] */ MatcherIndex(10),
- /* [181] */ MatcherIndex(13),
- /* [182] */ MatcherIndex(11),
- /* [183] */ MatcherIndex(42),
- /* [184] */ MatcherIndex(11),
- /* [185] */ MatcherIndex(12),
- /* [186] */ MatcherIndex(8),
- /* [187] */ MatcherIndex(43),
- /* [188] */ MatcherIndex(44),
- /* [189] */ MatcherIndex(18),
- /* [190] */ MatcherIndex(26),
- /* [191] */ MatcherIndex(27),
- /* [192] */ MatcherIndex(28),
- /* [193] */ MatcherIndex(29),
- /* [194] */ MatcherIndex(19),
- /* [195] */ MatcherIndex(30),
- /* [196] */ MatcherIndex(41),
- /* [197] */ MatcherIndex(37),
- /* [198] */ MatcherIndex(38),
- /* [199] */ MatcherIndex(39),
- /* [200] */ MatcherIndex(40),
- /* [201] */ MatcherIndex(45),
- /* [202] */ MatcherIndex(46),
- /* [203] */ MatcherIndex(47),
+ /* [180] */ MatcherIndex(11),
+ /* [181] */ MatcherIndex(12),
+ /* [182] */ MatcherIndex(8),
+ /* [183] */ MatcherIndex(43),
+ /* [184] */ MatcherIndex(44),
+ /* [185] */ MatcherIndex(18),
+ /* [186] */ MatcherIndex(26),
+ /* [187] */ MatcherIndex(27),
+ /* [188] */ MatcherIndex(28),
+ /* [189] */ MatcherIndex(29),
+ /* [190] */ MatcherIndex(19),
+ /* [191] */ MatcherIndex(30),
+ /* [192] */ MatcherIndex(41),
+ /* [193] */ MatcherIndex(37),
+ /* [194] */ MatcherIndex(38),
+ /* [195] */ MatcherIndex(39),
+ /* [196] */ MatcherIndex(40),
+ /* [197] */ MatcherIndex(45),
+ /* [198] */ MatcherIndex(46),
+ /* [199] */ MatcherIndex(47),
};
static_assert(MatcherIndicesIndex::CanIndex(kMatcherIndices),
@@ -1349,12 +1345,12 @@
{
/* [0] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(191),
+ /* matcher_indices */ MatcherIndicesIndex(187),
},
{
/* [1] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(194),
+ /* matcher_indices */ MatcherIndicesIndex(190),
},
{
/* [2] */
@@ -1374,7 +1370,7 @@
{
/* [5] */
/* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(196),
+ /* matcher_indices */ MatcherIndicesIndex(192),
},
{
/* [6] */
@@ -1389,7 +1385,7 @@
{
/* [8] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [9] */
@@ -1414,12 +1410,12 @@
{
/* [13] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(191),
+ /* matcher_indices */ MatcherIndicesIndex(187),
},
{
/* [14] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(194),
+ /* matcher_indices */ MatcherIndicesIndex(190),
},
{
/* [15] */
@@ -1449,7 +1445,7 @@
{
/* [20] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [21] */
@@ -1464,7 +1460,7 @@
{
/* [23] */
/* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(196),
+ /* matcher_indices */ MatcherIndicesIndex(192),
},
{
/* [24] */
@@ -1474,12 +1470,12 @@
{
/* [25] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(191),
+ /* matcher_indices */ MatcherIndicesIndex(187),
},
{
/* [26] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [27] */
@@ -1494,7 +1490,7 @@
{
/* [29] */
/* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(196),
+ /* matcher_indices */ MatcherIndicesIndex(192),
},
{
/* [30] */
@@ -1509,7 +1505,7 @@
{
/* [32] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [33] */
@@ -1524,7 +1520,7 @@
{
/* [35] */
/* usage */ core::ParameterUsage::kBias,
- /* matcher_indices */ MatcherIndicesIndex(197),
+ /* matcher_indices */ MatcherIndicesIndex(193),
},
{
/* [36] */
@@ -1539,7 +1535,7 @@
{
/* [38] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [39] */
@@ -1554,7 +1550,7 @@
{
/* [41] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(198),
+ /* matcher_indices */ MatcherIndicesIndex(194),
},
{
/* [42] */
@@ -1564,12 +1560,12 @@
{
/* [43] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(190),
+ /* matcher_indices */ MatcherIndicesIndex(186),
},
{
/* [44] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(194),
+ /* matcher_indices */ MatcherIndicesIndex(190),
},
{
/* [45] */
@@ -1584,7 +1580,7 @@
{
/* [47] */
/* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(196),
+ /* matcher_indices */ MatcherIndicesIndex(192),
},
{
/* [48] */
@@ -1594,12 +1590,12 @@
{
/* [49] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(193),
+ /* matcher_indices */ MatcherIndicesIndex(189),
},
{
/* [50] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(194),
+ /* matcher_indices */ MatcherIndicesIndex(190),
},
{
/* [51] */
@@ -1619,7 +1615,7 @@
{
/* [54] */
/* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(196),
+ /* matcher_indices */ MatcherIndicesIndex(192),
},
{
/* [55] */
@@ -1654,7 +1650,7 @@
{
/* [61] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [62] */
@@ -1679,7 +1675,7 @@
{
/* [66] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [67] */
@@ -1699,12 +1695,12 @@
{
/* [70] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(191),
+ /* matcher_indices */ MatcherIndicesIndex(187),
},
{
/* [71] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [72] */
@@ -1724,12 +1720,12 @@
{
/* [75] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(190),
+ /* matcher_indices */ MatcherIndicesIndex(186),
},
{
/* [76] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(194),
+ /* matcher_indices */ MatcherIndicesIndex(190),
},
{
/* [77] */
@@ -1754,7 +1750,7 @@
{
/* [81] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [82] */
@@ -1779,7 +1775,7 @@
{
/* [86] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [87] */
@@ -1789,7 +1785,7 @@
{
/* [88] */
/* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(196),
+ /* matcher_indices */ MatcherIndicesIndex(192),
},
{
/* [89] */
@@ -1804,7 +1800,7 @@
{
/* [91] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [92] */
@@ -1814,7 +1810,7 @@
{
/* [93] */
/* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(196),
+ /* matcher_indices */ MatcherIndicesIndex(192),
},
{
/* [94] */
@@ -1829,7 +1825,7 @@
{
/* [96] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [97] */
@@ -1844,17 +1840,17 @@
{
/* [99] */
/* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(196),
+ /* matcher_indices */ MatcherIndicesIndex(192),
},
{
/* [100] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(190),
+ /* matcher_indices */ MatcherIndicesIndex(186),
},
{
/* [101] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [102] */
@@ -1864,7 +1860,7 @@
{
/* [103] */
/* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(196),
+ /* matcher_indices */ MatcherIndicesIndex(192),
},
{
/* [104] */
@@ -1874,12 +1870,12 @@
{
/* [105] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(193),
+ /* matcher_indices */ MatcherIndicesIndex(189),
},
{
/* [106] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [107] */
@@ -1894,7 +1890,7 @@
{
/* [109] */
/* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(196),
+ /* matcher_indices */ MatcherIndicesIndex(192),
},
{
/* [110] */
@@ -1904,7 +1900,7 @@
{
/* [111] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [112] */
@@ -1914,7 +1910,7 @@
{
/* [113] */
/* usage */ core::ParameterUsage::kBias,
- /* matcher_indices */ MatcherIndicesIndex(197),
+ /* matcher_indices */ MatcherIndicesIndex(193),
},
{
/* [114] */
@@ -1929,7 +1925,7 @@
{
/* [116] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [117] */
@@ -1939,7 +1935,7 @@
{
/* [118] */
/* usage */ core::ParameterUsage::kBias,
- /* matcher_indices */ MatcherIndicesIndex(197),
+ /* matcher_indices */ MatcherIndicesIndex(193),
},
{
/* [119] */
@@ -1954,7 +1950,7 @@
{
/* [121] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [122] */
@@ -1969,7 +1965,7 @@
{
/* [124] */
/* usage */ core::ParameterUsage::kBias,
- /* matcher_indices */ MatcherIndicesIndex(197),
+ /* matcher_indices */ MatcherIndicesIndex(193),
},
{
/* [125] */
@@ -1979,7 +1975,7 @@
{
/* [126] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [127] */
@@ -1989,7 +1985,7 @@
{
/* [128] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(198),
+ /* matcher_indices */ MatcherIndicesIndex(194),
},
{
/* [129] */
@@ -2004,7 +2000,7 @@
{
/* [131] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [132] */
@@ -2014,7 +2010,7 @@
{
/* [133] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(199),
+ /* matcher_indices */ MatcherIndicesIndex(195),
},
{
/* [134] */
@@ -2029,7 +2025,7 @@
{
/* [136] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [137] */
@@ -2044,17 +2040,17 @@
{
/* [139] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(200),
+ /* matcher_indices */ MatcherIndicesIndex(196),
},
{
/* [140] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(192),
+ /* matcher_indices */ MatcherIndicesIndex(188),
},
{
/* [141] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(194),
+ /* matcher_indices */ MatcherIndicesIndex(190),
},
{
/* [142] */
@@ -2069,7 +2065,7 @@
{
/* [144] */
/* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(196),
+ /* matcher_indices */ MatcherIndicesIndex(192),
},
{
/* [145] */
@@ -2089,7 +2085,7 @@
{
/* [148] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(185),
+ /* matcher_indices */ MatcherIndicesIndex(181),
},
{
/* [149] */
@@ -2114,7 +2110,7 @@
{
/* [153] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(185),
+ /* matcher_indices */ MatcherIndicesIndex(181),
},
{
/* [154] */
@@ -2129,7 +2125,7 @@
{
/* [156] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [157] */
@@ -2144,12 +2140,12 @@
{
/* [159] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(190),
+ /* matcher_indices */ MatcherIndicesIndex(186),
},
{
/* [160] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [161] */
@@ -2184,7 +2180,7 @@
{
/* [167] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(191),
+ /* matcher_indices */ MatcherIndicesIndex(187),
},
{
/* [168] */
@@ -2209,7 +2205,7 @@
{
/* [172] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [173] */
@@ -2229,7 +2225,7 @@
{
/* [176] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [177] */
@@ -2249,7 +2245,7 @@
{
/* [180] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [181] */
@@ -2259,17 +2255,17 @@
{
/* [182] */
/* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(196),
+ /* matcher_indices */ MatcherIndicesIndex(192),
},
{
/* [183] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(192),
+ /* matcher_indices */ MatcherIndicesIndex(188),
},
{
/* [184] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [185] */
@@ -2279,7 +2275,7 @@
{
/* [186] */
/* usage */ core::ParameterUsage::kLevel,
- /* matcher_indices */ MatcherIndicesIndex(196),
+ /* matcher_indices */ MatcherIndicesIndex(192),
},
{
/* [187] */
@@ -2289,7 +2285,7 @@
{
/* [188] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [189] */
@@ -2299,7 +2295,7 @@
{
/* [190] */
/* usage */ core::ParameterUsage::kBias,
- /* matcher_indices */ MatcherIndicesIndex(197),
+ /* matcher_indices */ MatcherIndicesIndex(193),
},
{
/* [191] */
@@ -2309,7 +2305,7 @@
{
/* [192] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [193] */
@@ -2319,7 +2315,7 @@
{
/* [194] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(200),
+ /* matcher_indices */ MatcherIndicesIndex(196),
},
{
/* [195] */
@@ -2464,7 +2460,7 @@
{
/* [223] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(190),
+ /* matcher_indices */ MatcherIndicesIndex(186),
},
{
/* [224] */
@@ -2479,7 +2475,7 @@
{
/* [226] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(195),
+ /* matcher_indices */ MatcherIndicesIndex(191),
},
{
/* [227] */
@@ -2544,7 +2540,7 @@
{
/* [239] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(189),
+ /* matcher_indices */ MatcherIndicesIndex(185),
},
{
/* [240] */
@@ -2749,7 +2745,7 @@
{
/* [280] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(190),
+ /* matcher_indices */ MatcherIndicesIndex(186),
},
{
/* [281] */
@@ -2759,7 +2755,7 @@
{
/* [282] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(191),
+ /* matcher_indices */ MatcherIndicesIndex(187),
},
{
/* [283] */
@@ -2769,7 +2765,7 @@
{
/* [284] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(192),
+ /* matcher_indices */ MatcherIndicesIndex(188),
},
{
/* [285] */
@@ -2779,7 +2775,7 @@
{
/* [286] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(193),
+ /* matcher_indices */ MatcherIndicesIndex(189),
},
{
/* [287] */
@@ -2974,17 +2970,17 @@
{
/* [325] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(173),
+ /* matcher_indices */ MatcherIndicesIndex(171),
},
{
/* [326] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(175),
+ /* matcher_indices */ MatcherIndicesIndex(173),
},
{
/* [327] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(177),
+ /* matcher_indices */ MatcherIndicesIndex(175),
},
{
/* [328] */
@@ -2994,32 +2990,22 @@
{
/* [329] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(183),
+ /* matcher_indices */ MatcherIndicesIndex(151),
},
{
/* [330] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(171),
+ /* matcher_indices */ MatcherIndicesIndex(165),
},
{
/* [331] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(151),
+ /* matcher_indices */ MatcherIndicesIndex(137),
},
{
/* [332] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(165),
- },
- {
- /* [333] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(137),
- },
- {
- /* [334] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(181),
+ /* matcher_indices */ MatcherIndicesIndex(177),
},
};
@@ -3036,7 +3022,7 @@
{
/* [1] */
/* name */ "S",
- /* matcher_indices */ MatcherIndicesIndex(201),
+ /* matcher_indices */ MatcherIndicesIndex(197),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -3060,25 +3046,25 @@
{
/* [5] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(188),
+ /* matcher_indices */ MatcherIndicesIndex(184),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [6] */
/* name */ "A",
- /* matcher_indices */ MatcherIndicesIndex(187),
+ /* matcher_indices */ MatcherIndicesIndex(183),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [7] */
/* name */ "L",
- /* matcher_indices */ MatcherIndicesIndex(187),
+ /* matcher_indices */ MatcherIndicesIndex(183),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [8] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(187),
+ /* matcher_indices */ MatcherIndicesIndex(183),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -3102,25 +3088,25 @@
{
/* [12] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(188),
+ /* matcher_indices */ MatcherIndicesIndex(184),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [13] */
/* name */ "L",
- /* matcher_indices */ MatcherIndicesIndex(187),
+ /* matcher_indices */ MatcherIndicesIndex(183),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [14] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(188),
+ /* matcher_indices */ MatcherIndicesIndex(184),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [15] */
/* name */ "S",
- /* matcher_indices */ MatcherIndicesIndex(187),
+ /* matcher_indices */ MatcherIndicesIndex(183),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -3132,7 +3118,7 @@
{
/* [17] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(201),
+ /* matcher_indices */ MatcherIndicesIndex(197),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -3144,13 +3130,13 @@
{
/* [19] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(202),
+ /* matcher_indices */ MatcherIndicesIndex(198),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [20] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(203),
+ /* matcher_indices */ MatcherIndicesIndex(199),
/* kind */ TemplateInfo::Kind::kType,
},
};
@@ -4470,127 +4456,6 @@
},
{
/* [119] */
- /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 1,
- /* num_explicit_templates */ 0,
- /* num_templates */ 0,
- /* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(325),
- /* return_matcher_indices */ MatcherIndicesIndex(171),
- /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
- },
- {
- /* [120] */
- /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 1,
- /* num_explicit_templates */ 0,
- /* num_templates */ 0,
- /* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(326),
- /* return_matcher_indices */ MatcherIndicesIndex(151),
- /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
- },
- {
- /* [121] */
- /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 1,
- /* num_explicit_templates */ 0,
- /* num_templates */ 0,
- /* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(327),
- /* return_matcher_indices */ MatcherIndicesIndex(165),
- /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
- },
- {
- /* [122] */
- /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 1,
- /* num_explicit_templates */ 0,
- /* num_templates */ 0,
- /* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(328),
- /* return_matcher_indices */ MatcherIndicesIndex(137),
- /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
- },
- {
- /* [123] */
- /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 1,
- /* num_explicit_templates */ 0,
- /* num_templates */ 0,
- /* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(329),
- /* return_matcher_indices */ MatcherIndicesIndex(181),
- /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
- },
- {
- /* [124] */
- /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 1,
- /* num_explicit_templates */ 0,
- /* num_templates */ 0,
- /* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(330),
- /* return_matcher_indices */ MatcherIndicesIndex(173),
- /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
- },
- {
- /* [125] */
- /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 1,
- /* num_explicit_templates */ 0,
- /* num_templates */ 0,
- /* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(331),
- /* return_matcher_indices */ MatcherIndicesIndex(175),
- /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
- },
- {
- /* [126] */
- /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 1,
- /* num_explicit_templates */ 0,
- /* num_templates */ 0,
- /* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(332),
- /* return_matcher_indices */ MatcherIndicesIndex(177),
- /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
- },
- {
- /* [127] */
- /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 1,
- /* num_explicit_templates */ 0,
- /* num_templates */ 0,
- /* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(333),
- /* return_matcher_indices */ MatcherIndicesIndex(179),
- /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
- },
- {
- /* [128] */
- /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 1,
- /* num_explicit_templates */ 0,
- /* num_templates */ 0,
- /* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(334),
- /* return_matcher_indices */ MatcherIndicesIndex(183),
- /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
- },
- {
- /* [129] */
- /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 1,
- /* num_explicit_templates */ 0,
- /* num_templates */ 0,
- /* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(58),
- /* return_matcher_indices */ MatcherIndicesIndex(30),
- /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
- },
- {
- /* [130] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -4601,7 +4466,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [131] */
+ /* [120] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -4612,7 +4477,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [132] */
+ /* [121] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -4623,7 +4488,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [133] */
+ /* [122] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -4634,7 +4499,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [134] */
+ /* [123] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -4645,7 +4510,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [135] */
+ /* [124] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -4656,7 +4521,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [136] */
+ /* [125] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -4667,7 +4532,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [137] */
+ /* [126] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -4678,7 +4543,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [138] */
+ /* [127] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -4689,7 +4554,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [139] */
+ /* [128] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -4700,7 +4565,106 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [140] */
+ /* [129] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 0,
+ /* templates */ TemplateIndex(/* invalid */),
+ /* parameters */ ParameterIndex(325),
+ /* return_matcher_indices */ MatcherIndicesIndex(151),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [130] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 0,
+ /* templates */ TemplateIndex(/* invalid */),
+ /* parameters */ ParameterIndex(326),
+ /* return_matcher_indices */ MatcherIndicesIndex(165),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [131] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 0,
+ /* templates */ TemplateIndex(/* invalid */),
+ /* parameters */ ParameterIndex(327),
+ /* return_matcher_indices */ MatcherIndicesIndex(137),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [132] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 0,
+ /* templates */ TemplateIndex(/* invalid */),
+ /* parameters */ ParameterIndex(328),
+ /* return_matcher_indices */ MatcherIndicesIndex(177),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [133] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 0,
+ /* templates */ TemplateIndex(/* invalid */),
+ /* parameters */ ParameterIndex(329),
+ /* return_matcher_indices */ MatcherIndicesIndex(171),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [134] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 0,
+ /* templates */ TemplateIndex(/* invalid */),
+ /* parameters */ ParameterIndex(330),
+ /* return_matcher_indices */ MatcherIndicesIndex(173),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [135] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 0,
+ /* templates */ TemplateIndex(/* invalid */),
+ /* parameters */ ParameterIndex(331),
+ /* return_matcher_indices */ MatcherIndicesIndex(175),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [136] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 0,
+ /* templates */ TemplateIndex(/* invalid */),
+ /* parameters */ ParameterIndex(332),
+ /* return_matcher_indices */ MatcherIndicesIndex(179),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [137] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 0,
+ /* templates */ TemplateIndex(/* invalid */),
+ /* parameters */ ParameterIndex(58),
+ /* return_matcher_indices */ MatcherIndicesIndex(30),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [138] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 5,
/* num_explicit_templates */ 0,
@@ -4711,7 +4675,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [141] */
+ /* [139] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 6,
/* num_explicit_templates */ 0,
@@ -4722,7 +4686,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [142] */
+ /* [140] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 4,
/* num_explicit_templates */ 0,
@@ -4733,7 +4697,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [143] */
+ /* [141] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 5,
/* num_explicit_templates */ 0,
@@ -4744,7 +4708,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [144] */
+ /* [142] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 4,
/* num_explicit_templates */ 0,
@@ -4755,7 +4719,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [145] */
+ /* [143] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 5,
/* num_explicit_templates */ 0,
@@ -4766,7 +4730,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [146] */
+ /* [144] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -4777,7 +4741,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [147] */
+ /* [145] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 4,
/* num_explicit_templates */ 0,
@@ -4788,7 +4752,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [148] */
+ /* [146] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 4,
/* num_explicit_templates */ 0,
@@ -4799,7 +4763,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [149] */
+ /* [147] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 5,
/* num_explicit_templates */ 0,
@@ -4810,7 +4774,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [150] */
+ /* [148] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 5,
/* num_explicit_templates */ 0,
@@ -4821,7 +4785,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [151] */
+ /* [149] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 6,
/* num_explicit_templates */ 0,
@@ -4832,7 +4796,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [152] */
+ /* [150] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 4,
/* num_explicit_templates */ 0,
@@ -4843,7 +4807,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [153] */
+ /* [151] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 5,
/* num_explicit_templates */ 0,
@@ -4854,7 +4818,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [154] */
+ /* [152] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -4865,7 +4829,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [155] */
+ /* [153] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -4876,7 +4840,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [156] */
+ /* [154] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -4887,7 +4851,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [157] */
+ /* [155] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -4898,7 +4862,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [158] */
+ /* [156] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -4909,7 +4873,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [159] */
+ /* [157] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -4920,7 +4884,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [160] */
+ /* [158] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -4931,7 +4895,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [161] */
+ /* [159] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -4942,7 +4906,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [162] */
+ /* [160] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -4953,7 +4917,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [163] */
+ /* [161] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -4964,7 +4928,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [164] */
+ /* [162] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -4975,7 +4939,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [165] */
+ /* [163] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -4986,7 +4950,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [166] */
+ /* [164] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -4997,7 +4961,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [167] */
+ /* [165] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -5008,7 +4972,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [168] */
+ /* [166] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -5019,7 +4983,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [169] */
+ /* [167] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -5030,7 +4994,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [170] */
+ /* [168] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -5041,7 +5005,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [171] */
+ /* [169] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -5052,7 +5016,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [172] */
+ /* [170] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -5063,7 +5027,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [173] */
+ /* [171] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -5074,7 +5038,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [174] */
+ /* [172] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -5085,7 +5049,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [175] */
+ /* [173] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -5096,7 +5060,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [176] */
+ /* [174] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -5107,7 +5071,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [177] */
+ /* [175] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 5,
/* num_explicit_templates */ 0,
@@ -5118,7 +5082,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [178] */
+ /* [176] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -5129,7 +5093,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [179] */
+ /* [177] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -5140,7 +5104,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [180] */
+ /* [178] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -5151,7 +5115,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [181] */
+ /* [179] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -5162,7 +5126,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [182] */
+ /* [180] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -5173,7 +5137,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [183] */
+ /* [181] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -5184,7 +5148,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [184] */
+ /* [182] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -5195,7 +5159,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [185] */
+ /* [183] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 5,
/* num_explicit_templates */ 0,
@@ -5206,7 +5170,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [186] */
+ /* [184] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 5,
/* num_explicit_templates */ 0,
@@ -5217,7 +5181,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [187] */
+ /* [185] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -5228,7 +5192,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [188] */
+ /* [186] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 4,
/* num_explicit_templates */ 0,
@@ -5239,7 +5203,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [189] */
+ /* [187] */
/* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -5259,67 +5223,67 @@
/* [0] */
/* fn atomic_compare_exchange_weak_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, ptr<function, T, read_write>, T, u32, u32) -> bool */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(177),
+ /* overloads */ OverloadIndex(175),
},
{
/* [1] */
/* fn atomic_exchange_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(178),
+ /* overloads */ OverloadIndex(176),
},
{
/* [2] */
/* fn atomic_fetch_add_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(178),
+ /* overloads */ OverloadIndex(176),
},
{
/* [3] */
/* fn atomic_fetch_and_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(178),
+ /* overloads */ OverloadIndex(176),
},
{
/* [4] */
/* fn atomic_fetch_max_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(178),
+ /* overloads */ OverloadIndex(176),
},
{
/* [5] */
/* fn atomic_fetch_min_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(178),
+ /* overloads */ OverloadIndex(176),
},
{
/* [6] */
/* fn atomic_fetch_or_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(178),
+ /* overloads */ OverloadIndex(176),
},
{
/* [7] */
/* fn atomic_fetch_sub_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(178),
+ /* overloads */ OverloadIndex(176),
},
{
/* [8] */
/* fn atomic_fetch_xor_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(178),
+ /* overloads */ OverloadIndex(176),
},
{
/* [9] */
/* fn atomic_load_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(179),
+ /* overloads */ OverloadIndex(177),
},
{
/* [10] */
/* fn atomic_store_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(180),
+ /* overloads */ OverloadIndex(178),
},
{
/* [11] */
@@ -5328,7 +5292,7 @@
/* fn fence[F : texel_format, A : access](texture: texture_storage_2d_array<F, A>) */
/* fn fence[F : texel_format, A : access](texture: texture_storage_3d<F, A>) */
/* num overloads */ 4,
- /* overloads */ OverloadIndex(159),
+ /* overloads */ OverloadIndex(157),
},
{
/* [12] */
@@ -5341,7 +5305,7 @@
/* fn gather(texture: texture_depth_cube, sampler: sampler, coords: vec3<f32>) -> vec4<f32> */
/* fn gather[A : iu32](texture: texture_depth_cube_array, sampler: sampler, coords: vec3<f32>, array_index: A) -> vec4<f32> */
/* num overloads */ 8,
- /* overloads */ OverloadIndex(140),
+ /* overloads */ OverloadIndex(138),
},
{
/* [13] */
@@ -5352,7 +5316,7 @@
/* fn gather_compare(texture: texture_depth_cube, sampler: sampler_comparison, coords: vec3<f32>, depth_ref: f32) -> vec4<f32> */
/* fn gather_compare[A : iu32](texture: texture_depth_cube_array, sampler: sampler_comparison, coords: vec3<f32>, array_index: A, depth_ref: f32) -> vec4<f32> */
/* num overloads */ 6,
- /* overloads */ OverloadIndex(148),
+ /* overloads */ OverloadIndex(146),
},
{
/* [14] */
@@ -5399,7 +5363,7 @@
/* fn get_depth[T : fiu32](texture: texture_3d<T>, u32) -> u32 */
/* fn get_depth[F : texel_format, A : access](texture: texture_storage_3d<F, A>, u32) -> u32 */
/* num overloads */ 2,
- /* overloads */ OverloadIndex(167),
+ /* overloads */ OverloadIndex(165),
},
{
/* [17] */
@@ -5409,7 +5373,7 @@
/* fn get_array_size(texture: texture_depth_cube_array) -> u32 */
/* fn get_array_size[F : texel_format, A : access](texture: texture_storage_2d_array<F, A>) -> u32 */
/* num overloads */ 5,
- /* overloads */ OverloadIndex(154),
+ /* overloads */ OverloadIndex(152),
},
{
/* [18] */
@@ -5424,14 +5388,14 @@
/* fn get_num_mip_levels(texture: texture_depth_cube) -> u32 */
/* fn get_num_mip_levels(texture: texture_depth_cube_array) -> u32 */
/* num overloads */ 10,
- /* overloads */ OverloadIndex(130),
+ /* overloads */ OverloadIndex(119),
},
{
/* [19] */
/* fn get_num_samples[T : fiu32](texture: texture_multisampled_2d<T>) -> u32 */
/* fn get_num_samples(texture: texture_depth_multisampled_2d) -> u32 */
/* num overloads */ 2,
- /* overloads */ OverloadIndex(169),
+ /* overloads */ OverloadIndex(167),
},
{
/* [20] */
@@ -5546,13 +5510,13 @@
/* [24] */
/* fn distance[N : num, T : f32_f16](vec<N, T>, vec<N, T>) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(181),
+ /* overloads */ OverloadIndex(179),
},
{
/* [25] */
/* fn dot[N : num, T : f32_f16](vec<N, T>, vec<N, T>) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(181),
+ /* overloads */ OverloadIndex(179),
},
{
/* [26] */
@@ -5561,93 +5525,91 @@
/* fn fmod[N : num, T : f32_f16](T, vec<N, T>) -> vec<N, T> */
/* fn fmod[N : num, T : f32_f16](vec<N, T>, T) -> vec<N, T> */
/* num overloads */ 4,
- /* overloads */ OverloadIndex(163),
+ /* overloads */ OverloadIndex(161),
},
{
/* [27] */
/* fn frexp[T : f32_f16](T, i32) -> T */
/* fn frexp[N : num, T : f32_f16](vec<N, T>, vec<N, i32>) -> vec<N, T> */
/* num overloads */ 2,
- /* overloads */ OverloadIndex(171),
+ /* overloads */ OverloadIndex(169),
},
{
/* [28] */
/* fn length[N : num, T : f32_f16](vec<N, T>) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(182),
+ /* overloads */ OverloadIndex(180),
},
{
/* [29] */
/* fn modf[T : f32_f16](T, T) -> T */
/* fn modf[N : num, T : f32_f16](vec<N, T>, vec<N, T>) -> vec<N, T> */
/* num overloads */ 2,
- /* overloads */ OverloadIndex(163),
+ /* overloads */ OverloadIndex(161),
},
{
/* [30] */
/* fn sign[T : f32_f16](T) -> T */
/* fn sign[N : num, T : f32_f16](vec<N, T>) -> vec<N, T> */
/* num overloads */ 2,
- /* overloads */ OverloadIndex(173),
+ /* overloads */ OverloadIndex(171),
},
{
/* [31] */
/* fn threadgroup_barrier(u32) */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(183),
+ /* overloads */ OverloadIndex(181),
},
{
/* [32] */
/* fn simd_ballot(bool) -> vec2<u32> */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(184),
+ /* overloads */ OverloadIndex(182),
},
{
/* [33] */
/* fn quad_shuffle_xor[T : fiu32_f16](T, u32) -> T */
/* fn quad_shuffle_xor[N : num, T : fiu32_f16](vec<N, T>, u32) -> vec<N, T> */
/* num overloads */ 2,
- /* overloads */ OverloadIndex(175),
+ /* overloads */ OverloadIndex(173),
},
{
/* [34] */
- /* fn convert(__packed_vec3<bool>) -> vec3<bool> */
/* fn convert(__packed_vec3<u32>) -> vec3<u32> */
/* fn convert(__packed_vec3<i32>) -> vec3<i32> */
/* fn convert(__packed_vec3<f32>) -> vec3<f32> */
/* fn convert(__packed_vec3<f16>) -> vec3<f16> */
- /* fn convert(vec3<bool>) -> __packed_vec3<bool> */
/* fn convert(vec3<u32>) -> __packed_vec3<u32> */
/* fn convert(vec3<i32>) -> __packed_vec3<i32> */
/* fn convert(vec3<f32>) -> __packed_vec3<f32> */
/* fn convert(vec3<f16>) -> __packed_vec3<f16> */
/* fn convert(u32) -> u64 */
- /* num overloads */ 11,
- /* overloads */ OverloadIndex(119),
+ /* num overloads */ 9,
+ /* overloads */ OverloadIndex(129),
},
{
/* [35] */
/* fn simdgroup_load[K : subgroup_matrix_kind, S : f32_f16, C : num, R : num](subgroup_matrix<K, S, C, R>, ptr<workgroup_or_storage, S, readable>, u64, vec2<u64>, bool) */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(185),
+ /* overloads */ OverloadIndex(183),
},
{
/* [36] */
/* fn simdgroup_store[K : subgroup_matrix_kind, S : f32_f16, C : num, R : num](subgroup_matrix<K, S, C, R>, ptr<workgroup_or_storage, S, writable>, u64, vec2<u64>, bool) */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(186),
+ /* overloads */ OverloadIndex(184),
},
{
/* [37] */
/* fn simdgroup_multiply[S : f32_f16, C : num, R : num, K : num](subgroup_matrix<subgroup_matrix_kind_result, S, C, R>, subgroup_matrix<subgroup_matrix_kind_left, S, K, R>, subgroup_matrix<subgroup_matrix_kind_right, S, C, K>) */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(187),
+ /* overloads */ OverloadIndex(185),
},
{
/* [38] */
/* fn simdgroup_multiply_accumulate[S : f32_f16, C : num, R : num, K : num](subgroup_matrix<subgroup_matrix_kind_result, S, C, R>, subgroup_matrix<subgroup_matrix_kind_left, S, K, R>, subgroup_matrix<subgroup_matrix_kind_right, S, C, K>, subgroup_matrix<subgroup_matrix_kind_result, S, C, R>) */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(188),
+ /* overloads */ OverloadIndex(186),
},
};
@@ -5656,13 +5618,13 @@
/* [0] */
/* op +[T : iu8](T, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(189),
+ /* overloads */ OverloadIndex(187),
},
{
/* [1] */
/* op *[T : iu8](T, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(189),
+ /* overloads */ OverloadIndex(187),
},
};
constexpr uint8_t kBinaryOperatorPlus = 0;
diff --git a/src/tint/lang/msl/msl.def b/src/tint/lang/msl/msl.def
index 8142cac..5f5f8d8 100644
--- a/src/tint/lang/msl/msl.def
+++ b/src/tint/lang/msl/msl.def
@@ -354,12 +354,10 @@
@stage("fragment", "compute") implicit(T: fiu32_f16) fn quad_shuffle_xor(T, u32) -> T
@stage("fragment", "compute") implicit(N: num, T: fiu32_f16) fn quad_shuffle_xor(vec<N, T>, u32) -> vec<N, T>
-fn convert(__packed_vec3<bool>) -> vec3<bool>
fn convert(__packed_vec3<u32>) -> vec3<u32>
fn convert(__packed_vec3<i32>) -> vec3<i32>
fn convert(__packed_vec3<f32>) -> vec3<f32>
fn convert(__packed_vec3<f16>) -> vec3<f16>
-fn convert(vec3<bool>) -> __packed_vec3<bool>
fn convert(vec3<u32>) -> __packed_vec3<u32>
fn convert(vec3<i32>) -> __packed_vec3<i32>
fn convert(vec3<f32>) -> __packed_vec3<f32>
diff --git a/src/tint/lang/msl/writer/printer/printer.cc b/src/tint/lang/msl/writer/printer/printer.cc
index 401b00c..c359d5b 100644
--- a/src/tint/lang/msl/writer/printer/printer.cc
+++ b/src/tint/lang/msl/writer/printer/printer.cc
@@ -1416,6 +1416,9 @@
/// @param vec the vector to emit
void EmitVectorType(StringStream& out, const core::type::Vector* vec) {
if (vec->Packed()) {
+ // packed_bool* vectors are accepted by the MSL compiler but are reserved by the MSL
+ // spec, and cause issues with some drivers (see crbug.com/424772881).
+ TINT_ASSERT(!vec->IsBoolVector());
out << "packed_";
}
EmitType(out, vec->Type());
diff --git a/src/tint/lang/msl/writer/raise/packed_vec3.cc b/src/tint/lang/msl/writer/raise/packed_vec3.cc
index b58799c..8433492 100644
--- a/src/tint/lang/msl/writer/raise/packed_vec3.cc
+++ b/src/tint/lang/msl/writer/raise/packed_vec3.cc
@@ -155,7 +155,16 @@
},
[&](const core::type::Vector* vec) {
if (vec->Width() == 3) {
- return ty.packed_vec(vec->Type(), 3);
+ auto* el_ty = vec->Type();
+
+ // The packed_bool* types are reserved in MSL and cause issues in some
+ // drivers (see crbug.com/424772881), so use packed_uint3 instead. We then
+ // convert between `bool` and `u32` when we load or store the values.
+ if (el_ty->Is<core::type::Bool>()) {
+ el_ty = ty.u32();
+ }
+
+ return ty.packed_vec(el_ty, 3);
}
return vec;
},
@@ -302,15 +311,31 @@
});
load->Destroy();
},
- [&](core::ir::LoadVectorElement*) {
- // Nothing to do - packed vectors support component access.
+ [&](core::ir::LoadVectorElement* lve) {
+ // Packed vectors support component access so there is usually nothing to do.
+ // For vectors that were originally booleans we need to convert the u32 that we load
+ // to a bool.
+ if (unpacked_type->IsBoolVector()) {
+ auto* u32_load = b.InstructionResult<u32>();
+ auto* converted_to_bool = b.ConvertWithResult(lve->DetachResult(), u32_load);
+ converted_to_bool->InsertAfter(lve);
+ lve->SetResult(u32_load);
+ }
},
[&](core::ir::Store* store) {
b.InsertBefore(store, [&] { StoreUnpackedToPacked(store->To(), store->From()); });
store->Destroy();
},
- [&](core::ir::StoreVectorElement*) {
- // Nothing to do - packed vectors support component access.
+ [&](core::ir::StoreVectorElement* sve) {
+ // Packed vectors support component access so there is usually nothing to do.
+ // For vectors that were originally booleans we need to convert the bool to a u32
+ // before we store it.
+ if (unpacked_type->IsBoolVector()) {
+ auto* converted_to_u32 = b.Convert<u32>(sve->Value());
+ converted_to_u32->InsertBefore(sve);
+ sve->SetOperand(core::ir::StoreVectorElement::kValueOperandOffset,
+ converted_to_u32->Result());
+ }
},
[&](core::ir::UserCall*) {
// Nothing to do - pass the packed type to the function, which will be rewritten.
@@ -398,12 +423,22 @@
[&](const core::type::Struct* str) {
return b.Call(LoadPackedStructHelper(str, packed_ptr), from)->Result();
},
- [&](const core::type::Vector*) {
+ [&](const core::type::Vector* vec) {
// Load the packed vector and convert it to the unpacked equivalent.
- return b
- .Call<msl::ir::BuiltinCall>(unpacked_type, msl::BuiltinFn::kConvert,
- b.Load(from))
- ->Result();
+ auto* load = b.Load(from)->Result(0);
+ if (vec->Type()->Is<core::type::Bool>()) {
+ // The vector was originally a vec3<bool>, which will have been rewritten as a
+ // vec3<u32>. We need to unpack the packed_vec3<u32> to a vec3<u32> and then
+ // convert it to a vec3<bool>.
+ auto* unpacked =
+ b.Call<msl::ir::BuiltinCall>(ty.vec3<u32>(), msl::BuiltinFn::kConvert, load)
+ ->Result();
+ return b.Convert<vec3<bool>>(unpacked)->Result();
+ } else {
+ return b
+ .Call<msl::ir::BuiltinCall>(unpacked_type, msl::BuiltinFn::kConvert, load)
+ ->Result();
+ }
},
TINT_ICE_ON_NO_MATCH);
}
@@ -539,8 +574,13 @@
[&](const core::type::Struct* str) {
b.Call(StorePackedStructHelper(str, packed_ptr), to, value)->Result();
},
- [&](const core::type::Vector*) { //
+ [&](const core::type::Vector* vec) { //
// Convert the vector to the packed equivalent and store it.
+ // For vectors that were originally booleans we need to convert the value to a
+ // vec3<u32> before storing it.
+ if (vec->Type()->Is<core::type::Bool>()) {
+ value = b.Convert<vec3<u32>>(value)->Result();
+ }
b.Store(to,
b.Call<msl::ir::BuiltinCall>(packed_type, msl::BuiltinFn::kConvert, value)
->Result());
diff --git a/src/tint/lang/msl/writer/raise/packed_vec3_test.cc b/src/tint/lang/msl/writer/raise/packed_vec3_test.cc
index 86fbcc7..18ea20d 100644
--- a/src/tint/lang/msl/writer/raise/packed_vec3_test.cc
+++ b/src/tint/lang/msl/writer/raise/packed_vec3_test.cc
@@ -3762,6 +3762,7 @@
}
// Workgroup is the only address space that requires packed types that supports bool types.
+// These are rewritten as packed_vec3<u32> types since MSL does not support packed bool vectors.
TEST_F(MslWriter_PackedVec3Test, WorkgroupVar_Vec3_Bool) {
auto* var = b.Var<workgroup, vec3<bool>>("v");
mod.root_block->Append(var);
@@ -3789,16 +3790,67 @@
auto* expect = R"(
$B1: { # root
- %v:ptr<workgroup, __packed_vec3<bool>, read_write> = var undef
+ %v:ptr<workgroup, __packed_vec3<u32>, read_write> = var undef
}
%foo = func():vec3<bool> {
$B2: {
- %3:__packed_vec3<bool> = msl.convert vec3<bool>(false)
- store %v, %3
- %4:__packed_vec3<bool> = load %v
- %5:vec3<bool> = msl.convert %4
- ret %5
+ %3:vec3<u32> = convert vec3<bool>(false)
+ %4:__packed_vec3<u32> = msl.convert %3
+ store %v, %4
+ %5:__packed_vec3<u32> = load %v
+ %6:vec3<u32> = msl.convert %5
+ %7:vec3<bool> = convert %6
+ ret %7
+ }
+}
+)";
+
+ Run(PackedVec3);
+
+ EXPECT_EQ(expect, str());
+}
+
+// Workgroup is the only address space that requires packed types that supports bool types.
+// These are rewritten as packed_vec3<u32> types since MSL does not support packed bool vectors.
+TEST_F(MslWriter_PackedVec3Test, WorkgroupVar_Vec3_Bool_VectorElementLoadAndStore) {
+ auto* var = b.Var<workgroup, vec3<bool>>("v");
+ mod.root_block->Append(var);
+
+ auto* func = b.Function("foo", ty.void_());
+ b.Append(func->Block(), [&] { //
+ auto* el = b.LoadVectorElement(var, 0_u);
+ b.StoreVectorElement(var, 1_u, el);
+ b.Return(func);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %v:ptr<workgroup, vec3<bool>, read_write> = var undef
+}
+
+%foo = func():void {
+ $B2: {
+ %3:bool = load_vector_element %v, 0u
+ store_vector_element %v, 1u, %3
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+$B1: { # root
+ %v:ptr<workgroup, __packed_vec3<u32>, read_write> = var undef
+}
+
+%foo = func():void {
+ $B2: {
+ %3:u32 = load_vector_element %v, 0u
+ %4:bool = convert %3
+ %5:u32 = convert %4
+ store_vector_element %v, 1u, %5
+ ret
}
}
)";
diff --git a/test/tint/bug/tint/424772881.wgsl b/test/tint/bug/tint/424772881.wgsl
new file mode 100644
index 0000000..98cef1a
--- /dev/null
+++ b/test/tint/bug/tint/424772881.wgsl
@@ -0,0 +1,16 @@
+@group(0) @binding(0) var<storage, read_write> result: vec3<u32>;
+
+var<workgroup> wgvar: vec3<bool>;
+
+@compute @workgroup_size(1)
+fn main() {
+ let v = wgvar;
+ wgvar = v;
+
+ let e = wgvar[0];
+ wgvar[1] = e;
+
+ workgroupBarrier();
+
+ result = vec3<u32>(wgvar);
+}
diff --git a/test/tint/bug/tint/424772881.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/424772881.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..bc67700
--- /dev/null
+++ b/test/tint/bug/tint/424772881.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+groupshared bool3 wgvar;
+
+void tint_zero_workgroup_memory(uint local_idx) {
+ if ((local_idx < 1u)) {
+ wgvar = (false).xxx;
+ }
+ GroupMemoryBarrierWithGroupSync();
+}
+
+RWByteAddressBuffer result : register(u0);
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void main_inner(uint local_invocation_index) {
+ tint_zero_workgroup_memory(local_invocation_index);
+ bool3 v = wgvar;
+ wgvar = v;
+ bool e = wgvar[0];
+ wgvar[1] = e;
+ GroupMemoryBarrierWithGroupSync();
+ result.Store3(0u, asuint(uint3(wgvar)));
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+ main_inner(tint_symbol.local_invocation_index);
+ return;
+}
diff --git a/test/tint/bug/tint/424772881.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/424772881.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..bc67700
--- /dev/null
+++ b/test/tint/bug/tint/424772881.wgsl.expected.fxc.hlsl
@@ -0,0 +1,30 @@
+groupshared bool3 wgvar;
+
+void tint_zero_workgroup_memory(uint local_idx) {
+ if ((local_idx < 1u)) {
+ wgvar = (false).xxx;
+ }
+ GroupMemoryBarrierWithGroupSync();
+}
+
+RWByteAddressBuffer result : register(u0);
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void main_inner(uint local_invocation_index) {
+ tint_zero_workgroup_memory(local_invocation_index);
+ bool3 v = wgvar;
+ wgvar = v;
+ bool e = wgvar[0];
+ wgvar[1] = e;
+ GroupMemoryBarrierWithGroupSync();
+ result.Store3(0u, asuint(uint3(wgvar)));
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+ main_inner(tint_symbol.local_invocation_index);
+ return;
+}
diff --git a/test/tint/bug/tint/424772881.wgsl.expected.glsl b/test/tint/bug/tint/424772881.wgsl.expected.glsl
new file mode 100644
index 0000000..9dd1d20
--- /dev/null
+++ b/test/tint/bug/tint/424772881.wgsl.expected.glsl
@@ -0,0 +1,23 @@
+#version 310 es
+
+layout(binding = 0, std430)
+buffer result_block_1_ssbo {
+ uvec3 inner;
+} v_1;
+shared bvec3 wgvar;
+void main_inner(uint tint_local_index) {
+ if ((tint_local_index < 1u)) {
+ wgvar = bvec3(false);
+ }
+ barrier();
+ bvec3 v = wgvar;
+ wgvar = v;
+ bool e = wgvar.x;
+ wgvar.y = e;
+ barrier();
+ v_1.inner = uvec3(wgvar);
+}
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ main_inner(gl_LocalInvocationIndex);
+}
diff --git a/test/tint/bug/tint/424772881.wgsl.expected.ir.dxc.hlsl b/test/tint/bug/tint/424772881.wgsl.expected.ir.dxc.hlsl
new file mode 100644
index 0000000..140dcb6
--- /dev/null
+++ b/test/tint/bug/tint/424772881.wgsl.expected.ir.dxc.hlsl
@@ -0,0 +1,25 @@
+struct main_inputs {
+ uint tint_local_index : SV_GroupIndex;
+};
+
+
+RWByteAddressBuffer result : register(u0);
+groupshared bool3 wgvar;
+void main_inner(uint tint_local_index) {
+ if ((tint_local_index < 1u)) {
+ wgvar = (false).xxx;
+ }
+ GroupMemoryBarrierWithGroupSync();
+ bool3 v = wgvar;
+ wgvar = v;
+ bool e = wgvar.x;
+ wgvar.y = e;
+ GroupMemoryBarrierWithGroupSync();
+ result.Store3(0u, uint3(wgvar));
+}
+
+[numthreads(1, 1, 1)]
+void main(main_inputs inputs) {
+ main_inner(inputs.tint_local_index);
+}
+
diff --git a/test/tint/bug/tint/424772881.wgsl.expected.ir.fxc.hlsl b/test/tint/bug/tint/424772881.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..140dcb6
--- /dev/null
+++ b/test/tint/bug/tint/424772881.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,25 @@
+struct main_inputs {
+ uint tint_local_index : SV_GroupIndex;
+};
+
+
+RWByteAddressBuffer result : register(u0);
+groupshared bool3 wgvar;
+void main_inner(uint tint_local_index) {
+ if ((tint_local_index < 1u)) {
+ wgvar = (false).xxx;
+ }
+ GroupMemoryBarrierWithGroupSync();
+ bool3 v = wgvar;
+ wgvar = v;
+ bool e = wgvar.x;
+ wgvar.y = e;
+ GroupMemoryBarrierWithGroupSync();
+ result.Store3(0u, uint3(wgvar));
+}
+
+[numthreads(1, 1, 1)]
+void main(main_inputs inputs) {
+ main_inner(inputs.tint_local_index);
+}
+
diff --git a/test/tint/bug/tint/424772881.wgsl.expected.msl b/test/tint/bug/tint/424772881.wgsl.expected.msl
new file mode 100644
index 0000000..8552e06
--- /dev/null
+++ b/test/tint/bug/tint/424772881.wgsl.expected.msl
@@ -0,0 +1,29 @@
+#include <metal_stdlib>
+using namespace metal;
+
+struct tint_module_vars_struct {
+ device packed_uint3* result;
+ threadgroup packed_uint3* wgvar;
+};
+
+struct tint_symbol_1 {
+ packed_uint3 tint_symbol;
+};
+
+void main_inner(uint tint_local_index, tint_module_vars_struct tint_module_vars) {
+ if ((tint_local_index < 1u)) {
+ (*tint_module_vars.wgvar) = packed_uint3(uint3(bool3(false)));
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ bool3 const v = bool3(uint3((*tint_module_vars.wgvar)));
+ (*tint_module_vars.wgvar) = packed_uint3(uint3(v));
+ bool const e = bool((*tint_module_vars.wgvar).x);
+ (*tint_module_vars.wgvar).y = uint(e);
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ (*tint_module_vars.result) = packed_uint3(uint3(bool3(uint3((*tint_module_vars.wgvar)))));
+}
+
+kernel void v_1(uint tint_local_index [[thread_index_in_threadgroup]], device packed_uint3* result [[buffer(0)]], threadgroup tint_symbol_1* v_2 [[threadgroup(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.result=result, .wgvar=(&(*v_2).tint_symbol)};
+ main_inner(tint_local_index, tint_module_vars);
+}
diff --git a/test/tint/bug/tint/424772881.wgsl.expected.spvasm b/test/tint/bug/tint/424772881.wgsl.expected.spvasm
new file mode 100644
index 0000000..97ef47e
--- /dev/null
+++ b/test/tint/bug/tint/424772881.wgsl.expected.spvasm
@@ -0,0 +1,77 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 43
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main" %main_local_invocation_index_Input
+ OpExecutionMode %main LocalSize 1 1 1
+ OpMemberName %result_block 0 "inner"
+ OpName %result_block "result_block"
+ OpName %wgvar "wgvar"
+ OpName %main_local_invocation_index_Input "main_local_invocation_index_Input"
+ OpName %main_inner "main_inner"
+ OpName %tint_local_index "tint_local_index"
+ OpName %v "v"
+ OpName %e "e"
+ OpName %main "main"
+ OpMemberDecorate %result_block 0 Offset 0
+ OpDecorate %result_block Block
+ OpDecorate %1 DescriptorSet 0
+ OpDecorate %1 Binding 0
+ OpDecorate %1 Coherent
+ OpDecorate %main_local_invocation_index_Input BuiltIn LocalInvocationIndex
+ %uint = OpTypeInt 32 0
+ %v3uint = OpTypeVector %uint 3
+%result_block = OpTypeStruct %v3uint
+%_ptr_StorageBuffer_result_block = OpTypePointer StorageBuffer %result_block
+ %1 = OpVariable %_ptr_StorageBuffer_result_block StorageBuffer
+ %bool = OpTypeBool
+ %v3bool = OpTypeVector %bool 3
+%_ptr_Workgroup_v3bool = OpTypePointer Workgroup %v3bool
+ %wgvar = OpVariable %_ptr_Workgroup_v3bool Workgroup
+%_ptr_Input_uint = OpTypePointer Input %uint
+%main_local_invocation_index_Input = OpVariable %_ptr_Input_uint Input
+ %void = OpTypeVoid
+ %15 = OpTypeFunction %void %uint
+ %uint_1 = OpConstant %uint 1
+ %uint_2 = OpConstant %uint 2
+ %uint_264 = OpConstant %uint 264
+%_ptr_Workgroup_bool = OpTypePointer Workgroup %bool
+ %uint_0 = OpConstant %uint 0
+ %33 = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1
+ %34 = OpConstantNull %v3uint
+%_ptr_StorageBuffer_v3uint = OpTypePointer StorageBuffer %v3uint
+ %37 = OpConstantNull %v3bool
+ %39 = OpTypeFunction %void
+ %main_inner = OpFunction %void None %15
+%tint_local_index = OpFunctionParameter %uint
+ %16 = OpLabel
+ %17 = OpULessThan %bool %tint_local_index %uint_1
+ OpSelectionMerge %19 None
+ OpBranchConditional %17 %20 %19
+ %20 = OpLabel
+ OpStore %wgvar %37 None
+ OpBranch %19
+ %19 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %v = OpLoad %v3bool %wgvar None
+ OpStore %wgvar %v None
+ %25 = OpAccessChain %_ptr_Workgroup_bool %wgvar %uint_0
+ %e = OpLoad %bool %25 None
+ %29 = OpAccessChain %_ptr_Workgroup_bool %wgvar %uint_1
+ OpStore %29 %e None
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %31 = OpLoad %v3bool %wgvar None
+ %32 = OpSelect %v3uint %31 %33 %34
+ %35 = OpAccessChain %_ptr_StorageBuffer_v3uint %1 %uint_0
+ OpStore %35 %32 None
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %39
+ %40 = OpLabel
+ %41 = OpLoad %uint %main_local_invocation_index_Input None
+ %42 = OpFunctionCall %void %main_inner %41
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/bug/tint/424772881.wgsl.expected.wgsl b/test/tint/bug/tint/424772881.wgsl.expected.wgsl
new file mode 100644
index 0000000..fb6936e
--- /dev/null
+++ b/test/tint/bug/tint/424772881.wgsl.expected.wgsl
@@ -0,0 +1,13 @@
+@group(0) @binding(0) var<storage, read_write> result : vec3<u32>;
+
+var<workgroup> wgvar : vec3<bool>;
+
+@compute @workgroup_size(1)
+fn main() {
+ let v = wgvar;
+ wgvar = v;
+ let e = wgvar[0];
+ wgvar[1] = e;
+ workgroupBarrier();
+ result = vec3<u32>(wgvar);
+}
