39 #include <type_traits>
43 inline namespace _V1 {
86 template <
class F,
class T =
void>
88 detail::check_type_in_v<F, float, double, half, ext::oneapi::bfloat16>, T>;
91 sycl::access::target::device>;
100 sycl::access::target::device>;
114 return ((
static_cast<unsigned>(
static_cast<uint8_t
>(GlobalFlushBuf[WIOffset]))
116 static_cast<uint8_t
>(GlobalFlushBuf[WIOffset + 1]));
125 unsigned WIOffset,
unsigned Offset) {
126 GlobalFlushBuf[WIOffset] =
static_cast<char>((Offset >> 8) & 0xff);
127 GlobalFlushBuf[WIOffset + 1] =
static_cast<char>(Offset & 0xff);
131 unsigned WIOffset,
const char *Str,
unsigned Len,
132 unsigned Padding = 0) {
136 if ((Offset + Len + Padding > FlushBufferSize) ||
137 (WIOffset + Offset + Len + Padding > GlobalFlushBuf.size()))
142 for (
size_t I = 0; I < Padding; ++I, ++Offset)
143 GlobalFlushBuf[WIOffset + Offset] =
' ';
145 for (
size_t I = 0; I < Len; ++I, ++Offset) {
146 GlobalFlushBuf[WIOffset + Offset] = Str[I];
164 template <
typename T>
165 inline std::make_unsigned_t<T>
getAbsVal(
const T Val,
const int Base) {
166 return ((Base == 10) && (Val < 0)) ? -Val : Val;
173 return 'a' + Digit - 10;
177 template <
typename T>
178 inline typename std::enable_if_t<std::is_integral_v<T>,
unsigned>
180 unsigned NumDigits = 0;
191 template <
typename T>
192 inline typename std::enable_if_t<std::is_integral_v<T>,
unsigned>
193 ScalarToStr(
const T &Val,
char *Buf,
unsigned Flags,
int,
int Precision = -1) {
218 if ((Flags &
ShowPos) && Val >= 0)
225 if (Base != 10 && (Flags &
ShowBase)) {
233 const unsigned NumBuf =
integralToBase(AbsVal, Base, Buf + Offset);
236 return Offset + NumBuf;
239 inline unsigned append(
char *Dst,
const char *Src) {
241 for (; Src[Len] !=
'\0'; ++Len)
244 for (
unsigned I = 0; I < Len; ++I)
255 return Internal.I32Val;
258 inline unsigned long long D2I64(
double Val) {
261 unsigned long long I64Val;
264 return Internal.I64Val;
267 template <
typename T>
268 inline typename detail::enable_if_t<
269 std::is_same<T, float>::value || std::is_same<T, double>::value,
bool>
271 if constexpr (
sizeof(Val) == 4) {
272 return (
F2I32(Val) & 0x7fffffff) == 0x7f800000;
273 }
else if constexpr (
sizeof(Val) == 8) {
274 return (
D2I64(Val) & -1ULL >> 1) == 0x7ffULL << 52;
280 template <
typename T>
281 inline typename detail::enable_if_t<
282 std::is_same<T, float>::value || std::is_same<T, double>::value,
bool>
284 if constexpr (
sizeof(Val) == 4) {
285 return (
F2I32(Val) & 0x7fffffff) > 0x7f800000;
286 }
else if constexpr (
sizeof(Val) == 8) {
287 return (
D2I64(Val) & -1ULL >> 1) > 0x7ffULL << 52;
293 template <
typename T>
294 inline typename detail::enable_if_t<
295 std::is_same<T, float>::value || std::is_same<T, double>::value,
bool>
297 if constexpr (
sizeof(Val) == 4) {
298 return F2I32(Val) >> 31;
299 }
else if constexpr (
sizeof(Val) == 8) {
300 return D2I64(Val) >> 63;
306 template <
typename T>
307 typename detail::enable_if_t<
308 std::is_same<T, float>::value || std::is_same<T, double>::value,
unsigned>
315 return append(Buf,
"nan");
323 return append(Buf,
"-inf");
324 return append(Buf,
"inf");
329 template <
typename T>
330 inline typename std::enable_if_t<std::is_same_v<T, half>,
unsigned>
333 return append(Buf,
"nan");
336 const uint16_t Sign = sycl::bit_cast<uint16_t>(Val) & 0x8000;
338 const uint16_t Exp16 = (sycl::bit_cast<uint16_t>(Val) & 0x7c00) >> 10;
342 return append(Buf,
"-inf");
343 return append(Buf,
"inf");
348 template <
typename T>
349 inline typename std::enable_if_t<std::is_same_v<T, ext::oneapi::bfloat16>,
353 return append(Buf,
"nan");
356 const uint16_t Sign = sycl::bit_cast<uint16_t>(Val) & 0x8000;
358 const uint16_t Exp16 = (sycl::bit_cast<uint16_t>(Val) & 0x7f80) >> 7;
362 return append(Buf,
"-inf");
363 return append(Buf,
"inf");
368 template <
typename T>
370 int Precision,
bool IsSci) {
374 while (AbsVal >= T{10.0}) {
379 while (AbsVal > T{0.0} && AbsVal < T{1.0}) {
384 auto IntegralPart =
static_cast<int>(AbsVal);
385 auto FractionPart = AbsVal - IntegralPart;
390 int P = Precision > 0 ? Precision : 4;
391 size_t FractionLength = Exp + P;
400 for (
unsigned I = 0; I < FractionLength; ++I) {
401 FractionPart *= T{10.0};
402 FractionDigits[I] =
static_cast<int>(FractionPart);
403 FractionPart -=
static_cast<int>(FractionPart);
406 int Carry = FractionPart >
static_cast<T
>(0.5) ? 1 : 0;
409 for (
int I = FractionLength - 1; I >= 0 && Carry; --I) {
410 auto Digit = FractionDigits[I] + Carry;
411 FractionDigits[I] = Digit % 10;
416 IntegralPart += Carry;
417 if (IntegralPart == 10) {
428 Digits[Offset++] =
'.';
431 for (
unsigned I = 0; I < FractionLength; ++I)
434 auto AbsExp = Exp < 0 ? -Exp : Exp;
436 Digits[Offset++] =
'e';
437 Digits[Offset++] = Exp >= 0 ?
'+' :
'-';
442 Digits[Offset++] =
'0';
443 Digits[Offset++] =
'.';
445 Digits[Offset++] =
'0';
451 for (
unsigned I = 0; I < FractionLength; ++I)
458 for (; I < FractionLength && Exp--; ++I)
462 Digits[Offset++] =
'.';
465 for (; I < FractionLength; ++I)
470 while (Digits[Offset - 1] ==
'0')
474 if (Digits[Offset - 1] ==
'.')
481 template <
typename T>
482 inline EnableIfFP<T, unsigned>
483 ScalarToStr(
const T &Val,
char *Buf,
unsigned Flags,
int,
int Precision = -1) {
489 auto AbsVal = Val < 0 ? Neg : Val;
507 template <
typename T>
508 inline typename std::enable_if_t<std::is_integral_v<T>>
510 unsigned WIOffset,
unsigned Flags,
int Width,
const T &Val) {
512 unsigned Len =
ScalarToStr(Val, Digits, Flags, Width);
513 write(GlobalFlushBuf, FlushBufferSize, WIOffset, Digits, Len,
514 (Width > 0 &&
static_cast<unsigned>(Width) > Len)
515 ?
static_cast<unsigned>(Width) - Len
519 template <
typename T>
522 unsigned WIOffset,
unsigned Flags,
int Width,
int Precision,
525 unsigned Len =
ScalarToStr(Val, Digits, Flags, Width, Precision);
526 write(GlobalFlushBuf, FlushBufferSize, WIOffset, Digits, Len,
527 (Width > 0 &&
static_cast<unsigned>(Width) > Len)
528 ?
static_cast<unsigned>(Width) - Len
539 Cur = GlobalOffset[0].load();
541 if (GlobalBuf.get_range().size() - Cur < Size)
545 }
while (!GlobalOffset[0].compare_exchange_strong(Cur, New));
557 if (!
updateOffset(GlobalOffset, GlobalBuf, Offset, Cur))
561 for (
unsigned I = StmtOffset; I < StmtOffset + Offset; I++) {
562 GlobalBuf[Cur++] = GlobalFlushBuf[I];
568 template <
typename T,
int VecLength>
569 typename std::enable_if_t<(VecLength == 1),
unsigned>
572 return ScalarToStr(
static_cast<T
>(Vec.x()), VecStr, Flags, Width, Precision);
575 template <
typename T,
int VecLength>
576 typename std::enable_if_t<(VecLength == 2 || VecLength == 4 || VecLength == 8 ||
582 VecToStr<T, VecLength / 2>(Vec.lo(), VecStr, Flags, Width, Precision);
584 Len +=
VecToStr<T, VecLength / 2>(Vec.hi(), VecStr + Len, Flags, Width,
589 template <
typename T,
int VecLength>
590 typename std::enable_if_t<(VecLength == 3),
unsigned>
593 unsigned Len = VecToStr<T, 2>(Vec.lo(), VecStr, Flags, Width, Precision);
595 Len += VecToStr<T, 1>(Vec.z(), VecStr + Len, Flags, Width, Precision);
599 template <
typename T,
int VecLength>
601 unsigned WIOffset,
unsigned Flags,
int Width,
604 constexpr
size_t MAX_VEC_SIZE =
606 char VecStr[MAX_VEC_SIZE] = {0};
607 unsigned Len = VecToStr<T, VecLength>(Vec, VecStr, Flags, Width, Precision);
608 write(GlobalFlushBuf, FlushBufferSize, WIOffset, VecStr, Len,
609 (Width > 0 && Width > Len) ? Width - Len : 0);
612 template <
int ArrayLength>
617 for (
int I = 0; I < ArrayLength; ++I) {
619 if (I != ArrayLength - 1)
628 template <
int ArrayLength>
630 size_t FlushBufferSize,
unsigned WIOffset,
634 write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len);
637 template <
int Dimensions>
639 size_t FlushBufferSize,
unsigned WIOffset,
645 Len +=
append(Buf,
"item(");
646 Len +=
append(Buf + Len,
"range: ");
648 Len +=
append(Buf + Len,
", id: ");
650 Len +=
append(Buf + Len,
", offset: ");
653 write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len);
656 template <
int Dimensions>
658 size_t FlushBufferSize,
unsigned WIOffset,
664 Len +=
append(Buf,
"nd_range(");
665 Len +=
append(Buf + Len,
"global_range: ");
667 Len +=
append(Buf + Len,
", local_range: ");
669 Len +=
append(Buf + Len,
", offset: ");
672 write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len);
675 template <
int Dimensions>
677 size_t FlushBufferSize,
unsigned WIOffset,
683 Len +=
append(Buf,
"nd_item(");
684 Len +=
append(Buf + Len,
"global_id: ");
686 Len +=
append(Buf + Len,
", local_id: ");
689 write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len);
692 template <
int Dimensions>
694 size_t FlushBufferSize,
unsigned WIOffset,
700 Len +=
append(Buf,
"group(");
701 Len +=
append(Buf + Len,
"id: ");
703 Len +=
append(Buf + Len,
", global_range: ");
704 Len +=
ArrayToStr(Buf + Len, Group.get_global_range());
705 Len +=
append(Buf + Len,
", local_range: ");
706 Len +=
ArrayToStr(Buf + Len, Group.get_local_range());
707 Len +=
append(Buf + Len,
", group_range: ");
708 Len +=
ArrayToStr(Buf + Len, Group.get_group_range());
710 write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len);
717 template <
int Dimensions>
720 Len +=
append(Buf,
"item(");
721 for (
int I = 0; I < 2; ++I) {
722 Len +=
append(Buf + Len, I == 0 ?
"range: " :
", id: ");
729 template <
int Dimensions>
731 size_t FlushBufferSize,
unsigned WIOffset,
737 Len +=
append(Buf,
"h_item(");
738 for (
int I = 0; I < 3; ++I) {
739 Len +=
append(Buf + Len, I == 0 ?
"\n global "
740 : I == 1 ?
"\n logical local "
741 :
"\n physical local ");
746 Len +=
append(Buf + Len,
"\n)");
747 write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len);
752 template <
typename VecT,
typename OperationLeftT,
typename OperationRightT,
753 template <
typename>
class OperationCurrentT,
int... Indexes>
755 VecT, OperationLeftT, OperationRightT, OperationCurrentT, Indexes...>>
757 using T =
typename VecT::element_type;
761 template <
typename T>
763 typename std::enable_if_t<IsSwizzleOp<T>::value,
820 friend const stream &
operator<<(
const stream &,
830 int width()
const {
return Width_; }
832 friend const stream &
operator<<(
const stream &,
849 :
public detail::OwnerLessBase<stream> {
851 #ifndef __SYCL_DEVICE_ONLY__
853 stream(std::shared_ptr<detail::stream_impl> Impl,
857 : impl{Impl}, GlobalBuf{GlobalBuf}, GlobalOffset{GlobalOffset},
858 GlobalFlushBuf{GlobalFlushBuf} {}
862 #ifdef __SYCL_DEVICE_ONLY__
868 stream(
size_t BufferSize,
size_t MaxStatementSize, handler &CGH);
872 stream(
size_t BufferSize,
size_t MaxStatementSize, handler &CGH,
873 const property_list &PropList);
875 #ifdef __SYCL_DEVICE_ONLY__
879 size_t size() const
noexcept {
return GlobalBuf.size(); }
881 size_t get_work_item_buffer_size()
const {
886 "get_size() is deprecated since SYCL 2020. Please use size() instead.")
887 size_t get_size()
const {
return size(); }
890 "2020. Please use get_work_item_buffer_size() instead.")
891 size_t get_max_statement_size()
const {
return get_work_item_buffer_size(); }
895 size_t get_work_item_buffer_size() const;
898 "get_size() is deprecated since SYCL 2020. Please
use size()
instead.")
899 size_t get_size() const;
902 "2020. Please
use get_work_item_buffer_size()
instead.")
903 size_t get_max_statement_size() const;
906 size_t get_precision()
const {
return Precision; }
908 size_t get_width()
const {
return Width; }
918 template <typename propertyT> propertyT
get_property() const;
921 #ifdef __SYCL_DEVICE_ONLY__
922 char padding[
sizeof(std::shared_ptr<detail::stream_impl>)];
924 std::shared_ptr<detail::stream_impl> impl;
947 mutable unsigned WIOffset = 0;
949 mutable size_t FlushBufferSize;
957 mutable int Precision = -1;
958 mutable int Width = -1;
961 void set_flag(
FmtFlags FormatFlag)
const { Flags |= FormatFlag; }
963 void unset_flag(
FmtFlags FormatFlag)
const { Flags &= ~FormatFlag; }
965 FmtFlags get_flags()
const {
return Flags; }
972 Flags |= FormatFlag & Mask;
1017 #ifdef __SYCL_DEVICE_ONLY__
1019 range<detail::GlobalBufDim> GlobalBufAccRange,
1020 range<detail::GlobalBufDim> GlobalBufMemRange,
1021 id<detail::GlobalBufDim> GlobalBufId,
1023 range<detail::GlobalOffsetDim> GlobalOffsetAccRange,
1024 range<detail::GlobalOffsetDim> GlobalOffsetMemRange,
1025 id<detail::GlobalOffsetDim> GlobalOffsetId,
1027 range<detail::GlobalBufDim> GlobalFlushAccRange,
1028 range<detail::GlobalBufDim> GlobalFlushMemRange,
1029 id<detail::GlobalBufDim> GlobalFlushId,
size_t _FlushBufferSize) {
1030 GlobalBuf.__init(GlobalBufPtr, GlobalBufAccRange, GlobalBufMemRange,
1032 GlobalOffset.__init(GlobalOffsetPtr, GlobalOffsetAccRange,
1033 GlobalOffsetMemRange, GlobalOffsetId);
1034 GlobalFlushBuf.__init(GlobalFlushPtr, GlobalFlushAccRange,
1035 GlobalFlushMemRange, GlobalFlushId);
1036 FlushBufferSize = _FlushBufferSize;
1042 WIOffset = GlobalOffset[1].fetch_add(FlushBufferSize);
1058 flushBuffer(GlobalOffset, GlobalBuf, GlobalFlushBuf, WIOffset);
1062 friend class handler;
1064 template <
typename SYCLObjT>
friend class ext::oneapi::weak_object;
1066 friend const stream &
operator<<(
const stream &,
const char);
1067 friend const stream &
operator<<(
const stream &,
const char *);
1068 template <
typename ValueType>
1070 typename std::enable_if_t<std::is_integral_v<ValueType>,
const stream &>
1071 operator<<(
const stream &,
const ValueType &);
1072 friend const stream &
operator<<(
const stream &,
const float &);
1073 friend const stream &
operator<<(
const stream &,
const double &);
1075 friend const stream &
operator<<(
const stream &,
1076 const ext::oneapi::bfloat16 &);
1080 friend const stream &
operator<<(
const stream &Out,
1081 const __precision_manipulator__ &RHS);
1083 friend const stream &
operator<<(
const stream &Out,
1084 const __width_manipulator__ &RHS);
1085 template <
typename T,
int Dimensions>
1086 friend const stream &
operator<<(
const stream &Out,
1087 const vec<T, Dimensions> &RHS);
1088 template <
typename T>
1089 friend const stream &
operator<<(
const stream &Out,
const T *RHS);
1090 template <
int Dimensions>
1091 friend const stream &
operator<<(
const stream &Out,
const id<Dimensions> &RHS);
1093 template <
int Dimensions>
1094 friend const stream &
operator<<(
const stream &Out,
1095 const range<Dimensions> &RHS);
1097 template <
int Dimensions>
1098 friend const stream &
operator<<(
const stream &Out,
1099 const item<Dimensions> &RHS);
1101 template <
int Dimensions>
1102 friend const stream &
operator<<(
const stream &Out,
1103 const nd_range<Dimensions> &RHS);
1105 template <
int Dimensions>
1106 friend const stream &
operator<<(
const stream &Out,
1107 const nd_item<Dimensions> &RHS);
1109 template <
int Dimensions>
1110 friend const stream &
operator<<(
const stream &Out,
1111 const group<Dimensions> &RHS);
1113 template <
int Dimensions>
1114 friend const stream &
operator<<(
const stream &Out,
1115 const h_item<Dimensions> &RHS);
1118 #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
1120 template <
typename T>
1121 inline std::enable_if_t<std::is_same_v<T, std::byte>,
const stream &>
1123 static_assert(std::is_integral<T>(),
1124 "Convert the byte to a numeric value using std::to_integer");
1129 inline const stream &
operator<<(
const stream &Out,
const char C) {
1130 detail::write(Out.GlobalFlushBuf, Out.FlushBufferSize, Out.WIOffset, &C, 1);
1135 inline const stream &
operator<<(
const stream &Out,
const char *Str) {
1137 for (; Str[Len] !=
'\0'; Len++)
1140 detail::write(Out.GlobalFlushBuf, Out.FlushBufferSize, Out.WIOffset, Str,
1146 inline const stream &
operator<<(
const stream &Out,
const bool &RHS) {
1147 Out << (RHS ?
"true" :
"false");
1152 template <
typename ValueType>
1153 inline typename std::enable_if_t<std::is_integral_v<ValueType>,
const stream &>
1154 operator<<(
const stream &Out,
const ValueType &RHS) {
1156 Out.get_flags(), Out.get_width(), RHS);
1162 inline const stream &
operator<<(
const stream &Out,
const float &RHS) {
1163 detail::writeFloatingPoint<float>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1164 Out.WIOffset, Out.get_flags(),
1165 Out.get_width(), Out.get_precision(), RHS);
1169 inline const stream &
operator<<(
const stream &Out,
const double &RHS) {
1170 detail::writeFloatingPoint<double>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1171 Out.WIOffset, Out.get_flags(),
1172 Out.get_width(), Out.get_precision(), RHS);
1177 detail::writeFloatingPoint<half>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1178 Out.WIOffset, Out.get_flags(),
1179 Out.get_width(), Out.get_precision(), RHS);
1185 detail::writeFloatingPoint<ext::oneapi::bfloat16>(
1186 Out.GlobalFlushBuf, Out.FlushBufferSize, Out.WIOffset, Out.get_flags(),
1187 Out.get_width(), Out.get_precision(), RHS);
1195 inline const stream &
1202 template <
typename T>
1208 Flags, Out.get_width(),
reinterpret_cast<size_t>(RHS));
1222 Out.Width = RHS.
width();
1231 flushBuffer(Out.GlobalOffset, Out.GlobalBuf, Out.GlobalFlushBuf,
1235 flushBuffer(Out.GlobalOffset, Out.GlobalBuf, Out.GlobalFlushBuf,
1239 Out.set_manipulator(RHS);
1247 template <
typename T,
int VectorLength>
1249 detail::writeVec<T, VectorLength>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1250 Out.WIOffset, Out.get_flags(),
1251 Out.get_width(), Out.get_precision(), RHS);
1257 template <
int Dimensions>
1259 detail::writeArray<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1264 template <
int Dimensions>
1267 detail::writeArray<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1272 template <
int Dimensions>
1275 detail::writeItem<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1280 template <
int Dimensions>
1283 detail::writeNDRange<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1288 template <
int Dimensions>
1291 detail::writeNDItem<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1296 template <
int Dimensions>
1299 detail::writeGroup<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1304 template <
int Dimensions>
1307 detail::writeHItem<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1312 template <
typename T,
typename RT = detail::EnableIfSwizzleVec<T>>
1313 inline const stream &
operator<<(
const stream &Out,
const T &RHS) {
1322 template <>
struct hash<
sycl::stream> {
1324 #ifdef __SYCL_DEVICE_ONLY__
1328 return hash<std::shared_ptr<sycl::detail::stream_impl>>()(
The file contains implementations of accessor class.
__precision_manipulator__(int Precision)
__width_manipulator__(int Width)
Identifies an instance of a group::parallel_for_work_item function object executing at each point in ...
item< Dimensions, false > get_logical_local() const
item< Dimensions, false > get_global() const
item< Dimensions, false > get_physical_local() const
A unique identifier of an item in an index space.
Identifies an instance of the function object executing at each point in a range.
id< Dimensions > get_id() const
std::enable_if_t< has_offset, id< Dimensions > > get_offset() const
range< Dimensions > get_range() const
Identifies an instance of the function object executing at each point in an nd_range.
id< Dimensions > get_local_id() const
id< Dimensions > get_global_id() const
Defines the iteration domain of both the work-groups and the overall dispatch.
range< Dimensions > get_global_range() const
range< Dimensions > get_local_range() const
id< Dimensions > get_offset() const
Defines the iteration domain of either a single work-group in a parallel dispatch,...
class sycl::vec ///////////////////////// Provides a cross-patform vector class template that works e...
#define __SYCL_SPECIAL_CLASS
class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor< DataT
Image accessors.
static constexpr FmtFlags Fixed
typename detail::DecoratedType< unsigned, GlobalBufAS >::type * GlobalOffsetPtrType
static constexpr FmtFlags ShowPos
std::enable_if_t<(VecLength==1), unsigned > VecToStr(const vec< T, VecLength > &Vec, char *VecStr, unsigned Flags, int Width, int Precision)
constexpr size_t MAX_ARRAY_SIZE
std::enable_if_t< std::is_integral_v< T > > writeIntegral(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, unsigned Flags, int Width, const T &Val)
static constexpr FmtFlags Oct
constexpr static int GlobalOffsetDim
constexpr size_t MAX_FLOATING_POINT_DIGITS
detail::enable_if_t< std::is_same< T, float >::value||std::is_same< T, double >::value, bool > isFastMathSignBit(T Val)
static constexpr FmtFlags FloatField
std::enable_if_t< std::is_integral_v< T >, unsigned > ScalarToStr(const T &Val, char *Buf, unsigned Flags, int, int Precision=-1)
void writeHItem(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const h_item< Dimensions > &HItem)
typename std::enable_if_t< detail::check_type_in_v< F, float, double, half, ext::oneapi::bfloat16 >, T > EnableIfFP
void writeItem(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const item< Dimensions > &Item)
std::make_unsigned_t< T > getAbsVal(const T Val, const int Base)
void writeVec(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, unsigned Flags, int Width, int Precision, const vec< T, VecLength > &Vec)
typename detail::DecoratedType< char, GlobalBufAS >::type * GlobalBufPtrType
void writeArray(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const array< ArrayLength > &Arr)
void flushBuffer(GlobalOffsetAccessorT &GlobalOffset, GlobalBufAccessorT &GlobalBuf, GlobalBufAccessorT &GlobalFlushBuf, unsigned WIOffset)
EnableIfFP< T > writeFloatingPoint(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, unsigned Flags, int Width, int Precision, const T &Val)
detail::enable_if_t< std::is_same< T, float >::value||std::is_same< T, double >::value, unsigned > checkForInfNan(char *Buf, T Val)
void writeNDItem(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const nd_item< Dimensions > &ND_Item)
accessor< unsigned, 1, sycl::access::mode::atomic, sycl::access::target::device > GlobalOffsetAccessorT
constexpr char VEC_CLOSE_BRACE
accessor< char, 1, sycl::access::mode::read_write, sycl::access::target::device > GlobalBufAccessorT
constexpr const char * VEC_ELEMENT_DELIMITER
unsigned GetFlushBufOffset(const GlobalBufAccessorT &GlobalFlushBuf, unsigned WIOffset)
std::ostream & operator<<(std::ostream &os, std::optional< T > const &opt)
static constexpr FmtFlags ShowBase
constexpr size_t MAX_DIMENSIONS
void reverseBuf(char *Buf, unsigned Len)
decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject)
char digitToChar(const int Digit)
EnableIfFP< T, unsigned > floatingPointToDecStr(T AbsVal, char *Digits, int Precision, bool IsSci)
static constexpr FmtFlags Hex
constexpr char VEC_OPEN_BRACE
void writeGroup(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const group< Dimensions > &Group)
unsigned ItemToStr(char *Buf, const item< Dimensions, false > &Item)
unsigned long long D2I64(double Val)
constexpr static access::address_space GlobalOffsetAS
void SetFlushBufOffset(GlobalBufAccessorT &GlobalFlushBuf, unsigned WIOffset, unsigned Offset)
void writeNDRange(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const nd_range< Dimensions > &ND_Range)
static constexpr FmtFlags BaseField
constexpr size_t MAX_INTEGRAL_DIGITS
struct __SYCL2020_DEPRECATED("This type isn't device copyable in SYCL 2020") IsDeprecatedDeviceCopyable< T
constexpr static int GlobalBufDim
constexpr size_t MAX_ITEM_SIZE
constexpr unsigned FLUSH_BUF_OFFSET_SIZE
detail::enable_if_t< std::is_same< T, float >::value||std::is_same< T, double >::value, bool > isFastMathInf(T Val)
static constexpr FmtFlags Scientific
unsigned append(char *Dst, const char *Src)
void write(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const char *Str, unsigned Len, unsigned Padding=0)
static constexpr FmtFlags Dec
typename std::enable_if_t< IsSwizzleOp< T >::value, typename IsSwizzleOp< T >::Type > EnableIfSwizzleVec
std::enable_if_t< std::is_integral_v< T >, unsigned > integralToBase(T Val, int Base, char *Digits)
unsigned F2I32(float Val)
unsigned ArrayToStr(char *Buf, const array< ArrayLength > &Arr)
bool updateOffset(GlobalOffsetAccessorT &GlobalOffset, GlobalBufAccessorT &GlobalBuf, unsigned Size, unsigned &Cur)
constexpr static access::address_space GlobalBufAS
detail::enable_if_t< std::is_same< T, float >::value||std::is_same< T, double >::value, bool > isFastMathNan(T Val)
bool operator==(const cache_config &lhs, const cache_config &rhs)
bool operator!=(const cache_config &lhs, const cache_config &rhs)
static constexpr bool has_property()
std::enable_if_t< std::is_same_v< T, bfloat16 >, bool > isnan(T x)
static constexpr auto get_property()
constexpr stream_manipulator endl
__precision_manipulator__ setprecision(int Precision)
constexpr stream_manipulator fixed
constexpr stream_manipulator noshowbase
constexpr mode_tag_t< access_mode::read_write > read_write
constexpr stream_manipulator hex
constexpr stream_manipulator flush
__width_manipulator__ setw(int Width)
constexpr stream_manipulator scientific
constexpr stream_manipulator noshowpos
constexpr stream_manipulator showpos
constexpr stream_manipulator dec
constexpr stream_manipulator showbase
constexpr stream_manipulator oct
sycl::detail::half_impl::half half
constexpr stream_manipulator defaultfloat
constexpr stream_manipulator hexfloat
_Abi const simd< _Tp, _Abi > & noexcept
size_t operator()(const sycl::stream &S) const
typename VecT::element_type T
typename sycl::vec< T,(sizeof...(Indexes))> Type