[msl] Add polyfill for length builtin
If the argument is scalar, we need to use `abs()`.
Add an MSL builtin function for length that only supports vectors.
Bug: 42251016
Change-Id: I82930ebacdf5620eabb4c77ff8da124f0b930016
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/193920
Reviewed-by: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/msl/builtin_fn.cc b/src/tint/lang/msl/builtin_fn.cc
index 8b3ca14..e345a4d 100644
--- a/src/tint/lang/msl/builtin_fn.cc
+++ b/src/tint/lang/msl/builtin_fn.cc
@@ -90,6 +90,8 @@
return "sample_compare";
case BuiltinFn::kWrite:
return "write";
+ case BuiltinFn::kLength:
+ return "length";
case BuiltinFn::kThreadgroupBarrier:
return "threadgroup_barrier";
}
diff --git a/src/tint/lang/msl/builtin_fn.h b/src/tint/lang/msl/builtin_fn.h
index 6a6cef1..04b22b1 100644
--- a/src/tint/lang/msl/builtin_fn.h
+++ b/src/tint/lang/msl/builtin_fn.h
@@ -71,6 +71,7 @@
kSample,
kSampleCompare,
kWrite,
+ kLength,
kThreadgroupBarrier,
kNone,
};
diff --git a/src/tint/lang/msl/intrinsic/data.cc b/src/tint/lang/msl/intrinsic/data.cc
index 8928a86..eb585fd 100644
--- a/src/tint/lang/msl/intrinsic/data.cc
+++ b/src/tint/lang/msl/intrinsic/data.cc
@@ -131,6 +131,20 @@
};
+/// TypeMatcher for 'type f16'
+constexpr TypeMatcher kF16Matcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+ if (!MatchF16(state, ty)) {
+ return nullptr;
+ }
+ return BuildF16(state, ty);
+ },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
+ out << style::Type("f16");
+ }
+};
+
+
/// TypeMatcher for 'type vec2'
constexpr TypeMatcher kVec2Matcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
@@ -191,6 +205,32 @@
};
+/// TypeMatcher for 'type vec'
+constexpr TypeMatcher kVecMatcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+ Number N = Number::invalid;
+ const Type* T = nullptr;
+ if (!MatchVec(state, ty, N, T)) {
+ return nullptr;
+ }
+ N = state.Num(N);
+ if (!N.IsValid()) {
+ return nullptr;
+ }
+ T = state.Type(T);
+ if (T == nullptr) {
+ return nullptr;
+ }
+ return BuildVec(state, ty, N, T);
+ },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText N;
+ state->PrintNum(N);StyledText T;
+ state->PrintType(T);
+ out << style::Type("vec", N, "<", T, ">");
+ }
+};
+
+
/// TypeMatcher for 'type atomic'
constexpr TypeMatcher kAtomicMatcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
@@ -692,6 +732,23 @@
kF32Matcher.print(nullptr, out); out << style::Plain(", "); kI32Matcher.print(nullptr, out); out << style::Plain(" or "); kU32Matcher.print(nullptr, out);}
};
+/// TypeMatcher for 'match f32_f16'
+constexpr TypeMatcher kF32F16Matcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+ if (MatchF32(state, ty)) {
+ return BuildF32(state, ty);
+ }
+ if (MatchF16(state, ty)) {
+ return BuildF16(state, ty);
+ }
+ return nullptr;
+ },
+/* print */ [](MatchState*, StyledText& out) {
+ // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
+ // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
+ kF32Matcher.print(nullptr, out); out << style::Plain(" or "); kF16Matcher.print(nullptr, out);}
+};
+
/// EnumMatcher for 'match read_write'
constexpr NumberMatcher kReadWriteMatcher {
/* match */ [](MatchState&, Number number) -> Number {
@@ -837,36 +894,39 @@
/* [4] */ kI32Matcher,
/* [5] */ kU32Matcher,
/* [6] */ kF32Matcher,
- /* [7] */ kVec2Matcher,
- /* [8] */ kVec3Matcher,
- /* [9] */ kVec4Matcher,
- /* [10] */ kAtomicMatcher,
- /* [11] */ kPtrMatcher,
- /* [12] */ kSamplerMatcher,
- /* [13] */ kSamplerComparisonMatcher,
- /* [14] */ kTexture1DMatcher,
- /* [15] */ kTexture2DMatcher,
- /* [16] */ kTexture2DArrayMatcher,
- /* [17] */ kTexture3DMatcher,
- /* [18] */ kTextureCubeMatcher,
- /* [19] */ kTextureCubeArrayMatcher,
- /* [20] */ kTextureDepth2DMatcher,
- /* [21] */ kTextureDepth2DArrayMatcher,
- /* [22] */ kTextureDepthCubeMatcher,
- /* [23] */ kTextureDepthCubeArrayMatcher,
- /* [24] */ kTextureDepthMultisampled2DMatcher,
- /* [25] */ kTextureMultisampled2DMatcher,
- /* [26] */ kTextureStorage1DMatcher,
- /* [27] */ kTextureStorage2DMatcher,
- /* [28] */ kTextureStorage2DArrayMatcher,
- /* [29] */ kTextureStorage3DMatcher,
- /* [30] */ kBiasMatcher,
- /* [31] */ kGradient2DMatcher,
- /* [32] */ kGradient3DMatcher,
- /* [33] */ kGradientcubeMatcher,
- /* [34] */ kLevelMatcher,
- /* [35] */ kIu32Matcher,
- /* [36] */ kFiu32Matcher,
+ /* [7] */ kF16Matcher,
+ /* [8] */ kVec2Matcher,
+ /* [9] */ kVec3Matcher,
+ /* [10] */ kVec4Matcher,
+ /* [11] */ kVecMatcher,
+ /* [12] */ kAtomicMatcher,
+ /* [13] */ kPtrMatcher,
+ /* [14] */ kSamplerMatcher,
+ /* [15] */ kSamplerComparisonMatcher,
+ /* [16] */ kTexture1DMatcher,
+ /* [17] */ kTexture2DMatcher,
+ /* [18] */ kTexture2DArrayMatcher,
+ /* [19] */ kTexture3DMatcher,
+ /* [20] */ kTextureCubeMatcher,
+ /* [21] */ kTextureCubeArrayMatcher,
+ /* [22] */ kTextureDepth2DMatcher,
+ /* [23] */ kTextureDepth2DArrayMatcher,
+ /* [24] */ kTextureDepthCubeMatcher,
+ /* [25] */ kTextureDepthCubeArrayMatcher,
+ /* [26] */ kTextureDepthMultisampled2DMatcher,
+ /* [27] */ kTextureMultisampled2DMatcher,
+ /* [28] */ kTextureStorage1DMatcher,
+ /* [29] */ kTextureStorage2DMatcher,
+ /* [30] */ kTextureStorage2DArrayMatcher,
+ /* [31] */ kTextureStorage3DMatcher,
+ /* [32] */ kBiasMatcher,
+ /* [33] */ kGradient2DMatcher,
+ /* [34] */ kGradient3DMatcher,
+ /* [35] */ kGradientcubeMatcher,
+ /* [36] */ kLevelMatcher,
+ /* [37] */ kIu32Matcher,
+ /* [38] */ kFiu32Matcher,
+ /* [39] */ kF32F16Matcher,
};
/// The template numbers, and number matchers
@@ -885,152 +945,153 @@
};
constexpr MatcherIndex kMatcherIndices[] = {
- /* [0] */ MatcherIndex(11),
+ /* [0] */ MatcherIndex(13),
/* [1] */ MatcherIndex(1),
- /* [2] */ MatcherIndex(10),
+ /* [2] */ MatcherIndex(12),
/* [3] */ MatcherIndex(0),
/* [4] */ MatcherIndex(3),
- /* [5] */ MatcherIndex(11),
+ /* [5] */ MatcherIndex(13),
/* [6] */ MatcherIndex(6),
/* [7] */ MatcherIndex(0),
/* [8] */ MatcherIndex(3),
- /* [9] */ MatcherIndex(26),
+ /* [9] */ MatcherIndex(28),
/* [10] */ MatcherIndex(0),
/* [11] */ MatcherIndex(1),
- /* [12] */ MatcherIndex(27),
+ /* [12] */ MatcherIndex(29),
/* [13] */ MatcherIndex(0),
/* [14] */ MatcherIndex(1),
- /* [15] */ MatcherIndex(28),
+ /* [15] */ MatcherIndex(30),
/* [16] */ MatcherIndex(0),
/* [17] */ MatcherIndex(1),
- /* [18] */ MatcherIndex(29),
+ /* [18] */ MatcherIndex(31),
/* [19] */ MatcherIndex(0),
/* [20] */ MatcherIndex(1),
- /* [21] */ MatcherIndex(26),
+ /* [21] */ MatcherIndex(28),
/* [22] */ MatcherIndex(8),
/* [23] */ MatcherIndex(4),
- /* [24] */ MatcherIndex(27),
+ /* [24] */ MatcherIndex(29),
/* [25] */ MatcherIndex(8),
/* [26] */ MatcherIndex(4),
- /* [27] */ MatcherIndex(28),
+ /* [27] */ MatcherIndex(30),
/* [28] */ MatcherIndex(8),
/* [29] */ MatcherIndex(4),
- /* [30] */ MatcherIndex(29),
+ /* [30] */ MatcherIndex(31),
/* [31] */ MatcherIndex(8),
/* [32] */ MatcherIndex(4),
- /* [33] */ MatcherIndex(26),
+ /* [33] */ MatcherIndex(28),
/* [34] */ MatcherIndex(9),
/* [35] */ MatcherIndex(4),
- /* [36] */ MatcherIndex(27),
+ /* [36] */ MatcherIndex(29),
/* [37] */ MatcherIndex(9),
/* [38] */ MatcherIndex(4),
- /* [39] */ MatcherIndex(28),
+ /* [39] */ MatcherIndex(30),
/* [40] */ MatcherIndex(9),
/* [41] */ MatcherIndex(4),
- /* [42] */ MatcherIndex(29),
+ /* [42] */ MatcherIndex(31),
/* [43] */ MatcherIndex(9),
/* [44] */ MatcherIndex(4),
- /* [45] */ MatcherIndex(26),
+ /* [45] */ MatcherIndex(28),
/* [46] */ MatcherIndex(10),
/* [47] */ MatcherIndex(4),
- /* [48] */ MatcherIndex(27),
+ /* [48] */ MatcherIndex(29),
/* [49] */ MatcherIndex(10),
/* [50] */ MatcherIndex(4),
- /* [51] */ MatcherIndex(28),
+ /* [51] */ MatcherIndex(30),
/* [52] */ MatcherIndex(10),
/* [53] */ MatcherIndex(4),
- /* [54] */ MatcherIndex(29),
+ /* [54] */ MatcherIndex(31),
/* [55] */ MatcherIndex(10),
/* [56] */ MatcherIndex(4),
- /* [57] */ MatcherIndex(26),
+ /* [57] */ MatcherIndex(28),
/* [58] */ MatcherIndex(8),
/* [59] */ MatcherIndex(5),
- /* [60] */ MatcherIndex(27),
+ /* [60] */ MatcherIndex(29),
/* [61] */ MatcherIndex(8),
/* [62] */ MatcherIndex(5),
- /* [63] */ MatcherIndex(28),
+ /* [63] */ MatcherIndex(30),
/* [64] */ MatcherIndex(8),
/* [65] */ MatcherIndex(5),
- /* [66] */ MatcherIndex(29),
+ /* [66] */ MatcherIndex(31),
/* [67] */ MatcherIndex(8),
/* [68] */ MatcherIndex(5),
- /* [69] */ MatcherIndex(26),
+ /* [69] */ MatcherIndex(28),
/* [70] */ MatcherIndex(9),
/* [71] */ MatcherIndex(5),
- /* [72] */ MatcherIndex(27),
+ /* [72] */ MatcherIndex(29),
/* [73] */ MatcherIndex(9),
/* [74] */ MatcherIndex(5),
- /* [75] */ MatcherIndex(28),
+ /* [75] */ MatcherIndex(30),
/* [76] */ MatcherIndex(9),
/* [77] */ MatcherIndex(5),
- /* [78] */ MatcherIndex(29),
+ /* [78] */ MatcherIndex(31),
/* [79] */ MatcherIndex(9),
/* [80] */ MatcherIndex(5),
- /* [81] */ MatcherIndex(26),
+ /* [81] */ MatcherIndex(28),
/* [82] */ MatcherIndex(10),
/* [83] */ MatcherIndex(5),
- /* [84] */ MatcherIndex(27),
+ /* [84] */ MatcherIndex(29),
/* [85] */ MatcherIndex(10),
/* [86] */ MatcherIndex(5),
- /* [87] */ MatcherIndex(28),
+ /* [87] */ MatcherIndex(30),
/* [88] */ MatcherIndex(10),
/* [89] */ MatcherIndex(5),
- /* [90] */ MatcherIndex(29),
+ /* [90] */ MatcherIndex(31),
/* [91] */ MatcherIndex(10),
/* [92] */ MatcherIndex(5),
- /* [93] */ MatcherIndex(9),
+ /* [93] */ MatcherIndex(11),
/* [94] */ MatcherIndex(0),
- /* [95] */ MatcherIndex(15),
- /* [96] */ MatcherIndex(0),
- /* [97] */ MatcherIndex(7),
- /* [98] */ MatcherIndex(6),
- /* [99] */ MatcherIndex(7),
- /* [100] */ MatcherIndex(4),
- /* [101] */ MatcherIndex(16),
- /* [102] */ MatcherIndex(0),
- /* [103] */ MatcherIndex(18),
- /* [104] */ MatcherIndex(0),
- /* [105] */ MatcherIndex(8),
- /* [106] */ MatcherIndex(6),
- /* [107] */ MatcherIndex(19),
- /* [108] */ MatcherIndex(0),
- /* [109] */ MatcherIndex(9),
- /* [110] */ MatcherIndex(6),
- /* [111] */ MatcherIndex(14),
- /* [112] */ MatcherIndex(0),
- /* [113] */ MatcherIndex(17),
- /* [114] */ MatcherIndex(0),
- /* [115] */ MatcherIndex(25),
- /* [116] */ MatcherIndex(0),
- /* [117] */ MatcherIndex(7),
- /* [118] */ MatcherIndex(5),
- /* [119] */ MatcherIndex(14),
- /* [120] */ MatcherIndex(6),
- /* [121] */ MatcherIndex(15),
- /* [122] */ MatcherIndex(6),
- /* [123] */ MatcherIndex(16),
- /* [124] */ MatcherIndex(6),
- /* [125] */ MatcherIndex(17),
- /* [126] */ MatcherIndex(6),
- /* [127] */ MatcherIndex(18),
- /* [128] */ MatcherIndex(6),
- /* [129] */ MatcherIndex(19),
- /* [130] */ MatcherIndex(6),
- /* [131] */ MatcherIndex(35),
- /* [132] */ MatcherIndex(36),
- /* [133] */ MatcherIndex(12),
- /* [134] */ MatcherIndex(20),
- /* [135] */ MatcherIndex(21),
- /* [136] */ MatcherIndex(22),
- /* [137] */ MatcherIndex(23),
- /* [138] */ MatcherIndex(13),
- /* [139] */ MatcherIndex(24),
+ /* [95] */ MatcherIndex(1),
+ /* [96] */ MatcherIndex(10),
+ /* [97] */ MatcherIndex(0),
+ /* [98] */ MatcherIndex(17),
+ /* [99] */ MatcherIndex(0),
+ /* [100] */ MatcherIndex(8),
+ /* [101] */ MatcherIndex(6),
+ /* [102] */ MatcherIndex(18),
+ /* [103] */ MatcherIndex(0),
+ /* [104] */ MatcherIndex(20),
+ /* [105] */ MatcherIndex(0),
+ /* [106] */ MatcherIndex(9),
+ /* [107] */ MatcherIndex(6),
+ /* [108] */ MatcherIndex(21),
+ /* [109] */ MatcherIndex(0),
+ /* [110] */ MatcherIndex(10),
+ /* [111] */ MatcherIndex(6),
+ /* [112] */ MatcherIndex(16),
+ /* [113] */ MatcherIndex(0),
+ /* [114] */ MatcherIndex(19),
+ /* [115] */ MatcherIndex(0),
+ /* [116] */ MatcherIndex(27),
+ /* [117] */ MatcherIndex(0),
+ /* [118] */ MatcherIndex(16),
+ /* [119] */ MatcherIndex(6),
+ /* [120] */ MatcherIndex(17),
+ /* [121] */ MatcherIndex(6),
+ /* [122] */ MatcherIndex(18),
+ /* [123] */ MatcherIndex(6),
+ /* [124] */ MatcherIndex(19),
+ /* [125] */ MatcherIndex(6),
+ /* [126] */ MatcherIndex(20),
+ /* [127] */ MatcherIndex(6),
+ /* [128] */ MatcherIndex(21),
+ /* [129] */ MatcherIndex(6),
+ /* [130] */ MatcherIndex(37),
+ /* [131] */ MatcherIndex(7),
+ /* [132] */ MatcherIndex(38),
+ /* [133] */ MatcherIndex(14),
+ /* [134] */ MatcherIndex(22),
+ /* [135] */ MatcherIndex(23),
+ /* [136] */ MatcherIndex(24),
+ /* [137] */ MatcherIndex(25),
+ /* [138] */ MatcherIndex(15),
+ /* [139] */ MatcherIndex(26),
/* [140] */ MatcherIndex(2),
- /* [141] */ MatcherIndex(34),
- /* [142] */ MatcherIndex(30),
- /* [143] */ MatcherIndex(31),
- /* [144] */ MatcherIndex(32),
- /* [145] */ MatcherIndex(33),
+ /* [141] */ MatcherIndex(36),
+ /* [142] */ MatcherIndex(32),
+ /* [143] */ MatcherIndex(33),
+ /* [144] */ MatcherIndex(34),
+ /* [145] */ MatcherIndex(35),
+ /* [146] */ MatcherIndex(39),
};
static_assert(MatcherIndicesIndex::CanIndex(kMatcherIndices),
@@ -1050,7 +1111,7 @@
{
/* [2] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(97),
+ /* matcher_indices */ MatcherIndicesIndex(100),
},
{
/* [3] */
@@ -1070,12 +1131,12 @@
{
/* [6] */
/* usage */ core::ParameterUsage::kOffset,
- /* matcher_indices */ MatcherIndicesIndex(99),
+ /* matcher_indices */ MatcherIndicesIndex(22),
},
{
/* [7] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(101),
+ /* matcher_indices */ MatcherIndicesIndex(102),
},
{
/* [8] */
@@ -1085,7 +1146,7 @@
{
/* [9] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(97),
+ /* matcher_indices */ MatcherIndicesIndex(100),
},
{
/* [10] */
@@ -1095,7 +1156,7 @@
{
/* [11] */
/* usage */ core::ParameterUsage::kOffset,
- /* matcher_indices */ MatcherIndicesIndex(99),
+ /* matcher_indices */ MatcherIndicesIndex(22),
},
{
/* [12] */
@@ -1115,7 +1176,7 @@
{
/* [15] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(97),
+ /* matcher_indices */ MatcherIndicesIndex(100),
},
{
/* [16] */
@@ -1130,12 +1191,12 @@
{
/* [18] */
/* usage */ core::ParameterUsage::kOffset,
- /* matcher_indices */ MatcherIndicesIndex(99),
+ /* matcher_indices */ MatcherIndicesIndex(22),
},
{
/* [19] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(123),
+ /* matcher_indices */ MatcherIndicesIndex(122),
},
{
/* [20] */
@@ -1145,7 +1206,7 @@
{
/* [21] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(97),
+ /* matcher_indices */ MatcherIndicesIndex(100),
},
{
/* [22] */
@@ -1160,7 +1221,7 @@
{
/* [24] */
/* usage */ core::ParameterUsage::kOffset,
- /* matcher_indices */ MatcherIndicesIndex(99),
+ /* matcher_indices */ MatcherIndicesIndex(22),
},
{
/* [25] */
@@ -1175,7 +1236,7 @@
{
/* [27] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(97),
+ /* matcher_indices */ MatcherIndicesIndex(100),
},
{
/* [28] */
@@ -1190,12 +1251,12 @@
{
/* [30] */
/* usage */ core::ParameterUsage::kOffset,
- /* matcher_indices */ MatcherIndicesIndex(99),
+ /* matcher_indices */ MatcherIndicesIndex(22),
},
{
/* [31] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(123),
+ /* matcher_indices */ MatcherIndicesIndex(122),
},
{
/* [32] */
@@ -1205,7 +1266,7 @@
{
/* [33] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(97),
+ /* matcher_indices */ MatcherIndicesIndex(100),
},
{
/* [34] */
@@ -1220,12 +1281,12 @@
{
/* [36] */
/* usage */ core::ParameterUsage::kOffset,
- /* matcher_indices */ MatcherIndicesIndex(99),
+ /* matcher_indices */ MatcherIndicesIndex(22),
},
{
/* [37] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(123),
+ /* matcher_indices */ MatcherIndicesIndex(122),
},
{
/* [38] */
@@ -1235,7 +1296,7 @@
{
/* [39] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(97),
+ /* matcher_indices */ MatcherIndicesIndex(100),
},
{
/* [40] */
@@ -1250,7 +1311,7 @@
{
/* [42] */
/* usage */ core::ParameterUsage::kOffset,
- /* matcher_indices */ MatcherIndicesIndex(99),
+ /* matcher_indices */ MatcherIndicesIndex(22),
},
{
/* [43] */
@@ -1265,7 +1326,7 @@
{
/* [45] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(97),
+ /* matcher_indices */ MatcherIndicesIndex(100),
},
{
/* [46] */
@@ -1280,7 +1341,7 @@
{
/* [48] */
/* usage */ core::ParameterUsage::kOffset,
- /* matcher_indices */ MatcherIndicesIndex(99),
+ /* matcher_indices */ MatcherIndicesIndex(22),
},
{
/* [49] */
@@ -1295,7 +1356,7 @@
{
/* [51] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(105),
+ /* matcher_indices */ MatcherIndicesIndex(106),
},
{
/* [52] */
@@ -1340,7 +1401,7 @@
{
/* [60] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(95),
+ /* matcher_indices */ MatcherIndicesIndex(98),
},
{
/* [61] */
@@ -1350,12 +1411,12 @@
{
/* [62] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(97),
+ /* matcher_indices */ MatcherIndicesIndex(100),
},
{
/* [63] */
/* usage */ core::ParameterUsage::kOffset,
- /* matcher_indices */ MatcherIndicesIndex(99),
+ /* matcher_indices */ MatcherIndicesIndex(22),
},
{
/* [64] */
@@ -1365,7 +1426,7 @@
{
/* [65] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(107),
+ /* matcher_indices */ MatcherIndicesIndex(108),
},
{
/* [66] */
@@ -1375,7 +1436,7 @@
{
/* [67] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(105),
+ /* matcher_indices */ MatcherIndicesIndex(106),
},
{
/* [68] */
@@ -1400,7 +1461,7 @@
{
/* [72] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(97),
+ /* matcher_indices */ MatcherIndicesIndex(100),
},
{
/* [73] */
@@ -1410,7 +1471,7 @@
{
/* [74] */
/* usage */ core::ParameterUsage::kOffset,
- /* matcher_indices */ MatcherIndicesIndex(99),
+ /* matcher_indices */ MatcherIndicesIndex(22),
},
{
/* [75] */
@@ -1425,7 +1486,7 @@
{
/* [77] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(97),
+ /* matcher_indices */ MatcherIndicesIndex(100),
},
{
/* [78] */
@@ -1435,12 +1496,12 @@
{
/* [79] */
/* usage */ core::ParameterUsage::kOffset,
- /* matcher_indices */ MatcherIndicesIndex(99),
+ /* matcher_indices */ MatcherIndicesIndex(22),
},
{
/* [80] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(123),
+ /* matcher_indices */ MatcherIndicesIndex(122),
},
{
/* [81] */
@@ -1450,7 +1511,7 @@
{
/* [82] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(97),
+ /* matcher_indices */ MatcherIndicesIndex(100),
},
{
/* [83] */
@@ -1460,12 +1521,12 @@
{
/* [84] */
/* usage */ core::ParameterUsage::kOffset,
- /* matcher_indices */ MatcherIndicesIndex(99),
+ /* matcher_indices */ MatcherIndicesIndex(22),
},
{
/* [85] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(121),
+ /* matcher_indices */ MatcherIndicesIndex(120),
},
{
/* [86] */
@@ -1475,7 +1536,7 @@
{
/* [87] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(97),
+ /* matcher_indices */ MatcherIndicesIndex(100),
},
{
/* [88] */
@@ -1485,12 +1546,12 @@
{
/* [89] */
/* usage */ core::ParameterUsage::kOffset,
- /* matcher_indices */ MatcherIndicesIndex(99),
+ /* matcher_indices */ MatcherIndicesIndex(22),
},
{
/* [90] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(125),
+ /* matcher_indices */ MatcherIndicesIndex(124),
},
{
/* [91] */
@@ -1500,7 +1561,7 @@
{
/* [92] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(105),
+ /* matcher_indices */ MatcherIndicesIndex(106),
},
{
/* [93] */
@@ -1510,12 +1571,12 @@
{
/* [94] */
/* usage */ core::ParameterUsage::kOffset,
- /* matcher_indices */ MatcherIndicesIndex(22),
+ /* matcher_indices */ MatcherIndicesIndex(34),
},
{
/* [95] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(129),
+ /* matcher_indices */ MatcherIndicesIndex(128),
},
{
/* [96] */
@@ -1525,7 +1586,7 @@
{
/* [97] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(105),
+ /* matcher_indices */ MatcherIndicesIndex(106),
},
{
/* [98] */
@@ -1550,7 +1611,7 @@
{
/* [102] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(97),
+ /* matcher_indices */ MatcherIndicesIndex(100),
},
{
/* [103] */
@@ -1560,7 +1621,7 @@
{
/* [104] */
/* usage */ core::ParameterUsage::kOffset,
- /* matcher_indices */ MatcherIndicesIndex(99),
+ /* matcher_indices */ MatcherIndicesIndex(22),
},
{
/* [105] */
@@ -1575,7 +1636,7 @@
{
/* [107] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(105),
+ /* matcher_indices */ MatcherIndicesIndex(106),
},
{
/* [108] */
@@ -1590,7 +1651,7 @@
{
/* [110] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(121),
+ /* matcher_indices */ MatcherIndicesIndex(120),
},
{
/* [111] */
@@ -1600,7 +1661,7 @@
{
/* [112] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(97),
+ /* matcher_indices */ MatcherIndicesIndex(100),
},
{
/* [113] */
@@ -1610,12 +1671,12 @@
{
/* [114] */
/* usage */ core::ParameterUsage::kOffset,
- /* matcher_indices */ MatcherIndicesIndex(99),
+ /* matcher_indices */ MatcherIndicesIndex(22),
},
{
/* [115] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(125),
+ /* matcher_indices */ MatcherIndicesIndex(124),
},
{
/* [116] */
@@ -1625,7 +1686,7 @@
{
/* [117] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(105),
+ /* matcher_indices */ MatcherIndicesIndex(106),
},
{
/* [118] */
@@ -1635,12 +1696,12 @@
{
/* [119] */
/* usage */ core::ParameterUsage::kOffset,
- /* matcher_indices */ MatcherIndicesIndex(22),
+ /* matcher_indices */ MatcherIndicesIndex(34),
},
{
/* [120] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(129),
+ /* matcher_indices */ MatcherIndicesIndex(128),
},
{
/* [121] */
@@ -1650,7 +1711,7 @@
{
/* [122] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(105),
+ /* matcher_indices */ MatcherIndicesIndex(106),
},
{
/* [123] */
@@ -1665,7 +1726,7 @@
{
/* [125] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(121),
+ /* matcher_indices */ MatcherIndicesIndex(120),
},
{
/* [126] */
@@ -1675,7 +1736,7 @@
{
/* [127] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(97),
+ /* matcher_indices */ MatcherIndicesIndex(100),
},
{
/* [128] */
@@ -1685,12 +1746,12 @@
{
/* [129] */
/* usage */ core::ParameterUsage::kOffset,
- /* matcher_indices */ MatcherIndicesIndex(99),
+ /* matcher_indices */ MatcherIndicesIndex(22),
},
{
/* [130] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(125),
+ /* matcher_indices */ MatcherIndicesIndex(124),
},
{
/* [131] */
@@ -1700,7 +1761,7 @@
{
/* [132] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(105),
+ /* matcher_indices */ MatcherIndicesIndex(106),
},
{
/* [133] */
@@ -1710,12 +1771,12 @@
{
/* [134] */
/* usage */ core::ParameterUsage::kOffset,
- /* matcher_indices */ MatcherIndicesIndex(22),
+ /* matcher_indices */ MatcherIndicesIndex(34),
},
{
/* [135] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(129),
+ /* matcher_indices */ MatcherIndicesIndex(128),
},
{
/* [136] */
@@ -1725,7 +1786,7 @@
{
/* [137] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(105),
+ /* matcher_indices */ MatcherIndicesIndex(106),
},
{
/* [138] */
@@ -1750,7 +1811,7 @@
{
/* [142] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(105),
+ /* matcher_indices */ MatcherIndicesIndex(106),
},
{
/* [143] */
@@ -1765,7 +1826,7 @@
{
/* [145] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(103),
+ /* matcher_indices */ MatcherIndicesIndex(104),
},
{
/* [146] */
@@ -1775,7 +1836,7 @@
{
/* [147] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(105),
+ /* matcher_indices */ MatcherIndicesIndex(106),
},
{
/* [148] */
@@ -1795,22 +1856,22 @@
{
/* [151] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(97),
+ /* matcher_indices */ MatcherIndicesIndex(100),
},
{
/* [152] */
/* usage */ core::ParameterUsage::kOffset,
- /* matcher_indices */ MatcherIndicesIndex(99),
+ /* matcher_indices */ MatcherIndicesIndex(22),
},
{
/* [153] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(101),
+ /* matcher_indices */ MatcherIndicesIndex(102),
},
{
/* [154] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(117),
+ /* matcher_indices */ MatcherIndicesIndex(58),
},
{
/* [155] */
@@ -1830,7 +1891,7 @@
{
/* [158] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(117),
+ /* matcher_indices */ MatcherIndicesIndex(58),
},
{
/* [159] */
@@ -1845,7 +1906,7 @@
{
/* [161] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(121),
+ /* matcher_indices */ MatcherIndicesIndex(120),
},
{
/* [162] */
@@ -1855,17 +1916,17 @@
{
/* [163] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(97),
+ /* matcher_indices */ MatcherIndicesIndex(100),
},
{
/* [164] */
/* usage */ core::ParameterUsage::kOffset,
- /* matcher_indices */ MatcherIndicesIndex(99),
+ /* matcher_indices */ MatcherIndicesIndex(22),
},
{
/* [165] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(125),
+ /* matcher_indices */ MatcherIndicesIndex(124),
},
{
/* [166] */
@@ -1875,17 +1936,17 @@
{
/* [167] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(105),
+ /* matcher_indices */ MatcherIndicesIndex(106),
},
{
/* [168] */
/* usage */ core::ParameterUsage::kOffset,
- /* matcher_indices */ MatcherIndicesIndex(22),
+ /* matcher_indices */ MatcherIndicesIndex(34),
},
{
/* [169] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(127),
+ /* matcher_indices */ MatcherIndicesIndex(126),
},
{
/* [170] */
@@ -1895,7 +1956,7 @@
{
/* [171] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(105),
+ /* matcher_indices */ MatcherIndicesIndex(106),
},
{
/* [172] */
@@ -1915,7 +1976,7 @@
{
/* [175] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(105),
+ /* matcher_indices */ MatcherIndicesIndex(106),
},
{
/* [176] */
@@ -1925,7 +1986,7 @@
{
/* [177] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(127),
+ /* matcher_indices */ MatcherIndicesIndex(126),
},
{
/* [178] */
@@ -1935,7 +1996,7 @@
{
/* [179] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(105),
+ /* matcher_indices */ MatcherIndicesIndex(106),
},
{
/* [180] */
@@ -1945,7 +2006,7 @@
{
/* [181] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(127),
+ /* matcher_indices */ MatcherIndicesIndex(126),
},
{
/* [182] */
@@ -1955,7 +2016,7 @@
{
/* [183] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(105),
+ /* matcher_indices */ MatcherIndicesIndex(106),
},
{
/* [184] */
@@ -1970,12 +2031,12 @@
{
/* [186] */
/* usage */ core::ParameterUsage::kValue,
- /* matcher_indices */ MatcherIndicesIndex(109),
+ /* matcher_indices */ MatcherIndicesIndex(110),
},
{
/* [187] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(117),
+ /* matcher_indices */ MatcherIndicesIndex(58),
},
{
/* [188] */
@@ -1990,12 +2051,12 @@
{
/* [190] */
/* usage */ core::ParameterUsage::kValue,
- /* matcher_indices */ MatcherIndicesIndex(34),
+ /* matcher_indices */ MatcherIndicesIndex(46),
},
{
/* [191] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(117),
+ /* matcher_indices */ MatcherIndicesIndex(58),
},
{
/* [192] */
@@ -2010,12 +2071,12 @@
{
/* [194] */
/* usage */ core::ParameterUsage::kValue,
- /* matcher_indices */ MatcherIndicesIndex(70),
+ /* matcher_indices */ MatcherIndicesIndex(82),
},
{
/* [195] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(117),
+ /* matcher_indices */ MatcherIndicesIndex(58),
},
{
/* [196] */
@@ -2040,12 +2101,12 @@
{
/* [200] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(95),
+ /* matcher_indices */ MatcherIndicesIndex(98),
},
{
/* [201] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(117),
+ /* matcher_indices */ MatcherIndicesIndex(58),
},
{
/* [202] */
@@ -2055,12 +2116,12 @@
{
/* [203] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(113),
+ /* matcher_indices */ MatcherIndicesIndex(114),
},
{
/* [204] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(58),
+ /* matcher_indices */ MatcherIndicesIndex(70),
},
{
/* [205] */
@@ -2070,12 +2131,12 @@
{
/* [206] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(116),
},
{
/* [207] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(117),
+ /* matcher_indices */ MatcherIndicesIndex(58),
},
{
/* [208] */
@@ -2090,7 +2151,7 @@
{
/* [210] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(117),
+ /* matcher_indices */ MatcherIndicesIndex(58),
},
{
/* [211] */
@@ -2105,7 +2166,7 @@
{
/* [213] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(117),
+ /* matcher_indices */ MatcherIndicesIndex(58),
},
{
/* [214] */
@@ -2120,7 +2181,7 @@
{
/* [216] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(117),
+ /* matcher_indices */ MatcherIndicesIndex(58),
},
{
/* [217] */
@@ -2135,7 +2196,7 @@
{
/* [219] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(117),
+ /* matcher_indices */ MatcherIndicesIndex(58),
},
{
/* [220] */
@@ -2150,7 +2211,7 @@
{
/* [222] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(117),
+ /* matcher_indices */ MatcherIndicesIndex(58),
},
{
/* [223] */
@@ -2160,7 +2221,7 @@
{
/* [224] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(119),
+ /* matcher_indices */ MatcherIndicesIndex(118),
},
{
/* [225] */
@@ -2180,7 +2241,7 @@
{
/* [228] */
/* usage */ core::ParameterUsage::kValue,
- /* matcher_indices */ MatcherIndicesIndex(109),
+ /* matcher_indices */ MatcherIndicesIndex(110),
},
{
/* [229] */
@@ -2195,12 +2256,12 @@
{
/* [231] */
/* usage */ core::ParameterUsage::kValue,
- /* matcher_indices */ MatcherIndicesIndex(109),
+ /* matcher_indices */ MatcherIndicesIndex(110),
},
{
/* [232] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(117),
+ /* matcher_indices */ MatcherIndicesIndex(58),
},
{
/* [233] */
@@ -2210,12 +2271,12 @@
{
/* [234] */
/* usage */ core::ParameterUsage::kValue,
- /* matcher_indices */ MatcherIndicesIndex(109),
+ /* matcher_indices */ MatcherIndicesIndex(110),
},
{
/* [235] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(58),
+ /* matcher_indices */ MatcherIndicesIndex(70),
},
{
/* [236] */
@@ -2225,7 +2286,7 @@
{
/* [237] */
/* usage */ core::ParameterUsage::kValue,
- /* matcher_indices */ MatcherIndicesIndex(34),
+ /* matcher_indices */ MatcherIndicesIndex(46),
},
{
/* [238] */
@@ -2240,12 +2301,12 @@
{
/* [240] */
/* usage */ core::ParameterUsage::kValue,
- /* matcher_indices */ MatcherIndicesIndex(34),
+ /* matcher_indices */ MatcherIndicesIndex(46),
},
{
/* [241] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(117),
+ /* matcher_indices */ MatcherIndicesIndex(58),
},
{
/* [242] */
@@ -2255,12 +2316,12 @@
{
/* [243] */
/* usage */ core::ParameterUsage::kValue,
- /* matcher_indices */ MatcherIndicesIndex(34),
+ /* matcher_indices */ MatcherIndicesIndex(46),
},
{
/* [244] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(58),
+ /* matcher_indices */ MatcherIndicesIndex(70),
},
{
/* [245] */
@@ -2270,7 +2331,7 @@
{
/* [246] */
/* usage */ core::ParameterUsage::kValue,
- /* matcher_indices */ MatcherIndicesIndex(70),
+ /* matcher_indices */ MatcherIndicesIndex(82),
},
{
/* [247] */
@@ -2285,12 +2346,12 @@
{
/* [249] */
/* usage */ core::ParameterUsage::kValue,
- /* matcher_indices */ MatcherIndicesIndex(70),
+ /* matcher_indices */ MatcherIndicesIndex(82),
},
{
/* [250] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(117),
+ /* matcher_indices */ MatcherIndicesIndex(58),
},
{
/* [251] */
@@ -2300,12 +2361,12 @@
{
/* [252] */
/* usage */ core::ParameterUsage::kValue,
- /* matcher_indices */ MatcherIndicesIndex(70),
+ /* matcher_indices */ MatcherIndicesIndex(82),
},
{
/* [253] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(58),
+ /* matcher_indices */ MatcherIndicesIndex(70),
},
{
/* [254] */
@@ -2320,7 +2381,7 @@
{
/* [256] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(95),
+ /* matcher_indices */ MatcherIndicesIndex(98),
},
{
/* [257] */
@@ -2330,7 +2391,7 @@
{
/* [258] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(101),
+ /* matcher_indices */ MatcherIndicesIndex(102),
},
{
/* [259] */
@@ -2340,7 +2401,7 @@
{
/* [260] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(113),
+ /* matcher_indices */ MatcherIndicesIndex(114),
},
{
/* [261] */
@@ -2350,7 +2411,7 @@
{
/* [262] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(103),
+ /* matcher_indices */ MatcherIndicesIndex(104),
},
{
/* [263] */
@@ -2360,7 +2421,7 @@
{
/* [264] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(107),
+ /* matcher_indices */ MatcherIndicesIndex(108),
},
{
/* [265] */
@@ -2440,7 +2501,7 @@
{
/* [280] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(111),
+ /* matcher_indices */ MatcherIndicesIndex(112),
},
{
/* [281] */
@@ -2465,7 +2526,7 @@
{
/* [285] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(117),
+ /* matcher_indices */ MatcherIndicesIndex(58),
},
{
/* [286] */
@@ -2475,7 +2536,7 @@
{
/* [287] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(58),
+ /* matcher_indices */ MatcherIndicesIndex(70),
},
{
/* [288] */
@@ -2495,7 +2556,7 @@
{
/* [291] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(117),
+ /* matcher_indices */ MatcherIndicesIndex(58),
},
{
/* [292] */
@@ -2505,7 +2566,7 @@
{
/* [293] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(58),
+ /* matcher_indices */ MatcherIndicesIndex(70),
},
{
/* [294] */
@@ -2525,7 +2586,7 @@
{
/* [297] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(117),
+ /* matcher_indices */ MatcherIndicesIndex(58),
},
{
/* [298] */
@@ -2535,13 +2596,18 @@
{
/* [299] */
/* usage */ core::ParameterUsage::kCoords,
- /* matcher_indices */ MatcherIndicesIndex(58),
+ /* matcher_indices */ MatcherIndicesIndex(70),
},
{
/* [300] */
/* usage */ core::ParameterUsage::kTexture,
/* matcher_indices */ MatcherIndicesIndex(9),
},
+ {
+ /* [301] */
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(93),
+ },
};
static_assert(ParameterIndex::CanIndex(kParameters),
@@ -2557,25 +2623,25 @@
{
/* [1] */
/* name */ "A",
- /* matcher_indices */ MatcherIndicesIndex(131),
+ /* matcher_indices */ MatcherIndicesIndex(130),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [2] */
/* name */ "L",
- /* matcher_indices */ MatcherIndicesIndex(131),
+ /* matcher_indices */ MatcherIndicesIndex(130),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [3] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(131),
+ /* matcher_indices */ MatcherIndicesIndex(130),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [4] */
/* name */ "S",
- /* matcher_indices */ MatcherIndicesIndex(97),
+ /* matcher_indices */ MatcherIndicesIndex(131),
/* kind */ TemplateInfo::Kind::kNumber,
},
{
@@ -2599,7 +2665,7 @@
{
/* [8] */
/* name */ "L",
- /* matcher_indices */ MatcherIndicesIndex(131),
+ /* matcher_indices */ MatcherIndicesIndex(130),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -2611,7 +2677,19 @@
{
/* [10] */
/* name */ "S",
- /* matcher_indices */ MatcherIndicesIndex(131),
+ /* matcher_indices */ MatcherIndicesIndex(130),
+ /* kind */ TemplateInfo::Kind::kType,
+ },
+ {
+ /* [11] */
+ /* name */ "N",
+ /* matcher_indices */ MatcherIndicesIndex(/* invalid */),
+ /* kind */ TemplateInfo::Kind::kNumber,
+ },
+ {
+ /* [12] */
+ /* name */ "T",
+ /* matcher_indices */ MatcherIndicesIndex(146),
/* kind */ TemplateInfo::Kind::kType,
},
};
@@ -2628,7 +2706,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(224),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -2639,7 +2717,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(85),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -2650,7 +2728,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(161),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -2661,7 +2739,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(1),
/* parameters */ ParameterIndex(19),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -2672,7 +2750,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(1),
/* parameters */ ParameterIndex(80),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -2683,7 +2761,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(90),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -2694,7 +2772,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(165),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -2705,7 +2783,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(169),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -2716,7 +2794,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(1),
/* parameters */ ParameterIndex(95),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -2793,7 +2871,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(85),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -2804,7 +2882,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(85),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -2815,7 +2893,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(1),
/* parameters */ ParameterIndex(19),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -2826,7 +2904,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(1),
/* parameters */ ParameterIndex(19),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -2837,7 +2915,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(90),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -2848,7 +2926,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(90),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -2859,7 +2937,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(169),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -2870,7 +2948,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(1),
/* parameters */ ParameterIndex(95),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -2947,7 +3025,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(110),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -2958,7 +3036,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(110),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -2969,7 +3047,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(1),
/* parameters */ ParameterIndex(31),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -2980,7 +3058,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(1),
/* parameters */ ParameterIndex(31),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -2991,7 +3069,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(115),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3002,7 +3080,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(115),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3013,7 +3091,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(177),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3024,7 +3102,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(1),
/* parameters */ ParameterIndex(120),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3035,7 +3113,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(125),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3046,7 +3124,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(125),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3057,7 +3135,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(1),
/* parameters */ ParameterIndex(37),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3068,7 +3146,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(1),
/* parameters */ ParameterIndex(37),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3079,7 +3157,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(130),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3090,7 +3168,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(130),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3101,7 +3179,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(181),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3112,7 +3190,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(1),
/* parameters */ ParameterIndex(135),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3123,7 +3201,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(0),
/* parameters */ ParameterIndex(280),
- /* return_matcher_indices */ MatcherIndicesIndex(93),
+ /* return_matcher_indices */ MatcherIndicesIndex(96),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3134,7 +3212,7 @@
/* num_templates */ 2,
/* templates */ TemplateIndex(7),
/* parameters */ ParameterIndex(200),
- /* return_matcher_indices */ MatcherIndicesIndex(93),
+ /* return_matcher_indices */ MatcherIndicesIndex(96),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3145,7 +3223,7 @@
/* num_templates */ 3,
/* templates */ TemplateIndex(0),
/* parameters */ ParameterIndex(153),
- /* return_matcher_indices */ MatcherIndicesIndex(93),
+ /* return_matcher_indices */ MatcherIndicesIndex(96),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3156,7 +3234,7 @@
/* num_templates */ 2,
/* templates */ TemplateIndex(7),
/* parameters */ ParameterIndex(203),
- /* return_matcher_indices */ MatcherIndicesIndex(93),
+ /* return_matcher_indices */ MatcherIndicesIndex(96),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3167,7 +3245,7 @@
/* num_templates */ 2,
/* templates */ TemplateIndex(9),
/* parameters */ ParameterIndex(206),
- /* return_matcher_indices */ MatcherIndicesIndex(93),
+ /* return_matcher_indices */ MatcherIndicesIndex(96),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3211,7 +3289,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(282),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3222,7 +3300,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(284),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3233,7 +3311,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(1),
/* parameters */ ParameterIndex(215),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3244,7 +3322,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(286),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3255,7 +3333,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(288),
- /* return_matcher_indices */ MatcherIndicesIndex(34),
+ /* return_matcher_indices */ MatcherIndicesIndex(46),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3266,7 +3344,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(290),
- /* return_matcher_indices */ MatcherIndicesIndex(34),
+ /* return_matcher_indices */ MatcherIndicesIndex(46),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3277,7 +3355,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(1),
/* parameters */ ParameterIndex(218),
- /* return_matcher_indices */ MatcherIndicesIndex(34),
+ /* return_matcher_indices */ MatcherIndicesIndex(46),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3288,7 +3366,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(292),
- /* return_matcher_indices */ MatcherIndicesIndex(34),
+ /* return_matcher_indices */ MatcherIndicesIndex(46),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3299,7 +3377,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(294),
- /* return_matcher_indices */ MatcherIndicesIndex(70),
+ /* return_matcher_indices */ MatcherIndicesIndex(82),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3310,7 +3388,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(296),
- /* return_matcher_indices */ MatcherIndicesIndex(70),
+ /* return_matcher_indices */ MatcherIndicesIndex(82),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3321,7 +3399,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(1),
/* parameters */ ParameterIndex(221),
- /* return_matcher_indices */ MatcherIndicesIndex(70),
+ /* return_matcher_indices */ MatcherIndicesIndex(82),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3332,7 +3410,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(298),
- /* return_matcher_indices */ MatcherIndicesIndex(70),
+ /* return_matcher_indices */ MatcherIndicesIndex(82),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4047,7 +4125,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(0),
/* parameters */ ParameterIndex(60),
- /* return_matcher_indices */ MatcherIndicesIndex(93),
+ /* return_matcher_indices */ MatcherIndicesIndex(96),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4058,7 +4136,7 @@
/* num_templates */ 2,
/* templates */ TemplateIndex(0),
/* parameters */ ParameterIndex(7),
- /* return_matcher_indices */ MatcherIndicesIndex(93),
+ /* return_matcher_indices */ MatcherIndicesIndex(96),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4069,7 +4147,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(0),
/* parameters */ ParameterIndex(145),
- /* return_matcher_indices */ MatcherIndicesIndex(93),
+ /* return_matcher_indices */ MatcherIndicesIndex(96),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4080,7 +4158,7 @@
/* num_templates */ 2,
/* templates */ TemplateIndex(0),
/* parameters */ ParameterIndex(65),
- /* return_matcher_indices */ MatcherIndicesIndex(93),
+ /* return_matcher_indices */ MatcherIndicesIndex(96),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4091,7 +4169,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(149),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4102,7 +4180,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(1),
/* parameters */ ParameterIndex(70),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4113,7 +4191,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(173),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4124,7 +4202,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(1),
/* parameters */ ParameterIndex(105),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4135,7 +4213,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(43),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4146,7 +4224,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(75),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4157,7 +4235,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(1),
/* parameters */ ParameterIndex(0),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4168,7 +4246,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(1),
/* parameters */ ParameterIndex(13),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4179,7 +4257,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(140),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4190,7 +4268,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(1),
/* parameters */ ParameterIndex(49),
- /* return_matcher_indices */ MatcherIndicesIndex(109),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4382,6 +4460,17 @@
},
{
/* [160] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 2,
+ /* templates */ TemplateIndex(11),
+ /* parameters */ ParameterIndex(301),
+ /* return_matcher_indices */ MatcherIndicesIndex(1),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [161] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -4686,10 +4775,16 @@
},
{
/* [24] */
- /* fn threadgroup_barrier(u32) */
+ /* fn length[N : num, T : f32_f16](vec<N, T>) -> T */
/* num overloads */ 1,
/* overloads */ OverloadIndex(160),
},
+ {
+ /* [25] */
+ /* fn threadgroup_barrier(u32) */
+ /* num overloads */ 1,
+ /* overloads */ OverloadIndex(161),
+ },
};
// clang-format on
diff --git a/src/tint/lang/msl/msl.def b/src/tint/lang/msl/msl.def
index 9c39adb..bb569df 100644
--- a/src/tint/lang/msl/msl.def
+++ b/src/tint/lang/msl/msl.def
@@ -86,9 +86,11 @@
type i32
type u32
type f32
+type f16
type vec2<T>
type vec3<T>
type vec4<T>
+@display("vec{N}<{T}>") type vec<N: num, T>
type atomic<T>
type ptr<S: address_space, T, A: access>
type sampler
@@ -123,6 +125,7 @@
match iu32: i32 | u32
match fiu32: f32 | i32 | u32
+match f32_f16: f32 | f16
////////////////////////////////////////////////////////////////////////////////
// Builtin Functions //
@@ -318,5 +321,6 @@
@member_function fn write[A: iu32](texture: texture_storage_2d_array<u32_texel_format, writable>, value: vec4<u32>, coords: vec2<u32>, array_index: A)
@member_function fn write(texture: texture_storage_3d<u32_texel_format, writable>, value: vec4<u32>, coords: vec3<u32>)
+fn length[N: num, T: f32_f16](vec<N, T>) -> T
@stage("compute") fn threadgroup_barrier(u32)
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
index 286b8aa..2939e83 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
@@ -39,6 +39,7 @@
#include "src/tint/lang/core/ir/validator.h"
#include "src/tint/lang/core/type/depth_multisampled_texture.h"
#include "src/tint/lang/core/type/multisampled_texture.h"
+#include "src/tint/lang/core/type/scalar.h"
#include "src/tint/lang/core/type/storage_texture.h"
#include "src/tint/lang/core/type/texture.h"
#include "src/tint/lang/core/type/texture_dimension.h"
@@ -92,6 +93,7 @@
case core::BuiltinFn::kAtomicStore:
case core::BuiltinFn::kAtomicSub:
case core::BuiltinFn::kAtomicXor:
+ case core::BuiltinFn::kLength:
case core::BuiltinFn::kTextureDimensions:
case core::BuiltinFn::kTextureGather:
case core::BuiltinFn::kTextureGatherCompare:
@@ -156,6 +158,11 @@
AtomicCall(builtin, msl::BuiltinFn::kAtomicFetchXorExplicit);
break;
+ // Geometric builtins.
+ case core::BuiltinFn::kLength:
+ Length(builtin);
+ break;
+
// Texture builtins.
case core::BuiltinFn::kTextureDimensions:
TextureDimensions(builtin);
@@ -275,6 +282,22 @@
builtin->Destroy();
}
+ /// Polyfill a length call if necessary.
+ /// @param builtin the builtin call instruction
+ void Length(core::ir::CoreBuiltinCall* builtin) {
+ auto* arg = builtin->Args()[0];
+ if (arg->Type()->Is<core::type::Scalar>()) {
+ // Calls to `length` with a scalar argument are replaced with `abs`.
+ auto* call = b.CallWithResult(builtin->DetachResult(), core::BuiltinFn::kAbs, arg);
+ call->InsertBefore(builtin);
+ } else {
+ auto* call = b.CallWithResult<msl::ir::BuiltinCall>(builtin->DetachResult(),
+ msl::BuiltinFn::kLength, arg);
+ call->InsertBefore(builtin);
+ }
+ builtin->Destroy();
+ }
+
/// Replace a textureDimensions call with the equivalent MSL intrinsics.
/// @param builtin the builtin call instruction
void TextureDimensions(core::ir::CoreBuiltinCall* builtin) {
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
index 8c5fa9d..cff9ff5 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
@@ -758,6 +758,72 @@
EXPECT_EQ(expect, str());
}
+TEST_F(MslWriter_BuiltinPolyfillTest, Length_Scalar) {
+ auto* value = b.FunctionParam<f32>("value");
+ auto* func = b.Function("foo", ty.f32());
+ func->SetParams({value});
+ b.Append(func->Block(), [&] {
+ auto* result = b.Call<f32>(core::BuiltinFn::kLength, value);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%value:f32):f32 {
+ $B1: {
+ %3:f32 = length %value
+ ret %3
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%value:f32):f32 {
+ $B1: {
+ %3:f32 = abs %value
+ ret %3
+ }
+}
+)";
+
+ Run(BuiltinPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_BuiltinPolyfillTest, Length_Vector) {
+ auto* value = b.FunctionParam<vec4<f32>>("value");
+ auto* func = b.Function("foo", ty.f32());
+ func->SetParams({value});
+ b.Append(func->Block(), [&] {
+ auto* result = b.Call<f32>(core::BuiltinFn::kLength, value);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%value:vec4<f32>):f32 {
+ $B1: {
+ %3:f32 = length %value
+ ret %3
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%value:vec4<f32>):f32 {
+ $B1: {
+ %3:f32 = msl.length %value
+ ret %3
+ }
+}
+)";
+
+ Run(BuiltinPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
TEST_F(MslWriter_BuiltinPolyfillTest, TextureDimensions_1d) {
auto* t = b.FunctionParam(
"t", ty.Get<core::type::SampledTexture>(core::type::TextureDimension::k1d, ty.f32()));
diff --git a/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_builtin.wgsl.expected.ir.msl
index ce68c7d..3719645 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,42 +1,27 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, array<mat2x2<f32>, 4>, read> = var @binding_point(0, 0)
- %s:ptr<storage, f32, read_write> = var @binding_point(0, 1)
+struct tint_module_vars_struct {
+ const constant tint_array<float2x2, 4>* u;
+ device float* s;
+};
+
+kernel void f(const constant tint_array<float2x2, 4>* u [[buffer(0)]], device float* s [[buffer(1)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
+ float2x2 const t = transpose((*tint_module_vars.u)[2]);
+ float const l = length((*tint_module_vars.u)[0][1].yx);
+ float const a = abs((*tint_module_vars.u)[0][1].yx[0u]);
+ float const v = (t[0][0u] + float(l));
+ (*tint_module_vars.s) = (v + float(a));
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %4:ptr<uniform, mat2x2<f32>, read> = access %u, 2i
- %5:mat2x2<f32> = load %4
- %6:mat2x2<f32> = transpose %5
- %t:mat2x2<f32> = let %6
- %8:ptr<uniform, vec2<f32>, read> = access %u, 0i, 1i
- %9:vec2<f32> = load %8
- %10:vec2<f32> = swizzle %9, yx
- %11:f32 = length %10
- %l:f32 = let %11
- %13:ptr<uniform, vec2<f32>, read> = access %u, 0i, 1i
- %14:vec2<f32> = load %13
- %15:vec2<f32> = swizzle %14, yx
- %16:f32 = access %15, 0u
- %17:f32 = abs %16
- %a:f32 = let %17
- %19:f32 = access %t, 0i, 0u
- %20:f32 = construct %l
- %21:f32 = add %19, %20
- %22:f32 = let %21
- %23:f32 = construct %a
- %24:f32 = add %22, %23
- store %s, %24
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_builtin.wgsl.expected.ir.msl
index ced5761..7e4527f 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_builtin.wgsl.expected.ir.msl
@@ -1,42 +1,27 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, array<mat2x3<f16>, 4>, read> = var @binding_point(0, 0)
- %s:ptr<storage, f16, read_write> = var @binding_point(0, 1)
+struct tint_module_vars_struct {
+ const constant tint_array<half2x3, 4>* u;
+ device half* s;
+};
+
+kernel void f(const constant tint_array<half2x3, 4>* u [[buffer(0)]], device half* s [[buffer(1)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
+ half3x2 const t = transpose((*tint_module_vars.u)[2]);
+ half const l = length((*tint_module_vars.u)[0][1].zxy);
+ half const a = abs((*tint_module_vars.u)[0][1].zxy[0u]);
+ half const v = half(a);
+ (*tint_module_vars.s) = ((v + half(l)) + t[0][0u]);
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %4:ptr<uniform, mat2x3<f16>, read> = access %u, 2i
- %5:mat2x3<f16> = load %4
- %6:mat3x2<f16> = transpose %5
- %t:mat3x2<f16> = let %6
- %8:ptr<uniform, vec3<f16>, read> = access %u, 0i, 1i
- %9:vec3<f16> = load %8
- %10:vec3<f16> = swizzle %9, zxy
- %11:f16 = length %10
- %l:f16 = let %11
- %13:ptr<uniform, vec3<f16>, read> = access %u, 0i, 1i
- %14:vec3<f16> = load %13
- %15:vec3<f16> = swizzle %14, zxy
- %16:f16 = access %15, 0u
- %17:f16 = abs %16
- %a:f16 = let %17
- %19:f16 = construct %a
- %20:f16 = let %19
- %21:f16 = construct %l
- %22:f16 = add %20, %21
- %23:f16 = access %t, 0i, 0u
- %24:f16 = add %22, %23
- store %s, %24
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_builtin.wgsl.expected.ir.msl
index 47c1222..d8ac212 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,42 +1,27 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, array<mat2x3<f32>, 4>, read> = var @binding_point(0, 0)
- %s:ptr<storage, f32, read_write> = var @binding_point(0, 1)
+struct tint_module_vars_struct {
+ const constant tint_array<float2x3, 4>* u;
+ device float* s;
+};
+
+kernel void f(const constant tint_array<float2x3, 4>* u [[buffer(0)]], device float* s [[buffer(1)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
+ float3x2 const t = transpose((*tint_module_vars.u)[2]);
+ float const l = length((*tint_module_vars.u)[0][1].zxy);
+ float const a = abs((*tint_module_vars.u)[0][1].zxy[0u]);
+ float const v = (t[0][0u] + float(l));
+ (*tint_module_vars.s) = (v + float(a));
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %4:ptr<uniform, mat2x3<f32>, read> = access %u, 2i
- %5:mat2x3<f32> = load %4
- %6:mat3x2<f32> = transpose %5
- %t:mat3x2<f32> = let %6
- %8:ptr<uniform, vec3<f32>, read> = access %u, 0i, 1i
- %9:vec3<f32> = load %8
- %10:vec3<f32> = swizzle %9, zxy
- %11:f32 = length %10
- %l:f32 = let %11
- %13:ptr<uniform, vec3<f32>, read> = access %u, 0i, 1i
- %14:vec3<f32> = load %13
- %15:vec3<f32> = swizzle %14, zxy
- %16:f32 = access %15, 0u
- %17:f32 = abs %16
- %a:f32 = let %17
- %19:f32 = access %t, 0i, 0u
- %20:f32 = construct %l
- %21:f32 = add %19, %20
- %22:f32 = let %21
- %23:f32 = construct %a
- %24:f32 = add %22, %23
- store %s, %24
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_builtin.wgsl.expected.ir.msl
index 068f4c3..e5181d4 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_builtin.wgsl.expected.ir.msl
@@ -1,42 +1,27 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, array<mat2x4<f16>, 4>, read> = var @binding_point(0, 0)
- %s:ptr<storage, f16, read_write> = var @binding_point(0, 1)
+struct tint_module_vars_struct {
+ const constant tint_array<half2x4, 4>* u;
+ device half* s;
+};
+
+kernel void f(const constant tint_array<half2x4, 4>* u [[buffer(0)]], device half* s [[buffer(1)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
+ half4x2 const t = transpose((*tint_module_vars.u)[2]);
+ half const l = length((*tint_module_vars.u)[0][1].ywxz);
+ half const a = abs((*tint_module_vars.u)[0][1].ywxz[0u]);
+ half const v = (t[0][0u] + half(l));
+ (*tint_module_vars.s) = (v + half(a));
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %4:ptr<uniform, mat2x4<f16>, read> = access %u, 2i
- %5:mat2x4<f16> = load %4
- %6:mat4x2<f16> = transpose %5
- %t:mat4x2<f16> = let %6
- %8:ptr<uniform, vec4<f16>, read> = access %u, 0i, 1i
- %9:vec4<f16> = load %8
- %10:vec4<f16> = swizzle %9, ywxz
- %11:f16 = length %10
- %l:f16 = let %11
- %13:ptr<uniform, vec4<f16>, read> = access %u, 0i, 1i
- %14:vec4<f16> = load %13
- %15:vec4<f16> = swizzle %14, ywxz
- %16:f16 = access %15, 0u
- %17:f16 = abs %16
- %a:f16 = let %17
- %19:f16 = access %t, 0i, 0u
- %20:f16 = construct %l
- %21:f16 = add %19, %20
- %22:f16 = let %21
- %23:f16 = construct %a
- %24:f16 = add %22, %23
- store %s, %24
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_builtin.wgsl.expected.ir.msl
index 5cf3e84..65e11b3 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,42 +1,27 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, array<mat2x4<f32>, 4>, read> = var @binding_point(0, 0)
- %s:ptr<storage, f32, read_write> = var @binding_point(0, 1)
+struct tint_module_vars_struct {
+ const constant tint_array<float2x4, 4>* u;
+ device float* s;
+};
+
+kernel void f(const constant tint_array<float2x4, 4>* u [[buffer(0)]], device float* s [[buffer(1)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
+ float4x2 const t = transpose((*tint_module_vars.u)[2]);
+ float const l = length((*tint_module_vars.u)[0][1].ywxz);
+ float const a = abs((*tint_module_vars.u)[0][1].ywxz[0u]);
+ float const v = (t[0][0u] + float(l));
+ (*tint_module_vars.s) = (v + float(a));
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %4:ptr<uniform, mat2x4<f32>, read> = access %u, 2i
- %5:mat2x4<f32> = load %4
- %6:mat4x2<f32> = transpose %5
- %t:mat4x2<f32> = let %6
- %8:ptr<uniform, vec4<f32>, read> = access %u, 0i, 1i
- %9:vec4<f32> = load %8
- %10:vec4<f32> = swizzle %9, ywxz
- %11:f32 = length %10
- %l:f32 = let %11
- %13:ptr<uniform, vec4<f32>, read> = access %u, 0i, 1i
- %14:vec4<f32> = load %13
- %15:vec4<f32> = swizzle %14, ywxz
- %16:f32 = access %15, 0u
- %17:f32 = abs %16
- %a:f32 = let %17
- %19:f32 = access %t, 0i, 0u
- %20:f32 = construct %l
- %21:f32 = add %19, %20
- %22:f32 = let %21
- %23:f32 = construct %a
- %24:f32 = add %22, %23
- store %s, %24
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_builtin.wgsl.expected.ir.msl
index 246d346..96de885 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,42 +1,27 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, array<mat3x3<f32>, 4>, read> = var @binding_point(0, 0)
- %s:ptr<storage, f32, read_write> = var @binding_point(0, 1)
+struct tint_module_vars_struct {
+ const constant tint_array<float3x3, 4>* u;
+ device float* s;
+};
+
+kernel void f(const constant tint_array<float3x3, 4>* u [[buffer(0)]], device float* s [[buffer(1)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
+ float3x3 const t = transpose((*tint_module_vars.u)[2]);
+ float const l = length((*tint_module_vars.u)[0][1].zxy);
+ float const a = abs((*tint_module_vars.u)[0][1].zxy[0u]);
+ float const v = (t[0][0u] + float(l));
+ (*tint_module_vars.s) = (v + float(a));
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %4:ptr<uniform, mat3x3<f32>, read> = access %u, 2i
- %5:mat3x3<f32> = load %4
- %6:mat3x3<f32> = transpose %5
- %t:mat3x3<f32> = let %6
- %8:ptr<uniform, vec3<f32>, read> = access %u, 0i, 1i
- %9:vec3<f32> = load %8
- %10:vec3<f32> = swizzle %9, zxy
- %11:f32 = length %10
- %l:f32 = let %11
- %13:ptr<uniform, vec3<f32>, read> = access %u, 0i, 1i
- %14:vec3<f32> = load %13
- %15:vec3<f32> = swizzle %14, zxy
- %16:f32 = access %15, 0u
- %17:f32 = abs %16
- %a:f32 = let %17
- %19:f32 = access %t, 0i, 0u
- %20:f32 = construct %l
- %21:f32 = add %19, %20
- %22:f32 = let %21
- %23:f32 = construct %a
- %24:f32 = add %22, %23
- store %s, %24
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_builtin.wgsl.expected.ir.msl
index 270bf40..f55446a 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,42 +1,27 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, array<mat3x4<f32>, 4>, read> = var @binding_point(0, 0)
- %s:ptr<storage, f32, read_write> = var @binding_point(0, 1)
+struct tint_module_vars_struct {
+ const constant tint_array<float3x4, 4>* u;
+ device float* s;
+};
+
+kernel void f(const constant tint_array<float3x4, 4>* u [[buffer(0)]], device float* s [[buffer(1)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
+ float4x3 const t = transpose((*tint_module_vars.u)[2]);
+ float const l = length((*tint_module_vars.u)[0][1].ywxz);
+ float const a = abs((*tint_module_vars.u)[0][1].ywxz[0u]);
+ float const v = (t[0][0u] + float(l));
+ (*tint_module_vars.s) = (v + float(a));
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %4:ptr<uniform, mat3x4<f32>, read> = access %u, 2i
- %5:mat3x4<f32> = load %4
- %6:mat4x3<f32> = transpose %5
- %t:mat4x3<f32> = let %6
- %8:ptr<uniform, vec4<f32>, read> = access %u, 0i, 1i
- %9:vec4<f32> = load %8
- %10:vec4<f32> = swizzle %9, ywxz
- %11:f32 = length %10
- %l:f32 = let %11
- %13:ptr<uniform, vec4<f32>, read> = access %u, 0i, 1i
- %14:vec4<f32> = load %13
- %15:vec4<f32> = swizzle %14, ywxz
- %16:f32 = access %15, 0u
- %17:f32 = abs %16
- %a:f32 = let %17
- %19:f32 = access %t, 0i, 0u
- %20:f32 = construct %l
- %21:f32 = add %19, %20
- %22:f32 = let %21
- %23:f32 = construct %a
- %24:f32 = add %22, %23
- store %s, %24
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_builtin.wgsl.expected.ir.msl
index 652e7b9..af320a1 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_builtin.wgsl.expected.ir.msl
@@ -1,42 +1,27 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, array<mat4x2<f16>, 4>, read> = var @binding_point(0, 0)
- %s:ptr<storage, f16, read_write> = var @binding_point(0, 1)
+struct tint_module_vars_struct {
+ const constant tint_array<half4x2, 4>* u;
+ device half* s;
+};
+
+kernel void f(const constant tint_array<half4x2, 4>* u [[buffer(0)]], device half* s [[buffer(1)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
+ half2x4 const t = transpose((*tint_module_vars.u)[2]);
+ half const l = length((*tint_module_vars.u)[0][1].yx);
+ half const a = abs((*tint_module_vars.u)[0][1].yx[0u]);
+ half const v = (t[0][0u] + half(l));
+ (*tint_module_vars.s) = (v + half(a));
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %4:ptr<uniform, mat4x2<f16>, read> = access %u, 2i
- %5:mat4x2<f16> = load %4
- %6:mat2x4<f16> = transpose %5
- %t:mat2x4<f16> = let %6
- %8:ptr<uniform, vec2<f16>, read> = access %u, 0i, 1i
- %9:vec2<f16> = load %8
- %10:vec2<f16> = swizzle %9, yx
- %11:f16 = length %10
- %l:f16 = let %11
- %13:ptr<uniform, vec2<f16>, read> = access %u, 0i, 1i
- %14:vec2<f16> = load %13
- %15:vec2<f16> = swizzle %14, yx
- %16:f16 = access %15, 0u
- %17:f16 = abs %16
- %a:f16 = let %17
- %19:f16 = access %t, 0i, 0u
- %20:f16 = construct %l
- %21:f16 = add %19, %20
- %22:f16 = let %21
- %23:f16 = construct %a
- %24:f16 = add %22, %23
- store %s, %24
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_builtin.wgsl.expected.ir.msl
index 9459f95..2389e28 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,42 +1,27 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, array<mat4x2<f32>, 4>, read> = var @binding_point(0, 0)
- %s:ptr<storage, f32, read_write> = var @binding_point(0, 1)
+struct tint_module_vars_struct {
+ const constant tint_array<float4x2, 4>* u;
+ device float* s;
+};
+
+kernel void f(const constant tint_array<float4x2, 4>* u [[buffer(0)]], device float* s [[buffer(1)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
+ float2x4 const t = transpose((*tint_module_vars.u)[2]);
+ float const l = length((*tint_module_vars.u)[0][1].yx);
+ float const a = abs((*tint_module_vars.u)[0][1].yx[0u]);
+ float const v = (t[0][0u] + float(l));
+ (*tint_module_vars.s) = (v + float(a));
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %4:ptr<uniform, mat4x2<f32>, read> = access %u, 2i
- %5:mat4x2<f32> = load %4
- %6:mat2x4<f32> = transpose %5
- %t:mat2x4<f32> = let %6
- %8:ptr<uniform, vec2<f32>, read> = access %u, 0i, 1i
- %9:vec2<f32> = load %8
- %10:vec2<f32> = swizzle %9, yx
- %11:f32 = length %10
- %l:f32 = let %11
- %13:ptr<uniform, vec2<f32>, read> = access %u, 0i, 1i
- %14:vec2<f32> = load %13
- %15:vec2<f32> = swizzle %14, yx
- %16:f32 = access %15, 0u
- %17:f32 = abs %16
- %a:f32 = let %17
- %19:f32 = access %t, 0i, 0u
- %20:f32 = construct %l
- %21:f32 = add %19, %20
- %22:f32 = let %21
- %23:f32 = construct %a
- %24:f32 = add %22, %23
- store %s, %24
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_builtin.wgsl.expected.ir.msl
index f7e8030..a0a5f06 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_builtin.wgsl.expected.ir.msl
@@ -1,42 +1,27 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, array<mat4x3<f16>, 4>, read> = var @binding_point(0, 0)
- %s:ptr<storage, f16, read_write> = var @binding_point(0, 1)
+struct tint_module_vars_struct {
+ const constant tint_array<half4x3, 4>* u;
+ device half* s;
+};
+
+kernel void f(const constant tint_array<half4x3, 4>* u [[buffer(0)]], device half* s [[buffer(1)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
+ half3x4 const t = transpose((*tint_module_vars.u)[2]);
+ half const l = length((*tint_module_vars.u)[0][1].zxy);
+ half const a = abs((*tint_module_vars.u)[0][1].zxy[0u]);
+ half const v = (t[0][0u] + half(l));
+ (*tint_module_vars.s) = (v + half(a));
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %4:ptr<uniform, mat4x3<f16>, read> = access %u, 2i
- %5:mat4x3<f16> = load %4
- %6:mat3x4<f16> = transpose %5
- %t:mat3x4<f16> = let %6
- %8:ptr<uniform, vec3<f16>, read> = access %u, 0i, 1i
- %9:vec3<f16> = load %8
- %10:vec3<f16> = swizzle %9, zxy
- %11:f16 = length %10
- %l:f16 = let %11
- %13:ptr<uniform, vec3<f16>, read> = access %u, 0i, 1i
- %14:vec3<f16> = load %13
- %15:vec3<f16> = swizzle %14, zxy
- %16:f16 = access %15, 0u
- %17:f16 = abs %16
- %a:f16 = let %17
- %19:f16 = access %t, 0i, 0u
- %20:f16 = construct %l
- %21:f16 = add %19, %20
- %22:f16 = let %21
- %23:f16 = construct %a
- %24:f16 = add %22, %23
- store %s, %24
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_builtin.wgsl.expected.ir.msl
index 57a74ba..0da672e 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,42 +1,27 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, array<mat4x3<f32>, 4>, read> = var @binding_point(0, 0)
- %s:ptr<storage, f32, read_write> = var @binding_point(0, 1)
+struct tint_module_vars_struct {
+ const constant tint_array<float4x3, 4>* u;
+ device float* s;
+};
+
+kernel void f(const constant tint_array<float4x3, 4>* u [[buffer(0)]], device float* s [[buffer(1)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
+ float3x4 const t = transpose((*tint_module_vars.u)[2]);
+ float const l = length((*tint_module_vars.u)[0][1].zxy);
+ float const a = abs((*tint_module_vars.u)[0][1].zxy[0u]);
+ float const v = (t[0][0u] + float(l));
+ (*tint_module_vars.s) = (v + float(a));
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %4:ptr<uniform, mat4x3<f32>, read> = access %u, 2i
- %5:mat4x3<f32> = load %4
- %6:mat3x4<f32> = transpose %5
- %t:mat3x4<f32> = let %6
- %8:ptr<uniform, vec3<f32>, read> = access %u, 0i, 1i
- %9:vec3<f32> = load %8
- %10:vec3<f32> = swizzle %9, zxy
- %11:f32 = length %10
- %l:f32 = let %11
- %13:ptr<uniform, vec3<f32>, read> = access %u, 0i, 1i
- %14:vec3<f32> = load %13
- %15:vec3<f32> = swizzle %14, zxy
- %16:f32 = access %15, 0u
- %17:f32 = abs %16
- %a:f32 = let %17
- %19:f32 = access %t, 0i, 0u
- %20:f32 = construct %l
- %21:f32 = add %19, %20
- %22:f32 = let %21
- %23:f32 = construct %a
- %24:f32 = add %22, %23
- store %s, %24
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_builtin.wgsl.expected.ir.msl
index d99cf0f..669509d 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_builtin.wgsl.expected.ir.msl
@@ -1,42 +1,27 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, array<mat4x4<f16>, 4>, read> = var @binding_point(0, 0)
- %s:ptr<storage, f16, read_write> = var @binding_point(0, 1)
+struct tint_module_vars_struct {
+ const constant tint_array<half4x4, 4>* u;
+ device half* s;
+};
+
+kernel void f(const constant tint_array<half4x4, 4>* u [[buffer(0)]], device half* s [[buffer(1)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
+ half4x4 const t = transpose((*tint_module_vars.u)[2]);
+ half const l = length((*tint_module_vars.u)[0][1].ywxz);
+ half const a = abs((*tint_module_vars.u)[0][1].ywxz[0u]);
+ half const v = (t[0][0u] + half(l));
+ (*tint_module_vars.s) = (v + half(a));
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %4:ptr<uniform, mat4x4<f16>, read> = access %u, 2i
- %5:mat4x4<f16> = load %4
- %6:mat4x4<f16> = transpose %5
- %t:mat4x4<f16> = let %6
- %8:ptr<uniform, vec4<f16>, read> = access %u, 0i, 1i
- %9:vec4<f16> = load %8
- %10:vec4<f16> = swizzle %9, ywxz
- %11:f16 = length %10
- %l:f16 = let %11
- %13:ptr<uniform, vec4<f16>, read> = access %u, 0i, 1i
- %14:vec4<f16> = load %13
- %15:vec4<f16> = swizzle %14, ywxz
- %16:f16 = access %15, 0u
- %17:f16 = abs %16
- %a:f16 = let %17
- %19:f16 = access %t, 0i, 0u
- %20:f16 = construct %l
- %21:f16 = add %19, %20
- %22:f16 = let %21
- %23:f16 = construct %a
- %24:f16 = add %22, %23
- store %s, %24
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_builtin.wgsl.expected.ir.msl
index 7dbaef6..2b8d0c7 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,42 +1,27 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, array<mat4x4<f32>, 4>, read> = var @binding_point(0, 0)
- %s:ptr<storage, f32, read_write> = var @binding_point(0, 1)
+struct tint_module_vars_struct {
+ const constant tint_array<float4x4, 4>* u;
+ device float* s;
+};
+
+kernel void f(const constant tint_array<float4x4, 4>* u [[buffer(0)]], device float* s [[buffer(1)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
+ float4x4 const t = transpose((*tint_module_vars.u)[2]);
+ float const l = length((*tint_module_vars.u)[0][1].ywxz);
+ float const a = abs((*tint_module_vars.u)[0][1].ywxz[0u]);
+ float const v = (t[0][0u] + float(l));
+ (*tint_module_vars.s) = (v + float(a));
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %4:ptr<uniform, mat4x4<f32>, read> = access %u, 2i
- %5:mat4x4<f32> = load %4
- %6:mat4x4<f32> = transpose %5
- %t:mat4x4<f32> = let %6
- %8:ptr<uniform, vec4<f32>, read> = access %u, 0i, 1i
- %9:vec4<f32> = load %8
- %10:vec4<f32> = swizzle %9, ywxz
- %11:f32 = length %10
- %l:f32 = let %11
- %13:ptr<uniform, vec4<f32>, read> = access %u, 0i, 1i
- %14:vec4<f32> = load %13
- %15:vec4<f32> = swizzle %14, ywxz
- %16:f32 = access %15, 0u
- %17:f32 = abs %16
- %a:f32 = let %17
- %19:f32 = access %t, 0i, 0u
- %20:f32 = construct %l
- %21:f32 = add %19, %20
- %22:f32 = let %21
- %23:f32 = construct %a
- %24:f32 = add %22, %23
- store %s, %24
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_builtin.wgsl.expected.ir.msl
index c180b14..7c05d96 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_builtin.wgsl.expected.ir.msl
@@ -1,40 +1,29 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(64) {
- before:i32 @offset(0)
- m:mat2x2<f16> @offset(4)
- after:i32 @offset(64)
+struct S {
+ int before;
+ half2x2 m;
+ int after;
+};
+struct tint_module_vars_struct {
+ const constant tint_array<S, 4>* u;
+};
+
+kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ half2x2 const t = transpose((*tint_module_vars.u)[2].m);
+ half const l = length((*tint_module_vars.u)[0].m[1].yx);
+ half const a = abs((*tint_module_vars.u)[0].m[1].yx[0u]);
}
-
-$B1: { # root
- %u:ptr<uniform, array<S, 4>, read> = var @binding_point(0, 0)
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:ptr<uniform, mat2x2<f16>, read> = access %u, 2i, 1u
- %4:mat2x2<f16> = load %3
- %5:mat2x2<f16> = transpose %4
- %t:mat2x2<f16> = let %5
- %7:ptr<uniform, vec2<f16>, read> = access %u, 0i, 1u, 1i
- %8:vec2<f16> = load %7
- %9:vec2<f16> = swizzle %8, yx
- %10:f16 = length %9
- %l:f16 = let %10
- %12:ptr<uniform, vec2<f16>, read> = access %u, 0i, 1u, 1i
- %13:vec2<f16> = load %12
- %14:vec2<f16> = swizzle %13, yx
- %15:f16 = access %14, 0u
- %16:f16 = abs %15
- %a:f16 = let %16
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_builtin.wgsl.expected.ir.msl
index b0f4c1f..bcb8300 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,40 +1,29 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(64) {
- before:i32 @offset(0)
- m:mat2x2<f32> @offset(8)
- after:i32 @offset(64)
+struct S {
+ int before;
+ float2x2 m;
+ int after;
+};
+struct tint_module_vars_struct {
+ const constant tint_array<S, 4>* u;
+};
+
+kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ float2x2 const t = transpose((*tint_module_vars.u)[2].m);
+ float const l = length((*tint_module_vars.u)[0].m[1].yx);
+ float const a = abs((*tint_module_vars.u)[0].m[1].yx[0u]);
}
-
-$B1: { # root
- %u:ptr<uniform, array<S, 4>, read> = var @binding_point(0, 0)
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:ptr<uniform, mat2x2<f32>, read> = access %u, 2i, 1u
- %4:mat2x2<f32> = load %3
- %5:mat2x2<f32> = transpose %4
- %t:mat2x2<f32> = let %5
- %7:ptr<uniform, vec2<f32>, read> = access %u, 0i, 1u, 1i
- %8:vec2<f32> = load %7
- %9:vec2<f32> = swizzle %8, yx
- %10:f32 = length %9
- %l:f32 = let %10
- %12:ptr<uniform, vec2<f32>, read> = access %u, 0i, 1u, 1i
- %13:vec2<f32> = load %12
- %14:vec2<f32> = swizzle %13, yx
- %15:f32 = access %14, 0u
- %16:f32 = abs %15
- %a:f32 = let %16
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_builtin.wgsl.expected.ir.msl
index c34fcc2..bf6fdd2 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_builtin.wgsl.expected.ir.msl
@@ -1,40 +1,29 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(64) {
- before:i32 @offset(0)
- m:mat2x3<f16> @offset(8)
- after:i32 @offset(64)
+struct S {
+ int before;
+ half2x3 m;
+ int after;
+};
+struct tint_module_vars_struct {
+ const constant tint_array<S, 4>* u;
+};
+
+kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ half3x2 const t = transpose((*tint_module_vars.u)[2].m);
+ half const l = length((*tint_module_vars.u)[0].m[1].zxy);
+ half const a = abs((*tint_module_vars.u)[0].m[1].zxy[0u]);
}
-
-$B1: { # root
- %u:ptr<uniform, array<S, 4>, read> = var @binding_point(0, 0)
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:ptr<uniform, mat2x3<f16>, read> = access %u, 2i, 1u
- %4:mat2x3<f16> = load %3
- %5:mat3x2<f16> = transpose %4
- %t:mat3x2<f16> = let %5
- %7:ptr<uniform, vec3<f16>, read> = access %u, 0i, 1u, 1i
- %8:vec3<f16> = load %7
- %9:vec3<f16> = swizzle %8, zxy
- %10:f16 = length %9
- %l:f16 = let %10
- %12:ptr<uniform, vec3<f16>, read> = access %u, 0i, 1u, 1i
- %13:vec3<f16> = load %12
- %14:vec3<f16> = swizzle %13, zxy
- %15:f16 = access %14, 0u
- %16:f16 = abs %15
- %a:f16 = let %16
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_builtin.wgsl.expected.ir.msl
index 632586d..b221fb0 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,40 +1,29 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(64) {
- before:i32 @offset(0)
- m:mat2x3<f32> @offset(16)
- after:i32 @offset(64)
+struct S {
+ int before;
+ float2x3 m;
+ int after;
+};
+struct tint_module_vars_struct {
+ const constant tint_array<S, 4>* u;
+};
+
+kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ float3x2 const t = transpose((*tint_module_vars.u)[2].m);
+ float const l = length((*tint_module_vars.u)[0].m[1].zxy);
+ float const a = abs((*tint_module_vars.u)[0].m[1].zxy[0u]);
}
-
-$B1: { # root
- %u:ptr<uniform, array<S, 4>, read> = var @binding_point(0, 0)
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:ptr<uniform, mat2x3<f32>, read> = access %u, 2i, 1u
- %4:mat2x3<f32> = load %3
- %5:mat3x2<f32> = transpose %4
- %t:mat3x2<f32> = let %5
- %7:ptr<uniform, vec3<f32>, read> = access %u, 0i, 1u, 1i
- %8:vec3<f32> = load %7
- %9:vec3<f32> = swizzle %8, zxy
- %10:f32 = length %9
- %l:f32 = let %10
- %12:ptr<uniform, vec3<f32>, read> = access %u, 0i, 1u, 1i
- %13:vec3<f32> = load %12
- %14:vec3<f32> = swizzle %13, zxy
- %15:f32 = access %14, 0u
- %16:f32 = abs %15
- %a:f32 = let %16
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_builtin.wgsl.expected.ir.msl
index 53625ad..1ce691d 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_builtin.wgsl.expected.ir.msl
@@ -1,40 +1,29 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(64) {
- before:i32 @offset(0)
- m:mat2x4<f16> @offset(8)
- after:i32 @offset(64)
+struct S {
+ int before;
+ half2x4 m;
+ int after;
+};
+struct tint_module_vars_struct {
+ const constant tint_array<S, 4>* u;
+};
+
+kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ half4x2 const t = transpose((*tint_module_vars.u)[2].m);
+ half const l = length((*tint_module_vars.u)[0].m[1].ywxz);
+ half const a = abs((*tint_module_vars.u)[0].m[1].ywxz[0u]);
}
-
-$B1: { # root
- %u:ptr<uniform, array<S, 4>, read> = var @binding_point(0, 0)
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:ptr<uniform, mat2x4<f16>, read> = access %u, 2i, 1u
- %4:mat2x4<f16> = load %3
- %5:mat4x2<f16> = transpose %4
- %t:mat4x2<f16> = let %5
- %7:ptr<uniform, vec4<f16>, read> = access %u, 0i, 1u, 1i
- %8:vec4<f16> = load %7
- %9:vec4<f16> = swizzle %8, ywxz
- %10:f16 = length %9
- %l:f16 = let %10
- %12:ptr<uniform, vec4<f16>, read> = access %u, 0i, 1u, 1i
- %13:vec4<f16> = load %12
- %14:vec4<f16> = swizzle %13, ywxz
- %15:f16 = access %14, 0u
- %16:f16 = abs %15
- %a:f16 = let %16
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_builtin.wgsl.expected.ir.msl
index 26c5032..67fbc36 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,40 +1,29 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(64) {
- before:i32 @offset(0)
- m:mat2x4<f32> @offset(16)
- after:i32 @offset(64)
+struct S {
+ int before;
+ float2x4 m;
+ int after;
+};
+struct tint_module_vars_struct {
+ const constant tint_array<S, 4>* u;
+};
+
+kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ float4x2 const t = transpose((*tint_module_vars.u)[2].m);
+ float const l = length((*tint_module_vars.u)[0].m[1].ywxz);
+ float const a = abs((*tint_module_vars.u)[0].m[1].ywxz[0u]);
}
-
-$B1: { # root
- %u:ptr<uniform, array<S, 4>, read> = var @binding_point(0, 0)
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:ptr<uniform, mat2x4<f32>, read> = access %u, 2i, 1u
- %4:mat2x4<f32> = load %3
- %5:mat4x2<f32> = transpose %4
- %t:mat4x2<f32> = let %5
- %7:ptr<uniform, vec4<f32>, read> = access %u, 0i, 1u, 1i
- %8:vec4<f32> = load %7
- %9:vec4<f32> = swizzle %8, ywxz
- %10:f32 = length %9
- %l:f32 = let %10
- %12:ptr<uniform, vec4<f32>, read> = access %u, 0i, 1u, 1i
- %13:vec4<f32> = load %12
- %14:vec4<f32> = swizzle %13, ywxz
- %15:f32 = access %14, 0u
- %16:f32 = abs %15
- %a:f32 = let %16
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_builtin.wgsl.expected.ir.msl
index 379b718..c2c5965 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_builtin.wgsl.expected.ir.msl
@@ -1,40 +1,29 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(64) {
- before:i32 @offset(0)
- m:mat3x2<f16> @offset(4)
- after:i32 @offset(64)
+struct S {
+ int before;
+ half3x2 m;
+ int after;
+};
+struct tint_module_vars_struct {
+ const constant tint_array<S, 4>* u;
+};
+
+kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ half2x3 const t = transpose((*tint_module_vars.u)[2].m);
+ half const l = length((*tint_module_vars.u)[0].m[1].yx);
+ half const a = abs((*tint_module_vars.u)[0].m[1].yx[0u]);
}
-
-$B1: { # root
- %u:ptr<uniform, array<S, 4>, read> = var @binding_point(0, 0)
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:ptr<uniform, mat3x2<f16>, read> = access %u, 2i, 1u
- %4:mat3x2<f16> = load %3
- %5:mat2x3<f16> = transpose %4
- %t:mat2x3<f16> = let %5
- %7:ptr<uniform, vec2<f16>, read> = access %u, 0i, 1u, 1i
- %8:vec2<f16> = load %7
- %9:vec2<f16> = swizzle %8, yx
- %10:f16 = length %9
- %l:f16 = let %10
- %12:ptr<uniform, vec2<f16>, read> = access %u, 0i, 1u, 1i
- %13:vec2<f16> = load %12
- %14:vec2<f16> = swizzle %13, yx
- %15:f16 = access %14, 0u
- %16:f16 = abs %15
- %a:f16 = let %16
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_builtin.wgsl.expected.ir.msl
index 37ee1fe..4b79ab9 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,40 +1,29 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(64) {
- before:i32 @offset(0)
- m:mat3x2<f32> @offset(8)
- after:i32 @offset(64)
+struct S {
+ int before;
+ float3x2 m;
+ int after;
+};
+struct tint_module_vars_struct {
+ const constant tint_array<S, 4>* u;
+};
+
+kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ float2x3 const t = transpose((*tint_module_vars.u)[2].m);
+ float const l = length((*tint_module_vars.u)[0].m[1].yx);
+ float const a = abs((*tint_module_vars.u)[0].m[1].yx[0u]);
}
-
-$B1: { # root
- %u:ptr<uniform, array<S, 4>, read> = var @binding_point(0, 0)
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:ptr<uniform, mat3x2<f32>, read> = access %u, 2i, 1u
- %4:mat3x2<f32> = load %3
- %5:mat2x3<f32> = transpose %4
- %t:mat2x3<f32> = let %5
- %7:ptr<uniform, vec2<f32>, read> = access %u, 0i, 1u, 1i
- %8:vec2<f32> = load %7
- %9:vec2<f32> = swizzle %8, yx
- %10:f32 = length %9
- %l:f32 = let %10
- %12:ptr<uniform, vec2<f32>, read> = access %u, 0i, 1u, 1i
- %13:vec2<f32> = load %12
- %14:vec2<f32> = swizzle %13, yx
- %15:f32 = access %14, 0u
- %16:f32 = abs %15
- %a:f32 = let %16
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_builtin.wgsl.expected.ir.msl
index eecf916..6453786 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_builtin.wgsl.expected.ir.msl
@@ -1,40 +1,29 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(64) {
- before:i32 @offset(0)
- m:mat3x3<f16> @offset(8)
- after:i32 @offset(64)
+struct S {
+ int before;
+ half3x3 m;
+ int after;
+};
+struct tint_module_vars_struct {
+ const constant tint_array<S, 4>* u;
+};
+
+kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ half3x3 const t = transpose((*tint_module_vars.u)[2].m);
+ half const l = length((*tint_module_vars.u)[0].m[1].zxy);
+ half const a = abs((*tint_module_vars.u)[0].m[1].zxy[0u]);
}
-
-$B1: { # root
- %u:ptr<uniform, array<S, 4>, read> = var @binding_point(0, 0)
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:ptr<uniform, mat3x3<f16>, read> = access %u, 2i, 1u
- %4:mat3x3<f16> = load %3
- %5:mat3x3<f16> = transpose %4
- %t:mat3x3<f16> = let %5
- %7:ptr<uniform, vec3<f16>, read> = access %u, 0i, 1u, 1i
- %8:vec3<f16> = load %7
- %9:vec3<f16> = swizzle %8, zxy
- %10:f16 = length %9
- %l:f16 = let %10
- %12:ptr<uniform, vec3<f16>, read> = access %u, 0i, 1u, 1i
- %13:vec3<f16> = load %12
- %14:vec3<f16> = swizzle %13, zxy
- %15:f16 = access %14, 0u
- %16:f16 = abs %15
- %a:f16 = let %16
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_builtin.wgsl.expected.ir.msl
index 94e4040..1d1f669 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,40 +1,29 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(64) {
- before:i32 @offset(0)
- m:mat3x3<f32> @offset(16)
- after:i32 @offset(64)
+struct S {
+ int before;
+ float3x3 m;
+ int after;
+};
+struct tint_module_vars_struct {
+ const constant tint_array<S, 4>* u;
+};
+
+kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ float3x3 const t = transpose((*tint_module_vars.u)[2].m);
+ float const l = length((*tint_module_vars.u)[0].m[1].zxy);
+ float const a = abs((*tint_module_vars.u)[0].m[1].zxy[0u]);
}
-
-$B1: { # root
- %u:ptr<uniform, array<S, 4>, read> = var @binding_point(0, 0)
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:ptr<uniform, mat3x3<f32>, read> = access %u, 2i, 1u
- %4:mat3x3<f32> = load %3
- %5:mat3x3<f32> = transpose %4
- %t:mat3x3<f32> = let %5
- %7:ptr<uniform, vec3<f32>, read> = access %u, 0i, 1u, 1i
- %8:vec3<f32> = load %7
- %9:vec3<f32> = swizzle %8, zxy
- %10:f32 = length %9
- %l:f32 = let %10
- %12:ptr<uniform, vec3<f32>, read> = access %u, 0i, 1u, 1i
- %13:vec3<f32> = load %12
- %14:vec3<f32> = swizzle %13, zxy
- %15:f32 = access %14, 0u
- %16:f32 = abs %15
- %a:f32 = let %16
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_builtin.wgsl.expected.ir.msl
index 990cb8e..18840bc 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_builtin.wgsl.expected.ir.msl
@@ -1,40 +1,29 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(64) {
- before:i32 @offset(0)
- m:mat3x4<f16> @offset(8)
- after:i32 @offset(64)
+struct S {
+ int before;
+ half3x4 m;
+ int after;
+};
+struct tint_module_vars_struct {
+ const constant tint_array<S, 4>* u;
+};
+
+kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ half4x3 const t = transpose((*tint_module_vars.u)[2].m);
+ half const l = length((*tint_module_vars.u)[0].m[1].ywxz);
+ half const a = abs((*tint_module_vars.u)[0].m[1].ywxz[0u]);
}
-
-$B1: { # root
- %u:ptr<uniform, array<S, 4>, read> = var @binding_point(0, 0)
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:ptr<uniform, mat3x4<f16>, read> = access %u, 2i, 1u
- %4:mat3x4<f16> = load %3
- %5:mat4x3<f16> = transpose %4
- %t:mat4x3<f16> = let %5
- %7:ptr<uniform, vec4<f16>, read> = access %u, 0i, 1u, 1i
- %8:vec4<f16> = load %7
- %9:vec4<f16> = swizzle %8, ywxz
- %10:f16 = length %9
- %l:f16 = let %10
- %12:ptr<uniform, vec4<f16>, read> = access %u, 0i, 1u, 1i
- %13:vec4<f16> = load %12
- %14:vec4<f16> = swizzle %13, ywxz
- %15:f16 = access %14, 0u
- %16:f16 = abs %15
- %a:f16 = let %16
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_builtin.wgsl.expected.ir.msl
index 01d8863..b46a341 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,40 +1,29 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(64) {
- before:i32 @offset(0)
- m:mat3x4<f32> @offset(16)
- after:i32 @offset(64)
+struct S {
+ int before;
+ float3x4 m;
+ int after;
+};
+struct tint_module_vars_struct {
+ const constant tint_array<S, 4>* u;
+};
+
+kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ float4x3 const t = transpose((*tint_module_vars.u)[2].m);
+ float const l = length((*tint_module_vars.u)[0].m[1].ywxz);
+ float const a = abs((*tint_module_vars.u)[0].m[1].ywxz[0u]);
}
-
-$B1: { # root
- %u:ptr<uniform, array<S, 4>, read> = var @binding_point(0, 0)
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:ptr<uniform, mat3x4<f32>, read> = access %u, 2i, 1u
- %4:mat3x4<f32> = load %3
- %5:mat4x3<f32> = transpose %4
- %t:mat4x3<f32> = let %5
- %7:ptr<uniform, vec4<f32>, read> = access %u, 0i, 1u, 1i
- %8:vec4<f32> = load %7
- %9:vec4<f32> = swizzle %8, ywxz
- %10:f32 = length %9
- %l:f32 = let %10
- %12:ptr<uniform, vec4<f32>, read> = access %u, 0i, 1u, 1i
- %13:vec4<f32> = load %12
- %14:vec4<f32> = swizzle %13, ywxz
- %15:f32 = access %14, 0u
- %16:f32 = abs %15
- %a:f32 = let %16
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_builtin.wgsl.expected.ir.msl
index ae20072..6ff45a4 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_builtin.wgsl.expected.ir.msl
@@ -1,40 +1,29 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(64) {
- before:i32 @offset(0)
- m:mat4x2<f16> @offset(4)
- after:i32 @offset(64)
+struct S {
+ int before;
+ half4x2 m;
+ int after;
+};
+struct tint_module_vars_struct {
+ const constant tint_array<S, 4>* u;
+};
+
+kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ half2x4 const t = transpose((*tint_module_vars.u)[2].m);
+ half const l = length((*tint_module_vars.u)[0].m[1].yx);
+ half const a = abs((*tint_module_vars.u)[0].m[1].yx[0u]);
}
-
-$B1: { # root
- %u:ptr<uniform, array<S, 4>, read> = var @binding_point(0, 0)
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:ptr<uniform, mat4x2<f16>, read> = access %u, 2i, 1u
- %4:mat4x2<f16> = load %3
- %5:mat2x4<f16> = transpose %4
- %t:mat2x4<f16> = let %5
- %7:ptr<uniform, vec2<f16>, read> = access %u, 0i, 1u, 1i
- %8:vec2<f16> = load %7
- %9:vec2<f16> = swizzle %8, yx
- %10:f16 = length %9
- %l:f16 = let %10
- %12:ptr<uniform, vec2<f16>, read> = access %u, 0i, 1u, 1i
- %13:vec2<f16> = load %12
- %14:vec2<f16> = swizzle %13, yx
- %15:f16 = access %14, 0u
- %16:f16 = abs %15
- %a:f16 = let %16
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_builtin.wgsl.expected.ir.msl
index b1b4be2..1a2a0b6 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,40 +1,29 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(64) {
- before:i32 @offset(0)
- m:mat4x2<f32> @offset(8)
- after:i32 @offset(64)
+struct S {
+ int before;
+ float4x2 m;
+ int after;
+};
+struct tint_module_vars_struct {
+ const constant tint_array<S, 4>* u;
+};
+
+kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ float2x4 const t = transpose((*tint_module_vars.u)[2].m);
+ float const l = length((*tint_module_vars.u)[0].m[1].yx);
+ float const a = abs((*tint_module_vars.u)[0].m[1].yx[0u]);
}
-
-$B1: { # root
- %u:ptr<uniform, array<S, 4>, read> = var @binding_point(0, 0)
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:ptr<uniform, mat4x2<f32>, read> = access %u, 2i, 1u
- %4:mat4x2<f32> = load %3
- %5:mat2x4<f32> = transpose %4
- %t:mat2x4<f32> = let %5
- %7:ptr<uniform, vec2<f32>, read> = access %u, 0i, 1u, 1i
- %8:vec2<f32> = load %7
- %9:vec2<f32> = swizzle %8, yx
- %10:f32 = length %9
- %l:f32 = let %10
- %12:ptr<uniform, vec2<f32>, read> = access %u, 0i, 1u, 1i
- %13:vec2<f32> = load %12
- %14:vec2<f32> = swizzle %13, yx
- %15:f32 = access %14, 0u
- %16:f32 = abs %15
- %a:f32 = let %16
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_builtin.wgsl.expected.ir.msl
index 5cce2ca..50bb141 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_builtin.wgsl.expected.ir.msl
@@ -1,40 +1,29 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(64) {
- before:i32 @offset(0)
- m:mat4x3<f16> @offset(8)
- after:i32 @offset(64)
+struct S {
+ int before;
+ half4x3 m;
+ int after;
+};
+struct tint_module_vars_struct {
+ const constant tint_array<S, 4>* u;
+};
+
+kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ half3x4 const t = transpose((*tint_module_vars.u)[2].m);
+ half const l = length((*tint_module_vars.u)[0].m[1].zxy);
+ half const a = abs((*tint_module_vars.u)[0].m[1].zxy[0u]);
}
-
-$B1: { # root
- %u:ptr<uniform, array<S, 4>, read> = var @binding_point(0, 0)
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:ptr<uniform, mat4x3<f16>, read> = access %u, 2i, 1u
- %4:mat4x3<f16> = load %3
- %5:mat3x4<f16> = transpose %4
- %t:mat3x4<f16> = let %5
- %7:ptr<uniform, vec3<f16>, read> = access %u, 0i, 1u, 1i
- %8:vec3<f16> = load %7
- %9:vec3<f16> = swizzle %8, zxy
- %10:f16 = length %9
- %l:f16 = let %10
- %12:ptr<uniform, vec3<f16>, read> = access %u, 0i, 1u, 1i
- %13:vec3<f16> = load %12
- %14:vec3<f16> = swizzle %13, zxy
- %15:f16 = access %14, 0u
- %16:f16 = abs %15
- %a:f16 = let %16
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_builtin.wgsl.expected.ir.msl
index 43a93ce..28f0a53 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,40 +1,29 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(64) {
- before:i32 @offset(0)
- m:mat4x3<f32> @offset(16)
- after:i32 @offset(128)
+struct S {
+ int before;
+ float4x3 m;
+ int after;
+};
+struct tint_module_vars_struct {
+ const constant tint_array<S, 4>* u;
+};
+
+kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ float3x4 const t = transpose((*tint_module_vars.u)[2].m);
+ float const l = length((*tint_module_vars.u)[0].m[1].zxy);
+ float const a = abs((*tint_module_vars.u)[0].m[1].zxy[0u]);
}
-
-$B1: { # root
- %u:ptr<uniform, array<S, 4>, read> = var @binding_point(0, 0)
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:ptr<uniform, mat4x3<f32>, read> = access %u, 2i, 1u
- %4:mat4x3<f32> = load %3
- %5:mat3x4<f32> = transpose %4
- %t:mat3x4<f32> = let %5
- %7:ptr<uniform, vec3<f32>, read> = access %u, 0i, 1u, 1i
- %8:vec3<f32> = load %7
- %9:vec3<f32> = swizzle %8, zxy
- %10:f32 = length %9
- %l:f32 = let %10
- %12:ptr<uniform, vec3<f32>, read> = access %u, 0i, 1u, 1i
- %13:vec3<f32> = load %12
- %14:vec3<f32> = swizzle %13, zxy
- %15:f32 = access %14, 0u
- %16:f32 = abs %15
- %a:f32 = let %16
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_builtin.wgsl.expected.ir.msl
index a2c984e..028ac1c 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_builtin.wgsl.expected.ir.msl
@@ -1,40 +1,29 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(64) {
- before:i32 @offset(0)
- m:mat4x4<f16> @offset(8)
- after:i32 @offset(64)
+struct S {
+ int before;
+ half4x4 m;
+ int after;
+};
+struct tint_module_vars_struct {
+ const constant tint_array<S, 4>* u;
+};
+
+kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ half4x4 const t = transpose((*tint_module_vars.u)[2].m);
+ half const l = length((*tint_module_vars.u)[0].m[1].ywxz);
+ half const a = abs((*tint_module_vars.u)[0].m[1].ywxz[0u]);
}
-
-$B1: { # root
- %u:ptr<uniform, array<S, 4>, read> = var @binding_point(0, 0)
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:ptr<uniform, mat4x4<f16>, read> = access %u, 2i, 1u
- %4:mat4x4<f16> = load %3
- %5:mat4x4<f16> = transpose %4
- %t:mat4x4<f16> = let %5
- %7:ptr<uniform, vec4<f16>, read> = access %u, 0i, 1u, 1i
- %8:vec4<f16> = load %7
- %9:vec4<f16> = swizzle %8, ywxz
- %10:f16 = length %9
- %l:f16 = let %10
- %12:ptr<uniform, vec4<f16>, read> = access %u, 0i, 1u, 1i
- %13:vec4<f16> = load %12
- %14:vec4<f16> = swizzle %13, ywxz
- %15:f16 = access %14, 0u
- %16:f16 = abs %15
- %a:f16 = let %16
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_builtin.wgsl.expected.ir.msl
index 6a5a7de..e7f6b41 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,40 +1,29 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(64) {
- before:i32 @offset(0)
- m:mat4x4<f32> @offset(16)
- after:i32 @offset(128)
+struct S {
+ int before;
+ float4x4 m;
+ int after;
+};
+struct tint_module_vars_struct {
+ const constant tint_array<S, 4>* u;
+};
+
+kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ float4x4 const t = transpose((*tint_module_vars.u)[2].m);
+ float const l = length((*tint_module_vars.u)[0].m[1].ywxz);
+ float const a = abs((*tint_module_vars.u)[0].m[1].ywxz[0u]);
}
-
-$B1: { # root
- %u:ptr<uniform, array<S, 4>, read> = var @binding_point(0, 0)
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:ptr<uniform, mat4x4<f32>, read> = access %u, 2i, 1u
- %4:mat4x4<f32> = load %3
- %5:mat4x4<f32> = transpose %4
- %t:mat4x4<f32> = let %5
- %7:ptr<uniform, vec4<f32>, read> = access %u, 0i, 1u, 1i
- %8:vec4<f32> = load %7
- %9:vec4<f32> = swizzle %8, ywxz
- %10:f32 = length %9
- %l:f32 = let %10
- %12:ptr<uniform, vec4<f32>, read> = access %u, 0i, 1u, 1i
- %13:vec4<f32> = load %12
- %14:vec4<f32> = swizzle %13, ywxz
- %15:f32 = access %14, 0u
- %16:f32 = abs %15
- %a:f32 = let %16
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x2_f16/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x2_f16/to_builtin.wgsl.expected.ir.msl
index d3813373..8cc5e37 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x2_f16/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x2_f16/to_builtin.wgsl.expected.ir.msl
@@ -1,32 +1,12 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ const constant half2x2* u;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat2x2<f16>, read> = var @binding_point(0, 0)
+kernel void f(const constant half2x2* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ half2x2 const t = transpose((*tint_module_vars.u));
+ half const l = length((*tint_module_vars.u)[1]);
+ half const a = abs((*tint_module_vars.u)[0].yx[0u]);
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:mat2x2<f16> = load %u
- %4:mat2x2<f16> = transpose %3
- %t:mat2x2<f16> = let %4
- %6:ptr<uniform, vec2<f16>, read> = access %u, 1i
- %7:vec2<f16> = load %6
- %8:f16 = length %7
- %l:f16 = let %8
- %10:ptr<uniform, vec2<f16>, read> = access %u, 0i
- %11:vec2<f16> = load %10
- %12:vec2<f16> = swizzle %11, yx
- %13:f16 = access %12, 0u
- %14:f16 = abs %13
- %a:f16 = let %14
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x2_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x2_f32/to_builtin.wgsl.expected.ir.msl
index cc7db23..3be1168 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x2_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x2_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,32 +1,12 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ const constant float2x2* u;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat2x2<f32>, read> = var @binding_point(0, 0)
+kernel void f(const constant float2x2* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ float2x2 const t = transpose((*tint_module_vars.u));
+ float const l = length((*tint_module_vars.u)[1]);
+ float const a = abs((*tint_module_vars.u)[0].yx[0u]);
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:mat2x2<f32> = load %u
- %4:mat2x2<f32> = transpose %3
- %t:mat2x2<f32> = let %4
- %6:ptr<uniform, vec2<f32>, read> = access %u, 1i
- %7:vec2<f32> = load %6
- %8:f32 = length %7
- %l:f32 = let %8
- %10:ptr<uniform, vec2<f32>, read> = access %u, 0i
- %11:vec2<f32> = load %10
- %12:vec2<f32> = swizzle %11, yx
- %13:f32 = access %12, 0u
- %14:f32 = abs %13
- %a:f32 = let %14
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x3_f16/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x3_f16/to_builtin.wgsl.expected.ir.msl
index 07a3a53..7ec62ec 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x3_f16/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x3_f16/to_builtin.wgsl.expected.ir.msl
@@ -1,32 +1,12 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ const constant half2x3* u;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat2x3<f16>, read> = var @binding_point(0, 0)
+kernel void f(const constant half2x3* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ half3x2 const t = transpose((*tint_module_vars.u));
+ half const l = length((*tint_module_vars.u)[1]);
+ half const a = abs((*tint_module_vars.u)[0].zxy[0u]);
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:mat2x3<f16> = load %u
- %4:mat3x2<f16> = transpose %3
- %t:mat3x2<f16> = let %4
- %6:ptr<uniform, vec3<f16>, read> = access %u, 1i
- %7:vec3<f16> = load %6
- %8:f16 = length %7
- %l:f16 = let %8
- %10:ptr<uniform, vec3<f16>, read> = access %u, 0i
- %11:vec3<f16> = load %10
- %12:vec3<f16> = swizzle %11, zxy
- %13:f16 = access %12, 0u
- %14:f16 = abs %13
- %a:f16 = let %14
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x3_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x3_f32/to_builtin.wgsl.expected.ir.msl
index c2ff7b7..1c4cb3d 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x3_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x3_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,32 +1,12 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ const constant float2x3* u;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat2x3<f32>, read> = var @binding_point(0, 0)
+kernel void f(const constant float2x3* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ float3x2 const t = transpose((*tint_module_vars.u));
+ float const l = length((*tint_module_vars.u)[1]);
+ float const a = abs((*tint_module_vars.u)[0].zxy[0u]);
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:mat2x3<f32> = load %u
- %4:mat3x2<f32> = transpose %3
- %t:mat3x2<f32> = let %4
- %6:ptr<uniform, vec3<f32>, read> = access %u, 1i
- %7:vec3<f32> = load %6
- %8:f32 = length %7
- %l:f32 = let %8
- %10:ptr<uniform, vec3<f32>, read> = access %u, 0i
- %11:vec3<f32> = load %10
- %12:vec3<f32> = swizzle %11, zxy
- %13:f32 = access %12, 0u
- %14:f32 = abs %13
- %a:f32 = let %14
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x4_f16/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x4_f16/to_builtin.wgsl.expected.ir.msl
index 5042f28..9ff4dbc 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x4_f16/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x4_f16/to_builtin.wgsl.expected.ir.msl
@@ -1,32 +1,12 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ const constant half2x4* u;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat2x4<f16>, read> = var @binding_point(0, 0)
+kernel void f(const constant half2x4* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ half4x2 const t = transpose((*tint_module_vars.u));
+ half const l = length((*tint_module_vars.u)[1]);
+ half const a = abs((*tint_module_vars.u)[0].ywxz[0u]);
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:mat2x4<f16> = load %u
- %4:mat4x2<f16> = transpose %3
- %t:mat4x2<f16> = let %4
- %6:ptr<uniform, vec4<f16>, read> = access %u, 1i
- %7:vec4<f16> = load %6
- %8:f16 = length %7
- %l:f16 = let %8
- %10:ptr<uniform, vec4<f16>, read> = access %u, 0i
- %11:vec4<f16> = load %10
- %12:vec4<f16> = swizzle %11, ywxz
- %13:f16 = access %12, 0u
- %14:f16 = abs %13
- %a:f16 = let %14
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x4_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x4_f32/to_builtin.wgsl.expected.ir.msl
index 51feb64..edf04d6 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x4_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x4_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,32 +1,12 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ const constant float2x4* u;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat2x4<f32>, read> = var @binding_point(0, 0)
+kernel void f(const constant float2x4* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ float4x2 const t = transpose((*tint_module_vars.u));
+ float const l = length((*tint_module_vars.u)[1]);
+ float const a = abs((*tint_module_vars.u)[0].ywxz[0u]);
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:mat2x4<f32> = load %u
- %4:mat4x2<f32> = transpose %3
- %t:mat4x2<f32> = let %4
- %6:ptr<uniform, vec4<f32>, read> = access %u, 1i
- %7:vec4<f32> = load %6
- %8:f32 = length %7
- %l:f32 = let %8
- %10:ptr<uniform, vec4<f32>, read> = access %u, 0i
- %11:vec4<f32> = load %10
- %12:vec4<f32> = swizzle %11, ywxz
- %13:f32 = access %12, 0u
- %14:f32 = abs %13
- %a:f32 = let %14
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x2_f16/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x2_f16/to_builtin.wgsl.expected.ir.msl
index e904393..ffa4ad4 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x2_f16/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x2_f16/to_builtin.wgsl.expected.ir.msl
@@ -1,32 +1,12 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ const constant half3x2* u;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat3x2<f16>, read> = var @binding_point(0, 0)
+kernel void f(const constant half3x2* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ half2x3 const t = transpose((*tint_module_vars.u));
+ half const l = length((*tint_module_vars.u)[1]);
+ half const a = abs((*tint_module_vars.u)[0].yx[0u]);
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:mat3x2<f16> = load %u
- %4:mat2x3<f16> = transpose %3
- %t:mat2x3<f16> = let %4
- %6:ptr<uniform, vec2<f16>, read> = access %u, 1i
- %7:vec2<f16> = load %6
- %8:f16 = length %7
- %l:f16 = let %8
- %10:ptr<uniform, vec2<f16>, read> = access %u, 0i
- %11:vec2<f16> = load %10
- %12:vec2<f16> = swizzle %11, yx
- %13:f16 = access %12, 0u
- %14:f16 = abs %13
- %a:f16 = let %14
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x2_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x2_f32/to_builtin.wgsl.expected.ir.msl
index 3f6f8d2..b0c1468 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x2_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x2_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,32 +1,12 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ const constant float3x2* u;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat3x2<f32>, read> = var @binding_point(0, 0)
+kernel void f(const constant float3x2* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ float2x3 const t = transpose((*tint_module_vars.u));
+ float const l = length((*tint_module_vars.u)[1]);
+ float const a = abs((*tint_module_vars.u)[0].yx[0u]);
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:mat3x2<f32> = load %u
- %4:mat2x3<f32> = transpose %3
- %t:mat2x3<f32> = let %4
- %6:ptr<uniform, vec2<f32>, read> = access %u, 1i
- %7:vec2<f32> = load %6
- %8:f32 = length %7
- %l:f32 = let %8
- %10:ptr<uniform, vec2<f32>, read> = access %u, 0i
- %11:vec2<f32> = load %10
- %12:vec2<f32> = swizzle %11, yx
- %13:f32 = access %12, 0u
- %14:f32 = abs %13
- %a:f32 = let %14
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x3_f16/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x3_f16/to_builtin.wgsl.expected.ir.msl
index 921fbf7..81fff4f 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x3_f16/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x3_f16/to_builtin.wgsl.expected.ir.msl
@@ -1,32 +1,12 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ const constant half3x3* u;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat3x3<f16>, read> = var @binding_point(0, 0)
+kernel void f(const constant half3x3* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ half3x3 const t = transpose((*tint_module_vars.u));
+ half const l = length((*tint_module_vars.u)[1]);
+ half const a = abs((*tint_module_vars.u)[0].zxy[0u]);
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:mat3x3<f16> = load %u
- %4:mat3x3<f16> = transpose %3
- %t:mat3x3<f16> = let %4
- %6:ptr<uniform, vec3<f16>, read> = access %u, 1i
- %7:vec3<f16> = load %6
- %8:f16 = length %7
- %l:f16 = let %8
- %10:ptr<uniform, vec3<f16>, read> = access %u, 0i
- %11:vec3<f16> = load %10
- %12:vec3<f16> = swizzle %11, zxy
- %13:f16 = access %12, 0u
- %14:f16 = abs %13
- %a:f16 = let %14
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x3_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x3_f32/to_builtin.wgsl.expected.ir.msl
index 7a1127e..095eebe 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x3_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x3_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,32 +1,12 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ const constant float3x3* u;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat3x3<f32>, read> = var @binding_point(0, 0)
+kernel void f(const constant float3x3* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ float3x3 const t = transpose((*tint_module_vars.u));
+ float const l = length((*tint_module_vars.u)[1]);
+ float const a = abs((*tint_module_vars.u)[0].zxy[0u]);
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:mat3x3<f32> = load %u
- %4:mat3x3<f32> = transpose %3
- %t:mat3x3<f32> = let %4
- %6:ptr<uniform, vec3<f32>, read> = access %u, 1i
- %7:vec3<f32> = load %6
- %8:f32 = length %7
- %l:f32 = let %8
- %10:ptr<uniform, vec3<f32>, read> = access %u, 0i
- %11:vec3<f32> = load %10
- %12:vec3<f32> = swizzle %11, zxy
- %13:f32 = access %12, 0u
- %14:f32 = abs %13
- %a:f32 = let %14
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x4_f16/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x4_f16/to_builtin.wgsl.expected.ir.msl
index c4e57c6..2cac1b6 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x4_f16/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x4_f16/to_builtin.wgsl.expected.ir.msl
@@ -1,32 +1,12 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ const constant half3x4* u;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat3x4<f16>, read> = var @binding_point(0, 0)
+kernel void f(const constant half3x4* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ half4x3 const t = transpose((*tint_module_vars.u));
+ half const l = length((*tint_module_vars.u)[1]);
+ half const a = abs((*tint_module_vars.u)[0].ywxz[0u]);
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:mat3x4<f16> = load %u
- %4:mat4x3<f16> = transpose %3
- %t:mat4x3<f16> = let %4
- %6:ptr<uniform, vec4<f16>, read> = access %u, 1i
- %7:vec4<f16> = load %6
- %8:f16 = length %7
- %l:f16 = let %8
- %10:ptr<uniform, vec4<f16>, read> = access %u, 0i
- %11:vec4<f16> = load %10
- %12:vec4<f16> = swizzle %11, ywxz
- %13:f16 = access %12, 0u
- %14:f16 = abs %13
- %a:f16 = let %14
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x4_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x4_f32/to_builtin.wgsl.expected.ir.msl
index e13ff87..033ba27 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x4_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x4_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,32 +1,12 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ const constant float3x4* u;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat3x4<f32>, read> = var @binding_point(0, 0)
+kernel void f(const constant float3x4* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ float4x3 const t = transpose((*tint_module_vars.u));
+ float const l = length((*tint_module_vars.u)[1]);
+ float const a = abs((*tint_module_vars.u)[0].ywxz[0u]);
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:mat3x4<f32> = load %u
- %4:mat4x3<f32> = transpose %3
- %t:mat4x3<f32> = let %4
- %6:ptr<uniform, vec4<f32>, read> = access %u, 1i
- %7:vec4<f32> = load %6
- %8:f32 = length %7
- %l:f32 = let %8
- %10:ptr<uniform, vec4<f32>, read> = access %u, 0i
- %11:vec4<f32> = load %10
- %12:vec4<f32> = swizzle %11, ywxz
- %13:f32 = access %12, 0u
- %14:f32 = abs %13
- %a:f32 = let %14
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x2_f16/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x2_f16/to_builtin.wgsl.expected.ir.msl
index bdb7c20..e872e93 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x2_f16/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x2_f16/to_builtin.wgsl.expected.ir.msl
@@ -1,32 +1,12 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ const constant half4x2* u;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat4x2<f16>, read> = var @binding_point(0, 0)
+kernel void f(const constant half4x2* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ half2x4 const t = transpose((*tint_module_vars.u));
+ half const l = length((*tint_module_vars.u)[1]);
+ half const a = abs((*tint_module_vars.u)[0].yx[0u]);
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:mat4x2<f16> = load %u
- %4:mat2x4<f16> = transpose %3
- %t:mat2x4<f16> = let %4
- %6:ptr<uniform, vec2<f16>, read> = access %u, 1i
- %7:vec2<f16> = load %6
- %8:f16 = length %7
- %l:f16 = let %8
- %10:ptr<uniform, vec2<f16>, read> = access %u, 0i
- %11:vec2<f16> = load %10
- %12:vec2<f16> = swizzle %11, yx
- %13:f16 = access %12, 0u
- %14:f16 = abs %13
- %a:f16 = let %14
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x2_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x2_f32/to_builtin.wgsl.expected.ir.msl
index 585f2aa..4e2beec 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x2_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x2_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,32 +1,12 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ const constant float4x2* u;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat4x2<f32>, read> = var @binding_point(0, 0)
+kernel void f(const constant float4x2* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ float2x4 const t = transpose((*tint_module_vars.u));
+ float const l = length((*tint_module_vars.u)[1]);
+ float const a = abs((*tint_module_vars.u)[0].yx[0u]);
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:mat4x2<f32> = load %u
- %4:mat2x4<f32> = transpose %3
- %t:mat2x4<f32> = let %4
- %6:ptr<uniform, vec2<f32>, read> = access %u, 1i
- %7:vec2<f32> = load %6
- %8:f32 = length %7
- %l:f32 = let %8
- %10:ptr<uniform, vec2<f32>, read> = access %u, 0i
- %11:vec2<f32> = load %10
- %12:vec2<f32> = swizzle %11, yx
- %13:f32 = access %12, 0u
- %14:f32 = abs %13
- %a:f32 = let %14
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x3_f16/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x3_f16/to_builtin.wgsl.expected.ir.msl
index d600629..4520f30 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x3_f16/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x3_f16/to_builtin.wgsl.expected.ir.msl
@@ -1,32 +1,12 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ const constant half4x3* u;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat4x3<f16>, read> = var @binding_point(0, 0)
+kernel void f(const constant half4x3* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ half3x4 const t = transpose((*tint_module_vars.u));
+ half const l = length((*tint_module_vars.u)[1]);
+ half const a = abs((*tint_module_vars.u)[0].zxy[0u]);
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:mat4x3<f16> = load %u
- %4:mat3x4<f16> = transpose %3
- %t:mat3x4<f16> = let %4
- %6:ptr<uniform, vec3<f16>, read> = access %u, 1i
- %7:vec3<f16> = load %6
- %8:f16 = length %7
- %l:f16 = let %8
- %10:ptr<uniform, vec3<f16>, read> = access %u, 0i
- %11:vec3<f16> = load %10
- %12:vec3<f16> = swizzle %11, zxy
- %13:f16 = access %12, 0u
- %14:f16 = abs %13
- %a:f16 = let %14
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x3_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x3_f32/to_builtin.wgsl.expected.ir.msl
index 45acf1a..aa5a5b7 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x3_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x3_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,32 +1,12 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ const constant float4x3* u;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat4x3<f32>, read> = var @binding_point(0, 0)
+kernel void f(const constant float4x3* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ float3x4 const t = transpose((*tint_module_vars.u));
+ float const l = length((*tint_module_vars.u)[1]);
+ float const a = abs((*tint_module_vars.u)[0].zxy[0u]);
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:mat4x3<f32> = load %u
- %4:mat3x4<f32> = transpose %3
- %t:mat3x4<f32> = let %4
- %6:ptr<uniform, vec3<f32>, read> = access %u, 1i
- %7:vec3<f32> = load %6
- %8:f32 = length %7
- %l:f32 = let %8
- %10:ptr<uniform, vec3<f32>, read> = access %u, 0i
- %11:vec3<f32> = load %10
- %12:vec3<f32> = swizzle %11, zxy
- %13:f32 = access %12, 0u
- %14:f32 = abs %13
- %a:f32 = let %14
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x4_f16/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x4_f16/to_builtin.wgsl.expected.ir.msl
index 933a9d0..548a056 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x4_f16/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x4_f16/to_builtin.wgsl.expected.ir.msl
@@ -1,32 +1,12 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ const constant half4x4* u;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat4x4<f16>, read> = var @binding_point(0, 0)
+kernel void f(const constant half4x4* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ half4x4 const t = transpose((*tint_module_vars.u));
+ half const l = length((*tint_module_vars.u)[1]);
+ half const a = abs((*tint_module_vars.u)[0].ywxz[0u]);
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:mat4x4<f16> = load %u
- %4:mat4x4<f16> = transpose %3
- %t:mat4x4<f16> = let %4
- %6:ptr<uniform, vec4<f16>, read> = access %u, 1i
- %7:vec4<f16> = load %6
- %8:f16 = length %7
- %l:f16 = let %8
- %10:ptr<uniform, vec4<f16>, read> = access %u, 0i
- %11:vec4<f16> = load %10
- %12:vec4<f16> = swizzle %11, ywxz
- %13:f16 = access %12, 0u
- %14:f16 = abs %13
- %a:f16 = let %14
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x4_f32/to_builtin.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x4_f32/to_builtin.wgsl.expected.ir.msl
index 89826de..aa26e33 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x4_f32/to_builtin.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x4_f32/to_builtin.wgsl.expected.ir.msl
@@ -1,32 +1,12 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ const constant float4x4* u;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %u:ptr<uniform, mat4x4<f32>, read> = var @binding_point(0, 0)
+kernel void f(const constant float4x4* u [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u};
+ float4x4 const t = transpose((*tint_module_vars.u));
+ float const l = length((*tint_module_vars.u)[1]);
+ float const a = abs((*tint_module_vars.u)[0].ywxz[0u]);
}
-
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:mat4x4<f32> = load %u
- %4:mat4x4<f32> = transpose %3
- %t:mat4x4<f32> = let %4
- %6:ptr<uniform, vec4<f32>, read> = access %u, 1i
- %7:vec4<f32> = load %6
- %8:f32 = length %7
- %l:f32 = let %8
- %10:ptr<uniform, vec4<f32>, read> = access %u, 0i
- %11:vec4<f32> = load %10
- %12:vec4<f32> = swizzle %11, ywxz
- %13:f32 = access %12, 0u
- %14:f32 = abs %13
- %a:f32 = let %14
- ret
- }
-}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/length/056071.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/length/056071.wgsl.expected.ir.msl
index e95f9e4..2e5a291 100644
--- a/test/tint/builtins/gen/var/length/056071.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/length/056071.wgsl.expected.ir.msl
@@ -1,43 +1,37 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device float* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ float prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ float VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+float length_056071() {
+ float3 arg_0 = float3(0.0f);
+ float res = length(arg_0);
+ return res;
}
-
-%length_056071 = func():void {
- $B2: {
- %arg_0:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(0.0f)
- %4:vec3<f32> = load %arg_0
- %5:f32 = length %4
- %res:ptr<function, f32, read_write> = var, %5
- %7:f32 = load %res
- store %prevent_dce, %7
- ret
- }
+fragment void fragment_main(device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = length_056071();
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %length_056071
- ret vec4<f32>(0.0f)
- }
+kernel void compute_main(device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = length_056071();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %length_056071
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = length_056071();
+ return out;
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %length_056071
- ret
- }
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/length/3f0e13.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/length/3f0e13.wgsl.expected.ir.msl
index 35357a6..51b7fc3 100644
--- a/test/tint/builtins/gen/var/length/3f0e13.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/length/3f0e13.wgsl.expected.ir.msl
@@ -1,43 +1,37 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device half* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ half prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ half VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, f16, read_write> = var @binding_point(2, 0)
+half length_3f0e13() {
+ half2 arg_0 = half2(0.0h);
+ half res = length(arg_0);
+ return res;
}
-
-%length_3f0e13 = func():void {
- $B2: {
- %arg_0:ptr<function, vec2<f16>, read_write> = var, vec2<f16>(0.0h)
- %4:vec2<f16> = load %arg_0
- %5:f16 = length %4
- %res:ptr<function, f16, read_write> = var, %5
- %7:f16 = load %res
- store %prevent_dce, %7
- ret
- }
+fragment void fragment_main(device half* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = length_3f0e13();
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %length_3f0e13
- ret vec4<f32>(0.0f)
- }
+kernel void compute_main(device half* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = length_3f0e13();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %length_3f0e13
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = length_3f0e13();
+ return out;
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %length_3f0e13
- ret
- }
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/length/5b1a9b.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/length/5b1a9b.wgsl.expected.ir.msl
index 2c0d70d..c870877 100644
--- a/test/tint/builtins/gen/var/length/5b1a9b.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/length/5b1a9b.wgsl.expected.ir.msl
@@ -1,43 +1,37 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device half* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ half prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ half VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, f16, read_write> = var @binding_point(2, 0)
+half length_5b1a9b() {
+ half4 arg_0 = half4(0.0h);
+ half res = length(arg_0);
+ return res;
}
-
-%length_5b1a9b = func():void {
- $B2: {
- %arg_0:ptr<function, vec4<f16>, read_write> = var, vec4<f16>(0.0h)
- %4:vec4<f16> = load %arg_0
- %5:f16 = length %4
- %res:ptr<function, f16, read_write> = var, %5
- %7:f16 = load %res
- store %prevent_dce, %7
- ret
- }
+fragment void fragment_main(device half* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = length_5b1a9b();
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %length_5b1a9b
- ret vec4<f32>(0.0f)
- }
+kernel void compute_main(device half* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = length_5b1a9b();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %length_5b1a9b
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = length_5b1a9b();
+ return out;
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %length_5b1a9b
- ret
- }
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/length/602a17.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/length/602a17.wgsl.expected.ir.msl
index 7a2defb..44d3bdb 100644
--- a/test/tint/builtins/gen/var/length/602a17.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/length/602a17.wgsl.expected.ir.msl
@@ -1,43 +1,37 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device float* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ float prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ float VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+float length_602a17() {
+ float arg_0 = 0.0f;
+ float res = abs(arg_0);
+ return res;
}
-
-%length_602a17 = func():void {
- $B2: {
- %arg_0:ptr<function, f32, read_write> = var, 0.0f
- %4:f32 = load %arg_0
- %5:f32 = length %4
- %res:ptr<function, f32, read_write> = var, %5
- %7:f32 = load %res
- store %prevent_dce, %7
- ret
- }
+fragment void fragment_main(device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = length_602a17();
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %length_602a17
- ret vec4<f32>(0.0f)
- }
+kernel void compute_main(device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = length_602a17();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %length_602a17
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = length_602a17();
+ return out;
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %length_602a17
- ret
- }
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/length/afde8b.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/length/afde8b.wgsl.expected.ir.msl
index 85bcbc5d..76ad4ab 100644
--- a/test/tint/builtins/gen/var/length/afde8b.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/length/afde8b.wgsl.expected.ir.msl
@@ -1,43 +1,37 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device float* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ float prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ float VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+float length_afde8b() {
+ float2 arg_0 = float2(0.0f);
+ float res = length(arg_0);
+ return res;
}
-
-%length_afde8b = func():void {
- $B2: {
- %arg_0:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(0.0f)
- %4:vec2<f32> = load %arg_0
- %5:f32 = length %4
- %res:ptr<function, f32, read_write> = var, %5
- %7:f32 = load %res
- store %prevent_dce, %7
- ret
- }
+fragment void fragment_main(device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = length_afde8b();
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %length_afde8b
- ret vec4<f32>(0.0f)
- }
+kernel void compute_main(device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = length_afde8b();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %length_afde8b
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = length_afde8b();
+ return out;
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %length_afde8b
- ret
- }
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/length/ba16d6.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/length/ba16d6.wgsl.expected.ir.msl
index 3c41bbf..3796015 100644
--- a/test/tint/builtins/gen/var/length/ba16d6.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/length/ba16d6.wgsl.expected.ir.msl
@@ -1,43 +1,37 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device half* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ half prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ half VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, f16, read_write> = var @binding_point(2, 0)
+half length_ba16d6() {
+ half3 arg_0 = half3(0.0h);
+ half res = length(arg_0);
+ return res;
}
-
-%length_ba16d6 = func():void {
- $B2: {
- %arg_0:ptr<function, vec3<f16>, read_write> = var, vec3<f16>(0.0h)
- %4:vec3<f16> = load %arg_0
- %5:f16 = length %4
- %res:ptr<function, f16, read_write> = var, %5
- %7:f16 = load %res
- store %prevent_dce, %7
- ret
- }
+fragment void fragment_main(device half* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = length_ba16d6();
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %length_ba16d6
- ret vec4<f32>(0.0f)
- }
+kernel void compute_main(device half* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = length_ba16d6();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %length_ba16d6
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = length_ba16d6();
+ return out;
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %length_ba16d6
- ret
- }
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/length/becebf.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/length/becebf.wgsl.expected.ir.msl
index 642a661..c130236 100644
--- a/test/tint/builtins/gen/var/length/becebf.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/length/becebf.wgsl.expected.ir.msl
@@ -1,43 +1,37 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device float* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ float prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ float VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+float length_becebf() {
+ float4 arg_0 = float4(0.0f);
+ float res = length(arg_0);
+ return res;
}
-
-%length_becebf = func():void {
- $B2: {
- %arg_0:ptr<function, vec4<f32>, read_write> = var, vec4<f32>(0.0f)
- %4:vec4<f32> = load %arg_0
- %5:f32 = length %4
- %res:ptr<function, f32, read_write> = var, %5
- %7:f32 = load %res
- store %prevent_dce, %7
- ret
- }
+fragment void fragment_main(device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = length_becebf();
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %length_becebf
- ret vec4<f32>(0.0f)
- }
+kernel void compute_main(device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = length_becebf();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %length_becebf
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = length_becebf();
+ return out;
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %length_becebf
- ret
- }
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/length/c158da.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/length/c158da.wgsl.expected.ir.msl
index 016229a..0d63615 100644
--- a/test/tint/builtins/gen/var/length/c158da.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/length/c158da.wgsl.expected.ir.msl
@@ -1,43 +1,37 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device half* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ half prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ half VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, f16, read_write> = var @binding_point(2, 0)
+half length_c158da() {
+ half arg_0 = 0.0h;
+ half res = abs(arg_0);
+ return res;
}
-
-%length_c158da = func():void {
- $B2: {
- %arg_0:ptr<function, f16, read_write> = var, 0.0h
- %4:f16 = load %arg_0
- %5:f16 = length %4
- %res:ptr<function, f16, read_write> = var, %5
- %7:f16 = load %res
- store %prevent_dce, %7
- ret
- }
+fragment void fragment_main(device half* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = length_c158da();
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %length_c158da
- ret vec4<f32>(0.0f)
- }
+kernel void compute_main(device half* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = length_c158da();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %length_c158da
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = length_c158da();
+ return out;
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %length_c158da
- ret
- }
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
}
-
-unhandled variable address space
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
