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