DPC++ Runtime
Runtime libraries for oneAPI DPC++
stream.hpp
Go to the documentation of this file.
1 //==----------------- stream.hpp - SYCL standard header file ---------------==//
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 
9 #pragma once
10 
11 #include <sycl/builtins.hpp>
12 #include <sycl/detail/defines.hpp>
13 #include <sycl/detail/export.hpp>
16 #include <sycl/handler.hpp>
17 
18 namespace sycl {
20 
21 namespace detail {
22 
23 using FmtFlags = unsigned int;
24 
25 // Mapping from stream_manipulator to FmtFlags. Each manipulator corresponds
26 // to the bit in FmtFlags.
27 static constexpr FmtFlags Dec = 0x0001;
28 static constexpr FmtFlags Hex = 0x0002;
29 static constexpr FmtFlags Oct = 0x0004;
30 static constexpr FmtFlags ShowBase = 0x0008;
31 static constexpr FmtFlags ShowPos = 0x0010;
32 static constexpr FmtFlags Fixed = 0x0020;
33 static constexpr FmtFlags Scientific = 0x0040;
34 
35 // Bitmask made of the combination of the base flags. Base flags are mutually
36 // exclusive, this mask is used to clean base field before setting the new
37 // base flag.
38 static constexpr FmtFlags BaseField = Dec | Hex | Oct;
39 
40 // Bitmask made of the combination of the floating point value format flags.
41 // Thease flags are mutually exclusive, this mask is used to clean float field
42 // before setting the new float flag.
43 static constexpr FmtFlags FloatField = Scientific | Fixed;
44 
45 constexpr size_t MAX_FLOATING_POINT_DIGITS = 24;
46 constexpr size_t MAX_INTEGRAL_DIGITS = 23;
47 constexpr const char *VEC_ELEMENT_DELIMITER = ", ";
48 constexpr char VEC_OPEN_BRACE = '{';
49 constexpr char VEC_CLOSE_BRACE = '}';
50 
51 constexpr size_t MAX_DIMENSIONS = 3;
52 
53 // Space for integrals (up to 3), comma and space between the
54 // integrals and enclosing braces.
55 constexpr size_t MAX_ARRAY_SIZE =
57 
58 // First 2 bytes in each work item's flush buffer are reserved for saving
59 // statement offset.
60 constexpr unsigned FLUSH_BUF_OFFSET_SIZE = 2;
61 
62 template <class F, class T = void>
63 using EnableIfFP = typename std::enable_if_t<std::is_same_v<F, float> ||
64  std::is_same_v<F, double> ||
65  std::is_same_v<F, half>,
66  T>;
67 
69  sycl::access::target::global_buffer,
70  sycl::access::placeholder::false_t>;
71 
74 using GlobalBufPtrType =
76 constexpr static int GlobalBufDim = 1;
77 
78 using GlobalOffsetAccessorT = accessor<unsigned, 1, sycl::access::mode::atomic,
79  sycl::access::target::global_buffer,
80  sycl::access::placeholder::false_t>;
81 
84 using GlobalOffsetPtrType =
86 constexpr static int GlobalOffsetDim = 1;
87 
88 // Read first 2 bytes of flush buffer to get buffer offset.
89 // TODO: Should be optimized to the following:
90 // return *reinterpret_cast<uint16_t *>(&GlobalFlushBuf[WIOffset]);
91 // when an issue with device code compilation using this optimization is fixed.
92 inline unsigned GetFlushBufOffset(const GlobalBufAccessorT &GlobalFlushBuf,
93  unsigned WIOffset) {
94  return ((static_cast<unsigned>(static_cast<uint8_t>(GlobalFlushBuf[WIOffset]))
95  << 8) +
96  static_cast<uint8_t>(GlobalFlushBuf[WIOffset + 1]));
97 }
98 
99 // Write flush buffer's offset into first 2 bytes of that buffer.
100 // TODO: Should be optimized to the following:
101 // *reinterpret_cast<uint16_t *>(&GlobalFlushBuf[WIOffset]) =
102 // static_cast<uint16_t>(Offset);
103 // when an issue with device code compilation using this optimization is fixed.
104 inline void SetFlushBufOffset(GlobalBufAccessorT &GlobalFlushBuf,
105  unsigned WIOffset, unsigned Offset) {
106  GlobalFlushBuf[WIOffset] = static_cast<char>((Offset >> 8) & 0xff);
107  GlobalFlushBuf[WIOffset + 1] = static_cast<char>(Offset & 0xff);
108 }
109 
110 inline void write(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize,
111  unsigned WIOffset, const char *Str, unsigned Len,
112  unsigned Padding = 0) {
113  unsigned Offset =
114  GetFlushBufOffset(GlobalFlushBuf, WIOffset) + FLUSH_BUF_OFFSET_SIZE;
115 
116  if ((Offset + Len + Padding > FlushBufferSize) ||
117  (WIOffset + Offset + Len + Padding > GlobalFlushBuf.size()))
118  // TODO: flush here
119  return;
120 
121  // Write padding
122  for (size_t I = 0; I < Padding; ++I, ++Offset)
123  GlobalFlushBuf[WIOffset + Offset] = ' ';
124 
125  for (size_t I = 0; I < Len; ++I, ++Offset) {
126  GlobalFlushBuf[WIOffset + Offset] = Str[I];
127  }
128 
129  SetFlushBufOffset(GlobalFlushBuf, WIOffset, Offset - FLUSH_BUF_OFFSET_SIZE);
130 }
131 
132 inline void reverseBuf(char *Buf, unsigned Len) {
133  int I = Len - 1;
134  int J = 0;
135  while (I > J) {
136  int Temp = Buf[I];
137  Buf[I] = Buf[J];
138  Buf[J] = Temp;
139  I--;
140  J++;
141  }
142 }
143 
144 template <typename T>
145 inline std::make_unsigned_t<T> getAbsVal(const T Val, const int Base) {
146  return ((Base == 10) && (Val < 0)) ? -Val : Val;
147 }
148 
149 inline char digitToChar(const int Digit) {
150  if (Digit < 10) {
151  return '0' + Digit;
152  } else {
153  return 'a' + Digit - 10;
154  }
155 }
156 
157 template <typename T>
158 inline typename std::enable_if_t<std::is_integral_v<T>, unsigned>
159 integralToBase(T Val, int Base, char *Digits) {
160  unsigned NumDigits = 0;
161 
162  do {
163  Digits[NumDigits++] = digitToChar(Val % Base);
164  Val /= Base;
165  } while (Val);
166 
167  return NumDigits;
168 }
169 
170 // Returns number of symbols written to the buffer
171 template <typename T>
172 inline typename std::enable_if_t<std::is_integral_v<T>, unsigned>
173 ScalarToStr(const T &Val, char *Buf, unsigned Flags, int, int Precision = -1) {
174  (void)Precision;
175  int Base = 10;
176 
177  // append base manipulator
178  switch (Flags & BaseField) {
179  case Dec:
180  Base = 10;
181  break;
182  case Hex:
183  Base = 16;
184  break;
185  case Oct:
186  Base = 8;
187  break;
188  default:
189  // default value is 10
190  break;
191  }
192 
193  unsigned Offset = 0;
194 
195  // write '+' to the stream if the base is 10 and the value is non-negative
196  // or write '-' to stream if base is 10 and the value is negative
197  if (Base == 10) {
198  if ((Flags & ShowPos) && Val >= 0)
199  Buf[Offset++] = '+';
200  else if (Val < 0)
201  Buf[Offset++] = '-';
202  }
203 
204  // write 0 or 0x to the stream if base is not 10 and the manipulator is set
205  if (Base != 10 && (Flags & ShowBase)) {
206  Buf[Offset++] = '0';
207  if (Base == 16)
208  Buf[Offset++] = 'x';
209  }
210 
211  auto AbsVal = getAbsVal(Val, Base);
212 
213  const unsigned NumBuf = integralToBase(AbsVal, Base, Buf + Offset);
214 
215  reverseBuf(Buf + Offset, NumBuf);
216  return Offset + NumBuf;
217 }
218 
219 inline unsigned append(char *Dst, const char *Src) {
220  unsigned Len = 0;
221  for (; Src[Len] != '\0'; ++Len)
222  ;
223 
224  for (unsigned I = 0; I < Len; ++I)
225  Dst[I] = Src[I];
226  return Len;
227 }
228 
229 static inline unsigned F2I32(float Val) {
230  union {
231  float FVal;
232  unsigned I32Val;
233  } Internal;
234  Internal.FVal = Val;
235  return Internal.I32Val;
236 }
237 
238 static inline unsigned long long D2I64(double Val) {
239  union {
240  double DVal;
241  unsigned long long I64Val;
242  } Internal;
243  Internal.DVal = Val;
244  return Internal.I64Val;
245 }
246 
247 template <typename T>
248 inline typename detail::enable_if_t<
249  std::is_same<T, float>::value || std::is_same<T, double>::value, bool>
251  if constexpr (sizeof(Val) == 4) {
252  return (F2I32(Val) & 0x7fffffff) == 0x7f800000;
253  } else if constexpr (sizeof(Val) == 8) {
254  return (D2I64(Val) & -1ULL >> 1) == 0x7ffULL << 52;
255  }
256 
257  return false;
258 }
259 
260 template <typename T>
261 inline typename detail::enable_if_t<
262  std::is_same<T, float>::value || std::is_same<T, double>::value, bool>
264  if constexpr (sizeof(Val) == 4) {
265  return (F2I32(Val) & 0x7fffffff) > 0x7f800000;
266  } else if constexpr (sizeof(Val) == 8) {
267  return (D2I64(Val) & -1ULL >> 1) > 0x7ffULL << 52;
268  }
269 
270  return false;
271 }
272 
273 template <typename T>
274 inline typename detail::enable_if_t<
275  std::is_same<T, float>::value || std::is_same<T, double>::value, bool>
277  if constexpr (sizeof(Val) == 4) {
278  return F2I32(Val) >> 31;
279  } else if constexpr (sizeof(Val) == 8) {
280  return D2I64(Val) >> 63;
281  }
282 
283  return false;
284 }
285 
286 template <typename T>
287 typename detail::enable_if_t<
288  std::is_same<T, float>::value || std::is_same<T, double>::value, unsigned>
289 checkForInfNan(char *Buf, T Val) {
290 #ifdef __FAST_MATH__
291  if (isFastMathNan(Val))
292 #else
293  if (isnan(Val))
294 #endif
295  return append(Buf, "nan");
296 #ifdef __FAST_MATH__
297  if (isFastMathInf(Val)) {
298  if (isFastMathSignBit(Val))
299 #else
300  if (isinf(Val)) {
301  if (signbit(Val))
302 #endif
303  return append(Buf, "-inf");
304  return append(Buf, "inf");
305  }
306  return 0;
307 }
308 
309 template <typename T>
310 inline typename std::enable_if_t<std::is_same_v<T, half>, unsigned>
311 checkForInfNan(char *Buf, T Val) {
312  if (Val != Val)
313  return append(Buf, "nan");
314 
315  // Extract the sign from the bits
316  const uint16_t Sign = reinterpret_cast<uint16_t &>(Val) & 0x8000;
317  // Extract the exponent from the bits
318  const uint16_t Exp16 = (reinterpret_cast<uint16_t &>(Val) & 0x7c00) >> 10;
319 
320  if (Exp16 == 0x1f) {
321  if (Sign)
322  return append(Buf, "-inf");
323  return append(Buf, "inf");
324  }
325  return 0;
326 }
327 
328 template <typename T>
330  int Precision, bool IsSci) {
331  int Exp = 0;
332 
333  // For the case that the value is larger than 10.0
334  while (AbsVal >= T{10.0}) {
335  ++Exp;
336  AbsVal /= T{10.0};
337  }
338  // For the case that the value is less than 1.0
339  while (AbsVal > T{0.0} && AbsVal < T{1.0}) {
340  --Exp;
341  AbsVal *= T{10.0};
342  }
343 
344  auto IntegralPart = static_cast<int>(AbsVal);
345  auto FractionPart = AbsVal - IntegralPart;
346 
347  int FractionDigits[MAX_FLOATING_POINT_DIGITS] = {0};
348 
349  // Exponent
350  int P = Precision > 0 ? Precision : 4;
351  size_t FractionLength = Exp + P;
352 
353  // After normalization integral part contains 1 symbol, also there could be
354  // '.', 'e', sign of the exponent and sign of the number, overall 5 symbols.
355  // So, clamp fraction length if required according to maximum size of the
356  // buffer for floating point number.
357  if (FractionLength > MAX_FLOATING_POINT_DIGITS - 5)
358  FractionLength = MAX_FLOATING_POINT_DIGITS - 5;
359 
360  for (unsigned I = 0; I < FractionLength; ++I) {
361  FractionPart *= T{10.0};
362  FractionDigits[I] = static_cast<int>(FractionPart);
363  FractionPart -= static_cast<int>(FractionPart);
364  }
365 
366  int Carry = FractionPart > static_cast<T>(0.5) ? 1 : 0;
367 
368  // Propagate the Carry
369  for (int I = FractionLength - 1; I >= 0 && Carry; --I) {
370  auto Digit = FractionDigits[I] + Carry;
371  FractionDigits[I] = Digit % 10;
372  Carry = Digit / 10;
373  }
374 
375  // Carry from the fraction part is propagated to integral part
376  IntegralPart += Carry;
377  if (IntegralPart == 10) {
378  IntegralPart = 1;
379  ++Exp;
380  }
381 
382  unsigned Offset = 0;
383 
384  // Assemble the final string correspondingly
385  if (IsSci) { // scientific mode
386  // Append the integral part
387  Digits[Offset++] = digitToChar(IntegralPart);
388  Digits[Offset++] = '.';
389 
390  // Append all fraction
391  for (unsigned I = 0; I < FractionLength; ++I)
392  Digits[Offset++] = digitToChar(FractionDigits[I]);
393 
394  auto AbsExp = Exp < 0 ? -Exp : Exp;
395  // Exponent part
396  Digits[Offset++] = 'e';
397  Digits[Offset++] = Exp >= 0 ? '+' : '-';
398  Digits[Offset++] = digitToChar(AbsExp / 10);
399  Digits[Offset++] = digitToChar(AbsExp % 10);
400  } else { // normal mode
401  if (Exp < 0) {
402  Digits[Offset++] = '0';
403  Digits[Offset++] = '.';
404  while (++Exp)
405  Digits[Offset++] = '0';
406 
407  // Append the integral part
408  Digits[Offset++] = digitToChar(IntegralPart);
409 
410  // Append all fraction
411  for (unsigned I = 0; I < FractionLength; ++I)
412  Digits[Offset++] = digitToChar(FractionDigits[I]);
413  } else {
414  // Append the integral part
415  Digits[Offset++] = digitToChar(IntegralPart);
416  unsigned I = 0;
417  // Append the integral part first
418  for (; I < FractionLength && Exp--; ++I)
419  Digits[Offset++] = digitToChar(FractionDigits[I]);
420 
421  // Put the dot
422  Digits[Offset++] = '.';
423 
424  // Append the rest of fraction part, or the real fraction part
425  for (; I < FractionLength; ++I)
426  Digits[Offset++] = digitToChar(FractionDigits[I]);
427  }
428  // The normal mode requires no tailing zero digit, then we need to first
429  // find the first non-zero digit
430  while (Digits[Offset - 1] == '0')
431  Offset--;
432 
433  // If dot is the last digit, it should be stripped off as well
434  if (Digits[Offset - 1] == '.')
435  Offset--;
436  }
437  return Offset;
438 }
439 
440 // Returns number of symbols written to the buffer
441 template <typename T>
442 inline EnableIfFP<T, unsigned>
443 ScalarToStr(const T &Val, char *Buf, unsigned Flags, int, int Precision = -1) {
444  unsigned Offset = checkForInfNan(Buf, Val);
445  if (Offset)
446  return Offset;
447 
448  T Neg = -Val;
449  auto AbsVal = Val < 0 ? Neg : Val;
450 
451  if (Val < 0) {
452  Buf[Offset++] = '-';
453  } else if (Flags & ShowPos) {
454  Buf[Offset++] = '+';
455  }
456 
457  bool IsSci = false;
458  if (Flags & detail::Scientific)
459  IsSci = true;
460 
461  // TODO: manipulators for floating-point output - hexfloat, fixed
462  Offset += floatingPointToDecStr(AbsVal, Buf + Offset, Precision, IsSci);
463 
464  return Offset;
465 }
466 
467 template <typename T>
468 inline typename std::enable_if_t<std::is_integral_v<T>>
469 writeIntegral(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize,
470  unsigned WIOffset, unsigned Flags, int Width, const T &Val) {
471  char Digits[MAX_INTEGRAL_DIGITS] = {0};
472  unsigned Len = ScalarToStr(Val, Digits, Flags, Width);
473  write(GlobalFlushBuf, FlushBufferSize, WIOffset, Digits, Len,
474  (Width > 0 && static_cast<unsigned>(Width) > Len)
475  ? static_cast<unsigned>(Width) - Len
476  : 0);
477 }
478 
479 template <typename T>
480 inline EnableIfFP<T>
481 writeFloatingPoint(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize,
482  unsigned WIOffset, unsigned Flags, int Width, int Precision,
483  const T &Val) {
484  char Digits[MAX_FLOATING_POINT_DIGITS] = {0};
485  unsigned Len = ScalarToStr(Val, Digits, Flags, Width, Precision);
486  write(GlobalFlushBuf, FlushBufferSize, WIOffset, Digits, Len,
487  (Width > 0 && static_cast<unsigned>(Width) > Len)
488  ? static_cast<unsigned>(Width) - Len
489  : 0);
490 }
491 
492 // Helper method to update offset in the global buffer atomically according to
493 // the provided size of the data in the flush buffer. Return true if offset is
494 // updated and false in case of overflow.
495 inline bool updateOffset(GlobalOffsetAccessorT &GlobalOffset,
496  GlobalBufAccessorT &GlobalBuf, unsigned Size,
497  unsigned &Cur) {
498  unsigned New;
499  Cur = GlobalOffset[0].load();
500  do {
501  if (GlobalBuf.get_range().size() - Cur < Size)
502  // Overflow
503  return false;
504  New = Cur + Size;
505  } while (!GlobalOffset[0].compare_exchange_strong(Cur, New));
506  return true;
507 }
508 
509 inline void flushBuffer(GlobalOffsetAccessorT &GlobalOffset,
510  GlobalBufAccessorT &GlobalBuf,
511  GlobalBufAccessorT &GlobalFlushBuf, unsigned WIOffset) {
512  unsigned Offset = GetFlushBufOffset(GlobalFlushBuf, WIOffset);
513  if (Offset == 0)
514  return;
515 
516  unsigned Cur = 0;
517  if (!updateOffset(GlobalOffset, GlobalBuf, Offset, Cur))
518  return;
519 
520  unsigned StmtOffset = WIOffset + FLUSH_BUF_OFFSET_SIZE;
521  for (unsigned I = StmtOffset; I < StmtOffset + Offset; I++) {
522  GlobalBuf[Cur++] = GlobalFlushBuf[I];
523  }
524  // Reset the offset in the flush buffer
525  SetFlushBufOffset(GlobalFlushBuf, WIOffset, 0);
526 }
527 
528 template <typename T, int VecLength>
529 typename std::enable_if_t<(VecLength == 1), unsigned>
530 VecToStr(const vec<T, VecLength> &Vec, char *VecStr, unsigned Flags, int Width,
531  int Precision) {
532  return ScalarToStr(static_cast<T>(Vec.x()), VecStr, Flags, Width, Precision);
533 }
534 
535 template <typename T, int VecLength>
536 typename std::enable_if_t<(VecLength == 2 || VecLength == 4 || VecLength == 8 ||
537  VecLength == 16),
538  unsigned>
539 VecToStr(const vec<T, VecLength> &Vec, char *VecStr, unsigned Flags, int Width,
540  int Precision) {
541  unsigned Len =
542  VecToStr<T, VecLength / 2>(Vec.lo(), VecStr, Flags, Width, Precision);
543  Len += append(VecStr + Len, VEC_ELEMENT_DELIMITER);
544  Len += VecToStr<T, VecLength / 2>(Vec.hi(), VecStr + Len, Flags, Width,
545  Precision);
546  return Len;
547 }
548 
549 template <typename T, int VecLength>
550 typename std::enable_if_t<(VecLength == 3), unsigned>
551 VecToStr(const vec<T, VecLength> &Vec, char *VecStr, unsigned Flags, int Width,
552  int Precision) {
553  unsigned Len = VecToStr<T, 2>(Vec.lo(), VecStr, Flags, Width, Precision);
554  Len += append(VecStr + Len, VEC_ELEMENT_DELIMITER);
555  Len += VecToStr<T, 1>(Vec.z(), VecStr + Len, Flags, Width, Precision);
556  return Len;
557 }
558 
559 template <typename T, int VecLength>
560 inline void writeVec(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize,
561  unsigned WIOffset, unsigned Flags, int Width,
562  int Precision, const vec<T, VecLength> &Vec) {
563  // Reserve space for vector elements and delimiters
564  constexpr size_t MAX_VEC_SIZE =
565  MAX_FLOATING_POINT_DIGITS * VecLength + (VecLength - 1) * 2;
566  char VecStr[MAX_VEC_SIZE] = {0};
567  unsigned Len = VecToStr<T, VecLength>(Vec, VecStr, Flags, Width, Precision);
568  write(GlobalFlushBuf, FlushBufferSize, WIOffset, VecStr, Len,
569  (Width > 0 && Width > Len) ? Width - Len : 0);
570 }
571 
572 template <int ArrayLength>
573 inline unsigned ArrayToStr(char *Buf, const array<ArrayLength> &Arr) {
574  unsigned Len = 0;
575  Buf[Len++] = VEC_OPEN_BRACE;
576 
577  for (int I = 0; I < ArrayLength; ++I) {
578  Len += ScalarToStr(Arr[I], Buf + Len, 0 /* No flags */, -1, -1);
579  if (I != ArrayLength - 1)
580  Len += append(Buf + Len, VEC_ELEMENT_DELIMITER);
581  }
582 
583  Buf[Len++] = VEC_CLOSE_BRACE;
584 
585  return Len;
586 }
587 
588 template <int ArrayLength>
589 inline void writeArray(GlobalBufAccessorT &GlobalFlushBuf,
590  size_t FlushBufferSize, unsigned WIOffset,
591  const array<ArrayLength> &Arr) {
592  char Buf[MAX_ARRAY_SIZE];
593  unsigned Len = ArrayToStr(Buf, Arr);
594  write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len);
595 }
596 
597 template <int Dimensions>
598 inline void writeItem(GlobalBufAccessorT &GlobalFlushBuf,
599  size_t FlushBufferSize, unsigned WIOffset,
600  const item<Dimensions> &Item) {
601  // Reserve space for 3 arrays and additional place (40 symbols) for printing
602  // the text
603  char Buf[3 * MAX_ARRAY_SIZE + 40];
604  unsigned Len = 0;
605  Len += append(Buf, "item(");
606  Len += append(Buf + Len, "range: ");
607  Len += ArrayToStr(Buf + Len, Item.get_range());
608  Len += append(Buf + Len, ", id: ");
609  Len += ArrayToStr(Buf + Len, Item.get_id());
610  Len += append(Buf + Len, ", offset: ");
611  Len += ArrayToStr(Buf + Len, Item.get_offset());
612  Buf[Len++] = ')';
613  write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len);
614 }
615 
616 template <int Dimensions>
617 inline void writeNDRange(GlobalBufAccessorT &GlobalFlushBuf,
618  size_t FlushBufferSize, unsigned WIOffset,
619  const nd_range<Dimensions> &ND_Range) {
620  // Reserve space for 3 arrays and additional place (50 symbols) for printing
621  // the text
622  char Buf[3 * MAX_ARRAY_SIZE + 50];
623  unsigned Len = 0;
624  Len += append(Buf, "nd_range(");
625  Len += append(Buf + Len, "global_range: ");
626  Len += ArrayToStr(Buf + Len, ND_Range.get_global_range());
627  Len += append(Buf + Len, ", local_range: ");
628  Len += ArrayToStr(Buf + Len, ND_Range.get_local_range());
629  Len += append(Buf + Len, ", offset: ");
630  Len += ArrayToStr(Buf + Len, ND_Range.get_offset());
631  Buf[Len++] = ')';
632  write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len);
633 }
634 
635 template <int Dimensions>
636 inline void writeNDItem(GlobalBufAccessorT &GlobalFlushBuf,
637  size_t FlushBufferSize, unsigned WIOffset,
638  const nd_item<Dimensions> &ND_Item) {
639  // Reserve space for 2 arrays and additional place (40 symbols) for printing
640  // the text
641  char Buf[2 * MAX_ARRAY_SIZE + 40];
642  unsigned Len = 0;
643  Len += append(Buf, "nd_item(");
644  Len += append(Buf + Len, "global_id: ");
645  Len += ArrayToStr(Buf + Len, ND_Item.get_global_id());
646  Len += append(Buf + Len, ", local_id: ");
647  Len += ArrayToStr(Buf + Len, ND_Item.get_local_id());
648  Buf[Len++] = ')';
649  write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len);
650 }
651 
652 template <int Dimensions>
653 inline void writeGroup(GlobalBufAccessorT &GlobalFlushBuf,
654  size_t FlushBufferSize, unsigned WIOffset,
655  const group<Dimensions> &Group) {
656  // Reserve space for 4 arrays and additional place (60 symbols) for printing
657  // the text
658  char Buf[4 * MAX_ARRAY_SIZE + 60];
659  unsigned Len = 0;
660  Len += append(Buf, "group(");
661  Len += append(Buf + Len, "id: ");
662  Len += ArrayToStr(Buf + Len, Group.get_id());
663  Len += append(Buf + Len, ", global_range: ");
664  Len += ArrayToStr(Buf + Len, Group.get_global_range());
665  Len += append(Buf + Len, ", local_range: ");
666  Len += ArrayToStr(Buf + Len, Group.get_local_range());
667  Len += append(Buf + Len, ", group_range: ");
668  Len += ArrayToStr(Buf + Len, Group.get_group_range());
669  Buf[Len++] = ')';
670  write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len);
671 }
672 
673 // Space for 2 arrays and additional place (20 symbols) for printing
674 // the text
675 constexpr size_t MAX_ITEM_SIZE = 2 * MAX_ARRAY_SIZE + 20;
676 
677 template <int Dimensions>
678 inline unsigned ItemToStr(char *Buf, const item<Dimensions, false> &Item) {
679  unsigned Len = 0;
680  Len += append(Buf, "item(");
681  for (int I = 0; I < 2; ++I) {
682  Len += append(Buf + Len, I == 0 ? "range: " : ", id: ");
683  Len += ArrayToStr(Buf + Len, I == 0 ? Item.get_range() : Item.get_id());
684  }
685  Buf[Len++] = ')';
686  return Len;
687 }
688 
689 template <int Dimensions>
690 inline void writeHItem(GlobalBufAccessorT &GlobalFlushBuf,
691  size_t FlushBufferSize, unsigned WIOffset,
692  const h_item<Dimensions> &HItem) {
693  // Reserve space for 3 items and additional place (60 symbols) for printing
694  // the text
695  char Buf[3 * MAX_ITEM_SIZE + 60];
696  unsigned Len = 0;
697  Len += append(Buf, "h_item(");
698  for (int I = 0; I < 3; ++I) {
699  Len += append(Buf + Len, I == 0 ? "\n global "
700  : I == 1 ? "\n logical local "
701  : "\n physical local ");
702  Len += ItemToStr(Buf + Len, I == 0 ? HItem.get_global()
703  : I == 1 ? HItem.get_logical_local()
704  : HItem.get_physical_local());
705  }
706  Len += append(Buf + Len, "\n)");
707  write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len);
708 }
709 
710 template <typename> struct IsSwizzleOp : std::false_type {};
711 
712 template <typename VecT, typename OperationLeftT, typename OperationRightT,
713  template <typename> class OperationCurrentT, int... Indexes>
715  VecT, OperationLeftT, OperationRightT, OperationCurrentT, Indexes...>>
716  : std::true_type {
717  using T = typename VecT::element_type;
718  using Type = typename sycl::vec<T, (sizeof...(Indexes))>;
719 };
720 
721 template <typename T>
722 using EnableIfSwizzleVec =
723  typename std::enable_if_t<IsSwizzleOp<T>::value,
725 
726 } // namespace detail
727 
728 enum class stream_manipulator {
729  dec = 0,
730  hex = 1,
731  oct = 2,
732  noshowbase = 3,
733  showbase = 4,
734  noshowpos = 5,
735  showpos = 6,
736  endl = 7,
737  flush = 8,
738  fixed = 9,
739  scientific = 10,
740  hexfloat = 11,
741  defaultfloat = 12
742 };
743 
745 
747 
749 
751 
753 
755 
757 
759 
761 
763 
765 
767 
769 
770 class stream;
771 
773  int Precision_;
774 
775 public:
776  __precision_manipulator__(int Precision) : Precision_(Precision) {}
777 
778  int precision() const { return Precision_; }
779 
780  friend const stream &operator<<(const stream &,
781  const __precision_manipulator__ &);
782 };
783 
785  int Width_;
786 
787 public:
788  __width_manipulator__(int Width) : Width_(Width) {}
789 
790  int width() const { return Width_; }
791 
792  friend const stream &operator<<(const stream &,
793  const __width_manipulator__ &);
794 };
795 
797  return __precision_manipulator__(Precision);
798 }
799 
800 inline __width_manipulator__ setw(int Width) {
801  return __width_manipulator__(Width);
802 }
803 
808 class __SYCL_EXPORT __SYCL_SPECIAL_CLASS __SYCL_TYPE(stream) stream
809  : public detail::OwnerLessBase<stream> {
810 private:
811 #ifndef __SYCL_DEVICE_ONLY__
812  // Constructor for recreating a stream.
813  stream(std::shared_ptr<detail::stream_impl> Impl,
814  detail::GlobalBufAccessorT GlobalBuf,
815  detail::GlobalOffsetAccessorT GlobalOffset,
816  detail::GlobalBufAccessorT GlobalFlushBuf)
817  : impl{Impl}, GlobalBuf{GlobalBuf}, GlobalOffset{GlobalOffset},
818  GlobalFlushBuf{GlobalFlushBuf} {}
819 #endif
820 
821 public:
822 #ifdef __SYCL_DEVICE_ONLY__
823  // Default constructor for objects later initialized with __init member.
824  stream() = default;
825 #endif
826 
827  // Throws exception in case of invalid input parameters
828  stream(size_t BufferSize, size_t MaxStatementSize, handler &CGH);
829 
830  // Property-list constructor variant.
831  // TODO: Merge with other stream constructor and give PropList default value.
832  stream(size_t BufferSize, size_t MaxStatementSize, handler &CGH,
833  const property_list &PropList);
834 
835 #ifdef __SYCL_DEVICE_ONLY__
836  // We need the definitions of these functions in the header for device,
837  // otherwise they are not visible. Also, we cannot use `impl` because it's not
838  // there on the device, so we rely on GlobalBuf/GlobalFlushBuf.
839  size_t size() const noexcept { return GlobalBuf.size(); }
840 
841  size_t get_work_item_buffer_size() const {
842  return GlobalFlushBuf.size() - detail::FLUSH_BUF_OFFSET_SIZE;
843  }
844 
846  "get_size() is deprecated since SYCL 2020. Please use size() instead.")
847  size_t get_size() const { return size(); }
848 
849  __SYCL2020_DEPRECATED("get_max_statement_size() is deprecated since SYCL "
850  "2020. Please use get_work_item_buffer_size() instead.")
851  size_t get_max_statement_size() const { return get_work_item_buffer_size(); }
852 #else
853  size_t size() const noexcept;
854 
855  size_t get_work_item_buffer_size() const;
856 
857  __SYCL2020_DEPRECATED(
858  "get_size() is deprecated since SYCL 2020. Please use size() instead.")
859  size_t get_size() const;
860 
861  __SYCL2020_DEPRECATED("get_max_statement_size() is deprecated since SYCL "
862  "2020. Please use get_work_item_buffer_size() instead.")
863  size_t get_max_statement_size() const;
864 #endif
865 
866  size_t get_precision() const { return Precision; }
867 
868  size_t get_width() const { return Width; }
869 
870  stream_manipulator get_stream_mode() const { return Manipulator; }
871 
872  bool operator==(const stream &RHS) const;
873 
874  bool operator!=(const stream &LHS) const;
875 
876  template <typename propertyT> bool has_property() const noexcept;
877 
878  template <typename propertyT> propertyT get_property() const;
879 
880 private:
881 #ifdef __SYCL_DEVICE_ONLY__
882  char padding[sizeof(std::shared_ptr<detail::stream_impl>)];
883 #else
884  std::shared_ptr<detail::stream_impl> impl;
885  template <class Obj>
886  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
887 #endif
888 
889  // NOTE: Some members are required for reconstructing the stream, but are not
890  // part of the implementation class. If more members are added, they should
891  // also be added to the weak_object specialization for streams.
892 
893  // Accessor to the global stream buffer. Global buffer contains all output
894  // from the kernel.
895  mutable detail::GlobalBufAccessorT GlobalBuf;
896 
897  // Atomic accessor to the global offset variable. It represents an offset in
898  // the global stream buffer. Since work items will flush data to global buffer
899  // in parallel we need atomic access to this offset.
900  mutable detail::GlobalOffsetAccessorT GlobalOffset;
901 
902  // Accessor to the flush buffer. Each work item writes its
903  // output to a designated section of the flush buffer.
904  mutable detail::GlobalBufAccessorT GlobalFlushBuf;
905 
906  // Offset of the WI's flush buffer in the pool.
907  mutable unsigned WIOffset = 0;
908 
909  mutable size_t FlushBufferSize;
910 
911  // Fields and methods to work with manipulators
912  mutable stream_manipulator Manipulator = defaultfloat;
913 
914  // Type used for format flags
915  using FmtFlags = unsigned int;
916 
917  mutable int Precision = -1;
918  mutable int Width = -1;
919  mutable FmtFlags Flags = 0x0;
920 
921  void set_flag(FmtFlags FormatFlag) const { Flags |= FormatFlag; }
922 
923  void unset_flag(FmtFlags FormatFlag) const { Flags &= ~FormatFlag; }
924 
925  FmtFlags get_flags() const { return Flags; }
926 
927  // This method is used to set the flag for base and float manipulators. These
928  // flags are mutually exclusive and base/float field needs to be cleared
929  // before the setting new flag.
930  void set_flag(FmtFlags FormatFlag, FmtFlags Mask) const {
931  unset_flag(Mask);
932  Flags |= FormatFlag & Mask;
933  }
934 
935  // Set the flags which correspond to the input stream manipulator.
936  void set_manipulator(const stream_manipulator SM) const {
937  switch (SM) {
939  set_flag(detail::Dec, detail::BaseField);
940  break;
942  set_flag(detail::Hex, detail::BaseField);
943  break;
945  set_flag(detail::Oct, detail::BaseField);
946  break;
948  unset_flag(detail::ShowBase);
949  break;
951  set_flag(detail::ShowBase);
952  break;
954  unset_flag(detail::ShowPos);
955  break;
957  set_flag(detail::ShowPos);
958  break;
961  break;
964  break;
967  break;
969  unset_flag(detail::FloatField);
970  break;
971  default:
972  // Unknown manipulator
973  break;
974  }
975  }
976 
977 #ifdef __SYCL_DEVICE_ONLY__
978  void __init(detail::GlobalBufPtrType GlobalBufPtr,
979  range<detail::GlobalBufDim> GlobalBufAccRange,
980  range<detail::GlobalBufDim> GlobalBufMemRange,
981  id<detail::GlobalBufDim> GlobalBufId,
982  detail::GlobalOffsetPtrType GlobalOffsetPtr,
983  range<detail::GlobalOffsetDim> GlobalOffsetAccRange,
984  range<detail::GlobalOffsetDim> GlobalOffsetMemRange,
985  id<detail::GlobalOffsetDim> GlobalOffsetId,
986  detail::GlobalBufPtrType GlobalFlushPtr,
987  range<detail::GlobalBufDim> GlobalFlushAccRange,
988  range<detail::GlobalBufDim> GlobalFlushMemRange,
989  id<detail::GlobalBufDim> GlobalFlushId, size_t _FlushBufferSize) {
990  GlobalBuf.__init(GlobalBufPtr, GlobalBufAccRange, GlobalBufMemRange,
991  GlobalBufId);
992  GlobalOffset.__init(GlobalOffsetPtr, GlobalOffsetAccRange,
993  GlobalOffsetMemRange, GlobalOffsetId);
994  GlobalFlushBuf.__init(GlobalFlushPtr, GlobalFlushAccRange,
995  GlobalFlushMemRange, GlobalFlushId);
996  FlushBufferSize = _FlushBufferSize;
997  // Calculate offset in the flush buffer for each work item in the global
998  // work space. We need to avoid calling intrinsics to get global id because
999  // when stream is used in a single_task kernel this could cause some
1000  // overhead on FPGA target. That is why use global atomic variable to
1001  // calculate offsets.
1002  WIOffset = GlobalOffset[1].fetch_add(FlushBufferSize);
1003 
1004  // Initialize flush subbuffer's offset for each work item on device.
1005  // Initialization on host device is performed via submition of additional
1006  // host task.
1007  SetFlushBufOffset(GlobalFlushBuf, WIOffset, 0);
1008  }
1009 
1010  void __finalize() {
1011  // Flush data to global buffer if flush buffer is not empty. This could be
1012  // necessary if user hasn't yet flushed data on its own and kernel execution
1013  // is finished
1014  // NOTE: A call to this function will be generated by compiler
1015  // NOTE: In the current implementation user should explicitly flush data on
1016  // the host device. Data is not flushed automatically after kernel execution
1017  // because of the missing feature in scheduler.
1018  flushBuffer(GlobalOffset, GlobalBuf, GlobalFlushBuf, WIOffset);
1019  }
1020 #endif
1021 
1022  friend class handler;
1023 
1024  template <typename SYCLObjT> friend class ext::oneapi::weak_object;
1025 
1026  friend const stream &operator<<(const stream &, const char);
1027  friend const stream &operator<<(const stream &, const char *);
1028  template <typename ValueType>
1029  friend
1030  typename std::enable_if_t<std::is_integral_v<ValueType>, const stream &>
1031  operator<<(const stream &, const ValueType &);
1032  friend const stream &operator<<(const stream &, const float &);
1033  friend const stream &operator<<(const stream &, const double &);
1034  friend const stream &operator<<(const stream &, const half &);
1035 
1036  friend const stream &operator<<(const stream &, const stream_manipulator);
1037 
1038  friend const stream &operator<<(const stream &Out,
1039  const __precision_manipulator__ &RHS);
1040 
1041  friend const stream &operator<<(const stream &Out,
1042  const __width_manipulator__ &RHS);
1043  template <typename T, int Dimensions>
1044  friend const stream &operator<<(const stream &Out,
1045  const vec<T, Dimensions> &RHS);
1046  template <typename T>
1047  friend const stream &operator<<(const stream &Out, const T *RHS);
1048  template <int Dimensions>
1049  friend const stream &operator<<(const stream &Out, const id<Dimensions> &RHS);
1050 
1051  template <int Dimensions>
1052  friend const stream &operator<<(const stream &Out,
1053  const range<Dimensions> &RHS);
1054 
1055  template <int Dimensions>
1056  friend const stream &operator<<(const stream &Out,
1057  const item<Dimensions> &RHS);
1058 
1059  template <int Dimensions>
1060  friend const stream &operator<<(const stream &Out,
1061  const nd_range<Dimensions> &RHS);
1062 
1063  template <int Dimensions>
1064  friend const stream &operator<<(const stream &Out,
1065  const nd_item<Dimensions> &RHS);
1066 
1067  template <int Dimensions>
1068  friend const stream &operator<<(const stream &Out,
1069  const group<Dimensions> &RHS);
1070 
1071  template <int Dimensions>
1072  friend const stream &operator<<(const stream &Out,
1073  const h_item<Dimensions> &RHS);
1074 };
1075 
1076 #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
1077 // Byte (has to be converted to a numeric value)
1078 template <typename T>
1079 inline std::enable_if_t<std::is_same_v<T, std::byte>, const stream &>
1080 operator<<(const stream &, const T &) {
1081  static_assert(std::is_integral<T>(),
1082  "Convert the byte to a numeric value using std::to_integer");
1083 }
1084 #endif
1085 
1086 // Character
1087 inline const stream &operator<<(const stream &Out, const char C) {
1088  detail::write(Out.GlobalFlushBuf, Out.FlushBufferSize, Out.WIOffset, &C, 1);
1089  return Out;
1090 }
1091 
1092 // String
1093 inline const stream &operator<<(const stream &Out, const char *Str) {
1094  unsigned Len = 0;
1095  for (; Str[Len] != '\0'; Len++)
1096  ;
1097 
1098  detail::write(Out.GlobalFlushBuf, Out.FlushBufferSize, Out.WIOffset, Str,
1099  Len);
1100  return Out;
1101 }
1102 
1103 // Boolean
1104 inline const stream &operator<<(const stream &Out, const bool &RHS) {
1105  Out << (RHS ? "true" : "false");
1106  return Out;
1107 }
1108 
1109 // Integral
1110 template <typename ValueType>
1111 inline typename std::enable_if_t<std::is_integral_v<ValueType>, const stream &>
1112 operator<<(const stream &Out, const ValueType &RHS) {
1113  detail::writeIntegral(Out.GlobalFlushBuf, Out.FlushBufferSize, Out.WIOffset,
1114  Out.get_flags(), Out.get_width(), RHS);
1115  return Out;
1116 }
1117 
1118 // Floating points
1119 
1120 inline const stream &operator<<(const stream &Out, const float &RHS) {
1121  detail::writeFloatingPoint<float>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1122  Out.WIOffset, Out.get_flags(),
1123  Out.get_width(), Out.get_precision(), RHS);
1124  return Out;
1125 }
1126 
1127 inline const stream &operator<<(const stream &Out, const double &RHS) {
1128  detail::writeFloatingPoint<double>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1129  Out.WIOffset, Out.get_flags(),
1130  Out.get_width(), Out.get_precision(), RHS);
1131  return Out;
1132 }
1133 
1134 inline const stream &operator<<(const stream &Out, const half &RHS) {
1135  detail::writeFloatingPoint<half>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1136  Out.WIOffset, Out.get_flags(),
1137  Out.get_width(), Out.get_precision(), RHS);
1138  return Out;
1139 }
1140 
1141 // Pointer
1142 
1143 template <typename ElementType, access::address_space Space,
1144  access::decorated IsDecorated>
1145 inline const stream &
1146 operator<<(const stream &Out,
1148  Out << RHS.get();
1149  return Out;
1150 }
1151 
1152 template <typename T>
1153 const stream &operator<<(const stream &Out, const T *RHS) {
1154  detail::FmtFlags Flags = Out.get_flags();
1155  Flags &= ~detail::BaseField;
1156  Flags |= detail::Hex | detail::ShowBase;
1157  detail::writeIntegral(Out.GlobalFlushBuf, Out.FlushBufferSize, Out.WIOffset,
1158  Flags, Out.get_width(), reinterpret_cast<size_t>(RHS));
1159  return Out;
1160 }
1161 
1162 // Manipulators
1163 
1164 inline const stream &operator<<(const stream &Out,
1165  const __precision_manipulator__ &RHS) {
1166  Out.Precision = RHS.precision();
1167  return Out;
1168 }
1169 
1170 inline const stream &operator<<(const stream &Out,
1171  const __width_manipulator__ &RHS) {
1172  Out.Width = RHS.width();
1173  return Out;
1174 }
1175 
1176 inline const stream &operator<<(const stream &Out,
1177  const stream_manipulator RHS) {
1178  switch (RHS) {
1180  Out << '\n';
1181  flushBuffer(Out.GlobalOffset, Out.GlobalBuf, Out.GlobalFlushBuf,
1182  Out.WIOffset);
1183  break;
1185  flushBuffer(Out.GlobalOffset, Out.GlobalBuf, Out.GlobalFlushBuf,
1186  Out.WIOffset);
1187  break;
1188  default:
1189  Out.set_manipulator(RHS);
1190  break;
1191  }
1192  return Out;
1193 }
1194 
1195 // Vec
1196 
1197 template <typename T, int VectorLength>
1198 const stream &operator<<(const stream &Out, const vec<T, VectorLength> &RHS) {
1199  detail::writeVec<T, VectorLength>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1200  Out.WIOffset, Out.get_flags(),
1201  Out.get_width(), Out.get_precision(), RHS);
1202  return Out;
1203 }
1204 
1205 // SYCL types
1206 
1207 template <int Dimensions>
1208 inline const stream &operator<<(const stream &Out, const id<Dimensions> &RHS) {
1209  detail::writeArray<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1210  Out.WIOffset, RHS);
1211  return Out;
1212 }
1213 
1214 template <int Dimensions>
1215 inline const stream &operator<<(const stream &Out,
1216  const range<Dimensions> &RHS) {
1217  detail::writeArray<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1218  Out.WIOffset, RHS);
1219  return Out;
1220 }
1221 
1222 template <int Dimensions>
1223 inline const stream &operator<<(const stream &Out,
1224  const item<Dimensions> &RHS) {
1225  detail::writeItem<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1226  Out.WIOffset, RHS);
1227  return Out;
1228 }
1229 
1230 template <int Dimensions>
1231 inline const stream &operator<<(const stream &Out,
1232  const nd_range<Dimensions> &RHS) {
1233  detail::writeNDRange<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1234  Out.WIOffset, RHS);
1235  return Out;
1236 }
1237 
1238 template <int Dimensions>
1239 inline const stream &operator<<(const stream &Out,
1240  const nd_item<Dimensions> &RHS) {
1241  detail::writeNDItem<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1242  Out.WIOffset, RHS);
1243  return Out;
1244 }
1245 
1246 template <int Dimensions>
1247 inline const stream &operator<<(const stream &Out,
1248  const group<Dimensions> &RHS) {
1249  detail::writeGroup<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1250  Out.WIOffset, RHS);
1251  return Out;
1252 }
1253 
1254 template <int Dimensions>
1255 inline const stream &operator<<(const stream &Out,
1256  const h_item<Dimensions> &RHS) {
1257  detail::writeHItem<Dimensions>(Out.GlobalFlushBuf, Out.FlushBufferSize,
1258  Out.WIOffset, RHS);
1259  return Out;
1260 }
1261 
1262 template <typename T, typename RT = detail::EnableIfSwizzleVec<T>>
1263 inline const stream &operator<<(const stream &Out, const T &RHS) {
1264  RT V = RHS;
1265  Out << V;
1266  return Out;
1267 }
1268 
1269 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
1270 } // namespace sycl
1271 namespace std {
1272 template <> struct hash<sycl::stream> {
1273  size_t operator()(const sycl::stream &S) const {
1274 #ifdef __SYCL_DEVICE_ONLY__
1275  (void)S;
1276  return 0;
1277 #else
1278  return hash<std::shared_ptr<sycl::detail::stream_impl>>()(
1280 #endif
1281  }
1282 };
1283 } // namespace std
sycl::_V1::detail::writeGroup
void writeGroup(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const group< Dimensions > &Group)
Definition: stream.hpp:653
sycl::_V1::__width_manipulator__
Definition: stream.hpp:784
sycl::_V1::detail::IsSwizzleOp< sycl::detail::SwizzleOp< VecT, OperationLeftT, OperationRightT, OperationCurrentT, Indexes... > >::Type
typename sycl::vec< T,(sizeof...(Indexes))> Type
Definition: stream.hpp:718
sycl::_V1::detail::MAX_ARRAY_SIZE
constexpr size_t MAX_ARRAY_SIZE
Definition: stream.hpp:55
sycl::_V1::__width_manipulator__::width
int width() const
Definition: stream.hpp:790
builtins.hpp
sycl::_V1::detail::checkForInfNan
std::enable_if_t< std::is_same_v< T, half >, unsigned > checkForInfNan(char *Buf, T Val)
Definition: stream.hpp:311
sycl::_V1::endl
constexpr stream_manipulator endl
Definition: stream.hpp:758
sycl::_V1::detail::reverseBuf
void reverseBuf(char *Buf, unsigned Len)
Definition: stream.hpp:132
std::hash< sycl::stream >::operator()
size_t operator()(const sycl::stream &S) const
Definition: stream.hpp:1273
sycl::_V1::instead
std::uint8_t instead
Definition: aliases.hpp:95
sycl::_V1::detail::ScalarToStr
EnableIfFP< T, unsigned > ScalarToStr(const T &Val, char *Buf, unsigned Flags, int, int Precision=-1)
Definition: stream.hpp:443
sycl::_V1::defaultfloat
constexpr stream_manipulator defaultfloat
Definition: stream.hpp:768
T
sycl::_V1::h_item::get_physical_local
item< dimensions, false > get_physical_local() const
Definition: h_item.hpp:43
sycl::_V1::detail::updateOffset
bool updateOffset(GlobalOffsetAccessorT &GlobalOffset, GlobalBufAccessorT &GlobalBuf, unsigned Size, unsigned &Cur)
Definition: stream.hpp:495
sycl::_V1::nd_item::get_global_id
id< dimensions > get_global_id() const
Definition: nd_item.hpp:44
sycl::_V1::__precision_manipulator__::precision
int precision() const
Definition: stream.hpp:778
sycl::_V1::detail::EnableIfFP
typename std::enable_if_t< std::is_same_v< F, float >||std::is_same_v< F, double >||std::is_same_v< F, half >, T > EnableIfFP
Definition: stream.hpp:66
sycl::_V1::__precision_manipulator__::__precision_manipulator__
__precision_manipulator__(int Precision)
Definition: stream.hpp:776
sycl::_V1::detail::isFastMathInf
detail::enable_if_t< std::is_same< T, float >::value||std::is_same< T, double >::value, bool > isFastMathInf(T Val)
Definition: stream.hpp:250
sycl::_V1::detail::ItemToStr
unsigned ItemToStr(char *Buf, const item< Dimensions, false > &Item)
Definition: stream.hpp:678
__SYCL_INLINE_VER_NAMESPACE
#define __SYCL_INLINE_VER_NAMESPACE(X)
Definition: defines_elementary.hpp:11
sycl::_V1::detail::integralToBase
std::enable_if_t< std::is_integral_v< T >, unsigned > integralToBase(T Val, int Base, char *Digits)
Definition: stream.hpp:159
sycl::_V1::detail::SwizzleOp
Definition: types.hpp:136
sycl::_V1::detail::MAX_ITEM_SIZE
constexpr size_t MAX_ITEM_SIZE
Definition: stream.hpp:675
sycl::_V1::showpos
constexpr stream_manipulator showpos
Definition: stream.hpp:756
sycl::_V1::detail::EnableIfSwizzleVec
typename std::enable_if_t< IsSwizzleOp< T >::value, typename IsSwizzleOp< T >::Type > EnableIfSwizzleVec
Definition: stream.hpp:724
sycl::_V1::detail::digitToChar
char digitToChar(const int Digit)
Definition: stream.hpp:149
sycl::_V1::item::get_id
id< dimensions > get_id() const
Definition: item.hpp:53
sycl
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14
sycl::_V1::detail::GlobalOffsetAS
constexpr static access::address_space GlobalOffsetAS
Definition: stream.hpp:82
sycl::_V1::hexfloat
constexpr stream_manipulator hexfloat
Definition: stream.hpp:766
sycl::_V1::detail::writeIntegral
std::enable_if_t< std::is_integral_v< T > > writeIntegral(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, unsigned Flags, int Width, const T &Val)
Definition: stream.hpp:469
__SYCL_SPECIAL_CLASS
#define __SYCL_SPECIAL_CLASS
Definition: defines.hpp:30
operator==
bool operator==(const Slab &Lhs, const Slab &Rhs)
Definition: usm_allocator.cpp:313
sycl::_V1::id
A unique identifier of an item in an index space.
Definition: array.hpp:17
sycl::_V1::detail::writeVec
void writeVec(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, unsigned Flags, int Width, int Precision, const vec< T, VecLength > &Vec)
Definition: stream.hpp:560
owner_less_base.hpp
sycl::_V1::detail::FmtFlags
unsigned int FmtFlags
Definition: stream.hpp:23
sycl::_V1::detail::GlobalBufDim
constexpr static int GlobalBufDim
Definition: stream.hpp:76
sycl::_V1::detail::writeFloatingPoint
EnableIfFP< T > writeFloatingPoint(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, unsigned Flags, int Width, int Precision, const T &Val)
Definition: stream.hpp:481
sycl::_V1::detail::Dec
static constexpr FmtFlags Dec
Definition: stream.hpp:27
sycl::_V1::nd_range::get_offset
id< dimensions > get_offset() const
Definition: nd_range.hpp:47
sycl::_V1::detail::GlobalBufAccessorT
accessor< char, 1, sycl::access::mode::read_write, sycl::access::target::global_buffer, sycl::access::placeholder::false_t > GlobalBufAccessorT
Definition: stream.hpp:70
sycl::_V1::multi_ptr::get
pointer get() const
Definition: multi_ptr.hpp:244
sycl::_V1::detail::MAX_INTEGRAL_DIGITS
constexpr size_t MAX_INTEGRAL_DIGITS
Definition: stream.hpp:46
sycl::_V1::detail::floatingPointToDecStr
EnableIfFP< T, unsigned > floatingPointToDecStr(T AbsVal, char *Digits, int Precision, bool IsSci)
Definition: stream.hpp:329
sycl::_V1::item::get_offset
std::enable_if_t< has_offset, id< dimensions > > get_offset() const
Definition: item.hpp:79
sycl::_V1::flush
constexpr stream_manipulator flush
Definition: stream.hpp:760
sycl::_V1::detail::write
void write(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const char *Str, unsigned Len, unsigned Padding=0)
Definition: stream.hpp:110
sycl::_V1::range< Dimensions >
sycl::_V1::detail::writeHItem
void writeHItem(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const h_item< Dimensions > &HItem)
Definition: stream.hpp:690
export.hpp
sycl::_V1::detail::VEC_OPEN_BRACE
constexpr char VEC_OPEN_BRACE
Definition: stream.hpp:48
sycl::_V1::ext::oneapi::experimental::isnan
std::enable_if_t< std::is_same_v< T, bfloat16 >, bool > isnan(T x)
Definition: bfloat16_math.hpp:36
sycl::_V1::multi_ptr
Provides constructors for address space qualified and non address space qualified pointers to allow i...
Definition: atomic.hpp:34
sycl::_V1::nd_item
Identifies an instance of the function object executing at each point in an nd_range.
Definition: helpers.hpp:33
sycl::_V1::operator<<
const stream & operator<<(const stream &Out, const T &RHS)
Definition: stream.hpp:1263
sycl::_V1::ext::oneapi::experimental::has_property
static constexpr bool has_property()
Definition: annotated_arg.hpp:162
sycl::_V1::__width_manipulator__::__width_manipulator__
__width_manipulator__(int Width)
Definition: stream.hpp:788
sycl::_V1::detail::append
unsigned append(char *Dst, const char *Src)
Definition: stream.hpp:219
sycl::_V1::detail::GlobalBufAS
constexpr static access::address_space GlobalBufAS
Definition: stream.hpp:72
sycl::_V1::detail::SetFlushBufOffset
void SetFlushBufOffset(GlobalBufAccessorT &GlobalFlushBuf, unsigned WIOffset, unsigned Offset)
Definition: stream.hpp:104
sycl::_V1::item
Identifies an instance of the function object executing at each point in a range.
Definition: helpers.hpp:29
sycl::_V1::detail::MAX_DIMENSIONS
constexpr size_t MAX_DIMENSIONS
Definition: stream.hpp:51
sycl::_V1::nd_range::get_global_range
range< dimensions > get_global_range() const
Definition: nd_range.hpp:40
sycl::_V1::oct
constexpr stream_manipulator oct
Definition: stream.hpp:748
sycl::_V1::detail::writeArray
void writeArray(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const array< ArrayLength > &Arr)
Definition: stream.hpp:589
sycl::_V1::detail::IsSwizzleOp
Definition: stream.hpp:710
sycl::_V1::detail::MAX_FLOATING_POINT_DIGITS
constexpr size_t MAX_FLOATING_POINT_DIGITS
Definition: stream.hpp:45
sycl::_V1::detail::ArrayToStr
unsigned ArrayToStr(char *Buf, const array< ArrayLength > &Arr)
Definition: stream.hpp:573
defines.hpp
sycl::_V1::detail::VEC_ELEMENT_DELIMITER
constexpr const char * VEC_ELEMENT_DELIMITER
Definition: stream.hpp:47
sycl::_V1::noshowpos
constexpr stream_manipulator noshowpos
Definition: stream.hpp:754
sycl::_V1::detail::writeNDItem
void writeNDItem(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const nd_item< Dimensions > &ND_Item)
Definition: stream.hpp:636
sycl::_V1::detail::TargetToAS
Definition: access.hpp:131
sycl::_V1::nd_range
Defines the iteration domain of both the work-groups and the overall dispatch.
Definition: uniform.hpp:36
sycl::_V1::nd_item::get_local_id
id< dimensions > get_local_id() const
Definition: nd_item.hpp:58
sycl::_V1::item::get_range
range< dimensions > get_range() const
Definition: item.hpp:67
sycl::_V1::h_item
Identifies an instance of a group::parallel_for_work_item function object executing at each point in ...
Definition: helpers.hpp:34
sycl::_V1::read_write
constexpr mode_tag_t< access_mode::read_write > read_write
Definition: access.hpp:76
sycl::_V1::half
sycl::detail::half_impl::half half
Definition: aliases.hpp:103
sycl::_V1::detail::FLUSH_BUF_OFFSET_SIZE
constexpr unsigned FLUSH_BUF_OFFSET_SIZE
Definition: stream.hpp:60
sycl::_V1::ext::oneapi::experimental::get_property
static constexpr auto get_property()
Definition: annotated_arg.hpp:166
sycl::_V1::detail::GlobalOffsetDim
constexpr static int GlobalOffsetDim
Definition: stream.hpp:86
sycl::_V1::access::decorated
decorated
Definition: access.hpp:59
sycl::_V1::nd_range::get_local_range
range< dimensions > get_local_range() const
Definition: nd_range.hpp:42
sycl::_V1::detail::GlobalOffsetPtrType
typename detail::DecoratedType< unsigned, GlobalBufAS >::type * GlobalOffsetPtrType
Definition: stream.hpp:85
sycl::_V1::accessor
Definition: accessor.hpp:225
sycl::_V1::dec
constexpr stream_manipulator dec
Definition: stream.hpp:744
sycl::_V1::detail::F2I32
static unsigned F2I32(float Val)
Definition: stream.hpp:229
sycl::_V1::__precision_manipulator__
Definition: stream.hpp:772
sycl::_V1::detail::ShowPos
static constexpr FmtFlags ShowPos
Definition: stream.hpp:31
sycl::_V1::setw
__width_manipulator__ setw(int Width)
Definition: stream.hpp:800
sycl::_V1::setprecision
__precision_manipulator__ setprecision(int Precision)
Definition: stream.hpp:796
sycl::_V1::detail::Hex
static constexpr FmtFlags Hex
Definition: stream.hpp:28
sycl::_V1::detail::GlobalOffsetAccessorT
accessor< unsigned, 1, sycl::access::mode::atomic, sycl::access::target::global_buffer, sycl::access::placeholder::false_t > GlobalOffsetAccessorT
Definition: stream.hpp:80
sycl::_V1::detail::isFastMathNan
detail::enable_if_t< std::is_same< T, float >::value||std::is_same< T, double >::value, bool > isFastMathNan(T Val)
Definition: stream.hpp:263
handler.hpp
sycl::_V1::detail::getAbsVal
std::make_unsigned_t< T > getAbsVal(const T Val, const int Base)
Definition: stream.hpp:145
sycl::_V1::vec
Provides a cross-patform vector class template that works efficiently on SYCL devices as well as in h...
Definition: aliases.hpp:20
sycl::_V1::detail::Scientific
static constexpr FmtFlags Scientific
Definition: stream.hpp:33
sycl::_V1::detail::VEC_CLOSE_BRACE
constexpr char VEC_CLOSE_BRACE
Definition: stream.hpp:49
std
Definition: accessor.hpp:3910
sycl::_V1::detail::FloatField
static constexpr FmtFlags FloatField
Definition: stream.hpp:43
sycl::_V1::hex
constexpr stream_manipulator hex
Definition: stream.hpp:746
sycl::_V1::detail::writeNDRange
void writeNDRange(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const nd_range< Dimensions > &ND_Range)
Definition: stream.hpp:617
sycl::_V1::showbase
constexpr stream_manipulator showbase
Definition: stream.hpp:752
__SYCL_TYPE
#define __SYCL_TYPE(x)
Definition: defines.hpp:40
sycl::_V1::ext::oneapi::experimental::matrix::use
use
Definition: matrix-unified-utils.hpp:17
sycl::_V1::scientific
constexpr stream_manipulator scientific
Definition: stream.hpp:764
sycl::_V1::detail::DecoratedType
Definition: access.hpp:154
sycl::_V1::h_item::get_global
item< dimensions, false > get_global() const
Definition: h_item.hpp:37
sycl::_V1::h_item::get_logical_local
item< dimensions, false > get_logical_local() const
Definition: h_item.hpp:41
sycl::_V1::detail::flushBuffer
void flushBuffer(GlobalOffsetAccessorT &GlobalOffset, GlobalBufAccessorT &GlobalBuf, GlobalBufAccessorT &GlobalFlushBuf, unsigned WIOffset)
Definition: stream.hpp:509
sycl::_V1::detail::BaseField
static constexpr FmtFlags BaseField
Definition: stream.hpp:38
sycl::_V1::detail::D2I64
static unsigned long long D2I64(double Val)
Definition: stream.hpp:238
sycl::_V1::ext::intel::experimental::operator!=
bool operator!=(const cache_config &lhs, const cache_config &rhs)
Definition: kernel_execution_properties.hpp:37
sycl::_V1::detail::writeItem
void writeItem(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const item< Dimensions > &Item)
Definition: stream.hpp:598
sycl::_V1::detail::GetFlushBufOffset
unsigned GetFlushBufOffset(const GlobalBufAccessorT &GlobalFlushBuf, unsigned WIOffset)
Definition: stream.hpp:92
sycl::_V1::stream_manipulator
stream_manipulator
Definition: stream.hpp:728
weak_object_base.hpp
sycl::_V1::detail::GlobalBufPtrType
typename detail::DecoratedType< char, GlobalBufAS >::type * GlobalBufPtrType
Definition: stream.hpp:75
sycl::_V1::detail::array
Definition: array.hpp:21
sycl::_V1::detail::Oct
static constexpr FmtFlags Oct
Definition: stream.hpp:29
sycl::_V1::group
Definition: helpers.hpp:30
sycl::_V1::detail::VecToStr
std::enable_if_t<(VecLength==3), unsigned > VecToStr(const vec< T, VecLength > &Vec, char *VecStr, unsigned Flags, int Width, int Precision)
Definition: stream.hpp:551
sycl::_V1::detail::ShowBase
static constexpr FmtFlags ShowBase
Definition: stream.hpp:30
sycl::_V1::Space
Space
Definition: multi_ptr.hpp:1307
__SYCL2020_DEPRECATED
#define __SYCL2020_DEPRECATED(message)
Definition: defines_elementary.hpp:57
sycl::_V1::detail::getSyclObjImpl
decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject)
Definition: common.hpp:302
sycl::_V1::detail::Fixed
static constexpr FmtFlags Fixed
Definition: stream.hpp:32
sycl::_V1::noshowbase
constexpr stream_manipulator noshowbase
Definition: stream.hpp:750
sycl::_V1::detail::isFastMathSignBit
detail::enable_if_t< std::is_same< T, float >::value||std::is_same< T, double >::value, bool > isFastMathSignBit(T Val)
Definition: stream.hpp:276
sycl::_V1::fixed
constexpr stream_manipulator fixed
Definition: stream.hpp:762
sycl::_V1::access::address_space
address_space
Definition: access.hpp:47