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