[msl] Handle textureSample builtins
Replace them with texture.sample() member functions. Add support for
MemberFunctionCall to the printer.
Bug: 42251016
Change-Id: I24246253c9e3160ebe712d6e2e7f9d951e94a738
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/192305
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 c259530..46b324c 100644
--- a/src/tint/lang/msl/builtin_fn.cc
+++ b/src/tint/lang/msl/builtin_fn.cc
@@ -64,6 +64,8 @@
return "atomic_load_explicit";
case BuiltinFn::kAtomicStoreExplicit:
return "atomic_store_explicit";
+ case BuiltinFn::kSample:
+ return "sample";
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 dcdf48e..17703bf 100644
--- a/src/tint/lang/msl/builtin_fn.h
+++ b/src/tint/lang/msl/builtin_fn.h
@@ -58,6 +58,7 @@
kAtomicFetchXorExplicit,
kAtomicLoadExplicit,
kAtomicStoreExplicit,
+ kSample,
kThreadgroupBarrier,
kNone,
};
diff --git a/src/tint/lang/msl/intrinsic/data.cc b/src/tint/lang/msl/intrinsic/data.cc
index c632629..c3afbb3 100644
--- a/src/tint/lang/msl/intrinsic/data.cc
+++ b/src/tint/lang/msl/intrinsic/data.cc
@@ -116,6 +116,80 @@
};
+/// TypeMatcher for 'type f32'
+constexpr TypeMatcher kF32Matcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+ if (!MatchF32(state, ty)) {
+ return nullptr;
+ }
+ return BuildF32(state, ty);
+ },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
+ out << style::Type("f32");
+ }
+};
+
+
+/// TypeMatcher for 'type vec2'
+constexpr TypeMatcher kVec2Matcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+ const Type* T = nullptr;
+ if (!MatchVec2(state, ty, T)) {
+ return nullptr;
+ }
+ T = state.Type(T);
+ if (T == nullptr) {
+ return nullptr;
+ }
+ return BuildVec2(state, ty, T);
+ },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
+ state->PrintType(T);
+ out << style::Type("vec2", "<", T, ">");
+ }
+};
+
+
+/// TypeMatcher for 'type vec3'
+constexpr TypeMatcher kVec3Matcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+ const Type* T = nullptr;
+ if (!MatchVec3(state, ty, T)) {
+ return nullptr;
+ }
+ T = state.Type(T);
+ if (T == nullptr) {
+ return nullptr;
+ }
+ return BuildVec3(state, ty, T);
+ },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
+ state->PrintType(T);
+ out << style::Type("vec3", "<", T, ">");
+ }
+};
+
+
+/// TypeMatcher for 'type vec4'
+constexpr TypeMatcher kVec4Matcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+ const Type* T = nullptr;
+ if (!MatchVec4(state, ty, T)) {
+ return nullptr;
+ }
+ T = state.Type(T);
+ if (T == nullptr) {
+ return nullptr;
+ }
+ return BuildVec4(state, ty, T);
+ },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
+ state->PrintType(T);
+ out << style::Type("vec4", "<", T, ">");
+ }
+};
+
+
/// TypeMatcher for 'type atomic'
constexpr TypeMatcher kAtomicMatcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
@@ -168,6 +242,196 @@
};
+/// TypeMatcher for 'type sampler'
+constexpr TypeMatcher kSamplerMatcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+ if (!MatchSampler(state, ty)) {
+ return nullptr;
+ }
+ return BuildSampler(state, ty);
+ },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
+ out << style::Type("sampler");
+ }
+};
+
+
+/// TypeMatcher for 'type texture_1d'
+constexpr TypeMatcher kTexture1DMatcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+ const Type* T = nullptr;
+ if (!MatchTexture1D(state, ty, T)) {
+ return nullptr;
+ }
+ T = state.Type(T);
+ if (T == nullptr) {
+ return nullptr;
+ }
+ return BuildTexture1D(state, ty, T);
+ },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
+ state->PrintType(T);
+ out << style::Type("texture_1d", "<", T, ">");
+ }
+};
+
+
+/// TypeMatcher for 'type texture_2d'
+constexpr TypeMatcher kTexture2DMatcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+ const Type* T = nullptr;
+ if (!MatchTexture2D(state, ty, T)) {
+ return nullptr;
+ }
+ T = state.Type(T);
+ if (T == nullptr) {
+ return nullptr;
+ }
+ return BuildTexture2D(state, ty, T);
+ },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
+ state->PrintType(T);
+ out << style::Type("texture_2d", "<", T, ">");
+ }
+};
+
+
+/// TypeMatcher for 'type texture_2d_array'
+constexpr TypeMatcher kTexture2DArrayMatcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+ const Type* T = nullptr;
+ if (!MatchTexture2DArray(state, ty, T)) {
+ return nullptr;
+ }
+ T = state.Type(T);
+ if (T == nullptr) {
+ return nullptr;
+ }
+ return BuildTexture2DArray(state, ty, T);
+ },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
+ state->PrintType(T);
+ out << style::Type("texture_2d_array", "<", T, ">");
+ }
+};
+
+
+/// TypeMatcher for 'type texture_3d'
+constexpr TypeMatcher kTexture3DMatcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+ const Type* T = nullptr;
+ if (!MatchTexture3D(state, ty, T)) {
+ return nullptr;
+ }
+ T = state.Type(T);
+ if (T == nullptr) {
+ return nullptr;
+ }
+ return BuildTexture3D(state, ty, T);
+ },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
+ state->PrintType(T);
+ out << style::Type("texture_3d", "<", T, ">");
+ }
+};
+
+
+/// TypeMatcher for 'type texture_cube'
+constexpr TypeMatcher kTextureCubeMatcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+ const Type* T = nullptr;
+ if (!MatchTextureCube(state, ty, T)) {
+ return nullptr;
+ }
+ T = state.Type(T);
+ if (T == nullptr) {
+ return nullptr;
+ }
+ return BuildTextureCube(state, ty, T);
+ },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
+ state->PrintType(T);
+ out << style::Type("texture_cube", "<", T, ">");
+ }
+};
+
+
+/// TypeMatcher for 'type texture_cube_array'
+constexpr TypeMatcher kTextureCubeArrayMatcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+ const Type* T = nullptr;
+ if (!MatchTextureCubeArray(state, ty, T)) {
+ return nullptr;
+ }
+ T = state.Type(T);
+ if (T == nullptr) {
+ return nullptr;
+ }
+ return BuildTextureCubeArray(state, ty, T);
+ },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
+ state->PrintType(T);
+ out << style::Type("texture_cube_array", "<", T, ">");
+ }
+};
+
+
+/// TypeMatcher for 'type texture_depth_2d'
+constexpr TypeMatcher kTextureDepth2DMatcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+ if (!MatchTextureDepth2D(state, ty)) {
+ return nullptr;
+ }
+ return BuildTextureDepth2D(state, ty);
+ },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
+ out << style::Type("texture_depth_2d");
+ }
+};
+
+
+/// TypeMatcher for 'type texture_depth_2d_array'
+constexpr TypeMatcher kTextureDepth2DArrayMatcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+ if (!MatchTextureDepth2DArray(state, ty)) {
+ return nullptr;
+ }
+ return BuildTextureDepth2DArray(state, ty);
+ },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
+ out << style::Type("texture_depth_2d_array");
+ }
+};
+
+
+/// TypeMatcher for 'type texture_depth_cube'
+constexpr TypeMatcher kTextureDepthCubeMatcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+ if (!MatchTextureDepthCube(state, ty)) {
+ return nullptr;
+ }
+ return BuildTextureDepthCube(state, ty);
+ },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
+ out << style::Type("texture_depth_cube");
+ }
+};
+
+
+/// TypeMatcher for 'type texture_depth_cube_array'
+constexpr TypeMatcher kTextureDepthCubeArrayMatcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+ if (!MatchTextureDepthCubeArray(state, ty)) {
+ return nullptr;
+ }
+ return BuildTextureDepthCubeArray(state, ty);
+ },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
+ out << style::Type("texture_depth_cube_array");
+ }
+};
+
+
/// TypeMatcher for 'match iu32'
constexpr TypeMatcher kIu32Matcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
@@ -236,9 +500,24 @@
/* [2] */ kBoolMatcher,
/* [3] */ kI32Matcher,
/* [4] */ kU32Matcher,
- /* [5] */ kAtomicMatcher,
- /* [6] */ kPtrMatcher,
- /* [7] */ kIu32Matcher,
+ /* [5] */ kF32Matcher,
+ /* [6] */ kVec2Matcher,
+ /* [7] */ kVec3Matcher,
+ /* [8] */ kVec4Matcher,
+ /* [9] */ kAtomicMatcher,
+ /* [10] */ kPtrMatcher,
+ /* [11] */ kSamplerMatcher,
+ /* [12] */ kTexture1DMatcher,
+ /* [13] */ kTexture2DMatcher,
+ /* [14] */ kTexture2DArrayMatcher,
+ /* [15] */ kTexture3DMatcher,
+ /* [16] */ kTextureCubeMatcher,
+ /* [17] */ kTextureCubeArrayMatcher,
+ /* [18] */ kTextureDepth2DMatcher,
+ /* [19] */ kTextureDepth2DArrayMatcher,
+ /* [20] */ kTextureDepthCubeMatcher,
+ /* [21] */ kTextureDepthCubeArrayMatcher,
+ /* [22] */ kIu32Matcher,
};
/// The template numbers, and number matchers
@@ -251,17 +530,44 @@
};
constexpr MatcherIndex kMatcherIndices[] = {
- /* [0] */ MatcherIndex(6),
+ /* [0] */ MatcherIndex(10),
/* [1] */ MatcherIndex(1),
- /* [2] */ MatcherIndex(5),
+ /* [2] */ MatcherIndex(9),
/* [3] */ MatcherIndex(0),
/* [4] */ MatcherIndex(2),
- /* [5] */ MatcherIndex(6),
+ /* [5] */ MatcherIndex(10),
/* [6] */ MatcherIndex(3),
/* [7] */ MatcherIndex(0),
/* [8] */ MatcherIndex(2),
- /* [9] */ MatcherIndex(7),
- /* [10] */ MatcherIndex(4),
+ /* [9] */ MatcherIndex(8),
+ /* [10] */ MatcherIndex(5),
+ /* [11] */ MatcherIndex(12),
+ /* [12] */ MatcherIndex(5),
+ /* [13] */ MatcherIndex(13),
+ /* [14] */ MatcherIndex(5),
+ /* [15] */ MatcherIndex(6),
+ /* [16] */ MatcherIndex(5),
+ /* [17] */ MatcherIndex(6),
+ /* [18] */ MatcherIndex(3),
+ /* [19] */ MatcherIndex(14),
+ /* [20] */ MatcherIndex(5),
+ /* [21] */ MatcherIndex(15),
+ /* [22] */ MatcherIndex(5),
+ /* [23] */ MatcherIndex(7),
+ /* [24] */ MatcherIndex(5),
+ /* [25] */ MatcherIndex(7),
+ /* [26] */ MatcherIndex(3),
+ /* [27] */ MatcherIndex(16),
+ /* [28] */ MatcherIndex(5),
+ /* [29] */ MatcherIndex(17),
+ /* [30] */ MatcherIndex(5),
+ /* [31] */ MatcherIndex(22),
+ /* [32] */ MatcherIndex(4),
+ /* [33] */ MatcherIndex(11),
+ /* [34] */ MatcherIndex(18),
+ /* [35] */ MatcherIndex(19),
+ /* [36] */ MatcherIndex(20),
+ /* [37] */ MatcherIndex(21),
};
static_assert(MatcherIndicesIndex::CanIndex(kMatcherIndices),
@@ -286,37 +592,232 @@
{
/* [3] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(10),
+ /* matcher_indices */ MatcherIndicesIndex(32),
},
{
/* [4] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(10),
+ /* matcher_indices */ MatcherIndicesIndex(32),
},
{
/* [5] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(19),
+ },
+ {
+ /* [6] */
+ /* usage */ core::ParameterUsage::kSampler,
+ /* matcher_indices */ MatcherIndicesIndex(33),
+ },
+ {
+ /* [7] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(15),
+ },
+ {
+ /* [8] */
+ /* usage */ core::ParameterUsage::kArrayIndex,
+ /* matcher_indices */ MatcherIndicesIndex(3),
+ },
+ {
+ /* [9] */
+ /* usage */ core::ParameterUsage::kOffset,
+ /* matcher_indices */ MatcherIndicesIndex(17),
+ },
+ {
+ /* [10] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(35),
+ },
+ {
+ /* [11] */
+ /* usage */ core::ParameterUsage::kSampler,
+ /* matcher_indices */ MatcherIndicesIndex(33),
+ },
+ {
+ /* [12] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(15),
+ },
+ {
+ /* [13] */
+ /* usage */ core::ParameterUsage::kArrayIndex,
+ /* matcher_indices */ MatcherIndicesIndex(3),
+ },
+ {
+ /* [14] */
+ /* usage */ core::ParameterUsage::kOffset,
+ /* matcher_indices */ MatcherIndicesIndex(17),
+ },
+ {
+ /* [15] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(13),
+ },
+ {
+ /* [16] */
+ /* usage */ core::ParameterUsage::kSampler,
+ /* matcher_indices */ MatcherIndicesIndex(33),
+ },
+ {
+ /* [17] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(15),
+ },
+ {
+ /* [18] */
+ /* usage */ core::ParameterUsage::kOffset,
+ /* matcher_indices */ MatcherIndicesIndex(17),
+ },
+ {
+ /* [19] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(21),
+ },
+ {
+ /* [20] */
+ /* usage */ core::ParameterUsage::kSampler,
+ /* matcher_indices */ MatcherIndicesIndex(33),
+ },
+ {
+ /* [21] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(23),
+ },
+ {
+ /* [22] */
+ /* usage */ core::ParameterUsage::kOffset,
+ /* matcher_indices */ MatcherIndicesIndex(25),
+ },
+ {
+ /* [23] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(29),
+ },
+ {
+ /* [24] */
+ /* usage */ core::ParameterUsage::kSampler,
+ /* matcher_indices */ MatcherIndicesIndex(33),
+ },
+ {
+ /* [25] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(23),
+ },
+ {
+ /* [26] */
+ /* usage */ core::ParameterUsage::kArrayIndex,
+ /* matcher_indices */ MatcherIndicesIndex(3),
+ },
+ {
+ /* [27] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(34),
+ },
+ {
+ /* [28] */
+ /* usage */ core::ParameterUsage::kSampler,
+ /* matcher_indices */ MatcherIndicesIndex(33),
+ },
+ {
+ /* [29] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(15),
+ },
+ {
+ /* [30] */
+ /* usage */ core::ParameterUsage::kOffset,
+ /* matcher_indices */ MatcherIndicesIndex(17),
+ },
+ {
+ /* [31] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(37),
+ },
+ {
+ /* [32] */
+ /* usage */ core::ParameterUsage::kSampler,
+ /* matcher_indices */ MatcherIndicesIndex(33),
+ },
+ {
+ /* [33] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(23),
+ },
+ {
+ /* [34] */
+ /* usage */ core::ParameterUsage::kArrayIndex,
+ /* matcher_indices */ MatcherIndicesIndex(3),
+ },
+ {
+ /* [35] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(0),
},
{
- /* [6] */
+ /* [36] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(3),
},
{
- /* [7] */
+ /* [37] */
/* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(32),
+ },
+ {
+ /* [38] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(11),
+ },
+ {
+ /* [39] */
+ /* usage */ core::ParameterUsage::kSampler,
+ /* matcher_indices */ MatcherIndicesIndex(33),
+ },
+ {
+ /* [40] */
+ /* usage */ core::ParameterUsage::kCoords,
/* matcher_indices */ MatcherIndicesIndex(10),
},
{
- /* [8] */
+ /* [41] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(27),
+ },
+ {
+ /* [42] */
+ /* usage */ core::ParameterUsage::kSampler,
+ /* matcher_indices */ MatcherIndicesIndex(33),
+ },
+ {
+ /* [43] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(23),
+ },
+ {
+ /* [44] */
+ /* usage */ core::ParameterUsage::kTexture,
+ /* matcher_indices */ MatcherIndicesIndex(36),
+ },
+ {
+ /* [45] */
+ /* usage */ core::ParameterUsage::kSampler,
+ /* matcher_indices */ MatcherIndicesIndex(33),
+ },
+ {
+ /* [46] */
+ /* usage */ core::ParameterUsage::kCoords,
+ /* matcher_indices */ MatcherIndicesIndex(23),
+ },
+ {
+ /* [47] */
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(0),
},
{
- /* [9] */
+ /* [48] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(10),
+ /* matcher_indices */ MatcherIndicesIndex(32),
},
};
@@ -327,15 +828,21 @@
{
/* [0] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(9),
+ /* matcher_indices */ MatcherIndicesIndex(31),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [1] */
/* name */ "S",
- /* matcher_indices */ MatcherIndicesIndex(10),
+ /* matcher_indices */ MatcherIndicesIndex(32),
/* kind */ TemplateInfo::Kind::kNumber,
},
+ {
+ /* [2] */
+ /* name */ "A",
+ /* matcher_indices */ MatcherIndicesIndex(31),
+ /* kind */ TemplateInfo::Kind::kType,
+ },
};
static_assert(TemplateIndex::CanIndex(kTemplates),
@@ -344,6 +851,171 @@
constexpr OverloadInfo kOverloads[] = {
{
/* [0] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+ /* num_parameters */ 3,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 0,
+ /* templates */ TemplateIndex(/* invalid */),
+ /* parameters */ ParameterIndex(38),
+ /* return_matcher_indices */ MatcherIndicesIndex(9),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [1] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+ /* num_parameters */ 3,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 0,
+ /* templates */ TemplateIndex(/* invalid */),
+ /* parameters */ ParameterIndex(15),
+ /* return_matcher_indices */ MatcherIndicesIndex(9),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [2] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+ /* num_parameters */ 4,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 0,
+ /* templates */ TemplateIndex(/* invalid */),
+ /* parameters */ ParameterIndex(15),
+ /* return_matcher_indices */ MatcherIndicesIndex(9),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [3] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+ /* num_parameters */ 4,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 1,
+ /* templates */ TemplateIndex(2),
+ /* parameters */ ParameterIndex(5),
+ /* return_matcher_indices */ MatcherIndicesIndex(9),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [4] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+ /* num_parameters */ 5,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 1,
+ /* templates */ TemplateIndex(2),
+ /* parameters */ ParameterIndex(5),
+ /* return_matcher_indices */ MatcherIndicesIndex(9),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [5] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+ /* num_parameters */ 3,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 0,
+ /* templates */ TemplateIndex(/* invalid */),
+ /* parameters */ ParameterIndex(19),
+ /* return_matcher_indices */ MatcherIndicesIndex(9),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [6] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+ /* num_parameters */ 4,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 0,
+ /* templates */ TemplateIndex(/* invalid */),
+ /* parameters */ ParameterIndex(19),
+ /* return_matcher_indices */ MatcherIndicesIndex(9),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [7] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+ /* num_parameters */ 3,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 0,
+ /* templates */ TemplateIndex(/* invalid */),
+ /* parameters */ ParameterIndex(41),
+ /* return_matcher_indices */ MatcherIndicesIndex(9),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [8] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+ /* num_parameters */ 4,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 1,
+ /* templates */ TemplateIndex(2),
+ /* parameters */ ParameterIndex(23),
+ /* return_matcher_indices */ MatcherIndicesIndex(9),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [9] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+ /* num_parameters */ 3,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 0,
+ /* templates */ TemplateIndex(/* invalid */),
+ /* parameters */ ParameterIndex(27),
+ /* return_matcher_indices */ MatcherIndicesIndex(10),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [10] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+ /* num_parameters */ 4,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 0,
+ /* templates */ TemplateIndex(/* invalid */),
+ /* parameters */ ParameterIndex(27),
+ /* return_matcher_indices */ MatcherIndicesIndex(10),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [11] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+ /* num_parameters */ 4,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 1,
+ /* templates */ TemplateIndex(2),
+ /* parameters */ ParameterIndex(10),
+ /* return_matcher_indices */ MatcherIndicesIndex(10),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [12] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+ /* num_parameters */ 5,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 1,
+ /* templates */ TemplateIndex(2),
+ /* parameters */ ParameterIndex(10),
+ /* return_matcher_indices */ MatcherIndicesIndex(10),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [13] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+ /* num_parameters */ 3,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 0,
+ /* templates */ TemplateIndex(/* invalid */),
+ /* parameters */ ParameterIndex(44),
+ /* return_matcher_indices */ MatcherIndicesIndex(10),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [14] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+ /* num_parameters */ 4,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 1,
+ /* templates */ TemplateIndex(2),
+ /* parameters */ ParameterIndex(31),
+ /* return_matcher_indices */ MatcherIndicesIndex(10),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [15] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 5,
/* num_explicit_templates */ 0,
@@ -354,40 +1026,40 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [1] */
+ /* [16] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(0),
- /* parameters */ ParameterIndex(5),
+ /* parameters */ ParameterIndex(35),
/* return_matcher_indices */ MatcherIndicesIndex(3),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [2] */
+ /* [17] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(0),
- /* parameters */ ParameterIndex(8),
+ /* parameters */ ParameterIndex(47),
/* return_matcher_indices */ MatcherIndicesIndex(3),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [3] */
+ /* [18] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(0),
- /* parameters */ ParameterIndex(5),
+ /* parameters */ ParameterIndex(35),
/* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [4] */
+ /* [19] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -407,73 +1079,93 @@
/* [0] */
/* fn atomic_compare_exchange_weak_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, ptr<function, T, read_write>, T, u32, u32) -> bool */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(0),
+ /* overloads */ OverloadIndex(15),
},
{
/* [1] */
/* fn atomic_exchange_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(1),
+ /* overloads */ OverloadIndex(16),
},
{
/* [2] */
/* fn atomic_fetch_add_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(1),
+ /* overloads */ OverloadIndex(16),
},
{
/* [3] */
/* fn atomic_fetch_and_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(1),
+ /* overloads */ OverloadIndex(16),
},
{
/* [4] */
/* fn atomic_fetch_max_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(1),
+ /* overloads */ OverloadIndex(16),
},
{
/* [5] */
/* fn atomic_fetch_min_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(1),
+ /* overloads */ OverloadIndex(16),
},
{
/* [6] */
/* fn atomic_fetch_or_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(1),
+ /* overloads */ OverloadIndex(16),
},
{
/* [7] */
/* fn atomic_fetch_sub_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(1),
+ /* overloads */ OverloadIndex(16),
},
{
/* [8] */
/* fn atomic_fetch_xor_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(1),
+ /* overloads */ OverloadIndex(16),
},
{
/* [9] */
/* fn atomic_load_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(2),
+ /* overloads */ OverloadIndex(17),
},
{
/* [10] */
/* fn atomic_store_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(3),
+ /* overloads */ OverloadIndex(18),
},
{
/* [11] */
+ /* fn sample(texture: texture_1d<f32>, sampler: sampler, coords: f32) -> vec4<f32> */
+ /* fn sample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>) -> vec4<f32> */
+ /* fn sample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> vec4<f32> */
+ /* fn sample[A : iu32](texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: A) -> vec4<f32> */
+ /* fn sample[A : iu32](texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: A, @const offset: vec2<i32>) -> vec4<f32> */
+ /* fn sample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>) -> vec4<f32> */
+ /* fn sample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, @const offset: vec3<i32>) -> vec4<f32> */
+ /* fn sample(texture: texture_cube<f32>, sampler: sampler, coords: vec3<f32>) -> vec4<f32> */
+ /* fn sample[A : iu32](texture: texture_cube_array<f32>, sampler: sampler, coords: vec3<f32>, array_index: A) -> vec4<f32> */
+ /* fn sample(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>) -> f32 */
+ /* fn sample(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> f32 */
+ /* fn sample[A : iu32](texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: A) -> f32 */
+ /* fn sample[A : iu32](texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: A, @const offset: vec2<i32>) -> f32 */
+ /* fn sample(texture: texture_depth_cube, sampler: sampler, coords: vec3<f32>) -> f32 */
+ /* fn sample[A : iu32](texture: texture_depth_cube_array, sampler: sampler, coords: vec3<f32>, array_index: A) -> f32 */
+ /* num overloads */ 15,
+ /* overloads */ OverloadIndex(0),
+ },
+ {
+ /* [12] */
/* fn threadgroup_barrier(u32) */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(4),
+ /* overloads */ OverloadIndex(19),
},
};
diff --git a/src/tint/lang/msl/msl.def b/src/tint/lang/msl/msl.def
index 858ee55..235f697 100644
--- a/src/tint/lang/msl/msl.def
+++ b/src/tint/lang/msl/msl.def
@@ -56,8 +56,23 @@
type bool
type i32
type u32
+type f32
+type vec2<T>
+type vec3<T>
+type vec4<T>
type atomic<T>
type ptr<S: address_space, T, A: access>
+type sampler
+type texture_1d<T>
+type texture_2d<T>
+type texture_2d_array<T>
+type texture_3d<T>
+type texture_cube<T>
+type texture_cube_array<T>
+type texture_depth_2d
+type texture_depth_2d_array
+type texture_depth_cube
+type texture_depth_cube_array
////////////////////////////////////////////////////////////////////////////////
// Type matchers //
@@ -79,5 +94,22 @@
@stage("fragment", "compute") fn atomic_fetch_xor_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T
@stage("fragment", "compute") fn atomic_load_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, u32) -> T
@stage("fragment", "compute") fn atomic_store_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32)
+
+@member_function @stage("fragment") fn sample(texture: texture_1d<f32>, sampler: sampler, coords: f32) -> vec4<f32>
+@member_function @stage("fragment") fn sample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>) -> vec4<f32>
+@member_function @stage("fragment") fn sample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> vec4<f32>
+@member_function @stage("fragment") fn sample[A: iu32](texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: A) -> vec4<f32>
+@member_function @stage("fragment") fn sample[A: iu32](texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: A, @const offset: vec2<i32>) -> vec4<f32>
+@member_function @stage("fragment") fn sample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>) -> vec4<f32>
+@member_function @stage("fragment") fn sample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, @const offset: vec3<i32>) -> vec4<f32>
+@member_function @stage("fragment") fn sample(texture: texture_cube<f32>, sampler: sampler, coords: vec3<f32>) -> vec4<f32>
+@member_function @stage("fragment") fn sample[A: iu32](texture: texture_cube_array<f32>, sampler: sampler, coords: vec3<f32>, array_index: A) -> vec4<f32>
+@member_function @stage("fragment") fn sample(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>) -> f32
+@member_function @stage("fragment") fn sample(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> f32
+@member_function @stage("fragment") fn sample[A: iu32](texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: A) -> f32
+@member_function @stage("fragment") fn sample[A: iu32](texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: A, @const offset: vec2<i32>) -> f32
+@member_function @stage("fragment") fn sample(texture: texture_depth_cube, sampler: sampler, coords: vec3<f32>) -> f32
+@member_function @stage("fragment") fn sample[A: iu32](texture: texture_depth_cube_array, sampler: sampler, coords: vec3<f32>, array_index: A) -> f32
+
@stage("compute") fn threadgroup_barrier(u32)
diff --git a/src/tint/lang/msl/writer/printer/printer.cc b/src/tint/lang/msl/writer/printer/printer.cc
index 5d74063..2617d6d 100644
--- a/src/tint/lang/msl/writer/printer/printer.cc
+++ b/src/tint/lang/msl/writer/printer/printer.cc
@@ -89,6 +89,7 @@
#include "src/tint/lang/msl/barrier_type.h"
#include "src/tint/lang/msl/builtin_fn.h"
#include "src/tint/lang/msl/ir/builtin_call.h"
+#include "src/tint/lang/msl/ir/member_builtin_call.h"
#include "src/tint/lang/msl/ir/memory_order.h"
#include "src/tint/lang/msl/writer/common/printer_support.h"
#include "src/tint/utils/containers/map.h"
@@ -408,17 +409,20 @@
[&](const core::ir::Constant* c) { EmitConstant(out, c); }, //
[&](const core::ir::InstructionResult* r) {
Switch(
- r->Instruction(), //
- [&](const core::ir::CoreBinary* b) { EmitBinary(out, b); }, //
- [&](const core::ir::CoreUnary* u) { EmitUnary(out, u); }, //
- [&](const core::ir::Convert* b) { EmitConvert(out, b); }, //
- [&](const core::ir::Let* l) { out << NameOf(l->Result(0)); }, //
- [&](const core::ir::Load* l) { EmitLoad(out, l); }, //
- [&](const core::ir::Construct* c) { EmitConstruct(out, c); }, //
- [&](const core::ir::Var* var) { out << NameOf(var->Result(0)); }, //
- [&](const core::ir::Bitcast* b) { EmitBitcast(out, b); }, //
- [&](const core::ir::Access* a) { EmitAccess(out, a); }, //
- [&](const msl::ir::BuiltinCall* c) { EmitMslBuiltinCall(out, c); }, //
+ r->Instruction(), //
+ [&](const core::ir::CoreBinary* b) { EmitBinary(out, b); }, //
+ [&](const core::ir::CoreUnary* u) { EmitUnary(out, u); }, //
+ [&](const core::ir::Convert* b) { EmitConvert(out, b); }, //
+ [&](const core::ir::Let* l) { out << NameOf(l->Result(0)); }, //
+ [&](const core::ir::Load* l) { EmitLoad(out, l); }, //
+ [&](const core::ir::Construct* c) { EmitConstruct(out, c); }, //
+ [&](const core::ir::Var* var) { out << NameOf(var->Result(0)); }, //
+ [&](const core::ir::Bitcast* b) { EmitBitcast(out, b); }, //
+ [&](const core::ir::Access* a) { EmitAccess(out, a); }, //
+ [&](const msl::ir::BuiltinCall* c) { EmitMslBuiltinCall(out, c); }, //
+ [&](const msl::ir::MemberBuiltinCall* c) {
+ EmitMslMemberBuiltinCall(out, c);
+ }, //
[&](const core::ir::CoreBuiltinCall* c) { EmitCoreBuiltinCall(out, c); }, //
[&](const core::ir::UserCall* c) { EmitUserCall(out, c); }, //
[&](const core::ir::LoadVectorElement* e) {
@@ -854,6 +858,20 @@
out << ")";
}
+ void EmitMslMemberBuiltinCall(StringStream& out, const msl::ir::MemberBuiltinCall* c) {
+ EmitValue(out, c->Object());
+ out << "." << c->Func() << "(";
+ bool needs_comma = false;
+ for (const auto* arg : c->Args()) {
+ if (needs_comma) {
+ out << ", ";
+ }
+ EmitAndTakeAddressIfNeeded(out, arg);
+ needs_comma = true;
+ }
+ out << ")";
+ }
+
void EmitCoreBuiltinCall(StringStream& out, const core::ir::CoreBuiltinCall* c) {
EmitCoreBuiltinName(out, c->Func());
out << "(";
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
index f5c36f4..b5bbaf1 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/msl/barrier_type.h"
#include "src/tint/lang/msl/builtin_fn.h"
#include "src/tint/lang/msl/ir/builtin_call.h"
+#include "src/tint/lang/msl/ir/member_builtin_call.h"
#include "src/tint/lang/msl/ir/memory_order.h"
#include "src/tint/utils/containers/hashmap.h"
@@ -79,6 +80,7 @@
case core::BuiltinFn::kAtomicStore:
case core::BuiltinFn::kAtomicSub:
case core::BuiltinFn::kAtomicXor:
+ case core::BuiltinFn::kTextureSample:
case core::BuiltinFn::kStorageBarrier:
case core::BuiltinFn::kWorkgroupBarrier:
case core::BuiltinFn::kTextureBarrier:
@@ -93,6 +95,7 @@
// Replace the builtins that we found.
for (auto* builtin : worklist) {
switch (builtin->Func()) {
+ // Atomics.
case core::BuiltinFn::kAtomicAdd:
AtomicCall(builtin, msl::BuiltinFn::kAtomicFetchAddExplicit);
break;
@@ -126,6 +129,13 @@
case core::BuiltinFn::kAtomicXor:
AtomicCall(builtin, msl::BuiltinFn::kAtomicFetchXorExplicit);
break;
+
+ // Texture builtins.
+ case core::BuiltinFn::kTextureSample:
+ TextureSample(builtin);
+ break;
+
+ // Barriers.
case core::BuiltinFn::kStorageBarrier:
ThreadgroupBarrier(builtin, BarrierType::kDevice);
break;
@@ -135,6 +145,7 @@
case core::BuiltinFn::kTextureBarrier:
ThreadgroupBarrier(builtin, BarrierType::kTexture);
break;
+
default:
break;
}
@@ -194,6 +205,17 @@
builtin->Destroy();
}
+ /// Replace a textureSample call with the equivalent MSL intrinsic.
+ /// @param builtin the builtin call instruction
+ void TextureSample(core::ir::CoreBuiltinCall* builtin) {
+ // The MSL intrinsic is a member function, so we split the first argument off as the object.
+ auto args = Vector<core::ir::Value*, 4>(builtin->Args().Offset(1));
+ auto* call = b.MemberCallWithResult<msl::ir::MemberBuiltinCall>(
+ builtin->DetachResult(), msl::BuiltinFn::kSample, builtin->Args()[0], std::move(args));
+ call->InsertBefore(builtin);
+ builtin->Destroy();
+ }
+
/// Replace a barrier builtin with the `threadgroupBarrier()` intrinsic.
/// @param builtin the builtin call instruction
/// @param type the barrier type
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 527c7c7..904e1fa 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
@@ -35,6 +35,7 @@
#include "src/tint/lang/core/number.h"
#include "src/tint/lang/core/type/atomic.h"
#include "src/tint/lang/core/type/builtin_structs.h"
+#include "src/tint/lang/core/type/sampled_texture.h"
using namespace tint::core::fluent_types; // NOLINT
using namespace tint::core::number_suffixes; // NOLINT
@@ -751,6 +752,42 @@
EXPECT_EQ(expect, str());
}
+TEST_F(MslWriter_BuiltinPolyfillTest, TextureSample) {
+ auto* t = b.FunctionParam(
+ "t", ty.Get<core::type::SampledTexture>(core::type::TextureDimension::k2d, ty.f32()));
+ auto* s = b.FunctionParam("s", ty.sampler());
+ auto* coords = b.FunctionParam("coords", ty.vec2<f32>());
+ auto* func = b.Function("foo", ty.vec4<f32>());
+ func->SetParams({t, s, coords});
+ b.Append(func->Block(), [&] {
+ auto* result = b.Call<vec4<f32>>(core::BuiltinFn::kTextureSample, t, s, coords);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%t:texture_2d<f32>, %s:sampler, %coords:vec2<f32>):vec4<f32> {
+ $B1: {
+ %5:vec4<f32> = textureSample %t, %s, %coords
+ ret %5
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%t:texture_2d<f32>, %s:sampler, %coords:vec2<f32>):vec4<f32> {
+ $B1: {
+ %5:vec4<f32> = %t.sample %s, %coords
+ ret %5
+ }
+}
+)";
+
+ Run(BuiltinPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
TEST_F(MslWriter_BuiltinPolyfillTest, WorkgroupBarrier) {
auto* func = b.Function("foo", ty.void_());
func->SetStage(core::ir::Function::PipelineStage::kCompute);
diff --git a/test/tint/bug/tint/1183.wgsl.expected.ir.msl b/test/tint/bug/tint/1183.wgsl.expected.ir.msl
index 3543919..c6e4754 100644
--- a/test/tint/bug/tint/1183.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1183.wgsl.expected.ir.msl
@@ -1,9 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture2d<float, access::sample> t;
+ sampler s;
+};
+struct f_outputs {
+ float4 tint_symbol [[color(0)]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-* 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. *
-********************************************************************
+float4 f_inner(tint_module_vars_struct tint_module_vars) {
+ return tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f), int2(4, 6));
+}
+fragment f_outputs f(texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s};
+ return f_outputs{.tint_symbol=f_inner(tint_module_vars)};
+}
diff --git a/test/tint/bug/tint/978.wgsl.expected.ir.msl b/test/tint/bug/tint/978.wgsl.expected.ir.msl
index 414ed77..e9fbfff 100644
--- a/test/tint/bug/tint/978.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/978.wgsl.expected.ir.msl
@@ -1,40 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct FragmentOutput {
+ float4 color;
+};
+struct FragmentInput {
+ float2 vUv;
+};
+struct tint_module_vars_struct {
+ depth2d<float, access::sample> depthMap;
+ sampler texSampler;
+};
+struct tint_symbol_outputs {
+ float4 FragmentOutput_color [[color(0)]];
+};
+struct tint_symbol_inputs {
+ float2 FragmentInput_vUv [[user(locn2)]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:1257 internal compiler error: FragmentOutput = struct @align(16) {
- color:vec4<f32> @offset(0), @location(0)
+FragmentOutput tint_symbol_inner(FragmentInput fIn, tint_module_vars_struct tint_module_vars) {
+ float const sample = tint_module_vars.depthMap.sample(tint_module_vars.texSampler, fIn.vUv);
+ float3 const color = float3(sample, sample, sample);
+ FragmentOutput fOut = {};
+ fOut.color = float4(color, 1.0f);
+ return fOut;
}
-
-FragmentInput = struct @align(8) {
- vUv:vec2<f32> @offset(0), @location(2)
+fragment tint_symbol_outputs tint_symbol(tint_symbol_inputs inputs [[stage_in]], depth2d<float, access::sample> depthMap [[texture(0)]], sampler texSampler [[sampler(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.depthMap=depthMap, .texSampler=texSampler};
+ return tint_symbol_outputs{.FragmentOutput_color=tint_symbol_inner(FragmentInput{.vUv=inputs.FragmentInput_vUv}, tint_module_vars).color};
}
-
-$B1: { # root
- %depthMap:ptr<handle, texture_depth_2d, read> = var @binding_point(1, 5)
- %texSampler:ptr<handle, sampler, read> = var @binding_point(1, 3)
-}
-
-%tint_symbol = @fragment func(%fIn:FragmentInput):FragmentOutput {
- $B2: {
- %5:texture_depth_2d = load %depthMap
- %6:sampler = load %texSampler
- %7:vec2<f32> = access %fIn, 0u
- %8:f32 = textureSample %5, %6, %7
- %sample:f32 = let %8
- %10:vec3<f32> = construct %sample, %sample, %sample
- %color:vec3<f32> = let %10
- %fOut:ptr<function, FragmentOutput, read_write> = var
- %13:ptr<function, vec4<f32>, read_write> = access %fOut, 0u
- %14:vec4<f32> = construct %color, 1.0f
- store %13, %14
- %15:FragmentOutput = load %fOut
- ret %15
- }
-}
-
-invalid entry point IO struct uses
-********************************************************************
-* 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/literal/textureSample/0dff6c.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/0dff6c.wgsl.expected.ir.msl
index 55819ad..adce6d3 100644
--- a/test/tint/builtins/gen/literal/textureSample/0dff6c.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/0dff6c.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ depth2d<float, access::sample> arg_0;
+ sampler arg_1;
+ device float* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_depth_2d, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_0dff6c(tint_module_vars_struct tint_module_vars) {
+ float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float2(1.0f), int2(1));
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_0dff6c = func():void {
- $B2: {
- %5:texture_depth_2d = load %arg_0
- %6:sampler = load %arg_1
- %7:f32 = textureSample %5, %6, vec2<f32>(1.0f), vec2<i32>(1i)
- %res:ptr<function, f32, read_write> = var, %7
- %9:f32 = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(depth2d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_0dff6c(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %11:void = call %textureSample_0dff6c
- 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/literal/textureSample/17e988.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/17e988.wgsl.expected.ir.msl
index 1da17ad..754c9df 100644
--- a/test/tint/builtins/gen/literal/textureSample/17e988.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/17e988.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture2d_array<float, access::sample> arg_0;
+ sampler arg_1;
+ device float4* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_2d_array<f32>, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_17e988(tint_module_vars_struct tint_module_vars) {
+ float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float2(1.0f), 1, int2(1));
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_17e988 = func():void {
- $B2: {
- %5:texture_2d_array<f32> = load %arg_0
- %6:sampler = load %arg_1
- %7:vec4<f32> = textureSample %5, %6, vec2<f32>(1.0f), 1i, vec2<i32>(1i)
- %res:ptr<function, vec4<f32>, read_write> = var, %7
- %9:vec4<f32> = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(texture2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_17e988(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %11:void = call %textureSample_17e988
- 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/literal/textureSample/193203.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/193203.wgsl.expected.ir.msl
index ae6b1e6..7441785 100644
--- a/test/tint/builtins/gen/literal/textureSample/193203.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/193203.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture2d_array<float, access::sample> arg_0;
+ sampler arg_1;
+ device float4* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_2d_array<f32>, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_193203(tint_module_vars_struct tint_module_vars) {
+ float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float2(1.0f), 1u, int2(1));
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_193203 = func():void {
- $B2: {
- %5:texture_2d_array<f32> = load %arg_0
- %6:sampler = load %arg_1
- %7:vec4<f32> = textureSample %5, %6, vec2<f32>(1.0f), 1u, vec2<i32>(1i)
- %res:ptr<function, vec4<f32>, read_write> = var, %7
- %9:vec4<f32> = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(texture2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_193203(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %11:void = call %textureSample_193203
- 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/literal/textureSample/1a4e1b.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/1a4e1b.wgsl.expected.ir.msl
index a0e7fe3..4d76dd5 100644
--- a/test/tint/builtins/gen/literal/textureSample/1a4e1b.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/1a4e1b.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ depth2d_array<float, access::sample> arg_0;
+ sampler arg_1;
+ device float* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_depth_2d_array, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_1a4e1b(tint_module_vars_struct tint_module_vars) {
+ float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float2(1.0f), 1u);
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_1a4e1b = func():void {
- $B2: {
- %5:texture_depth_2d_array = load %arg_0
- %6:sampler = load %arg_1
- %7:f32 = textureSample %5, %6, vec2<f32>(1.0f), 1u
- %res:ptr<function, f32, read_write> = var, %7
- %9:f32 = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(depth2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_1a4e1b(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %11:void = call %textureSample_1a4e1b
- 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/literal/textureSample/2149ec.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/2149ec.wgsl.expected.ir.msl
index 8b191d3..fa08bbe 100644
--- a/test/tint/builtins/gen/literal/textureSample/2149ec.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/2149ec.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture3d<float, access::sample> arg_0;
+ sampler arg_1;
+ device float4* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_3d<f32>, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_2149ec(tint_module_vars_struct tint_module_vars) {
+ float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float3(1.0f), int3(1));
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_2149ec = func():void {
- $B2: {
- %5:texture_3d<f32> = load %arg_0
- %6:sampler = load %arg_1
- %7:vec4<f32> = textureSample %5, %6, vec3<f32>(1.0f), vec3<i32>(1i)
- %res:ptr<function, vec4<f32>, read_write> = var, %7
- %9:vec4<f32> = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(texture3d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_2149ec(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %11:void = call %textureSample_2149ec
- 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/literal/textureSample/38bbb9.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/38bbb9.wgsl.expected.ir.msl
index a4a4204..f5743b7 100644
--- a/test/tint/builtins/gen/literal/textureSample/38bbb9.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/38bbb9.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ depth2d<float, access::sample> arg_0;
+ sampler arg_1;
+ device float* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_depth_2d, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_38bbb9(tint_module_vars_struct tint_module_vars) {
+ float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float2(1.0f));
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_38bbb9 = func():void {
- $B2: {
- %5:texture_depth_2d = load %arg_0
- %6:sampler = load %arg_1
- %7:f32 = textureSample %5, %6, vec2<f32>(1.0f)
- %res:ptr<function, f32, read_write> = var, %7
- %9:f32 = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(depth2d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_38bbb9(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %11:void = call %textureSample_38bbb9
- 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/literal/textureSample/3b50bd.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/3b50bd.wgsl.expected.ir.msl
index 3f682f9..bd92d77 100644
--- a/test/tint/builtins/gen/literal/textureSample/3b50bd.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/3b50bd.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture3d<float, access::sample> arg_0;
+ sampler arg_1;
+ device float4* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_3d<f32>, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_3b50bd(tint_module_vars_struct tint_module_vars) {
+ float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float3(1.0f));
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_3b50bd = func():void {
- $B2: {
- %5:texture_3d<f32> = load %arg_0
- %6:sampler = load %arg_1
- %7:vec4<f32> = textureSample %5, %6, vec3<f32>(1.0f)
- %res:ptr<function, vec4<f32>, read_write> = var, %7
- %9:vec4<f32> = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(texture3d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_3b50bd(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %11:void = call %textureSample_3b50bd
- 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/literal/textureSample/4703d0.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/4703d0.wgsl.expected.ir.msl
index 979bf90..83b2790 100644
--- a/test/tint/builtins/gen/literal/textureSample/4703d0.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/4703d0.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ depth2d_array<float, access::sample> arg_0;
+ sampler arg_1;
+ device float* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_depth_2d_array, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_4703d0(tint_module_vars_struct tint_module_vars) {
+ float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float2(1.0f), 1u, int2(1));
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_4703d0 = func():void {
- $B2: {
- %5:texture_depth_2d_array = load %arg_0
- %6:sampler = load %arg_1
- %7:f32 = textureSample %5, %6, vec2<f32>(1.0f), 1u, vec2<i32>(1i)
- %res:ptr<function, f32, read_write> = var, %7
- %9:f32 = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(depth2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_4703d0(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %11:void = call %textureSample_4703d0
- 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/literal/textureSample/4dd1bf.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/4dd1bf.wgsl.expected.ir.msl
index 9ebb5c2..b7252ad 100644
--- a/test/tint/builtins/gen/literal/textureSample/4dd1bf.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/4dd1bf.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texturecube_array<float, access::sample> arg_0;
+ sampler arg_1;
+ device float4* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_cube_array<f32>, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_4dd1bf(tint_module_vars_struct tint_module_vars) {
+ float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float3(1.0f), 1);
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_4dd1bf = func():void {
- $B2: {
- %5:texture_cube_array<f32> = load %arg_0
- %6:sampler = load %arg_1
- %7:vec4<f32> = textureSample %5, %6, vec3<f32>(1.0f), 1i
- %res:ptr<function, vec4<f32>, read_write> = var, %7
- %9:vec4<f32> = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(texturecube_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_4dd1bf(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %11:void = call %textureSample_4dd1bf
- 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/literal/textureSample/51b514.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/51b514.wgsl.expected.ir.msl
index e019e2b..850b230 100644
--- a/test/tint/builtins/gen/literal/textureSample/51b514.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/51b514.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture2d<float, access::sample> arg_0;
+ sampler arg_1;
+ device float4* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_2d<f32>, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_51b514(tint_module_vars_struct tint_module_vars) {
+ float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float2(1.0f));
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_51b514 = func():void {
- $B2: {
- %5:texture_2d<f32> = load %arg_0
- %6:sampler = load %arg_1
- %7:vec4<f32> = textureSample %5, %6, vec2<f32>(1.0f)
- %res:ptr<function, vec4<f32>, read_write> = var, %7
- %9:vec4<f32> = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(texture2d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_51b514(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %11:void = call %textureSample_51b514
- 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/literal/textureSample/60bf45.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/60bf45.wgsl.expected.ir.msl
index 03697ed..90a9283 100644
--- a/test/tint/builtins/gen/literal/textureSample/60bf45.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/60bf45.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ depth2d_array<float, access::sample> arg_0;
+ sampler arg_1;
+ device float* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_depth_2d_array, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_60bf45(tint_module_vars_struct tint_module_vars) {
+ float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float2(1.0f), 1, int2(1));
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_60bf45 = func():void {
- $B2: {
- %5:texture_depth_2d_array = load %arg_0
- %6:sampler = load %arg_1
- %7:f32 = textureSample %5, %6, vec2<f32>(1.0f), 1i, vec2<i32>(1i)
- %res:ptr<function, f32, read_write> = var, %7
- %9:f32 = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(depth2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_60bf45(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %11:void = call %textureSample_60bf45
- 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/literal/textureSample/6717ca.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/6717ca.wgsl.expected.ir.msl
index c5d638e..34f57aa 100644
--- a/test/tint/builtins/gen/literal/textureSample/6717ca.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/6717ca.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture2d_array<float, access::sample> arg_0;
+ sampler arg_1;
+ device float4* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_2d_array<f32>, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_6717ca(tint_module_vars_struct tint_module_vars) {
+ float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float2(1.0f), 1);
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_6717ca = func():void {
- $B2: {
- %5:texture_2d_array<f32> = load %arg_0
- %6:sampler = load %arg_1
- %7:vec4<f32> = textureSample %5, %6, vec2<f32>(1.0f), 1i
- %res:ptr<function, vec4<f32>, read_write> = var, %7
- %9:vec4<f32> = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(texture2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_6717ca(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %11:void = call %textureSample_6717ca
- 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/literal/textureSample/6e64fb.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/6e64fb.wgsl.expected.ir.msl
index a3c0665..971c6fb 100644
--- a/test/tint/builtins/gen/literal/textureSample/6e64fb.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/6e64fb.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture1d<float, access::sample> arg_0;
+ sampler arg_1;
+ device float4* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_1d<f32>, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_6e64fb(tint_module_vars_struct tint_module_vars) {
+ float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, 1.0f);
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_6e64fb = func():void {
- $B2: {
- %5:texture_1d<f32> = load %arg_0
- %6:sampler = load %arg_1
- %7:vec4<f32> = textureSample %5, %6, 1.0f
- %res:ptr<function, vec4<f32>, read_write> = var, %7
- %9:vec4<f32> = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(texture1d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_6e64fb(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %11:void = call %textureSample_6e64fb
- 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/literal/textureSample/7e9ffd.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/7e9ffd.wgsl.expected.ir.msl
index 203f61c..ccd1290 100644
--- a/test/tint/builtins/gen/literal/textureSample/7e9ffd.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/7e9ffd.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ depth2d_array<float, access::sample> arg_0;
+ sampler arg_1;
+ device float* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_depth_2d_array, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_7e9ffd(tint_module_vars_struct tint_module_vars) {
+ float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float2(1.0f), 1);
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_7e9ffd = func():void {
- $B2: {
- %5:texture_depth_2d_array = load %arg_0
- %6:sampler = load %arg_1
- %7:f32 = textureSample %5, %6, vec2<f32>(1.0f), 1i
- %res:ptr<function, f32, read_write> = var, %7
- %9:f32 = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(depth2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_7e9ffd(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %11:void = call %textureSample_7e9ffd
- 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/literal/textureSample/7fd8cb.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/7fd8cb.wgsl.expected.ir.msl
index f6d8f84..b24c4de 100644
--- a/test/tint/builtins/gen/literal/textureSample/7fd8cb.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/7fd8cb.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ depthcube_array<float, access::sample> arg_0;
+ sampler arg_1;
+ device float* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_depth_cube_array, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_7fd8cb(tint_module_vars_struct tint_module_vars) {
+ float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float3(1.0f), 1u);
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_7fd8cb = func():void {
- $B2: {
- %5:texture_depth_cube_array = load %arg_0
- %6:sampler = load %arg_1
- %7:f32 = textureSample %5, %6, vec3<f32>(1.0f), 1u
- %res:ptr<function, f32, read_write> = var, %7
- %9:f32 = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(depthcube_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_7fd8cb(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %11:void = call %textureSample_7fd8cb
- 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/literal/textureSample/85c4ba.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/85c4ba.wgsl.expected.ir.msl
index df38079..38b8474 100644
--- a/test/tint/builtins/gen/literal/textureSample/85c4ba.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/85c4ba.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture2d<float, access::sample> arg_0;
+ sampler arg_1;
+ device float4* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_2d<f32>, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_85c4ba(tint_module_vars_struct tint_module_vars) {
+ float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float2(1.0f), int2(1));
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_85c4ba = func():void {
- $B2: {
- %5:texture_2d<f32> = load %arg_0
- %6:sampler = load %arg_1
- %7:vec4<f32> = textureSample %5, %6, vec2<f32>(1.0f), vec2<i32>(1i)
- %res:ptr<function, vec4<f32>, read_write> = var, %7
- %9:vec4<f32> = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(texture2d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_85c4ba(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %11:void = call %textureSample_85c4ba
- 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/literal/textureSample/bc7477.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/bc7477.wgsl.expected.ir.msl
index 63be6d6..dfeb0c3 100644
--- a/test/tint/builtins/gen/literal/textureSample/bc7477.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/bc7477.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texturecube_array<float, access::sample> arg_0;
+ sampler arg_1;
+ device float4* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_cube_array<f32>, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_bc7477(tint_module_vars_struct tint_module_vars) {
+ float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float3(1.0f), 1u);
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_bc7477 = func():void {
- $B2: {
- %5:texture_cube_array<f32> = load %arg_0
- %6:sampler = load %arg_1
- %7:vec4<f32> = textureSample %5, %6, vec3<f32>(1.0f), 1u
- %res:ptr<function, vec4<f32>, read_write> = var, %7
- %9:vec4<f32> = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(texturecube_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_bc7477(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %11:void = call %textureSample_bc7477
- 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/literal/textureSample/c2f4e8.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/c2f4e8.wgsl.expected.ir.msl
index cc9e0ac..5d3648c 100644
--- a/test/tint/builtins/gen/literal/textureSample/c2f4e8.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/c2f4e8.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ depthcube_array<float, access::sample> arg_0;
+ sampler arg_1;
+ device float* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_depth_cube_array, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_c2f4e8(tint_module_vars_struct tint_module_vars) {
+ float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float3(1.0f), 1);
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_c2f4e8 = func():void {
- $B2: {
- %5:texture_depth_cube_array = load %arg_0
- %6:sampler = load %arg_1
- %7:f32 = textureSample %5, %6, vec3<f32>(1.0f), 1i
- %res:ptr<function, f32, read_write> = var, %7
- %9:f32 = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(depthcube_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_c2f4e8(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %11:void = call %textureSample_c2f4e8
- 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/literal/textureSample/d6b281.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/d6b281.wgsl.expected.ir.msl
index d663331..c729426 100644
--- a/test/tint/builtins/gen/literal/textureSample/d6b281.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/d6b281.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture2d_array<float, access::sample> arg_0;
+ sampler arg_1;
+ device float4* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_2d_array<f32>, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_d6b281(tint_module_vars_struct tint_module_vars) {
+ float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float2(1.0f), 1u);
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_d6b281 = func():void {
- $B2: {
- %5:texture_2d_array<f32> = load %arg_0
- %6:sampler = load %arg_1
- %7:vec4<f32> = textureSample %5, %6, vec2<f32>(1.0f), 1u
- %res:ptr<function, vec4<f32>, read_write> = var, %7
- %9:vec4<f32> = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(texture2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_d6b281(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %11:void = call %textureSample_d6b281
- 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/literal/textureSample/e53267.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/e53267.wgsl.expected.ir.msl
index 3babe54..6be1416 100644
--- a/test/tint/builtins/gen/literal/textureSample/e53267.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/e53267.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texturecube<float, access::sample> arg_0;
+ sampler arg_1;
+ device float4* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_cube<f32>, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_e53267(tint_module_vars_struct tint_module_vars) {
+ float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float3(1.0f));
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_e53267 = func():void {
- $B2: {
- %5:texture_cube<f32> = load %arg_0
- %6:sampler = load %arg_1
- %7:vec4<f32> = textureSample %5, %6, vec3<f32>(1.0f)
- %res:ptr<function, vec4<f32>, read_write> = var, %7
- %9:vec4<f32> = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(texturecube<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_e53267(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %11:void = call %textureSample_e53267
- 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/literal/textureSample/ea7030.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/ea7030.wgsl.expected.ir.msl
index 62669b6..db29db5 100644
--- a/test/tint/builtins/gen/literal/textureSample/ea7030.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/ea7030.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ depthcube<float, access::sample> arg_0;
+ sampler arg_1;
+ device float* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_depth_cube, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_ea7030(tint_module_vars_struct tint_module_vars) {
+ float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float3(1.0f));
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_ea7030 = func():void {
- $B2: {
- %5:texture_depth_cube = load %arg_0
- %6:sampler = load %arg_1
- %7:f32 = textureSample %5, %6, vec3<f32>(1.0f)
- %res:ptr<function, f32, read_write> = var, %7
- %9:f32 = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(depthcube<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_ea7030(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %11:void = call %textureSample_ea7030
- 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/textureSample/0dff6c.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/0dff6c.wgsl.expected.ir.msl
index 3c40846..c730055 100644
--- a/test/tint/builtins/gen/var/textureSample/0dff6c.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/0dff6c.wgsl.expected.ir.msl
@@ -1,35 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ depth2d<float, access::sample> arg_0;
+ sampler arg_1;
+ device float* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_depth_2d, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_0dff6c(tint_module_vars_struct tint_module_vars) {
+ float2 arg_2 = float2(1.0f);
+ float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, int2(1));
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_0dff6c = func():void {
- $B2: {
- %arg_2:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
- %6:texture_depth_2d = load %arg_0
- %7:sampler = load %arg_1
- %8:vec2<f32> = load %arg_2
- %9:f32 = textureSample %6, %7, %8, vec2<i32>(1i)
- %res:ptr<function, f32, read_write> = var, %9
- %11:f32 = load %res
- store %prevent_dce, %11
- ret
- }
+fragment void fragment_main(depth2d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_0dff6c(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %13:void = call %textureSample_0dff6c
- 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/textureSample/17e988.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/17e988.wgsl.expected.ir.msl
index 5f62935..759875c9 100644
--- a/test/tint/builtins/gen/var/textureSample/17e988.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/17e988.wgsl.expected.ir.msl
@@ -1,37 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture2d_array<float, access::sample> arg_0;
+ sampler arg_1;
+ device float4* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_2d_array<f32>, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_17e988(tint_module_vars_struct tint_module_vars) {
+ float2 arg_2 = float2(1.0f);
+ int arg_3 = 1;
+ float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, arg_3, int2(1));
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_17e988 = func():void {
- $B2: {
- %arg_2:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
- %arg_3:ptr<function, i32, read_write> = var, 1i
- %7:texture_2d_array<f32> = load %arg_0
- %8:sampler = load %arg_1
- %9:vec2<f32> = load %arg_2
- %10:i32 = load %arg_3
- %11:vec4<f32> = textureSample %7, %8, %9, %10, vec2<i32>(1i)
- %res:ptr<function, vec4<f32>, read_write> = var, %11
- %13:vec4<f32> = load %res
- store %prevent_dce, %13
- ret
- }
+fragment void fragment_main(texture2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_17e988(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %15:void = call %textureSample_17e988
- 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/textureSample/193203.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/193203.wgsl.expected.ir.msl
index b0393af..08d4cf4 100644
--- a/test/tint/builtins/gen/var/textureSample/193203.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/193203.wgsl.expected.ir.msl
@@ -1,37 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture2d_array<float, access::sample> arg_0;
+ sampler arg_1;
+ device float4* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_2d_array<f32>, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_193203(tint_module_vars_struct tint_module_vars) {
+ float2 arg_2 = float2(1.0f);
+ uint arg_3 = 1u;
+ float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, arg_3, int2(1));
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_193203 = func():void {
- $B2: {
- %arg_2:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
- %arg_3:ptr<function, u32, read_write> = var, 1u
- %7:texture_2d_array<f32> = load %arg_0
- %8:sampler = load %arg_1
- %9:vec2<f32> = load %arg_2
- %10:u32 = load %arg_3
- %11:vec4<f32> = textureSample %7, %8, %9, %10, vec2<i32>(1i)
- %res:ptr<function, vec4<f32>, read_write> = var, %11
- %13:vec4<f32> = load %res
- store %prevent_dce, %13
- ret
- }
+fragment void fragment_main(texture2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_193203(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %15:void = call %textureSample_193203
- 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/textureSample/1a4e1b.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/1a4e1b.wgsl.expected.ir.msl
index ddca6b6..0a84e65 100644
--- a/test/tint/builtins/gen/var/textureSample/1a4e1b.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/1a4e1b.wgsl.expected.ir.msl
@@ -1,37 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ depth2d_array<float, access::sample> arg_0;
+ sampler arg_1;
+ device float* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_depth_2d_array, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_1a4e1b(tint_module_vars_struct tint_module_vars) {
+ float2 arg_2 = float2(1.0f);
+ uint arg_3 = 1u;
+ float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, arg_3);
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_1a4e1b = func():void {
- $B2: {
- %arg_2:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
- %arg_3:ptr<function, u32, read_write> = var, 1u
- %7:texture_depth_2d_array = load %arg_0
- %8:sampler = load %arg_1
- %9:vec2<f32> = load %arg_2
- %10:u32 = load %arg_3
- %11:f32 = textureSample %7, %8, %9, %10
- %res:ptr<function, f32, read_write> = var, %11
- %13:f32 = load %res
- store %prevent_dce, %13
- ret
- }
+fragment void fragment_main(depth2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_1a4e1b(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %15:void = call %textureSample_1a4e1b
- 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/textureSample/2149ec.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/2149ec.wgsl.expected.ir.msl
index 4ad8af4..f25e8dd 100644
--- a/test/tint/builtins/gen/var/textureSample/2149ec.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/2149ec.wgsl.expected.ir.msl
@@ -1,35 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture3d<float, access::sample> arg_0;
+ sampler arg_1;
+ device float4* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_3d<f32>, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_2149ec(tint_module_vars_struct tint_module_vars) {
+ float3 arg_2 = float3(1.0f);
+ float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, int3(1));
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_2149ec = func():void {
- $B2: {
- %arg_2:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(1.0f)
- %6:texture_3d<f32> = load %arg_0
- %7:sampler = load %arg_1
- %8:vec3<f32> = load %arg_2
- %9:vec4<f32> = textureSample %6, %7, %8, vec3<i32>(1i)
- %res:ptr<function, vec4<f32>, read_write> = var, %9
- %11:vec4<f32> = load %res
- store %prevent_dce, %11
- ret
- }
+fragment void fragment_main(texture3d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_2149ec(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %13:void = call %textureSample_2149ec
- 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/textureSample/38bbb9.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/38bbb9.wgsl.expected.ir.msl
index a113cd4..dc1f22e 100644
--- a/test/tint/builtins/gen/var/textureSample/38bbb9.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/38bbb9.wgsl.expected.ir.msl
@@ -1,35 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ depth2d<float, access::sample> arg_0;
+ sampler arg_1;
+ device float* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_depth_2d, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_38bbb9(tint_module_vars_struct tint_module_vars) {
+ float2 arg_2 = float2(1.0f);
+ float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2);
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_38bbb9 = func():void {
- $B2: {
- %arg_2:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
- %6:texture_depth_2d = load %arg_0
- %7:sampler = load %arg_1
- %8:vec2<f32> = load %arg_2
- %9:f32 = textureSample %6, %7, %8
- %res:ptr<function, f32, read_write> = var, %9
- %11:f32 = load %res
- store %prevent_dce, %11
- ret
- }
+fragment void fragment_main(depth2d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_38bbb9(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %13:void = call %textureSample_38bbb9
- 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/textureSample/3b50bd.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/3b50bd.wgsl.expected.ir.msl
index 6e29909..7d4dccf 100644
--- a/test/tint/builtins/gen/var/textureSample/3b50bd.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/3b50bd.wgsl.expected.ir.msl
@@ -1,35 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture3d<float, access::sample> arg_0;
+ sampler arg_1;
+ device float4* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_3d<f32>, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_3b50bd(tint_module_vars_struct tint_module_vars) {
+ float3 arg_2 = float3(1.0f);
+ float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2);
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_3b50bd = func():void {
- $B2: {
- %arg_2:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(1.0f)
- %6:texture_3d<f32> = load %arg_0
- %7:sampler = load %arg_1
- %8:vec3<f32> = load %arg_2
- %9:vec4<f32> = textureSample %6, %7, %8
- %res:ptr<function, vec4<f32>, read_write> = var, %9
- %11:vec4<f32> = load %res
- store %prevent_dce, %11
- ret
- }
+fragment void fragment_main(texture3d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_3b50bd(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %13:void = call %textureSample_3b50bd
- 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/textureSample/4703d0.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/4703d0.wgsl.expected.ir.msl
index 17b17fa..8713856 100644
--- a/test/tint/builtins/gen/var/textureSample/4703d0.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/4703d0.wgsl.expected.ir.msl
@@ -1,37 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ depth2d_array<float, access::sample> arg_0;
+ sampler arg_1;
+ device float* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_depth_2d_array, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_4703d0(tint_module_vars_struct tint_module_vars) {
+ float2 arg_2 = float2(1.0f);
+ uint arg_3 = 1u;
+ float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, arg_3, int2(1));
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_4703d0 = func():void {
- $B2: {
- %arg_2:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
- %arg_3:ptr<function, u32, read_write> = var, 1u
- %7:texture_depth_2d_array = load %arg_0
- %8:sampler = load %arg_1
- %9:vec2<f32> = load %arg_2
- %10:u32 = load %arg_3
- %11:f32 = textureSample %7, %8, %9, %10, vec2<i32>(1i)
- %res:ptr<function, f32, read_write> = var, %11
- %13:f32 = load %res
- store %prevent_dce, %13
- ret
- }
+fragment void fragment_main(depth2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_4703d0(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %15:void = call %textureSample_4703d0
- 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/textureSample/4dd1bf.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/4dd1bf.wgsl.expected.ir.msl
index b021b59..78bf83a 100644
--- a/test/tint/builtins/gen/var/textureSample/4dd1bf.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/4dd1bf.wgsl.expected.ir.msl
@@ -1,37 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texturecube_array<float, access::sample> arg_0;
+ sampler arg_1;
+ device float4* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_cube_array<f32>, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_4dd1bf(tint_module_vars_struct tint_module_vars) {
+ float3 arg_2 = float3(1.0f);
+ int arg_3 = 1;
+ float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, arg_3);
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_4dd1bf = func():void {
- $B2: {
- %arg_2:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(1.0f)
- %arg_3:ptr<function, i32, read_write> = var, 1i
- %7:texture_cube_array<f32> = load %arg_0
- %8:sampler = load %arg_1
- %9:vec3<f32> = load %arg_2
- %10:i32 = load %arg_3
- %11:vec4<f32> = textureSample %7, %8, %9, %10
- %res:ptr<function, vec4<f32>, read_write> = var, %11
- %13:vec4<f32> = load %res
- store %prevent_dce, %13
- ret
- }
+fragment void fragment_main(texturecube_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_4dd1bf(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %15:void = call %textureSample_4dd1bf
- 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/textureSample/51b514.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/51b514.wgsl.expected.ir.msl
index aa66ed2..f3bd051 100644
--- a/test/tint/builtins/gen/var/textureSample/51b514.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/51b514.wgsl.expected.ir.msl
@@ -1,35 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture2d<float, access::sample> arg_0;
+ sampler arg_1;
+ device float4* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_2d<f32>, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_51b514(tint_module_vars_struct tint_module_vars) {
+ float2 arg_2 = float2(1.0f);
+ float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2);
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_51b514 = func():void {
- $B2: {
- %arg_2:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
- %6:texture_2d<f32> = load %arg_0
- %7:sampler = load %arg_1
- %8:vec2<f32> = load %arg_2
- %9:vec4<f32> = textureSample %6, %7, %8
- %res:ptr<function, vec4<f32>, read_write> = var, %9
- %11:vec4<f32> = load %res
- store %prevent_dce, %11
- ret
- }
+fragment void fragment_main(texture2d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_51b514(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %13:void = call %textureSample_51b514
- 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/textureSample/60bf45.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/60bf45.wgsl.expected.ir.msl
index 31b95cf..55c40ab 100644
--- a/test/tint/builtins/gen/var/textureSample/60bf45.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/60bf45.wgsl.expected.ir.msl
@@ -1,37 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ depth2d_array<float, access::sample> arg_0;
+ sampler arg_1;
+ device float* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_depth_2d_array, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_60bf45(tint_module_vars_struct tint_module_vars) {
+ float2 arg_2 = float2(1.0f);
+ int arg_3 = 1;
+ float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, arg_3, int2(1));
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_60bf45 = func():void {
- $B2: {
- %arg_2:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
- %arg_3:ptr<function, i32, read_write> = var, 1i
- %7:texture_depth_2d_array = load %arg_0
- %8:sampler = load %arg_1
- %9:vec2<f32> = load %arg_2
- %10:i32 = load %arg_3
- %11:f32 = textureSample %7, %8, %9, %10, vec2<i32>(1i)
- %res:ptr<function, f32, read_write> = var, %11
- %13:f32 = load %res
- store %prevent_dce, %13
- ret
- }
+fragment void fragment_main(depth2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_60bf45(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %15:void = call %textureSample_60bf45
- 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/textureSample/6717ca.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/6717ca.wgsl.expected.ir.msl
index 3adc780..199640d 100644
--- a/test/tint/builtins/gen/var/textureSample/6717ca.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/6717ca.wgsl.expected.ir.msl
@@ -1,37 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture2d_array<float, access::sample> arg_0;
+ sampler arg_1;
+ device float4* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_2d_array<f32>, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_6717ca(tint_module_vars_struct tint_module_vars) {
+ float2 arg_2 = float2(1.0f);
+ int arg_3 = 1;
+ float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, arg_3);
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_6717ca = func():void {
- $B2: {
- %arg_2:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
- %arg_3:ptr<function, i32, read_write> = var, 1i
- %7:texture_2d_array<f32> = load %arg_0
- %8:sampler = load %arg_1
- %9:vec2<f32> = load %arg_2
- %10:i32 = load %arg_3
- %11:vec4<f32> = textureSample %7, %8, %9, %10
- %res:ptr<function, vec4<f32>, read_write> = var, %11
- %13:vec4<f32> = load %res
- store %prevent_dce, %13
- ret
- }
+fragment void fragment_main(texture2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_6717ca(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %15:void = call %textureSample_6717ca
- 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/textureSample/6e64fb.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/6e64fb.wgsl.expected.ir.msl
index 2d7a629..2de338d 100644
--- a/test/tint/builtins/gen/var/textureSample/6e64fb.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/6e64fb.wgsl.expected.ir.msl
@@ -1,35 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture1d<float, access::sample> arg_0;
+ sampler arg_1;
+ device float4* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_1d<f32>, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_6e64fb(tint_module_vars_struct tint_module_vars) {
+ float arg_2 = 1.0f;
+ float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2);
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_6e64fb = func():void {
- $B2: {
- %arg_2:ptr<function, f32, read_write> = var, 1.0f
- %6:texture_1d<f32> = load %arg_0
- %7:sampler = load %arg_1
- %8:f32 = load %arg_2
- %9:vec4<f32> = textureSample %6, %7, %8
- %res:ptr<function, vec4<f32>, read_write> = var, %9
- %11:vec4<f32> = load %res
- store %prevent_dce, %11
- ret
- }
+fragment void fragment_main(texture1d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_6e64fb(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %13:void = call %textureSample_6e64fb
- 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/textureSample/7e9ffd.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/7e9ffd.wgsl.expected.ir.msl
index 2ff2807..d26d806 100644
--- a/test/tint/builtins/gen/var/textureSample/7e9ffd.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/7e9ffd.wgsl.expected.ir.msl
@@ -1,37 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ depth2d_array<float, access::sample> arg_0;
+ sampler arg_1;
+ device float* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_depth_2d_array, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_7e9ffd(tint_module_vars_struct tint_module_vars) {
+ float2 arg_2 = float2(1.0f);
+ int arg_3 = 1;
+ float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, arg_3);
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_7e9ffd = func():void {
- $B2: {
- %arg_2:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
- %arg_3:ptr<function, i32, read_write> = var, 1i
- %7:texture_depth_2d_array = load %arg_0
- %8:sampler = load %arg_1
- %9:vec2<f32> = load %arg_2
- %10:i32 = load %arg_3
- %11:f32 = textureSample %7, %8, %9, %10
- %res:ptr<function, f32, read_write> = var, %11
- %13:f32 = load %res
- store %prevent_dce, %13
- ret
- }
+fragment void fragment_main(depth2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_7e9ffd(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %15:void = call %textureSample_7e9ffd
- 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/textureSample/7fd8cb.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/7fd8cb.wgsl.expected.ir.msl
index cba8a48..d2fde33 100644
--- a/test/tint/builtins/gen/var/textureSample/7fd8cb.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/7fd8cb.wgsl.expected.ir.msl
@@ -1,37 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ depthcube_array<float, access::sample> arg_0;
+ sampler arg_1;
+ device float* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_depth_cube_array, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_7fd8cb(tint_module_vars_struct tint_module_vars) {
+ float3 arg_2 = float3(1.0f);
+ uint arg_3 = 1u;
+ float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, arg_3);
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_7fd8cb = func():void {
- $B2: {
- %arg_2:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(1.0f)
- %arg_3:ptr<function, u32, read_write> = var, 1u
- %7:texture_depth_cube_array = load %arg_0
- %8:sampler = load %arg_1
- %9:vec3<f32> = load %arg_2
- %10:u32 = load %arg_3
- %11:f32 = textureSample %7, %8, %9, %10
- %res:ptr<function, f32, read_write> = var, %11
- %13:f32 = load %res
- store %prevent_dce, %13
- ret
- }
+fragment void fragment_main(depthcube_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_7fd8cb(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %15:void = call %textureSample_7fd8cb
- 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/textureSample/85c4ba.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/85c4ba.wgsl.expected.ir.msl
index 4515efb..8eed4cb 100644
--- a/test/tint/builtins/gen/var/textureSample/85c4ba.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/85c4ba.wgsl.expected.ir.msl
@@ -1,35 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture2d<float, access::sample> arg_0;
+ sampler arg_1;
+ device float4* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_2d<f32>, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_85c4ba(tint_module_vars_struct tint_module_vars) {
+ float2 arg_2 = float2(1.0f);
+ float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, int2(1));
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_85c4ba = func():void {
- $B2: {
- %arg_2:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
- %6:texture_2d<f32> = load %arg_0
- %7:sampler = load %arg_1
- %8:vec2<f32> = load %arg_2
- %9:vec4<f32> = textureSample %6, %7, %8, vec2<i32>(1i)
- %res:ptr<function, vec4<f32>, read_write> = var, %9
- %11:vec4<f32> = load %res
- store %prevent_dce, %11
- ret
- }
+fragment void fragment_main(texture2d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_85c4ba(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %13:void = call %textureSample_85c4ba
- 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/textureSample/bc7477.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/bc7477.wgsl.expected.ir.msl
index 6f235ed..1344e67 100644
--- a/test/tint/builtins/gen/var/textureSample/bc7477.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/bc7477.wgsl.expected.ir.msl
@@ -1,37 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texturecube_array<float, access::sample> arg_0;
+ sampler arg_1;
+ device float4* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_cube_array<f32>, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_bc7477(tint_module_vars_struct tint_module_vars) {
+ float3 arg_2 = float3(1.0f);
+ uint arg_3 = 1u;
+ float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, arg_3);
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_bc7477 = func():void {
- $B2: {
- %arg_2:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(1.0f)
- %arg_3:ptr<function, u32, read_write> = var, 1u
- %7:texture_cube_array<f32> = load %arg_0
- %8:sampler = load %arg_1
- %9:vec3<f32> = load %arg_2
- %10:u32 = load %arg_3
- %11:vec4<f32> = textureSample %7, %8, %9, %10
- %res:ptr<function, vec4<f32>, read_write> = var, %11
- %13:vec4<f32> = load %res
- store %prevent_dce, %13
- ret
- }
+fragment void fragment_main(texturecube_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_bc7477(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %15:void = call %textureSample_bc7477
- 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/textureSample/c2f4e8.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/c2f4e8.wgsl.expected.ir.msl
index 5e96887..ac62210 100644
--- a/test/tint/builtins/gen/var/textureSample/c2f4e8.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/c2f4e8.wgsl.expected.ir.msl
@@ -1,37 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ depthcube_array<float, access::sample> arg_0;
+ sampler arg_1;
+ device float* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_depth_cube_array, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_c2f4e8(tint_module_vars_struct tint_module_vars) {
+ float3 arg_2 = float3(1.0f);
+ int arg_3 = 1;
+ float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, arg_3);
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_c2f4e8 = func():void {
- $B2: {
- %arg_2:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(1.0f)
- %arg_3:ptr<function, i32, read_write> = var, 1i
- %7:texture_depth_cube_array = load %arg_0
- %8:sampler = load %arg_1
- %9:vec3<f32> = load %arg_2
- %10:i32 = load %arg_3
- %11:f32 = textureSample %7, %8, %9, %10
- %res:ptr<function, f32, read_write> = var, %11
- %13:f32 = load %res
- store %prevent_dce, %13
- ret
- }
+fragment void fragment_main(depthcube_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_c2f4e8(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %15:void = call %textureSample_c2f4e8
- 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/textureSample/d6b281.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/d6b281.wgsl.expected.ir.msl
index 465b0e7..1f78789 100644
--- a/test/tint/builtins/gen/var/textureSample/d6b281.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/d6b281.wgsl.expected.ir.msl
@@ -1,37 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture2d_array<float, access::sample> arg_0;
+ sampler arg_1;
+ device float4* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_2d_array<f32>, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_d6b281(tint_module_vars_struct tint_module_vars) {
+ float2 arg_2 = float2(1.0f);
+ uint arg_3 = 1u;
+ float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, arg_3);
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_d6b281 = func():void {
- $B2: {
- %arg_2:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
- %arg_3:ptr<function, u32, read_write> = var, 1u
- %7:texture_2d_array<f32> = load %arg_0
- %8:sampler = load %arg_1
- %9:vec2<f32> = load %arg_2
- %10:u32 = load %arg_3
- %11:vec4<f32> = textureSample %7, %8, %9, %10
- %res:ptr<function, vec4<f32>, read_write> = var, %11
- %13:vec4<f32> = load %res
- store %prevent_dce, %13
- ret
- }
+fragment void fragment_main(texture2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_d6b281(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %15:void = call %textureSample_d6b281
- 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/textureSample/e53267.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/e53267.wgsl.expected.ir.msl
index 8c1cd90..a3053c5 100644
--- a/test/tint/builtins/gen/var/textureSample/e53267.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/e53267.wgsl.expected.ir.msl
@@ -1,35 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texturecube<float, access::sample> arg_0;
+ sampler arg_1;
+ device float4* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_cube<f32>, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_e53267(tint_module_vars_struct tint_module_vars) {
+ float3 arg_2 = float3(1.0f);
+ float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2);
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_e53267 = func():void {
- $B2: {
- %arg_2:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(1.0f)
- %6:texture_cube<f32> = load %arg_0
- %7:sampler = load %arg_1
- %8:vec3<f32> = load %arg_2
- %9:vec4<f32> = textureSample %6, %7, %8
- %res:ptr<function, vec4<f32>, read_write> = var, %9
- %11:vec4<f32> = load %res
- store %prevent_dce, %11
- ret
- }
+fragment void fragment_main(texturecube<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_e53267(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %13:void = call %textureSample_e53267
- 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/textureSample/ea7030.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/ea7030.wgsl.expected.ir.msl
index 0962884..964f1a1 100644
--- a/test/tint/builtins/gen/var/textureSample/ea7030.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/ea7030.wgsl.expected.ir.msl
@@ -1,35 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ depthcube<float, access::sample> arg_0;
+ sampler arg_1;
+ device float* prevent_dce;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_depth_cube, read> = var @binding_point(1, 0)
- %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_ea7030(tint_module_vars_struct tint_module_vars) {
+ float3 arg_2 = float3(1.0f);
+ float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2);
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%textureSample_ea7030 = func():void {
- $B2: {
- %arg_2:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(1.0f)
- %6:texture_depth_cube = load %arg_0
- %7:sampler = load %arg_1
- %8:vec3<f32> = load %arg_2
- %9:f32 = textureSample %6, %7, %8
- %res:ptr<function, f32, read_write> = var, %9
- %11:f32 = load %res
- store %prevent_dce, %11
- ret
- }
+fragment void fragment_main(depthcube<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+ textureSample_ea7030(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B3: {
- %13:void = call %textureSample_ea7030
- 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/diagnostic_filtering/case_body_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/case_body_attribute.wgsl.expected.ir.msl
index 14f9fde..61c2299 100644
--- a/test/tint/diagnostic_filtering/case_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/case_body_attribute.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
<dawn>/test/tint/diagnostic_filtering/case_body_attribute.wgsl:8:11 warning: 'textureSample' must only be called from uniform control flow
_ = textureSample(t, s, vec2(0, 0));
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -12,10 +10,33 @@
switch (i32(x)) {
^
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-* 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. *
-********************************************************************
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture2d<float, access::sample> t;
+ sampler s;
+};
+struct tint_symbol_inputs {
+ float x [[user(locn0)]];
+};
+
+int tint_f32_to_i32(float value) {
+ return select(2147483647, select((-2147483647 - 1), int(value), (value >= -2147483648.0f)), (value <= 2147483520.0f));
+}
+void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
+ switch(tint_f32_to_i32(x)) {
+ case 0:
+ {
+ tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f));
+ break;
+ }
+ default:
+ {
+ break;
+ }
+ }
+}
+fragment void tint_symbol(tint_symbol_inputs inputs [[stage_in]], texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s};
+ tint_symbol_inner(inputs.x, tint_module_vars);
+}
diff --git a/test/tint/diagnostic_filtering/compound_statement_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/compound_statement_attribute.wgsl.expected.ir.msl
index 36bb0939..a83e3b5 100644
--- a/test/tint/diagnostic_filtering/compound_statement_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/compound_statement_attribute.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
<dawn>/test/tint/diagnostic_filtering/compound_statement_attribute.wgsl:8:11 warning: 'textureSample' must only be called from uniform control flow
_ = textureSample(t, s, vec2(0, 0));
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -12,10 +10,22 @@
if (x > 0) {
^
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-* 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. *
-********************************************************************
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture2d<float, access::sample> t;
+ sampler s;
+};
+struct tint_symbol_inputs {
+ float x [[user(locn0)]];
+};
+
+void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
+ if ((x > 0.0f)) {
+ tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f));
+ }
+}
+fragment void tint_symbol(tint_symbol_inputs inputs [[stage_in]], texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s};
+ tint_symbol_inner(inputs.x, tint_module_vars);
+}
diff --git a/test/tint/diagnostic_filtering/default_case_body_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/default_case_body_attribute.wgsl.expected.ir.msl
index 3790f50..b9f6609 100644
--- a/test/tint/diagnostic_filtering/default_case_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/default_case_body_attribute.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
<dawn>/test/tint/diagnostic_filtering/default_case_body_attribute.wgsl:8:11 warning: 'textureSample' must only be called from uniform control flow
_ = textureSample(t, s, vec2(0, 0));
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -12,10 +10,29 @@
switch (i32(x)) {
^
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-* 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. *
-********************************************************************
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture2d<float, access::sample> t;
+ sampler s;
+};
+struct tint_symbol_inputs {
+ float x [[user(locn0)]];
+};
+
+int tint_f32_to_i32(float value) {
+ return select(2147483647, select((-2147483647 - 1), int(value), (value >= -2147483648.0f)), (value <= 2147483520.0f));
+}
+void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
+ switch(tint_f32_to_i32(x)) {
+ default:
+ {
+ tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f));
+ break;
+ }
+ }
+}
+fragment void tint_symbol(tint_symbol_inputs inputs [[stage_in]], texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s};
+ tint_symbol_inner(inputs.x, tint_module_vars);
+}
diff --git a/test/tint/diagnostic_filtering/directive.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/directive.wgsl.expected.ir.msl
index 5a1ef83..9e92edd 100644
--- a/test/tint/diagnostic_filtering/directive.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/directive.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
<dawn>/test/tint/diagnostic_filtering/directive.wgsl:9:9 warning: 'textureSample' must only be called from uniform control flow
_ = textureSample(t, s, vec2(0, 0));
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -12,10 +10,22 @@
if (x > 0) {
^
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-* 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. *
-********************************************************************
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture2d<float, access::sample> t;
+ sampler s;
+};
+struct tint_symbol_inputs {
+ float x [[user(locn0)]];
+};
+
+void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
+ if ((x > 0.0f)) {
+ tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f));
+ }
+}
+fragment void tint_symbol(tint_symbol_inputs inputs [[stage_in]], texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s};
+ tint_symbol_inner(inputs.x, tint_module_vars);
+}
diff --git a/test/tint/diagnostic_filtering/else_body_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/else_body_attribute.wgsl.expected.ir.msl
index a49460b..432c17b 100644
--- a/test/tint/diagnostic_filtering/else_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/else_body_attribute.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
<dawn>/test/tint/diagnostic_filtering/else_body_attribute.wgsl:8:9 warning: 'textureSample' must only be called from uniform control flow
_ = textureSample(t, s, vec2(0, 0));
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -12,10 +10,23 @@
if (x > 0) {
^
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-* 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. *
-********************************************************************
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture2d<float, access::sample> t;
+ sampler s;
+};
+struct tint_symbol_inputs {
+ float x [[user(locn0)]];
+};
+
+void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
+ if ((x > 0.0f)) {
+ } else {
+ tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f));
+ }
+}
+fragment void tint_symbol(tint_symbol_inputs inputs [[stage_in]], texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s};
+ tint_symbol_inner(inputs.x, tint_module_vars);
+}
diff --git a/test/tint/diagnostic_filtering/else_if_body_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/else_if_body_attribute.wgsl.expected.ir.msl
index 74ecb5a..0890905 100644
--- a/test/tint/diagnostic_filtering/else_if_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/else_if_body_attribute.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
<dawn>/test/tint/diagnostic_filtering/else_if_body_attribute.wgsl:8:9 warning: 'textureSample' must only be called from uniform control flow
_ = textureSample(t, s, vec2(0, 0));
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -12,10 +10,25 @@
if (x > 0) {
^
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-* 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. *
-********************************************************************
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture2d<float, access::sample> t;
+ sampler s;
+};
+struct tint_symbol_inputs {
+ float x [[user(locn0)]];
+};
+
+void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
+ if ((x > 0.0f)) {
+ } else {
+ if ((x < 0.0f)) {
+ tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f));
+ }
+ }
+}
+fragment void tint_symbol(tint_symbol_inputs inputs [[stage_in]], texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s};
+ tint_symbol_inner(inputs.x, tint_module_vars);
+}
diff --git a/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.ir.msl
index b0a68a3..d13f01f 100644
--- a/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
<dawn>/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl:8:9 warning: 'textureSample' must only be called from uniform control flow
v = textureSample(t, s, vec2(0, 0));
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -12,10 +10,32 @@
v = textureSample(t, s, vec2(0, 0));
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-* 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. *
-********************************************************************
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture2d<float, access::sample> t;
+ sampler s;
+};
+struct tint_symbol_inputs {
+ float x [[user(locn0)]];
+};
+
+void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
+ float4 v = float4(0.0f);
+ {
+ while(true) {
+ if ((x > v[0u])) {
+ } else {
+ break;
+ }
+ v = tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f));
+ {
+ }
+ continue;
+ }
+ }
+}
+fragment void tint_symbol(tint_symbol_inputs inputs [[stage_in]], texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s};
+ tint_symbol_inner(inputs.x, tint_module_vars);
+}
diff --git a/test/tint/diagnostic_filtering/function_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/function_attribute.wgsl.expected.ir.msl
index 47e965a..b654cba 100644
--- a/test/tint/diagnostic_filtering/function_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/function_attribute.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
<dawn>/test/tint/diagnostic_filtering/function_attribute.wgsl:7:9 warning: 'textureSample' must only be called from uniform control flow
_ = textureSample(t, s, vec2(0, 0));
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -12,10 +10,22 @@
if (x > 0) {
^
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-* 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. *
-********************************************************************
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture2d<float, access::sample> t;
+ sampler s;
+};
+struct tint_symbol_inputs {
+ float x [[user(locn0)]];
+};
+
+void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
+ if ((x > 0.0f)) {
+ tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f));
+ }
+}
+fragment void tint_symbol(tint_symbol_inputs inputs [[stage_in]], texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s};
+ tint_symbol_inner(inputs.x, tint_module_vars);
+}
diff --git a/test/tint/diagnostic_filtering/function_body_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/function_body_attribute.wgsl.expected.ir.msl
index 2c424b1..86b399d 100644
--- a/test/tint/diagnostic_filtering/function_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/function_body_attribute.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
<dawn>/test/tint/diagnostic_filtering/function_body_attribute.wgsl:7:9 warning: 'textureSample' must only be called from uniform control flow
_ = textureSample(t, s, vec2(0, 0));
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -12,10 +10,22 @@
if (x > 0) {
^
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-* 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. *
-********************************************************************
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture2d<float, access::sample> t;
+ sampler s;
+};
+struct tint_symbol_inputs {
+ float x [[user(locn0)]];
+};
+
+void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
+ if ((x > 0.0f)) {
+ tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f));
+ }
+}
+fragment void tint_symbol(tint_symbol_inputs inputs [[stage_in]], texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s};
+ tint_symbol_inner(inputs.x, tint_module_vars);
+}
diff --git a/test/tint/diagnostic_filtering/if_body_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/if_body_attribute.wgsl.expected.ir.msl
index 4e582b6..dde7eb0 100644
--- a/test/tint/diagnostic_filtering/if_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/if_body_attribute.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
<dawn>/test/tint/diagnostic_filtering/if_body_attribute.wgsl:7:9 warning: 'textureSample' must only be called from uniform control flow
_ = textureSample(t, s, vec2(0, 0));
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -12,10 +10,22 @@
if (x > 0) @diagnostic(warning, derivative_uniformity) {
^
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-* 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. *
-********************************************************************
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture2d<float, access::sample> t;
+ sampler s;
+};
+struct tint_symbol_inputs {
+ float x [[user(locn0)]];
+};
+
+void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
+ if ((x > 0.0f)) {
+ tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f));
+ }
+}
+fragment void tint_symbol(tint_symbol_inputs inputs [[stage_in]], texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s};
+ tint_symbol_inner(inputs.x, tint_module_vars);
+}
diff --git a/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.ir.msl
index 930218e..689a0bf 100644
--- a/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
<dawn>/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl:8:9 warning: 'textureSample' must only be called from uniform control flow
v = textureSample(t, s, vec2(0, 0));
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -12,10 +10,32 @@
v = textureSample(t, s, vec2(0, 0));
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-* 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. *
-********************************************************************
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture2d<float, access::sample> t;
+ sampler s;
+};
+struct tint_symbol_inputs {
+ float x [[user(locn0)]];
+};
+
+void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
+ float4 v = float4(0.0f);
+ {
+ while(true) {
+ if ((x > v[0u])) {
+ } else {
+ break;
+ }
+ v = tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f));
+ {
+ }
+ continue;
+ }
+ }
+}
+fragment void tint_symbol(tint_symbol_inputs inputs [[stage_in]], texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s};
+ tint_symbol_inner(inputs.x, tint_module_vars);
+}
diff --git a/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.ir.msl b/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.ir.msl
index 4a13f56..935a3a0 100644
--- a/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.ir.msl
+++ b/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.ir.msl
@@ -1,88 +1,52 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture2d<float, access::sample> t;
+ sampler s;
+ device atomic_int* a;
+ thread bool* continue_execution;
+};
+struct foo_outputs {
+ int tint_symbol [[color(0)]];
+};
+struct foo_inputs {
+ float in [[user(locn0)]];
+ float2 coord [[user(locn1)]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %t:ptr<handle, texture_2d<f32>, read> = var @binding_point(0, 0)
- %s:ptr<handle, sampler, read> = var @binding_point(0, 1)
- %a:ptr<storage, atomic<i32>, read_write> = var @binding_point(0, 2)
- %continue_execution:ptr<private, bool, read_write> = var, true
+int tint_f32_to_i32(float value) {
+ return select(2147483647, select((-2147483647 - 1), int(value), (value >= -2147483648.0f)), (value <= 2147483520.0f));
}
-
-%foo = @fragment func(%in:f32 [@location(0)], %coord:vec2<f32> [@location(1)]):i32 [@location(0)] {
- $B2: {
- %8:bool = eq %in, 0.0f
- if %8 [t: $B3] { # if_1
- $B3: { # true
- store %continue_execution, false
- exit_if # if_1
- }
- }
- %9:texture_2d<f32> = load %t
- %10:sampler = load %s
- %11:vec4<f32> = textureSample %9, %10, %coord
- %12:f32 = access %11, 0u
- %13:i32 = call %tint_f32_to_i32, %12
- %result:ptr<function, i32, read_write> = var, %13
- loop [i: $B4, b: $B5, c: $B6] { # loop_1
- $B4: { # initializer
- %i:ptr<function, i32, read_write> = var, 0i
- next_iteration # -> $B5
- }
- $B5: { # body
- %17:i32 = load %i
- %18:bool = lt %17, 10i
- if %18 [t: $B7, f: $B8] { # if_2
- $B7: { # true
- exit_if # if_2
- }
- $B8: { # false
- exit_loop # loop_1
- }
- }
- %19:i32 = load %i
- %20:i32 = load %result
- %21:i32 = add %20, %19
- store %result, %21
- continue # -> $B6
- }
- $B6: { # continuing
- %22:bool = load %continue_execution
- %23:i32 = if %22 [t: $B9] { # if_3
- $B9: { # true
- %24:i32 = atomicAdd %a, 1i
- exit_if %24 # if_3
- }
- # implicit false block: exit_if undef
- }
- store %i, %23
- next_iteration # -> $B5
- }
- }
- %25:i32 = load %result
- %26:bool = load %continue_execution
- %27:bool = eq %26, false
- if %27 [t: $B10] { # if_4
- $B10: { # true
- terminate_invocation
- }
- }
- ret %25
+int foo_inner(float in, float2 coord, tint_module_vars_struct tint_module_vars) {
+ if ((in == 0.0f)) {
+ (*tint_module_vars.continue_execution) = false;
}
-}
-%tint_f32_to_i32 = func(%value:f32):i32 {
- $B11: {
- %29:i32 = convert %value
- %30:bool = gte %value, -2147483648.0f
- %31:i32 = select -2147483648i, %29, %30
- %32:bool = lte %value, 2147483520.0f
- %33:i32 = select 2147483647i, %31, %32
- ret %33
+ int result = tint_f32_to_i32(tint_module_vars.t.sample(tint_module_vars.s, coord)[0u]);
+ {
+ int i = 0;
+ while(true) {
+ if ((i < 10)) {
+ } else {
+ break;
+ }
+ result = (result + i);
+ {
+ int v = 0;
+ if ((*tint_module_vars.continue_execution)) {
+ v = atomic_fetch_add_explicit(tint_module_vars.a, 1, memory_order_relaxed);
+ }
+ i = v;
+ }
+ continue;
+ }
}
+ if (!((*tint_module_vars.continue_execution))) {
+ discard_fragment();
+ }
+ return result;
}
-
-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. *
-********************************************************************
+fragment foo_outputs foo(foo_inputs inputs [[stage_in]], texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]], device atomic_int* a [[buffer(0)]]) {
+ thread bool continue_execution = true;
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s, .a=a, .continue_execution=(&continue_execution)};
+ return foo_outputs{.tint_symbol=foo_inner(inputs.in, inputs.coord, tint_module_vars)};
+}
diff --git a/test/tint/unittest/reader/spirv/Samples_SpvParserHandleTest_RegisterHandleUsage_SampledImage_Variable_2.spvasm.expected.ir.msl b/test/tint/unittest/reader/spirv/Samples_SpvParserHandleTest_RegisterHandleUsage_SampledImage_Variable_2.spvasm.expected.ir.msl
deleted file mode 100644
index 3543919..0000000
--- a/test/tint/unittest/reader/spirv/Samples_SpvParserHandleTest_RegisterHandleUsage_SampledImage_Variable_2.spvasm.expected.ir.msl
+++ /dev/null
@@ -1,9 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-* 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/unittest/reader/spirv/Samples_SpvParserHandleTest_RegisterHandleUsage_SampledImage_Variable_6.spvasm.expected.ir.msl b/test/tint/unittest/reader/spirv/Samples_SpvParserHandleTest_RegisterHandleUsage_SampledImage_Variable_6.spvasm.expected.ir.msl
deleted file mode 100644
index 3543919..0000000
--- a/test/tint/unittest/reader/spirv/Samples_SpvParserHandleTest_RegisterHandleUsage_SampledImage_Variable_6.spvasm.expected.ir.msl
+++ /dev/null
@@ -1,9 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-* 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/unittest/reader/spirv/SpvParserHandleTest_NeverGenerateConstDeclForHandle_UseVariableDirectly.spvasm.expected.ir.msl b/test/tint/unittest/reader/spirv/SpvParserHandleTest_NeverGenerateConstDeclForHandle_UseVariableDirectly.spvasm.expected.ir.msl
deleted file mode 100644
index 3543919..0000000
--- a/test/tint/unittest/reader/spirv/SpvParserHandleTest_NeverGenerateConstDeclForHandle_UseVariableDirectly.spvasm.expected.ir.msl
+++ /dev/null
@@ -1,9 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-* 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. *
-********************************************************************
