38 #include <type_traits>
42 inline namespace _V1 {
87 template <
class F,
class T =
void>
89 detail::check_type_in_v<F, float, double, half, ext::oneapi::bfloat16>, T>;
92 sycl::access::target::device>;
101 sycl::access::target::device>;
115 return ((
static_cast<unsigned>(
static_cast<uint8_t
>(GlobalFlushBuf[WIOffset]))
117 static_cast<uint8_t
>(GlobalFlushBuf[WIOffset + 1]));
126 unsigned WIOffset,
unsigned Offset) {
127 GlobalFlushBuf[WIOffset] =
static_cast<char>((Offset >> 8) & 0xff);
128 GlobalFlushBuf[WIOffset + 1] =
static_cast<char>(Offset & 0xff);
132 unsigned WIOffset,
const char *Str,
unsigned Len,
133 unsigned Padding = 0) {
137 if ((Offset + Len + Padding > FlushBufferSize) ||
138 (WIOffset + Offset + Len + Padding > GlobalFlushBuf.size()))
143 for (
size_t I = 0; I < Padding; ++I, ++Offset)
144 GlobalFlushBuf[WIOffset + Offset] =
' ';
146 for (
size_t I = 0; I < Len; ++I, ++Offset) {
147 GlobalFlushBuf[WIOffset + Offset] = Str[I];
165 template <
typename T>
166 inline std::make_unsigned_t<T>
getAbsVal(
const T Val,
const int Base) {
167 return ((Base == 10) && (Val < 0)) ? -Val : Val;
174 return 'a' + Digit - 10;
178 template <
typename T>
179 inline typename std::enable_if_t<std::is_integral_v<T>,
unsigned>
181 unsigned NumDigits = 0;
192 template <
typename T>
193 inline typename std::enable_if_t<std::is_integral_v<T>,
unsigned>
194 ScalarToStr(
const T &Val,
char *Buf,
unsigned Flags,
int,
int Precision = -1) {
219 if ((Flags &
ShowPos) && Val >= 0)
226 if (Base != 10 && (Flags &
ShowBase)) {
234 const unsigned NumBuf =
integralToBase(AbsVal, Base, Buf + Offset);
237 return Offset + NumBuf;
240 inline unsigned append(
char *Dst,
const char *Src) {
242 for (; Src[Len] !=
'\0'; ++Len)
245 for (
unsigned I = 0; I < Len; ++I)
256 return Internal.I32Val;
259 inline unsigned long long D2I64(
double Val) {
262 unsigned long long I64Val;
265 return Internal.I64Val;
268 template <
typename T>
269 inline typename detail::enable_if_t<
270 std::is_same<T, float>::value || std::is_same<T, double>::value,
bool>
272 if constexpr (
sizeof(Val) == 4) {
273 return (
F2I32(Val) & 0x7fffffff) == 0x7f800000;
274 }
else if constexpr (
sizeof(Val) == 8) {
275 return (
D2I64(Val) & -1ULL >> 1) == 0x7ffULL << 52;
281 template <
typename T>
282 inline typename detail::enable_if_t<
283 std::is_same<T, float>::value || std::is_same<T, double>::value,
bool>
285 if constexpr (
sizeof(Val) == 4) {
286 return (
F2I32(Val) & 0x7fffffff) > 0x7f800000;
287 }
else if constexpr (
sizeof(Val) == 8) {
288 return (
D2I64(Val) & -1ULL >> 1) > 0x7ffULL << 52;
294 template <
typename T>
295 inline typename detail::enable_if_t<
296 std::is_same<T, float>::value || std::is_same<T, double>::value,
bool>
298 if constexpr (
sizeof(Val) == 4) {
299 return F2I32(Val) >> 31;
300 }
else if constexpr (
sizeof(Val) == 8) {
301 return D2I64(Val) >> 63;
307 template <
typename T>
308 typename detail::enable_if_t<
309 std::is_same<T, float>::value || std::is_same<T, double>::value,
unsigned>
316 return append(Buf,
"nan");
324 return append(Buf,
"-inf");
325 return append(Buf,
"inf");
330 template <
typename T>
331 inline typename std::enable_if_t<std::is_same_v<T, half>,
unsigned>
334 return append(Buf,
"nan");
337 const uint16_t Sign = sycl::bit_cast<uint16_t>(Val) & 0x8000;
339 const uint16_t Exp16 = (sycl::bit_cast<uint16_t>(Val) & 0x7c00) >> 10;
343 return append(Buf,
"-inf");
344 return append(Buf,
"inf");
349 template <
typename T>
350 inline typename std::enable_if_t<std::is_same_v<T, ext::oneapi::bfloat16>,
354 return append(Buf,
"nan");
357 const uint16_t Sign = sycl::bit_cast<uint16_t>(Val) & 0x8000;
359 const uint16_t Exp16 = (sycl::bit_cast<uint16_t>(Val) & 0x7f80) >> 7;
363 return append(Buf,
"-inf");
364 return append(Buf,
"inf");
369 template <
typename T>
371 int Precision,
bool IsSci) {
375 while (AbsVal >= T{10.0}) {
380 while (AbsVal > T{0.0} && AbsVal < T{1.0}) {
385 auto IntegralPart =
static_cast<int>(AbsVal);
386 auto FractionPart = AbsVal - IntegralPart;
391 int P = Precision > 0 ? Precision : 4;
392 size_t FractionLength = Exp + P;
401 for (
unsigned I = 0; I < FractionLength; ++I) {
402 FractionPart *= T{10.0};
403 FractionDigits[I] =
static_cast<int>(FractionPart);
404 FractionPart -=
static_cast<int>(FractionPart);
407 int Carry = FractionPart >
static_cast<T
>(0.5) ? 1 : 0;
410 for (
int I = FractionLength - 1; I >= 0 && Carry; --I) {
411 auto Digit = FractionDigits[I] + Carry;
412 FractionDigits[I] = Digit % 10;
417 IntegralPart += Carry;
418 if (IntegralPart == 10) {
429 Digits[Offset++] =
'.';
432 for (
unsigned I = 0; I < FractionLength; ++I)
435 auto AbsExp = Exp < 0 ? -Exp : Exp;
437 Digits[Offset++] =
'e';
438 Digits[Offset++] = Exp >= 0 ?
'+' :
'-';
443 Digits[Offset++] =
'0';
444 Digits[Offset++] =
'.';
446 Digits[Offset++] =
'0';
452 for (
unsigned I = 0; I < FractionLength; ++I)
459 for (; I < FractionLength && Exp--; ++I)
463 Digits[Offset++] =
'.';
466 for (; I < FractionLength; ++I)
471 while (Digits[Offset - 1] ==
'0')
475 if (Digits[Offset - 1] ==
'.')
482 template <
typename T>
483 inline EnableIfFP<T, unsigned>
484 ScalarToStr(
const T &Val,
char *Buf,
unsigned Flags,
int,
int Precision = -1) {
490 auto AbsVal = Val < 0 ? Neg : Val;
508 template <
typename T>
509 inline typename std::enable_if_t<std::is_integral_v<T>>
511 unsigned WIOffset,
unsigned Flags,
int Width,
const T &Val) {
513 unsigned Len =
ScalarToStr(Val, Digits, Flags, Width);
514 write(GlobalFlushBuf, FlushBufferSize, WIOffset, Digits, Len,
515 (Width > 0 &&
static_cast<unsigned>(Width) > Len)
516 ?
static_cast<unsigned>(Width) - Len
520 template <
typename T>
523 unsigned WIOffset,
unsigned Flags,
int Width,
int Precision,
526 unsigned Len =
ScalarToStr(Val, Digits, Flags, Width, Precision);
527 write(GlobalFlushBuf, FlushBufferSize, WIOffset, Digits, Len,
528 (Width > 0 &&
static_cast<unsigned>(Width) > Len)
529 ?
static_cast<unsigned>(Width) - Len
540 Cur = GlobalOffset[0].load();
542 if (GlobalBuf.get_range().size() - Cur < Size)
546 }
while (!GlobalOffset[0].compare_exchange_strong(Cur, New));
558 if (!
updateOffset(GlobalOffset, GlobalBuf, Offset, Cur))
562 for (
unsigned I = StmtOffset; I < StmtOffset + Offset; I++) {
563 GlobalBuf[Cur++] = GlobalFlushBuf[I];
569 template <
typename T,
int VecLength>
570 typename std::enable_if_t<(VecLength == 1),
unsigned>
573 return ScalarToStr(
static_cast<T
>(Vec.x()), VecStr, Flags, Width, Precision);
576 template <
typename T,
int VecLength>
577 typename std::enable_if_t<(VecLength == 2 || VecLength == 4 || VecLength == 8 ||
583 VecToStr<T, VecLength / 2>(Vec.lo(), VecStr, Flags, Width, Precision);
585 Len +=
VecToStr<T, VecLength / 2>(Vec.hi(), VecStr + Len, Flags, Width,
590 template <
typename T,
int VecLength>
591 typename std::enable_if_t<(VecLength == 3),
unsigned>
594 unsigned Len = VecToStr<T, 2>(Vec.lo(), VecStr, Flags, Width, Precision);
596 Len += VecToStr<T, 1>(Vec.z(), VecStr + Len, Flags, Width, Precision);
600 template <
typename T,
int VecLength>
602 unsigned WIOffset,
unsigned Flags,
int Width,
605 constexpr
size_t MAX_VEC_SIZE =
607 char VecStr[MAX_VEC_SIZE] = {0};
608 unsigned Len = VecToStr<T, VecLength>(Vec, VecStr, Flags, Width, Precision);
609 write(GlobalFlushBuf, FlushBufferSize, WIOffset, VecStr, Len,
610 (Width > 0 && Width > Len) ? Width - Len : 0);
613 template <
int ArrayLength>
618 for (
int I = 0; I < ArrayLength; ++I) {
620 if (I != ArrayLength - 1)
629 template <
int ArrayLength>
631 size_t FlushBufferSize,
unsigned WIOffset,
635 write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len);
638 template <
int Dimensions>
640 size_t FlushBufferSize,
unsigned WIOffset,
646 Len +=
append(Buf,
"item(");
647 Len +=
append(Buf + Len,
"range: ");
649 Len +=
append(Buf + Len,
", id: ");
651 Len +=
append(Buf + Len,
", offset: ");
654 write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len);
657 template <
int Dimensions>
659 size_t FlushBufferSize,
unsigned WIOffset,
665 Len +=
append(Buf,
"nd_range(");
666 Len +=
append(Buf + Len,
"global_range: ");
668 Len +=
append(Buf + Len,
", local_range: ");
670 Len +=
append(Buf + Len,
", offset: ");
673 write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len);
676 template <
int Dimensions>
678 size_t FlushBufferSize,
unsigned WIOffset,
684 Len +=
append(Buf,
"nd_item(");
685 Len +=
append(Buf + Len,
"global_id: ");
687 Len +=
append(Buf + Len,
", local_id: ");
690 write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len);
693 template <
int Dimensions>
695 size_t FlushBufferSize,
unsigned WIOffset,
701 Len +=
append(Buf,
"group(");
702 Len +=
append(Buf + Len,
"id: ");
704 Len +=
append(Buf + Len,
", global_range: ");
705 Len +=
ArrayToStr(Buf + Len, Group.get_global_range());
706 Len +=
append(Buf + Len,
", local_range: ");
707 Len +=
ArrayToStr(Buf + Len, Group.get_local_range());
708 Len +=
append(Buf + Len,
", group_range: ");
709 Len +=
ArrayToStr(Buf + Len, Group.get_group_range());
711 write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len);
718 template <
int Dimensions>
721 Len +=
append(Buf,
"item(");
722 for (
int I = 0; I < 2; ++I) {
723 Len +=
append(Buf + Len, I == 0 ?
"range: " :
", id: ");
730 template <
int Dimensions>
732 size_t FlushBufferSize,
unsigned WIOffset,
738 Len +=
append(Buf,
"h_item(");
739 for (
int I = 0; I < 3; ++I) {
740 Len +=
append(Buf + Len, I == 0 ?
"\n global "
741 : I == 1 ?
"\n logical local "
742 :
"\n physical local ");
747 Len +=
append(Buf + Len,
"\n)");
748 write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len);
753 template <
typename VecT,
typename OperationLeftT,
typename OperationRightT,
754 template <
typename>
class OperationCurrentT,
int... Indexes>
756 VecT, OperationLeftT, OperationRightT, OperationCurrentT, Indexes...>>
758 using T =
typename VecT::element_type;
762 template <
typename T>
764 typename std::enable_if_t<IsSwizzleOp<T>::value,
821 friend const stream &
operator<<(
const stream &,
831 int width()
const {
return Width_; }
833 friend const stream &
operator<<(
const stream &,
850 :
public detail::OwnerLessBase<stream> {
852 #ifndef __SYCL_DEVICE_ONLY__
854 stream(std::shared_ptr<detail::stream_impl> Impl,
858 : impl{Impl}, GlobalBuf{GlobalBuf}, GlobalOffset{GlobalOffset},
859 GlobalFlushBuf{GlobalFlushBuf} {}
863 #ifdef __SYCL_DEVICE_ONLY__
869 stream(
size_t BufferSize,
size_t MaxStatementSize, handler &CGH);
873 stream(
size_t BufferSize,
size_t MaxStatementSize, handler &CGH,
874 const property_list &PropList);
876 #ifdef __SYCL_DEVICE_ONLY__
880 size_t size() const
noexcept {
return GlobalBuf.size(); }
882 size_t get_work_item_buffer_size()
const {
887 "get_size() is deprecated since SYCL 2020. Please use size() instead.")
888 size_t get_size()
const {
return size(); }
891 "2020. Please use get_work_item_buffer_size() instead.")
892 size_t get_max_statement_size()
const {
return get_work_item_buffer_size(); }
896 size_t get_work_item_buffer_size() const;
899 "get_size() is deprecated since SYCL 2020. Please
use size()
instead.")
900 size_t get_size() const;
903 "2020. Please
use get_work_item_buffer_size()
instead.")
904 size_t get_max_statement_size() const;
907 size_t get_precision()
const {
return Precision; }
909 size_t get_width()
const {
return Width; }
918 return getPropList().template has_property<propertyT>();
921 template <
typename propertyT> propertyT
get_property()
const {
922 return getPropList().template get_property<propertyT>();
926 #ifdef __SYCL_DEVICE_ONLY__
927 char padding[
sizeof(std::shared_ptr<detail::stream_impl>)];
929 std::shared_ptr<detail::stream_impl> impl;
931 friend const decltype(Obj::impl) &
953 mutable unsigned WIOffset = 0;
955 mutable size_t FlushBufferSize;
963 mutable int Precision = -1;
964 mutable int Width = -1;
967 void set_flag(
FmtFlags FormatFlag)
const { Flags |= FormatFlag; }
969 void unset_flag(
FmtFlags FormatFlag)
const { Flags &= ~FormatFlag; }
971 FmtFlags get_flags()
const {
return Flags; }
978 Flags |= FormatFlag & Mask;
1023 #ifdef __SYCL_DEVICE_ONLY__
1025 range<detail::GlobalBufDim> GlobalBufAccRange,
1026 range<detail::GlobalBufDim> GlobalBufMemRange,
1027 id<detail::GlobalBufDim> GlobalBufId,
1029 range<detail::GlobalOffsetDim> GlobalOffsetAccRange,
1030 range<detail::GlobalOffsetDim> GlobalOffsetMemRange,
1031 id<detail::GlobalOffsetDim> GlobalOffsetId,
1033 range<detail::GlobalBufDim> GlobalFlushAccRange,
1034 range<detail::GlobalBufDim> GlobalFlushMemRange,
1035 id<detail::GlobalBufDim> GlobalFlushId,
size_t _FlushBufferSize) {
1036 GlobalBuf.__init(GlobalBufPtr, GlobalBufAccRange, GlobalBufMemRange,
1038 GlobalOffset.__init(GlobalOffsetPtr, GlobalOffsetAccRange,
1039 GlobalOffsetMemRange, GlobalOffsetId);
1040 GlobalFlushBuf.__init(GlobalFlushPtr, GlobalFlushAccRange,
1041 GlobalFlushMemRange, GlobalFlushId);
1042 FlushBufferSize = _FlushBufferSize;
1048 WIOffset = GlobalOffset[1].fetch_add(FlushBufferSize);
1059 flushBuffer(GlobalOffset, GlobalBuf, GlobalFlushBuf, WIOffset);
1063 friend class handler;
1065 template <
typename SYCLObjT>
friend class ext::oneapi::weak_object;
1067 friend const stream &
operator<<(
const stream &,
const char);
1068 friend const stream &
operator<<(
const stream &,
const char *);
1069 template <
typename ValueType>
1071 typename std::enable_if_t<std::is_integral_v<ValueType>,
const stream &>
1072 operator<<(
const stream &,
const ValueType &);
1073 friend const stream &
operator<<(
const stream &,
const float &);
1074 friend const stream &
operator<<(
const stream &,
const double &);
1076 friend const stream &
operator<<(
const stream &,
1081 friend const stream &
operator<<(
const stream &Out,
1082 const __precision_manipulator__ &RHS);
1084 friend const stream &
operator<<(
const stream &Out,
1085 const __width_manipulator__ &RHS);
1086 template <
typename T,
int Dimensions>
1087 friend const stream &
operator<<(
const stream &Out,
1088 const vec<T, Dimensions> &RHS);
1089 template <
typename T>
1090 friend const stream &
operator<<(
const stream &Out,
const T *RHS);
1091 template <
int Dimensions>
1092 friend const stream &
operator<<(
const stream &Out,
const id<Dimensions> &RHS);
1094 template <
int Dimensions>
1095 friend const stream &
operator<<(
const stream &Out,
1096 const range<Dimensions> &RHS);
1098 template <
int Dimensions>
1099 friend const stream &
operator<<(
const stream &Out,
1100 const item<Dimensions> &RHS);
1102 template <
int Dimensions>
1103 friend const stream &
operator<<(
const stream &Out,
1104 const nd_range<Dimensions> &RHS);
1106 template <
int Dimensions>
1107 friend const stream &
operator<<(
const stream &Out,
1108 const nd_item<Dimensions> &RHS);
1110 template <
int Dimensions>
1111 friend const stream &
operator<<(
const stream &Out,
1112 const group<Dimensions> &RHS);
1114 template <
int Dimensions>
1115 friend const stream &
operator<<(
const stream &Out,
1116 const h_item<Dimensions> &RHS);
1118 const property_list &getPropList()
const;
1121 #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
1123 template <
typename T>
1124 inline std::enable_if_t<std::is_same_v<T, std::byte>,
const stream &>
1126 static_assert(std::is_integral<T>(),
1127 "Convert the byte to a numeric value using std::to_integer");
1132 inline const stream &
operator<<(
const stream &Out,
const char C) {
1133 detail::write(Out.GlobalFlushBuf, Out.FlushBufferSize, Out.WIOffset, &C, 1);
1138 inline const stream &
operator<<(
const stream &Out,
const char *Str) {
1140 for (; Str[Len] !=
'\0'; Len++)
1143 detail::write(Out.GlobalFlushBuf, Out.FlushBufferSize, Out.WIOffset, Str,
1149 inline const stream &
operator<<(
const stream &Out,
const bool &RHS) {
1150 Out << (RHS ?
"true" :
"false");
1155 template <
typename ValueType>
1156 inline typename std::enable_if_t<std::is_integral_v<ValueType>,
const stream &>
1157 operator<<(
const stream &Out,
const ValueType &RHS) {
1159 Out.get_flags(), Out.get_width(), RHS);
1165 inline const stream &
operator<<(
const stream &Out,
const float &RHS) {
1166 detail::writeFloatingPoint<float>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1167 Out.WIOffset, Out.get_flags(),
1168 Out.get_width(), Out.get_precision(), RHS);
1172 inline const stream &
operator<<(
const stream &Out,
const double &RHS) {
1173 detail::writeFloatingPoint<double>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1174 Out.WIOffset, Out.get_flags(),
1175 Out.get_width(), Out.get_precision(), RHS);
1180 detail::writeFloatingPoint<half>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1181 Out.WIOffset, Out.get_flags(),
1182 Out.get_width(), Out.get_precision(), RHS);
1188 detail::writeFloatingPoint<ext::oneapi::bfloat16>(
1189 Out.GlobalFlushBuf, Out.FlushBufferSize, Out.WIOffset, Out.get_flags(),
1190 Out.get_width(), Out.get_precision(), RHS);
1198 inline const stream &
1205 template <
typename T>
1211 Flags, Out.get_width(),
reinterpret_cast<size_t>(RHS));
1225 Out.Width = RHS.
width();
1234 flushBuffer(Out.GlobalOffset, Out.GlobalBuf, Out.GlobalFlushBuf,
1238 flushBuffer(Out.GlobalOffset, Out.GlobalBuf, Out.GlobalFlushBuf,
1242 Out.set_manipulator(RHS);
1250 template <
typename T,
int VectorLength>
1252 detail::writeVec<T, VectorLength>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1253 Out.WIOffset, Out.get_flags(),
1254 Out.get_width(), Out.get_precision(), RHS);
1260 template <
int Dimensions>
1262 detail::writeArray<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1267 template <
int Dimensions>
1270 detail::writeArray<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1275 template <
int Dimensions>
1278 detail::writeItem<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1283 template <
int Dimensions>
1286 detail::writeNDRange<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1291 template <
int Dimensions>
1294 detail::writeNDItem<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1299 template <
int Dimensions>
1302 detail::writeGroup<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1307 template <
int Dimensions>
1310 detail::writeHItem<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1315 template <
typename T,
typename RT = detail::EnableIfSwizzleVec<T>>
1316 inline const stream &
operator<<(
const stream &Out,
const T &RHS) {
1325 template <>
struct hash<
sycl::stream> {
1327 #ifdef __SYCL_DEVICE_ONLY__
1331 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,...
#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)
decltype(Obj::impl) const & getSyclObjImpl(const Obj &SyclObject)
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)
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
sycl::ext::oneapi::bfloat16 bfloat16
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