diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc index 78d4dcb3bdd..91d42375257 100644 --- a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc +++ b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc @@ -57,6 +57,7 @@ #include "src/tint/lang/msl/writer/ast_raise/module_scope_var_to_entry_point_param.h" #include "src/tint/lang/msl/writer/ast_raise/packed_vec3.h" #include "src/tint/lang/msl/writer/ast_raise/pixel_local.h" +#include "src/tint/lang/msl/writer/ast_raise/quad_swap.h" #include "src/tint/lang/msl/writer/ast_raise/subgroup_ballot.h" #include "src/tint/lang/msl/writer/common/option_helpers.h" #include "src/tint/lang/msl/writer/common/printer_support.h" @@ -245,6 +246,9 @@ SanitizedResult Sanitize(const Program& in, const Options& options) { // SubgroupBallot() must come after CanonicalizeEntryPointIO. manager.Add(); + // QuadSwap() must come after CanonicalizeEntryPointIO. + manager.Add(); + // ArrayLengthFromUniform must come after SimplifyPointers, as // it assumes that the form of the array length argument is &var.array. manager.Add(); @@ -676,6 +680,19 @@ bool ASTPrinter::EmitFunctionCall(StringStream& out, return true; } + if (ast::GetAttribute(fn->Declaration()->attributes) != nullptr) { + out << "quad_shuffle("; + if (!EmitExpression(out, call->Arguments()[0]->Declaration())) { + return false; + } + out << ","; + if (!EmitExpression(out, call->Arguments()[1]->Declaration())) { + return false; + } + out << ")"; + return true; + } + out << fn->Declaration()->name->symbol.Name() << "("; bool first = true; @@ -2252,6 +2269,12 @@ bool ASTPrinter::EmitEntryPointFunction(const ast::Function* func) { auto& attrs = param->attributes; bool builtin_found = false; for (auto* attr : attrs) { + if (attr->Is()) { + out << " [[thread_index_in_quadgroup]]"; + builtin_found = true; + continue; + } + auto* builtin_attr = attr->As(); if (!builtin_attr) { continue; @@ -2264,6 +2287,7 @@ bool ASTPrinter::EmitEntryPointFunction(const ast::Function* func) { diagnostics_.AddError(Source{}) << "unknown builtin"; return false; } + out << " [[" << name << "]]"; } if (TINT_UNLIKELY(!builtin_found)) { diff --git a/src/tint/lang/msl/writer/ast_raise/BUILD.bazel b/src/tint/lang/msl/writer/ast_raise/BUILD.bazel index 3764a2a870b..94baf895e2b 100644 --- a/src/tint/lang/msl/writer/ast_raise/BUILD.bazel +++ b/src/tint/lang/msl/writer/ast_raise/BUILD.bazel @@ -42,12 +42,14 @@ cc_library( "module_scope_var_to_entry_point_param.cc", "packed_vec3.cc", "pixel_local.cc", + "quad_swap.cc", "subgroup_ballot.cc", ], hdrs = [ "module_scope_var_to_entry_point_param.h", "packed_vec3.h", "pixel_local.h", + "quad_swap.h", "subgroup_ballot.h", ], deps = [ @@ -87,6 +89,7 @@ cc_library( "module_scope_var_to_entry_point_param_test.cc", "packed_vec3_test.cc", "pixel_local_test.cc", + "quad_swap_test.cc", "subgroup_ballot_test.cc", ], deps = [ diff --git a/src/tint/lang/msl/writer/ast_raise/BUILD.cmake b/src/tint/lang/msl/writer/ast_raise/BUILD.cmake index 4366d7cc75c..5d417f935ca 100644 --- a/src/tint/lang/msl/writer/ast_raise/BUILD.cmake +++ b/src/tint/lang/msl/writer/ast_raise/BUILD.cmake @@ -47,6 +47,8 @@ tint_add_target(tint_lang_msl_writer_ast_raise lib lang/msl/writer/ast_raise/packed_vec3.h lang/msl/writer/ast_raise/pixel_local.cc lang/msl/writer/ast_raise/pixel_local.h + lang/msl/writer/ast_raise/quad_swap.cc + lang/msl/writer/ast_raise/quad_swap.h lang/msl/writer/ast_raise/subgroup_ballot.cc lang/msl/writer/ast_raise/subgroup_ballot.h ) @@ -90,6 +92,7 @@ tint_add_target(tint_lang_msl_writer_ast_raise_test test lang/msl/writer/ast_raise/module_scope_var_to_entry_point_param_test.cc lang/msl/writer/ast_raise/packed_vec3_test.cc lang/msl/writer/ast_raise/pixel_local_test.cc + lang/msl/writer/ast_raise/quad_swap_test.cc lang/msl/writer/ast_raise/subgroup_ballot_test.cc ) diff --git a/src/tint/lang/msl/writer/ast_raise/BUILD.gn b/src/tint/lang/msl/writer/ast_raise/BUILD.gn index c3f012345f3..b708177b781 100644 --- a/src/tint/lang/msl/writer/ast_raise/BUILD.gn +++ b/src/tint/lang/msl/writer/ast_raise/BUILD.gn @@ -50,6 +50,8 @@ if (tint_build_msl_writer) { "packed_vec3.h", "pixel_local.cc", "pixel_local.h", + "quad_swap.cc", + "quad_swap.h", "subgroup_ballot.cc", "subgroup_ballot.h", ] @@ -90,6 +92,7 @@ if (tint_build_unittests) { "module_scope_var_to_entry_point_param_test.cc", "packed_vec3_test.cc", "pixel_local_test.cc", + "quad_swap_test.cc", "subgroup_ballot_test.cc", ] deps = [ diff --git a/src/tint/lang/msl/writer/ast_raise/quad_swap.cc b/src/tint/lang/msl/writer/ast_raise/quad_swap.cc new file mode 100644 index 00000000000..0afe1a07f34 --- /dev/null +++ b/src/tint/lang/msl/writer/ast_raise/quad_swap.cc @@ -0,0 +1,253 @@ +// Copyright 2024 The Dawn & Tint Authors +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// +// 1. Redistributions of source code must retain the above copyright notice, this +// list of conditions and the following disclaimer. +// +// 2. Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// 3. Neither the name of the copyright holder nor the names of its +// contributors may be used to endorse or promote products derived from +// this software without specific prior written permission. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE +// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE +// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL +// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR +// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER +// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, +// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + +#include "src/tint/lang/msl/writer/ast_raise/quad_swap.h" + +#include + +#include "src/tint/lang/core/builtin_fn.h" +#include "src/tint/lang/core/fluent_types.h" +#include "src/tint/lang/wgsl/ast/transform/transform.h" +#include "src/tint/lang/wgsl/builtin_fn.h" +#include "src/tint/lang/wgsl/program/clone_context.h" +#include "src/tint/lang/wgsl/program/program_builder.h" +#include "src/tint/lang/wgsl/resolver/resolve.h" +#include "src/tint/lang/wgsl/sem/call.h" +#include "src/tint/lang/wgsl/sem/function.h" +#include "src/tint/lang/wgsl/sem/statement.h" + +TINT_INSTANTIATE_TYPEINFO(tint::msl::writer::QuadSwap); +TINT_INSTANTIATE_TYPEINFO(tint::msl::writer::QuadSwap::QuadShuffle); +TINT_INSTANTIATE_TYPEINFO(tint::msl::writer::QuadSwap::ThreadIndexInQuadgroup); + +using namespace tint::core::number_suffixes; // NOLINT +using namespace tint::core::fluent_types; // NOLINT + +namespace tint::msl::writer { + +/// PIMPL state for the transform +struct QuadSwap::State { + /// The source program + const Program& src; + /// The target program builder + ProgramBuilder b; + /// The clone context + program::CloneContext ctx = {&b, &src, /* auto_clone_symbols */ true}; + + /// The set of names for the tint_msl_quadSwap* helper functions. + Hashmap, Symbol, 8> quad_swap_helpers; + + /// The set of names for the tint_msl_quad_shuffle* intrinsic functions. + Hashmap quad_shuffle_intrinsics; + + /// The name of the `tint_msl_thread_index_in_quadgroup` global variable. + Symbol thread_index_in_quadgroup{}; + + /// The set of a functions that directly call `quadSwap*()` builtin functions. + Hashset quad_swap_callers; + + /// Constructor + /// @param program the source program + explicit State(const Program& program) : src(program) {} + + /// Runs the transform + /// @returns the new program or SkipTransform if the transform is not required + ApplyResult Run() { + auto& sem = src.Sem(); + + bool made_changes = false; + for (auto* node : ctx.src->ASTNodes().Objects()) { + auto* call = sem.Get(node); + if (call) { + // If this is a call to a `quadSwap*()` builtin, replace it with a call to the + // helper function and make a note of the function that we are in. + auto* builtin = call->Target()->As(); + if (builtin && builtin->IsQuadSwap()) { + auto* expr = ctx.Clone(call->Arguments()[0]->Declaration()); + ctx.Replace(call->Declaration(), + b.Call(GetHelper(builtin->Fn(), call->Type()), expr)); + quad_swap_callers.Add(call->Stmt()->Function()); + made_changes = true; + } + } + } + if (!made_changes) { + return SkipTransform; + } + + // Set the thread_index_in_quadgroup global variable at the start of each entry point that + // transitively calls the `quadSwap*()` builtin function. + for (auto* global : src.AST().GlobalDeclarations()) { + auto* func = global->As(); + if (func && func->IsEntryPoint() && TransitvelyCallsQuadSwap(sem.Get(func))) { + SetThreadIndexInQuadgroup(func); + } + } + + ctx.Clone(); + return resolver::Resolve(b); + } + + /// Get or create the appropriate `tint_msl_quadSwap*` helper function for the given builtin + /// function and return type. + /// @returns the name of the helper function + Symbol GetHelper(wgsl::BuiltinFn func, const core::type::Type* type) { + return quad_swap_helpers.GetOrAdd(std::make_pair(func, type), [&] { + Symbol quad_swap_helper = b.Symbols().New(std::string("tint_msl_") + str(func)); + Symbol intrinsic = GetIntrinsic(type); + + // Declare the `tint_msl_quadSwap*` helper function as follows: + // fn tint_msl_quadSwapFnName(e : T) -> T { + // tint_msl_quad_shuffle(e, tint_msl_thread_index_in_quadgroup ^ rhs) + // } + // where rhs is determined based on the builtin function call: + // +------------------+------+ + // | func | rhs | + // +------------------+------+ + // | quadSwapX | 0b1 | + // | quadSwapY | 0b10 | + // | quadSwapDiagonal | 0b11 | + // +------------------+------+ + auto* expr = b.Param("e", CreateASTTypeFor(ctx, type)); + core::u32 rhs; + switch (func) { + case wgsl::BuiltinFn::kQuadSwapX: + rhs = 0b1; + break; + case wgsl::BuiltinFn::kQuadSwapY: + rhs = 0b10; + break; + case wgsl::BuiltinFn::kQuadSwapDiagonal: + rhs = 0b11; + break; + default: + TINT_UNREACHABLE() << "unsupported builtin function"; + } + b.Func(quad_swap_helper, Vector{expr}, CreateASTTypeFor(ctx, type), + Vector{ + b.Return(b.Call(intrinsic, expr, b.Xor(GetGlobalVar(), rhs))), + }); + return quad_swap_helper; + }); + } + + /// Get or create the `tint_msl_quad_shuffle` intrinsic placeholder function for the given + /// return type. + /// @returns the name of the function + Symbol GetIntrinsic(const core::type::Type* type) { + return quad_shuffle_intrinsics.GetOrAdd(type, [&] { + auto intrinsic = b.Symbols().New(std::string("tint_msl_quad_shuffle")); + + // Declare the `tint_msl_quad_shuffle` function, which will be replaced by the MSL + // `quad_shuffle` intrinsic function to perform the swap. + { + auto* data = b.Param("data", CreateASTTypeFor(ctx, type)); + auto* quad_lane_id = b.Param("quad_lane_id", b.ty.u32()); + b.Func(intrinsic, Vector{data, quad_lane_id}, CreateASTTypeFor(ctx, type), nullptr, + Vector{b.ASTNodes().Create(b.ID(), b.AllocateNodeID()), + b.Disable(ast::DisabledValidation::kFunctionHasNoBody)}); + } + + return intrinsic; + }); + } + + /// Get or create the `tint_msl_thread_index_in_quadgroup` global variable. + /// @returns the name of the variable + Symbol GetGlobalVar() { + if (!thread_index_in_quadgroup) { + thread_index_in_quadgroup = b.Symbols().New("tint_msl_thread_index_in_quadgroup"); + + // Declare the `tint_msl_thread_index_in_quadgroup` variable. + b.GlobalVar(thread_index_in_quadgroup, core::AddressSpace::kPrivate, b.ty.u32()); + } + + return thread_index_in_quadgroup; + } + + /// Check if a function directly or transitively calls a `quadSwap*()` builtin. + /// @param func the function to check + /// @returns true if the function transitively calls `quadSwap*()` + bool TransitvelyCallsQuadSwap(const sem::Function* func) { + if (quad_swap_callers.Contains(func)) { + return true; + } + for (auto* called : func->TransitivelyCalledFunctions()) { + if (quad_swap_callers.Contains(called)) { + return true; + } + } + return false; + } + + /// Add code to set the `thread_index_in_quadgroup` variable at the start of an entry point. + /// @param ep the entry point + void SetThreadIndexInQuadgroup(const ast::Function* ep) { + // Add the entry point parameter with an attribute to indicate to the the MSL backend that + // it should be annotated with [[thread_index_in_quadgroup]]. + Symbol thread_index_in_quadgroup_param = b.Symbols().New("tint_thread_index_in_quadgroup"); + ctx.InsertBack(ep->params, b.Param(thread_index_in_quadgroup_param, b.ty.u32(), + Vector{b.ASTNodes().Create( + b.ID(), b.AllocateNodeID())})); + + // Add the following to the top of the entry point: + // { + // tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + // } + auto* block = b.Block(Vector{ + b.Assign(thread_index_in_quadgroup, thread_index_in_quadgroup_param), + }); + ctx.InsertFront(ep->body->statements, block); + } +}; + +QuadSwap::QuadSwap() = default; + +QuadSwap::~QuadSwap() = default; + +ast::transform::Transform::ApplyResult QuadSwap::Apply(const Program& src, + const ast::transform::DataMap&, + ast::transform::DataMap&) const { + return State(src).Run(); +} + +QuadSwap::QuadShuffle::~QuadShuffle() = default; + +const QuadSwap::QuadShuffle* QuadSwap::QuadShuffle::Clone(ast::CloneContext& ctx) const { + return ctx.dst->ASTNodes().Create(ctx.dst->ID(), + ctx.dst->AllocateNodeID()); +} + +QuadSwap::ThreadIndexInQuadgroup::~ThreadIndexInQuadgroup() = default; + +const QuadSwap::ThreadIndexInQuadgroup* QuadSwap::ThreadIndexInQuadgroup::Clone( + ast::CloneContext& ctx) const { + return ctx.dst->ASTNodes().Create(ctx.dst->ID(), + ctx.dst->AllocateNodeID()); +} + +} // namespace tint::msl::writer diff --git a/src/tint/lang/msl/writer/ast_raise/quad_swap.h b/src/tint/lang/msl/writer/ast_raise/quad_swap.h new file mode 100644 index 00000000000..1654ad43855 --- /dev/null +++ b/src/tint/lang/msl/writer/ast_raise/quad_swap.h @@ -0,0 +1,113 @@ +// Copyright 2024 The Dawn & Tint Authors +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// +// 1. Redistributions of source code must retain the above copyright notice, this +// list of conditions and the following disclaimer. +// +// 2. Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// 3. Neither the name of the copyright holder nor the names of its +// contributors may be used to endorse or promote products derived from +// this software without specific prior written permission. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE +// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE +// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL +// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR +// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER +// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, +// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + +#ifndef SRC_TINT_LANG_MSL_WRITER_AST_RAISE_QUAD_SWAP_H_ +#define SRC_TINT_LANG_MSL_WRITER_AST_RAISE_QUAD_SWAP_H_ + +#include + +#include "src/tint/lang/wgsl/ast/internal_attribute.h" +#include "src/tint/lang/wgsl/ast/transform/transform.h" + +namespace tint::msl::writer { + +/// QuadSwap is a transform that replaces calls to `quadSwap{X, Y, Diagonal}()` with an +/// implementation that uses MSL's `quad_shuffle()` using the following mapping: +/// +------------------+------------------------------------------------------------------+ +/// | WGSL | MSL | +/// +------------------+------------------------------------------------------------------+ +/// | quadSwapX | quad_shuffle with quad_lane_id=thread_index_in_quad_group ^ 0b1 | +/// | quadSwapY | quad_shuffle with quad_lane_id=thread_index_in_quad_group ^ 0b10 | +/// | quadSwapDiagonal | quad_shuffle with quad_lane_id=thread_index_in_quad_group ^ 0b11 | +/// +------------------+------------------------------------------------------------------+ +/// +/// @note Depends on the following transforms to have been run first: +/// * CanonicalizeEntryPointIO +class QuadSwap final : public Castable { + public: + /// Constructor + QuadSwap(); + + /// Destructor + ~QuadSwap() override; + + /// @copydoc ast::transform::Transform::Apply + ApplyResult Apply(const Program& program, + const ast::transform::DataMap& inputs, + ast::transform::DataMap& outputs) const override; + + /// QuadShuffle is an InternalAttribute that is used to decorate a stub function so + /// that the MSL backend transforms this into calls to the `quad_shuffle` function. + class QuadShuffle final : public Castable { + public: + /// Constructor + /// @param pid the identifier of the program that owns this node + /// @param nid the unique node identifier + QuadShuffle(GenerationID pid, ast::NodeID nid) : Base(pid, nid, Empty) {} + + /// Destructor + ~QuadShuffle() override; + + /// @copydoc ast::InternalAttribute::InternalName + std::string InternalName() const override { return "quad_shuffle"; } + + /// Performs a deep clone of this object using the program::CloneContext `ctx`. + /// @param ctx the clone context + /// @return the newly cloned object + const QuadShuffle* Clone(ast::CloneContext& ctx) const override; + }; + + /// ThreadIndexInQuadgroup is an InternalAttribute that is used to decorate an entrypoint + /// parameter so that the MSL backend transforms this into a `[[thread_index_in_quadgroup]]` + /// attribute. + class ThreadIndexInQuadgroup final + : public Castable { + public: + /// Constructor + /// @param pid the identifier of the program that owns this node + /// @param nid the unique node identifier + ThreadIndexInQuadgroup(GenerationID pid, ast::NodeID nid) : Base(pid, nid, Empty) {} + + /// Destructor + ~ThreadIndexInQuadgroup() override; + + /// @copydoc ast::InternalAttribute::InternalName + std::string InternalName() const override { return "thread_index_in_quadgroup"; } + + /// Performs a deep clone of this object using the program::CloneContext `ctx`. + /// @param ctx the clone context + /// @return the newly cloned object + const ThreadIndexInQuadgroup* Clone(ast::CloneContext& ctx) const override; + }; + + private: + struct State; +}; + +} // namespace tint::msl::writer + +#endif // SRC_TINT_LANG_MSL_WRITER_AST_RAISE_QUAD_SWAP_H_ diff --git a/src/tint/lang/msl/writer/ast_raise/quad_swap_test.cc b/src/tint/lang/msl/writer/ast_raise/quad_swap_test.cc new file mode 100644 index 00000000000..f24f48ec3f4 --- /dev/null +++ b/src/tint/lang/msl/writer/ast_raise/quad_swap_test.cc @@ -0,0 +1,376 @@ +// Copyright 2024 The Dawn & Tint Authors +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// +// 1. Redistributions of source code must retain the above copyright notice, this +// list of conditions and the following disclaimer. +// +// 2. Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// 3. Neither the name of the copyright holder nor the names of its +// contributors may be used to endorse or promote products derived from +// this software without specific prior written permission. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE +// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE +// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL +// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR +// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER +// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, +// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + +#include "src/tint/lang/msl/writer/ast_raise/quad_swap.h" + +#include "src/tint/lang/wgsl/ast/transform/helper_test.h" + +namespace tint::msl::writer { +namespace { + +using QuadSwapTest = ast::transform::TransformTest; + +TEST_F(QuadSwapTest, EmptyModule) { + auto* src = ""; + + EXPECT_FALSE(ShouldRun(src)); +} + +TEST_F(QuadSwapTest, DirectUseQuadSwapX) { + auto* src = R"( +enable subgroups; + +@compute @workgroup_size(64) +fn foo() { + let x: f32 = quadSwapX(1.f); +} +)"; + + auto* expect = + R"( +enable subgroups; + +@internal(quad_shuffle) @internal(disable_validation__function_has_no_body) +fn tint_msl_quad_shuffle(data : f32, quad_lane_id : u32) -> f32 + +var tint_msl_thread_index_in_quadgroup : u32; + +fn tint_msl_quadSwapX(e : f32) -> f32 { + return tint_msl_quad_shuffle(e, (tint_msl_thread_index_in_quadgroup ^ 1u)); +} + +@compute @workgroup_size(64) +fn foo(@internal(thread_index_in_quadgroup) tint_thread_index_in_quadgroup : u32) { + { + tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + let x : f32 = tint_msl_quadSwapX(1.0f); +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(QuadSwapTest, IndirectUseQuadSwapX) { + auto* src = R"( +enable subgroups; + +fn bar() -> vec4u { + let expr = vec4u(1u, 1u, 1u, 1u); + return quadSwapX(expr); +} + +@compute @workgroup_size(64) +fn foo() { + let x: vec4u = bar(); +} +)"; + + auto* expect = + R"( +enable subgroups; + +@internal(quad_shuffle) @internal(disable_validation__function_has_no_body) +fn tint_msl_quad_shuffle(data : vec4, quad_lane_id : u32) -> vec4 + +var tint_msl_thread_index_in_quadgroup : u32; + +fn tint_msl_quadSwapX(e : vec4) -> vec4 { + return tint_msl_quad_shuffle(e, (tint_msl_thread_index_in_quadgroup ^ 1u)); +} + +fn bar() -> vec4u { + let expr = vec4u(1u, 1u, 1u, 1u); + return tint_msl_quadSwapX(expr); +} + +@compute @workgroup_size(64) +fn foo(@internal(thread_index_in_quadgroup) tint_thread_index_in_quadgroup : u32) { + { + tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + let x : vec4u = bar(); +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(QuadSwapTest, DirectUseQuadSwapY) { + auto* src = R"( +enable subgroups; + +@compute @workgroup_size(64) +fn foo() { + let x: i32 = quadSwapY(1i); +} +)"; + + auto* expect = + R"( +enable subgroups; + +@internal(quad_shuffle) @internal(disable_validation__function_has_no_body) +fn tint_msl_quad_shuffle(data : i32, quad_lane_id : u32) -> i32 + +var tint_msl_thread_index_in_quadgroup : u32; + +fn tint_msl_quadSwapY(e : i32) -> i32 { + return tint_msl_quad_shuffle(e, (tint_msl_thread_index_in_quadgroup ^ 2u)); +} + +@compute @workgroup_size(64) +fn foo(@internal(thread_index_in_quadgroup) tint_thread_index_in_quadgroup : u32) { + { + tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + let x : i32 = tint_msl_quadSwapY(1i); +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(QuadSwapTest, DirectUseQuadSwapDiagonal) { + auto* src = R"( +enable subgroups; + +@compute @workgroup_size(64) +fn foo() { + let x: i32 = quadSwapDiagonal(1i); +} +)"; + + auto* expect = + R"( +enable subgroups; + +@internal(quad_shuffle) @internal(disable_validation__function_has_no_body) +fn tint_msl_quad_shuffle(data : i32, quad_lane_id : u32) -> i32 + +var tint_msl_thread_index_in_quadgroup : u32; + +fn tint_msl_quadSwapDiagonal(e : i32) -> i32 { + return tint_msl_quad_shuffle(e, (tint_msl_thread_index_in_quadgroup ^ 3u)); +} + +@compute @workgroup_size(64) +fn foo(@internal(thread_index_in_quadgroup) tint_thread_index_in_quadgroup : u32) { + { + tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + let x : i32 = tint_msl_quadSwapDiagonal(1i); +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(QuadSwapTest, MultipleCallsSameFunction) { + auto* src = R"( +enable subgroups; + +@compute @workgroup_size(64) +fn foo() { + let x: u32 = quadSwapX(1u); + let y: u32 = quadSwapX(1u); +} +)"; + + auto* expect = + R"( +enable subgroups; + +@internal(quad_shuffle) @internal(disable_validation__function_has_no_body) +fn tint_msl_quad_shuffle(data : u32, quad_lane_id : u32) -> u32 + +var tint_msl_thread_index_in_quadgroup : u32; + +fn tint_msl_quadSwapX(e : u32) -> u32 { + return tint_msl_quad_shuffle(e, (tint_msl_thread_index_in_quadgroup ^ 1u)); +} + +@compute @workgroup_size(64) +fn foo(@internal(thread_index_in_quadgroup) tint_thread_index_in_quadgroup : u32) { + { + tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + let x : u32 = tint_msl_quadSwapX(1u); + let y : u32 = tint_msl_quadSwapX(1u); +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(QuadSwapTest, MultipleCallsSameFunctionDifferentType) { + auto* src = R"( +enable subgroups; + +@compute @workgroup_size(64) +fn foo() { + let x: u32 = quadSwapX(1u); + let y: i32 = quadSwapX(1i); +} +)"; + + auto* expect = + R"( +enable subgroups; + +@internal(quad_shuffle) @internal(disable_validation__function_has_no_body) +fn tint_msl_quad_shuffle(data : u32, quad_lane_id : u32) -> u32 + +var tint_msl_thread_index_in_quadgroup : u32; + +fn tint_msl_quadSwapX(e : u32) -> u32 { + return tint_msl_quad_shuffle(e, (tint_msl_thread_index_in_quadgroup ^ 1u)); +} + +@internal(quad_shuffle) @internal(disable_validation__function_has_no_body) +fn tint_msl_quad_shuffle_1(data : i32, quad_lane_id : u32) -> i32 + +fn tint_msl_quadSwapX_1(e : i32) -> i32 { + return tint_msl_quad_shuffle_1(e, (tint_msl_thread_index_in_quadgroup ^ 1u)); +} + +@compute @workgroup_size(64) +fn foo(@internal(thread_index_in_quadgroup) tint_thread_index_in_quadgroup : u32) { + { + tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + let x : u32 = tint_msl_quadSwapX(1u); + let y : i32 = tint_msl_quadSwapX_1(1i); +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(QuadSwapTest, MultipleCallsSameTypeDifferentFunction) { + auto* src = R"( +enable subgroups; + +@compute @workgroup_size(64) +fn foo() { + let x: u32 = quadSwapX(1u); + let y: u32 = quadSwapY(1u); +} +)"; + + auto* expect = + R"( +enable subgroups; + +@internal(quad_shuffle) @internal(disable_validation__function_has_no_body) +fn tint_msl_quad_shuffle(data : u32, quad_lane_id : u32) -> u32 + +var tint_msl_thread_index_in_quadgroup : u32; + +fn tint_msl_quadSwapX(e : u32) -> u32 { + return tint_msl_quad_shuffle(e, (tint_msl_thread_index_in_quadgroup ^ 1u)); +} + +fn tint_msl_quadSwapY(e : u32) -> u32 { + return tint_msl_quad_shuffle(e, (tint_msl_thread_index_in_quadgroup ^ 2u)); +} + +@compute @workgroup_size(64) +fn foo(@internal(thread_index_in_quadgroup) tint_thread_index_in_quadgroup : u32) { + { + tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + let x : u32 = tint_msl_quadSwapX(1u); + let y : u32 = tint_msl_quadSwapY(1u); +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(QuadSwapTest, MultipleCallsDifferentTypeDifferentFunction) { + auto* src = R"( +enable subgroups; + +@compute @workgroup_size(64) +fn foo() { + let x: i32 = quadSwapX(1i); + let y: u32 = quadSwapY(1u); +} +)"; + + auto* expect = + R"( +enable subgroups; + +@internal(quad_shuffle) @internal(disable_validation__function_has_no_body) +fn tint_msl_quad_shuffle(data : i32, quad_lane_id : u32) -> i32 + +var tint_msl_thread_index_in_quadgroup : u32; + +fn tint_msl_quadSwapX(e : i32) -> i32 { + return tint_msl_quad_shuffle(e, (tint_msl_thread_index_in_quadgroup ^ 1u)); +} + +@internal(quad_shuffle) @internal(disable_validation__function_has_no_body) +fn tint_msl_quad_shuffle_1(data : u32, quad_lane_id : u32) -> u32 + +fn tint_msl_quadSwapY(e : u32) -> u32 { + return tint_msl_quad_shuffle_1(e, (tint_msl_thread_index_in_quadgroup ^ 2u)); +} + +@compute @workgroup_size(64) +fn foo(@internal(thread_index_in_quadgroup) tint_thread_index_in_quadgroup : u32) { + { + tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + let x : i32 = tint_msl_quadSwapX(1i); + let y : u32 = tint_msl_quadSwapY(1u); +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +} // namespace +} // namespace tint::msl::writer diff --git a/src/tint/lang/msl/writer/printer/printer.cc b/src/tint/lang/msl/writer/printer/printer.cc index ab72cf30621..7595cf2b357 100644 --- a/src/tint/lang/msl/writer/printer/printer.cc +++ b/src/tint/lang/msl/writer/printer/printer.cc @@ -877,7 +877,8 @@ class Printer : public tint::TextGenerator { } void EmitCoreBuiltinName(StringStream& out, core::BuiltinFn func) { - // TODO(crbug.com/354738715): Implement subgroupBallot in MSL IR backend. + // TODO(crbug.com/360166776): Implement subgroupBallot and quadSwap{X,Y,Diagonal} builtins + // in MSL IR backend. switch (func) { case core::BuiltinFn::kAbs: case core::BuiltinFn::kAcos: diff --git a/src/tint/lang/wgsl/builtin_fn.cc b/src/tint/lang/wgsl/builtin_fn.cc index 68efef3d2bf..75ea131110b 100644 --- a/src/tint/lang/wgsl/builtin_fn.cc +++ b/src/tint/lang/wgsl/builtin_fn.cc @@ -877,6 +877,17 @@ bool IsSubgroup(BuiltinFn f) { } } +bool IsQuadSwap(BuiltinFn f) { + switch (f) { + case BuiltinFn::kQuadSwapX: + case BuiltinFn::kQuadSwapY: + case BuiltinFn::kQuadSwapDiagonal: + return true; + default: + return false; + } +} + bool HasSideEffects(BuiltinFn f) { switch (f) { case BuiltinFn::kAtomicAdd: diff --git a/src/tint/lang/wgsl/builtin_fn.cc.tmpl b/src/tint/lang/wgsl/builtin_fn.cc.tmpl index 7a88242784f..6fb225668e2 100644 --- a/src/tint/lang/wgsl/builtin_fn.cc.tmpl +++ b/src/tint/lang/wgsl/builtin_fn.cc.tmpl @@ -138,6 +138,17 @@ bool IsSubgroup(BuiltinFn f) { } } +bool IsQuadSwap(BuiltinFn f) { + switch (f) { + case BuiltinFn::kQuadSwapX: + case BuiltinFn::kQuadSwapY: + case BuiltinFn::kQuadSwapDiagonal: + return true; + default: + return false; + } +} + bool HasSideEffects(BuiltinFn f) { switch (f) { case BuiltinFn::kAtomicAdd: diff --git a/src/tint/lang/wgsl/builtin_fn.h b/src/tint/lang/wgsl/builtin_fn.h index 75753539db8..d39cd405908 100644 --- a/src/tint/lang/wgsl/builtin_fn.h +++ b/src/tint/lang/wgsl/builtin_fn.h @@ -570,6 +570,11 @@ bool IsPacked4x8IntegerDotProductBuiltin(BuiltinFn f); /// @returns true if the given `f` is a subgroup builtin bool IsSubgroup(BuiltinFn f); +/// Determines if the given `f` is a quadSwap* builtin. +/// @param f the builtin type +/// @returns true if the given `f` is a quadSwap* builtin +bool IsQuadSwap(BuiltinFn f); + /// Determines if the given `f` may have side-effects (i.e. writes to at least one of its inputs) /// @returns true if intrinsic may have side-effects bool HasSideEffects(BuiltinFn f); diff --git a/src/tint/lang/wgsl/builtin_fn.h.tmpl b/src/tint/lang/wgsl/builtin_fn.h.tmpl index 57d0fe8f4c4..b5b1602b9e9 100644 --- a/src/tint/lang/wgsl/builtin_fn.h.tmpl +++ b/src/tint/lang/wgsl/builtin_fn.h.tmpl @@ -120,6 +120,11 @@ bool IsPacked4x8IntegerDotProductBuiltin(BuiltinFn f); /// @returns true if the given `f` is a subgroup builtin bool IsSubgroup(BuiltinFn f); +/// Determines if the given `f` is a quadSwap* builtin. +/// @param f the builtin type +/// @returns true if the given `f` is a quadSwap* builtin +bool IsQuadSwap(BuiltinFn f); + /// Determines if the given `f` may have side-effects (i.e. writes to at least one of its inputs) /// @returns true if intrinsic may have side-effects bool HasSideEffects(BuiltinFn f); diff --git a/src/tint/lang/wgsl/resolver/validator.cc b/src/tint/lang/wgsl/resolver/validator.cc index a873f921548..16c46923521 100644 --- a/src/tint/lang/wgsl/resolver/validator.cc +++ b/src/tint/lang/wgsl/resolver/validator.cc @@ -1432,6 +1432,10 @@ bool Validator::EntryPoint(const sem::Function* func, ast::PipelineStage stage) invariant_attribute = invariant; return InvariantAttribute(invariant, stage); }, + [&](const ast::InternalAttribute* internal) { + pipeline_io_attribute = internal; + return true; + }, [&](Default) { return true; }); if (!ok) { diff --git a/src/tint/lang/wgsl/sem/builtin_fn.cc b/src/tint/lang/wgsl/sem/builtin_fn.cc index c602fc51357..fc6f001a9a9 100644 --- a/src/tint/lang/wgsl/sem/builtin_fn.cc +++ b/src/tint/lang/wgsl/sem/builtin_fn.cc @@ -108,6 +108,10 @@ bool BuiltinFn::IsSubgroup() const { return wgsl::IsSubgroup(fn_); } +bool BuiltinFn::IsQuadSwap() const { + return wgsl::IsQuadSwap(fn_); +} + bool BuiltinFn::HasSideEffects() const { return wgsl::HasSideEffects(fn_); } diff --git a/src/tint/lang/wgsl/sem/builtin_fn.h b/src/tint/lang/wgsl/sem/builtin_fn.h index f9abb9f0096..69b542e9a63 100644 --- a/src/tint/lang/wgsl/sem/builtin_fn.h +++ b/src/tint/lang/wgsl/sem/builtin_fn.h @@ -113,6 +113,9 @@ class BuiltinFn final : public Castable { /// chromium_experimental_subgroups) bool IsSubgroup() const; + /// @returns true if builtin is a quadSwap builtin + bool IsQuadSwap() const; + /// @returns true if intrinsic may have side-effects (i.e. writes to at least /// one of its inputs) bool HasSideEffects() const; diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/15ac75.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapDiagonal/15ac75.wgsl.expected.msl index 4d52eed7448..e1e4fa66550 100644 --- a/test/tint/builtins/gen/literal/quadSwapDiagonal/15ac75.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/15ac75.wgsl.expected.msl @@ -1,25 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -enable subgroups_f16; -enable f16; -@group(0) @binding(0) var prevent_dce : vec2; +half2 tint_msl_quadSwapDiagonal(half2 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_15ac75() -> vec2 { - var res : vec2 = quadSwapDiagonal(vec2(1.0h)); +half2 quadSwapDiagonal_15ac75(thread tint_private_vars_struct* const tint_private_vars) { + half2 res = tint_msl_quadSwapDiagonal(half2(1.0h), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_15ac75(); +fragment void fragment_main(device half2* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapDiagonal_15ac75(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_15ac75(); +kernel void compute_main(device half2* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapDiagonal_15ac75(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/2be5e7.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapDiagonal/2be5e7.wgsl.expected.msl index 223e3542cc1..8702036ff02 100644 --- a/test/tint/builtins/gen/literal/quadSwapDiagonal/2be5e7.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/2be5e7.wgsl.expected.msl @@ -1,25 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -enable subgroups_f16; -enable f16; -@group(0) @binding(0) var prevent_dce : f16; +half tint_msl_quadSwapDiagonal(half e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_2be5e7() -> f16 { - var res : f16 = quadSwapDiagonal(1.0h); +half quadSwapDiagonal_2be5e7(thread tint_private_vars_struct* const tint_private_vars) { + half res = tint_msl_quadSwapDiagonal(1.0h, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_2be5e7(); +fragment void fragment_main(device half* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapDiagonal_2be5e7(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_2be5e7(); +kernel void compute_main(device half* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapDiagonal_2be5e7(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/331804.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapDiagonal/331804.wgsl.expected.msl index 45c204425df..96041e09079 100644 --- a/test/tint/builtins/gen/literal/quadSwapDiagonal/331804.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/331804.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec4; +float4 tint_msl_quadSwapDiagonal(float4 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_331804() -> vec4 { - var res : vec4 = quadSwapDiagonal(vec4(1.0f)); +float4 quadSwapDiagonal_331804(thread tint_private_vars_struct* const tint_private_vars) { + float4 res = tint_msl_quadSwapDiagonal(float4(1.0f), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_331804(); +fragment void fragment_main(device float4* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapDiagonal_331804(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_331804(); +kernel void compute_main(device float4* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapDiagonal_331804(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/348173.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapDiagonal/348173.wgsl.expected.msl index 8f21d6e7fac..dc2f8fd493d 100644 --- a/test/tint/builtins/gen/literal/quadSwapDiagonal/348173.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/348173.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec2; +uint2 tint_msl_quadSwapDiagonal(uint2 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_348173() -> vec2 { - var res : vec2 = quadSwapDiagonal(vec2(1u)); +uint2 quadSwapDiagonal_348173(thread tint_private_vars_struct* const tint_private_vars) { + uint2 res = tint_msl_quadSwapDiagonal(uint2(1u), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_348173(); +fragment void fragment_main(device uint2* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapDiagonal_348173(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_348173(); +kernel void compute_main(device uint2* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapDiagonal_348173(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/486196.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapDiagonal/486196.wgsl.expected.msl index bf00d285e3e..68b93660b55 100644 --- a/test/tint/builtins/gen/literal/quadSwapDiagonal/486196.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/486196.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : f32; +float tint_msl_quadSwapDiagonal(float e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_486196() -> f32 { - var res : f32 = quadSwapDiagonal(1.0f); +float quadSwapDiagonal_486196(thread tint_private_vars_struct* const tint_private_vars) { + float res = tint_msl_quadSwapDiagonal(1.0f, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_486196(); +fragment void fragment_main(device float* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapDiagonal_486196(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_486196(); +kernel void compute_main(device float* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapDiagonal_486196(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/730e40.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapDiagonal/730e40.wgsl.expected.msl index e7b96b46cac..33d371fa125 100644 --- a/test/tint/builtins/gen/literal/quadSwapDiagonal/730e40.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/730e40.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : u32; +uint tint_msl_quadSwapDiagonal(uint e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_730e40() -> u32 { - var res : u32 = quadSwapDiagonal(1u); +uint quadSwapDiagonal_730e40(thread tint_private_vars_struct* const tint_private_vars) { + uint res = tint_msl_quadSwapDiagonal(1u, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_730e40(); +fragment void fragment_main(device uint* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapDiagonal_730e40(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_730e40(); +kernel void compute_main(device uint* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapDiagonal_730e40(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/8077c8.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapDiagonal/8077c8.wgsl.expected.msl index ce5a8e68eac..b940540913a 100644 --- a/test/tint/builtins/gen/literal/quadSwapDiagonal/8077c8.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/8077c8.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec2; +float2 tint_msl_quadSwapDiagonal(float2 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_8077c8() -> vec2 { - var res : vec2 = quadSwapDiagonal(vec2(1.0f)); +float2 quadSwapDiagonal_8077c8(thread tint_private_vars_struct* const tint_private_vars) { + float2 res = tint_msl_quadSwapDiagonal(float2(1.0f), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_8077c8(); +fragment void fragment_main(device float2* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapDiagonal_8077c8(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_8077c8(); +kernel void compute_main(device float2* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapDiagonal_8077c8(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/856536.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapDiagonal/856536.wgsl.expected.msl index cbc6975a1c2..3d8c5fefdb8 100644 --- a/test/tint/builtins/gen/literal/quadSwapDiagonal/856536.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/856536.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec3; +uint3 tint_msl_quadSwapDiagonal(uint3 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_856536() -> vec3 { - var res : vec3 = quadSwapDiagonal(vec3(1u)); +uint3 quadSwapDiagonal_856536(thread tint_private_vars_struct* const tint_private_vars) { + uint3 res = tint_msl_quadSwapDiagonal(uint3(1u), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_856536(); +fragment void fragment_main(device packed_uint3* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = packed_uint3(quadSwapDiagonal_856536(&(tint_private_vars))); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_856536(); +kernel void compute_main(device packed_uint3* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = packed_uint3(quadSwapDiagonal_856536(&(tint_private_vars))); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/9ccb38.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapDiagonal/9ccb38.wgsl.expected.msl index c65ead41d7e..da89e0881a7 100644 --- a/test/tint/builtins/gen/literal/quadSwapDiagonal/9ccb38.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/9ccb38.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : i32; +int tint_msl_quadSwapDiagonal(int e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_9ccb38() -> i32 { - var res : i32 = quadSwapDiagonal(1i); +int quadSwapDiagonal_9ccb38(thread tint_private_vars_struct* const tint_private_vars) { + int res = tint_msl_quadSwapDiagonal(1, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_9ccb38(); +fragment void fragment_main(device int* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapDiagonal_9ccb38(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_9ccb38(); +kernel void compute_main(device int* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapDiagonal_9ccb38(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/a090b0.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapDiagonal/a090b0.wgsl.expected.msl index 2ae25f733bf..1176481409f 100644 --- a/test/tint/builtins/gen/literal/quadSwapDiagonal/a090b0.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/a090b0.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec2; +int2 tint_msl_quadSwapDiagonal(int2 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_a090b0() -> vec2 { - var res : vec2 = quadSwapDiagonal(vec2(1i)); +int2 quadSwapDiagonal_a090b0(thread tint_private_vars_struct* const tint_private_vars) { + int2 res = tint_msl_quadSwapDiagonal(int2(1), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_a090b0(); +fragment void fragment_main(device int2* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapDiagonal_a090b0(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_a090b0(); +kernel void compute_main(device int2* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapDiagonal_a090b0(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/a665b1.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapDiagonal/a665b1.wgsl.expected.msl index 10957e0a8d2..051427d9d67 100644 --- a/test/tint/builtins/gen/literal/quadSwapDiagonal/a665b1.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/a665b1.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec4; +int4 tint_msl_quadSwapDiagonal(int4 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_a665b1() -> vec4 { - var res : vec4 = quadSwapDiagonal(vec4(1i)); +int4 quadSwapDiagonal_a665b1(thread tint_private_vars_struct* const tint_private_vars) { + int4 res = tint_msl_quadSwapDiagonal(int4(1), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_a665b1(); +fragment void fragment_main(device int4* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapDiagonal_a665b1(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_a665b1(); +kernel void compute_main(device int4* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapDiagonal_a665b1(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/a82e1d.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapDiagonal/a82e1d.wgsl.expected.msl index 7dba196ddf2..673a5b3f148 100644 --- a/test/tint/builtins/gen/literal/quadSwapDiagonal/a82e1d.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/a82e1d.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec3; +int3 tint_msl_quadSwapDiagonal(int3 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_a82e1d() -> vec3 { - var res : vec3 = quadSwapDiagonal(vec3(1i)); +int3 quadSwapDiagonal_a82e1d(thread tint_private_vars_struct* const tint_private_vars) { + int3 res = tint_msl_quadSwapDiagonal(int3(1), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_a82e1d(); +fragment void fragment_main(device packed_int3* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = packed_int3(quadSwapDiagonal_a82e1d(&(tint_private_vars))); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_a82e1d(); +kernel void compute_main(device packed_int3* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = packed_int3(quadSwapDiagonal_a82e1d(&(tint_private_vars))); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/af19a5.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapDiagonal/af19a5.wgsl.expected.msl index 0ece294e5eb..a48763a5fdc 100644 --- a/test/tint/builtins/gen/literal/quadSwapDiagonal/af19a5.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/af19a5.wgsl.expected.msl @@ -1,25 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -enable subgroups_f16; -enable f16; -@group(0) @binding(0) var prevent_dce : vec4; +half4 tint_msl_quadSwapDiagonal(half4 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_af19a5() -> vec4 { - var res : vec4 = quadSwapDiagonal(vec4(1.0h)); +half4 quadSwapDiagonal_af19a5(thread tint_private_vars_struct* const tint_private_vars) { + half4 res = tint_msl_quadSwapDiagonal(half4(1.0h), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_af19a5(); +fragment void fragment_main(device half4* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapDiagonal_af19a5(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_af19a5(); +kernel void compute_main(device half4* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapDiagonal_af19a5(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/b905fc.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapDiagonal/b905fc.wgsl.expected.msl index 69f807c9fd6..c08ab79fa8d 100644 --- a/test/tint/builtins/gen/literal/quadSwapDiagonal/b905fc.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/b905fc.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec3; +float3 tint_msl_quadSwapDiagonal(float3 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_b905fc() -> vec3 { - var res : vec3 = quadSwapDiagonal(vec3(1.0f)); +float3 quadSwapDiagonal_b905fc(thread tint_private_vars_struct* const tint_private_vars) { + float3 res = tint_msl_quadSwapDiagonal(float3(1.0f), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_b905fc(); +fragment void fragment_main(device packed_float3* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = packed_float3(quadSwapDiagonal_b905fc(&(tint_private_vars))); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_b905fc(); +kernel void compute_main(device packed_float3* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = packed_float3(quadSwapDiagonal_b905fc(&(tint_private_vars))); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/c31636.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapDiagonal/c31636.wgsl.expected.msl index 0ff8731592f..ba9ced85538 100644 --- a/test/tint/builtins/gen/literal/quadSwapDiagonal/c31636.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/c31636.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec4; +uint4 tint_msl_quadSwapDiagonal(uint4 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_c31636() -> vec4 { - var res : vec4 = quadSwapDiagonal(vec4(1u)); +uint4 quadSwapDiagonal_c31636(thread tint_private_vars_struct* const tint_private_vars) { + uint4 res = tint_msl_quadSwapDiagonal(uint4(1u), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_c31636(); +fragment void fragment_main(device uint4* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapDiagonal_c31636(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_c31636(); +kernel void compute_main(device uint4* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapDiagonal_c31636(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/e4bec8.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapDiagonal/e4bec8.wgsl.expected.msl index b761433f5c6..004641d6cfe 100644 --- a/test/tint/builtins/gen/literal/quadSwapDiagonal/e4bec8.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/e4bec8.wgsl.expected.msl @@ -1,25 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -enable subgroups_f16; -enable f16; -@group(0) @binding(0) var prevent_dce : vec3; +half3 tint_msl_quadSwapDiagonal(half3 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_e4bec8() -> vec3 { - var res : vec3 = quadSwapDiagonal(vec3(1.0h)); +half3 quadSwapDiagonal_e4bec8(thread tint_private_vars_struct* const tint_private_vars) { + half3 res = tint_msl_quadSwapDiagonal(half3(1.0h), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_e4bec8(); +fragment void fragment_main(device packed_half3* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = packed_half3(quadSwapDiagonal_e4bec8(&(tint_private_vars))); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_e4bec8(); +kernel void compute_main(device packed_half3* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = packed_half3(quadSwapDiagonal_e4bec8(&(tint_private_vars))); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/literal/quadSwapX/02834c.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapX/02834c.wgsl.expected.msl index f6afa819db7..eb010e39231 100644 --- a/test/tint/builtins/gen/literal/quadSwapX/02834c.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapX/02834c.wgsl.expected.msl @@ -1,25 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -enable subgroups_f16; -enable f16; -@group(0) @binding(0) var prevent_dce : vec4; +half4 tint_msl_quadSwapX(half4 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_02834c() -> vec4 { - var res : vec4 = quadSwapX(vec4(1.0h)); +half4 quadSwapX_02834c(thread tint_private_vars_struct* const tint_private_vars) { + half4 res = tint_msl_quadSwapX(half4(1.0h), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_02834c(); +fragment void fragment_main(device half4* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapX_02834c(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_02834c(); +kernel void compute_main(device half4* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapX_02834c(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/literal/quadSwapX/053f3b.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapX/053f3b.wgsl.expected.msl index 80e841248ca..254c2d5872d 100644 --- a/test/tint/builtins/gen/literal/quadSwapX/053f3b.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapX/053f3b.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec3; +int3 tint_msl_quadSwapX(int3 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_053f3b() -> vec3 { - var res : vec3 = quadSwapX(vec3(1i)); +int3 quadSwapX_053f3b(thread tint_private_vars_struct* const tint_private_vars) { + int3 res = tint_msl_quadSwapX(int3(1), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_053f3b(); +fragment void fragment_main(device packed_int3* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = packed_int3(quadSwapX_053f3b(&(tint_private_vars))); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_053f3b(); +kernel void compute_main(device packed_int3* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = packed_int3(quadSwapX_053f3b(&(tint_private_vars))); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/literal/quadSwapX/07f1fc.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapX/07f1fc.wgsl.expected.msl index 4ff3b627d7b..4f79a183af8 100644 --- a/test/tint/builtins/gen/literal/quadSwapX/07f1fc.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapX/07f1fc.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec4; +uint4 tint_msl_quadSwapX(uint4 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_07f1fc() -> vec4 { - var res : vec4 = quadSwapX(vec4(1u)); +uint4 quadSwapX_07f1fc(thread tint_private_vars_struct* const tint_private_vars) { + uint4 res = tint_msl_quadSwapX(uint4(1u), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_07f1fc(); +fragment void fragment_main(device uint4* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapX_07f1fc(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_07f1fc(); +kernel void compute_main(device uint4* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapX_07f1fc(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/literal/quadSwapX/150d6f.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapX/150d6f.wgsl.expected.msl index 6bcdedee59b..605c3f5652a 100644 --- a/test/tint/builtins/gen/literal/quadSwapX/150d6f.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapX/150d6f.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec3; +float3 tint_msl_quadSwapX(float3 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_150d6f() -> vec3 { - var res : vec3 = quadSwapX(vec3(1.0f)); +float3 quadSwapX_150d6f(thread tint_private_vars_struct* const tint_private_vars) { + float3 res = tint_msl_quadSwapX(float3(1.0f), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_150d6f(); +fragment void fragment_main(device packed_float3* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = packed_float3(quadSwapX_150d6f(&(tint_private_vars))); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_150d6f(); +kernel void compute_main(device packed_float3* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = packed_float3(quadSwapX_150d6f(&(tint_private_vars))); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/literal/quadSwapX/19f8ce.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapX/19f8ce.wgsl.expected.msl index e8e6ccbae84..7b2d862c0cb 100644 --- a/test/tint/builtins/gen/literal/quadSwapX/19f8ce.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapX/19f8ce.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec2; +uint2 tint_msl_quadSwapX(uint2 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_19f8ce() -> vec2 { - var res : vec2 = quadSwapX(vec2(1u)); +uint2 quadSwapX_19f8ce(thread tint_private_vars_struct* const tint_private_vars) { + uint2 res = tint_msl_quadSwapX(uint2(1u), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_19f8ce(); +fragment void fragment_main(device uint2* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapX_19f8ce(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_19f8ce(); +kernel void compute_main(device uint2* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapX_19f8ce(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/literal/quadSwapX/1e1086.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapX/1e1086.wgsl.expected.msl index c270f4bda37..16c8924bf68 100644 --- a/test/tint/builtins/gen/literal/quadSwapX/1e1086.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapX/1e1086.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : i32; +int tint_msl_quadSwapX(int e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_1e1086() -> i32 { - var res : i32 = quadSwapX(1i); +int quadSwapX_1e1086(thread tint_private_vars_struct* const tint_private_vars) { + int res = tint_msl_quadSwapX(1, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_1e1086(); +fragment void fragment_main(device int* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapX_1e1086(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_1e1086(); +kernel void compute_main(device int* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapX_1e1086(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/literal/quadSwapX/69af6a.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapX/69af6a.wgsl.expected.msl index 30a122b84f5..91b4d0c5c4a 100644 --- a/test/tint/builtins/gen/literal/quadSwapX/69af6a.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapX/69af6a.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec4; +float4 tint_msl_quadSwapX(float4 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_69af6a() -> vec4 { - var res : vec4 = quadSwapX(vec4(1.0f)); +float4 quadSwapX_69af6a(thread tint_private_vars_struct* const tint_private_vars) { + float4 res = tint_msl_quadSwapX(float4(1.0f), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_69af6a(); +fragment void fragment_main(device float4* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapX_69af6a(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_69af6a(); +kernel void compute_main(device float4* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapX_69af6a(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/literal/quadSwapX/8203ad.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapX/8203ad.wgsl.expected.msl index 0a862c5ce5a..58cff8e5c4d 100644 --- a/test/tint/builtins/gen/literal/quadSwapX/8203ad.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapX/8203ad.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : u32; +uint tint_msl_quadSwapX(uint e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_8203ad() -> u32 { - var res : u32 = quadSwapX(1u); +uint quadSwapX_8203ad(thread tint_private_vars_struct* const tint_private_vars) { + uint res = tint_msl_quadSwapX(1u, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_8203ad(); +fragment void fragment_main(device uint* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapX_8203ad(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_8203ad(); +kernel void compute_main(device uint* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapX_8203ad(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/literal/quadSwapX/879738.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapX/879738.wgsl.expected.msl index dd948dedc8b..6a1c8c47332 100644 --- a/test/tint/builtins/gen/literal/quadSwapX/879738.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapX/879738.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec2; +float2 tint_msl_quadSwapX(float2 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_879738() -> vec2 { - var res : vec2 = quadSwapX(vec2(1.0f)); +float2 quadSwapX_879738(thread tint_private_vars_struct* const tint_private_vars) { + float2 res = tint_msl_quadSwapX(float2(1.0f), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_879738(); +fragment void fragment_main(device float2* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapX_879738(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_879738(); +kernel void compute_main(device float2* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapX_879738(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/literal/quadSwapX/9bea80.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapX/9bea80.wgsl.expected.msl index 54d97832fa7..7fe2a6440c0 100644 --- a/test/tint/builtins/gen/literal/quadSwapX/9bea80.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapX/9bea80.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : f32; +float tint_msl_quadSwapX(float e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_9bea80() -> f32 { - var res : f32 = quadSwapX(1.0f); +float quadSwapX_9bea80(thread tint_private_vars_struct* const tint_private_vars) { + float res = tint_msl_quadSwapX(1.0f, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_9bea80(); +fragment void fragment_main(device float* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapX_9bea80(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_9bea80(); +kernel void compute_main(device float* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapX_9bea80(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/literal/quadSwapX/a4e103.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapX/a4e103.wgsl.expected.msl index a16bb6f5909..bd1f64dc213 100644 --- a/test/tint/builtins/gen/literal/quadSwapX/a4e103.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapX/a4e103.wgsl.expected.msl @@ -1,25 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -enable subgroups_f16; -enable f16; -@group(0) @binding(0) var prevent_dce : f16; +half tint_msl_quadSwapX(half e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_a4e103() -> f16 { - var res : f16 = quadSwapX(1.0h); +half quadSwapX_a4e103(thread tint_private_vars_struct* const tint_private_vars) { + half res = tint_msl_quadSwapX(1.0h, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_a4e103(); +fragment void fragment_main(device half* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapX_a4e103(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_a4e103(); +kernel void compute_main(device half* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapX_a4e103(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/literal/quadSwapX/b1a5fe.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapX/b1a5fe.wgsl.expected.msl index a28b3957796..3aa30d921b7 100644 --- a/test/tint/builtins/gen/literal/quadSwapX/b1a5fe.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapX/b1a5fe.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec2; +int2 tint_msl_quadSwapX(int2 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_b1a5fe() -> vec2 { - var res : vec2 = quadSwapX(vec2(1i)); +int2 quadSwapX_b1a5fe(thread tint_private_vars_struct* const tint_private_vars) { + int2 res = tint_msl_quadSwapX(int2(1), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_b1a5fe(); +fragment void fragment_main(device int2* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapX_b1a5fe(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_b1a5fe(); +kernel void compute_main(device int2* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapX_b1a5fe(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/literal/quadSwapX/bc2013.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapX/bc2013.wgsl.expected.msl index 6fd1aaee305..ba281763ffd 100644 --- a/test/tint/builtins/gen/literal/quadSwapX/bc2013.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapX/bc2013.wgsl.expected.msl @@ -1,25 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -enable subgroups_f16; -enable f16; -@group(0) @binding(0) var prevent_dce : vec3; +half3 tint_msl_quadSwapX(half3 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_bc2013() -> vec3 { - var res : vec3 = quadSwapX(vec3(1.0h)); +half3 quadSwapX_bc2013(thread tint_private_vars_struct* const tint_private_vars) { + half3 res = tint_msl_quadSwapX(half3(1.0h), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_bc2013(); +fragment void fragment_main(device packed_half3* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = packed_half3(quadSwapX_bc2013(&(tint_private_vars))); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_bc2013(); +kernel void compute_main(device packed_half3* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = packed_half3(quadSwapX_bc2013(&(tint_private_vars))); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/literal/quadSwapX/bddb9f.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapX/bddb9f.wgsl.expected.msl index 8873d125a80..dd1bb129fc6 100644 --- a/test/tint/builtins/gen/literal/quadSwapX/bddb9f.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapX/bddb9f.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec3; +uint3 tint_msl_quadSwapX(uint3 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_bddb9f() -> vec3 { - var res : vec3 = quadSwapX(vec3(1u)); +uint3 quadSwapX_bddb9f(thread tint_private_vars_struct* const tint_private_vars) { + uint3 res = tint_msl_quadSwapX(uint3(1u), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_bddb9f(); +fragment void fragment_main(device packed_uint3* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = packed_uint3(quadSwapX_bddb9f(&(tint_private_vars))); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_bddb9f(); +kernel void compute_main(device packed_uint3* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = packed_uint3(quadSwapX_bddb9f(&(tint_private_vars))); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/literal/quadSwapX/d60cec.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapX/d60cec.wgsl.expected.msl index c26fac8241b..8f066e2f400 100644 --- a/test/tint/builtins/gen/literal/quadSwapX/d60cec.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapX/d60cec.wgsl.expected.msl @@ -1,25 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -enable subgroups_f16; -enable f16; -@group(0) @binding(0) var prevent_dce : vec2; +half2 tint_msl_quadSwapX(half2 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_d60cec() -> vec2 { - var res : vec2 = quadSwapX(vec2(1.0h)); +half2 quadSwapX_d60cec(thread tint_private_vars_struct* const tint_private_vars) { + half2 res = tint_msl_quadSwapX(half2(1.0h), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_d60cec(); +fragment void fragment_main(device half2* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapX_d60cec(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_d60cec(); +kernel void compute_main(device half2* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapX_d60cec(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/literal/quadSwapX/edfa1f.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapX/edfa1f.wgsl.expected.msl index e0903978606..423ea938fc3 100644 --- a/test/tint/builtins/gen/literal/quadSwapX/edfa1f.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapX/edfa1f.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec4; +int4 tint_msl_quadSwapX(int4 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_edfa1f() -> vec4 { - var res : vec4 = quadSwapX(vec4(1i)); +int4 quadSwapX_edfa1f(thread tint_private_vars_struct* const tint_private_vars) { + int4 res = tint_msl_quadSwapX(int4(1), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_edfa1f(); +fragment void fragment_main(device int4* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapX_edfa1f(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_edfa1f(); +kernel void compute_main(device int4* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapX_edfa1f(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/literal/quadSwapY/06a67c.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapY/06a67c.wgsl.expected.msl index 0c46bd75c74..6237c44a3f1 100644 --- a/test/tint/builtins/gen/literal/quadSwapY/06a67c.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapY/06a67c.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec3; +uint3 tint_msl_quadSwapY(uint3 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_06a67c() -> vec3 { - var res : vec3 = quadSwapY(vec3(1u)); +uint3 quadSwapY_06a67c(thread tint_private_vars_struct* const tint_private_vars) { + uint3 res = tint_msl_quadSwapY(uint3(1u), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_06a67c(); +fragment void fragment_main(device packed_uint3* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = packed_uint3(quadSwapY_06a67c(&(tint_private_vars))); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_06a67c(); +kernel void compute_main(device packed_uint3* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = packed_uint3(quadSwapY_06a67c(&(tint_private_vars))); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/literal/quadSwapY/0c4938.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapY/0c4938.wgsl.expected.msl index 94e9bb60944..3e83adacc23 100644 --- a/test/tint/builtins/gen/literal/quadSwapY/0c4938.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapY/0c4938.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : u32; +uint tint_msl_quadSwapY(uint e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_0c4938() -> u32 { - var res : u32 = quadSwapY(1u); +uint quadSwapY_0c4938(thread tint_private_vars_struct* const tint_private_vars) { + uint res = tint_msl_quadSwapY(1u, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_0c4938(); +fragment void fragment_main(device uint* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapY_0c4938(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_0c4938(); +kernel void compute_main(device uint* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapY_0c4938(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/literal/quadSwapY/0d05a8.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapY/0d05a8.wgsl.expected.msl index 43f99dfa138..64b13e95fc7 100644 --- a/test/tint/builtins/gen/literal/quadSwapY/0d05a8.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapY/0d05a8.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec2; +int2 tint_msl_quadSwapY(int2 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_0d05a8() -> vec2 { - var res : vec2 = quadSwapY(vec2(1i)); +int2 quadSwapY_0d05a8(thread tint_private_vars_struct* const tint_private_vars) { + int2 res = tint_msl_quadSwapY(int2(1), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_0d05a8(); +fragment void fragment_main(device int2* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapY_0d05a8(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_0d05a8(); +kernel void compute_main(device int2* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapY_0d05a8(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/literal/quadSwapY/14bb9a.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapY/14bb9a.wgsl.expected.msl index bd05a671bbd..6f4905dfd7c 100644 --- a/test/tint/builtins/gen/literal/quadSwapY/14bb9a.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapY/14bb9a.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec4; +int4 tint_msl_quadSwapY(int4 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_14bb9a() -> vec4 { - var res : vec4 = quadSwapY(vec4(1i)); +int4 quadSwapY_14bb9a(thread tint_private_vars_struct* const tint_private_vars) { + int4 res = tint_msl_quadSwapY(int4(1), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_14bb9a(); +fragment void fragment_main(device int4* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapY_14bb9a(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_14bb9a(); +kernel void compute_main(device int4* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapY_14bb9a(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/literal/quadSwapY/1f1a06.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapY/1f1a06.wgsl.expected.msl index f7b0870d715..a97a0613b74 100644 --- a/test/tint/builtins/gen/literal/quadSwapY/1f1a06.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapY/1f1a06.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec2; +float2 tint_msl_quadSwapY(float2 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_1f1a06() -> vec2 { - var res : vec2 = quadSwapY(vec2(1.0f)); +float2 quadSwapY_1f1a06(thread tint_private_vars_struct* const tint_private_vars) { + float2 res = tint_msl_quadSwapY(float2(1.0f), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_1f1a06(); +fragment void fragment_main(device float2* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapY_1f1a06(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_1f1a06(); +kernel void compute_main(device float2* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapY_1f1a06(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/literal/quadSwapY/264908.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapY/264908.wgsl.expected.msl index 09f7425c031..f9ca6b43e3d 100644 --- a/test/tint/builtins/gen/literal/quadSwapY/264908.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapY/264908.wgsl.expected.msl @@ -1,25 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -enable subgroups_f16; -enable f16; -@group(0) @binding(0) var prevent_dce : vec3; +half3 tint_msl_quadSwapY(half3 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_264908() -> vec3 { - var res : vec3 = quadSwapY(vec3(1.0h)); +half3 quadSwapY_264908(thread tint_private_vars_struct* const tint_private_vars) { + half3 res = tint_msl_quadSwapY(half3(1.0h), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_264908(); +fragment void fragment_main(device packed_half3* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = packed_half3(quadSwapY_264908(&(tint_private_vars))); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_264908(); +kernel void compute_main(device packed_half3* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = packed_half3(quadSwapY_264908(&(tint_private_vars))); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/literal/quadSwapY/5b2e67.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapY/5b2e67.wgsl.expected.msl index f0c2a075516..a170dcd7725 100644 --- a/test/tint/builtins/gen/literal/quadSwapY/5b2e67.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapY/5b2e67.wgsl.expected.msl @@ -1,25 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -enable subgroups_f16; -enable f16; -@group(0) @binding(0) var prevent_dce : vec4; +half4 tint_msl_quadSwapY(half4 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_5b2e67() -> vec4 { - var res : vec4 = quadSwapY(vec4(1.0h)); +half4 quadSwapY_5b2e67(thread tint_private_vars_struct* const tint_private_vars) { + half4 res = tint_msl_quadSwapY(half4(1.0h), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_5b2e67(); +fragment void fragment_main(device half4* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapY_5b2e67(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_5b2e67(); +kernel void compute_main(device half4* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapY_5b2e67(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/literal/quadSwapY/6f6bc9.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapY/6f6bc9.wgsl.expected.msl index c5bdbd58e3e..3e24c0183c2 100644 --- a/test/tint/builtins/gen/literal/quadSwapY/6f6bc9.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapY/6f6bc9.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : f32; +float tint_msl_quadSwapY(float e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_6f6bc9() -> f32 { - var res : f32 = quadSwapY(1.0f); +float quadSwapY_6f6bc9(thread tint_private_vars_struct* const tint_private_vars) { + float res = tint_msl_quadSwapY(1.0f, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_6f6bc9(); +fragment void fragment_main(device float* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapY_6f6bc9(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_6f6bc9(); +kernel void compute_main(device float* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapY_6f6bc9(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/literal/quadSwapY/9277e9.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapY/9277e9.wgsl.expected.msl index 05607cc19df..afefef8bdd4 100644 --- a/test/tint/builtins/gen/literal/quadSwapY/9277e9.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapY/9277e9.wgsl.expected.msl @@ -1,25 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -enable subgroups_f16; -enable f16; -@group(0) @binding(0) var prevent_dce : f16; +half tint_msl_quadSwapY(half e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_9277e9() -> f16 { - var res : f16 = quadSwapY(1.0h); +half quadSwapY_9277e9(thread tint_private_vars_struct* const tint_private_vars) { + half res = tint_msl_quadSwapY(1.0h, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_9277e9(); +fragment void fragment_main(device half* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapY_9277e9(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_9277e9(); +kernel void compute_main(device half* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapY_9277e9(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/literal/quadSwapY/94ab6d.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapY/94ab6d.wgsl.expected.msl index 7089b1adc36..8a55dfa1b07 100644 --- a/test/tint/builtins/gen/literal/quadSwapY/94ab6d.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapY/94ab6d.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : i32; +int tint_msl_quadSwapY(int e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_94ab6d() -> i32 { - var res : i32 = quadSwapY(1i); +int quadSwapY_94ab6d(thread tint_private_vars_struct* const tint_private_vars) { + int res = tint_msl_quadSwapY(1, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_94ab6d(); +fragment void fragment_main(device int* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapY_94ab6d(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_94ab6d(); +kernel void compute_main(device int* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapY_94ab6d(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/literal/quadSwapY/a27e1c.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapY/a27e1c.wgsl.expected.msl index 74ca974f476..f515fc8851f 100644 --- a/test/tint/builtins/gen/literal/quadSwapY/a27e1c.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapY/a27e1c.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec2; +uint2 tint_msl_quadSwapY(uint2 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_a27e1c() -> vec2 { - var res : vec2 = quadSwapY(vec2(1u)); +uint2 quadSwapY_a27e1c(thread tint_private_vars_struct* const tint_private_vars) { + uint2 res = tint_msl_quadSwapY(uint2(1u), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_a27e1c(); +fragment void fragment_main(device uint2* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapY_a27e1c(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_a27e1c(); +kernel void compute_main(device uint2* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapY_a27e1c(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/literal/quadSwapY/a50fcb.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapY/a50fcb.wgsl.expected.msl index 2adf640195d..4ead5aaa2af 100644 --- a/test/tint/builtins/gen/literal/quadSwapY/a50fcb.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapY/a50fcb.wgsl.expected.msl @@ -1,25 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -enable subgroups_f16; -enable f16; -@group(0) @binding(0) var prevent_dce : vec2; +half2 tint_msl_quadSwapY(half2 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_a50fcb() -> vec2 { - var res : vec2 = quadSwapY(vec2(1.0h)); +half2 quadSwapY_a50fcb(thread tint_private_vars_struct* const tint_private_vars) { + half2 res = tint_msl_quadSwapY(half2(1.0h), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_a50fcb(); +fragment void fragment_main(device half2* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapY_a50fcb(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_a50fcb(); +kernel void compute_main(device half2* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapY_a50fcb(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/literal/quadSwapY/b9d9e7.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapY/b9d9e7.wgsl.expected.msl index 21bb21992c6..8c523a795a2 100644 --- a/test/tint/builtins/gen/literal/quadSwapY/b9d9e7.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapY/b9d9e7.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec4; +float4 tint_msl_quadSwapY(float4 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_b9d9e7() -> vec4 { - var res : vec4 = quadSwapY(vec4(1.0f)); +float4 quadSwapY_b9d9e7(thread tint_private_vars_struct* const tint_private_vars) { + float4 res = tint_msl_quadSwapY(float4(1.0f), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_b9d9e7(); +fragment void fragment_main(device float4* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapY_b9d9e7(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_b9d9e7(); +kernel void compute_main(device float4* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapY_b9d9e7(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/literal/quadSwapY/bb697b.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapY/bb697b.wgsl.expected.msl index 642394044e8..c96b55c4668 100644 --- a/test/tint/builtins/gen/literal/quadSwapY/bb697b.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapY/bb697b.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec4; +uint4 tint_msl_quadSwapY(uint4 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_bb697b() -> vec4 { - var res : vec4 = quadSwapY(vec4(1u)); +uint4 quadSwapY_bb697b(thread tint_private_vars_struct* const tint_private_vars) { + uint4 res = tint_msl_quadSwapY(uint4(1u), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_bb697b(); +fragment void fragment_main(device uint4* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapY_bb697b(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_bb697b(); +kernel void compute_main(device uint4* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapY_bb697b(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/literal/quadSwapY/be4e72.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapY/be4e72.wgsl.expected.msl index d048f225812..c83dc8debf7 100644 --- a/test/tint/builtins/gen/literal/quadSwapY/be4e72.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapY/be4e72.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec3; +int3 tint_msl_quadSwapY(int3 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_be4e72() -> vec3 { - var res : vec3 = quadSwapY(vec3(1i)); +int3 quadSwapY_be4e72(thread tint_private_vars_struct* const tint_private_vars) { + int3 res = tint_msl_quadSwapY(int3(1), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_be4e72(); +fragment void fragment_main(device packed_int3* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = packed_int3(quadSwapY_be4e72(&(tint_private_vars))); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_be4e72(); +kernel void compute_main(device packed_int3* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = packed_int3(quadSwapY_be4e72(&(tint_private_vars))); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/literal/quadSwapY/d1ab4d.wgsl.expected.msl b/test/tint/builtins/gen/literal/quadSwapY/d1ab4d.wgsl.expected.msl index af21f34dc9e..0bb656bce74 100644 --- a/test/tint/builtins/gen/literal/quadSwapY/d1ab4d.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/quadSwapY/d1ab4d.wgsl.expected.msl @@ -1,23 +1,35 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec3; +float3 tint_msl_quadSwapY(float3 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_d1ab4d() -> vec3 { - var res : vec3 = quadSwapY(vec3(1.0f)); +float3 quadSwapY_d1ab4d(thread tint_private_vars_struct* const tint_private_vars) { + float3 res = tint_msl_quadSwapY(float3(1.0f), tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_d1ab4d(); +fragment void fragment_main(device packed_float3* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = packed_float3(quadSwapY_d1ab4d(&(tint_private_vars))); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_d1ab4d(); +kernel void compute_main(device packed_float3* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = packed_float3(quadSwapY_d1ab4d(&(tint_private_vars))); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/15ac75.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapDiagonal/15ac75.wgsl.expected.msl index b58667621ce..0470ad153e4 100644 --- a/test/tint/builtins/gen/var/quadSwapDiagonal/15ac75.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapDiagonal/15ac75.wgsl.expected.msl @@ -1,26 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -enable subgroups_f16; -enable f16; -@group(0) @binding(0) var prevent_dce : vec2; +half2 tint_msl_quadSwapDiagonal(half2 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_15ac75() -> vec2 { - var arg_0 = vec2(1.0h); - var res : vec2 = quadSwapDiagonal(arg_0); +half2 quadSwapDiagonal_15ac75(thread tint_private_vars_struct* const tint_private_vars) { + half2 arg_0 = half2(1.0h); + half2 res = tint_msl_quadSwapDiagonal(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_15ac75(); +fragment void fragment_main(device half2* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapDiagonal_15ac75(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_15ac75(); +kernel void compute_main(device half2* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapDiagonal_15ac75(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/2be5e7.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapDiagonal/2be5e7.wgsl.expected.msl index caff0895a1a..5fddd3000f1 100644 --- a/test/tint/builtins/gen/var/quadSwapDiagonal/2be5e7.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapDiagonal/2be5e7.wgsl.expected.msl @@ -1,26 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -enable subgroups_f16; -enable f16; -@group(0) @binding(0) var prevent_dce : f16; +half tint_msl_quadSwapDiagonal(half e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_2be5e7() -> f16 { - var arg_0 = 1.0h; - var res : f16 = quadSwapDiagonal(arg_0); +half quadSwapDiagonal_2be5e7(thread tint_private_vars_struct* const tint_private_vars) { + half arg_0 = 1.0h; + half res = tint_msl_quadSwapDiagonal(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_2be5e7(); +fragment void fragment_main(device half* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapDiagonal_2be5e7(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_2be5e7(); +kernel void compute_main(device half* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapDiagonal_2be5e7(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/331804.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapDiagonal/331804.wgsl.expected.msl index f6d82359b5d..da850035761 100644 --- a/test/tint/builtins/gen/var/quadSwapDiagonal/331804.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapDiagonal/331804.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec4; +float4 tint_msl_quadSwapDiagonal(float4 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_331804() -> vec4 { - var arg_0 = vec4(1.0f); - var res : vec4 = quadSwapDiagonal(arg_0); +float4 quadSwapDiagonal_331804(thread tint_private_vars_struct* const tint_private_vars) { + float4 arg_0 = float4(1.0f); + float4 res = tint_msl_quadSwapDiagonal(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_331804(); +fragment void fragment_main(device float4* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapDiagonal_331804(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_331804(); +kernel void compute_main(device float4* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapDiagonal_331804(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/348173.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapDiagonal/348173.wgsl.expected.msl index 0c54c461144..301456346a6 100644 --- a/test/tint/builtins/gen/var/quadSwapDiagonal/348173.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapDiagonal/348173.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec2; +uint2 tint_msl_quadSwapDiagonal(uint2 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_348173() -> vec2 { - var arg_0 = vec2(1u); - var res : vec2 = quadSwapDiagonal(arg_0); +uint2 quadSwapDiagonal_348173(thread tint_private_vars_struct* const tint_private_vars) { + uint2 arg_0 = uint2(1u); + uint2 res = tint_msl_quadSwapDiagonal(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_348173(); +fragment void fragment_main(device uint2* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapDiagonal_348173(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_348173(); +kernel void compute_main(device uint2* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapDiagonal_348173(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/486196.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapDiagonal/486196.wgsl.expected.msl index 2b6bc331ddc..f0b8efff5be 100644 --- a/test/tint/builtins/gen/var/quadSwapDiagonal/486196.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapDiagonal/486196.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : f32; +float tint_msl_quadSwapDiagonal(float e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_486196() -> f32 { - var arg_0 = 1.0f; - var res : f32 = quadSwapDiagonal(arg_0); +float quadSwapDiagonal_486196(thread tint_private_vars_struct* const tint_private_vars) { + float arg_0 = 1.0f; + float res = tint_msl_quadSwapDiagonal(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_486196(); +fragment void fragment_main(device float* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapDiagonal_486196(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_486196(); +kernel void compute_main(device float* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapDiagonal_486196(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/730e40.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapDiagonal/730e40.wgsl.expected.msl index 99a562ad6a1..c35dd152c98 100644 --- a/test/tint/builtins/gen/var/quadSwapDiagonal/730e40.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapDiagonal/730e40.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : u32; +uint tint_msl_quadSwapDiagonal(uint e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_730e40() -> u32 { - var arg_0 = 1u; - var res : u32 = quadSwapDiagonal(arg_0); +uint quadSwapDiagonal_730e40(thread tint_private_vars_struct* const tint_private_vars) { + uint arg_0 = 1u; + uint res = tint_msl_quadSwapDiagonal(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_730e40(); +fragment void fragment_main(device uint* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapDiagonal_730e40(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_730e40(); +kernel void compute_main(device uint* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapDiagonal_730e40(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/8077c8.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapDiagonal/8077c8.wgsl.expected.msl index a4bb005e5af..be2113fbd16 100644 --- a/test/tint/builtins/gen/var/quadSwapDiagonal/8077c8.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapDiagonal/8077c8.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec2; +float2 tint_msl_quadSwapDiagonal(float2 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_8077c8() -> vec2 { - var arg_0 = vec2(1.0f); - var res : vec2 = quadSwapDiagonal(arg_0); +float2 quadSwapDiagonal_8077c8(thread tint_private_vars_struct* const tint_private_vars) { + float2 arg_0 = float2(1.0f); + float2 res = tint_msl_quadSwapDiagonal(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_8077c8(); +fragment void fragment_main(device float2* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapDiagonal_8077c8(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_8077c8(); +kernel void compute_main(device float2* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapDiagonal_8077c8(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/856536.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapDiagonal/856536.wgsl.expected.msl index 12966e17b67..fbc0f595634 100644 --- a/test/tint/builtins/gen/var/quadSwapDiagonal/856536.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapDiagonal/856536.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec3; +uint3 tint_msl_quadSwapDiagonal(uint3 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_856536() -> vec3 { - var arg_0 = vec3(1u); - var res : vec3 = quadSwapDiagonal(arg_0); +uint3 quadSwapDiagonal_856536(thread tint_private_vars_struct* const tint_private_vars) { + uint3 arg_0 = uint3(1u); + uint3 res = tint_msl_quadSwapDiagonal(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_856536(); +fragment void fragment_main(device packed_uint3* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = packed_uint3(quadSwapDiagonal_856536(&(tint_private_vars))); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_856536(); +kernel void compute_main(device packed_uint3* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = packed_uint3(quadSwapDiagonal_856536(&(tint_private_vars))); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/9ccb38.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapDiagonal/9ccb38.wgsl.expected.msl index d4e90a732ab..e79e74040a5 100644 --- a/test/tint/builtins/gen/var/quadSwapDiagonal/9ccb38.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapDiagonal/9ccb38.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : i32; +int tint_msl_quadSwapDiagonal(int e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_9ccb38() -> i32 { - var arg_0 = 1i; - var res : i32 = quadSwapDiagonal(arg_0); +int quadSwapDiagonal_9ccb38(thread tint_private_vars_struct* const tint_private_vars) { + int arg_0 = 1; + int res = tint_msl_quadSwapDiagonal(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_9ccb38(); +fragment void fragment_main(device int* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapDiagonal_9ccb38(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_9ccb38(); +kernel void compute_main(device int* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapDiagonal_9ccb38(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/a090b0.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapDiagonal/a090b0.wgsl.expected.msl index f43fdc7e4fa..bef17982db7 100644 --- a/test/tint/builtins/gen/var/quadSwapDiagonal/a090b0.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapDiagonal/a090b0.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec2; +int2 tint_msl_quadSwapDiagonal(int2 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_a090b0() -> vec2 { - var arg_0 = vec2(1i); - var res : vec2 = quadSwapDiagonal(arg_0); +int2 quadSwapDiagonal_a090b0(thread tint_private_vars_struct* const tint_private_vars) { + int2 arg_0 = int2(1); + int2 res = tint_msl_quadSwapDiagonal(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_a090b0(); +fragment void fragment_main(device int2* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapDiagonal_a090b0(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_a090b0(); +kernel void compute_main(device int2* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapDiagonal_a090b0(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/a665b1.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapDiagonal/a665b1.wgsl.expected.msl index 6faeab04f0e..0ad15d97e5d 100644 --- a/test/tint/builtins/gen/var/quadSwapDiagonal/a665b1.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapDiagonal/a665b1.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec4; +int4 tint_msl_quadSwapDiagonal(int4 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_a665b1() -> vec4 { - var arg_0 = vec4(1i); - var res : vec4 = quadSwapDiagonal(arg_0); +int4 quadSwapDiagonal_a665b1(thread tint_private_vars_struct* const tint_private_vars) { + int4 arg_0 = int4(1); + int4 res = tint_msl_quadSwapDiagonal(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_a665b1(); +fragment void fragment_main(device int4* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapDiagonal_a665b1(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_a665b1(); +kernel void compute_main(device int4* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapDiagonal_a665b1(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/a82e1d.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapDiagonal/a82e1d.wgsl.expected.msl index f3764a1e33e..16daf07e580 100644 --- a/test/tint/builtins/gen/var/quadSwapDiagonal/a82e1d.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapDiagonal/a82e1d.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec3; +int3 tint_msl_quadSwapDiagonal(int3 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_a82e1d() -> vec3 { - var arg_0 = vec3(1i); - var res : vec3 = quadSwapDiagonal(arg_0); +int3 quadSwapDiagonal_a82e1d(thread tint_private_vars_struct* const tint_private_vars) { + int3 arg_0 = int3(1); + int3 res = tint_msl_quadSwapDiagonal(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_a82e1d(); +fragment void fragment_main(device packed_int3* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = packed_int3(quadSwapDiagonal_a82e1d(&(tint_private_vars))); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_a82e1d(); +kernel void compute_main(device packed_int3* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = packed_int3(quadSwapDiagonal_a82e1d(&(tint_private_vars))); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/af19a5.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapDiagonal/af19a5.wgsl.expected.msl index 031f680b877..9079d678a40 100644 --- a/test/tint/builtins/gen/var/quadSwapDiagonal/af19a5.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapDiagonal/af19a5.wgsl.expected.msl @@ -1,26 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -enable subgroups_f16; -enable f16; -@group(0) @binding(0) var prevent_dce : vec4; +half4 tint_msl_quadSwapDiagonal(half4 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_af19a5() -> vec4 { - var arg_0 = vec4(1.0h); - var res : vec4 = quadSwapDiagonal(arg_0); +half4 quadSwapDiagonal_af19a5(thread tint_private_vars_struct* const tint_private_vars) { + half4 arg_0 = half4(1.0h); + half4 res = tint_msl_quadSwapDiagonal(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_af19a5(); +fragment void fragment_main(device half4* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapDiagonal_af19a5(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_af19a5(); +kernel void compute_main(device half4* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapDiagonal_af19a5(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/b905fc.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapDiagonal/b905fc.wgsl.expected.msl index c26ec22f22d..759fd096c32 100644 --- a/test/tint/builtins/gen/var/quadSwapDiagonal/b905fc.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapDiagonal/b905fc.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec3; +float3 tint_msl_quadSwapDiagonal(float3 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_b905fc() -> vec3 { - var arg_0 = vec3(1.0f); - var res : vec3 = quadSwapDiagonal(arg_0); +float3 quadSwapDiagonal_b905fc(thread tint_private_vars_struct* const tint_private_vars) { + float3 arg_0 = float3(1.0f); + float3 res = tint_msl_quadSwapDiagonal(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_b905fc(); +fragment void fragment_main(device packed_float3* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = packed_float3(quadSwapDiagonal_b905fc(&(tint_private_vars))); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_b905fc(); +kernel void compute_main(device packed_float3* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = packed_float3(quadSwapDiagonal_b905fc(&(tint_private_vars))); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/c31636.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapDiagonal/c31636.wgsl.expected.msl index 58dc97cc44e..80017e18372 100644 --- a/test/tint/builtins/gen/var/quadSwapDiagonal/c31636.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapDiagonal/c31636.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec4; +uint4 tint_msl_quadSwapDiagonal(uint4 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_c31636() -> vec4 { - var arg_0 = vec4(1u); - var res : vec4 = quadSwapDiagonal(arg_0); +uint4 quadSwapDiagonal_c31636(thread tint_private_vars_struct* const tint_private_vars) { + uint4 arg_0 = uint4(1u); + uint4 res = tint_msl_quadSwapDiagonal(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_c31636(); +fragment void fragment_main(device uint4* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapDiagonal_c31636(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_c31636(); +kernel void compute_main(device uint4* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapDiagonal_c31636(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/e4bec8.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapDiagonal/e4bec8.wgsl.expected.msl index 5ae38ad041a..ea185841a3c 100644 --- a/test/tint/builtins/gen/var/quadSwapDiagonal/e4bec8.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapDiagonal/e4bec8.wgsl.expected.msl @@ -1,26 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -enable subgroups_f16; -enable f16; -@group(0) @binding(0) var prevent_dce : vec3; +half3 tint_msl_quadSwapDiagonal(half3 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u)); +} -fn quadSwapDiagonal_e4bec8() -> vec3 { - var arg_0 = vec3(1.0h); - var res : vec3 = quadSwapDiagonal(arg_0); +half3 quadSwapDiagonal_e4bec8(thread tint_private_vars_struct* const tint_private_vars) { + half3 arg_0 = half3(1.0h); + half3 res = tint_msl_quadSwapDiagonal(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapDiagonal_e4bec8(); +fragment void fragment_main(device packed_half3* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = packed_half3(quadSwapDiagonal_e4bec8(&(tint_private_vars))); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapDiagonal_e4bec8(); +kernel void compute_main(device packed_half3* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = packed_half3(quadSwapDiagonal_e4bec8(&(tint_private_vars))); + return; } -Failed to generate: error: Unknown import method: quadSwapDiagonal diff --git a/test/tint/builtins/gen/var/quadSwapX/02834c.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapX/02834c.wgsl.expected.msl index 2a0e22d4f62..3d9aba3192e 100644 --- a/test/tint/builtins/gen/var/quadSwapX/02834c.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapX/02834c.wgsl.expected.msl @@ -1,26 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -enable subgroups_f16; -enable f16; -@group(0) @binding(0) var prevent_dce : vec4; +half4 tint_msl_quadSwapX(half4 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_02834c() -> vec4 { - var arg_0 = vec4(1.0h); - var res : vec4 = quadSwapX(arg_0); +half4 quadSwapX_02834c(thread tint_private_vars_struct* const tint_private_vars) { + half4 arg_0 = half4(1.0h); + half4 res = tint_msl_quadSwapX(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_02834c(); +fragment void fragment_main(device half4* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapX_02834c(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_02834c(); +kernel void compute_main(device half4* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapX_02834c(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/var/quadSwapX/053f3b.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapX/053f3b.wgsl.expected.msl index c782b8d17a4..9999fa8b1a7 100644 --- a/test/tint/builtins/gen/var/quadSwapX/053f3b.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapX/053f3b.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec3; +int3 tint_msl_quadSwapX(int3 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_053f3b() -> vec3 { - var arg_0 = vec3(1i); - var res : vec3 = quadSwapX(arg_0); +int3 quadSwapX_053f3b(thread tint_private_vars_struct* const tint_private_vars) { + int3 arg_0 = int3(1); + int3 res = tint_msl_quadSwapX(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_053f3b(); +fragment void fragment_main(device packed_int3* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = packed_int3(quadSwapX_053f3b(&(tint_private_vars))); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_053f3b(); +kernel void compute_main(device packed_int3* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = packed_int3(quadSwapX_053f3b(&(tint_private_vars))); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/var/quadSwapX/07f1fc.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapX/07f1fc.wgsl.expected.msl index ae48f10f39c..6a654fff163 100644 --- a/test/tint/builtins/gen/var/quadSwapX/07f1fc.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapX/07f1fc.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec4; +uint4 tint_msl_quadSwapX(uint4 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_07f1fc() -> vec4 { - var arg_0 = vec4(1u); - var res : vec4 = quadSwapX(arg_0); +uint4 quadSwapX_07f1fc(thread tint_private_vars_struct* const tint_private_vars) { + uint4 arg_0 = uint4(1u); + uint4 res = tint_msl_quadSwapX(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_07f1fc(); +fragment void fragment_main(device uint4* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapX_07f1fc(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_07f1fc(); +kernel void compute_main(device uint4* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapX_07f1fc(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/var/quadSwapX/150d6f.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapX/150d6f.wgsl.expected.msl index c98ab2305df..d10d9349216 100644 --- a/test/tint/builtins/gen/var/quadSwapX/150d6f.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapX/150d6f.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec3; +float3 tint_msl_quadSwapX(float3 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_150d6f() -> vec3 { - var arg_0 = vec3(1.0f); - var res : vec3 = quadSwapX(arg_0); +float3 quadSwapX_150d6f(thread tint_private_vars_struct* const tint_private_vars) { + float3 arg_0 = float3(1.0f); + float3 res = tint_msl_quadSwapX(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_150d6f(); +fragment void fragment_main(device packed_float3* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = packed_float3(quadSwapX_150d6f(&(tint_private_vars))); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_150d6f(); +kernel void compute_main(device packed_float3* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = packed_float3(quadSwapX_150d6f(&(tint_private_vars))); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/var/quadSwapX/19f8ce.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapX/19f8ce.wgsl.expected.msl index 3028af2c51d..5eab471aaa6 100644 --- a/test/tint/builtins/gen/var/quadSwapX/19f8ce.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapX/19f8ce.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec2; +uint2 tint_msl_quadSwapX(uint2 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_19f8ce() -> vec2 { - var arg_0 = vec2(1u); - var res : vec2 = quadSwapX(arg_0); +uint2 quadSwapX_19f8ce(thread tint_private_vars_struct* const tint_private_vars) { + uint2 arg_0 = uint2(1u); + uint2 res = tint_msl_quadSwapX(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_19f8ce(); +fragment void fragment_main(device uint2* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapX_19f8ce(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_19f8ce(); +kernel void compute_main(device uint2* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapX_19f8ce(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/var/quadSwapX/1e1086.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapX/1e1086.wgsl.expected.msl index 301123e1320..82c1362ab66 100644 --- a/test/tint/builtins/gen/var/quadSwapX/1e1086.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapX/1e1086.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : i32; +int tint_msl_quadSwapX(int e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_1e1086() -> i32 { - var arg_0 = 1i; - var res : i32 = quadSwapX(arg_0); +int quadSwapX_1e1086(thread tint_private_vars_struct* const tint_private_vars) { + int arg_0 = 1; + int res = tint_msl_quadSwapX(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_1e1086(); +fragment void fragment_main(device int* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapX_1e1086(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_1e1086(); +kernel void compute_main(device int* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapX_1e1086(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/var/quadSwapX/69af6a.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapX/69af6a.wgsl.expected.msl index 4befc2ea170..4fa83948c01 100644 --- a/test/tint/builtins/gen/var/quadSwapX/69af6a.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapX/69af6a.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec4; +float4 tint_msl_quadSwapX(float4 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_69af6a() -> vec4 { - var arg_0 = vec4(1.0f); - var res : vec4 = quadSwapX(arg_0); +float4 quadSwapX_69af6a(thread tint_private_vars_struct* const tint_private_vars) { + float4 arg_0 = float4(1.0f); + float4 res = tint_msl_quadSwapX(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_69af6a(); +fragment void fragment_main(device float4* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapX_69af6a(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_69af6a(); +kernel void compute_main(device float4* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapX_69af6a(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/var/quadSwapX/8203ad.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapX/8203ad.wgsl.expected.msl index e13c58747de..fc5e07cd76d 100644 --- a/test/tint/builtins/gen/var/quadSwapX/8203ad.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapX/8203ad.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : u32; +uint tint_msl_quadSwapX(uint e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_8203ad() -> u32 { - var arg_0 = 1u; - var res : u32 = quadSwapX(arg_0); +uint quadSwapX_8203ad(thread tint_private_vars_struct* const tint_private_vars) { + uint arg_0 = 1u; + uint res = tint_msl_quadSwapX(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_8203ad(); +fragment void fragment_main(device uint* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapX_8203ad(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_8203ad(); +kernel void compute_main(device uint* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapX_8203ad(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/var/quadSwapX/879738.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapX/879738.wgsl.expected.msl index f6b6fb9cab7..9952d7d2e8b 100644 --- a/test/tint/builtins/gen/var/quadSwapX/879738.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapX/879738.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec2; +float2 tint_msl_quadSwapX(float2 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_879738() -> vec2 { - var arg_0 = vec2(1.0f); - var res : vec2 = quadSwapX(arg_0); +float2 quadSwapX_879738(thread tint_private_vars_struct* const tint_private_vars) { + float2 arg_0 = float2(1.0f); + float2 res = tint_msl_quadSwapX(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_879738(); +fragment void fragment_main(device float2* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapX_879738(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_879738(); +kernel void compute_main(device float2* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapX_879738(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/var/quadSwapX/9bea80.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapX/9bea80.wgsl.expected.msl index ce807fa2551..0c0d7d62366 100644 --- a/test/tint/builtins/gen/var/quadSwapX/9bea80.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapX/9bea80.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : f32; +float tint_msl_quadSwapX(float e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_9bea80() -> f32 { - var arg_0 = 1.0f; - var res : f32 = quadSwapX(arg_0); +float quadSwapX_9bea80(thread tint_private_vars_struct* const tint_private_vars) { + float arg_0 = 1.0f; + float res = tint_msl_quadSwapX(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_9bea80(); +fragment void fragment_main(device float* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapX_9bea80(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_9bea80(); +kernel void compute_main(device float* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapX_9bea80(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/var/quadSwapX/a4e103.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapX/a4e103.wgsl.expected.msl index b3c2afef5bc..2ac8c90597c 100644 --- a/test/tint/builtins/gen/var/quadSwapX/a4e103.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapX/a4e103.wgsl.expected.msl @@ -1,26 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -enable subgroups_f16; -enable f16; -@group(0) @binding(0) var prevent_dce : f16; +half tint_msl_quadSwapX(half e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_a4e103() -> f16 { - var arg_0 = 1.0h; - var res : f16 = quadSwapX(arg_0); +half quadSwapX_a4e103(thread tint_private_vars_struct* const tint_private_vars) { + half arg_0 = 1.0h; + half res = tint_msl_quadSwapX(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_a4e103(); +fragment void fragment_main(device half* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapX_a4e103(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_a4e103(); +kernel void compute_main(device half* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapX_a4e103(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/var/quadSwapX/b1a5fe.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapX/b1a5fe.wgsl.expected.msl index 44f175cc4dd..2259fb89053 100644 --- a/test/tint/builtins/gen/var/quadSwapX/b1a5fe.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapX/b1a5fe.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec2; +int2 tint_msl_quadSwapX(int2 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_b1a5fe() -> vec2 { - var arg_0 = vec2(1i); - var res : vec2 = quadSwapX(arg_0); +int2 quadSwapX_b1a5fe(thread tint_private_vars_struct* const tint_private_vars) { + int2 arg_0 = int2(1); + int2 res = tint_msl_quadSwapX(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_b1a5fe(); +fragment void fragment_main(device int2* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapX_b1a5fe(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_b1a5fe(); +kernel void compute_main(device int2* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapX_b1a5fe(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/var/quadSwapX/bc2013.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapX/bc2013.wgsl.expected.msl index 2a68f319e6e..167fde42a7e 100644 --- a/test/tint/builtins/gen/var/quadSwapX/bc2013.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapX/bc2013.wgsl.expected.msl @@ -1,26 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -enable subgroups_f16; -enable f16; -@group(0) @binding(0) var prevent_dce : vec3; +half3 tint_msl_quadSwapX(half3 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_bc2013() -> vec3 { - var arg_0 = vec3(1.0h); - var res : vec3 = quadSwapX(arg_0); +half3 quadSwapX_bc2013(thread tint_private_vars_struct* const tint_private_vars) { + half3 arg_0 = half3(1.0h); + half3 res = tint_msl_quadSwapX(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_bc2013(); +fragment void fragment_main(device packed_half3* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = packed_half3(quadSwapX_bc2013(&(tint_private_vars))); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_bc2013(); +kernel void compute_main(device packed_half3* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = packed_half3(quadSwapX_bc2013(&(tint_private_vars))); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/var/quadSwapX/bddb9f.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapX/bddb9f.wgsl.expected.msl index 50b5921d2ca..a324d484bba 100644 --- a/test/tint/builtins/gen/var/quadSwapX/bddb9f.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapX/bddb9f.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec3; +uint3 tint_msl_quadSwapX(uint3 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_bddb9f() -> vec3 { - var arg_0 = vec3(1u); - var res : vec3 = quadSwapX(arg_0); +uint3 quadSwapX_bddb9f(thread tint_private_vars_struct* const tint_private_vars) { + uint3 arg_0 = uint3(1u); + uint3 res = tint_msl_quadSwapX(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_bddb9f(); +fragment void fragment_main(device packed_uint3* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = packed_uint3(quadSwapX_bddb9f(&(tint_private_vars))); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_bddb9f(); +kernel void compute_main(device packed_uint3* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = packed_uint3(quadSwapX_bddb9f(&(tint_private_vars))); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/var/quadSwapX/d60cec.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapX/d60cec.wgsl.expected.msl index 4dcc1b096e7..6ef0a725351 100644 --- a/test/tint/builtins/gen/var/quadSwapX/d60cec.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapX/d60cec.wgsl.expected.msl @@ -1,26 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -enable subgroups_f16; -enable f16; -@group(0) @binding(0) var prevent_dce : vec2; +half2 tint_msl_quadSwapX(half2 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_d60cec() -> vec2 { - var arg_0 = vec2(1.0h); - var res : vec2 = quadSwapX(arg_0); +half2 quadSwapX_d60cec(thread tint_private_vars_struct* const tint_private_vars) { + half2 arg_0 = half2(1.0h); + half2 res = tint_msl_quadSwapX(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_d60cec(); +fragment void fragment_main(device half2* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapX_d60cec(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_d60cec(); +kernel void compute_main(device half2* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapX_d60cec(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/var/quadSwapX/edfa1f.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapX/edfa1f.wgsl.expected.msl index cd38d054f1d..3adb0bc2b63 100644 --- a/test/tint/builtins/gen/var/quadSwapX/edfa1f.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapX/edfa1f.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec4; +int4 tint_msl_quadSwapX(int4 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 1u)); +} -fn quadSwapX_edfa1f() -> vec4 { - var arg_0 = vec4(1i); - var res : vec4 = quadSwapX(arg_0); +int4 quadSwapX_edfa1f(thread tint_private_vars_struct* const tint_private_vars) { + int4 arg_0 = int4(1); + int4 res = tint_msl_quadSwapX(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapX_edfa1f(); +fragment void fragment_main(device int4* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapX_edfa1f(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapX_edfa1f(); +kernel void compute_main(device int4* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapX_edfa1f(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapX diff --git a/test/tint/builtins/gen/var/quadSwapY/06a67c.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapY/06a67c.wgsl.expected.msl index 37268e81c5c..8997ccd91ac 100644 --- a/test/tint/builtins/gen/var/quadSwapY/06a67c.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapY/06a67c.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec3; +uint3 tint_msl_quadSwapY(uint3 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_06a67c() -> vec3 { - var arg_0 = vec3(1u); - var res : vec3 = quadSwapY(arg_0); +uint3 quadSwapY_06a67c(thread tint_private_vars_struct* const tint_private_vars) { + uint3 arg_0 = uint3(1u); + uint3 res = tint_msl_quadSwapY(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_06a67c(); +fragment void fragment_main(device packed_uint3* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = packed_uint3(quadSwapY_06a67c(&(tint_private_vars))); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_06a67c(); +kernel void compute_main(device packed_uint3* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = packed_uint3(quadSwapY_06a67c(&(tint_private_vars))); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/var/quadSwapY/0c4938.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapY/0c4938.wgsl.expected.msl index e9a080c6cf3..1873357c4e9 100644 --- a/test/tint/builtins/gen/var/quadSwapY/0c4938.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapY/0c4938.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : u32; +uint tint_msl_quadSwapY(uint e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_0c4938() -> u32 { - var arg_0 = 1u; - var res : u32 = quadSwapY(arg_0); +uint quadSwapY_0c4938(thread tint_private_vars_struct* const tint_private_vars) { + uint arg_0 = 1u; + uint res = tint_msl_quadSwapY(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_0c4938(); +fragment void fragment_main(device uint* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapY_0c4938(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_0c4938(); +kernel void compute_main(device uint* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapY_0c4938(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/var/quadSwapY/0d05a8.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapY/0d05a8.wgsl.expected.msl index 44a7e215d9a..de3ab68f1ae 100644 --- a/test/tint/builtins/gen/var/quadSwapY/0d05a8.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapY/0d05a8.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec2; +int2 tint_msl_quadSwapY(int2 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_0d05a8() -> vec2 { - var arg_0 = vec2(1i); - var res : vec2 = quadSwapY(arg_0); +int2 quadSwapY_0d05a8(thread tint_private_vars_struct* const tint_private_vars) { + int2 arg_0 = int2(1); + int2 res = tint_msl_quadSwapY(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_0d05a8(); +fragment void fragment_main(device int2* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapY_0d05a8(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_0d05a8(); +kernel void compute_main(device int2* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapY_0d05a8(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/var/quadSwapY/14bb9a.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapY/14bb9a.wgsl.expected.msl index 3351af65a6b..3a3368541ba 100644 --- a/test/tint/builtins/gen/var/quadSwapY/14bb9a.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapY/14bb9a.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec4; +int4 tint_msl_quadSwapY(int4 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_14bb9a() -> vec4 { - var arg_0 = vec4(1i); - var res : vec4 = quadSwapY(arg_0); +int4 quadSwapY_14bb9a(thread tint_private_vars_struct* const tint_private_vars) { + int4 arg_0 = int4(1); + int4 res = tint_msl_quadSwapY(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_14bb9a(); +fragment void fragment_main(device int4* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapY_14bb9a(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_14bb9a(); +kernel void compute_main(device int4* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapY_14bb9a(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/var/quadSwapY/1f1a06.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapY/1f1a06.wgsl.expected.msl index 8b3b68cc1ae..4b33cba2824 100644 --- a/test/tint/builtins/gen/var/quadSwapY/1f1a06.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapY/1f1a06.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec2; +float2 tint_msl_quadSwapY(float2 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_1f1a06() -> vec2 { - var arg_0 = vec2(1.0f); - var res : vec2 = quadSwapY(arg_0); +float2 quadSwapY_1f1a06(thread tint_private_vars_struct* const tint_private_vars) { + float2 arg_0 = float2(1.0f); + float2 res = tint_msl_quadSwapY(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_1f1a06(); +fragment void fragment_main(device float2* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapY_1f1a06(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_1f1a06(); +kernel void compute_main(device float2* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapY_1f1a06(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/var/quadSwapY/264908.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapY/264908.wgsl.expected.msl index 7d2b700ac3e..58271e24cb8 100644 --- a/test/tint/builtins/gen/var/quadSwapY/264908.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapY/264908.wgsl.expected.msl @@ -1,26 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -enable subgroups_f16; -enable f16; -@group(0) @binding(0) var prevent_dce : vec3; +half3 tint_msl_quadSwapY(half3 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_264908() -> vec3 { - var arg_0 = vec3(1.0h); - var res : vec3 = quadSwapY(arg_0); +half3 quadSwapY_264908(thread tint_private_vars_struct* const tint_private_vars) { + half3 arg_0 = half3(1.0h); + half3 res = tint_msl_quadSwapY(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_264908(); +fragment void fragment_main(device packed_half3* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = packed_half3(quadSwapY_264908(&(tint_private_vars))); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_264908(); +kernel void compute_main(device packed_half3* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = packed_half3(quadSwapY_264908(&(tint_private_vars))); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/var/quadSwapY/5b2e67.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapY/5b2e67.wgsl.expected.msl index a89da2fbb94..3d05bb5011b 100644 --- a/test/tint/builtins/gen/var/quadSwapY/5b2e67.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapY/5b2e67.wgsl.expected.msl @@ -1,26 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -enable subgroups_f16; -enable f16; -@group(0) @binding(0) var prevent_dce : vec4; +half4 tint_msl_quadSwapY(half4 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_5b2e67() -> vec4 { - var arg_0 = vec4(1.0h); - var res : vec4 = quadSwapY(arg_0); +half4 quadSwapY_5b2e67(thread tint_private_vars_struct* const tint_private_vars) { + half4 arg_0 = half4(1.0h); + half4 res = tint_msl_quadSwapY(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_5b2e67(); +fragment void fragment_main(device half4* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapY_5b2e67(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_5b2e67(); +kernel void compute_main(device half4* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapY_5b2e67(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/var/quadSwapY/6f6bc9.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapY/6f6bc9.wgsl.expected.msl index badf5af6a69..4691f9fccd7 100644 --- a/test/tint/builtins/gen/var/quadSwapY/6f6bc9.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapY/6f6bc9.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : f32; +float tint_msl_quadSwapY(float e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_6f6bc9() -> f32 { - var arg_0 = 1.0f; - var res : f32 = quadSwapY(arg_0); +float quadSwapY_6f6bc9(thread tint_private_vars_struct* const tint_private_vars) { + float arg_0 = 1.0f; + float res = tint_msl_quadSwapY(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_6f6bc9(); +fragment void fragment_main(device float* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapY_6f6bc9(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_6f6bc9(); +kernel void compute_main(device float* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapY_6f6bc9(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/var/quadSwapY/9277e9.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapY/9277e9.wgsl.expected.msl index e7d5a05e2a0..e2129f7cd94 100644 --- a/test/tint/builtins/gen/var/quadSwapY/9277e9.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapY/9277e9.wgsl.expected.msl @@ -1,26 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -enable subgroups_f16; -enable f16; -@group(0) @binding(0) var prevent_dce : f16; +half tint_msl_quadSwapY(half e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_9277e9() -> f16 { - var arg_0 = 1.0h; - var res : f16 = quadSwapY(arg_0); +half quadSwapY_9277e9(thread tint_private_vars_struct* const tint_private_vars) { + half arg_0 = 1.0h; + half res = tint_msl_quadSwapY(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_9277e9(); +fragment void fragment_main(device half* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapY_9277e9(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_9277e9(); +kernel void compute_main(device half* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapY_9277e9(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/var/quadSwapY/94ab6d.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapY/94ab6d.wgsl.expected.msl index 5d86cb13897..22010aa1fca 100644 --- a/test/tint/builtins/gen/var/quadSwapY/94ab6d.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapY/94ab6d.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : i32; +int tint_msl_quadSwapY(int e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_94ab6d() -> i32 { - var arg_0 = 1i; - var res : i32 = quadSwapY(arg_0); +int quadSwapY_94ab6d(thread tint_private_vars_struct* const tint_private_vars) { + int arg_0 = 1; + int res = tint_msl_quadSwapY(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_94ab6d(); +fragment void fragment_main(device int* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapY_94ab6d(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_94ab6d(); +kernel void compute_main(device int* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapY_94ab6d(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/var/quadSwapY/a27e1c.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapY/a27e1c.wgsl.expected.msl index 28b56a75f9f..83f36badf87 100644 --- a/test/tint/builtins/gen/var/quadSwapY/a27e1c.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapY/a27e1c.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec2; +uint2 tint_msl_quadSwapY(uint2 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_a27e1c() -> vec2 { - var arg_0 = vec2(1u); - var res : vec2 = quadSwapY(arg_0); +uint2 quadSwapY_a27e1c(thread tint_private_vars_struct* const tint_private_vars) { + uint2 arg_0 = uint2(1u); + uint2 res = tint_msl_quadSwapY(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_a27e1c(); +fragment void fragment_main(device uint2* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapY_a27e1c(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_a27e1c(); +kernel void compute_main(device uint2* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapY_a27e1c(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/var/quadSwapY/a50fcb.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapY/a50fcb.wgsl.expected.msl index 671decae376..91123f898f7 100644 --- a/test/tint/builtins/gen/var/quadSwapY/a50fcb.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapY/a50fcb.wgsl.expected.msl @@ -1,26 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -enable subgroups_f16; -enable f16; -@group(0) @binding(0) var prevent_dce : vec2; +half2 tint_msl_quadSwapY(half2 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_a50fcb() -> vec2 { - var arg_0 = vec2(1.0h); - var res : vec2 = quadSwapY(arg_0); +half2 quadSwapY_a50fcb(thread tint_private_vars_struct* const tint_private_vars) { + half2 arg_0 = half2(1.0h); + half2 res = tint_msl_quadSwapY(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_a50fcb(); +fragment void fragment_main(device half2* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapY_a50fcb(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_a50fcb(); +kernel void compute_main(device half2* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapY_a50fcb(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/var/quadSwapY/b9d9e7.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapY/b9d9e7.wgsl.expected.msl index e11967eddf4..7fedc0ea2c9 100644 --- a/test/tint/builtins/gen/var/quadSwapY/b9d9e7.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapY/b9d9e7.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec4; +float4 tint_msl_quadSwapY(float4 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_b9d9e7() -> vec4 { - var arg_0 = vec4(1.0f); - var res : vec4 = quadSwapY(arg_0); +float4 quadSwapY_b9d9e7(thread tint_private_vars_struct* const tint_private_vars) { + float4 arg_0 = float4(1.0f); + float4 res = tint_msl_quadSwapY(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_b9d9e7(); +fragment void fragment_main(device float4* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapY_b9d9e7(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_b9d9e7(); +kernel void compute_main(device float4* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapY_b9d9e7(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/var/quadSwapY/bb697b.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapY/bb697b.wgsl.expected.msl index 3602fc93a65..7cd8059c49f 100644 --- a/test/tint/builtins/gen/var/quadSwapY/bb697b.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapY/bb697b.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec4; +uint4 tint_msl_quadSwapY(uint4 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_bb697b() -> vec4 { - var arg_0 = vec4(1u); - var res : vec4 = quadSwapY(arg_0); +uint4 quadSwapY_bb697b(thread tint_private_vars_struct* const tint_private_vars) { + uint4 arg_0 = uint4(1u); + uint4 res = tint_msl_quadSwapY(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_bb697b(); +fragment void fragment_main(device uint4* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = quadSwapY_bb697b(&(tint_private_vars)); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_bb697b(); +kernel void compute_main(device uint4* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = quadSwapY_bb697b(&(tint_private_vars)); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/var/quadSwapY/be4e72.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapY/be4e72.wgsl.expected.msl index 8ecd5e8cb95..a7e1b110c93 100644 --- a/test/tint/builtins/gen/var/quadSwapY/be4e72.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapY/be4e72.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec3; +int3 tint_msl_quadSwapY(int3 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_be4e72() -> vec3 { - var arg_0 = vec3(1i); - var res : vec3 = quadSwapY(arg_0); +int3 quadSwapY_be4e72(thread tint_private_vars_struct* const tint_private_vars) { + int3 arg_0 = int3(1); + int3 res = tint_msl_quadSwapY(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_be4e72(); +fragment void fragment_main(device packed_int3* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = packed_int3(quadSwapY_be4e72(&(tint_private_vars))); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_be4e72(); +kernel void compute_main(device packed_int3* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = packed_int3(quadSwapY_be4e72(&(tint_private_vars))); + return; } -Failed to generate: error: Unknown import method: quadSwapY diff --git a/test/tint/builtins/gen/var/quadSwapY/d1ab4d.wgsl.expected.msl b/test/tint/builtins/gen/var/quadSwapY/d1ab4d.wgsl.expected.msl index 1d82fefcf4b..723750c47e1 100644 --- a/test/tint/builtins/gen/var/quadSwapY/d1ab4d.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/quadSwapY/d1ab4d.wgsl.expected.msl @@ -1,24 +1,36 @@ -SKIP: FAILED +#include +using namespace metal; +struct tint_private_vars_struct { + uint tint_msl_thread_index_in_quadgroup; +}; -enable subgroups; -@group(0) @binding(0) var prevent_dce : vec3; +float3 tint_msl_quadSwapY(float3 e, thread tint_private_vars_struct* const tint_private_vars) { + return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 2u)); +} -fn quadSwapY_d1ab4d() -> vec3 { - var arg_0 = vec3(1.0f); - var res : vec3 = quadSwapY(arg_0); +float3 quadSwapY_d1ab4d(thread tint_private_vars_struct* const tint_private_vars) { + float3 arg_0 = float3(1.0f); + float3 res = tint_msl_quadSwapY(arg_0, tint_private_vars); return res; } -@fragment -fn fragment_main() { - prevent_dce = quadSwapY_d1ab4d(); +fragment void fragment_main(device packed_float3* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup; + } + *(tint_symbol) = packed_float3(quadSwapY_d1ab4d(&(tint_private_vars))); + return; } -@compute @workgroup_size(1) -fn compute_main() { - prevent_dce = quadSwapY_d1ab4d(); +kernel void compute_main(device packed_float3* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) { + thread tint_private_vars_struct tint_private_vars = {}; + { + tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1; + } + *(tint_symbol_1) = packed_float3(quadSwapY_d1ab4d(&(tint_private_vars))); + return; } -Failed to generate: error: Unknown import method: quadSwapY