23 #include <type_traits>
26 inline namespace _V1 {
28 template <
int N>
struct Boolean;
32 is_contained_v<T, gtl::scalar_vector_float_list>;
36 is_contained_v<T, gtl::scalar_vector_double_list>;
39 inline constexpr
bool is_half_v = is_contained_v<T, gtl::scalar_half_list>;
43 is_contained_v<T, gtl::scalar_bfloat16_list>;
47 is_contained_v<T, gtl::half_bfloat16_list>;
51 is_contained_v<T, gtl::scalar_vector_half_list>;
54 inline constexpr
bool is_genfloat_v = is_contained_v<T, gtl::floating_list>;
58 is_contained_v<T, gtl::scalar_floating_list>;
62 is_contained_v<T, gtl::vector_floating_list>;
66 is_contained_v<T, gtl::scalar_vector_floating_list>;
70 is_marray_v<T> && is_svgenfloat_v<get_elem_type_t<T>>;
77 is_contained_v<T, gtl::geo_double_list>;
81 is_contained_v<T, gtl::marray_geo_float_list>;
85 is_contained_v<T, gtl::marray_geo_list>;
92 is_contained_v<T, gtl::vector_geo_float_list>;
96 is_contained_v<T, gtl::vector_geo_double_list>;
100 is_contained_v<T, gtl::vector_geo_half_list>;
102 template <
typename T>
103 inline constexpr
bool is_sgengeo_v = is_contained_v<T, gtl::scalar_geo_list>;
105 template <
typename T>
106 inline constexpr
bool is_vgengeo_v = is_contained_v<T, gtl::vector_geo_list>;
108 template <
typename T>
110 is_contained_v<T, gtl::cross_float_list>;
112 template <
typename T>
114 is_contained_v<T, gtl::cross_double_list>;
116 template <
typename T>
118 is_contained_v<T, gtl::cross_half_list>;
120 template <
typename T>
122 is_contained_v<T, gtl::cross_floating_list>;
124 template <
typename T>
126 is_contained_v<T, gtl::cross_marray_list>;
128 template <
typename T>
129 inline constexpr
bool is_ugenint_v = is_contained_v<T, gtl::unsigned_int_list>;
131 template <
typename T>
133 is_contained_v<T, gtl::vector_signed_int_list>;
135 template <
typename T>
136 inline constexpr
bool is_genint_v = is_contained_v<T, gtl::signed_int_list>;
138 template <
typename T>
141 template <
typename T>
144 template <
typename T>
146 is_contained_v<T, gtl::signed_integer_list>;
148 template <
typename T>
151 template <
typename T>
153 is_contained_v<T, gtl::unsigned_integer_list>;
155 template <
typename T>
158 template <
typename T>
160 is_contained_v<T, gtl::scalar_integer_list>;
162 template <
typename T>
164 is_contained_v<T, gtl::vector_integer_list>;
166 template <
typename T>
168 is_contained_v<T, gtl::scalar_signed_integer_list>;
170 template <
typename T>
172 is_contained_v<T, gtl::scalar_unsigned_integer_list>;
174 template <
typename T>
176 is_contained_v<T, gtl::vector_signed_integer_list>;
178 template <
typename T>
180 is_contained_v<T, gtl::vector_unsigned_integer_list>;
182 template <
typename T>
183 inline constexpr
bool is_genbool_v = is_contained_v<T, gtl::bool_list>;
185 template <
typename T>
186 inline constexpr
bool is_gentype_v = is_contained_v<T, gtl::basic_list>;
188 template <
typename T>
189 inline constexpr
bool is_vgentype_v = is_contained_v<T, gtl::vector_basic_list>;
191 template <
typename T>
192 inline constexpr
bool is_sgentype_v = is_contained_v<T, gtl::scalar_basic_list>;
194 template <
typename T>
196 is_gen_based_on_type_sizeof_v<T, 1, is_igeninteger>;
198 template <
typename T>
200 is_gen_based_on_type_sizeof_v<T, 2, is_igeninteger>;
202 template <
typename T>
204 is_gen_based_on_type_sizeof_v<T, 4, is_igeninteger>;
206 template <
typename T>
208 is_gen_based_on_type_sizeof_v<T, 8, is_igeninteger>;
210 template <
typename T>
212 is_gen_based_on_type_sizeof_v<T, 1, is_ugeninteger>;
214 template <
typename T>
216 is_gen_based_on_type_sizeof_v<T, 2, is_ugeninteger>;
218 template <
typename T>
220 is_gen_based_on_type_sizeof_v<T, 4, is_ugeninteger>;
222 template <
typename T>
224 is_gen_based_on_type_sizeof_v<T, 8, is_ugeninteger>;
226 template <
typename T>
228 is_pointer_v<T> && is_genint_v<remove_pointer_t<T>> &&
229 is_address_space_compliant_v<T, gvl::nonconst_address_space_list>;
234 std::is_same_v<T, sycl::marray<marray_element_t<T>, T::size()>> &&
235 is_genint_v<marray_element_t<remove_pointer_t<T>>> &&
236 is_address_space_compliant_v<multi_ptr<T, AddressSpace, IsDecorated>,
241 template <
typename T>
243 is_pointer_v<T> && is_genfloat_v<remove_pointer_t<T>> &&
244 is_address_space_compliant_v<T, gvl::nonconst_address_space_list>;
250 is_address_space_compliant_v<multi_ptr<T, AddressSpace, IsDecorated>,
255 template <
typename T>
258 template <
typename T>
261 template <
typename T>
267 template <
typename T,
typename T8,
typename T16,
typename T32,
typename T64>
270 std::conditional_t<
sizeof(T) == 2, T16,
271 std::conditional_t<
sizeof(T) == 4, T32, T64>>>;
274 template <
typename T>
279 template <
typename T>
293 template <
typename T>
298 using no_ref = std::remove_reference_t<T>;
299 if constexpr (is_multi_ptr_v<no_ref>) {
301 }
else if constexpr (std::is_pointer_v<no_ref>) {
305 std::declval<std::remove_const_t<elem_type>>()));
306 using converted_elem_type =
307 std::conditional_t<std::is_const_v<elem_type>,
308 const converted_elem_type_no_cv,
309 converted_elem_type_no_cv>;
310 #ifdef __SYCL_DEVICE_ONLY__
313 deduce_AS<no_ref>::value>::type *;
315 using result_type = converted_elem_type *;
317 return reinterpret_cast<result_type
>(
x);
318 }
else if constexpr (is_vec_v<no_ref>) {
319 using ElemTy =
typename no_ref::element_type;
325 std::declval<ElemTy>()))>,
327 #ifdef __SYCL_DEVICE_ONLY__
331 if constexpr (std::is_same_v<MatchingVec, no_ref>)
332 return static_cast<typename MatchingVec::vector_t
>(
x);
334 return static_cast<typename MatchingVec::vector_t
>(
335 x.template as<MatchingVec>());
337 return x.template as<MatchingVec>();
339 }
else if constexpr (is_boolean_v<no_ref>) {
340 #ifdef __SYCL_DEVICE_ONLY__
341 if constexpr (std::is_same_v<
Boolean<1>, no_ref>) {
343 return std::forward<T>(
x);
345 return static_cast<typename no_ref::vector_t
>(
x);
348 return std::forward<T>(
x);
350 #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
351 }
else if constexpr (std::is_same_v<no_ref, std::byte>) {
352 return static_cast<uint8_t
>(
x);
354 }
else if constexpr (std::is_integral_v<no_ref>) {
356 std::conditional_t<std::is_signed_v<no_ref>,
359 static_assert(
sizeof(OpenCLType) ==
sizeof(T));
360 return static_cast<OpenCLType
>(
x);
361 }
else if constexpr (is_half_v<no_ref>) {
363 static_assert(
sizeof(OpenCLType) ==
sizeof(T));
364 return static_cast<OpenCLType
>(
x);
365 }
else if constexpr (is_bfloat16_v<no_ref>) {
366 return std::forward<T>(
x);
367 }
else if constexpr (std::is_floating_point_v<no_ref>) {
368 static_assert(std::is_same_v<no_ref, float> ||
369 std::is_same_v<no_ref, double>,
370 "Other FP types are not expected/supported (yet?)");
371 static_assert(std::is_same_v<float, sycl::opencl::cl_float> &&
372 std::is_same_v<double, sycl::opencl::cl_double>);
373 return std::forward<T>(
x);
376 static_assert(
sizeof(OpenCLType) ==
sizeof(T));
377 return static_cast<OpenCLType
>(
x);
381 template <
typename T>
385 if constexpr (std::is_same_v<To, bool> &&
386 std::is_same_v<std::remove_reference_t<From>,
bool>) {
391 static_assert(std::is_same_v<std::remove_reference_t<From>, OpenCLType>);
392 static_assert(
sizeof(OpenCLType) ==
sizeof(To));
393 using To_noref = std::remove_reference_t<To>;
394 using From_noref = std::remove_reference_t<From>;
395 if constexpr (is_vec_v<To_noref> && is_vec_v<From_noref>)
396 return x.template as<To_noref>();
397 else if constexpr (is_vec_v<To_noref> && is_ext_vector_v<From_noref>)
398 return To_noref{bit_cast<typename To_noref::vector_t>(
x)};
400 return static_cast<To
>(
x);
405 template <
typename T>
inline constexpr T
msbMask(T) {
407 return T(UT(1) << (
sizeof(T) * 8 - 1));
410 template <
typename T>
inline constexpr
bool msbIsSet(
const T x) {
418 template <
typename Type,
int NumElements>
420 static constexpr
int value = NumElements;
430 template <
typename A>
static typename A::element_type check(
const A &);
431 template <
typename A,
typename = std::enable_if_t<!std::is_same_v<
436 using type = decltype(check(T()));
437 static constexpr
bool value = !std::is_same_v<T, type>;
440 template <
typename T>
static constexpr T
max_v() {
444 template <
typename T>
static constexpr T
min_v() {
decltype(check(T())) type
static constexpr bool value
Provides a cross-patform vector class template that works efficiently on SYCL devices as well as in h...
auto convertToOpenCLType(T &&x)
constexpr bool is_vgengeo_v
constexpr bool is_sgeninteger_v
constexpr bool is_ugeninteger8bit_v
constexpr bool is_svgenfloatd_v
constexpr bool is_gentype_v
std::bool_constant< is_ugeninteger_v< T > > is_ugeninteger
constexpr bool is_ugeninteger16bit_v
constexpr bool is_ugeninteger64bit_v
constexpr bool is_sugeninteger_v
constexpr bool is_ugeninteger_v
make_type_t< T, gtl::scalar_unsigned_integer_list > make_unsinged_integer_t
constexpr bool is_mgenfloat_v
constexpr bool is_igeninteger8bit_v
constexpr bool is_gengeofloat_v
static constexpr T min_v()
constexpr bool is_vugeninteger_v
constexpr bool is_sgengeo_v
std::conditional_t< sizeof(T)==1, T8, std::conditional_t< sizeof(T)==2, T16, std::conditional_t< sizeof(T)==4, T32, T64 > >> select_apply_cl_scalar_t
typename make_unsigned< T >::type make_unsigned_t
constexpr bool is_genfloat_v
constexpr bool is_vgengeofloat_v
constexpr bool is_igeninteger_v
make_type_t< T, gtl::scalar_signed_integer_list > make_singed_integer_t
constexpr bool is_ugenint_v
constexpr bool is_bfloat16_v
auto convertFromOpenCLTypeFor(From &&x)
select_apply_cl_scalar_t< T, sycl::opencl::cl_uchar, sycl::opencl::cl_ushort, sycl::opencl::cl_uint, sycl::opencl::cl_ulong > select_cl_scalar_integral_unsigned_t
constexpr bool is_genfloatptr_v
constexpr bool is_svgenfloath_v
constexpr bool is_sgenfloat_v
constexpr bool is_vgeninteger_v
constexpr bool is_sgentype_v
constexpr bool is_vgentype_v
constexpr bool is_gencrossmarray_v
constexpr bool is_ugeninteger32bit_v
constexpr bool is_igeninteger32bit_v
constexpr bool is_gencrosshalf_v
static constexpr T max_v()
constexpr bool is_igeninteger64bit_v
constexpr bool is_gengeodouble_v
constexpr bool is_gencross_v
constexpr bool is_svgenfloatf_v
constexpr bool is_vgenfloat_v
std::bool_constant< is_geninteger_v< T > > is_geninteger
constexpr bool is_geninteger_v
constexpr bool msbIsSet(const T x)
decltype(convertToOpenCLType(std::declval< T >())) ConvertToOpenCLType_t
constexpr bool is_gencrossfloat_v
constexpr bool is_genintptr_v
constexpr bool is_genintptr_marray_v
constexpr bool is_vgengeohalf_v
constexpr bool is_svgenfloat_v
constexpr bool is_vigeninteger_v
constexpr bool is_half_or_bf16_v
constexpr bool is_gengeomarray_v
std::bool_constant< is_igeninteger_v< T > > is_igeninteger
make_type_t< T, gtl::scalar_floating_list > make_floating_point_t
constexpr bool is_genfloatptr_marray_v
typename select_cl_scalar_complex_or_T< T >::type select_cl_scalar_complex_or_T_t
constexpr bool is_vgengeodouble_v
constexpr bool is_gengeohalf_v
constexpr bool is_gencrossdouble_v
typename make_type_impl< T, TL >::type make_type_t
select_apply_cl_scalar_t< T, sycl::opencl::cl_char, sycl::opencl::cl_short, sycl::opencl::cl_int, sycl::opencl::cl_long > select_cl_scalar_integral_signed_t
constexpr bool is_genint_v
constexpr bool is_gengeomarrayfloat_v
constexpr bool is_igeninteger16bit_v
constexpr bool is_sigeninteger_v
constexpr bool is_genbool_v
typename remove_decoration< T >::type remove_decoration_t
static constexpr int value