Skip to content

Commit 1f37b5e

Browse files
[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.
1 parent 6098a75 commit 1f37b5e

File tree

9 files changed

+181
-10
lines changed

9 files changed

+181
-10
lines changed

sycl/include/sycl/builtins_preview.hpp

Lines changed: 8 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -137,7 +137,14 @@ auto builtin_marray_impl(FuncTy F, const Ts &...x) {
137137
marray<ret_elem_type, T::size()> Res;
138138
constexpr auto N = T::size();
139139
for (size_t I = 0; I < N / 2; ++I) {
140-
auto PartialRes = F(to_vec2(x, I * 2)...);
140+
auto PartialRes = [&]() {
141+
using elem_ty = get_elem_type_t<T>;
142+
if constexpr (std::is_integral_v<elem_ty>)
143+
return F(to_vec2(x, I * 2)
144+
.template as<vec<get_fixed_sized_int_t<elem_ty>, 2>>()...);
145+
else
146+
return F(to_vec2(x, I * 2)...);
147+
}();
141148
std::memcpy(&Res[I * 2], &PartialRes, sizeof(decltype(PartialRes)));
142149
}
143150
if (N % 2)

sycl/include/sycl/builtins_utils_scalar.hpp

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -128,6 +128,17 @@ template <size_t Size> struct get_unsigned_int_by_size {
128128
template <typename T> struct same_size_unsigned_int {
129129
using type = typename get_unsigned_int_by_size<sizeof(T)>::type;
130130
};
131+
template <typename T>
132+
using same_size_unsigned_int_t = typename same_size_unsigned_int<T>::type;
133+
134+
template <typename T> struct get_fixed_sized_int {
135+
static_assert(std::is_integral_v<T>);
136+
using type =
137+
std::conditional_t<std::is_signed_v<T>, same_size_signed_int_t<T>,
138+
same_size_unsigned_int_t<T>>;
139+
};
140+
template <typename T>
141+
using get_fixed_sized_int_t = typename get_fixed_sized_int<T>::type;
131142

132143
// Utility trait for getting an upsampled integer type.
133144
// NOTE: For upsampling we look for an integer of double the size of the

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

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -48,6 +48,11 @@
4848
FOR_EACH4_A6(BASE_CASE, FIXED1, FIXED2, FIXED3, FIXED4, ARG1, ARG2, ARG3, \
4949
ARG4, ARG5, ARG6) \
5050
BASE_CASE(FIXED1, FIXED2, FIXED3, FIXED4, ARG7)
51+
#define FOR_EACH4_A8(BASE_CASE, FIXED1, FIXED2, FIXED3, FIXED4, ARG1, ARG2, \
52+
ARG3, ARG4, ARG5, ARG6, ARG7, ARG8) \
53+
FOR_EACH4_A7(BASE_CASE, FIXED1, FIXED2, FIXED3, FIXED4, ARG1, ARG2, ARG3, \
54+
ARG4, ARG5, ARG6, ARG7) \
55+
BASE_CASE(FIXED1, FIXED2, FIXED3, FIXED4, ARG8)
5156
#define FOR_EACH4_A11(BASE_CASE, FIXED1, FIXED2, FIXED3, FIXED4, ARG1, ARG2, \
5257
ARG3, ARG4, ARG5, ARG6, ARG7, ARG8, ARG9, ARG10, ARG11) \
5358
FOR_EACH4_A7(BASE_CASE, FIXED1, FIXED2, FIXED3, FIXED4, ARG1, ARG2, ARG3, \
@@ -169,6 +174,9 @@
169174
unsigned char, unsigned short, unsigned int, unsigned long, unsigned long long
170175
// 11 types
171176
#define INTEGER_TYPES SIGNED_TYPES, UNSIGNED_TYPES
177+
// 8 types
178+
#define FIXED_WIDTH_INTEGER_TYPES \
179+
int8_t, int16_t, int32_t, int64_t, uint8_t, uint16_t, uint32_t, uint64_t
172180

173181
#define DEVICE_IMPL_TEMPLATE_CUSTOM_DELEGATE( \
174182
NUM_ARGS, NAME, ENABLER, DELEGATOR, NS, /*SCALAR_VEC_IMPL*/...) \

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

Lines changed: 4 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -16,9 +16,10 @@ namespace detail {
1616
template <typename T>
1717
struct integer_elem_type
1818
: std::bool_constant<
19-
check_type_in_v<get_elem_type_t<T>, char, signed char, short, int,
20-
long, long long, unsigned char, unsigned short,
21-
unsigned int, unsigned long, unsigned long long>> {};
19+
(is_vec_or_swizzle_v<T> &&
20+
check_type_in_v<get_elem_type_t<T>, FIXED_WIDTH_INTEGER_TYPES>) ||
21+
(!is_vec_or_swizzle_v<T> &&
22+
check_type_in_v<get_elem_type_t<T>, INTEGER_TYPES>)> {};
2223
template <typename T>
2324
struct suint32_elem_type
2425
: std::bool_constant<

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

Lines changed: 6 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -15,10 +15,12 @@ inline namespace _V1 {
1515
namespace detail {
1616
template <typename T>
1717
struct bitselect_elem_type
18-
: std::bool_constant<check_type_in_v<
19-
get_elem_type_t<T>, float, double, half, char, signed char, short,
20-
int, long, long long, unsigned char, unsigned short, unsigned int,
21-
unsigned long, unsigned long long>> {};
18+
: std::bool_constant<
19+
check_type_in_v<get_elem_type_t<T>, FP_TYPES> ||
20+
(is_vec_or_swizzle_v<T> &&
21+
check_type_in_v<get_elem_type_t<T>, FIXED_WIDTH_INTEGER_TYPES>) ||
22+
(!is_vec_or_swizzle_v<T> &&
23+
check_type_in_v<get_elem_type_t<T>, INTEGER_TYPES>)> {};
2224

2325
template <typename T>
2426
struct rel_ret_traits

sycl/source/builtins/host_helper_macros.hpp

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -56,6 +56,9 @@
5656
#define EXPORT_VEC(NUM_ARGS, NAME, TYPE, VL) \
5757
EXPORT_VEC_NS(NUM_ARGS, NAME, sycl, TYPE, VL)
5858

59+
#define EXPORT_VEC_1_16_IMPL(NUM_ARGS, NAME, NS, TYPE) \
60+
FOR_VEC_1_16(EXPORT_VEC_NS, NUM_ARGS, NAME, NS, TYPE)
61+
5962
#define EXPORT_SCALAR_AND_VEC_1_16_IMPL(NUM_ARGS, NAME, NS, TYPE) \
6063
EXPORT_SCALAR_NS(NUM_ARGS, NAME, NS, TYPE) \
6164
FOR_VEC_1_16(EXPORT_VEC_NS, NUM_ARGS, NAME, NS, TYPE)
@@ -69,8 +72,12 @@
6972

7073
#define EXPORT_SCALAR_AND_VEC_1_16_NS(NUM_ARGS, NAME, NS, ...) \
7174
FOR_EACH3(EXPORT_SCALAR_AND_VEC_1_16_IMPL, NUM_ARGS, NAME, NS, __VA_ARGS__)
75+
#define EXPORT_VEC_1_16_NS(NUM_ARGS, NAME, NS, ...) \
76+
FOR_EACH3(EXPORT_VEC_1_16_IMPL, NUM_ARGS, NAME, NS, __VA_ARGS__)
7277
#define EXPORT_SCALAR_AND_VEC_1_16(NUM_ARGS, NAME, ...) \
7378
EXPORT_SCALAR_AND_VEC_1_16_NS(NUM_ARGS, NAME, sycl, __VA_ARGS__)
79+
#define EXPORT_VEC_1_16(NUM_ARGS, NAME, ...) \
80+
EXPORT_VEC_1_16_NS(NUM_ARGS, NAME, sycl, __VA_ARGS__)
7481

7582
#define EXPORT_SCALAR_AND_VEC_2_4(NUM_ARGS, NAME, ...) \
7683
FOR_EACH2(EXPORT_SCALAR_AND_VEC_2_4_IMPL, NUM_ARGS, NAME, __VA_ARGS__)

sycl/source/builtins/integer_functions.cpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -76,7 +76,8 @@ namespace sycl {
7676
inline namespace _V1 {
7777
#define BUILTIN_GENINT(NUM_ARGS, NAME, IMPL) \
7878
HOST_IMPL(NAME, IMPL) \
79-
EXPORT_SCALAR_AND_VEC_1_16(NUM_ARGS, NAME, INTEGER_TYPES)
79+
FOR_EACH2(EXPORT_SCALAR, NUM_ARGS, NAME, INTEGER_TYPES) \
80+
EXPORT_VEC_1_16(NUM_ARGS, NAME, FIXED_WIDTH_INTEGER_TYPES)
8081
#define BUILTIN_GENINT_SU(NUM_ARGS, NAME, IMPL) \
8182
BUILTIN_GENINT(NUM_ARGS, NAME, IMPL)
8283

sycl/source/builtins/relational_functions.cpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -103,6 +103,7 @@ HOST_IMPL(bitselect, [](auto x, auto y, auto z) {
103103
assert((ures & std::numeric_limits<utype>::max()) == ures);
104104
return bit_cast<T0>(static_cast<utype>(ures));
105105
})
106-
EXPORT_SCALAR_AND_VEC_1_16(THREE_ARGS, bitselect, INTEGER_TYPES, FP_TYPES)
106+
FOR_EACH2(EXPORT_SCALAR, THREE_ARGS, bitselect, INTEGER_TYPES, FP_TYPES)
107+
EXPORT_VEC_1_16(THREE_ARGS, bitselect, FIXED_WIDTH_INTEGER_TYPES, FP_TYPES)
107108
} // namespace _V1
108109
} // namespace sycl
Lines changed: 133 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,133 @@
1+
// RUN: %clangxx -fsycl -fpreview-breaking-changes -fsyntax-only %s -Xclang -verify
2+
// REQUIRES: preview-breaking-changes-supported
3+
4+
#include <sycl/sycl.hpp>
5+
6+
using namespace sycl;
7+
using namespace sycl::detail;
8+
9+
namespace builtin_same_shape_v_tests {
10+
using swizzle1 = decltype(std::declval<vec<float, 2>>().swizzle<0>());
11+
using swizzle2 = decltype(std::declval<vec<float, 2>>().swizzle<0, 0>());
12+
using swizzle3 = decltype(std::declval<vec<float, 2>>().swizzle<0, 0, 1>());
13+
14+
static_assert(builtin_same_shape_v<float>);
15+
static_assert(builtin_same_shape_v<int, float>);
16+
static_assert(builtin_same_shape_v<marray<int, 2>>);
17+
static_assert(builtin_same_shape_v<marray<int, 2>, marray<float, 2>>);
18+
static_assert(builtin_same_shape_v<vec<int, 2>>);
19+
static_assert(builtin_same_shape_v<vec<int, 2>, vec<float, 2>>);
20+
static_assert(builtin_same_shape_v<vec<int, 2>, swizzle2>);
21+
22+
static_assert(!builtin_same_shape_v<float, marray<float, 1>>);
23+
static_assert(!builtin_same_shape_v<float, vec<float, 1>>);
24+
static_assert(!builtin_same_shape_v<marray<float, 1>, vec<float, 1>>);
25+
static_assert(!builtin_same_shape_v<float, swizzle1>);
26+
static_assert(!builtin_same_shape_v<marray<float, 1>, swizzle1>);
27+
static_assert(!builtin_same_shape_v<swizzle2, swizzle1>);
28+
} // namespace builtin_same_shape_v_tests
29+
30+
namespace builtin_marray_impl_tests {
31+
// Integer functions/relational bitselect only accept fixed-width integer
32+
// element types for vector/swizzle elements. Make sure that our marray->vec
33+
// delegator can handle that.
34+
35+
auto foo(char x) { return x; }
36+
auto foo(signed char x) { return x; }
37+
auto foo(unsigned char x) { return x; }
38+
auto foo(vec<int8_t, 2> x) { return x; }
39+
auto foo(vec<uint8_t, 2> x) { return x; }
40+
41+
auto test() {
42+
marray<char, 2> x;
43+
marray<signed char, 2> y;
44+
marray<unsigned char, 2> z;
45+
auto TestOne = [](auto x) {
46+
std::ignore = builtin_marray_impl([](auto x) { return foo(x); }, x);
47+
};
48+
TestOne(x);
49+
TestOne(y);
50+
TestOne(z);
51+
}
52+
} // namespace builtin_marray_impl_tests
53+
54+
namespace builtin_enable_integer_tests {
55+
using swizzle1 = decltype(std::declval<vec<int8_t, 2>>().swizzle<0>());
56+
using swizzle2 = decltype(std::declval<vec<int8_t, 2>>().swizzle<0, 0>());
57+
template <typename... Ts> void ignore() {}
58+
59+
void test() {
60+
// clang-format off
61+
ignore<builtin_enable_integer_t<char>,
62+
builtin_enable_integer_t<signed char>,
63+
builtin_enable_integer_t<unsigned char>>();
64+
// clang-format on
65+
66+
ignore<builtin_enable_integer_t<vec<int8_t, 2>>,
67+
builtin_enable_integer_t<vec<uint8_t, 2>>>();
68+
69+
ignore<builtin_enable_integer_t<char, char>>();
70+
ignore<builtin_enable_integer_t<vec<int8_t, 2>, vec<int8_t, 2>>>();
71+
ignore<builtin_enable_integer_t<vec<int8_t, 2>, swizzle2>>();
72+
ignore<builtin_enable_integer_t<swizzle2, swizzle2>>();
73+
74+
{
75+
// Only one of char/signed char maps onto int8_t. The other type isn't a
76+
// valid vector element type for integer builtins.
77+
78+
static_assert(std::is_signed_v<char>);
79+
80+
// clang-format off
81+
// 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>>'}}
82+
// expected-note@+1 {{in instantiation of template type alias 'builtin_enable_integer_t' requested here}}
83+
ignore<builtin_enable_integer_t<vec<signed char, 2>>, builtin_enable_integer_t<vec<char, 2>>>();
84+
// clang-format on
85+
}
86+
87+
// 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>'}}
88+
// expected-note@+1 {{in instantiation of template type alias 'builtin_enable_integer_t' requested here}}
89+
ignore<builtin_enable_integer_t<char, signed char>>();
90+
}
91+
} // namespace builtin_enable_integer_tests
92+
93+
namespace builtin_enable_bitselect_tests {
94+
// Essentially the same as builtin_enable_integer_t + FP types support.
95+
using swizzle1 = decltype(std::declval<vec<int8_t, 2>>().swizzle<0>());
96+
using swizzle2 = decltype(std::declval<vec<int8_t, 2>>().swizzle<0, 0>());
97+
template <typename... Ts> void ignore() {}
98+
99+
void test() {
100+
// clang-format off
101+
ignore<builtin_enable_bitselect_t<char>,
102+
builtin_enable_bitselect_t<signed char>,
103+
builtin_enable_bitselect_t<unsigned char>,
104+
builtin_enable_bitselect_t<float>>();
105+
// clang-format on
106+
107+
ignore<builtin_enable_bitselect_t<vec<int8_t, 2>>,
108+
builtin_enable_bitselect_t<vec<uint8_t, 2>>,
109+
builtin_enable_bitselect_t<vec<float, 2>>>();
110+
111+
ignore<builtin_enable_bitselect_t<char, char>>();
112+
ignore<builtin_enable_bitselect_t<vec<int8_t, 2>, vec<int8_t, 2>>>();
113+
ignore<builtin_enable_bitselect_t<vec<int8_t, 2>, swizzle2>>();
114+
ignore<builtin_enable_bitselect_t<swizzle2, swizzle2>>();
115+
116+
{
117+
// Only one of char/signed char maps onto int8_t. The other type isn't a
118+
// valid vector element type for integer builtins.
119+
120+
static_assert(std::is_signed_v<char>);
121+
122+
// clang-format off
123+
// 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>>'}}
124+
// expected-note@+1 {{in instantiation of template type alias 'builtin_enable_bitselect_t' requested here}}
125+
ignore<builtin_enable_bitselect_t<vec<signed char, 2>>, builtin_enable_bitselect_t<vec<char, 2>>>();
126+
// clang-format on
127+
}
128+
129+
// 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>'}}
130+
// expected-note@+1 {{in instantiation of template type alias 'builtin_enable_bitselect_t' requested here}}
131+
ignore<builtin_enable_bitselect_t<char, signed char>>();
132+
}
133+
} // namespace builtin_enable_bitselect_tests

0 commit comments

Comments
 (0)