62 template <
class F,
class T =
void>
65 std::is_same<F, double>::value ||
66 std::is_same<F, half>::value,
70 sycl::access::target::global_buffer,
71 sycl::access::placeholder::false_t>;
80 sycl::access::target::global_buffer,
81 sycl::access::placeholder::false_t>;
95 return ((
static_cast<unsigned>(
static_cast<uint8_t
>(GlobalFlushBuf[WIOffset]))
97 static_cast<uint8_t
>(GlobalFlushBuf[WIOffset + 1]));
106 unsigned WIOffset,
unsigned Offset) {
107 GlobalFlushBuf[WIOffset] =
static_cast<char>((Offset >> 8) & 0xff);
108 GlobalFlushBuf[WIOffset + 1] =
static_cast<char>(Offset & 0xff);
112 unsigned WIOffset,
const char *Str,
unsigned Len,
113 unsigned Padding = 0) {
117 if ((Offset + Len + Padding > FlushBufferSize) ||
118 (WIOffset + Offset + Len + Padding > GlobalFlushBuf.size()))
123 for (
size_t I = 0; I < Padding; ++I, ++Offset)
124 GlobalFlushBuf[WIOffset + Offset] =
' ';
126 for (
size_t I = 0; I < Len; ++I, ++Offset) {
127 GlobalFlushBuf[WIOffset + Offset] = Str[I];
145 template <
typename T>
146 inline typename std::make_unsigned<T>::type
getAbsVal(
const T Val,
148 return ((Base == 10) && (Val < 0)) ? -Val : Val;
155 return 'a' + Digit - 10;
159 template <
typename T>
162 unsigned NumDigits = 0;
173 template <
typename T>
175 ScalarToStr(
const T &Val,
char *Buf,
unsigned Flags,
int,
int Precision = -1) {
200 if ((Flags &
ShowPos) && Val >= 0)
207 if (Base != 10 && (Flags &
ShowBase)) {
215 const unsigned NumBuf =
integralToBase(AbsVal, Base, Buf + Offset);
218 return Offset + NumBuf;
221 inline unsigned append(
char *Dst,
const char *Src) {
223 for (; Src[Len] !=
'\0'; ++Len)
226 for (
unsigned I = 0; I < Len; ++I)
231 template <
typename T>
233 std::is_same<T, float>::value || std::is_same<T, double>::value,
unsigned>
236 return append(Buf,
"nan");
239 return append(Buf,
"-inf");
240 return append(Buf,
"inf");
245 template <
typename T>
249 return append(Buf,
"nan");
252 const uint16_t Sign =
reinterpret_cast<uint16_t &
>(Val) & 0x8000;
254 const uint16_t Exp16 = (
reinterpret_cast<uint16_t &
>(Val) & 0x7c00) >> 10;
258 return append(Buf,
"-inf");
259 return append(Buf,
"inf");
264 template <
typename T>
266 int Precision,
bool IsSci) {
270 while (AbsVal >= T{10.0}) {
275 while (AbsVal > T{0.0} && AbsVal < T{1.0}) {
280 auto IntegralPart =
static_cast<int>(AbsVal);
281 auto FractionPart = AbsVal - IntegralPart;
286 int P = Precision > 0 ? Precision : 4;
287 size_t FractionLength = Exp + P;
296 for (
unsigned I = 0; I < FractionLength; ++I) {
297 FractionPart *= T{10.0};
298 FractionDigits[I] =
static_cast<int>(FractionPart);
299 FractionPart -=
static_cast<int>(FractionPart);
302 int Carry = FractionPart >
static_cast<T
>(0.5) ? 1 : 0;
305 for (
int I = FractionLength - 1; I >= 0 && Carry; --I) {
306 auto Digit = FractionDigits[I] + Carry;
307 FractionDigits[I] = Digit % 10;
312 IntegralPart += Carry;
313 if (IntegralPart == 10) {
324 Digits[Offset++] =
'.';
327 for (
unsigned I = 0; I < FractionLength; ++I)
330 auto AbsExp = Exp < 0 ? -Exp : Exp;
332 Digits[Offset++] =
'e';
333 Digits[Offset++] = Exp >= 0 ?
'+' :
'-';
338 Digits[Offset++] =
'0';
339 Digits[Offset++] =
'.';
341 Digits[Offset++] =
'0';
347 for (
unsigned I = 0; I < FractionLength; ++I)
354 for (; I < FractionLength && Exp--; ++I)
358 Digits[Offset++] =
'.';
361 for (; I < FractionLength; ++I)
366 while (Digits[Offset - 1] ==
'0')
370 if (Digits[Offset - 1] ==
'.')
377 template <
typename T>
378 inline EnableIfFP<T, unsigned>
379 ScalarToStr(
const T &Val,
char *Buf,
unsigned Flags,
int,
int Precision = -1) {
385 auto AbsVal = Val < 0 ? Neg : Val;
403 template <
typename T>
406 unsigned WIOffset,
unsigned Flags,
int Width,
const T &Val) {
408 unsigned Len =
ScalarToStr(Val, Digits, Flags, Width);
409 write(GlobalFlushBuf, FlushBufferSize, WIOffset, Digits, Len,
410 (Width > 0 &&
static_cast<unsigned>(Width) > Len)
411 ?
static_cast<unsigned>(Width) - Len
415 template <
typename T>
418 unsigned WIOffset,
unsigned Flags,
int Width,
int Precision,
421 unsigned Len =
ScalarToStr(Val, Digits, Flags, Width, Precision);
422 write(GlobalFlushBuf, FlushBufferSize, WIOffset, Digits, Len,
423 (Width > 0 &&
static_cast<unsigned>(Width) > Len)
424 ?
static_cast<unsigned>(Width) - Len
435 Cur = GlobalOffset[0].load();
437 if (GlobalBuf.get_range().size() - Cur < Size)
441 }
while (!GlobalOffset[0].compare_exchange_strong(Cur, New));
453 if (!
updateOffset(GlobalOffset, GlobalBuf, Offset, Cur))
457 for (
unsigned I = StmtOffset; I < StmtOffset + Offset; I++) {
458 GlobalBuf[Cur++] = GlobalFlushBuf[I];
464 template <
typename T,
int VecLength>
468 return ScalarToStr(
static_cast<T
>(Vec.x()), VecStr, Flags, Width, Precision);
471 template <
typename T,
int VecLength>
473 VecLength == 8 || VecLength == 16),
478 VecToStr<T, VecLength / 2>(Vec.lo(), VecStr, Flags, Width, Precision);
480 Len +=
VecToStr<T, VecLength / 2>(Vec.hi(), VecStr + Len, Flags, Width,
485 template <
typename T,
int VecLength>
489 unsigned Len = VecToStr<T, 2>(Vec.lo(), VecStr, Flags, Width, Precision);
491 Len += VecToStr<T, 1>(Vec.z(), VecStr + Len, Flags, Width, Precision);
495 template <
typename T,
int VecLength>
497 unsigned WIOffset,
unsigned Flags,
int Width,
500 constexpr
size_t MAX_VEC_SIZE =
502 char VecStr[MAX_VEC_SIZE] = {0};
503 unsigned Len = VecToStr<T, VecLength>(Vec, VecStr, Flags, Width, Precision);
504 write(GlobalFlushBuf, FlushBufferSize, WIOffset, VecStr, Len,
505 (Width > 0 && Width > Len) ? Width - Len : 0);
508 template <
int ArrayLength>
513 for (
int I = 0; I < ArrayLength; ++I) {
515 if (I != ArrayLength - 1)
524 template <
int ArrayLength>
526 size_t FlushBufferSize,
unsigned WIOffset,
530 write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len);
533 template <
int Dimensions>
535 size_t FlushBufferSize,
unsigned WIOffset,
541 Len +=
append(Buf,
"item(");
542 Len +=
append(Buf + Len,
"range: ");
544 Len +=
append(Buf + Len,
", id: ");
546 Len +=
append(Buf + Len,
", offset: ");
549 write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len);
552 template <
int Dimensions>
554 size_t FlushBufferSize,
unsigned WIOffset,
560 Len +=
append(Buf,
"nd_range(");
561 Len +=
append(Buf + Len,
"global_range: ");
563 Len +=
append(Buf + Len,
", local_range: ");
565 Len +=
append(Buf + Len,
", offset: ");
568 write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len);
571 template <
int Dimensions>
573 size_t FlushBufferSize,
unsigned WIOffset,
579 Len +=
append(Buf,
"nd_item(");
580 Len +=
append(Buf + Len,
"global_id: ");
582 Len +=
append(Buf + Len,
", local_id: ");
585 write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len);
588 template <
int Dimensions>
590 size_t FlushBufferSize,
unsigned WIOffset,
596 Len +=
append(Buf,
"group(");
597 Len +=
append(Buf + Len,
"id: ");
599 Len +=
append(Buf + Len,
", global_range: ");
600 Len +=
ArrayToStr(Buf + Len, Group.get_global_range());
601 Len +=
append(Buf + Len,
", local_range: ");
602 Len +=
ArrayToStr(Buf + Len, Group.get_local_range());
603 Len +=
append(Buf + Len,
", group_range: ");
604 Len +=
ArrayToStr(Buf + Len, Group.get_group_range());
606 write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len);
613 template <
int Dimensions>
616 Len +=
append(Buf,
"item(");
617 for (
int I = 0; I < 2; ++I) {
618 Len +=
append(Buf + Len, I == 0 ?
"range: " :
", id: ");
625 template <
int Dimensions>
627 size_t FlushBufferSize,
unsigned WIOffset,
633 Len +=
append(Buf,
"h_item(");
634 for (
int I = 0; I < 3; ++I) {
635 Len +=
append(Buf + Len, I == 0 ?
"\n global "
636 : I == 1 ?
"\n logical local "
637 :
"\n physical local ");
642 Len +=
append(Buf + Len,
"\n)");
643 write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len);
648 template <
typename VecT,
typename OperationLeftT,
typename OperationRightT,
649 template <
typename>
class OperationCurrentT,
int... Indexes>
651 VecT, OperationLeftT, OperationRightT, OperationCurrentT, Indexes...>>
653 using T =
typename VecT::element_type;
654 using Type =
typename sycl::vec<
T, (
sizeof...(Indexes))>;
657 template <
typename T>
716 friend const stream &
operator<<(
const stream &,
726 int width()
const {
return Width_; }
728 friend const stream &
operator<<(
const stream &,
745 :
public detail::OwnerLessBase<stream> {
747 #ifdef __SYCL_DEVICE_ONLY__
753 stream(
size_t BufferSize,
size_t MaxStatementSize, handler &CGH);
757 stream(
size_t BufferSize,
size_t MaxStatementSize, handler &CGH,
758 const property_list &PropList);
760 size_t get_size()
const;
762 size_t get_max_statement_size()
const;
764 size_t get_precision()
const {
return Precision; }
766 size_t get_width()
const {
return Width; }
774 template <
typename propertyT>
bool has_property() const noexcept;
776 template <typename propertyT> propertyT
get_property() const;
779 #ifdef __SYCL_DEVICE_ONLY__
780 char padding[
sizeof(std::shared_ptr<detail::stream_impl>)];
782 std::shared_ptr<detail::stream_impl> impl;
801 mutable unsigned WIOffset = 0;
803 mutable size_t FlushBufferSize;
811 mutable int Precision = -1;
812 mutable int Width = -1;
815 void set_flag(
FmtFlags FormatFlag)
const { Flags |= FormatFlag; }
817 void unset_flag(
FmtFlags FormatFlag)
const { Flags &= ~FormatFlag; }
819 FmtFlags get_flags()
const {
return Flags; }
826 Flags |= FormatFlag & Mask;
871 #ifdef __SYCL_DEVICE_ONLY__
873 range<detail::GlobalBufDim> GlobalBufAccRange,
874 range<detail::GlobalBufDim> GlobalBufMemRange,
875 id<detail::GlobalBufDim> GlobalBufId,
877 range<detail::GlobalOffsetDim> GlobalOffsetAccRange,
878 range<detail::GlobalOffsetDim> GlobalOffsetMemRange,
879 id<detail::GlobalOffsetDim> GlobalOffsetId,
881 range<detail::GlobalBufDim> GlobalFlushAccRange,
882 range<detail::GlobalBufDim> GlobalFlushMemRange,
883 id<detail::GlobalBufDim> GlobalFlushId,
size_t _FlushBufferSize) {
884 GlobalBuf.__init(GlobalBufPtr, GlobalBufAccRange, GlobalBufMemRange,
886 GlobalOffset.__init(GlobalOffsetPtr, GlobalOffsetAccRange,
887 GlobalOffsetMemRange, GlobalOffsetId);
888 GlobalFlushBuf.__init(GlobalFlushPtr, GlobalFlushAccRange,
889 GlobalFlushMemRange, GlobalFlushId);
890 FlushBufferSize = _FlushBufferSize;
896 WIOffset = GlobalOffset[1].fetch_add(FlushBufferSize);
912 flushBuffer(GlobalOffset, GlobalBuf, GlobalFlushBuf, WIOffset);
916 friend class handler;
918 friend const stream &
operator<<(
const stream &,
const char);
919 friend const stream &
operator<<(
const stream &,
const char *);
920 template <
typename ValueType>
921 friend typename detail::enable_if_t<std::is_integral<ValueType>::value,
923 operator<<(
const stream &,
const ValueType &);
924 friend const stream &
operator<<(
const stream &,
const float &);
925 friend const stream &
operator<<(
const stream &,
const double &);
930 friend const stream &
operator<<(
const stream &Out,
931 const __precision_manipulator__ &RHS);
933 friend const stream &
operator<<(
const stream &Out,
934 const __width_manipulator__ &RHS);
935 template <
typename T,
int Dimensions>
936 friend const stream &
operator<<(
const stream &Out,
937 const vec<T, Dimensions> &RHS);
938 template <
typename T>
939 friend const stream &
operator<<(
const stream &Out,
const T *RHS);
940 template <
int Dimensions>
941 friend const stream &
operator<<(
const stream &Out,
const id<Dimensions> &RHS);
943 template <
int Dimensions>
944 friend const stream &
operator<<(
const stream &Out,
945 const range<Dimensions> &RHS);
947 template <
int Dimensions>
948 friend const stream &
operator<<(
const stream &Out,
949 const item<Dimensions> &RHS);
951 template <
int Dimensions>
952 friend const stream &
operator<<(
const stream &Out,
953 const nd_range<Dimensions> &RHS);
955 template <
int Dimensions>
956 friend const stream &
operator<<(
const stream &Out,
957 const nd_item<Dimensions> &RHS);
959 template <
int Dimensions>
960 friend const stream &
operator<<(
const stream &Out,
961 const group<Dimensions> &RHS);
963 template <
int Dimensions>
964 friend const stream &
operator<<(
const stream &Out,
965 const h_item<Dimensions> &RHS);
968 #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
970 template <
typename T>
971 inline std::enable_if_t<std::is_same<T, std::byte>::value,
const stream &>
973 static_assert(std::is_integral<T>(),
974 "Convert the byte to a numeric value using std::to_integer");
979 inline const stream &
operator<<(
const stream &Out,
const char C) {
980 detail::write(Out.GlobalFlushBuf, Out.FlushBufferSize, Out.WIOffset, &C, 1);
985 inline const stream &
operator<<(
const stream &Out,
const char *Str) {
987 for (; Str[Len] !=
'\0'; Len++)
990 detail::write(Out.GlobalFlushBuf, Out.FlushBufferSize, Out.WIOffset, Str,
996 inline const stream &
operator<<(
const stream &Out,
const bool &RHS) {
997 Out << (RHS ?
"true" :
"false");
1002 template <
typename ValueType>
1003 inline typename detail::enable_if_t<std::is_integral<ValueType>::value,
1005 operator<<(
const stream &Out,
const ValueType &RHS) {
1007 Out.get_flags(), Out.get_width(), RHS);
1013 inline const stream &
operator<<(
const stream &Out,
const float &RHS) {
1014 detail::writeFloatingPoint<float>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1015 Out.WIOffset, Out.get_flags(),
1016 Out.get_width(), Out.get_precision(), RHS);
1020 inline const stream &
operator<<(
const stream &Out,
const double &RHS) {
1021 detail::writeFloatingPoint<double>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1022 Out.WIOffset, Out.get_flags(),
1023 Out.get_width(), Out.get_precision(), RHS);
1028 detail::writeFloatingPoint<half>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1029 Out.WIOffset, Out.get_flags(),
1030 Out.get_width(), Out.get_precision(), RHS);
1038 inline const stream &
1045 template <
typename T>
1051 Flags, Out.get_width(),
reinterpret_cast<size_t>(RHS));
1065 Out.Width = RHS.
width();
1074 flushBuffer(Out.GlobalOffset, Out.GlobalBuf, Out.GlobalFlushBuf,
1078 flushBuffer(Out.GlobalOffset, Out.GlobalBuf, Out.GlobalFlushBuf,
1082 Out.set_manipulator(RHS);
1090 template <
typename T,
int VectorLength>
1092 detail::writeVec<T, VectorLength>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1093 Out.WIOffset, Out.get_flags(),
1094 Out.get_width(), Out.get_precision(), RHS);
1100 template <
int Dimensions>
1102 detail::writeArray<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1107 template <
int Dimensions>
1110 detail::writeArray<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1115 template <
int Dimensions>
1118 detail::writeItem<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1123 template <
int Dimensions>
1126 detail::writeNDRange<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1131 template <
int Dimensions>
1134 detail::writeNDItem<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1139 template <
int Dimensions>
1142 detail::writeGroup<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1147 template <
int Dimensions>
1150 detail::writeHItem<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1155 template <
typename T,
typename RT = detail::EnableIfSwizzleVec<T>>
1156 inline const stream &
operator<<(
const stream &Out,
const T &RHS) {
1165 template <>
struct hash<
sycl::stream> {
1167 #ifdef __SYCL_DEVICE_ONLY__
1171 return hash<std::shared_ptr<sycl::detail::stream_impl>>()(
__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_physical_local() const
item< dimensions, false > get_global() const
item< dimensions, false > get_logical_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.
range< dimensions > get_range() const
id< dimensions > get_id() const
detail::enable_if_t< has_offset, id< dimensions > > get_offset() const
Provides constructors for address space qualified and non address space qualified pointers to allow i...
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_local_range() const
id< dimensions > get_offset() const
range< dimensions > get_global_range() const
Provides a cross-patform vector class template that works efficiently on SYCL devices as well as in h...
#define __SYCL_SPECIAL_CLASS
#define __SYCL_INLINE_VER_NAMESPACE(X)
static constexpr FmtFlags Fixed
typename detail::DecoratedType< unsigned, GlobalBufAS >::type * GlobalOffsetPtrType
static constexpr FmtFlags ShowPos
EnableIfFP< T, unsigned > ScalarToStr(const T &Val, char *Buf, unsigned Flags, int, int Precision=-1)
constexpr size_t MAX_ARRAY_SIZE
static constexpr FmtFlags Oct
constexpr static int GlobalOffsetDim
constexpr size_t MAX_FLOATING_POINT_DIGITS
detail::enable_if_t< std::is_same< T, half >::value, unsigned > checkForInfNan(char *Buf, T Val)
static constexpr FmtFlags FloatField
void writeHItem(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const h_item< Dimensions > &HItem)
typename detail::enable_if_t< std::is_same< F, float >::value||std::is_same< F, double >::value||std::is_same< F, half >::value, T > EnableIfFP
std::make_unsigned< T >::type getAbsVal(const T Val, const int Base)
void writeItem(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const item< Dimensions > &Item)
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)
void writeNDItem(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const nd_item< Dimensions > &ND_Item)
accessor< char, 1, sycl::access::mode::read_write, sycl::access::target::global_buffer, sycl::access::placeholder::false_t > GlobalBufAccessorT
constexpr char VEC_CLOSE_BRACE
accessor< unsigned, 1, sycl::access::mode::atomic, sycl::access::target::global_buffer, sycl::access::placeholder::false_t > GlobalOffsetAccessorT
constexpr const char * VEC_ELEMENT_DELIMITER
unsigned GetFlushBufOffset(const GlobalBufAccessorT &GlobalFlushBuf, unsigned WIOffset)
static constexpr FmtFlags ShowBase
constexpr size_t MAX_DIMENSIONS
void reverseBuf(char *Buf, unsigned Len)
decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject)
typename detail::enable_if_t< IsSwizzleOp< T >::value, typename IsSwizzleOp< T >::Type > EnableIfSwizzleVec
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)
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)
detail::enable_if_t<(VecLength==3), unsigned > VecToStr(const vec< T, VecLength > &Vec, char *VecStr, unsigned Flags, int Width, int Precision)
static constexpr FmtFlags BaseField
constexpr size_t MAX_INTEGRAL_DIGITS
constexpr static int GlobalBufDim
constexpr size_t MAX_ITEM_SIZE
constexpr unsigned FLUSH_BUF_OFFSET_SIZE
detail::enable_if_t< std::is_integral< T >::value, unsigned > integralToBase(T Val, int Base, char *Digits)
static constexpr FmtFlags Scientific
typename std::enable_if< B, T >::type enable_if_t
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
detail::enable_if_t< std::is_integral< T >::value > writeIntegral(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, unsigned Flags, int Width, const T &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
static constexpr bool has_property()
std::enable_if_t< std::is_same< T, bfloat16 >::value, bool > isnan(T x)
static constexpr auto get_property()
constexpr std::enable_if_t< detail::IsCompileTimeProperty< PropertyT >::value, bool > operator!=(const property_value< PropertyT, A... > &, const property_value< PropertyT, B... > &)
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
const stream & operator<<(const stream &Out, const T &RHS)
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
---— Error handling, matching OpenCL plugin semantics.
size_t operator()(const sycl::stream &S) const
typename VecT::element_type T
typename sycl::vec< T,(sizeof...(Indexes))> Type
bool operator==(const Slab &Lhs, const Slab &Rhs)