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