From 8357b6bd1f0291bac3ad01d3fedda4b69ebb2054 Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Thu, 12 Dec 2024 14:39:00 -0800 Subject: [PATCH 1/2] [NFCI][SYCL] Refactor named swizzles implementation Instead of using preprocessor via `swizzles.def` file that is included multiple times, employ mixin-based approach which is more C++'ish. --- .../sycl/detail/named_swizzles_mixin.hpp | 823 ++++++++++++++++++ sycl/include/sycl/swizzles.def | 823 ------------------ sycl/include/sycl/vector.hpp | 62 +- .../vector/vector_bf16_builtins.cpp | 10 +- .../vector/vector_convert_bfloat.cpp | 8 +- .../vector/vector_math_ops.cpp | 44 +- 6 files changed, 872 insertions(+), 898 deletions(-) create mode 100644 sycl/include/sycl/detail/named_swizzles_mixin.hpp delete mode 100644 sycl/include/sycl/swizzles.def diff --git a/sycl/include/sycl/detail/named_swizzles_mixin.hpp b/sycl/include/sycl/detail/named_swizzles_mixin.hpp new file mode 100644 index 0000000000000..c7bee0d1b7bb1 --- /dev/null +++ b/sycl/include/sycl/detail/named_swizzles_mixin.hpp @@ -0,0 +1,823 @@ +//==---------------- named_swizzles_mixin.hpp ------------------------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +// This files implements two mixins +// `NamedSwizzlesMixinConst`/`NamedSwizzlesMixinBoth` that abstract away named +// swizzles implementation for SYCL vector and swizzles classes + +#include + +namespace sycl { +inline namespace _V1 { +namespace detail { + +#ifndef SYCL_SIMPLE_SWIZZLES +#define __SYCL_SWIZZLE_MIXIN_SIMPLE_SWIZZLES +#else +// TODO: It might be beneficial to use partial specializations for different Ns, +// instead of making all the named swizzles templates with SFINAE conditions. +#define __SYCL_SWIZZLE_MIXIN_SIMPLE_SWIZZLES \ + /* __swizzled_vec__ XYZW_SWIZZLE() const; */ \ + __SYCL_SWIZZLE_MIXIN_METHOD(N <= 4, xx, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(2 <= N && N <= 4, xy, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, xz, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xw, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(2 <= N && N <= 4, yx, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(2 <= N && N <= 4, yy, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, yz, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yw, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zx, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zy, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zz, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zw, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wx, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wy, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wz, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ww, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N <= 4, xxx, 0, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(2 <= N && N <= 4, xxy, 0, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, xxz, 0, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xxw, 0, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(2 <= N && N <= 4, xyx, 0, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(2 <= N && N <= 4, xyy, 0, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, xyz, 0, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xyw, 0, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, xzx, 0, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, xzy, 0, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, xzz, 0, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xzw, 0, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xwx, 0, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xwy, 0, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xwz, 0, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xww, 0, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(2 <= N && N <= 4, yxx, 1, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(2 <= N && N <= 4, yxy, 1, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, yxz, 1, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yxw, 1, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(2 <= N && N <= 4, yyx, 1, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(2 <= N && N <= 4, yyy, 1, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, yyz, 1, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yyw, 1, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, yzx, 1, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, yzy, 1, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, yzz, 1, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yzw, 1, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ywx, 1, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ywy, 1, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ywz, 1, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yww, 1, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zxx, 2, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zxy, 2, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zxz, 2, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zxw, 2, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zyx, 2, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zyy, 2, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zyz, 2, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zyw, 2, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zzx, 2, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zzy, 2, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zzz, 2, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zzw, 2, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zwx, 2, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zwy, 2, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zwz, 2, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zww, 2, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wxx, 3, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wxy, 3, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wxz, 3, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wxw, 3, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wyx, 3, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wyy, 3, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wyz, 3, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wyw, 3, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wzx, 3, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wzy, 3, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wzz, 3, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wzw, 3, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wwx, 3, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wwy, 3, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wwz, 3, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, www, 3, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N <= 4, xxxx, 0, 0, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(2 <= N && N <= 4, xxxy, 0, 0, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, xxxz, 0, 0, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xxxw, 0, 0, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(2 <= N && N <= 4, xxyx, 0, 0, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(2 <= N && N <= 4, xxyy, 0, 0, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, xxyz, 0, 0, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xxyw, 0, 0, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, xxzx, 0, 0, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, xxzy, 0, 0, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, xxzz, 0, 0, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xxzw, 0, 0, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xxwx, 0, 0, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xxwy, 0, 0, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xxwz, 0, 0, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xxww, 0, 0, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(2 <= N && N <= 4, xyxx, 0, 1, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(2 <= N && N <= 4, xyxy, 0, 1, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, xyxz, 0, 1, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xyxw, 0, 1, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(2 <= N && N <= 4, xyyx, 0, 1, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(2 <= N && N <= 4, xyyy, 0, 1, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, xyyz, 0, 1, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xyyw, 0, 1, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, xyzx, 0, 1, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, xyzy, 0, 1, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, xyzz, 0, 1, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xyzw, 0, 1, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xywx, 0, 1, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xywy, 0, 1, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xywz, 0, 1, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xyww, 0, 1, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, xzxx, 0, 2, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, xzxy, 0, 2, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, xzxz, 0, 2, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xzxw, 0, 2, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, xzyx, 0, 2, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, xzyy, 0, 2, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, xzyz, 0, 2, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xzyw, 0, 2, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, xzzx, 0, 2, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, xzzy, 0, 2, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, xzzz, 0, 2, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xzzw, 0, 2, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xzwx, 0, 2, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xzwy, 0, 2, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xzwz, 0, 2, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xzww, 0, 2, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xwxx, 0, 3, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xwxy, 0, 3, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xwxz, 0, 3, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xwxw, 0, 3, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xwyx, 0, 3, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xwyy, 0, 3, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xwyz, 0, 3, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xwyw, 0, 3, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xwzx, 0, 3, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xwzy, 0, 3, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xwzz, 0, 3, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xwzw, 0, 3, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xwwx, 0, 3, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xwwy, 0, 3, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xwwz, 0, 3, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, xwww, 0, 3, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(2 <= N && N <= 4, yxxx, 1, 0, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(2 <= N && N <= 4, yxxy, 1, 0, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, yxxz, 1, 0, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yxxw, 1, 0, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(2 <= N && N <= 4, yxyx, 1, 0, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(2 <= N && N <= 4, yxyy, 1, 0, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, yxyz, 1, 0, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yxyw, 1, 0, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, yxzx, 1, 0, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, yxzy, 1, 0, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, yxzz, 1, 0, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yxzw, 1, 0, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yxwx, 1, 0, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yxwy, 1, 0, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yxwz, 1, 0, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yxww, 1, 0, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(2 <= N && N <= 4, yyxx, 1, 1, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(2 <= N && N <= 4, yyxy, 1, 1, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, yyxz, 1, 1, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yyxw, 1, 1, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(2 <= N && N <= 4, yyyx, 1, 1, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(2 <= N && N <= 4, yyyy, 1, 1, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, yyyz, 1, 1, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yyyw, 1, 1, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, yyzx, 1, 1, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, yyzy, 1, 1, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, yyzz, 1, 1, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yyzw, 1, 1, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yywx, 1, 1, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yywy, 1, 1, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yywz, 1, 1, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yyww, 1, 1, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, yzxx, 1, 2, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, yzxy, 1, 2, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, yzxz, 1, 2, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yzxw, 1, 2, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, yzyx, 1, 2, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, yzyy, 1, 2, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, yzyz, 1, 2, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yzyw, 1, 2, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, yzzx, 1, 2, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, yzzy, 1, 2, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, yzzz, 1, 2, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yzzw, 1, 2, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yzwx, 1, 2, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yzwy, 1, 2, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yzwz, 1, 2, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, yzww, 1, 2, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ywxx, 1, 3, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ywxy, 1, 3, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ywxz, 1, 3, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ywxw, 1, 3, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ywyx, 1, 3, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ywyy, 1, 3, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ywyz, 1, 3, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ywyw, 1, 3, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ywzx, 1, 3, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ywzy, 1, 3, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ywzz, 1, 3, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ywzw, 1, 3, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ywwx, 1, 3, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ywwy, 1, 3, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ywwz, 1, 3, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ywww, 1, 3, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zxxx, 2, 0, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zxxy, 2, 0, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zxxz, 2, 0, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zxxw, 2, 0, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zxyx, 2, 0, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zxyy, 2, 0, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zxyz, 2, 0, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zxyw, 2, 0, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zxzx, 2, 0, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zxzy, 2, 0, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zxzz, 2, 0, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zxzw, 2, 0, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zxwx, 2, 0, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zxwy, 2, 0, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zxwz, 2, 0, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zxww, 2, 0, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zyxx, 2, 1, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zyxy, 2, 1, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zyxz, 2, 1, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zyxw, 2, 1, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zyyx, 2, 1, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zyyy, 2, 1, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zyyz, 2, 1, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zyyw, 2, 1, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zyzx, 2, 1, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zyzy, 2, 1, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zyzz, 2, 1, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zyzw, 2, 1, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zywx, 2, 1, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zywy, 2, 1, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zywz, 2, 1, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zyww, 2, 1, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zzxx, 2, 2, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zzxy, 2, 2, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zzxz, 2, 2, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zzxw, 2, 2, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zzyx, 2, 2, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zzyy, 2, 2, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zzyz, 2, 2, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zzyw, 2, 2, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zzzx, 2, 2, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zzzy, 2, 2, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3 || N == 4, zzzz, 2, 2, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zzzw, 2, 2, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zzwx, 2, 2, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zzwy, 2, 2, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zzwz, 2, 2, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zzww, 2, 2, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zwxx, 2, 3, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zwxy, 2, 3, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zwxz, 2, 3, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zwxw, 2, 3, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zwyx, 2, 3, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zwyy, 2, 3, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zwyz, 2, 3, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zwyw, 2, 3, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zwzx, 2, 3, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zwzy, 2, 3, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zwzz, 2, 3, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zwzw, 2, 3, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zwwx, 2, 3, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zwwy, 2, 3, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zwwz, 2, 3, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, zwww, 2, 3, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wxxx, 3, 0, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wxxy, 3, 0, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wxxz, 3, 0, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wxxw, 3, 0, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wxyx, 3, 0, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wxyy, 3, 0, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wxyz, 3, 0, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wxyw, 3, 0, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wxzx, 3, 0, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wxzy, 3, 0, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wxzz, 3, 0, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wxzw, 3, 0, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wxwx, 3, 0, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wxwy, 3, 0, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wxwz, 3, 0, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wxww, 3, 0, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wyxx, 3, 1, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wyxy, 3, 1, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wyxz, 3, 1, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wyxw, 3, 1, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wyyx, 3, 1, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wyyy, 3, 1, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wyyz, 3, 1, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wyyw, 3, 1, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wyzx, 3, 1, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wyzy, 3, 1, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wyzz, 3, 1, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wyzw, 3, 1, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wywx, 3, 1, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wywy, 3, 1, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wywz, 3, 1, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wyww, 3, 1, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wzxx, 3, 2, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wzxy, 3, 2, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wzxz, 3, 2, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wzxw, 3, 2, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wzyx, 3, 2, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wzyy, 3, 2, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wzyz, 3, 2, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wzyw, 3, 2, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wzzx, 3, 2, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wzzy, 3, 2, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wzzz, 3, 2, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wzzw, 3, 2, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wzwx, 3, 2, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wzwy, 3, 2, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wzwz, 3, 2, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wzww, 3, 2, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wwxx, 3, 3, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wwxy, 3, 3, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wwxz, 3, 3, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wwxw, 3, 3, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wwyx, 3, 3, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wwyy, 3, 3, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wwyz, 3, 3, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wwyw, 3, 3, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wwzx, 3, 3, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wwzy, 3, 3, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wwzz, 3, 3, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wwzw, 3, 3, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wwwx, 3, 3, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wwwy, 3, 3, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wwwz, 3, 3, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, wwww, 3, 3, 3, 3) \ + \ + /* __swizzled_vec__ RGBA_SWIZZLE() const; */ \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rr, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rg, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rb, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ra, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gr, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gg, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gb, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ga, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, br, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bg, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bb, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ba, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ar, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ag, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ab, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aa, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rrr, 0, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rrg, 0, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rrb, 0, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rra, 0, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rgr, 0, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rgg, 0, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rgb, 0, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rga, 0, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rbr, 0, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rbg, 0, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rbb, 0, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rba, 0, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rar, 0, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rag, 0, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rab, 0, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, raa, 0, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, grr, 1, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, grg, 1, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, grb, 1, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gra, 1, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ggr, 1, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ggg, 1, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ggb, 1, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gga, 1, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gbr, 1, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gbg, 1, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gbb, 1, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gba, 1, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gar, 1, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gag, 1, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gab, 1, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gaa, 1, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, brr, 2, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, brg, 2, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, brb, 2, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bra, 2, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bgr, 2, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bgg, 2, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bgb, 2, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bga, 2, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bbr, 2, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bbg, 2, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bbb, 2, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bba, 2, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bar, 2, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bag, 2, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bab, 2, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, baa, 2, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, arr, 3, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, arg, 3, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, arb, 3, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ara, 3, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, agr, 3, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, agg, 3, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, agb, 3, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aga, 3, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, abr, 3, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, abg, 3, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, abb, 3, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aba, 3, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aar, 3, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aag, 3, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aab, 3, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aaa, 3, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rrrr, 0, 0, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rrrg, 0, 0, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rrrb, 0, 0, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rrra, 0, 0, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rrgr, 0, 0, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rrgg, 0, 0, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rrgb, 0, 0, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rrga, 0, 0, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rrbr, 0, 0, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rrbg, 0, 0, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rrbb, 0, 0, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rrba, 0, 0, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rrar, 0, 0, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rrag, 0, 0, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rrab, 0, 0, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rraa, 0, 0, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rgrr, 0, 1, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rgrg, 0, 1, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rgrb, 0, 1, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rgra, 0, 1, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rggr, 0, 1, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rggg, 0, 1, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rggb, 0, 1, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rgga, 0, 1, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rgbr, 0, 1, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rgbg, 0, 1, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rgbb, 0, 1, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rgba, 0, 1, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rgar, 0, 1, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rgag, 0, 1, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rgab, 0, 1, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rgaa, 0, 1, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rbrr, 0, 2, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rbrg, 0, 2, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rbrb, 0, 2, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rbra, 0, 2, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rbgr, 0, 2, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rbgg, 0, 2, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rbgb, 0, 2, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rbga, 0, 2, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rbbr, 0, 2, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rbbg, 0, 2, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rbbb, 0, 2, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rbba, 0, 2, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rbar, 0, 2, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rbag, 0, 2, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rbab, 0, 2, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rbaa, 0, 2, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rarr, 0, 3, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rarg, 0, 3, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rarb, 0, 3, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rara, 0, 3, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ragr, 0, 3, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ragg, 0, 3, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ragb, 0, 3, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, raga, 0, 3, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rabr, 0, 3, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rabg, 0, 3, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, rabb, 0, 3, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, raba, 0, 3, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, raar, 0, 3, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, raag, 0, 3, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, raab, 0, 3, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, raaa, 0, 3, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, grrr, 1, 0, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, grrg, 1, 0, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, grrb, 1, 0, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, grra, 1, 0, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, grgr, 1, 0, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, grgg, 1, 0, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, grgb, 1, 0, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, grga, 1, 0, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, grbr, 1, 0, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, grbg, 1, 0, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, grbb, 1, 0, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, grba, 1, 0, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, grar, 1, 0, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, grag, 1, 0, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, grab, 1, 0, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, graa, 1, 0, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ggrr, 1, 1, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ggrg, 1, 1, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ggrb, 1, 1, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ggra, 1, 1, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gggr, 1, 1, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gggg, 1, 1, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gggb, 1, 1, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ggga, 1, 1, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ggbr, 1, 1, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ggbg, 1, 1, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ggbb, 1, 1, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ggba, 1, 1, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ggar, 1, 1, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ggag, 1, 1, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ggab, 1, 1, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, ggaa, 1, 1, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gbrr, 1, 2, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gbrg, 1, 2, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gbrb, 1, 2, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gbra, 1, 2, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gbgr, 1, 2, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gbgg, 1, 2, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gbgb, 1, 2, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gbga, 1, 2, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gbbr, 1, 2, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gbbg, 1, 2, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gbbb, 1, 2, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gbba, 1, 2, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gbar, 1, 2, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gbag, 1, 2, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gbab, 1, 2, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gbaa, 1, 2, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, garr, 1, 3, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, garg, 1, 3, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, garb, 1, 3, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gara, 1, 3, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gagr, 1, 3, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gagg, 1, 3, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gagb, 1, 3, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gaga, 1, 3, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gabr, 1, 3, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gabg, 1, 3, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gabb, 1, 3, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gaba, 1, 3, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gaar, 1, 3, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gaag, 1, 3, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gaab, 1, 3, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, gaaa, 1, 3, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, brrr, 2, 0, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, brrg, 2, 0, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, brrb, 2, 0, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, brra, 2, 0, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, brgr, 2, 0, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, brgg, 2, 0, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, brgb, 2, 0, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, brga, 2, 0, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, brbr, 2, 0, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, brbg, 2, 0, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, brbb, 2, 0, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, brba, 2, 0, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, brar, 2, 0, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, brag, 2, 0, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, brab, 2, 0, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, braa, 2, 0, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bgrr, 2, 1, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bgrg, 2, 1, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bgrb, 2, 1, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bgra, 2, 1, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bggr, 2, 1, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bggg, 2, 1, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bggb, 2, 1, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bgga, 2, 1, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bgbr, 2, 1, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bgbg, 2, 1, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bgbb, 2, 1, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bgba, 2, 1, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bgar, 2, 1, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bgag, 2, 1, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bgab, 2, 1, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bgaa, 2, 1, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bbrr, 2, 2, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bbrg, 2, 2, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bbrb, 2, 2, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bbra, 2, 2, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bbgr, 2, 2, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bbgg, 2, 2, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bbgb, 2, 2, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bbga, 2, 2, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bbbr, 2, 2, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bbbg, 2, 2, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bbbb, 2, 2, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bbba, 2, 2, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bbar, 2, 2, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bbag, 2, 2, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bbab, 2, 2, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bbaa, 2, 2, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, barr, 2, 3, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, barg, 2, 3, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, barb, 2, 3, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bara, 2, 3, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bagr, 2, 3, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bagg, 2, 3, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, bagb, 2, 3, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, baga, 2, 3, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, babr, 2, 3, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, babg, 2, 3, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, babb, 2, 3, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, baba, 2, 3, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, baar, 2, 3, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, baag, 2, 3, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, baab, 2, 3, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, baaa, 2, 3, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, arrr, 3, 0, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, arrg, 3, 0, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, arrb, 3, 0, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, arra, 3, 0, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, argr, 3, 0, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, argg, 3, 0, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, argb, 3, 0, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, arga, 3, 0, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, arbr, 3, 0, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, arbg, 3, 0, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, arbb, 3, 0, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, arba, 3, 0, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, arar, 3, 0, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, arag, 3, 0, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, arab, 3, 0, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, araa, 3, 0, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, agrr, 3, 1, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, agrg, 3, 1, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, agrb, 3, 1, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, agra, 3, 1, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aggr, 3, 1, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aggg, 3, 1, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aggb, 3, 1, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, agga, 3, 1, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, agbr, 3, 1, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, agbg, 3, 1, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, agbb, 3, 1, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, agba, 3, 1, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, agar, 3, 1, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, agag, 3, 1, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, agab, 3, 1, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, agaa, 3, 1, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, abrr, 3, 2, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, abrg, 3, 2, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, abrb, 3, 2, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, abra, 3, 2, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, abgr, 3, 2, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, abgg, 3, 2, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, abgb, 3, 2, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, abga, 3, 2, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, abbr, 3, 2, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, abbg, 3, 2, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, abbb, 3, 2, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, abba, 3, 2, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, abar, 3, 2, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, abag, 3, 2, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, abab, 3, 2, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, abaa, 3, 2, 3, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aarr, 3, 3, 0, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aarg, 3, 3, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aarb, 3, 3, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aara, 3, 3, 0, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aagr, 3, 3, 1, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aagg, 3, 3, 1, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aagb, 3, 3, 1, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aaga, 3, 3, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aabr, 3, 3, 2, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aabg, 3, 3, 2, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aabb, 3, 3, 2, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aaba, 3, 3, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aaar, 3, 3, 3, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aaag, 3, 3, 3, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aaab, 3, 3, 3, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, aaaa, 3, 3, 3, 3) +#endif + +#define __SYCL_SWIZZLE_MIXIN_ALL_SWIZZLES \ + /* __swizzled_vec__ XYZW_ACCESS() const; */ \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(N <= 4, x, 0) \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(N == 2 || N == 3 || N == 4, y, 1) \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(N == 3 || N == 4, z, 2) \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(N == 4, w, 3) \ + \ + /* __swizzled_vec__ RGBA_ACCESS() const; */ \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(N == 4, r, 0) \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(N == 4, g, 1) \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(N == 4, b, 2) \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(N == 4, a, 3) \ + \ + /* __swizzled_vec__ INDEX_ACCESS() const; */ \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(N > 0, s0, 0) \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(N > 1, s1, 1) \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(N > 2, s2, 2) \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(N > 2, s3, 3) \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(N > 4, s4, 4) \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(N > 4, s5, 5) \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(N > 4, s6, 6) \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(N > 4, s7, 7) \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(N == 16, s8, 8) \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(N == 16, s9, 9) \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(N == 16, sA, 10) \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(N == 16, sB, 11) \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(N == 16, sC, 12) \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(N == 16, sD, 13) \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(N == 16, sE, 14) \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(N == 16, sF, 15) \ + \ + /* __swizzled_vec__ lo()/hi() const; */ \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 2, lo, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3, lo, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, lo, 0, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 8, lo, 0, 1, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 16, lo, 0, 1, 2, 3, 4, 5, 6, 7) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 2, hi, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3, hi, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, hi, 2, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 8, hi, 4, 5, 6, 7) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 16, hi, 8, 9, 10, 11, 12, 13, 14, 15) \ + /* __swizzled_vec__ odd()/even() const; */ \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 2, odd, 1) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3, odd, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, odd, 1, 3) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 8, odd, 1, 3, 5, 7) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 16, odd, 1, 3, 5, 7, 9, 11, 13, 15) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 2, even, 0) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 3, even, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 4, even, 0, 2) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 8, even, 0, 2, 4, 6) \ + __SYCL_SWIZZLE_MIXIN_METHOD(N == 16, even, 0, 2, 4, 6, 8, 10, 12, 14) \ + /* SYCL_SIMPLE_SWIZZLES */ \ + __SYCL_SWIZZLE_MIXIN_SIMPLE_SWIZZLES + +#define __SYCL_SWIZZLE_MIXIN_METHOD_NON_CONST(COND, NAME, ...) \ + template \ + std::enable_if_t< \ + (COND), decltype(std::declval().template swizzle<__VA_ARGS__>())> \ + NAME() { \ + return static_cast(this)->template swizzle<__VA_ARGS__>(); \ + } + +#define __SYCL_SWIZZLE_MIXIN_METHOD_CONST(COND, NAME, ...) \ + template \ + std::enable_if_t<(COND), decltype(std::declval() \ + .template swizzle<__VA_ARGS__>())> \ + NAME() const { \ + return static_cast(this)->template swizzle<__VA_ARGS__>(); \ + } + +#define __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS_NON_CONST(COND, NAME, INDEX) \ + template \ + std::enable_if_t<(COND), decltype(std::declval()[0])> NAME() { \ + return (*static_cast(this))[INDEX]; \ + } +#define __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS_CONST(COND, NAME, INDEX) \ + template \ + std::enable_if_t<(COND), decltype(std::declval()[0])> NAME() \ + const { \ + return (*static_cast(this))[INDEX]; \ + } + +template struct NamedSwizzlesMixinConst { +#define __SYCL_SWIZZLE_MIXIN_METHOD(COND, NAME, ...) \ + __SYCL_SWIZZLE_MIXIN_METHOD_CONST(COND, NAME, __VA_ARGS__) + +#define __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(COND, NAME, INDEX) \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS_CONST(COND, NAME, INDEX) + + __SYCL_SWIZZLE_MIXIN_ALL_SWIZZLES + +#undef __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS +#undef __SYCL_SWIZZLE_MIXIN_METHOD +}; + +template struct NamedSwizzlesMixinBoth { +#define __SYCL_SWIZZLE_MIXIN_METHOD(COND, NAME, ...) \ + __SYCL_SWIZZLE_MIXIN_METHOD_NON_CONST(COND, NAME, __VA_ARGS__) \ + __SYCL_SWIZZLE_MIXIN_METHOD_CONST(COND, NAME, __VA_ARGS__) + +#define __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS(COND, NAME, INDEX) \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS_NON_CONST(COND, NAME, INDEX) \ + __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS_CONST(COND, NAME, INDEX) + + __SYCL_SWIZZLE_MIXIN_ALL_SWIZZLES + +#undef __SYCL_SWIZLLE_MIXIN_SCALAR_ACCESS +#undef __SYCL_SWIZZLE_MIXIN_METHOD +}; + +#undef __SYCL_SWIZZLE_MIXIN_METHOD_CONST +#undef __SYCL_SWIZZLE_MIXIN_METHOD_NON_CONST + +#undef __SYCL_SWIZZLE_MIXIN_SIMPLE_SWIZZLES + +} // namespace detail +} // namespace _V1 +} // namespace sycl diff --git a/sycl/include/sycl/swizzles.def b/sycl/include/sycl/swizzles.def deleted file mode 100644 index fd503e445d7fe..0000000000000 --- a/sycl/include/sycl/swizzles.def +++ /dev/null @@ -1,823 +0,0 @@ -//==---------------- swizzles.def --- SYCL types ---------------------------==// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -// Included by types.hpp twice, once for vec<> and once for SwizzleOp<>. - -// TODO: exclude L-Value swizzle like vec.xxxx() - -#if defined(__SYCL_ACCESS) || defined(__SYCL_INDEXER) || \ - defined(__SYCL_EXPAND) || defined(__SYCL_NTH_ARG) || \ - defined(__SYCL_E0) || defined(__SYCL_E1) || defined(__SYCL_E2) || \ - defined(__SYCL_E3) || defined(__SYCL_E4) || defined(__SYCL_E5) || \ - defined(__SYCL_E6) || defined(__SYCL_E7) || defined(__SYCL_E8) -#error "Undefine __SYCL_{ACCESS, INDEXER, EXPAND, NTH_ARG, E[0-8]} macros." -#endif - -#define __SYCL_INDEXER(_X) Indexer<_X>::value - -// Accepts any number of args >= _N, but expands to just the _N-th one. -// Make N equal to the max number of args handled + 1. Here, N == 9. -#define __SYCL_NTH_ARG(_1, _2, _3, _4, _5, _6, _7, _8, _N, ...) _N - -// Some compilers for such defines as '#define M1(_A, ...) M2(__VA_ARGS__)' and -// expression M1(1, 2, 3, 4) may pass (2, 3, 4) to M2 as one argument, -// not as three. To avoid that the following macro is needed. -// So, defining M1 as '#define M1(_A, ...) __SYCL_EXPAND( M2(__VA_ARGS__) )' -// solves the problem. -#define __SYCL_EXPAND(_X) _X - -// These are helper macros to create overrides based on the arity. -#define __SYCL_E0(_M, ...) -#define __SYCL_E1(_M, _X) _M(_X) -#define __SYCL_E2(_M, _X, ...) _M(_X), __SYCL_EXPAND(__SYCL_E1(_M, __VA_ARGS__)) -#define __SYCL_E3(_M, _X, ...) _M(_X), __SYCL_EXPAND(__SYCL_E2(_M, __VA_ARGS__)) -#define __SYCL_E4(_M, _X, ...) _M(_X), __SYCL_EXPAND(__SYCL_E3(_M, __VA_ARGS__)) -#define __SYCL_E5(_M, _X, ...) _M(_X), __SYCL_EXPAND(__SYCL_E4(_M, __VA_ARGS__)) -#define __SYCL_E6(_M, _X, ...) _M(_X), __SYCL_EXPAND(__SYCL_E5(_M, __VA_ARGS__)) -#define __SYCL_E7(_M, _X, ...) _M(_X), __SYCL_EXPAND(__SYCL_E6(_M, __VA_ARGS__)) -#define __SYCL_E8(_M, _X, ...) _M(_X), __SYCL_EXPAND(__SYCL_E7(_M, __VA_ARGS__)) - -// Creates template functions with required number of parameters. -// Each of the parameters X from __VA_ARGS__ is transformed into -// __SYCL_INDEXER(X). -// Currently the max number of parameters handled by this macro is 8. -#define __SYCL_ACCESS(_COND, _NAME, ...) \ - template \ - typename std::enable_if<(_COND), Swizzle< \ - __SYCL_EXPAND(__SYCL_NTH_ARG(__VA_ARGS__, \ - __SYCL_E8, __SYCL_E7, __SYCL_E6, __SYCL_E5, __SYCL_E4, \ - __SYCL_E3, __SYCL_E2, __SYCL_E1, __SYCL_E0) \ - (__SYCL_INDEXER, __VA_ARGS__))>>::type _NAME() { \ - return __SYCL_ACCESS_RETURN; \ - } \ - template \ - typename std::enable_if<(_COND), ConstSwizzle< \ - __SYCL_EXPAND(__SYCL_NTH_ARG(__VA_ARGS__, \ - __SYCL_E8, __SYCL_E7, __SYCL_E6, __SYCL_E5, __SYCL_E4, \ - __SYCL_E3, __SYCL_E2, __SYCL_E1, __SYCL_E0) \ - (__SYCL_INDEXER, __VA_ARGS__))>>::type _NAME() const { \ - return __SYCL_ACCESS_RETURN; \ - } - -#define __SYCL_SCALAR_ACCESS(_COND, _NAME, _INDEX) \ - template \ - typename std::enable_if<(_COND), DataT &>::type _NAME() { \ - return (*__SYCL_ACCESS_RETURN)[_INDEX]; \ - } \ - template \ - typename std::enable_if<(_COND), const DataT &>::type _NAME() const { \ - return (*__SYCL_ACCESS_RETURN)[_INDEX]; \ - } - -//__swizzled_vec__ XYZW_ACCESS() const; -__SYCL_SCALAR_ACCESS(N <= 4, x, 0) -__SYCL_SCALAR_ACCESS(N == 2 || N == 3 || N == 4, y, 1) -__SYCL_SCALAR_ACCESS(N == 3 || N == 4, z, 2) -__SYCL_SCALAR_ACCESS(N == 4, w, 3) - -//__swizzled_vec__ RGBA_ACCESS() const; -__SYCL_SCALAR_ACCESS(N == 4, r, 0) -__SYCL_SCALAR_ACCESS(N == 4, g, 1) -__SYCL_SCALAR_ACCESS(N == 4, b, 2) -__SYCL_SCALAR_ACCESS(N == 4, a, 3) - -//__swizzled_vec__ INDEX_ACCESS() const; -__SYCL_SCALAR_ACCESS(N > 0, s0, 0) -__SYCL_SCALAR_ACCESS(N > 1, s1, 1) -__SYCL_SCALAR_ACCESS(N > 2, s2, 2) -__SYCL_SCALAR_ACCESS(N > 2, s3, 3) -__SYCL_SCALAR_ACCESS(N > 4, s4, 4) -__SYCL_SCALAR_ACCESS(N > 4, s5, 5) -__SYCL_SCALAR_ACCESS(N > 4, s6, 6) -__SYCL_SCALAR_ACCESS(N > 4, s7, 7) -__SYCL_SCALAR_ACCESS(N == 16, s8, 8) -__SYCL_SCALAR_ACCESS(N == 16, s9, 9) -__SYCL_SCALAR_ACCESS(N == 16, sA, 10) -__SYCL_SCALAR_ACCESS(N == 16, sB, 11) -__SYCL_SCALAR_ACCESS(N == 16, sC, 12) -__SYCL_SCALAR_ACCESS(N == 16, sD, 13) -__SYCL_SCALAR_ACCESS(N == 16, sE, 14) -__SYCL_SCALAR_ACCESS(N == 16, sF, 15) - -#ifdef SYCL_SIMPLE_SWIZZLES -//__swizzled_vec__ XYZW_SWIZZLE() const; -__SYCL_ACCESS(N <= 4, xx, 0, 0) -__SYCL_ACCESS(2 <= N && N <= 4, xy, 0, 1) -__SYCL_ACCESS(N == 3 || N == 4, xz, 0, 2) -__SYCL_ACCESS(N == 4, xw, 0, 3) -__SYCL_ACCESS(2 <= N && N <= 4, yx, 1, 0) -__SYCL_ACCESS(2 <= N && N <= 4, yy, 1, 1) -__SYCL_ACCESS(N == 3 || N == 4, yz, 1, 2) -__SYCL_ACCESS(N == 4, yw, 1, 3) -__SYCL_ACCESS(N == 3 || N == 4, zx, 2, 0) -__SYCL_ACCESS(N == 3 || N == 4, zy, 2, 1) -__SYCL_ACCESS(N == 3 || N == 4, zz, 2, 2) -__SYCL_ACCESS(N == 4, zw, 2, 3) -__SYCL_ACCESS(N == 4, wx, 3, 0) -__SYCL_ACCESS(N == 4, wy, 3, 1) -__SYCL_ACCESS(N == 4, wz, 3, 2) -__SYCL_ACCESS(N == 4, ww, 3, 3) -__SYCL_ACCESS(N <= 4, xxx, 0, 0, 0) -__SYCL_ACCESS(2 <= N && N <= 4, xxy, 0, 0, 1) -__SYCL_ACCESS(N == 3 || N == 4, xxz, 0, 0, 2) -__SYCL_ACCESS(N == 4, xxw, 0, 0, 3) -__SYCL_ACCESS(2 <= N && N <= 4, xyx, 0, 1, 0) -__SYCL_ACCESS(2 <= N && N <= 4, xyy, 0, 1, 1) -__SYCL_ACCESS(N == 3 || N == 4, xyz, 0, 1, 2) -__SYCL_ACCESS(N == 4, xyw, 0, 1, 3) -__SYCL_ACCESS(N == 3 || N == 4, xzx, 0, 2, 0) -__SYCL_ACCESS(N == 3 || N == 4, xzy, 0, 2, 1) -__SYCL_ACCESS(N == 3 || N == 4, xzz, 0, 2, 2) -__SYCL_ACCESS(N == 4, xzw, 0, 2, 3) -__SYCL_ACCESS(N == 4, xwx, 0, 3, 0) -__SYCL_ACCESS(N == 4, xwy, 0, 3, 1) -__SYCL_ACCESS(N == 4, xwz, 0, 3, 2) -__SYCL_ACCESS(N == 4, xww, 0, 3, 3) -__SYCL_ACCESS(2 <= N && N <= 4, yxx, 1, 0, 0) -__SYCL_ACCESS(2 <= N && N <= 4, yxy, 1, 0, 1) -__SYCL_ACCESS(N == 3 || N == 4, yxz, 1, 0, 2) -__SYCL_ACCESS(N == 4, yxw, 1, 0, 3) -__SYCL_ACCESS(2 <= N && N <= 4, yyx, 1, 1, 0) -__SYCL_ACCESS(2 <= N && N <= 4, yyy, 1, 1, 1) -__SYCL_ACCESS(N == 3 || N == 4, yyz, 1, 1, 2) -__SYCL_ACCESS(N == 4, yyw, 1, 1, 3) -__SYCL_ACCESS(N == 3 || N == 4, yzx, 1, 2, 0) -__SYCL_ACCESS(N == 3 || N == 4, yzy, 1, 2, 1) -__SYCL_ACCESS(N == 3 || N == 4, yzz, 1, 2, 2) -__SYCL_ACCESS(N == 4, yzw, 1, 2, 3) -__SYCL_ACCESS(N == 4, ywx, 1, 3, 0) -__SYCL_ACCESS(N == 4, ywy, 1, 3, 1) -__SYCL_ACCESS(N == 4, ywz, 1, 3, 2) -__SYCL_ACCESS(N == 4, yww, 1, 3, 3) -__SYCL_ACCESS(N == 3 || N == 4, zxx, 2, 0, 0) -__SYCL_ACCESS(N == 3 || N == 4, zxy, 2, 0, 1) -__SYCL_ACCESS(N == 3 || N == 4, zxz, 2, 0, 2) -__SYCL_ACCESS(N == 4, zxw, 2, 0, 3) -__SYCL_ACCESS(N == 3 || N == 4, zyx, 2, 1, 0) -__SYCL_ACCESS(N == 3 || N == 4, zyy, 2, 1, 1) -__SYCL_ACCESS(N == 3 || N == 4, zyz, 2, 1, 2) -__SYCL_ACCESS(N == 4, zyw, 2, 1, 3) -__SYCL_ACCESS(N == 3 || N == 4, zzx, 2, 2, 0) -__SYCL_ACCESS(N == 3 || N == 4, zzy, 2, 2, 1) -__SYCL_ACCESS(N == 3 || N == 4, zzz, 2, 2, 2) -__SYCL_ACCESS(N == 4, zzw, 2, 2, 3) -__SYCL_ACCESS(N == 4, zwx, 2, 3, 0) -__SYCL_ACCESS(N == 4, zwy, 2, 3, 1) -__SYCL_ACCESS(N == 4, zwz, 2, 3, 2) -__SYCL_ACCESS(N == 4, zww, 2, 3, 3) -__SYCL_ACCESS(N == 4, wxx, 3, 0, 0) -__SYCL_ACCESS(N == 4, wxy, 3, 0, 1) -__SYCL_ACCESS(N == 4, wxz, 3, 0, 2) -__SYCL_ACCESS(N == 4, wxw, 3, 0, 3) -__SYCL_ACCESS(N == 4, wyx, 3, 1, 0) -__SYCL_ACCESS(N == 4, wyy, 3, 1, 1) -__SYCL_ACCESS(N == 4, wyz, 3, 1, 2) -__SYCL_ACCESS(N == 4, wyw, 3, 1, 3) -__SYCL_ACCESS(N == 4, wzx, 3, 2, 0) -__SYCL_ACCESS(N == 4, wzy, 3, 2, 1) -__SYCL_ACCESS(N == 4, wzz, 3, 2, 2) -__SYCL_ACCESS(N == 4, wzw, 3, 2, 3) -__SYCL_ACCESS(N == 4, wwx, 3, 3, 0) -__SYCL_ACCESS(N == 4, wwy, 3, 3, 1) -__SYCL_ACCESS(N == 4, wwz, 3, 3, 2) -__SYCL_ACCESS(N == 4, www, 3, 3, 3) -__SYCL_ACCESS(N <= 4, xxxx, 0, 0, 0, 0) -__SYCL_ACCESS(2 <= N && N <= 4, xxxy, 0, 0, 0, 1) -__SYCL_ACCESS(N == 3 || N == 4, xxxz, 0, 0, 0, 2) -__SYCL_ACCESS(N == 4, xxxw, 0, 0, 0, 3) -__SYCL_ACCESS(2 <= N && N <= 4, xxyx, 0, 0, 1, 0) -__SYCL_ACCESS(2 <= N && N <= 4, xxyy, 0, 0, 1, 1) -__SYCL_ACCESS(N == 3 || N == 4, xxyz, 0, 0, 1, 2) -__SYCL_ACCESS(N == 4, xxyw, 0, 0, 1, 3) -__SYCL_ACCESS(N == 3 || N == 4, xxzx, 0, 0, 2, 0) -__SYCL_ACCESS(N == 3 || N == 4, xxzy, 0, 0, 2, 1) -__SYCL_ACCESS(N == 3 || N == 4, xxzz, 0, 0, 2, 2) -__SYCL_ACCESS(N == 4, xxzw, 0, 0, 2, 3) -__SYCL_ACCESS(N == 4, xxwx, 0, 0, 3, 0) -__SYCL_ACCESS(N == 4, xxwy, 0, 0, 3, 1) -__SYCL_ACCESS(N == 4, xxwz, 0, 0, 3, 2) -__SYCL_ACCESS(N == 4, xxww, 0, 0, 3, 3) -__SYCL_ACCESS(2 <= N && N <= 4, xyxx, 0, 1, 0, 0) -__SYCL_ACCESS(2 <= N && N <= 4, xyxy, 0, 1, 0, 1) -__SYCL_ACCESS(N == 3 || N == 4, xyxz, 0, 1, 0, 2) -__SYCL_ACCESS(N == 4, xyxw, 0, 1, 0, 3) -__SYCL_ACCESS(2 <= N && N <= 4, xyyx, 0, 1, 1, 0) -__SYCL_ACCESS(2 <= N && N <= 4, xyyy, 0, 1, 1, 1) -__SYCL_ACCESS(N == 3 || N == 4, xyyz, 0, 1, 1, 2) -__SYCL_ACCESS(N == 4, xyyw, 0, 1, 1, 3) -__SYCL_ACCESS(N == 3 || N == 4, xyzx, 0, 1, 2, 0) -__SYCL_ACCESS(N == 3 || N == 4, xyzy, 0, 1, 2, 1) -__SYCL_ACCESS(N == 3 || N == 4, xyzz, 0, 1, 2, 2) -__SYCL_ACCESS(N == 4, xyzw, 0, 1, 2, 3) -__SYCL_ACCESS(N == 4, xywx, 0, 1, 3, 0) -__SYCL_ACCESS(N == 4, xywy, 0, 1, 3, 1) -__SYCL_ACCESS(N == 4, xywz, 0, 1, 3, 2) -__SYCL_ACCESS(N == 4, xyww, 0, 1, 3, 3) -__SYCL_ACCESS(N == 3 || N == 4, xzxx, 0, 2, 0, 0) -__SYCL_ACCESS(N == 3 || N == 4, xzxy, 0, 2, 0, 1) -__SYCL_ACCESS(N == 3 || N == 4, xzxz, 0, 2, 0, 2) -__SYCL_ACCESS(N == 4, xzxw, 0, 2, 0, 3) -__SYCL_ACCESS(N == 3 || N == 4, xzyx, 0, 2, 1, 0) -__SYCL_ACCESS(N == 3 || N == 4, xzyy, 0, 2, 1, 1) -__SYCL_ACCESS(N == 3 || N == 4, xzyz, 0, 2, 1, 2) -__SYCL_ACCESS(N == 4, xzyw, 0, 2, 1, 3) -__SYCL_ACCESS(N == 3 || N == 4, xzzx, 0, 2, 2, 0) -__SYCL_ACCESS(N == 3 || N == 4, xzzy, 0, 2, 2, 1) -__SYCL_ACCESS(N == 3 || N == 4, xzzz, 0, 2, 2, 2) -__SYCL_ACCESS(N == 4, xzzw, 0, 2, 2, 3) -__SYCL_ACCESS(N == 4, xzwx, 0, 2, 3, 0) -__SYCL_ACCESS(N == 4, xzwy, 0, 2, 3, 1) -__SYCL_ACCESS(N == 4, xzwz, 0, 2, 3, 2) -__SYCL_ACCESS(N == 4, xzww, 0, 2, 3, 3) -__SYCL_ACCESS(N == 4, xwxx, 0, 3, 0, 0) -__SYCL_ACCESS(N == 4, xwxy, 0, 3, 0, 1) -__SYCL_ACCESS(N == 4, xwxz, 0, 3, 0, 2) -__SYCL_ACCESS(N == 4, xwxw, 0, 3, 0, 3) -__SYCL_ACCESS(N == 4, xwyx, 0, 3, 1, 0) -__SYCL_ACCESS(N == 4, xwyy, 0, 3, 1, 1) -__SYCL_ACCESS(N == 4, xwyz, 0, 3, 1, 2) -__SYCL_ACCESS(N == 4, xwyw, 0, 3, 1, 3) -__SYCL_ACCESS(N == 4, xwzx, 0, 3, 2, 0) -__SYCL_ACCESS(N == 4, xwzy, 0, 3, 2, 1) -__SYCL_ACCESS(N == 4, xwzz, 0, 3, 2, 2) -__SYCL_ACCESS(N == 4, xwzw, 0, 3, 2, 3) -__SYCL_ACCESS(N == 4, xwwx, 0, 3, 3, 0) -__SYCL_ACCESS(N == 4, xwwy, 0, 3, 3, 1) -__SYCL_ACCESS(N == 4, xwwz, 0, 3, 3, 2) -__SYCL_ACCESS(N == 4, xwww, 0, 3, 3, 3) -__SYCL_ACCESS(2 <= N && N <= 4, yxxx, 1, 0, 0, 0) -__SYCL_ACCESS(2 <= N && N <= 4, yxxy, 1, 0, 0, 1) -__SYCL_ACCESS(N == 3 || N == 4, yxxz, 1, 0, 0, 2) -__SYCL_ACCESS(N == 4, yxxw, 1, 0, 0, 3) -__SYCL_ACCESS(2 <= N && N <= 4, yxyx, 1, 0, 1, 0) -__SYCL_ACCESS(2 <= N && N <= 4, yxyy, 1, 0, 1, 1) -__SYCL_ACCESS(N == 3 || N == 4, yxyz, 1, 0, 1, 2) -__SYCL_ACCESS(N == 4, yxyw, 1, 0, 1, 3) -__SYCL_ACCESS(N == 3 || N == 4, yxzx, 1, 0, 2, 0) -__SYCL_ACCESS(N == 3 || N == 4, yxzy, 1, 0, 2, 1) -__SYCL_ACCESS(N == 3 || N == 4, yxzz, 1, 0, 2, 2) -__SYCL_ACCESS(N == 4, yxzw, 1, 0, 2, 3) -__SYCL_ACCESS(N == 4, yxwx, 1, 0, 3, 0) -__SYCL_ACCESS(N == 4, yxwy, 1, 0, 3, 1) -__SYCL_ACCESS(N == 4, yxwz, 1, 0, 3, 2) -__SYCL_ACCESS(N == 4, yxww, 1, 0, 3, 3) -__SYCL_ACCESS(2 <= N && N <= 4, yyxx, 1, 1, 0, 0) -__SYCL_ACCESS(2 <= N && N <= 4, yyxy, 1, 1, 0, 1) -__SYCL_ACCESS(N == 3 || N == 4, yyxz, 1, 1, 0, 2) -__SYCL_ACCESS(N == 4, yyxw, 1, 1, 0, 3) -__SYCL_ACCESS(2 <= N && N <= 4, yyyx, 1, 1, 1, 0) -__SYCL_ACCESS(2 <= N && N <= 4, yyyy, 1, 1, 1, 1) -__SYCL_ACCESS(N == 3 || N == 4, yyyz, 1, 1, 1, 2) -__SYCL_ACCESS(N == 4, yyyw, 1, 1, 1, 3) -__SYCL_ACCESS(N == 3 || N == 4, yyzx, 1, 1, 2, 0) -__SYCL_ACCESS(N == 3 || N == 4, yyzy, 1, 1, 2, 1) -__SYCL_ACCESS(N == 3 || N == 4, yyzz, 1, 1, 2, 2) -__SYCL_ACCESS(N == 4, yyzw, 1, 1, 2, 3) -__SYCL_ACCESS(N == 4, yywx, 1, 1, 3, 0) -__SYCL_ACCESS(N == 4, yywy, 1, 1, 3, 1) -__SYCL_ACCESS(N == 4, yywz, 1, 1, 3, 2) -__SYCL_ACCESS(N == 4, yyww, 1, 1, 3, 3) -__SYCL_ACCESS(N == 3 || N == 4, yzxx, 1, 2, 0, 0) -__SYCL_ACCESS(N == 3 || N == 4, yzxy, 1, 2, 0, 1) -__SYCL_ACCESS(N == 3 || N == 4, yzxz, 1, 2, 0, 2) -__SYCL_ACCESS(N == 4, yzxw, 1, 2, 0, 3) -__SYCL_ACCESS(N == 3 || N == 4, yzyx, 1, 2, 1, 0) -__SYCL_ACCESS(N == 3 || N == 4, yzyy, 1, 2, 1, 1) -__SYCL_ACCESS(N == 3 || N == 4, yzyz, 1, 2, 1, 2) -__SYCL_ACCESS(N == 4, yzyw, 1, 2, 1, 3) -__SYCL_ACCESS(N == 3 || N == 4, yzzx, 1, 2, 2, 0) -__SYCL_ACCESS(N == 3 || N == 4, yzzy, 1, 2, 2, 1) -__SYCL_ACCESS(N == 3 || N == 4, yzzz, 1, 2, 2, 2) -__SYCL_ACCESS(N == 4, yzzw, 1, 2, 2, 3) -__SYCL_ACCESS(N == 4, yzwx, 1, 2, 3, 0) -__SYCL_ACCESS(N == 4, yzwy, 1, 2, 3, 1) -__SYCL_ACCESS(N == 4, yzwz, 1, 2, 3, 2) -__SYCL_ACCESS(N == 4, yzww, 1, 2, 3, 3) -__SYCL_ACCESS(N == 4, ywxx, 1, 3, 0, 0) -__SYCL_ACCESS(N == 4, ywxy, 1, 3, 0, 1) -__SYCL_ACCESS(N == 4, ywxz, 1, 3, 0, 2) -__SYCL_ACCESS(N == 4, ywxw, 1, 3, 0, 3) -__SYCL_ACCESS(N == 4, ywyx, 1, 3, 1, 0) -__SYCL_ACCESS(N == 4, ywyy, 1, 3, 1, 1) -__SYCL_ACCESS(N == 4, ywyz, 1, 3, 1, 2) -__SYCL_ACCESS(N == 4, ywyw, 1, 3, 1, 3) -__SYCL_ACCESS(N == 4, ywzx, 1, 3, 2, 0) -__SYCL_ACCESS(N == 4, ywzy, 1, 3, 2, 1) -__SYCL_ACCESS(N == 4, ywzz, 1, 3, 2, 2) -__SYCL_ACCESS(N == 4, ywzw, 1, 3, 2, 3) -__SYCL_ACCESS(N == 4, ywwx, 1, 3, 3, 0) -__SYCL_ACCESS(N == 4, ywwy, 1, 3, 3, 1) -__SYCL_ACCESS(N == 4, ywwz, 1, 3, 3, 2) -__SYCL_ACCESS(N == 4, ywww, 1, 3, 3, 3) -__SYCL_ACCESS(N == 3 || N == 4, zxxx, 2, 0, 0, 0) -__SYCL_ACCESS(N == 3 || N == 4, zxxy, 2, 0, 0, 1) -__SYCL_ACCESS(N == 3 || N == 4, zxxz, 2, 0, 0, 2) -__SYCL_ACCESS(N == 4, zxxw, 2, 0, 0, 3) -__SYCL_ACCESS(N == 3 || N == 4, zxyx, 2, 0, 1, 0) -__SYCL_ACCESS(N == 3 || N == 4, zxyy, 2, 0, 1, 1) -__SYCL_ACCESS(N == 3 || N == 4, zxyz, 2, 0, 1, 2) -__SYCL_ACCESS(N == 4, zxyw, 2, 0, 1, 3) -__SYCL_ACCESS(N == 3 || N == 4, zxzx, 2, 0, 2, 0) -__SYCL_ACCESS(N == 3 || N == 4, zxzy, 2, 0, 2, 1) -__SYCL_ACCESS(N == 3 || N == 4, zxzz, 2, 0, 2, 2) -__SYCL_ACCESS(N == 4, zxzw, 2, 0, 2, 3) -__SYCL_ACCESS(N == 4, zxwx, 2, 0, 3, 0) -__SYCL_ACCESS(N == 4, zxwy, 2, 0, 3, 1) -__SYCL_ACCESS(N == 4, zxwz, 2, 0, 3, 2) -__SYCL_ACCESS(N == 4, zxww, 2, 0, 3, 3) -__SYCL_ACCESS(N == 3 || N == 4, zyxx, 2, 1, 0, 0) -__SYCL_ACCESS(N == 3 || N == 4, zyxy, 2, 1, 0, 1) -__SYCL_ACCESS(N == 3 || N == 4, zyxz, 2, 1, 0, 2) -__SYCL_ACCESS(N == 4, zyxw, 2, 1, 0, 3) -__SYCL_ACCESS(N == 3 || N == 4, zyyx, 2, 1, 1, 0) -__SYCL_ACCESS(N == 3 || N == 4, zyyy, 2, 1, 1, 1) -__SYCL_ACCESS(N == 3 || N == 4, zyyz, 2, 1, 1, 2) -__SYCL_ACCESS(N == 4, zyyw, 2, 1, 1, 3) -__SYCL_ACCESS(N == 3 || N == 4, zyzx, 2, 1, 2, 0) -__SYCL_ACCESS(N == 3 || N == 4, zyzy, 2, 1, 2, 1) -__SYCL_ACCESS(N == 3 || N == 4, zyzz, 2, 1, 2, 2) -__SYCL_ACCESS(N == 4, zyzw, 2, 1, 2, 3) -__SYCL_ACCESS(N == 4, zywx, 2, 1, 3, 0) -__SYCL_ACCESS(N == 4, zywy, 2, 1, 3, 1) -__SYCL_ACCESS(N == 4, zywz, 2, 1, 3, 2) -__SYCL_ACCESS(N == 4, zyww, 2, 1, 3, 3) -__SYCL_ACCESS(N == 3 || N == 4, zzxx, 2, 2, 0, 0) -__SYCL_ACCESS(N == 3 || N == 4, zzxy, 2, 2, 0, 1) -__SYCL_ACCESS(N == 3 || N == 4, zzxz, 2, 2, 0, 2) -__SYCL_ACCESS(N == 4, zzxw, 2, 2, 0, 3) -__SYCL_ACCESS(N == 3 || N == 4, zzyx, 2, 2, 1, 0) -__SYCL_ACCESS(N == 3 || N == 4, zzyy, 2, 2, 1, 1) -__SYCL_ACCESS(N == 3 || N == 4, zzyz, 2, 2, 1, 2) -__SYCL_ACCESS(N == 4, zzyw, 2, 2, 1, 3) -__SYCL_ACCESS(N == 3 || N == 4, zzzx, 2, 2, 2, 0) -__SYCL_ACCESS(N == 3 || N == 4, zzzy, 2, 2, 2, 1) -__SYCL_ACCESS(N == 3 || N == 4, zzzz, 2, 2, 2, 2) -__SYCL_ACCESS(N == 4, zzzw, 2, 2, 2, 3) -__SYCL_ACCESS(N == 4, zzwx, 2, 2, 3, 0) -__SYCL_ACCESS(N == 4, zzwy, 2, 2, 3, 1) -__SYCL_ACCESS(N == 4, zzwz, 2, 2, 3, 2) -__SYCL_ACCESS(N == 4, zzww, 2, 2, 3, 3) -__SYCL_ACCESS(N == 4, zwxx, 2, 3, 0, 0) -__SYCL_ACCESS(N == 4, zwxy, 2, 3, 0, 1) -__SYCL_ACCESS(N == 4, zwxz, 2, 3, 0, 2) -__SYCL_ACCESS(N == 4, zwxw, 2, 3, 0, 3) -__SYCL_ACCESS(N == 4, zwyx, 2, 3, 1, 0) -__SYCL_ACCESS(N == 4, zwyy, 2, 3, 1, 1) -__SYCL_ACCESS(N == 4, zwyz, 2, 3, 1, 2) -__SYCL_ACCESS(N == 4, zwyw, 2, 3, 1, 3) -__SYCL_ACCESS(N == 4, zwzx, 2, 3, 2, 0) -__SYCL_ACCESS(N == 4, zwzy, 2, 3, 2, 1) -__SYCL_ACCESS(N == 4, zwzz, 2, 3, 2, 2) -__SYCL_ACCESS(N == 4, zwzw, 2, 3, 2, 3) -__SYCL_ACCESS(N == 4, zwwx, 2, 3, 3, 0) -__SYCL_ACCESS(N == 4, zwwy, 2, 3, 3, 1) -__SYCL_ACCESS(N == 4, zwwz, 2, 3, 3, 2) -__SYCL_ACCESS(N == 4, zwww, 2, 3, 3, 3) -__SYCL_ACCESS(N == 4, wxxx, 3, 0, 0, 0) -__SYCL_ACCESS(N == 4, wxxy, 3, 0, 0, 1) -__SYCL_ACCESS(N == 4, wxxz, 3, 0, 0, 2) -__SYCL_ACCESS(N == 4, wxxw, 3, 0, 0, 3) -__SYCL_ACCESS(N == 4, wxyx, 3, 0, 1, 0) -__SYCL_ACCESS(N == 4, wxyy, 3, 0, 1, 1) -__SYCL_ACCESS(N == 4, wxyz, 3, 0, 1, 2) -__SYCL_ACCESS(N == 4, wxyw, 3, 0, 1, 3) -__SYCL_ACCESS(N == 4, wxzx, 3, 0, 2, 0) -__SYCL_ACCESS(N == 4, wxzy, 3, 0, 2, 1) -__SYCL_ACCESS(N == 4, wxzz, 3, 0, 2, 2) -__SYCL_ACCESS(N == 4, wxzw, 3, 0, 2, 3) -__SYCL_ACCESS(N == 4, wxwx, 3, 0, 3, 0) -__SYCL_ACCESS(N == 4, wxwy, 3, 0, 3, 1) -__SYCL_ACCESS(N == 4, wxwz, 3, 0, 3, 2) -__SYCL_ACCESS(N == 4, wxww, 3, 0, 3, 3) -__SYCL_ACCESS(N == 4, wyxx, 3, 1, 0, 0) -__SYCL_ACCESS(N == 4, wyxy, 3, 1, 0, 1) -__SYCL_ACCESS(N == 4, wyxz, 3, 1, 0, 2) -__SYCL_ACCESS(N == 4, wyxw, 3, 1, 0, 3) -__SYCL_ACCESS(N == 4, wyyx, 3, 1, 1, 0) -__SYCL_ACCESS(N == 4, wyyy, 3, 1, 1, 1) -__SYCL_ACCESS(N == 4, wyyz, 3, 1, 1, 2) -__SYCL_ACCESS(N == 4, wyyw, 3, 1, 1, 3) -__SYCL_ACCESS(N == 4, wyzx, 3, 1, 2, 0) -__SYCL_ACCESS(N == 4, wyzy, 3, 1, 2, 1) -__SYCL_ACCESS(N == 4, wyzz, 3, 1, 2, 2) -__SYCL_ACCESS(N == 4, wyzw, 3, 1, 2, 3) -__SYCL_ACCESS(N == 4, wywx, 3, 1, 3, 0) -__SYCL_ACCESS(N == 4, wywy, 3, 1, 3, 1) -__SYCL_ACCESS(N == 4, wywz, 3, 1, 3, 2) -__SYCL_ACCESS(N == 4, wyww, 3, 1, 3, 3) -__SYCL_ACCESS(N == 4, wzxx, 3, 2, 0, 0) -__SYCL_ACCESS(N == 4, wzxy, 3, 2, 0, 1) -__SYCL_ACCESS(N == 4, wzxz, 3, 2, 0, 2) -__SYCL_ACCESS(N == 4, wzxw, 3, 2, 0, 3) -__SYCL_ACCESS(N == 4, wzyx, 3, 2, 1, 0) -__SYCL_ACCESS(N == 4, wzyy, 3, 2, 1, 1) -__SYCL_ACCESS(N == 4, wzyz, 3, 2, 1, 2) -__SYCL_ACCESS(N == 4, wzyw, 3, 2, 1, 3) -__SYCL_ACCESS(N == 4, wzzx, 3, 2, 2, 0) -__SYCL_ACCESS(N == 4, wzzy, 3, 2, 2, 1) -__SYCL_ACCESS(N == 4, wzzz, 3, 2, 2, 2) -__SYCL_ACCESS(N == 4, wzzw, 3, 2, 2, 3) -__SYCL_ACCESS(N == 4, wzwx, 3, 2, 3, 0) -__SYCL_ACCESS(N == 4, wzwy, 3, 2, 3, 1) -__SYCL_ACCESS(N == 4, wzwz, 3, 2, 3, 2) -__SYCL_ACCESS(N == 4, wzww, 3, 2, 3, 3) -__SYCL_ACCESS(N == 4, wwxx, 3, 3, 0, 0) -__SYCL_ACCESS(N == 4, wwxy, 3, 3, 0, 1) -__SYCL_ACCESS(N == 4, wwxz, 3, 3, 0, 2) -__SYCL_ACCESS(N == 4, wwxw, 3, 3, 0, 3) -__SYCL_ACCESS(N == 4, wwyx, 3, 3, 1, 0) -__SYCL_ACCESS(N == 4, wwyy, 3, 3, 1, 1) -__SYCL_ACCESS(N == 4, wwyz, 3, 3, 1, 2) -__SYCL_ACCESS(N == 4, wwyw, 3, 3, 1, 3) -__SYCL_ACCESS(N == 4, wwzx, 3, 3, 2, 0) -__SYCL_ACCESS(N == 4, wwzy, 3, 3, 2, 1) -__SYCL_ACCESS(N == 4, wwzz, 3, 3, 2, 2) -__SYCL_ACCESS(N == 4, wwzw, 3, 3, 2, 3) -__SYCL_ACCESS(N == 4, wwwx, 3, 3, 3, 0) -__SYCL_ACCESS(N == 4, wwwy, 3, 3, 3, 1) -__SYCL_ACCESS(N == 4, wwwz, 3, 3, 3, 2) -__SYCL_ACCESS(N == 4, wwww, 3, 3, 3, 3) - -//__swizzled_vec__ RGBA_SWIZZLE() const; -__SYCL_ACCESS(N == 4, rr, 0, 0) -__SYCL_ACCESS(N == 4, rg, 0, 1) -__SYCL_ACCESS(N == 4, rb, 0, 2) -__SYCL_ACCESS(N == 4, ra, 0, 3) -__SYCL_ACCESS(N == 4, gr, 1, 0) -__SYCL_ACCESS(N == 4, gg, 1, 1) -__SYCL_ACCESS(N == 4, gb, 1, 2) -__SYCL_ACCESS(N == 4, ga, 1, 3) -__SYCL_ACCESS(N == 4, br, 2, 0) -__SYCL_ACCESS(N == 4, bg, 2, 1) -__SYCL_ACCESS(N == 4, bb, 2, 2) -__SYCL_ACCESS(N == 4, ba, 2, 3) -__SYCL_ACCESS(N == 4, ar, 3, 0) -__SYCL_ACCESS(N == 4, ag, 3, 1) -__SYCL_ACCESS(N == 4, ab, 3, 2) -__SYCL_ACCESS(N == 4, aa, 3, 3) -__SYCL_ACCESS(N == 4, rrr, 0, 0, 0) -__SYCL_ACCESS(N == 4, rrg, 0, 0, 1) -__SYCL_ACCESS(N == 4, rrb, 0, 0, 2) -__SYCL_ACCESS(N == 4, rra, 0, 0, 3) -__SYCL_ACCESS(N == 4, rgr, 0, 1, 0) -__SYCL_ACCESS(N == 4, rgg, 0, 1, 1) -__SYCL_ACCESS(N == 4, rgb, 0, 1, 2) -__SYCL_ACCESS(N == 4, rga, 0, 1, 3) -__SYCL_ACCESS(N == 4, rbr, 0, 2, 0) -__SYCL_ACCESS(N == 4, rbg, 0, 2, 1) -__SYCL_ACCESS(N == 4, rbb, 0, 2, 2) -__SYCL_ACCESS(N == 4, rba, 0, 2, 3) -__SYCL_ACCESS(N == 4, rar, 0, 3, 0) -__SYCL_ACCESS(N == 4, rag, 0, 3, 1) -__SYCL_ACCESS(N == 4, rab, 0, 3, 2) -__SYCL_ACCESS(N == 4, raa, 0, 3, 3) -__SYCL_ACCESS(N == 4, grr, 1, 0, 0) -__SYCL_ACCESS(N == 4, grg, 1, 0, 1) -__SYCL_ACCESS(N == 4, grb, 1, 0, 2) -__SYCL_ACCESS(N == 4, gra, 1, 0, 3) -__SYCL_ACCESS(N == 4, ggr, 1, 1, 0) -__SYCL_ACCESS(N == 4, ggg, 1, 1, 1) -__SYCL_ACCESS(N == 4, ggb, 1, 1, 2) -__SYCL_ACCESS(N == 4, gga, 1, 1, 3) -__SYCL_ACCESS(N == 4, gbr, 1, 2, 0) -__SYCL_ACCESS(N == 4, gbg, 1, 2, 1) -__SYCL_ACCESS(N == 4, gbb, 1, 2, 2) -__SYCL_ACCESS(N == 4, gba, 1, 2, 3) -__SYCL_ACCESS(N == 4, gar, 1, 3, 0) -__SYCL_ACCESS(N == 4, gag, 1, 3, 1) -__SYCL_ACCESS(N == 4, gab, 1, 3, 2) -__SYCL_ACCESS(N == 4, gaa, 1, 3, 3) -__SYCL_ACCESS(N == 4, brr, 2, 0, 0) -__SYCL_ACCESS(N == 4, brg, 2, 0, 1) -__SYCL_ACCESS(N == 4, brb, 2, 0, 2) -__SYCL_ACCESS(N == 4, bra, 2, 0, 3) -__SYCL_ACCESS(N == 4, bgr, 2, 1, 0) -__SYCL_ACCESS(N == 4, bgg, 2, 1, 1) -__SYCL_ACCESS(N == 4, bgb, 2, 1, 2) -__SYCL_ACCESS(N == 4, bga, 2, 1, 3) -__SYCL_ACCESS(N == 4, bbr, 2, 2, 0) -__SYCL_ACCESS(N == 4, bbg, 2, 2, 1) -__SYCL_ACCESS(N == 4, bbb, 2, 2, 2) -__SYCL_ACCESS(N == 4, bba, 2, 2, 3) -__SYCL_ACCESS(N == 4, bar, 2, 3, 0) -__SYCL_ACCESS(N == 4, bag, 2, 3, 1) -__SYCL_ACCESS(N == 4, bab, 2, 3, 2) -__SYCL_ACCESS(N == 4, baa, 2, 3, 3) -__SYCL_ACCESS(N == 4, arr, 3, 0, 0) -__SYCL_ACCESS(N == 4, arg, 3, 0, 1) -__SYCL_ACCESS(N == 4, arb, 3, 0, 2) -__SYCL_ACCESS(N == 4, ara, 3, 0, 3) -__SYCL_ACCESS(N == 4, agr, 3, 1, 0) -__SYCL_ACCESS(N == 4, agg, 3, 1, 1) -__SYCL_ACCESS(N == 4, agb, 3, 1, 2) -__SYCL_ACCESS(N == 4, aga, 3, 1, 3) -__SYCL_ACCESS(N == 4, abr, 3, 2, 0) -__SYCL_ACCESS(N == 4, abg, 3, 2, 1) -__SYCL_ACCESS(N == 4, abb, 3, 2, 2) -__SYCL_ACCESS(N == 4, aba, 3, 2, 3) -__SYCL_ACCESS(N == 4, aar, 3, 3, 0) -__SYCL_ACCESS(N == 4, aag, 3, 3, 1) -__SYCL_ACCESS(N == 4, aab, 3, 3, 2) -__SYCL_ACCESS(N == 4, aaa, 3, 3, 3) -__SYCL_ACCESS(N == 4, rrrr, 0, 0, 0, 0) -__SYCL_ACCESS(N == 4, rrrg, 0, 0, 0, 1) -__SYCL_ACCESS(N == 4, rrrb, 0, 0, 0, 2) -__SYCL_ACCESS(N == 4, rrra, 0, 0, 0, 3) -__SYCL_ACCESS(N == 4, rrgr, 0, 0, 1, 0) -__SYCL_ACCESS(N == 4, rrgg, 0, 0, 1, 1) -__SYCL_ACCESS(N == 4, rrgb, 0, 0, 1, 2) -__SYCL_ACCESS(N == 4, rrga, 0, 0, 1, 3) -__SYCL_ACCESS(N == 4, rrbr, 0, 0, 2, 0) -__SYCL_ACCESS(N == 4, rrbg, 0, 0, 2, 1) -__SYCL_ACCESS(N == 4, rrbb, 0, 0, 2, 2) -__SYCL_ACCESS(N == 4, rrba, 0, 0, 2, 3) -__SYCL_ACCESS(N == 4, rrar, 0, 0, 3, 0) -__SYCL_ACCESS(N == 4, rrag, 0, 0, 3, 1) -__SYCL_ACCESS(N == 4, rrab, 0, 0, 3, 2) -__SYCL_ACCESS(N == 4, rraa, 0, 0, 3, 3) -__SYCL_ACCESS(N == 4, rgrr, 0, 1, 0, 0) -__SYCL_ACCESS(N == 4, rgrg, 0, 1, 0, 1) -__SYCL_ACCESS(N == 4, rgrb, 0, 1, 0, 2) -__SYCL_ACCESS(N == 4, rgra, 0, 1, 0, 3) -__SYCL_ACCESS(N == 4, rggr, 0, 1, 1, 0) -__SYCL_ACCESS(N == 4, rggg, 0, 1, 1, 1) -__SYCL_ACCESS(N == 4, rggb, 0, 1, 1, 2) -__SYCL_ACCESS(N == 4, rgga, 0, 1, 1, 3) -__SYCL_ACCESS(N == 4, rgbr, 0, 1, 2, 0) -__SYCL_ACCESS(N == 4, rgbg, 0, 1, 2, 1) -__SYCL_ACCESS(N == 4, rgbb, 0, 1, 2, 2) -__SYCL_ACCESS(N == 4, rgba, 0, 1, 2, 3) -__SYCL_ACCESS(N == 4, rgar, 0, 1, 3, 0) -__SYCL_ACCESS(N == 4, rgag, 0, 1, 3, 1) -__SYCL_ACCESS(N == 4, rgab, 0, 1, 3, 2) -__SYCL_ACCESS(N == 4, rgaa, 0, 1, 3, 3) -__SYCL_ACCESS(N == 4, rbrr, 0, 2, 0, 0) -__SYCL_ACCESS(N == 4, rbrg, 0, 2, 0, 1) -__SYCL_ACCESS(N == 4, rbrb, 0, 2, 0, 2) -__SYCL_ACCESS(N == 4, rbra, 0, 2, 0, 3) -__SYCL_ACCESS(N == 4, rbgr, 0, 2, 1, 0) -__SYCL_ACCESS(N == 4, rbgg, 0, 2, 1, 1) -__SYCL_ACCESS(N == 4, rbgb, 0, 2, 1, 2) -__SYCL_ACCESS(N == 4, rbga, 0, 2, 1, 3) -__SYCL_ACCESS(N == 4, rbbr, 0, 2, 2, 0) -__SYCL_ACCESS(N == 4, rbbg, 0, 2, 2, 1) -__SYCL_ACCESS(N == 4, rbbb, 0, 2, 2, 2) -__SYCL_ACCESS(N == 4, rbba, 0, 2, 2, 3) -__SYCL_ACCESS(N == 4, rbar, 0, 2, 3, 0) -__SYCL_ACCESS(N == 4, rbag, 0, 2, 3, 1) -__SYCL_ACCESS(N == 4, rbab, 0, 2, 3, 2) -__SYCL_ACCESS(N == 4, rbaa, 0, 2, 3, 3) -__SYCL_ACCESS(N == 4, rarr, 0, 3, 0, 0) -__SYCL_ACCESS(N == 4, rarg, 0, 3, 0, 1) -__SYCL_ACCESS(N == 4, rarb, 0, 3, 0, 2) -__SYCL_ACCESS(N == 4, rara, 0, 3, 0, 3) -__SYCL_ACCESS(N == 4, ragr, 0, 3, 1, 0) -__SYCL_ACCESS(N == 4, ragg, 0, 3, 1, 1) -__SYCL_ACCESS(N == 4, ragb, 0, 3, 1, 2) -__SYCL_ACCESS(N == 4, raga, 0, 3, 1, 3) -__SYCL_ACCESS(N == 4, rabr, 0, 3, 2, 0) -__SYCL_ACCESS(N == 4, rabg, 0, 3, 2, 1) -__SYCL_ACCESS(N == 4, rabb, 0, 3, 2, 2) -__SYCL_ACCESS(N == 4, raba, 0, 3, 2, 3) -__SYCL_ACCESS(N == 4, raar, 0, 3, 3, 0) -__SYCL_ACCESS(N == 4, raag, 0, 3, 3, 1) -__SYCL_ACCESS(N == 4, raab, 0, 3, 3, 2) -__SYCL_ACCESS(N == 4, raaa, 0, 3, 3, 3) -__SYCL_ACCESS(N == 4, grrr, 1, 0, 0, 0) -__SYCL_ACCESS(N == 4, grrg, 1, 0, 0, 1) -__SYCL_ACCESS(N == 4, grrb, 1, 0, 0, 2) -__SYCL_ACCESS(N == 4, grra, 1, 0, 0, 3) -__SYCL_ACCESS(N == 4, grgr, 1, 0, 1, 0) -__SYCL_ACCESS(N == 4, grgg, 1, 0, 1, 1) -__SYCL_ACCESS(N == 4, grgb, 1, 0, 1, 2) -__SYCL_ACCESS(N == 4, grga, 1, 0, 1, 3) -__SYCL_ACCESS(N == 4, grbr, 1, 0, 2, 0) -__SYCL_ACCESS(N == 4, grbg, 1, 0, 2, 1) -__SYCL_ACCESS(N == 4, grbb, 1, 0, 2, 2) -__SYCL_ACCESS(N == 4, grba, 1, 0, 2, 3) -__SYCL_ACCESS(N == 4, grar, 1, 0, 3, 0) -__SYCL_ACCESS(N == 4, grag, 1, 0, 3, 1) -__SYCL_ACCESS(N == 4, grab, 1, 0, 3, 2) -__SYCL_ACCESS(N == 4, graa, 1, 0, 3, 3) -__SYCL_ACCESS(N == 4, ggrr, 1, 1, 0, 0) -__SYCL_ACCESS(N == 4, ggrg, 1, 1, 0, 1) -__SYCL_ACCESS(N == 4, ggrb, 1, 1, 0, 2) -__SYCL_ACCESS(N == 4, ggra, 1, 1, 0, 3) -__SYCL_ACCESS(N == 4, gggr, 1, 1, 1, 0) -__SYCL_ACCESS(N == 4, gggg, 1, 1, 1, 1) -__SYCL_ACCESS(N == 4, gggb, 1, 1, 1, 2) -__SYCL_ACCESS(N == 4, ggga, 1, 1, 1, 3) -__SYCL_ACCESS(N == 4, ggbr, 1, 1, 2, 0) -__SYCL_ACCESS(N == 4, ggbg, 1, 1, 2, 1) -__SYCL_ACCESS(N == 4, ggbb, 1, 1, 2, 2) -__SYCL_ACCESS(N == 4, ggba, 1, 1, 2, 3) -__SYCL_ACCESS(N == 4, ggar, 1, 1, 3, 0) -__SYCL_ACCESS(N == 4, ggag, 1, 1, 3, 1) -__SYCL_ACCESS(N == 4, ggab, 1, 1, 3, 2) -__SYCL_ACCESS(N == 4, ggaa, 1, 1, 3, 3) -__SYCL_ACCESS(N == 4, gbrr, 1, 2, 0, 0) -__SYCL_ACCESS(N == 4, gbrg, 1, 2, 0, 1) -__SYCL_ACCESS(N == 4, gbrb, 1, 2, 0, 2) -__SYCL_ACCESS(N == 4, gbra, 1, 2, 0, 3) -__SYCL_ACCESS(N == 4, gbgr, 1, 2, 1, 0) -__SYCL_ACCESS(N == 4, gbgg, 1, 2, 1, 1) -__SYCL_ACCESS(N == 4, gbgb, 1, 2, 1, 2) -__SYCL_ACCESS(N == 4, gbga, 1, 2, 1, 3) -__SYCL_ACCESS(N == 4, gbbr, 1, 2, 2, 0) -__SYCL_ACCESS(N == 4, gbbg, 1, 2, 2, 1) -__SYCL_ACCESS(N == 4, gbbb, 1, 2, 2, 2) -__SYCL_ACCESS(N == 4, gbba, 1, 2, 2, 3) -__SYCL_ACCESS(N == 4, gbar, 1, 2, 3, 0) -__SYCL_ACCESS(N == 4, gbag, 1, 2, 3, 1) -__SYCL_ACCESS(N == 4, gbab, 1, 2, 3, 2) -__SYCL_ACCESS(N == 4, gbaa, 1, 2, 3, 3) -__SYCL_ACCESS(N == 4, garr, 1, 3, 0, 0) -__SYCL_ACCESS(N == 4, garg, 1, 3, 0, 1) -__SYCL_ACCESS(N == 4, garb, 1, 3, 0, 2) -__SYCL_ACCESS(N == 4, gara, 1, 3, 0, 3) -__SYCL_ACCESS(N == 4, gagr, 1, 3, 1, 0) -__SYCL_ACCESS(N == 4, gagg, 1, 3, 1, 1) -__SYCL_ACCESS(N == 4, gagb, 1, 3, 1, 2) -__SYCL_ACCESS(N == 4, gaga, 1, 3, 1, 3) -__SYCL_ACCESS(N == 4, gabr, 1, 3, 2, 0) -__SYCL_ACCESS(N == 4, gabg, 1, 3, 2, 1) -__SYCL_ACCESS(N == 4, gabb, 1, 3, 2, 2) -__SYCL_ACCESS(N == 4, gaba, 1, 3, 2, 3) -__SYCL_ACCESS(N == 4, gaar, 1, 3, 3, 0) -__SYCL_ACCESS(N == 4, gaag, 1, 3, 3, 1) -__SYCL_ACCESS(N == 4, gaab, 1, 3, 3, 2) -__SYCL_ACCESS(N == 4, gaaa, 1, 3, 3, 3) -__SYCL_ACCESS(N == 4, brrr, 2, 0, 0, 0) -__SYCL_ACCESS(N == 4, brrg, 2, 0, 0, 1) -__SYCL_ACCESS(N == 4, brrb, 2, 0, 0, 2) -__SYCL_ACCESS(N == 4, brra, 2, 0, 0, 3) -__SYCL_ACCESS(N == 4, brgr, 2, 0, 1, 0) -__SYCL_ACCESS(N == 4, brgg, 2, 0, 1, 1) -__SYCL_ACCESS(N == 4, brgb, 2, 0, 1, 2) -__SYCL_ACCESS(N == 4, brga, 2, 0, 1, 3) -__SYCL_ACCESS(N == 4, brbr, 2, 0, 2, 0) -__SYCL_ACCESS(N == 4, brbg, 2, 0, 2, 1) -__SYCL_ACCESS(N == 4, brbb, 2, 0, 2, 2) -__SYCL_ACCESS(N == 4, brba, 2, 0, 2, 3) -__SYCL_ACCESS(N == 4, brar, 2, 0, 3, 0) -__SYCL_ACCESS(N == 4, brag, 2, 0, 3, 1) -__SYCL_ACCESS(N == 4, brab, 2, 0, 3, 2) -__SYCL_ACCESS(N == 4, braa, 2, 0, 3, 3) -__SYCL_ACCESS(N == 4, bgrr, 2, 1, 0, 0) -__SYCL_ACCESS(N == 4, bgrg, 2, 1, 0, 1) -__SYCL_ACCESS(N == 4, bgrb, 2, 1, 0, 2) -__SYCL_ACCESS(N == 4, bgra, 2, 1, 0, 3) -__SYCL_ACCESS(N == 4, bggr, 2, 1, 1, 0) -__SYCL_ACCESS(N == 4, bggg, 2, 1, 1, 1) -__SYCL_ACCESS(N == 4, bggb, 2, 1, 1, 2) -__SYCL_ACCESS(N == 4, bgga, 2, 1, 1, 3) -__SYCL_ACCESS(N == 4, bgbr, 2, 1, 2, 0) -__SYCL_ACCESS(N == 4, bgbg, 2, 1, 2, 1) -__SYCL_ACCESS(N == 4, bgbb, 2, 1, 2, 2) -__SYCL_ACCESS(N == 4, bgba, 2, 1, 2, 3) -__SYCL_ACCESS(N == 4, bgar, 2, 1, 3, 0) -__SYCL_ACCESS(N == 4, bgag, 2, 1, 3, 1) -__SYCL_ACCESS(N == 4, bgab, 2, 1, 3, 2) -__SYCL_ACCESS(N == 4, bgaa, 2, 1, 3, 3) -__SYCL_ACCESS(N == 4, bbrr, 2, 2, 0, 0) -__SYCL_ACCESS(N == 4, bbrg, 2, 2, 0, 1) -__SYCL_ACCESS(N == 4, bbrb, 2, 2, 0, 2) -__SYCL_ACCESS(N == 4, bbra, 2, 2, 0, 3) -__SYCL_ACCESS(N == 4, bbgr, 2, 2, 1, 0) -__SYCL_ACCESS(N == 4, bbgg, 2, 2, 1, 1) -__SYCL_ACCESS(N == 4, bbgb, 2, 2, 1, 2) -__SYCL_ACCESS(N == 4, bbga, 2, 2, 1, 3) -__SYCL_ACCESS(N == 4, bbbr, 2, 2, 2, 0) -__SYCL_ACCESS(N == 4, bbbg, 2, 2, 2, 1) -__SYCL_ACCESS(N == 4, bbbb, 2, 2, 2, 2) -__SYCL_ACCESS(N == 4, bbba, 2, 2, 2, 3) -__SYCL_ACCESS(N == 4, bbar, 2, 2, 3, 0) -__SYCL_ACCESS(N == 4, bbag, 2, 2, 3, 1) -__SYCL_ACCESS(N == 4, bbab, 2, 2, 3, 2) -__SYCL_ACCESS(N == 4, bbaa, 2, 2, 3, 3) -__SYCL_ACCESS(N == 4, barr, 2, 3, 0, 0) -__SYCL_ACCESS(N == 4, barg, 2, 3, 0, 1) -__SYCL_ACCESS(N == 4, barb, 2, 3, 0, 2) -__SYCL_ACCESS(N == 4, bara, 2, 3, 0, 3) -__SYCL_ACCESS(N == 4, bagr, 2, 3, 1, 0) -__SYCL_ACCESS(N == 4, bagg, 2, 3, 1, 1) -__SYCL_ACCESS(N == 4, bagb, 2, 3, 1, 2) -__SYCL_ACCESS(N == 4, baga, 2, 3, 1, 3) -__SYCL_ACCESS(N == 4, babr, 2, 3, 2, 0) -__SYCL_ACCESS(N == 4, babg, 2, 3, 2, 1) -__SYCL_ACCESS(N == 4, babb, 2, 3, 2, 2) -__SYCL_ACCESS(N == 4, baba, 2, 3, 2, 3) -__SYCL_ACCESS(N == 4, baar, 2, 3, 3, 0) -__SYCL_ACCESS(N == 4, baag, 2, 3, 3, 1) -__SYCL_ACCESS(N == 4, baab, 2, 3, 3, 2) -__SYCL_ACCESS(N == 4, baaa, 2, 3, 3, 3) -__SYCL_ACCESS(N == 4, arrr, 3, 0, 0, 0) -__SYCL_ACCESS(N == 4, arrg, 3, 0, 0, 1) -__SYCL_ACCESS(N == 4, arrb, 3, 0, 0, 2) -__SYCL_ACCESS(N == 4, arra, 3, 0, 0, 3) -__SYCL_ACCESS(N == 4, argr, 3, 0, 1, 0) -__SYCL_ACCESS(N == 4, argg, 3, 0, 1, 1) -__SYCL_ACCESS(N == 4, argb, 3, 0, 1, 2) -__SYCL_ACCESS(N == 4, arga, 3, 0, 1, 3) -__SYCL_ACCESS(N == 4, arbr, 3, 0, 2, 0) -__SYCL_ACCESS(N == 4, arbg, 3, 0, 2, 1) -__SYCL_ACCESS(N == 4, arbb, 3, 0, 2, 2) -__SYCL_ACCESS(N == 4, arba, 3, 0, 2, 3) -__SYCL_ACCESS(N == 4, arar, 3, 0, 3, 0) -__SYCL_ACCESS(N == 4, arag, 3, 0, 3, 1) -__SYCL_ACCESS(N == 4, arab, 3, 0, 3, 2) -__SYCL_ACCESS(N == 4, araa, 3, 0, 3, 3) -__SYCL_ACCESS(N == 4, agrr, 3, 1, 0, 0) -__SYCL_ACCESS(N == 4, agrg, 3, 1, 0, 1) -__SYCL_ACCESS(N == 4, agrb, 3, 1, 0, 2) -__SYCL_ACCESS(N == 4, agra, 3, 1, 0, 3) -__SYCL_ACCESS(N == 4, aggr, 3, 1, 1, 0) -__SYCL_ACCESS(N == 4, aggg, 3, 1, 1, 1) -__SYCL_ACCESS(N == 4, aggb, 3, 1, 1, 2) -__SYCL_ACCESS(N == 4, agga, 3, 1, 1, 3) -__SYCL_ACCESS(N == 4, agbr, 3, 1, 2, 0) -__SYCL_ACCESS(N == 4, agbg, 3, 1, 2, 1) -__SYCL_ACCESS(N == 4, agbb, 3, 1, 2, 2) -__SYCL_ACCESS(N == 4, agba, 3, 1, 2, 3) -__SYCL_ACCESS(N == 4, agar, 3, 1, 3, 0) -__SYCL_ACCESS(N == 4, agag, 3, 1, 3, 1) -__SYCL_ACCESS(N == 4, agab, 3, 1, 3, 2) -__SYCL_ACCESS(N == 4, agaa, 3, 1, 3, 3) -__SYCL_ACCESS(N == 4, abrr, 3, 2, 0, 0) -__SYCL_ACCESS(N == 4, abrg, 3, 2, 0, 1) -__SYCL_ACCESS(N == 4, abrb, 3, 2, 0, 2) -__SYCL_ACCESS(N == 4, abra, 3, 2, 0, 3) -__SYCL_ACCESS(N == 4, abgr, 3, 2, 1, 0) -__SYCL_ACCESS(N == 4, abgg, 3, 2, 1, 1) -__SYCL_ACCESS(N == 4, abgb, 3, 2, 1, 2) -__SYCL_ACCESS(N == 4, abga, 3, 2, 1, 3) -__SYCL_ACCESS(N == 4, abbr, 3, 2, 2, 0) -__SYCL_ACCESS(N == 4, abbg, 3, 2, 2, 1) -__SYCL_ACCESS(N == 4, abbb, 3, 2, 2, 2) -__SYCL_ACCESS(N == 4, abba, 3, 2, 2, 3) -__SYCL_ACCESS(N == 4, abar, 3, 2, 3, 0) -__SYCL_ACCESS(N == 4, abag, 3, 2, 3, 1) -__SYCL_ACCESS(N == 4, abab, 3, 2, 3, 2) -__SYCL_ACCESS(N == 4, abaa, 3, 2, 3, 3) -__SYCL_ACCESS(N == 4, aarr, 3, 3, 0, 0) -__SYCL_ACCESS(N == 4, aarg, 3, 3, 0, 1) -__SYCL_ACCESS(N == 4, aarb, 3, 3, 0, 2) -__SYCL_ACCESS(N == 4, aara, 3, 3, 0, 3) -__SYCL_ACCESS(N == 4, aagr, 3, 3, 1, 0) -__SYCL_ACCESS(N == 4, aagg, 3, 3, 1, 1) -__SYCL_ACCESS(N == 4, aagb, 3, 3, 1, 2) -__SYCL_ACCESS(N == 4, aaga, 3, 3, 1, 3) -__SYCL_ACCESS(N == 4, aabr, 3, 3, 2, 0) -__SYCL_ACCESS(N == 4, aabg, 3, 3, 2, 1) -__SYCL_ACCESS(N == 4, aabb, 3, 3, 2, 2) -__SYCL_ACCESS(N == 4, aaba, 3, 3, 2, 3) -__SYCL_ACCESS(N == 4, aaar, 3, 3, 3, 0) -__SYCL_ACCESS(N == 4, aaag, 3, 3, 3, 1) -__SYCL_ACCESS(N == 4, aaab, 3, 3, 3, 2) -__SYCL_ACCESS(N == 4, aaaa, 3, 3, 3, 3) - -#endif // #ifdef SYCL_SIMPLE_SWIZZLES - -//__swizzled_vec__ lo()/hi() const; -__SYCL_ACCESS(N == 2, lo, 0) -__SYCL_ACCESS(N == 3, lo, 0, 1) -__SYCL_ACCESS(N == 4, lo, 0, 1) -__SYCL_ACCESS(N == 8, lo, 0, 1, 2, 3) -__SYCL_ACCESS(N == 16, lo, 0, 1, 2, 3, 4, 5, 6, 7) -__SYCL_ACCESS(N == 2, hi, 1) -__SYCL_ACCESS(N == 3, hi, 2, 3) -__SYCL_ACCESS(N == 4, hi, 2, 3) -__SYCL_ACCESS(N == 8, hi, 4, 5, 6, 7) -__SYCL_ACCESS(N == 16, hi, 8, 9, 10, 11, 12, 13, 14, 15) -//__swizzled_vec__ odd()/even() const; -__SYCL_ACCESS(N == 2, odd, 1) -__SYCL_ACCESS(N == 3, odd, 1, 3) -__SYCL_ACCESS(N == 4, odd, 1, 3) -__SYCL_ACCESS(N == 8, odd, 1, 3, 5, 7) -__SYCL_ACCESS(N == 16, odd, 1, 3, 5, 7, 9, 11, 13, 15) -__SYCL_ACCESS(N == 2, even, 0) -__SYCL_ACCESS(N == 3, even, 0, 2) -__SYCL_ACCESS(N == 4, even, 0, 2) -__SYCL_ACCESS(N == 8, even, 0, 2, 4, 6) -__SYCL_ACCESS(N == 16, even, 0, 2, 4, 6, 8, 10, 12, 14) - -#undef __SYCL_E0 -#undef __SYCL_E1 -#undef __SYCL_E2 -#undef __SYCL_E3 -#undef __SYCL_E4 -#undef __SYCL_E5 -#undef __SYCL_E6 -#undef __SYCL_E7 -#undef __SYCL_E8 - -#undef __SYCL_NTH_ARG -#undef __SYCL_EXPAND -#undef __SYCL_INDEXER -#undef __SYCL_ACCESS -#undef __SYCL_SCALAR_ACCESS diff --git a/sycl/include/sycl/vector.hpp b/sycl/include/sycl/vector.hpp index 75eb80cac829c..82ce6b676eb75 100644 --- a/sycl/include/sycl/vector.hpp +++ b/sycl/include/sycl/vector.hpp @@ -32,9 +32,10 @@ #include // for __SYCL2020_DEPRECATED #include // for is_sigeninteger, is_s... #include // for memcpy -#include // for is_floating_point +#include +#include // for is_floating_point #include -#include // for StorageT, half, Vec16... +#include // for StorageT, half, Vec16... #include // bfloat16 @@ -134,7 +135,9 @@ template class __SYCL_EBO vec : public detail::vec_arith, public detail::ScalarConversionOperatorMixIn, - DataT, NumElements> { + DataT, NumElements>, + public detail::NamedSwizzlesMixinBoth, + NumElements> { static_assert(std::is_same_v>, "DataT must be cv-unqualified"); @@ -379,26 +382,6 @@ class __SYCL_EBO vec DataT &operator[](int i) { return m_Data[i]; } - // Begin hi/lo, even/odd, xyzw, and rgba swizzles. @{ -private: - // Indexer used in the swizzles.def - // Currently it is defined as a template struct. Replacing it with a constexpr - // function would activate a bug in MSVC that is fixed only in v19.20. - // Until then MSVC does not recognize such constexpr functions as const and - // thus does not let using them in template parameters inside swizzle.def. - template struct Indexer { - static constexpr int value = Index; - }; - -public: -#ifdef __SYCL_ACCESS_RETURN -#error "Undefine __SYCL_ACCESS_RETURN macro" -#endif -#define __SYCL_ACCESS_RETURN this -#include "swizzles.def" -#undef __SYCL_ACCESS_RETURN - // }@ End of hi/lo, even/odd, xyzw, and rgba swizzles. - template void load(size_t Offset, multi_ptr Ptr) { for (int I = 0; I < NumElements; I++) { @@ -560,7 +543,10 @@ template struct LShift { // Actual computation performed on conversion or assignment operators. template class OperationCurrentT, int... Indexes> -class SwizzleOp { +class SwizzleOp : public detail::NamedSwizzlesMixinBoth< + SwizzleOp, + sizeof...(Indexes)> { using DataT = typename VecT::element_type; // Certain operators return a vector with a different element type. Also, the // left and right operand types may differ. CommonDataT selects a result type @@ -1194,37 +1180,25 @@ class SwizzleOp { return NewRelOp(NULL, *this, Rhs); } - // Begin hi/lo, even/odd, xyzw, and rgba swizzles. private: - // Indexer used in the swizzles.def. - // Currently it is defined as a template struct. Replacing it with a constexpr - // function would activate a bug in MSVC that is fixed only in v19.20. - // Until then MSVC does not recognize such constexpr functions as const and - // thus does not let using them in template parameters inside swizzle.def. - template struct Indexer { - static constexpr int IDXs[sizeof...(Indexes)] = {Indexes...}; - static constexpr int value = IDXs[Index >= getNumElements() ? 0 : Index]; - }; + static constexpr auto get_vec_idx(int idx) { + int counter = 0; + int result = -1; + ((result = counter++ == idx ? Indexes : result), ...); + return result; + } public: template - ConstSwizzle::value...> swizzle() const { + ConstSwizzle swizzle() const { return m_Vector; } template - Swizzle::value...> swizzle() { + Swizzle swizzle() { return m_Vector; } -#ifdef __SYCL_ACCESS_RETURN -#error "Undefine __SYCL_ACCESS_RETURN macro" -#endif -#define __SYCL_ACCESS_RETURN m_Vector -#include "swizzles.def" -#undef __SYCL_ACCESS_RETURN - // End of hi/lo, even/odd, xyzw, and rgba swizzles. - // Leave store() interface to automatic conversion to vec<>. // Load to vec_t and then assign to swizzle. template diff --git a/sycl/test/check_device_code/vector/vector_bf16_builtins.cpp b/sycl/test/check_device_code/vector/vector_bf16_builtins.cpp index 891c79a1a8f94..a3c385eae35f0 100644 --- a/sycl/test/check_device_code/vector/vector_bf16_builtins.cpp +++ b/sycl/test/check_device_code/vector/vector_bf16_builtins.cpp @@ -69,7 +69,7 @@ SYCL_EXTERNAL auto TestFMin(vec a, vec b) { } // CHECK-LABEL: define dso_local spir_func void @_Z8TestFMaxN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi3EEES5_( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.5") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.5") align 8 [[A:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.5") align 8 [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META24:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.6") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.6") align 8 [[A:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.6") align 8 [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META24:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: // CHECK-NEXT: [[VEC_ADDR_I_I_I_I12_I:%.*]] = alloca <3 x float>, align 16 // CHECK-NEXT: [[DST_I_I_I_I13_I:%.*]] = alloca [4 x %"class.sycl::_V1::ext::oneapi::bfloat16"], align 2 @@ -123,7 +123,7 @@ SYCL_EXTERNAL auto TestFMax(vec a, vec b) { } // CHECK-LABEL: define dso_local spir_func void @_Z9TestIsNanN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi4EEE( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.15") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.20") align 8 [[A:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META38:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.18") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.24") align 8 [[A:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META38:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: // CHECK-NEXT: [[VEC_ADDR_I_I_I_I_I:%.*]] = alloca <4 x i16>, align 8 // CHECK-NEXT: [[DST_I_I_I_I_I:%.*]] = alloca [4 x float], align 4 @@ -149,7 +149,7 @@ SYCL_EXTERNAL auto TestIsNan(vec a) { } // CHECK-LABEL: define dso_local spir_func void @_Z8TestFabsN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi8EEE( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.38") align 16 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.38") align 16 [[A:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META48:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.46") align 16 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.46") align 16 [[A:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META48:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: // CHECK-NEXT: [[VEC_ADDR_I_I_I_I2_I:%.*]] = alloca <8 x float>, align 32 // CHECK-NEXT: [[DST_I_I_I_I3_I:%.*]] = alloca [8 x %"class.sycl::_V1::ext::oneapi::bfloat16"], align 2 @@ -185,7 +185,7 @@ SYCL_EXTERNAL auto TestFabs(vec a) { } // CHECK-LABEL: define dso_local spir_func void @_Z8TestCeilN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi8EEE( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.38") align 16 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.38") align 16 [[A:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META59:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.46") align 16 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.46") align 16 [[A:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META59:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: // CHECK-NEXT: [[VEC_ADDR_I_I_I_I2_I:%.*]] = alloca <8 x float>, align 32 // CHECK-NEXT: [[DST_I_I_I_I3_I:%.*]] = alloca [8 x %"class.sycl::_V1::ext::oneapi::bfloat16"], align 2 @@ -221,7 +221,7 @@ SYCL_EXTERNAL auto TestCeil(vec a) { } // CHECK-LABEL: define dso_local spir_func void @_Z7TestFMAN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi16EEES5_S5_( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.48") align 32 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.48") align 32 [[A:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.48") align 32 [[B:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.48") align 32 [[C:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META70:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.58") align 32 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.58") align 32 [[A:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.58") align 32 [[B:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.58") align 32 [[C:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META70:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: // CHECK-NEXT: [[VEC_ADDR_I_I_I_I14_I:%.*]] = alloca <16 x float>, align 64 // CHECK-NEXT: [[DST_I_I_I_I15_I:%.*]] = alloca [16 x %"class.sycl::_V1::ext::oneapi::bfloat16"], align 2 diff --git a/sycl/test/check_device_code/vector/vector_convert_bfloat.cpp b/sycl/test/check_device_code/vector/vector_convert_bfloat.cpp index 3166aacc915e3..e531e2e79df0d 100644 --- a/sycl/test/check_device_code/vector/vector_convert_bfloat.cpp +++ b/sycl/test/check_device_code/vector/vector_convert_bfloat.cpp @@ -63,7 +63,7 @@ SYCL_EXTERNAL auto TestBFtoFDeviceRZ(vec &inp) { } // CHECK-LABEL: define dso_local spir_func void @_Z19TestBFtointDeviceRZRN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi3EEE( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.5") align 16 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readonly align 8 dereferenceable(8) [[INP:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META18:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.6") align 16 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readonly align 8 dereferenceable(8) [[INP:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META18:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META19:![0-9]+]]) // CHECK-NEXT: [[LOADVEC4_I_I:%.*]] = load <4 x i16>, ptr addrspace(4) [[INP]], align 8, !noalias [[META19]] @@ -93,7 +93,7 @@ SYCL_EXTERNAL auto TestBFtointDeviceRZ(vec &inp) { } // CHECK-LABEL: define dso_local spir_func void @_Z20TestBFtointDeviceRNERN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi1EEE( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.10") align 4 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readonly align 2 dereferenceable(2) [[INP:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META24:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.12") align 4 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readonly align 2 dereferenceable(2) [[INP:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META24:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META25:![0-9]+]]) // CHECK-NEXT: [[TMP0:%.*]] = load i16, ptr addrspace(4) [[INP]], align 2, !tbaa [[TBAA11]], !noalias [[META25]] @@ -191,7 +191,7 @@ SYCL_EXTERNAL auto TestInttoBFDeviceRZ(vec &inp) { } // CHECK-LABEL: define dso_local spir_func void @_Z19TestLLtoBFDeviceRTPRN4sycl3_V13vecIxLi1EEE( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.15") align 2 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readonly align 8 dereferenceable(8) [[INP:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META42:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.18") align 2 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readonly align 8 dereferenceable(8) [[INP:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META42:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META43:![0-9]+]]) // CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr addrspace(4) [[INP]], align 8, !tbaa [[TBAA46:![0-9]+]], !noalias [[META43]] @@ -204,7 +204,7 @@ SYCL_EXTERNAL auto TestLLtoBFDeviceRTP(vec &inp) { } // CHECK-LABEL: define dso_local spir_func void @_Z22TestShorttoBFDeviceRTNRN4sycl3_V13vecIsLi2EEE( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.25") align 4 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readonly align 4 dereferenceable(4) [[INP:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META48:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.30") align 4 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readonly align 4 dereferenceable(4) [[INP:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META48:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META49:![0-9]+]]) // CHECK-NEXT: [[TMP0:%.*]] = load <2 x i16>, ptr addrspace(4) [[INP]], align 4, !tbaa [[TBAA11]], !noalias [[META49]] diff --git a/sycl/test/check_device_code/vector/vector_math_ops.cpp b/sycl/test/check_device_code/vector/vector_math_ops.cpp index 527a6bee05831..946c30e17d647 100644 --- a/sycl/test/check_device_code/vector/vector_math_ops.cpp +++ b/sycl/test/check_device_code/vector/vector_math_ops.cpp @@ -46,7 +46,7 @@ SYCL_EXTERNAL auto TestAdd(vec a, vec b) { return a + b; } SYCL_EXTERNAL auto TestAdd(vec a, vec b) { return a + b; } // CHECK-LABEL: define dso_local spir_func void @_Z7TestAddN4sycl3_V13vecIcLi16EEES2_( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.8") align 16 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.8") align 16 [[A:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.8") align 16 [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META26:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.9") align 16 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.9") align 16 [[A:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.9") align 16 [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META26:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META27:![0-9]+]]) // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META30:![0-9]+]]) @@ -60,7 +60,7 @@ SYCL_EXTERNAL auto TestAdd(vec a, vec b) { return a + b; } // std::byte does not support '+'. Therefore, using bitwise XOR as a substitute. // CHECK-LABEL: define dso_local spir_func void @_Z7TestXorN4sycl3_V13vecISt4byteLi8EEES3_( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.15") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.15") align 8 [[A:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.15") align 8 [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META34:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.17") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.17") align 8 [[A:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.17") align 8 [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META34:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META35:![0-9]+]]) // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META38:![0-9]+]]) @@ -75,7 +75,7 @@ SYCL_EXTERNAL auto TestXor(vec a, vec b) { } // CHECK-LABEL: define dso_local spir_func void @_Z7TestAddN4sycl3_V13vecIbLi4EEES2_( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.22") align 4 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.22") align 4 [[A:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.22") align 4 [[B:%.*]]) local_unnamed_addr #[[ATTR1:[0-9]+]] !srcloc [[META42:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.25") align 4 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.25") align 4 [[A:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.25") align 4 [[B:%.*]]) local_unnamed_addr #[[ATTR1:[0-9]+]] !srcloc [[META42:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META43:![0-9]+]]) // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META46:![0-9]+]]) @@ -102,7 +102,7 @@ SYCL_EXTERNAL auto TestXor(vec a, vec b) { SYCL_EXTERNAL auto TestAdd(vec a, vec b) { return a + b; } // CHECK-LABEL: define dso_local spir_func void @_Z7TestAddN4sycl3_V13vecINS0_6detail9half_impl4halfELi3EEES5_( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.29") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.29") align 8 [[A:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.29") align 8 [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META52:![0-9]+]] !sycl_used_aspects [[META53:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.33") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.33") align 8 [[A:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.33") align 8 [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META52:![0-9]+]] !sycl_used_aspects [[META53:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META55:![0-9]+]]) // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META58:![0-9]+]]) @@ -116,10 +116,10 @@ SYCL_EXTERNAL auto TestAdd(vec a, vec b) { return a + b; } SYCL_EXTERNAL auto TestAdd(vec a, vec b) { return a + b; } // CHECK-LABEL: define dso_local spir_func void @_Z7TestAddN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi3EEES5_( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.36") align 8 [[AGG_RESULT:%.*]], ptr noundef byval(%"class.sycl::_V1::vec.36") align 8 [[A:%.*]], ptr noundef byval(%"class.sycl::_V1::vec.36") align 8 [[B:%.*]]) local_unnamed_addr #[[ATTR2:[0-9]+]] !srcloc [[META62:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.41") align 8 [[AGG_RESULT:%.*]], ptr noundef byval(%"class.sycl::_V1::vec.41") align 8 [[A:%.*]], ptr noundef byval(%"class.sycl::_V1::vec.41") align 8 [[B:%.*]]) local_unnamed_addr #[[ATTR2:[0-9]+]] !srcloc [[META62:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: // CHECK-NEXT: [[REF_TMP_I_I_I_I:%.*]] = alloca float, align 4 -// CHECK-NEXT: [[RES_I_I:%.*]] = alloca %"class.sycl::_V1::vec.36", align 8 +// CHECK-NEXT: [[RES_I_I:%.*]] = alloca %"class.sycl::_V1::vec.41", align 8 // CHECK-NEXT: [[A_ASCAST:%.*]] = addrspacecast ptr [[A]] to ptr addrspace(4) // CHECK-NEXT: [[B_ASCAST:%.*]] = addrspacecast ptr [[B]] to ptr addrspace(4) // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META63:![0-9]+]]) @@ -160,7 +160,7 @@ SYCL_EXTERNAL auto TestAdd(vec a, /***************** Binary Logical Ops *******************/ // CHECK-LABEL: define dso_local spir_func void @_Z15TestGreaterThanN4sycl3_V13vecIiLi16EEES2_( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.41") align 64 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.41") align 64 [[A:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.41") align 64 [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META81:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.47") align 64 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.47") align 64 [[A:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.47") align 64 [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META81:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META82:![0-9]+]]) // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META85:![0-9]+]]) @@ -176,7 +176,7 @@ SYCL_EXTERNAL auto TestGreaterThan(vec a, vec b) { } // CHECK-LABEL: define dso_local spir_func noundef range(i8 -1, 1) <3 x i8> @_Z15TestGreaterThanN4sycl3_V13vecISt4byteLi3EEES3_( -// CHECK-SAME: ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.49") align 4 [[A:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.49") align 4 [[B:%.*]]) local_unnamed_addr #[[ATTR4:[0-9]+]] !srcloc [[META89:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.56") align 4 [[A:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.56") align 4 [[B:%.*]]) local_unnamed_addr #[[ATTR4:[0-9]+]] !srcloc [[META89:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: // CHECK-NEXT: [[LOADVEC4_I_I:%.*]] = load <4 x i8>, ptr [[A]], align 1 // CHECK-NEXT: [[LOADVEC4_I_I2:%.*]] = load <4 x i8>, ptr [[B]], align 1 @@ -190,7 +190,7 @@ SYCL_EXTERNAL auto TestGreaterThan(vec a, vec b) { } // CHECK-LABEL: define dso_local spir_func void @_Z15TestGreaterThanN4sycl3_V13vecIbLi2EEES2_( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.54") align 2 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.59") align 2 [[A:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.59") align 2 [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META90:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.62") align 2 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.68") align 2 [[A:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.68") align 2 [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META90:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META91:![0-9]+]]) // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META94:![0-9]+]]) @@ -206,7 +206,7 @@ SYCL_EXTERNAL auto TestGreaterThan(vec a, vec b) { } // CHECK-LABEL: define dso_local spir_func void @_Z15TestGreaterThanN4sycl3_V13vecINS0_6detail9half_impl4halfELi8EEES5_( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.68") align 16 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.73") align 16 [[A:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.73") align 16 [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META98:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.78") align 16 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.84") align 16 [[A:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.84") align 16 [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META98:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META99:![0-9]+]]) // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META102:![0-9]+]]) @@ -222,9 +222,9 @@ SYCL_EXTERNAL auto TestGreaterThan(vec a, vec b) { } // CHECK-LABEL: define dso_local spir_func void @_Z15TestGreaterThanN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi4EEES5_( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.82") align 8 [[AGG_RESULT:%.*]], ptr noundef byval(%"class.sycl::_V1::vec.87") align 8 [[A:%.*]], ptr noundef byval(%"class.sycl::_V1::vec.87") align 8 [[B:%.*]]) local_unnamed_addr #[[ATTR2]] !srcloc [[META106:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.94") align 8 [[AGG_RESULT:%.*]], ptr noundef byval(%"class.sycl::_V1::vec.100") align 8 [[A:%.*]], ptr noundef byval(%"class.sycl::_V1::vec.100") align 8 [[B:%.*]]) local_unnamed_addr #[[ATTR2]] !srcloc [[META106:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: -// CHECK-NEXT: [[RES_I_I:%.*]] = alloca %"class.sycl::_V1::vec.82", align 8 +// CHECK-NEXT: [[RES_I_I:%.*]] = alloca %"class.sycl::_V1::vec.94", align 8 // CHECK-NEXT: [[A_ASCAST:%.*]] = addrspacecast ptr [[A]] to ptr addrspace(4) // CHECK-NEXT: [[B_ASCAST:%.*]] = addrspacecast ptr [[B]] to ptr addrspace(4) // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META107:![0-9]+]]) @@ -261,7 +261,7 @@ SYCL_EXTERNAL auto TestGreaterThan(vec a, /********************** Unary Ops **********************/ // CHECK-LABEL: define dso_local spir_func void @_Z12TestNegationN4sycl3_V13vecIiLi3EEE( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.91") align 16 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.91") align 16 [[A:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META115:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.105") align 16 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.105") align 16 [[A:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META115:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META116:![0-9]+]]) // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META119:![0-9]+]]) @@ -276,7 +276,7 @@ SYCL_EXTERNAL auto TestGreaterThan(vec a, SYCL_EXTERNAL auto TestNegation(vec a) { return !a; } // CHECK-LABEL: define dso_local spir_func void @_Z9TestMinusN4sycl3_V13vecIiLi4EEE( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.97") align 16 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.97") align 16 [[A:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META123:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.112") align 16 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.112") align 16 [[A:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META123:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META124:![0-9]+]]) // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META127:![0-9]+]]) @@ -289,7 +289,7 @@ SYCL_EXTERNAL auto TestMinus(vec a) { return -a; } // Negation is not valid for std::byte. Therefore, using bitwise negation. // CHECK-LABEL: define dso_local spir_func void @_Z19TestBitwiseNegationN4sycl3_V13vecISt4byteLi16EEE( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.102") align 16 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.102") align 16 [[A:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META131:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.118") align 16 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.118") align 16 [[A:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META131:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META132:![0-9]+]]) // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META135:![0-9]+]]) @@ -301,7 +301,7 @@ SYCL_EXTERNAL auto TestMinus(vec a) { return -a; } SYCL_EXTERNAL auto TestBitwiseNegation(vec a) { return ~a; } // CHECK-LABEL: define dso_local spir_func void @_Z12TestNegationN4sycl3_V13vecIbLi4EEE( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.108") align 4 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.22") align 4 [[A:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META139:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.125") align 4 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.25") align 4 [[A:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META139:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META140:![0-9]+]]) // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META143:![0-9]+]]) @@ -314,7 +314,7 @@ SYCL_EXTERNAL auto TestBitwiseNegation(vec a) { return ~a; } SYCL_EXTERNAL auto TestNegation(vec a) { return !a; } // CHECK-LABEL: define dso_local spir_func void @_Z12TestNegationN4sycl3_V13vecINS0_6detail9half_impl4halfELi2EEE( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.114") align 4 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.119") align 4 [[A:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META147:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.132") align 4 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.138") align 4 [[A:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META147:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META148:![0-9]+]]) // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META151:![0-9]+]]) @@ -327,7 +327,7 @@ SYCL_EXTERNAL auto TestNegation(vec a) { return !a; } SYCL_EXTERNAL auto TestNegation(vec a) { return !a; } // CHECK-LABEL: define dso_local spir_func void @_Z9TestMinusN4sycl3_V13vecINS0_6detail9half_impl4halfELi8EEE( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.73") align 16 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.73") align 16 [[A:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META155:![0-9]+]] !sycl_used_aspects [[META53]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.84") align 16 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.84") align 16 [[A:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META155:![0-9]+]] !sycl_used_aspects [[META53]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META156:![0-9]+]]) // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META159:![0-9]+]]) @@ -339,9 +339,9 @@ SYCL_EXTERNAL auto TestNegation(vec a) { return !a; } SYCL_EXTERNAL auto TestMinus(vec a) { return -a; } // CHECK-LABEL: define dso_local spir_func void @_Z12TestNegationN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi3EEE( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.126") align 8 [[AGG_RESULT:%.*]], ptr noundef byval(%"class.sycl::_V1::vec.36") align 8 [[A:%.*]]) local_unnamed_addr #[[ATTR2]] !srcloc [[META163:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.146") align 8 [[AGG_RESULT:%.*]], ptr noundef byval(%"class.sycl::_V1::vec.41") align 8 [[A:%.*]]) local_unnamed_addr #[[ATTR2]] !srcloc [[META163:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: -// CHECK-NEXT: [[RES_I_I:%.*]] = alloca %"class.sycl::_V1::vec.126", align 8 +// CHECK-NEXT: [[RES_I_I:%.*]] = alloca %"class.sycl::_V1::vec.146", align 8 // CHECK-NEXT: [[A_ASCAST:%.*]] = addrspacecast ptr [[A]] to ptr addrspace(4) // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META164:![0-9]+]]) // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META167:![0-9]+]]) @@ -370,10 +370,10 @@ SYCL_EXTERNAL auto TestMinus(vec a) { return -a; } SYCL_EXTERNAL auto TestNegation(vec a) { return !a; } // CHECK-LABEL: define dso_local spir_func void @_Z9TestMinusN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi16EEE( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.130") align 32 [[AGG_RESULT:%.*]], ptr noundef byval(%"class.sycl::_V1::vec.130") align 32 [[A:%.*]]) local_unnamed_addr #[[ATTR2]] !srcloc [[META172:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.151") align 32 [[AGG_RESULT:%.*]], ptr noundef byval(%"class.sycl::_V1::vec.151") align 32 [[A:%.*]]) local_unnamed_addr #[[ATTR2]] !srcloc [[META172:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: entry: // CHECK-NEXT: [[REF_TMP_I_I_I_I:%.*]] = alloca float, align 4 -// CHECK-NEXT: [[RES_I_I:%.*]] = alloca %"class.sycl::_V1::vec.130", align 32 +// CHECK-NEXT: [[RES_I_I:%.*]] = alloca %"class.sycl::_V1::vec.151", align 32 // CHECK-NEXT: [[A_ASCAST:%.*]] = addrspacecast ptr [[A]] to ptr addrspace(4) // CHECK-NEXT: call void @llvm.lifetime.start.p0(i64 32, ptr nonnull [[RES_I_I]]), !noalias [[META173:![0-9]+]] // CHECK-NEXT: call void @llvm.memset.p0.i64(ptr align 32 [[RES_I_I]], i8 0, i64 32, i1 false), !noalias [[META176:![0-9]+]] From b2a2864d24e0fbf7d1228b1df9491a81e2e6c669 Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Wed, 18 Dec 2024 09:41:12 -0800 Subject: [PATCH 2/2] undef __SYCL_SWIZZLE_MIXIN_ALL_SWIZZLES --- sycl/include/sycl/detail/named_swizzles_mixin.hpp | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/include/sycl/detail/named_swizzles_mixin.hpp b/sycl/include/sycl/detail/named_swizzles_mixin.hpp index c7bee0d1b7bb1..f35656b2e6381 100644 --- a/sycl/include/sycl/detail/named_swizzles_mixin.hpp +++ b/sycl/include/sycl/detail/named_swizzles_mixin.hpp @@ -816,6 +816,7 @@ template struct NamedSwizzlesMixinBoth { #undef __SYCL_SWIZZLE_MIXIN_METHOD_CONST #undef __SYCL_SWIZZLE_MIXIN_METHOD_NON_CONST +#undef __SYCL_SWIZZLE_MIXIN_ALL_SWIZZLES #undef __SYCL_SWIZZLE_MIXIN_SIMPLE_SWIZZLES } // namespace detail