[SYCL] Fix SFINAE rules for integer builtins/bitselect (#12671)

In case of vectors/swizzles of integer types only fixed width types are
allowed per SYCL 2020 revision 8. Update the implementation to match
that.

sycl/include/sycl/builtins_preview.hpp

@@ -137,7 +137,14 @@ auto builtin_marray_impl(FuncTy F, const Ts &...x) {
  marray<ret_elem_type, T::size()> Res;
  constexpr auto N = T::size();
  for (size_t I = 0; I < N / 2; ++I) {
- auto PartialRes = F(to_vec2(x, I * 2)...);
+ auto PartialRes = [&]() {
+ using elem_ty = get_elem_type_t<T>;
+ if constexpr (std::is_integral_v<elem_ty>)
+ return F(to_vec2(x, I * 2)
+ .template as<vec<get_fixed_sized_int_t<elem_ty>, 2>>()...);
+ else
+ return F(to_vec2(x, I * 2)...);
+ }();
  std::memcpy(&Res[I * 2], &PartialRes, sizeof(decltype(PartialRes)));
  }
  if (N % 2)

sycl/include/sycl/builtins_utils_scalar.hpp

@@ -128,6 +128,17 @@ template <size_t Size> struct get_unsigned_int_by_size {
 template <typename T> struct same_size_unsigned_int {
  using type = typename get_unsigned_int_by_size<sizeof(T)>::type;
 };
+template <typename T>
+using same_size_unsigned_int_t = typename same_size_unsigned_int<T>::type;
+
+template <typename T> struct get_fixed_sized_int {
+ static_assert(std::is_integral_v<T>);
+ using type =
+ std::conditional_t<std::is_signed_v<T>, same_size_signed_int_t<T>,
+ same_size_unsigned_int_t<T>>;
+};
+template <typename T>
+using get_fixed_sized_int_t = typename get_fixed_sized_int<T>::type;
 
 // Utility trait for getting an upsampled integer type.
 // NOTE: For upsampling we look for an integer of double the size of the

sycl/include/sycl/detail/builtins/helper_macros.hpp

@@ -48,6 +48,11 @@
  FOR_EACH4_A6(BASE_CASE, FIXED1, FIXED2, FIXED3, FIXED4, ARG1, ARG2, ARG3, \
  ARG4, ARG5, ARG6) \
  BASE_CASE(FIXED1, FIXED2, FIXED3, FIXED4, ARG7)
+#define FOR_EACH4_A8(BASE_CASE, FIXED1, FIXED2, FIXED3, FIXED4, ARG1, ARG2, \
+ ARG3, ARG4, ARG5, ARG6, ARG7, ARG8) \
+ FOR_EACH4_A7(BASE_CASE, FIXED1, FIXED2, FIXED3, FIXED4, ARG1, ARG2, ARG3, \
+ ARG4, ARG5, ARG6, ARG7) \
+ BASE_CASE(FIXED1, FIXED2, FIXED3, FIXED4, ARG8)
 #define FOR_EACH4_A11(BASE_CASE, FIXED1, FIXED2, FIXED3, FIXED4, ARG1, ARG2, \
  ARG3, ARG4, ARG5, ARG6, ARG7, ARG8, ARG9, ARG10, ARG11) \
  FOR_EACH4_A7(BASE_CASE, FIXED1, FIXED2, FIXED3, FIXED4, ARG1, ARG2, ARG3, \
@@ -169,6 +174,9 @@
  unsigned char, unsigned short, unsigned int, unsigned long, unsigned long long
 // 11 types
 #define INTEGER_TYPES SIGNED_TYPES, UNSIGNED_TYPES
+// 8 types
+#define FIXED_WIDTH_INTEGER_TYPES \
+ int8_t, int16_t, int32_t, int64_t, uint8_t, uint16_t, uint32_t, uint64_t
 
 #define DEVICE_IMPL_TEMPLATE_CUSTOM_DELEGATE( \
  NUM_ARGS, NAME, ENABLER, DELEGATOR, NS, /*SCALAR_VEC_IMPL*/...) \

sycl/include/sycl/detail/builtins/integer_functions.inc

@@ -16,9 +16,10 @@ namespace detail {
 template <typename T>
 struct integer_elem_type
  : std::bool_constant<
- check_type_in_v<get_elem_type_t<T>, char, signed char, short, int,
- long, long long, unsigned char, unsigned short,
- unsigned int, unsigned long, unsigned long long>> {};
+ (is_vec_or_swizzle_v<T> &&
+ check_type_in_v<get_elem_type_t<T>, FIXED_WIDTH_INTEGER_TYPES>) ||
+ (!is_vec_or_swizzle_v<T> &&
+ check_type_in_v<get_elem_type_t<T>, INTEGER_TYPES>)> {};
 template <typename T>
 struct suint32_elem_type
  : std::bool_constant<

sycl/include/sycl/detail/builtins/relational_functions.inc

@@ -15,10 +15,12 @@ inline namespace _V1 {
 namespace detail {
 template <typename T>
 struct bitselect_elem_type
- : std::bool_constant<check_type_in_v<
- get_elem_type_t<T>, float, double, half, char, signed char, short,
- int, long, long long, unsigned char, unsigned short, unsigned int,
- unsigned long, unsigned long long>> {};
+ : std::bool_constant<
+ check_type_in_v<get_elem_type_t<T>, FP_TYPES> ||
+ (is_vec_or_swizzle_v<T> &&
+ check_type_in_v<get_elem_type_t<T>, FIXED_WIDTH_INTEGER_TYPES>) ||
+ (!is_vec_or_swizzle_v<T> &&
+ check_type_in_v<get_elem_type_t<T>, INTEGER_TYPES>)> {};
 
 template <typename T>
 struct rel_ret_traits

sycl/source/builtins/host_helper_macros.hpp

@@ -56,6 +56,9 @@
 #define EXPORT_VEC(NUM_ARGS, NAME, TYPE, VL) \
  EXPORT_VEC_NS(NUM_ARGS, NAME, sycl, TYPE, VL)
 
+#define EXPORT_VEC_1_16_IMPL(NUM_ARGS, NAME, NS, TYPE) \
+ FOR_VEC_1_16(EXPORT_VEC_NS, NUM_ARGS, NAME, NS, TYPE)
+
 #define EXPORT_SCALAR_AND_VEC_1_16_IMPL(NUM_ARGS, NAME, NS, TYPE) \
  EXPORT_SCALAR_NS(NUM_ARGS, NAME, NS, TYPE) \
  FOR_VEC_1_16(EXPORT_VEC_NS, NUM_ARGS, NAME, NS, TYPE)
@@ -69,8 +72,12 @@
 
 #define EXPORT_SCALAR_AND_VEC_1_16_NS(NUM_ARGS, NAME, NS, ...) \
  FOR_EACH3(EXPORT_SCALAR_AND_VEC_1_16_IMPL, NUM_ARGS, NAME, NS, __VA_ARGS__)
+#define EXPORT_VEC_1_16_NS(NUM_ARGS, NAME, NS, ...) \
+ FOR_EACH3(EXPORT_VEC_1_16_IMPL, NUM_ARGS, NAME, NS, __VA_ARGS__)
 #define EXPORT_SCALAR_AND_VEC_1_16(NUM_ARGS, NAME, ...) \
  EXPORT_SCALAR_AND_VEC_1_16_NS(NUM_ARGS, NAME, sycl, __VA_ARGS__)
+#define EXPORT_VEC_1_16(NUM_ARGS, NAME, ...) \
+ EXPORT_VEC_1_16_NS(NUM_ARGS, NAME, sycl, __VA_ARGS__)
 
 #define EXPORT_SCALAR_AND_VEC_2_4(NUM_ARGS, NAME, ...) \
  FOR_EACH2(EXPORT_SCALAR_AND_VEC_2_4_IMPL, NUM_ARGS, NAME, __VA_ARGS__)

sycl/source/builtins/integer_functions.cpp

@@ -76,7 +76,8 @@ namespace sycl {
 inline namespace _V1 {
 #define BUILTIN_GENINT(NUM_ARGS, NAME, IMPL) \
  HOST_IMPL(NAME, IMPL) \
- EXPORT_SCALAR_AND_VEC_1_16(NUM_ARGS, NAME, INTEGER_TYPES)
+ FOR_EACH2(EXPORT_SCALAR, NUM_ARGS, NAME, INTEGER_TYPES) \
+ EXPORT_VEC_1_16(NUM_ARGS, NAME, FIXED_WIDTH_INTEGER_TYPES)
 #define BUILTIN_GENINT_SU(NUM_ARGS, NAME, IMPL) \
  BUILTIN_GENINT(NUM_ARGS, NAME, IMPL)
 

sycl/source/builtins/relational_functions.cpp

@@ -103,6 +103,7 @@ HOST_IMPL(bitselect, [](auto x, auto y, auto z) {
  assert((ures & std::numeric_limits<utype>::max()) == ures);
  return bit_cast<T0>(static_cast<utype>(ures));
 })
-EXPORT_SCALAR_AND_VEC_1_16(THREE_ARGS, bitselect, INTEGER_TYPES, FP_TYPES)
+FOR_EACH2(EXPORT_SCALAR, THREE_ARGS, bitselect, INTEGER_TYPES, FP_TYPES)
+EXPORT_VEC_1_16(THREE_ARGS, bitselect, FIXED_WIDTH_INTEGER_TYPES, FP_TYPES)
 } // namespace _V1
 } // namespace sycl

sycl/test/basic_tests/builtins/builtin_unit_tests.cpp

@@ -0,0 +1,133 @@
+// RUN: %clangxx -fsycl -fpreview-breaking-changes -fsyntax-only %s -Xclang -verify
+// REQUIRES: preview-breaking-changes-supported
+
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+using namespace sycl::detail;
+
+namespace builtin_same_shape_v_tests {
+using swizzle1 = decltype(std::declval<vec<float, 2>>().swizzle<0>());
+using swizzle2 = decltype(std::declval<vec<float, 2>>().swizzle<0, 0>());
+using swizzle3 = decltype(std::declval<vec<float, 2>>().swizzle<0, 0, 1>());
+
+static_assert(builtin_same_shape_v<float>);
+static_assert(builtin_same_shape_v<int, float>);
+static_assert(builtin_same_shape_v<marray<int, 2>>);
+static_assert(builtin_same_shape_v<marray<int, 2>, marray<float, 2>>);
+static_assert(builtin_same_shape_v<vec<int, 2>>);
+static_assert(builtin_same_shape_v<vec<int, 2>, vec<float, 2>>);
+static_assert(builtin_same_shape_v<vec<int, 2>, swizzle2>);
+
+static_assert(!builtin_same_shape_v<float, marray<float, 1>>);
+static_assert(!builtin_same_shape_v<float, vec<float, 1>>);
+static_assert(!builtin_same_shape_v<marray<float, 1>, vec<float, 1>>);
+static_assert(!builtin_same_shape_v<float, swizzle1>);
+static_assert(!builtin_same_shape_v<marray<float, 1>, swizzle1>);
+static_assert(!builtin_same_shape_v<swizzle2, swizzle1>);
+} // namespace builtin_same_shape_v_tests
+
+namespace builtin_marray_impl_tests {
+// Integer functions/relational bitselect only accept fixed-width integer
+// element types for vector/swizzle elements. Make sure that our marray->vec
+// delegator can handle that.
+
+auto foo(char x) { return x; }
+auto foo(signed char x) { return x; }
+auto foo(unsigned char x) { return x; }
+auto foo(vec<int8_t, 2> x) { return x; }
+auto foo(vec<uint8_t, 2> x) { return x; }
+
+auto test() {
+ marray<char, 2> x;
+ marray<signed char, 2> y;
+ marray<unsigned char, 2> z;
+ auto TestOne = [](auto x) {
+ std::ignore = builtin_marray_impl([](auto x) { return foo(x); }, x);
+ };
+ TestOne(x);
+ TestOne(y);
+ TestOne(z);
+}
+} // namespace builtin_marray_impl_tests
+
+namespace builtin_enable_integer_tests {
+using swizzle1 = decltype(std::declval<vec<int8_t, 2>>().swizzle<0>());
+using swizzle2 = decltype(std::declval<vec<int8_t, 2>>().swizzle<0, 0>());
+template <typename... Ts> void ignore() {}
+
+void test() {
+ // clang-format off
+ ignore<builtin_enable_integer_t<char>,
+ builtin_enable_integer_t<signed char>,
+ builtin_enable_integer_t<unsigned char>>();
+ // clang-format on
+
+ ignore<builtin_enable_integer_t<vec<int8_t, 2>>,
+ builtin_enable_integer_t<vec<uint8_t, 2>>>();
+
+ ignore<builtin_enable_integer_t<char, char>>();
+ ignore<builtin_enable_integer_t<vec<int8_t, 2>, vec<int8_t, 2>>>();
+ ignore<builtin_enable_integer_t<vec<int8_t, 2>, swizzle2>>();
+ ignore<builtin_enable_integer_t<swizzle2, swizzle2>>();
+
+ {
+ // Only one of char/signed char maps onto int8_t. The other type isn't a
+ // valid vector element type for integer builtins.
+
+ static_assert(std::is_signed_v<char>);
+
+ // clang-format off
+ // expected-error-re@*:* {{no type named 'type' in 'sycl::detail::builtin_enable<sycl::detail::default_ret_type, sycl::detail::integer_elem_type, sycl::detail::any_shape, sycl::detail::same_elem_type, sycl::vec<{{.*}}, 2>>'}}
+ // expected-note@+1 {{in instantiation of template type alias 'builtin_enable_integer_t' requested here}}
+ ignore<builtin_enable_integer_t<vec<signed char, 2>>, builtin_enable_integer_t<vec<char, 2>>>();
+ // clang-format on
+ }
+
+ // expected-error@*:* {{no type named 'type' in 'sycl::detail::builtin_enable<sycl::detail::default_ret_type, sycl::detail::integer_elem_type, sycl::detail::any_shape, sycl::detail::same_elem_type, char, signed char>'}}
+ // expected-note@+1 {{in instantiation of template type alias 'builtin_enable_integer_t' requested here}}
+ ignore<builtin_enable_integer_t<char, signed char>>();
+}
+} // namespace builtin_enable_integer_tests
+
+namespace builtin_enable_bitselect_tests {
+// Essentially the same as builtin_enable_integer_t + FP types support.
+using swizzle1 = decltype(std::declval<vec<int8_t, 2>>().swizzle<0>());
+using swizzle2 = decltype(std::declval<vec<int8_t, 2>>().swizzle<0, 0>());
+template <typename... Ts> void ignore() {}
+
+void test() {
+ // clang-format off
+ ignore<builtin_enable_bitselect_t<char>,
+ builtin_enable_bitselect_t<signed char>,
+ builtin_enable_bitselect_t<unsigned char>,
+ builtin_enable_bitselect_t<float>>();
+ // clang-format on
+
+ ignore<builtin_enable_bitselect_t<vec<int8_t, 2>>,
+ builtin_enable_bitselect_t<vec<uint8_t, 2>>,
+ builtin_enable_bitselect_t<vec<float, 2>>>();
+
+ ignore<builtin_enable_bitselect_t<char, char>>();
+ ignore<builtin_enable_bitselect_t<vec<int8_t, 2>, vec<int8_t, 2>>>();
+ ignore<builtin_enable_bitselect_t<vec<int8_t, 2>, swizzle2>>();
+ ignore<builtin_enable_bitselect_t<swizzle2, swizzle2>>();
+
+ {
+ // Only one of char/signed char maps onto int8_t. The other type isn't a
+ // valid vector element type for integer builtins.
+
+ static_assert(std::is_signed_v<char>);
+
+ // clang-format off
+ // expected-error-re@*:* {{no type named 'type' in 'sycl::detail::builtin_enable<sycl::detail::default_ret_type, sycl::detail::bitselect_elem_type, sycl::detail::any_shape, sycl::detail::same_elem_type, sycl::vec<{{.*}}, 2>>'}}
+ // expected-note@+1 {{in instantiation of template type alias 'builtin_enable_bitselect_t' requested here}}
+ ignore<builtin_enable_bitselect_t<vec<signed char, 2>>, builtin_enable_bitselect_t<vec<char, 2>>>();
+ // clang-format on
+ }
+
+ // expected-error@*:* {{no type named 'type' in 'sycl::detail::builtin_enable<sycl::detail::default_ret_type, sycl::detail::bitselect_elem_type, sycl::detail::any_shape, sycl::detail::same_elem_type, char, signed char>'}}
+ // expected-note@+1 {{in instantiation of template type alias 'builtin_enable_bitselect_t' requested here}}
+ ignore<builtin_enable_bitselect_t<char, signed char>>();
+}
+} // namespace builtin_enable_bitselect_tests