DPC++ Runtime
Runtime libraries for oneAPI DPC++
vector_convert.hpp
Go to the documentation of this file.
1 //==-- vector_convert.hpp --- vec::convert implementation ------------------==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 
54 
55 #pragma once
56 
57 #include <sycl/detail/generic_type_traits.hpp> // for is_sigeninteger, is_s...
58 #include <sycl/exception.hpp> // for errc
59 
60 #include <sycl/ext/oneapi/bfloat16.hpp> // bfloat16
61 
62 #ifndef __SYCL_DEVICE_ONLY__
63 #include <cfenv> // for fesetround, fegetround
64 #endif
65 
66 #include <type_traits>
67 
68 // Enable on only intel devices.
69 #if defined(__SYCL_DEVICE_ONLY__) && (defined(__SPIR__) || defined(__SPIRV__))
70 extern "C" {
71 // For converting BF16 to other types.
72 extern __DPCPP_SYCL_EXTERNAL float __imf_bfloat162float(uint16_t x);
73 extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_bfloat162uint_rd(uint16_t x);
74 extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_bfloat162uint_rn(uint16_t x);
75 extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_bfloat162uint_ru(uint16_t x);
76 extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_bfloat162uint_rz(uint16_t x);
77 extern __DPCPP_SYCL_EXTERNAL unsigned short
78 __imf_bfloat162ushort_rd(uint16_t x);
79 extern __DPCPP_SYCL_EXTERNAL unsigned short
80 __imf_bfloat162ushort_rn(uint16_t x);
81 extern __DPCPP_SYCL_EXTERNAL unsigned short
82 __imf_bfloat162ushort_ru(uint16_t x);
83 extern __DPCPP_SYCL_EXTERNAL unsigned short
84 __imf_bfloat162ushort_rz(uint16_t x);
85 extern __DPCPP_SYCL_EXTERNAL unsigned long long
86 __imf_bfloat162ull_rd(uint16_t x);
87 extern __DPCPP_SYCL_EXTERNAL unsigned long long
88 __imf_bfloat162ull_rn(uint16_t x);
89 extern __DPCPP_SYCL_EXTERNAL unsigned long long
90 __imf_bfloat162ull_ru(uint16_t x);
91 extern __DPCPP_SYCL_EXTERNAL unsigned long long
92 __imf_bfloat162ull_rz(uint16_t x);
93 extern __DPCPP_SYCL_EXTERNAL int __imf_bfloat162int_rd(uint16_t x);
94 extern __DPCPP_SYCL_EXTERNAL int __imf_bfloat162int_rn(uint16_t x);
95 extern __DPCPP_SYCL_EXTERNAL int __imf_bfloat162int_ru(uint16_t x);
96 extern __DPCPP_SYCL_EXTERNAL int __imf_bfloat162int_rz(uint16_t x);
97 extern __DPCPP_SYCL_EXTERNAL short __imf_bfloat162short_rd(uint16_t x);
98 extern __DPCPP_SYCL_EXTERNAL short __imf_bfloat162short_rn(uint16_t x);
99 extern __DPCPP_SYCL_EXTERNAL short __imf_bfloat162short_ru(uint16_t x);
100 extern __DPCPP_SYCL_EXTERNAL short __imf_bfloat162short_rz(uint16_t x);
101 extern __DPCPP_SYCL_EXTERNAL long long __imf_bfloat162ll_rd(uint16_t x);
102 extern __DPCPP_SYCL_EXTERNAL long long __imf_bfloat162ll_rn(uint16_t x);
103 extern __DPCPP_SYCL_EXTERNAL long long __imf_bfloat162ll_ru(uint16_t x);
104 extern __DPCPP_SYCL_EXTERNAL long long __imf_bfloat162ll_rz(uint16_t x);
105 extern __DPCPP_SYCL_EXTERNAL short __imf_bfloat16_as_short(uint16_t x);
106 extern __DPCPP_SYCL_EXTERNAL unsigned short
107 __imf_bfloat16_as_ushort(uint16_t x);
108 
109 // For converting other types to BF16.
110 extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_float2bfloat16(float x);
111 extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_float2bfloat16_rd(float x);
112 extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_float2bfloat16_rn(float x);
113 extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_float2bfloat16_ru(float x);
114 extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_float2bfloat16_rz(float x);
115 extern __DPCPP_SYCL_EXTERNAL uint16_t
116 __imf_ushort2bfloat16_rd(unsigned short x);
117 extern __DPCPP_SYCL_EXTERNAL uint16_t
118 __imf_ushort2bfloat16_rn(unsigned short x);
119 extern __DPCPP_SYCL_EXTERNAL uint16_t
120 __imf_ushort2bfloat16_ru(unsigned short x);
121 extern __DPCPP_SYCL_EXTERNAL uint16_t
122 __imf_ushort2bfloat16_rz(unsigned short x);
123 extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_uint2bfloat16_rd(unsigned int x);
124 extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_uint2bfloat16_rn(unsigned int x);
125 extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_uint2bfloat16_ru(unsigned int x);
126 extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_uint2bfloat16_rz(unsigned int x);
127 extern __DPCPP_SYCL_EXTERNAL uint16_t
128 __imf_ull2bfloat16_rd(unsigned long long x);
129 extern __DPCPP_SYCL_EXTERNAL uint16_t
130 __imf_ull2bfloat16_rn(unsigned long long x);
131 extern __DPCPP_SYCL_EXTERNAL uint16_t
132 __imf_ull2bfloat16_ru(unsigned long long x);
133 extern __DPCPP_SYCL_EXTERNAL uint16_t
134 __imf_ull2bfloat16_rz(unsigned long long x);
135 extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_short2bfloat16_rd(short x);
136 extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_short2bfloat16_rn(short x);
137 extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_short2bfloat16_ru(short x);
138 extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_short2bfloat16_rz(short x);
139 extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_int2bfloat16_rd(int x);
140 extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_int2bfloat16_rn(int x);
141 extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_int2bfloat16_ru(int x);
142 extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_int2bfloat16_rz(int x);
143 extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_ll2bfloat16_rd(long long x);
144 extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_ll2bfloat16_rn(long long x);
145 extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_ll2bfloat16_ru(long long x);
146 extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_ll2bfloat16_rz(long long x);
147 extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_double2bfloat16(double x);
148 extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_short_as_bfloat16(short x);
149 extern __DPCPP_SYCL_EXTERNAL uint16_t
150 __imf_ushort_as_bfloat16(unsigned short x);
151 }
152 #endif // __SYCL_DEVICE_ONLY__ && (defined(__SPIR__) || defined(__SPIRV__))
153 
154 namespace sycl {
155 
156 enum class rounding_mode { automatic = 0, rte = 1, rtz = 2, rtp = 3, rtn = 4 };
157 
158 inline namespace _V1 {
159 #ifndef __SYCL_DEVICE_ONLY__
160 // TODO: Refactor includes so we can just "#include".
161 inline float ceil(float);
162 inline double ceil(double);
163 inline float floor(float);
164 inline double floor(double);
165 inline float rint(float);
166 inline double rint(double);
167 inline float trunc(float);
168 inline double trunc(double);
169 #endif
170 namespace detail {
171 
172 template <typename FromT, typename ToT, sycl::rounding_mode RoundingMode,
173  int VecSize, typename NativeFromT, typename NativeToT>
174 NativeToT convertImpl(NativeFromT);
175 
176 template <typename T, typename R>
177 using is_sint_to_sint =
178  std::bool_constant<is_sigeninteger_v<T> && is_sigeninteger_v<R>>;
179 
180 template <typename T, typename R>
181 using is_uint_to_uint =
182  std::bool_constant<is_sugeninteger_v<T> && is_sugeninteger_v<R>>;
183 
184 template <typename T, typename R>
185 using is_sint_to_from_uint = std::bool_constant<
186  (detail::is_sigeninteger_v<T> && detail::is_sugeninteger_v<R>) ||
187  (detail::is_sugeninteger_v<T> && detail::is_sigeninteger_v<R>)>;
188 
189 template <typename T, typename R>
190 using is_sint_to_float =
191  std::bool_constant<std::is_integral_v<T> && !std::is_unsigned_v<T> &&
192  detail::is_floating_point<R>::value>;
193 
194 template <typename T, typename R>
195 using is_uint_to_float =
196  std::bool_constant<std::is_unsigned_v<T> &&
197  detail::is_floating_point<R>::value>;
198 
199 template <typename T, typename R>
200 using is_int_to_float = std::bool_constant<std::is_integral_v<T> &&
201  detail::is_floating_point<R>::value>;
202 
203 template <typename T, typename R>
204 using is_float_to_uint =
205  std::bool_constant<detail::is_floating_point<T>::value &&
206  std::is_unsigned_v<R>>;
207 
208 template <typename T, typename R>
209 using is_float_to_sint =
210  std::bool_constant<detail::is_floating_point<T>::value &&
211  std::is_integral_v<R> && !std::is_unsigned_v<R>>;
212 
213 template <typename T, typename R>
214 using is_float_to_float =
215  std::bool_constant<detail::is_floating_point<T>::value &&
216  detail::is_floating_point<R>::value>;
217 
218 using bfloat16 = sycl::ext::oneapi::bfloat16;
219 
220 #ifndef __SYCL_DEVICE_ONLY__
221 template <typename From, typename To, int VecSize,
222  typename Enable = std::enable_if_t<VecSize == 1>>
223 To SConvert(From Value) {
224  return static_cast<To>(Value);
225 }
226 
227 template <typename From, typename To, int VecSize,
228  typename Enable = std::enable_if_t<VecSize == 1>>
229 To UConvert(From Value) {
230  return static_cast<To>(Value);
231 }
232 
233 template <typename From, typename To, int VecSize,
234  typename Enable = std::enable_if_t<VecSize == 1>>
235 To ConvertSToF(From Value) {
236  return static_cast<To>(Value);
237 }
238 
239 template <typename From, typename To, int VecSize,
240  typename Enable = std::enable_if_t<VecSize == 1>>
241 To ConvertUToF(From Value) {
242  return static_cast<To>(Value);
243 }
244 
245 template <typename From, typename To, int VecSize,
246  typename Enable = std::enable_if_t<VecSize == 1>,
247  sycl::rounding_mode RM>
248 To FConvert(From Value) {
249  return static_cast<To>(Value);
250 }
251 
252 template <typename From, typename To, int VecSize,
253  typename Enable = std::enable_if_t<VecSize == 1>,
254  sycl::rounding_mode roundingMode>
255 To ConvertFToS(From Value) {
256  switch (roundingMode) {
257  // Round to nearest even is default rounding mode for floating-point types
258  case rounding_mode::automatic:
259  // Round to nearest even.
260  case rounding_mode::rte: {
261  int OldRoundingDirection = std::fegetround();
262  int Err = std::fesetround(FE_TONEAREST);
263  if (Err)
264  throw sycl::exception(make_error_code(errc::runtime),
265  "Unable to set rounding mode to FE_TONEAREST");
266  To Result = sycl::rint(Value);
267  Err = std::fesetround(OldRoundingDirection);
268  if (Err)
269  throw sycl::exception(make_error_code(errc::runtime),
270  "Unable to restore rounding mode.");
271  return Result;
272  }
273  // Round toward zero.
274  case rounding_mode::rtz:
275  return sycl::trunc(Value);
276  // Round toward positive infinity.
277  case rounding_mode::rtp:
278  return sycl::ceil(Value);
279  // Round toward negative infinity.
280  case rounding_mode::rtn:
281  return sycl::floor(Value);
282  };
283  assert(false && "Unsupported rounding mode!");
284  return static_cast<To>(Value);
285 }
286 
287 template <typename From, typename To, int VecSize,
288  typename Enable = std::enable_if_t<VecSize == 1>,
289  sycl::rounding_mode roundingMode>
290 To ConvertFToU(From Value) {
291  return ConvertFToS<From, To, VecSize, Enable, roundingMode>(Value);
292 }
293 
294 template <typename NativeToT, sycl::rounding_mode RoundingMode>
295 inline NativeToT ConvertFromBF16Scalar(bfloat16 val) {
296  // On host, NativeBF16T is bfloat16. Convert BF16 to float losslessly.
297  float fval = static_cast<float>(val);
298 
299  if constexpr (std::is_same_v<NativeToT, float>)
300  return fval;
301  else
302  // Convert float to the desired type.
303  return convertImpl<float, NativeToT, RoundingMode, 1, float, NativeToT>(
304  fval);
305 }
306 
307 template <typename NativeFromT, sycl::rounding_mode RoundingMode>
308 bfloat16 ConvertToBF16Scalar(NativeFromT val) {
309 
310  constexpr int rm = static_cast<int>(RoundingMode);
311  return sycl::ext::oneapi::detail::ConvertToBfloat16::
312  getBfloat16WithRoundingMode<NativeFromT, rm>(val);
313 }
314 
315 #else
316 // Bunch of helpers to "specialize" each template for its own destination type
317 // and vector size.
318 
319 // Added for unification, to be able to have single enable_if-like trait for all
320 // cases regardless of whether rounding mode is actually applicable or not.
321 template <rounding_mode Mode> using AnyRM = std::bool_constant<true>;
322 
323 template <rounding_mode Mode>
324 using RteOrAutomatic = std::bool_constant<Mode == rounding_mode::automatic ||
325  Mode == rounding_mode::rte>;
326 
327 template <rounding_mode Mode>
328 using Rtz = std::bool_constant<Mode == rounding_mode::rtz>;
329 
330 template <rounding_mode Mode>
331 using Rtp = std::bool_constant<Mode == rounding_mode::rtp>;
332 
333 template <rounding_mode Mode>
334 using Rtn = std::bool_constant<Mode == rounding_mode::rtn>;
335 
336 template <int VecSize> using IsScalar = std::bool_constant<VecSize == 1>;
337 
338 template <int ExpectedVecSize, int ActualVecSize>
339 using IsVectorOf = std::bool_constant<ActualVecSize == ExpectedVecSize>;
340 
341 // This is a key condition for "specializations" below: it helps restrict each
342 // "specialization" to a (mostly) single type with one exception for
343 // signed char -> char case.
344 template <typename ExpectedType, typename ActualType>
345 using IsExpectedIntType =
346  std::bool_constant<std::is_same_v<ExpectedType, ActualType> ||
347  (std::is_same_v<ExpectedType, sycl::opencl::cl_char> &&
348  std::is_same_v<ActualType, signed char>)>;
349 
350 // Helpers which are used for conversions to an integer type
351 template <typename ExpectedType, typename ActualType, int VecSize,
352  typename ReturnType,
353  template <sycl::rounding_mode> typename RoundingModeCondition,
354  sycl::rounding_mode RoundingMode>
355 struct enable_if_to_int_scalar
356  : std::enable_if<IsExpectedIntType<ExpectedType, ActualType>::value &&
357  IsScalar<VecSize>::value &&
358  RoundingModeCondition<RoundingMode>::value,
359  ReturnType> {};
360 
361 template <typename ExpectedType, typename ActualType, int VecSize,
362  typename ReturnType,
363  template <sycl::rounding_mode> typename RoundingModeCondition = AnyRM,
364  sycl::rounding_mode RoundingMode = sycl::rounding_mode::automatic>
365 using enable_if_to_int_scalar_t =
366  typename enable_if_to_int_scalar<ExpectedType, ActualType, VecSize,
367  ReturnType, RoundingModeCondition,
368  RoundingMode>::type;
369 
370 template <typename ExpectedType, typename ActualType, int ExpectedVecSize,
371  int ActualVecSize, typename ReturnType,
372  template <sycl::rounding_mode> typename RoundingModeCondition,
373  sycl::rounding_mode RoundingMode>
374 struct enable_if_to_int_vector
375  : std::enable_if<IsExpectedIntType<ExpectedType, ActualType>::value &&
376  IsVectorOf<ExpectedVecSize, ActualVecSize>::value &&
377  RoundingModeCondition<RoundingMode>::value,
378  ReturnType> {};
379 
380 template <typename ExpectedType, typename ActualType, int ExpectedVecSize,
381  int ActualVecSize, typename ReturnType,
382  template <sycl::rounding_mode> typename RoundingModeCondition = AnyRM,
383  sycl::rounding_mode RoundingMode = sycl::rounding_mode::automatic>
384 using enable_if_to_int_vector_t =
385  typename enable_if_to_int_vector<ExpectedType, ActualType, ExpectedVecSize,
386  ActualVecSize, ReturnType,
387  RoundingModeCondition, RoundingMode>::type;
388 
389 // signed to signed, unsigned to unsigned conversions
390 #define __SYCL_SCALAR_INT_INT_CONVERT(Op, DestType) \
391  template <typename From, typename To, int VecSize, typename Enable> \
392  enable_if_to_int_scalar_t<sycl::opencl::cl_##DestType, Enable, VecSize, To> \
393  Op##Convert(From value) { \
394  return __spirv_##Op##Convert_R##DestType(value); \
395  }
396 
397 #define __SYCL_VECTOR_INT_INT_CONVERT(Op, N, DestType) \
398  template <typename From, typename To, int VecSize, typename Enable> \
399  enable_if_to_int_vector_t<sycl::opencl::cl_##DestType, Enable, N, VecSize, \
400  To> \
401  Op##Convert(From value) { \
402  return __spirv_##Op##Convert_R##DestType##N(value); \
403  }
404 
405 #define __SYCL_INT_INT_CONVERT(Op, DestType) \
406  __SYCL_SCALAR_INT_INT_CONVERT(Op, DestType) \
407  __SYCL_VECTOR_INT_INT_CONVERT(Op, 2, DestType) \
408  __SYCL_VECTOR_INT_INT_CONVERT(Op, 3, DestType) \
409  __SYCL_VECTOR_INT_INT_CONVERT(Op, 4, DestType) \
410  __SYCL_VECTOR_INT_INT_CONVERT(Op, 8, DestType) \
411  __SYCL_VECTOR_INT_INT_CONVERT(Op, 16, DestType)
412 
413 __SYCL_INT_INT_CONVERT(S, char)
414 __SYCL_INT_INT_CONVERT(S, short)
415 __SYCL_INT_INT_CONVERT(S, int)
416 __SYCL_INT_INT_CONVERT(S, long)
417 
418 __SYCL_INT_INT_CONVERT(U, uchar)
419 __SYCL_INT_INT_CONVERT(U, ushort)
420 __SYCL_INT_INT_CONVERT(U, uint)
421 __SYCL_INT_INT_CONVERT(U, ulong)
422 
423 #undef __SYCL_SCALAR_INT_INT_CONVERT
424 #undef __SYCL_VECTOR_INT_INT_CONVERT
425 #undef __SYCL_INT_INT_CONVERT
426 
427 // float to signed, float to unsigned conversion
428 #define __SYCL_SCALAR_FLOAT_INT_CONVERT(Op, DestType, RoundingMode, \
429  RoundingModeCondition) \
430  template <typename From, typename To, int VecSize, typename Enable, \
431  sycl::rounding_mode RM> \
432  enable_if_to_int_scalar_t<sycl::opencl::cl_##DestType, Enable, VecSize, To, \
433  RoundingModeCondition, RM> \
434  Convert##Op(From Value) { \
435  return __spirv_Convert##Op##_R##DestType##_##RoundingMode(Value); \
436  }
437 
438 #define __SYCL_VECTOR_FLOAT_INT_CONVERT(Op, N, DestType, RoundingMode, \
439  RoundingModeCondition) \
440  template <typename From, typename To, int VecSize, typename Enable, \
441  sycl::rounding_mode RM> \
442  enable_if_to_int_vector_t<sycl::opencl::cl_##DestType, Enable, N, VecSize, \
443  To, RoundingModeCondition, RM> \
444  Convert##Op(From Value) { \
445  return __spirv_Convert##Op##_R##DestType##N##_##RoundingMode(Value); \
446  }
447 
448 #define __SYCL_FLOAT_INT_CONVERT(Op, DestType, RoundingMode, \
449  RoundingModeCondition) \
450  __SYCL_SCALAR_FLOAT_INT_CONVERT(Op, DestType, RoundingMode, \
451  RoundingModeCondition) \
452  __SYCL_VECTOR_FLOAT_INT_CONVERT(Op, 2, DestType, RoundingMode, \
453  RoundingModeCondition) \
454  __SYCL_VECTOR_FLOAT_INT_CONVERT(Op, 3, DestType, RoundingMode, \
455  RoundingModeCondition) \
456  __SYCL_VECTOR_FLOAT_INT_CONVERT(Op, 4, DestType, RoundingMode, \
457  RoundingModeCondition) \
458  __SYCL_VECTOR_FLOAT_INT_CONVERT(Op, 8, DestType, RoundingMode, \
459  RoundingModeCondition) \
460  __SYCL_VECTOR_FLOAT_INT_CONVERT(Op, 16, DestType, RoundingMode, \
461  RoundingModeCondition)
462 
463 #define __SYCL_FLOAT_INT_CONVERT_FOR_TYPE(Op, DestType) \
464  __SYCL_FLOAT_INT_CONVERT(Op, DestType, rte, RteOrAutomatic) \
465  __SYCL_FLOAT_INT_CONVERT(Op, DestType, rtz, Rtz) \
466  __SYCL_FLOAT_INT_CONVERT(Op, DestType, rtp, Rtp) \
467  __SYCL_FLOAT_INT_CONVERT(Op, DestType, rtn, Rtn)
468 
469 __SYCL_FLOAT_INT_CONVERT_FOR_TYPE(FToS, char)
470 __SYCL_FLOAT_INT_CONVERT_FOR_TYPE(FToS, short)
471 __SYCL_FLOAT_INT_CONVERT_FOR_TYPE(FToS, int)
472 __SYCL_FLOAT_INT_CONVERT_FOR_TYPE(FToS, long)
473 
474 __SYCL_FLOAT_INT_CONVERT_FOR_TYPE(FToU, uchar)
475 __SYCL_FLOAT_INT_CONVERT_FOR_TYPE(FToU, ushort)
476 __SYCL_FLOAT_INT_CONVERT_FOR_TYPE(FToU, uint)
477 __SYCL_FLOAT_INT_CONVERT_FOR_TYPE(FToU, ulong)
478 
479 #undef __SYCL_SCALAR_FLOAT_INT_CONVERT
480 #undef __SYCL_VECTOR_FLOAT_INT_CONVERT
481 #undef __SYCL_FLOAT_INT_CONVERT
482 #undef __SYCL_FLOAT_INT_CONVERT_FOR_TYPE
483 
484 // Helpers which are used for conversions to a floating-point type
485 template <typename ExpectedType, typename ActualType>
486 using IsExpectedFloatType =
487  std::bool_constant<std::is_same_v<ExpectedType, ActualType> ||
488  (std::is_same_v<ExpectedType, sycl::opencl::cl_half> &&
489  std::is_same_v<ActualType, _Float16>)>;
490 
491 template <typename ExpectedType, typename ActualType, int VecSize,
492  typename ReturnType,
493  template <sycl::rounding_mode> typename RoundingModeCondition,
494  sycl::rounding_mode RoundingMode>
495 struct enable_if_to_float_scalar
496  : std::enable_if<IsExpectedFloatType<ExpectedType, ActualType>::value &&
497  IsScalar<VecSize>::value &&
498  RoundingModeCondition<RoundingMode>::value,
499  ReturnType> {};
500 
501 template <typename ExpectedType, typename ActualType, int VecSize,
502  typename ReturnType,
503  template <sycl::rounding_mode> typename RoundingModeCondition = AnyRM,
504  sycl::rounding_mode RoundingMode = sycl::rounding_mode::automatic>
505 using enable_if_to_float_scalar_t =
506  typename enable_if_to_float_scalar<ExpectedType, ActualType, VecSize,
507  ReturnType, RoundingModeCondition,
508  RoundingMode>::type;
509 
510 template <typename ExpectedType, typename ActualType, int ExpectedVecSize,
511  int ActualVecSize, typename ReturnType,
512  template <sycl::rounding_mode> typename RoundingModeCondition,
513  sycl::rounding_mode RoundingMode>
514 struct enable_if_to_float_vector
515  : std::enable_if<IsExpectedFloatType<ExpectedType, ActualType>::value &&
516  IsVectorOf<ExpectedVecSize, ActualVecSize>::value &&
517  RoundingModeCondition<RoundingMode>::value,
518  ReturnType> {};
519 
520 template <typename ExpectedType, typename ActualType, int ExpectedVecSize,
521  int ActualVecSize, typename ReturnType,
522  template <sycl::rounding_mode> typename RoundingModeCondition = AnyRM,
523  sycl::rounding_mode RoundingMode = sycl::rounding_mode::automatic>
524 using enable_if_to_float_vector_t = typename enable_if_to_float_vector<
525  ExpectedType, ActualType, ExpectedVecSize, ActualVecSize, ReturnType,
526  RoundingModeCondition, RoundingMode>::type;
527 
528 // signed to float, unsigned to float conversions
529 #define __SYCL_SCALAR_INT_FLOAT_CONVERT(Op, DestType) \
530  template <typename From, typename To, int VecSize, typename Enable> \
531  enable_if_to_float_scalar_t<sycl::opencl::cl_##DestType, Enable, VecSize, \
532  To> \
533  Convert##Op(From value) { \
534  return __spirv_Convert##Op##_R##DestType(value); \
535  }
536 
537 #define __SYCL_VECTOR_INT_FLOAT_CONVERT(Op, N, DestType) \
538  template <typename From, typename To, int VecSize, typename Enable> \
539  enable_if_to_float_vector_t<sycl::opencl::cl_##DestType, Enable, N, VecSize, \
540  To> \
541  Convert##Op(From value) { \
542  return __spirv_Convert##Op##_R##DestType##N(value); \
543  }
544 
545 #define __SYCL_INT_FLOAT_CONVERT(Op, DestType) \
546  __SYCL_SCALAR_INT_FLOAT_CONVERT(Op, DestType) \
547  __SYCL_VECTOR_INT_FLOAT_CONVERT(Op, 2, DestType) \
548  __SYCL_VECTOR_INT_FLOAT_CONVERT(Op, 3, DestType) \
549  __SYCL_VECTOR_INT_FLOAT_CONVERT(Op, 4, DestType) \
550  __SYCL_VECTOR_INT_FLOAT_CONVERT(Op, 8, DestType) \
551  __SYCL_VECTOR_INT_FLOAT_CONVERT(Op, 16, DestType)
552 
553 __SYCL_INT_FLOAT_CONVERT(SToF, half)
554 __SYCL_INT_FLOAT_CONVERT(SToF, float)
555 __SYCL_INT_FLOAT_CONVERT(SToF, double)
556 
557 __SYCL_INT_FLOAT_CONVERT(UToF, half)
558 __SYCL_INT_FLOAT_CONVERT(UToF, float)
559 __SYCL_INT_FLOAT_CONVERT(UToF, double)
560 
561 #undef __SYCL_SCALAR_INT_FLOAT_CONVERT
562 #undef __SYCL_VECTOR_INT_FLOAT_CONVERT
563 #undef __SYCL_INT_FLOAT_CONVERT
564 
565 // float to float conversions
566 #define __SYCL_SCALAR_FLOAT_FLOAT_CONVERT(DestType, RoundingMode, \
567  RoundingModeCondition) \
568  template <typename From, typename To, int VecSize, typename Enable, \
569  sycl::rounding_mode RM> \
570  enable_if_to_float_scalar_t<sycl::opencl::cl_##DestType, Enable, VecSize, \
571  To, RoundingModeCondition, RM> \
572  FConvert(From Value) { \
573  return __spirv_FConvert_R##DestType##_##RoundingMode(Value); \
574  }
575 
576 #define __SYCL_VECTOR_FLOAT_FLOAT_CONVERT(N, DestType, RoundingMode, \
577  RoundingModeCondition) \
578  template <typename From, typename To, int VecSize, typename Enable, \
579  sycl::rounding_mode RM> \
580  enable_if_to_float_vector_t<sycl::opencl::cl_##DestType, Enable, N, VecSize, \
581  To, RoundingModeCondition, RM> \
582  FConvert(From Value) { \
583  return __spirv_FConvert_R##DestType##N##_##RoundingMode(Value); \
584  }
585 
586 #define __SYCL_FLOAT_FLOAT_CONVERT(DestType, RoundingMode, \
587  RoundingModeCondition) \
588  __SYCL_SCALAR_FLOAT_FLOAT_CONVERT(DestType, RoundingMode, \
589  RoundingModeCondition) \
590  __SYCL_VECTOR_FLOAT_FLOAT_CONVERT(2, DestType, RoundingMode, \
591  RoundingModeCondition) \
592  __SYCL_VECTOR_FLOAT_FLOAT_CONVERT(3, DestType, RoundingMode, \
593  RoundingModeCondition) \
594  __SYCL_VECTOR_FLOAT_FLOAT_CONVERT(4, DestType, RoundingMode, \
595  RoundingModeCondition) \
596  __SYCL_VECTOR_FLOAT_FLOAT_CONVERT(8, DestType, RoundingMode, \
597  RoundingModeCondition) \
598  __SYCL_VECTOR_FLOAT_FLOAT_CONVERT(16, DestType, RoundingMode, \
599  RoundingModeCondition)
600 
601 #define __SYCL_FLOAT_FLOAT_CONVERT_FOR_TYPE(DestType) \
602  __SYCL_FLOAT_FLOAT_CONVERT(DestType, rte, RteOrAutomatic) \
603  __SYCL_FLOAT_FLOAT_CONVERT(DestType, rtz, Rtz) \
604  __SYCL_FLOAT_FLOAT_CONVERT(DestType, rtp, Rtp) \
605  __SYCL_FLOAT_FLOAT_CONVERT(DestType, rtn, Rtn)
606 
607 __SYCL_FLOAT_FLOAT_CONVERT_FOR_TYPE(half)
608 __SYCL_FLOAT_FLOAT_CONVERT_FOR_TYPE(float)
609 __SYCL_FLOAT_FLOAT_CONVERT_FOR_TYPE(double)
610 
611 #undef __SYCL_SCALAR_FLOAT_FLOAT_CONVERT
612 #undef __SYCL_VECTOR_FLOAT_FLOAT_CONVERT
613 #undef __SYCL_FLOAT_FLOAT_CONVERT
614 #undef __SYCL_FLOAT_FLOAT_CONVERT_FOR_TYPE
615 
616 template <typename NativeBFT, typename NativeFloatT, int VecSize>
617 inline NativeFloatT ConvertBF16ToFVec(NativeBFT vec) {
618  bfloat16 *src = sycl::bit_cast<bfloat16 *>(&vec);
619 
620  // OpenCL vector of 3 elements is aligned to 4 multiplied by
621  // the size of data type.
622  constexpr int AdjustedSize = (VecSize == 3) ? 4 : VecSize;
623  float dst[AdjustedSize];
624  sycl::ext::oneapi::detail::BF16VecToFloatVec<VecSize>(src, dst);
625 
626  return sycl::bit_cast<NativeFloatT>(dst);
627 }
628 
629 template <typename NativeFloatT, typename NativeBFT, int VecSize>
630 inline NativeBFT ConvertFToBF16Vec(NativeFloatT vec) {
631  float *src = sycl::bit_cast<float *>(&vec);
632 
633  // OpenCL vector of 3 elements is aligned to 4 multiplied by
634  // the size of data type.
635  constexpr int AdjustedSize = (VecSize == 3) ? 4 : VecSize;
636  bfloat16 dst[AdjustedSize];
637 
638  sycl::ext::oneapi::detail::FloatVecToBF16Vec<VecSize>(src, dst);
639  return sycl::bit_cast<NativeBFT>(dst);
640 }
641 
642 /* Emit _imf_* funcs only on Intel hardware. */
643 #if defined(__SPIR__) || defined(__SPIRV__)
644 #define EXPAND_BF16_ROUNDING_MODE(type, type_str, rmode, rmode_str) \
645  template <typename NativeToT, sycl::rounding_mode RoundingMode> \
646  std::enable_if_t<(std::is_same_v<NativeToT, type> && RoundingMode == rmode), \
647  NativeToT> \
648  ConvertFromBF16Scalar(uint16_t val) { \
649  return __imf_bfloat162##type_str##_##rmode_str(val); \
650  } \
651  template <typename NativeFromT, sycl::rounding_mode RoundingMode> \
652  std::enable_if_t< \
653  (std::is_same_v<NativeFromT, type> && RoundingMode == rmode), uint16_t> \
654  ConvertToBF16Scalar(NativeFromT val) { \
655  return __imf_##type_str##2bfloat16_##rmode_str(val); \
656  }
657 
658 #else // __SYCL_DEVICE_ONLY__ && (defined(__SPIR__) || defined(__SPIRV__))
659 // On non-Intel HWs, convert BF16 to float (losslessly) and convert float
660 // to the desired type.
661 #define EXPAND_BF16_ROUNDING_MODE(type, type_str, rmode, rmode_str) \
662  template <typename NativeToT, sycl::rounding_mode RoundingMode> \
663  std::enable_if_t<(std::is_same_v<NativeToT, type> && RoundingMode == rmode), \
664  NativeToT> \
665  ConvertFromBF16Scalar(uint16_t val) { \
666  bfloat16 bfval = sycl::bit_cast<bfloat16>(val); \
667  float fval = static_cast<float>(bfval); \
668  return convertImpl<fval, NativeToT, RoundingMode, 1, float, NativeToT>( \
669  fval); \
670  } \
671  template <typename NativeFromT, sycl::rounding_mode RoundingMode> \
672  std::enable_if_t< \
673  (std::is_same_v<NativeFromT, type> && RoundingMode == rmode), uint16_t> \
674  ConvertToBF16Scalar(NativeFromT val) { \
675  constexpr int rm = static_cast<int>(RoundingMode); \
676  bfloat16 bfval = sycl::ext::oneapi::detail::ConvertToBfloat16:: \
677  getBfloat16WithRoundingMode<NativeFromT, rm>(val); \
678  return sycl::bit_cast<uint16_t>(bfval); \
679  }
680 #endif // __SYCL_DEVICE_ONLY__ && (defined(__SPIR__) || defined(__SPIRV__))
681 
682 #define EXPAND_BF16_TYPE(type, type_str) \
683  EXPAND_BF16_ROUNDING_MODE(type, type_str, sycl::rounding_mode::automatic, \
684  rn) \
685  EXPAND_BF16_ROUNDING_MODE(type, type_str, sycl::rounding_mode::rte, rn) \
686  EXPAND_BF16_ROUNDING_MODE(type, type_str, sycl::rounding_mode::rtp, ru) \
687  EXPAND_BF16_ROUNDING_MODE(type, type_str, sycl::rounding_mode::rtn, rd) \
688  EXPAND_BF16_ROUNDING_MODE(type, type_str, sycl::rounding_mode::rtz, rz)
689 
690 EXPAND_BF16_TYPE(uint, uint)
691 EXPAND_BF16_TYPE(int, int)
692 EXPAND_BF16_TYPE(ushort, ushort)
693 EXPAND_BF16_TYPE(short, short)
694 EXPAND_BF16_TYPE(long, ll)
695 EXPAND_BF16_TYPE(unsigned long long, ull)
696 
697 #undef EXPAND_BF16_TYPE
698 #undef EXPAND_BF16_ROUNDING_MODE
699 
700 // Mapping from BF16 to float is 1:1, lossless, so we accept all
701 // rounding modes.
702 template <typename NativeToT, sycl::rounding_mode RoundingMode>
703 std::enable_if_t<std::is_same_v<NativeToT, float>, NativeToT>
704 ConvertFromBF16Scalar(uint16_t val) {
705  bfloat16 bfval = sycl::bit_cast<bfloat16>(val);
706  return static_cast<float>(bfval);
707 }
708 
709 template <typename NativeFromT, sycl::rounding_mode RoundingMode>
710 std::enable_if_t<std::is_same_v<NativeFromT, double>, uint16_t>
711 ConvertToBF16Scalar(NativeFromT val) {
712 #if defined(__SPIR__) || defined(__SPIRV__)
713  return __imf_double2bfloat16(val);
714 #else
715  constexpr int rm = static_cast<int>(RoundingMode);
716  bfloat16 bfval =
717  sycl::ext::oneapi::detail::ConvertToBfloat16::getBfloat16WithRoundingMode<
718  NativeFromT, rm>(val);
719  return sycl::bit_cast<uint16_t>(bfval);
720 #endif
721 }
722 
723 template <typename NativeFromT, sycl::rounding_mode RoundingMode>
724 std::enable_if_t<std::is_same_v<NativeFromT, float>, uint16_t>
725 ConvertToBF16Scalar(NativeFromT val) {
726 
727 #if defined(__SPIR__) || defined(__SPIRV__)
728  if constexpr (RoundingMode == sycl::rounding_mode::automatic ||
729  RoundingMode == sycl::rounding_mode::rte)
730  return __imf_float2bfloat16_rn(val);
731  else if constexpr (RoundingMode == sycl::rounding_mode::rtp)
732  return __imf_float2bfloat16_ru(val);
733  else if constexpr (RoundingMode == sycl::rounding_mode::rtn)
734  return __imf_float2bfloat16_rd(val);
735  else if constexpr (RoundingMode == sycl::rounding_mode::rtz)
736  return __imf_float2bfloat16_rz(val);
737  else
738  static_assert(false, "Invalid rounding mode.");
739 #else
740  constexpr int rm = static_cast<int>(RoundingMode);
741  bfloat16 bfval =
742  sycl::ext::oneapi::detail::ConvertToBfloat16::getBfloat16WithRoundingMode<
743  float, rm>(val);
744  return sycl::bit_cast<uint16_t>(bfval);
745 #endif
746 }
747 
748 #endif // __SYCL_DEVICE_ONLY__
749 
750 // Wrapper function for scalar and vector conversions from BF16 type.
751 template <typename ToT, typename NativeFromT, typename NativeToT,
752  sycl::rounding_mode RoundingMode, int VecSize>
753 NativeToT ConvertFromBF16(NativeFromT val) {
754 #ifdef __SYCL_DEVICE_ONLY__
755  // Use vector conversion from BF16 to float for all rounding modes.
756  if constexpr (std::is_same_v<ToT, float> && VecSize > 1)
757  return ConvertBF16ToFVec<NativeFromT, NativeToT, VecSize>(val);
758  else
759 #endif
760  // For VecSize > 1. Only for device.
761  if constexpr (VecSize > 1) {
762  NativeToT retval;
763  for (int i = 0; i < VecSize; i++) {
764  retval[i] = ConvertFromBF16Scalar<ToT, RoundingMode>(val[i]);
765  }
766  return retval;
767  }
768  // For VecSize == 1.
769  else
770  return ConvertFromBF16Scalar<NativeToT, RoundingMode>(val);
771 }
772 
773 // Wrapper function for scalar and vector conversions to BF16 type.
774 template <typename FromT, typename NativeFromT, typename NativeToT,
775  sycl::rounding_mode RoundingMode, int VecSize>
776 NativeToT ConvertToBF16(NativeFromT val) {
777 #ifdef __SYCL_DEVICE_ONLY__
778  // Use vector conversion to BF16 from float for RNE rounding mode.
779  if constexpr (std::is_same_v<FromT, float> && VecSize > 1 &&
780  (RoundingMode == sycl::rounding_mode::automatic ||
781  RoundingMode == sycl::rounding_mode::rte))
782  return ConvertFToBF16Vec<NativeFromT, NativeToT, VecSize>(val);
783  else
784 #endif
785  // For VecSize > 1. Only for device.
786  if constexpr (VecSize > 1) {
787  NativeToT retval;
788  for (int i = 0; i < VecSize; i++) {
789  retval[i] = ConvertToBF16Scalar<FromT, RoundingMode>(val[i]);
790  }
791  return retval;
792  }
793  // For VecSize == 1.
794  else
795  return ConvertToBF16Scalar<NativeFromT, RoundingMode>(val);
796 }
797 
818 template <typename FromT, typename ToT, sycl::rounding_mode RoundingMode,
819  int VecSize, typename NativeFromT, typename NativeToT>
820 NativeToT convertImpl(NativeFromT Value) {
821  static_assert(!std::is_same_v<FromT, ToT>);
822  static_assert(!std::is_same_v<NativeFromT, NativeToT>);
823  using ElemTy = typename detail::ConvertToOpenCLType_t<ToT>;
824  if constexpr (is_sint_to_sint<FromT, ToT>::value)
825  return SConvert<NativeFromT, NativeToT, VecSize, ElemTy>(Value);
826  else if constexpr (is_uint_to_uint<FromT, ToT>::value)
827  return UConvert<NativeFromT, NativeToT, VecSize, ElemTy>(Value);
828  else if constexpr (is_sint_to_float<FromT, ToT>::value)
829  return ConvertSToF<NativeFromT, NativeToT, VecSize, ElemTy>(Value);
830  else if constexpr (is_uint_to_float<FromT, ToT>::value)
831  return ConvertUToF<NativeFromT, NativeToT, VecSize, ElemTy>(Value);
832  else if constexpr (is_float_to_float<FromT, ToT>::value)
833  return FConvert<NativeFromT, NativeToT, VecSize, ElemTy, RoundingMode>(
834  Value);
835  // BF16 conversion to other types.
836  else if constexpr (std::is_same_v<FromT, bfloat16>)
837  return ConvertFromBF16<ToT, NativeFromT, NativeToT, RoundingMode, VecSize>(
838  Value);
839  // conversion from other types to BF16.
840  else if constexpr (std::is_same_v<ToT, bfloat16>)
841  return ConvertToBF16<FromT, NativeFromT, NativeToT, RoundingMode, VecSize>(
842  Value);
843  else if constexpr (is_float_to_sint<FromT, ToT>::value)
844  return ConvertFToS<NativeFromT, NativeToT, VecSize, ElemTy, RoundingMode>(
845  Value);
846  else if constexpr (is_float_to_uint<FromT, ToT>::value)
847  return ConvertFToU<NativeFromT, NativeToT, VecSize, ElemTy, RoundingMode>(
848  Value);
849  else {
850  static_assert(is_sint_to_from_uint<FromT, ToT>::value,
851  "Unexpected conversion type");
852  static_assert(VecSize == 1, "Conversion between signed and unsigned data "
853  "types is only available for scalars");
854  // vec::convert is underspecified and therefore it is not entirely clear
855  // what to do here. 'static_cast' implementation matches SYCL CTS and it
856  // matches our old implementation. Unfortunately, OpSetConvertUToS and
857  // OpSatConvertSToU behave differently and we can't use them here until the
858  // behavior of conversions is well-defined by the SYCL 2020 specificiation.
859  // See https://github.com/KhronosGroup/SYCL-Docs/issues/492
860  return static_cast<NativeToT>(Value);
861  }
862 }
863 
864 #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
865 template <typename FromT, typename ToT, sycl::rounding_mode RoundingMode,
866  int VecSize, typename NativeFromT, typename NativeToT>
867 auto ConvertImpl(std::byte val) {
868  return convertImpl<FromT, ToT, RoundingMode, VecSize, NativeFromT, NativeToT>(
869  (std::int8_t)val);
870 }
871 #endif
872 
873 } // namespace detail
874 } // namespace _V1
875 } // namespace sycl
sycl::exception
Definition: exception.hpp:75
sycl::_V1::ext::oneapi::bfloat16
Definition: bfloat16.hpp:83
sycl::_V1::ext::oneapi::detail::ConvertToBfloat16::getBfloat16WithRoundingMode
static bfloat16 getBfloat16WithRoundingMode(const Ty &a)
Definition: bfloat16.hpp:579
__DPCPP_SYCL_EXTERNAL
#define __DPCPP_SYCL_EXTERNAL
Definition: defines_elementary.hpp:33
generic_type_traits.hpp
sycl::_V1::ext::intel::esimd::uchar
unsigned char uchar
Definition: common.hpp:41
sycl::_V1::ext::intel::esimd::ushort
unsigned short ushort
Definition: common.hpp:42
sycl::_V1::ext::intel::esimd::uint
unsigned int uint
Definition: common.hpp:43
__imf_short_as_bfloat16
uint16_t __imf_short_as_bfloat16(short)
__imf_bfloat162ull_rz
unsigned long long __imf_bfloat162ull_rz(uint16_t)
__imf_bfloat162ushort_rz
unsigned short __imf_bfloat162ushort_rz(uint16_t)
__imf_float2bfloat16_rd
uint16_t __imf_float2bfloat16_rd(float)
__imf_ull2bfloat16_rz
uint16_t __imf_ull2bfloat16_rz(unsigned long long)
__imf_ushort2bfloat16_rz
uint16_t __imf_ushort2bfloat16_rz(unsigned short)
__imf_bfloat162short_rn
short __imf_bfloat162short_rn(uint16_t)
__imf_int2bfloat16_ru
uint16_t __imf_int2bfloat16_ru(int)
__imf_ushort2bfloat16_rn
uint16_t __imf_ushort2bfloat16_rn(unsigned short)
__imf_ll2bfloat16_rd
uint16_t __imf_ll2bfloat16_rd(long long)
__imf_bfloat162uint_ru
unsigned int __imf_bfloat162uint_ru(uint16_t)
__imf_float2bfloat16
uint16_t __imf_float2bfloat16(float)
__imf_short2bfloat16_rd
uint16_t __imf_short2bfloat16_rd(short)
__imf_bfloat162ull_rn
unsigned long long __imf_bfloat162ull_rn(uint16_t)
__imf_bfloat162ll_rz
long long __imf_bfloat162ll_rz(uint16_t)
__imf_int2bfloat16_rd
uint16_t __imf_int2bfloat16_rd(int)
__imf_uint2bfloat16_rd
uint16_t __imf_uint2bfloat16_rd(unsigned int)
__imf_float2bfloat16_ru
uint16_t __imf_float2bfloat16_ru(float)
__imf_bfloat162int_ru
int __imf_bfloat162int_ru(uint16_t)
__imf_bfloat162uint_rn
unsigned int __imf_bfloat162uint_rn(uint16_t)
__imf_bfloat162ll_rn
long long __imf_bfloat162ll_rn(uint16_t)
__imf_ushort2bfloat16_ru
uint16_t __imf_ushort2bfloat16_ru(unsigned short)
__imf_bfloat162int_rz
int __imf_bfloat162int_rz(uint16_t)
__imf_bfloat162ull_ru
unsigned long long __imf_bfloat162ull_ru(uint16_t)
__imf_bfloat162ushort_rn
unsigned short __imf_bfloat162ushort_rn(uint16_t)
__imf_float2bfloat16_rn
uint16_t __imf_float2bfloat16_rn(float)
__imf_ll2bfloat16_rn
uint16_t __imf_ll2bfloat16_rn(long long)
__imf_ll2bfloat16_rz
uint16_t __imf_ll2bfloat16_rz(long long)
__imf_bfloat162int_rd
int __imf_bfloat162int_rd(uint16_t)
__imf_bfloat162ushort_ru
unsigned short __imf_bfloat162ushort_ru(uint16_t)
__imf_bfloat162uint_rd
unsigned int __imf_bfloat162uint_rd(uint16_t)
__imf_int2bfloat16_rn
uint16_t __imf_int2bfloat16_rn(int)
__imf_float2bfloat16_rz
uint16_t __imf_float2bfloat16_rz(float)
__imf_ull2bfloat16_rd
uint16_t __imf_ull2bfloat16_rd(unsigned long long)
__imf_bfloat162short_ru
short __imf_bfloat162short_ru(uint16_t)
__imf_short2bfloat16_ru
uint16_t __imf_short2bfloat16_ru(short)
__imf_int2bfloat16_rz
uint16_t __imf_int2bfloat16_rz(int)
__imf_ull2bfloat16_ru
uint16_t __imf_ull2bfloat16_ru(unsigned long long)
__imf_bfloat162float
float __imf_bfloat162float(uint16_t)
__imf_bfloat162int_rn
int __imf_bfloat162int_rn(uint16_t)
__imf_ull2bfloat16_rn
uint16_t __imf_ull2bfloat16_rn(unsigned long long)
__imf_double2bfloat16
uint16_t __imf_double2bfloat16(double)
__imf_ushort_as_bfloat16
uint16_t __imf_ushort_as_bfloat16(unsigned short)
__imf_uint2bfloat16_ru
uint16_t __imf_uint2bfloat16_ru(unsigned int)
__imf_short2bfloat16_rn
uint16_t __imf_short2bfloat16_rn(short)
__imf_bfloat162ll_ru
long long __imf_bfloat162ll_ru(uint16_t)
__imf_ll2bfloat16_ru
uint16_t __imf_ll2bfloat16_ru(long long)
__imf_bfloat162short_rd
short __imf_bfloat162short_rd(uint16_t)
__imf_bfloat162uint_rz
unsigned int __imf_bfloat162uint_rz(uint16_t)
__imf_bfloat16_as_short
short __imf_bfloat16_as_short(uint16_t)
__imf_bfloat16_as_ushort
unsigned short __imf_bfloat16_as_ushort(uint16_t)
__imf_bfloat162ll_rd
long long __imf_bfloat162ll_rd(uint16_t)
__imf_uint2bfloat16_rn
uint16_t __imf_uint2bfloat16_rn(unsigned int)
__imf_bfloat162ushort_rd
unsigned short __imf_bfloat162ushort_rd(uint16_t)
__imf_uint2bfloat16_rz
uint16_t __imf_uint2bfloat16_rz(unsigned int)
__imf_ushort2bfloat16_rd
uint16_t __imf_ushort2bfloat16_rd(unsigned short)
__imf_short2bfloat16_rz
uint16_t __imf_short2bfloat16_rz(short)
__imf_bfloat162short_rz
short __imf_bfloat162short_rz(uint16_t)
__imf_bfloat162ull_rd
unsigned long long __imf_bfloat162ull_rd(uint16_t)
sycl::_V1::detail::ConvertUToF
To ConvertUToF(From Value)
Definition: vector_convert.hpp:241
sycl::_V1::detail::is_sint_to_from_uint
std::bool_constant<(detail::is_sigeninteger_v< T > &&detail::is_sugeninteger_v< R >)||(detail::is_sugeninteger_v< T > &&detail::is_sigeninteger_v< R >)> is_sint_to_from_uint
Definition: vector_convert.hpp:187
sycl::_V1::detail::is_sint_to_sint
std::bool_constant< is_sigeninteger_v< T > &&is_sigeninteger_v< R > > is_sint_to_sint
Definition: vector_convert.hpp:178
sycl::_V1::detail::is_float_to_sint
std::bool_constant< detail::is_floating_point< T >::value &&std::is_integral_v< R > &&!std::is_unsigned_v< R > > is_float_to_sint
Definition: vector_convert.hpp:211
sycl::_V1::detail::is_sint_to_float
std::bool_constant< std::is_integral_v< T > &&!std::is_unsigned_v< T > &&detail::is_floating_point< R >::value > is_sint_to_float
Definition: vector_convert.hpp:192
sycl::_V1::detail::ConvertFromBF16Scalar
NativeToT ConvertFromBF16Scalar(bfloat16 val)
Definition: vector_convert.hpp:295
sycl::_V1::detail::convertImpl
NativeToT convertImpl(NativeFromT)
Entry point helper for all kinds of converts between scalars and vectors, it dispatches to a right fu...
Definition: vector_convert.hpp:820
sycl::_V1::detail::is_uint_to_float
std::bool_constant< std::is_unsigned_v< T > &&detail::is_floating_point< R >::value > is_uint_to_float
Definition: vector_convert.hpp:197
sycl::_V1::detail::UConvert
To UConvert(From Value)
Definition: vector_convert.hpp:229
sycl::_V1::detail::is_uint_to_uint
std::bool_constant< is_sugeninteger_v< T > &&is_sugeninteger_v< R > > is_uint_to_uint
Definition: vector_convert.hpp:182
sycl::_V1::detail::ConvertSToF
To ConvertSToF(From Value)
Definition: vector_convert.hpp:235
sycl::_V1::detail::ConvertImpl
auto ConvertImpl(std::byte val)
Definition: vector_convert.hpp:867
sycl::_V1::detail::ConvertToBF16
NativeToT ConvertToBF16(NativeFromT val)
Definition: vector_convert.hpp:776
sycl::_V1::detail::ConvertFToU
To ConvertFToU(From Value)
Definition: vector_convert.hpp:290
sycl::_V1::detail::ConvertToBF16Scalar
bfloat16 ConvertToBF16Scalar(NativeFromT val)
Definition: vector_convert.hpp:308
sycl::_V1::detail::is_float_to_uint
std::bool_constant< detail::is_floating_point< T >::value &&std::is_unsigned_v< R > > is_float_to_uint
Definition: vector_convert.hpp:206
sycl::_V1::detail::ConvertFToS
To ConvertFToS(From Value)
Definition: vector_convert.hpp:255
sycl::_V1::detail::ConvertFromBF16
NativeToT ConvertFromBF16(NativeFromT val)
Definition: vector_convert.hpp:753
sycl::_V1::detail::bfloat16
sycl::ext::oneapi::bfloat16 bfloat16
Definition: vector_convert.hpp:218
sycl::_V1::detail::ConvertToOpenCLType_t
decltype(convertToOpenCLType(std::declval< T >())) ConvertToOpenCLType_t
Definition: generic_type_traits.hpp:400
sycl::_V1::detail::is_int_to_float
std::bool_constant< std::is_integral_v< T > &&detail::is_floating_point< R >::value > is_int_to_float
Definition: vector_convert.hpp:201
sycl::_V1::detail::SConvert
To SConvert(From Value)
Definition: vector_convert.hpp:223
sycl::_V1::detail::FConvert
To FConvert(From Value)
Definition: vector_convert.hpp:248
sycl::_V1::detail::is_float_to_float
std::bool_constant< detail::is_floating_point< T >::value &&detail::is_floating_point< R >::value > is_float_to_float
Definition: vector_convert.hpp:216
sycl::_V1::ext::intel::experimental::byte
unsigned char byte
Definition: online_compiler.hpp:22
sycl::_V1::ceil
float ceil(float)
sycl::_V1::Mode
Mode
Definition: multi_ptr.hpp:504
sycl::_V1::errc::runtime
@ runtime
sycl::_V1::vec
class __SYCL_EBO vec
Definition: aliases.hpp:18
sycl::_V1::floor
float floor(float)
sycl::_V1::rint
float rint(float)
sycl::_V1::make_error_code
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
Definition: exception.cpp:65
sycl::_V1::half
sycl::detail::half_impl::half half
Definition: aliases.hpp:101
sycl::_V1::x
autodecltype(x) x
Definition: common_functions.cpp:33
sycl::_V1::trunc
float trunc(float)
sycl
Definition: access.hpp:18
sycl::_V1::detail::is_floating_point
Definition: type_traits.hpp:334