DPC++ Runtime
Runtime libraries for oneAPI DPC++
memory_intrin.hpp
Go to the documentation of this file.
1 //==------------ memory_intrin.hpp - DPC++ Explicit SIMD API ---------------==//
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 // Declares experimental memory Explicit SIMD intrinsics.
9 //===----------------------------------------------------------------------===//
10 
12 
13 #pragma once
14 
16 
17 // generic work-group split barrier
18 __ESIMD_INTRIN void __esimd_sbarrier(__ESIMD_ENS::split_barrier_action flag)
19 #ifdef __SYCL_DEVICE_ONLY__
20  ;
21 #else
22 {
23  sycl::detail::getESIMDDeviceInterface()->cm_sbarrier_ptr((uint32_t)flag);
24 }
25 #endif // __SYCL_DEVICE_ONLY__
26 
27 // \brief Raw sends load.
28 //
29 // @param modifier the send message flags (Bit-0: isSendc, Bit-1: isEOT).
30 //
31 // @param execSize the execution size, which must be a compile time constant.
32 //
33 // @param pred the predicate to specify enabled channels.
34 //
35 // @param numSrc0 the number of GRFs for source-0, which must be a compile time
36 // constant.
37 //
38 // @param numSrc1 the number of GRFs for source-1, which must be a compile time
39 // constant.
40 //
41 // @param numDst the number of GRFs for destination, which must be a compile
42 // time constant.
43 //
44 // @param sfid the shared function ID, which must be a compile time constant.
45 //
46 // @param exDesc the extended message descriptor.
47 //
48 // @param msgDesc the message descriptor.
49 //
50 // @param msgSrc0 the first source operand of send message.
51 //
52 // @param msgSrc1 the second source operand of send message.
53 //
54 // @param msgDst the destination operand of send message.
55 //
56 // Returns a simd vector of type Ty1 and size N1.
57 //
58 template <typename Ty1, int N1, typename Ty2, int N2, typename Ty3, int N3,
59  int N = 16>
60 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<Ty1, N1>
61 __esimd_raw_sends2(uint8_t modifier, uint8_t execSize,
62  __ESIMD_DNS::simd_mask_storage_t<N> pred, uint8_t numSrc0,
63  uint8_t numSrc1, uint8_t numDst, uint8_t sfid,
64  uint32_t exDesc, uint32_t msgDesc,
65  __ESIMD_DNS::vector_type_t<Ty2, N2> msgSrc0,
66  __ESIMD_DNS::vector_type_t<Ty3, N3> msgSrc1,
67  __ESIMD_DNS::vector_type_t<Ty1, N1> msgDst)
68 #ifdef __SYCL_DEVICE_ONLY__
69  ;
70 #else
71 {
72  __ESIMD_UNSUPPORTED_ON_HOST;
73 }
74 #endif // __SYCL_DEVICE_ONLY__
75 
76 // \brief Raw send load.
77 //
78 // @param modifier the send message flags (Bit-0: isSendc, Bit-1: isEOT).
79 //
80 // @param execSize the execution size, which must be a compile time constant.
81 //
82 // @param pred the predicate to specify enabled channels.
83 //
84 // @param numSrc0 the number of GRFs for source-0, which must be a compile time
85 // constant.
86 //
87 // @param numDst the number of GRFs for destination, which must be a compile
88 // time constant.
89 //
90 // @param sfid the shared function ID, which must be a compile time constant.
91 //
92 // @param exDesc the extended message descriptor.
93 //
94 // @param msgDesc the message descriptor.
95 //
96 // @param msgSrc0 the first source operand of send message.
97 //
98 // @param msgDst the destination operand of send message.
99 //
100 // Returns a simd vector of type Ty1 and size N1.
101 //
102 template <typename Ty1, int N1, typename Ty2, int N2, int N = 16>
103 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<Ty1, N1>
104 __esimd_raw_send2(uint8_t modifier, uint8_t execSize,
105  __ESIMD_DNS::simd_mask_storage_t<N> pred, uint8_t numSrc0,
106  uint8_t numDst, uint8_t sfid, uint32_t exDesc,
107  uint32_t msgDesc, __ESIMD_DNS::vector_type_t<Ty2, N2> msgSrc0,
108  __ESIMD_DNS::vector_type_t<Ty1, N1> msgDst)
109 #ifdef __SYCL_DEVICE_ONLY__
110  ;
111 #else
112 {
113  __ESIMD_UNSUPPORTED_ON_HOST;
114 }
115 #endif // __SYCL_DEVICE_ONLY__
116 
117 // \brief Raw sends store.
118 //
119 // @param modifier the send message flags (Bit-0: isSendc, Bit-1: isEOT).
120 //
121 // @param execSize the execution size, which must be a compile time constant.
122 //
123 // @param pred the predicate to specify enabled channels.
124 //
125 // @param numSrc0 the number of GRFs for source-0, which must be a compile time
126 // constant.
127 //
128 // @param numSrc1 the number of GRFs for source-1, which must be a compile time
129 // constant.
130 //
131 // @param sfid the shared function ID, which must be a compile time constant.
132 //
133 // @param exDesc the extended message descriptor.
134 //
135 // @param msgDesc the message descriptor.
136 //
137 // @param msgSrc0 the first source operand of send message.
138 //
139 // @param msgSrc1 the second source operand of send message.
140 //
141 template <typename Ty1, int N1, typename Ty2, int N2, int N = 16>
142 __ESIMD_INTRIN void
143 __esimd_raw_sends2_noresult(uint8_t modifier, uint8_t execSize,
144  __ESIMD_DNS::simd_mask_storage_t<N> pred,
145  uint8_t numSrc0, uint8_t numSrc1, uint8_t sfid,
146  uint32_t exDesc, uint32_t msgDesc,
147  __ESIMD_DNS::vector_type_t<Ty1, N1> msgSrc0,
148  __ESIMD_DNS::vector_type_t<Ty2, N2> msgSrc1)
149 #ifdef __SYCL_DEVICE_ONLY__
150  ;
151 #else
152 {
153  __ESIMD_UNSUPPORTED_ON_HOST;
154 }
155 #endif // __SYCL_DEVICE_ONLY__
156 
157 // \brief Raw send store.
158 //
159 // @param modifier the send message flags (Bit-0: isSendc, Bit-1: isEOT).
160 //
161 // @param execSize the execution size, which must be a compile time constant.
162 //
163 // @param pred the predicate to specify enabled channels.
164 //
165 // @param numSrc0 the number of GRFs for source-0, which must be a compile time
166 // constant.
167 //
168 // @param sfid the shared function ID, which must be a compile time constant.
169 //
170 // @param exDesc the extended message descriptor.
171 //
172 // @param msgDesc the message descriptor.
173 //
174 // @param msgSrc0 the first source operand of send message.
175 //
176 template <typename Ty1, int N1, int N = 16>
177 __ESIMD_INTRIN void
178 __esimd_raw_send2_noresult(uint8_t modifier, uint8_t execSize,
179  __ESIMD_DNS::simd_mask_storage_t<N> pred,
180  uint8_t numSrc0, uint8_t sfid, uint32_t exDesc,
181  uint32_t msgDesc,
182  __ESIMD_DNS::vector_type_t<Ty1, N1> msgSrc0)
183 #ifdef __SYCL_DEVICE_ONLY__
184  ;
185 #else
186 {
187  __ESIMD_UNSUPPORTED_ON_HOST;
188 }
189 #endif // __SYCL_DEVICE_ONLY__
190 
199 __ESIMD_INTRIN void __esimd_nbarrier(uint8_t mode, uint8_t id,
200  uint8_t thread_count)
201 #ifdef __SYCL_DEVICE_ONLY__
202  ;
203 #else // __SYCL_DEVICE_ONLY__
204 {
205  __ESIMD_UNSUPPORTED_ON_HOST;
206 }
207 #endif // __SYCL_DEVICE_ONLY__
208 
213 __ESIMD_INTRIN void __esimd_nbarrier_init(uint8_t count)
214 #ifdef __SYCL_DEVICE_ONLY__
215  ;
216 #else // __SYCL_DEVICE_ONLY__
217 {
218  __ESIMD_UNSUPPORTED_ON_HOST;
219 }
220 #endif // __SYCL_DEVICE_ONLY__
221 
237 template <typename Ty, int N>
238 __ESIMD_INTRIN void __esimd_raw_send_nbarrier_signal(
239  uint32_t is_sendc, uint32_t extended_descriptor, uint32_t descriptor,
240  __ESIMD_DNS::vector_type_t<Ty, N> msg_var, uint16_t pred = 1)
241 #ifdef __SYCL_DEVICE_ONLY__
242  ;
243 #else // __SYCL_DEVICE_ONLY__
244 {
245  __ESIMD_UNSUPPORTED_ON_HOST;
246 }
247 #endif // __SYCL_DEVICE_ONLY__
248 
249 #ifndef __SYCL_DEVICE_ONLY__
250 // Shared utility/helper functions for LSC support under emulation
251 // (ESIMD_EMULATOR backend)
252 
253 // Raw-address increment function for u8u32 and u16u32
254 template <typename Ty, __ESIMD_ENS::lsc_data_size DS>
255 constexpr uint32_t rawAddressIncrement() {
256  if constexpr (DS == __ESIMD_ENS::lsc_data_size::u8u32) {
257  return 1;
258  } else if constexpr (DS == __ESIMD_ENS::lsc_data_size::u16u32) {
259  return 2;
260  } else {
261  return (uint32_t)sizeof(Ty);
262  }
263 }
264 
265 // Vector index increment function for 'Transposed' 2D-surface access
266 template <int N, __ESIMD_EDNS::lsc_data_order _Transposed>
267 constexpr int vectorIndexIncrement() {
268  if constexpr (_Transposed == __ESIMD_EDNS::lsc_data_order::transpose) {
269  return 1;
270  } else {
271  return N;
272  }
273 }
274 
275 // Load/Store align bitmask generator for 1-D vector load/store
276 //
277 // Not only generates address-align bitmask, but also checks
278 // legitimacy of load/store operation with respect to vector size,
279 // data size
284 template <typename Ty, __ESIMD_EDNS::lsc_vector_size VS,
285  __ESIMD_ENS::lsc_data_size DS, int N>
286 constexpr unsigned loadstoreAlignMask() {
287  constexpr __ESIMD_ENS::lsc_data_size _DS =
288  __ESIMD_EDNS::finalize_data_size<Ty, DS>(); // Actual data_size
289 
290  if constexpr (VS == __ESIMD_EDNS::lsc_vector_size::n1) {
291  static_assert(((_DS == __ESIMD_ENS::lsc_data_size::u32) ||
292  (_DS == __ESIMD_ENS::lsc_data_size::u64) ||
293  (_DS == __ESIMD_ENS::lsc_data_size::u8) ||
294  (_DS == __ESIMD_ENS::lsc_data_size::u16) ||
295  (_DS == __ESIMD_ENS::lsc_data_size::u8u32) ||
296  (_DS == __ESIMD_ENS::lsc_data_size::u16u32)) &&
297  "Wrong __ESIMD_EDNS::lsc_data_size for "
298  "__ESIMD_EDNS::lsc_vector_size == 1\n"
299  "(loadstoreAlignMask)");
300  return 0x0;
301  } else if constexpr ((VS == __ESIMD_EDNS::lsc_vector_size::n2) ||
302  (VS == __ESIMD_EDNS::lsc_vector_size::n3) ||
303  (VS == __ESIMD_EDNS::lsc_vector_size::n4) ||
304  (VS == __ESIMD_EDNS::lsc_vector_size::n8)) {
305  static_assert(
306  ((_DS == __ESIMD_ENS::lsc_data_size::u32) ||
307  (_DS == __ESIMD_ENS::lsc_data_size::u64)) &&
308  "Wrong Data Size for __ESIMD_EDNS::lsc_vector_size == 2/3/4/8\n"
309  "(loadstoreAlignMask)");
310  // 0x3 for u32 / 0x7 for u64
311  if constexpr (_DS == __ESIMD_ENS::lsc_data_size::u32)
312  return 0x3;
313  else
314  return 0x7;
315  } else if constexpr ((VS == __ESIMD_EDNS::lsc_vector_size::n16) ||
316  (VS == __ESIMD_EDNS::lsc_vector_size::n32) ||
317  (VS == __ESIMD_EDNS::lsc_vector_size::n64)) {
318  static_assert(
319  (N == 1) &&
320  "Unsupported Size for __ESIMD_EDNS::lsc_vector_size = 16/32/64\n"
321  "(loadstoreAlignMask)");
322  // 0x3 for u32 / 0x7 for u64
323  if constexpr (_DS == __ESIMD_ENS::lsc_data_size::u32)
324  return 0x3;
325  else
326  return 0x7;
327  } else {
328  static_assert((N != N) && "Wrong Vector Size!!");
329  }
330 }
331 
332 // Helper function for loading from indexed-surface and SLM
333 // INT_MAX is for SLM
334 template <typename Ty, uint16_t AddressScale, int ImmOffset,
336  __ESIMD_EDNS::lsc_data_order _Transposed, int N, uint32_t MASK>
337 auto __esimd_emu_lsc_offset_read(
338  __ESIMD_DNS::simd_mask_storage_t<N> Pred,
339  __ESIMD_DNS::vector_type_t<uint32_t, N> Offsets, char *ReadBase,
340  int BufByteWidth = INT_MAX) {
341  // TODO : Support AddressScale, ImmOffset
342  static_assert(AddressScale == 1);
343  static_assert(ImmOffset == 0);
344  static_assert(DS != __ESIMD_ENS::lsc_data_size::u16u32h);
345 
346  __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()> Output = 0;
347 
348  constexpr int ChanlCount = __ESIMD_EDNS::to_int<VS>();
349 
350  for (int OffsetIdx = 0; OffsetIdx < N; OffsetIdx += 1) {
351  if (Pred[OffsetIdx] == 0) {
352  // Skip Output vector elements correpsonding to
353  // predicates whose value is zero
354  continue;
355  }
356 
357  assert(((Offsets[OffsetIdx] & MASK)) == 0 && "Offset Alignment Error!!");
358 
359  // ByteDistance : byte-distance from buffer-read base
360  int ByteDistance = Offsets[OffsetIdx];
361 
362  for (int ChanelIdx = 0, VecIdx = OffsetIdx; ChanelIdx < ChanlCount;
363  ChanelIdx += 1, ByteDistance += rawAddressIncrement<Ty, DS>(),
364  VecIdx += vectorIndexIncrement<N, _Transposed>()) {
365 
366  if ((ByteDistance >= 0) && (ByteDistance < BufByteWidth)) {
367  Output[VecIdx] = *((Ty *)(ReadBase + ByteDistance));
368  }
369  }
370  }
371  return Output;
372 }
373 
374 // Helper function for storing to indexed-surface and SLM. INT_MAX is
375 // for SLM
376 template <typename Ty, uint16_t AddressScale, int ImmOffset,
378  __ESIMD_EDNS::lsc_data_order _Transposed, int N, uint32_t MASK>
379 void __esimd_emu_lsc_offset_write(
380  __ESIMD_DNS::simd_mask_storage_t<N> Pred,
381  __ESIMD_DNS::vector_type_t<uint32_t, N> Offsets,
382  __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()> vals,
383  char *WriteBase, int BufByteWidth = INT_MAX) {
384  // TODO : Support AddressScale, ImmOffset
385  static_assert(AddressScale == 1);
386  static_assert(ImmOffset == 0);
387  static_assert(DS != __ESIMD_ENS::lsc_data_size::u16u32h);
388 
389  using StoreType = typename std::conditional_t<
390  DS == __ESIMD_ENS::lsc_data_size::u8, uint8_t,
392  DS == __ESIMD_ENS::lsc_data_size::u16, uint16_t,
394  DS == __ESIMD_ENS::lsc_data_size::u32, uint32_t,
396  DS == __ESIMD_ENS::lsc_data_size::u64, uint64_t,
398  DS == __ESIMD_ENS::lsc_data_size::u8u32, uint8_t,
399  std::conditional_t<DS ==
400  __ESIMD_ENS::lsc_data_size::u16u32,
401  uint16_t, void>>>>>>;
402 
403  for (int OffsetIdx = 0; OffsetIdx < N; OffsetIdx += 1) {
404  if (Pred[OffsetIdx] == 0) {
405  // Skip input vector elements correpsonding to
406  // predicates whose value is zero
407  continue;
408  }
409 
410  assert(((Offsets[OffsetIdx] & MASK)) == 0 && "Offset Alignment Error!!");
411 
412  // ByteDistance : byte-distance from buffer-write base
413  int ByteDistance = Offsets[OffsetIdx];
414  constexpr int ChanlCount = __ESIMD_EDNS::to_int<VS>();
415 
416  for (int ChanelIdx = 0, VecIdx = OffsetIdx; ChanelIdx < ChanlCount;
417  ChanelIdx += 1, ByteDistance += rawAddressIncrement<Ty, DS>(),
418  VecIdx += vectorIndexIncrement<N, _Transposed>()) {
419 
420  if ((ByteDistance >= 0) && (ByteDistance < BufByteWidth)) {
421  *((StoreType *)(WriteBase + ByteDistance)) = vals[VecIdx];
422  }
423  }
424  }
425 }
426 
429 template <typename Ty, int N>
430 __ESIMD_DNS::vector_type_t<Ty, N>
431 __esimd_emu_read_2d(__ESIMD_DNS::simd_mask_storage_t<N> Pred, uintptr_t Ptr,
432  unsigned SurfaceWidth, unsigned SurfaceHeight,
433  unsigned SurfacePitch, int X, int Y, int Width, int Height,
434  int NBlks, __ESIMD_EDNS::lsc_data_order _Transposed,
435  bool Transformed) {
436  assert(SurfaceHeight >= 0);
437  assert(SurfaceWidth >= 0);
438  assert(SurfaceWidth <= SurfacePitch);
439 
440  SurfaceHeight += 1;
441  SurfaceWidth += 1;
442  SurfacePitch += 1;
443 
444  constexpr unsigned sizeofTy = sizeof(Ty);
445 
446  __ESIMD_DNS::vector_type_t<Ty, N> Output = 0;
447 
448  char *buff = (char *)Ptr;
449  assert(buff != NULL);
450 
451  int vecIdx = 0;
452  int blkCount = 0;
453 
454  for (int xBase = X * sizeofTy; blkCount < NBlks; xBase += sizeofTy * Width) {
455  if (Transformed == true) {
456  constexpr int elems_per_DW = (sizeofTy == 1) ? 4 : 2;
457  int yRead = Y * SurfacePitch;
458  for (int u = 0; u < Height;
459  u += elems_per_DW, yRead += SurfacePitch * elems_per_DW) {
460  vecIdx = u * sycl::detail::getNextPowerOfTwo(Width) +
461  blkCount * Height * sycl::detail::getNextPowerOfTwo(Width);
462  if ((yRead < 0) || (yRead >= SurfacePitch * SurfaceHeight)) {
464  vecIdx += Width * elems_per_DW;
465  continue;
466  }
467 
468  int xRead = xBase;
469  for (int v = 0; v < Width; v += 1, xRead += sizeofTy) {
470  if ((xRead < 0) || (xRead >= SurfaceWidth)) {
472  vecIdx += elems_per_DW;
473  continue;
474  }
475 
476  char *base = buff + xRead;
477  int offset = yRead;
478  for (int k = 0; k < elems_per_DW; k++, vecIdx += 1) {
479  if (Pred[vecIdx] != 0) {
480  if (offset >= 0 && offset < SurfacePitch * SurfaceHeight) {
481  Output[vecIdx] = *((Ty *)(base + offset));
482  }
483  }
484  // Increasing in Y-direction
485  offset += SurfacePitch;
486  } // k loop
487  } // v loop
488  } // u loop
489  } // (Transformed == true)
490  else if (_Transposed == __ESIMD_EDNS::lsc_data_order::transpose) {
491  int xRead = xBase;
492  for (int v = 0; v < Width; v += 1, xRead += sizeofTy) {
493  if ((xRead < 0) || (xRead >= SurfaceWidth)) {
494  // Horizontally out-of-bound, skip corresponding vector elements
495  vecIdx += Height;
496  continue;
497  }
498 
499  int yRead = Y * SurfacePitch;
500  for (int u = 0; u < Height;
501  u += 1, yRead += SurfacePitch, vecIdx += 1) {
502  if (Pred[vecIdx] != 0) {
503  if ((yRead >= 0) && (yRead < SurfacePitch * SurfaceHeight)) {
504  Output[vecIdx] = *((Ty *)(buff + yRead + xRead));
505  }
506  }
507  } // u loop
508  } // v loop
509  } // (_Transposed == __ESIMD_EDNS::lsc_data_order::transpose)
510  else {
511  int yRead = Y * SurfacePitch;
512  for (int u = 0; u < Height; u += 1, yRead += SurfacePitch) {
513  if ((yRead < 0) || (yRead >= SurfacePitch * SurfaceHeight)) {
514  // Vertically Out-of-bound, skip corresponding vector elements
515  vecIdx += Width;
516  continue;
517  }
518 
519  int xRead = xBase;
520  for (int v = 0; v < Width; v += 1, xRead += sizeofTy, vecIdx += 1) {
521  if (Pred[vecIdx] != 0) {
522  if ((xRead >= 0) && (xRead < SurfaceWidth)) {
523  Output[vecIdx] = *((Ty *)(buff + yRead + xRead));
524  }
525  }
526  } // v loop
527  } // u loop
528  } // Linear loading
529  blkCount += 1;
530  vecIdx = blkCount * sycl::detail::getNextPowerOfTwo(Width) * Height;
531  } // xBase loop
532 
533  return Output;
534 }
535 
538 template <typename Ty, int N>
539 void __esimd_emu_write_2d(__ESIMD_DNS::simd_mask_storage_t<N> Pred,
540  uintptr_t Ptr, unsigned SurfaceWidth,
541  unsigned SurfaceHeight, unsigned SurfacePitch, int X,
542  int Y, __ESIMD_DNS::vector_type_t<Ty, N> vals,
543  int Width, int Height) {
544  assert(SurfaceHeight >= 0);
545  assert(SurfaceWidth >= 0);
546  assert(SurfaceWidth <= SurfacePitch);
547 
548  SurfaceHeight += 1;
549  SurfaceWidth += 1;
550  SurfacePitch += 1;
551 
552  constexpr unsigned sizeofTy = sizeof(Ty);
553 
554  char *buff = (char *)Ptr;
555  assert(buff != NULL);
556 
557  int vecIdx = 0;
558  int rowCount = 0;
559  for (int yWrite = Y * SurfacePitch; rowCount < Height;
560  yWrite += SurfacePitch) {
561  if (yWrite == SurfacePitch * SurfaceHeight) {
562  // Vertically Out-of-bound
563  break;
564  }
565  int writeCount = 0;
566  for (int xWrite = X * sizeofTy; writeCount < Width;
567  xWrite += sizeofTy, vecIdx += 1, writeCount += 1) {
568  if (xWrite >= 0 && xWrite < SurfaceWidth && Pred[vecIdx] != 0) {
569  *((Ty *)(buff + yWrite + xWrite)) = vals[vecIdx];
570  }
571  } // xWrite loop
572  rowCount += 1;
573  } // yWrite loop
574 }
575 
576 #endif
577 
596 template <typename Ty, __ESIMD_ENS::cache_hint L1H, __ESIMD_ENS::cache_hint L3H,
597  uint16_t AddressScale, int ImmOffset, __ESIMD_ENS::lsc_data_size DS,
599  __ESIMD_EDNS::lsc_data_order _Transposed, int N>
600 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()>
601 __esimd_lsc_load_slm(__ESIMD_DNS::simd_mask_storage_t<N> pred,
602  __ESIMD_DNS::vector_type_t<uint32_t, N> offsets)
603 #ifdef __SYCL_DEVICE_ONLY__
604  ;
605 #else // __SYCL_DEVICE_ONLY__
606 {
607  sycl::detail::ESIMDDeviceInterface *I =
609 
610  return __esimd_emu_lsc_offset_read<Ty, AddressScale, ImmOffset, DS, VS,
611  _Transposed, N,
612  loadstoreAlignMask<Ty, VS, DS, N>()>(
613  pred, offsets, I->__cm_emu_get_slm_ptr());
614 }
615 #endif // __SYCL_DEVICE_ONLY__
616 
637 template <typename Ty, __ESIMD_ENS::cache_hint L1H, __ESIMD_ENS::cache_hint L3H,
638  uint16_t AddressScale, int ImmOffset, __ESIMD_ENS::lsc_data_size DS,
640  __ESIMD_EDNS::lsc_data_order _Transposed, int N,
641  typename SurfIndAliasTy>
642 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()>
643 __esimd_lsc_load_bti(__ESIMD_DNS::simd_mask_storage_t<N> pred,
644  __ESIMD_DNS::vector_type_t<uint32_t, N> offsets,
645  SurfIndAliasTy surf_ind)
646 #ifdef __SYCL_DEVICE_ONLY__
647  ;
648 #else // __SYCL_DEVICE_ONLY__
649 {
650  char *readBase;
651  uint32_t width;
652  std::mutex *mutexLock;
653 
654  sycl::detail::ESIMDDeviceInterface *I =
656 
657  I->sycl_get_cm_buffer_params_ptr(surf_ind, &readBase, &width, &mutexLock);
658 
659  std::lock_guard<std::mutex> lock(*mutexLock);
660 
661  return __esimd_emu_lsc_offset_read<Ty, AddressScale, ImmOffset, DS, VS,
662  _Transposed, N,
663  loadstoreAlignMask<Ty, VS, DS, N>()>(
664  pred, offsets, readBase, width);
665 }
666 #endif // __SYCL_DEVICE_ONLY__
667 
686 template <typename Ty, __ESIMD_ENS::cache_hint L1H, __ESIMD_ENS::cache_hint L3H,
687  uint16_t AddressScale, int ImmOffset, __ESIMD_ENS::lsc_data_size DS,
689  __ESIMD_EDNS::lsc_data_order _Transposed, int N>
690 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()>
691 __esimd_lsc_load_stateless(__ESIMD_DNS::simd_mask_storage_t<N> pred,
692  __ESIMD_DNS::vector_type_t<uintptr_t, N> addrs)
693 #ifdef __SYCL_DEVICE_ONLY__
694  ;
695 #else // __SYCL_DEVICE_ONLY__
696 {
697  // TODO : Support AddressScale, ImmOffset
698  static_assert(AddressScale == 1);
699  static_assert(ImmOffset == 0);
700  static_assert(DS != __ESIMD_ENS::lsc_data_size::u16u32h);
701 
702  __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()> Output = 0;
703 
704  for (int AddrIdx = 0; AddrIdx < N; AddrIdx += 1) {
705  if (pred[AddrIdx] == 0) {
706  // Skip Output vector elements correpsonding to
707  // predicates whose value is zero
708  continue;
709  }
710 
711  constexpr uint MASK = loadstoreAlignMask<Ty, VS, DS, N>();
712  constexpr int ChanlCount = __ESIMD_EDNS::to_int<VS>();
713 
714  int ByteDistance = 0;
715  uintptr_t BaseAddr = addrs[AddrIdx];
716 
717  assert(((BaseAddr & MASK)) == 0 && "Address Alignment Error!!");
718 
719  for (int ChanelIdx = 0, VecIdx = AddrIdx; ChanelIdx < ChanlCount;
720  ChanelIdx += 1, ByteDistance += rawAddressIncrement<Ty, DS>(),
721  VecIdx += vectorIndexIncrement<N, _Transposed>()) {
722 
723  Output[VecIdx] = *((Ty *)(BaseAddr + ByteDistance));
724  }
725  }
726  return Output;
727 }
728 #endif // __SYCL_DEVICE_ONLY__
729 
748 template <typename Ty, __ESIMD_ENS::cache_hint L1H, __ESIMD_ENS::cache_hint L3H,
749  uint16_t AddressScale, int ImmOffset, __ESIMD_ENS::lsc_data_size DS,
751  __ESIMD_EDNS::lsc_data_order _Transposed, int N,
752  typename SurfIndAliasTy>
753 __ESIMD_INTRIN void
754 __esimd_lsc_prefetch_bti(__ESIMD_DNS::simd_mask_storage_t<N> pred,
755  __ESIMD_DNS::vector_type_t<uint32_t, N> offsets,
756  SurfIndAliasTy surf_ind)
757 #ifdef __SYCL_DEVICE_ONLY__
758  ;
759 #else // __SYCL_DEVICE_ONLY__
760 {
761  // Prefetch is NOP under ESIMD_EMULATOR
762  return;
763 }
764 #endif // __SYCL_DEVICE_ONLY__
765 
782 template <typename Ty, __ESIMD_ENS::cache_hint L1H, __ESIMD_ENS::cache_hint L3H,
783  uint16_t AddressScale, int ImmOffset, __ESIMD_ENS::lsc_data_size DS,
785  __ESIMD_EDNS::lsc_data_order _Transposed, int N>
786 __ESIMD_INTRIN void
787 __esimd_lsc_prefetch_stateless(__ESIMD_DNS::simd_mask_storage_t<N> pred,
788  __ESIMD_DNS::vector_type_t<uintptr_t, N> addrs)
789 #ifdef __SYCL_DEVICE_ONLY__
790  ;
791 #else // __SYCL_DEVICE_ONLY__
792 {
793  // Prefetch is NOP under ESIMD_EMULATOR
794  return;
795 }
796 #endif // __SYCL_DEVICE_ONLY__
797 
815 template <typename Ty, __ESIMD_ENS::cache_hint L1H, __ESIMD_ENS::cache_hint L3H,
816  uint16_t AddressScale, int ImmOffset, __ESIMD_ENS::lsc_data_size DS,
818  __ESIMD_EDNS::lsc_data_order _Transposed, int N>
819 __ESIMD_INTRIN void __esimd_lsc_store_slm(
820  __ESIMD_DNS::simd_mask_storage_t<N> pred,
821  __ESIMD_DNS::vector_type_t<uint32_t, N> offsets,
822  __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()> vals)
823 #ifdef __SYCL_DEVICE_ONLY__
824  ;
825 #else // __SYCL_DEVICE_ONLY__
826 {
827  sycl::detail::ESIMDDeviceInterface *I =
829 
830  __esimd_emu_lsc_offset_write<Ty, AddressScale, ImmOffset, DS, VS, _Transposed,
831  N, loadstoreAlignMask<Ty, VS, DS, N>()>(
832  pred, offsets, vals, I->__cm_emu_get_slm_ptr());
833 }
834 #endif // __SYCL_DEVICE_ONLY__
835 
855 template <typename Ty, __ESIMD_ENS::cache_hint L1H, __ESIMD_ENS::cache_hint L3H,
856  uint16_t AddressScale, int ImmOffset, __ESIMD_ENS::lsc_data_size DS,
858  __ESIMD_EDNS::lsc_data_order _Transposed, int N,
859  typename SurfIndAliasTy>
860 __ESIMD_INTRIN void __esimd_lsc_store_bti(
861  __ESIMD_DNS::simd_mask_storage_t<N> pred,
862  __ESIMD_DNS::vector_type_t<uint32_t, N> offsets,
863  __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()> vals,
864  SurfIndAliasTy surf_ind)
865 #ifdef __SYCL_DEVICE_ONLY__
866  ;
867 #else // __SYCL_DEVICE_ONLY__
868 {
869  char *writeBase;
870  uint32_t width;
871  std::mutex *mutexLock;
872 
873  sycl::detail::ESIMDDeviceInterface *I =
875 
876  I->sycl_get_cm_buffer_params_ptr(surf_ind, &writeBase, &width, &mutexLock);
877 
878  std::lock_guard<std::mutex> lock(*mutexLock);
879 
880  __esimd_emu_lsc_offset_write<Ty, AddressScale, ImmOffset, DS, VS, _Transposed,
881  N, loadstoreAlignMask<Ty, VS, DS, N>()>(
882  pred, offsets, vals, writeBase, width);
883 }
884 #endif // __SYCL_DEVICE_ONLY__
885 
903 template <typename Ty, __ESIMD_ENS::cache_hint L1H, __ESIMD_ENS::cache_hint L3H,
904  uint16_t AddressScale, int ImmOffset, __ESIMD_ENS::lsc_data_size DS,
906  __ESIMD_EDNS::lsc_data_order _Transposed, int N>
907 __ESIMD_INTRIN void __esimd_lsc_store_stateless(
908  __ESIMD_DNS::simd_mask_storage_t<N> pred,
909  __ESIMD_DNS::vector_type_t<uintptr_t, N> addrs,
910  __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()> vals)
911 #ifdef __SYCL_DEVICE_ONLY__
912  ;
913 #else // __SYCL_DEVICE_ONLY__
914 {
915  // TODO : Support AddressScale, ImmOffset
916  static_assert(AddressScale == 1);
917  static_assert(ImmOffset == 0);
918  static_assert(DS != __ESIMD_ENS::lsc_data_size::u16u32h);
919 
920  using StoreType = typename std::conditional_t<
921  DS == __ESIMD_ENS::lsc_data_size::u8, uint8_t,
923  DS == __ESIMD_ENS::lsc_data_size::u16, uint16_t,
925  DS == __ESIMD_ENS::lsc_data_size::u32, uint32_t,
927  DS == __ESIMD_ENS::lsc_data_size::u64, uint64_t,
929  DS == __ESIMD_ENS::lsc_data_size::u8u32, uint8_t,
930  std::conditional_t<DS ==
931  __ESIMD_ENS::lsc_data_size::u16u32,
932  uint16_t, void>>>>>>;
933 
934  for (int AddrIdx = 0; AddrIdx < N; AddrIdx += 1) {
935  if (pred[AddrIdx] == 0) {
936  // Skip Output vector elements correpsonding to
937  // predicates whose value is zero
938  continue;
939  }
940 
941  constexpr uint MASK = loadstoreAlignMask<Ty, VS, DS, N>();
942  constexpr int ChanlCount = __ESIMD_EDNS::to_int<VS>();
943 
944  int ByteDistance = 0;
945  uintptr_t BaseAddr = addrs[AddrIdx];
946 
947  assert(((BaseAddr & MASK)) == 0 && "Address Alignment Error!!");
948 
949  for (int ChanelIdx = 0, VecIdx = AddrIdx; ChanelIdx < ChanlCount;
950  ChanelIdx += 1, ByteDistance += rawAddressIncrement<Ty, DS>(),
951  VecIdx += vectorIndexIncrement<N, _Transposed>()) {
952  *((StoreType *)(BaseAddr + ByteDistance)) = vals[VecIdx];
953  }
954  }
955 }
956 #endif // __SYCL_DEVICE_ONLY__
957 
988 template <typename Ty, __ESIMD_ENS::cache_hint L1H, __ESIMD_ENS::cache_hint L3H,
990  __ESIMD_EDNS::lsc_data_order _Transposed, uint8_t NBlocks,
991  int BlockWidth, int BlockHeight, bool Transformed, int N>
992 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<Ty, N>
993 __esimd_lsc_load2d_stateless(__ESIMD_DNS::simd_mask_storage_t<N> Pred,
994  uintptr_t Ptr, int SurfaceWidth, int SurfaceHeight,
995  int SurfacePitch, int X, int Y)
996 #ifdef __SYCL_DEVICE_ONLY__
997  ;
998 #else // __SYCL_DEVICE_ONLY__
999 {
1000  // Template arguments are already checked by
1001  // check_lsc_block_2d_restrictions()
1002  return __esimd_emu_read_2d<Ty, N>(Pred, Ptr, SurfaceWidth, SurfaceHeight,
1003  SurfacePitch, X, Y, BlockWidth, BlockHeight,
1004  NBlocks, _Transposed, Transformed);
1005 }
1006 #endif // __SYCL_DEVICE_ONLY__
1007 
1032 template <typename Ty, __ESIMD_ENS::cache_hint L1H, __ESIMD_ENS::cache_hint L3H,
1034  __ESIMD_EDNS::lsc_data_order _Transposed, uint8_t NBlocks,
1035  int BlockWidth, int BlockHeight, bool Transformed, int N>
1036 __ESIMD_INTRIN void __esimd_lsc_prefetch2d_stateless(
1037  __ESIMD_DNS::simd_mask_storage_t<N> Pred, uintptr_t Ptr, int SurfaceWidth,
1038  int SurfaceHeight, int SurfacePitch, int X, int Y)
1039 #ifdef __SYCL_DEVICE_ONLY__
1040  ;
1041 #else // __SYCL_DEVICE_ONLY__
1042 {
1043  // Prefetch is NOP under ESIMD_EMULATOR
1044  return;
1045 }
1046 #endif // __SYCL_DEVICE_ONLY__
1047 
1077 template <typename Ty, __ESIMD_ENS::cache_hint L1H, __ESIMD_ENS::cache_hint L3H,
1079  __ESIMD_EDNS::lsc_data_order _Transposed, uint8_t NBlocks,
1080  int BlockWidth, int BlockHeight, bool Transformed, int N>
1081 __ESIMD_INTRIN void
1082 __esimd_lsc_store2d_stateless(__ESIMD_DNS::simd_mask_storage_t<N> Pred,
1083  uintptr_t Ptr, int SurfaceWidth,
1084  int SurfaceHeight, int SurfacePitch, int X, int Y,
1085  __ESIMD_DNS::vector_type_t<Ty, N> vals)
1086 #ifdef __SYCL_DEVICE_ONLY__
1087  ;
1088 #else // __SYCL_DEVICE_ONLY__
1089 {
1090  // Template arguments are already checked by
1091  // check_lsc_block_2d_restrictions()
1092  __esimd_emu_write_2d<Ty, N>(Pred, Ptr, SurfaceWidth, SurfaceHeight,
1093  SurfacePitch, X, Y, vals, BlockWidth,
1094  BlockHeight);
1095 }
1096 #endif // __SYCL_DEVICE_ONLY__
1097 
1113 template <typename Ty, __ESIMD_EDNS::lsc_atomic_op Op,
1115  uint16_t AddressScale, int ImmOffset, __ESIMD_ENS::lsc_data_size DS,
1117  __ESIMD_EDNS::lsc_data_order _Transposed, int N>
1118 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()>
1119 __esimd_lsc_xatomic_slm_0(__ESIMD_DNS::simd_mask_storage_t<N> pred,
1120  __ESIMD_DNS::vector_type_t<uint32_t, N> offsets)
1121 #ifdef __SYCL_DEVICE_ONLY__
1122  ;
1123 #else // __SYCL_DEVICE_ONLY__
1124 {
1125  __ESIMD_UNSUPPORTED_ON_HOST;
1126  return 0;
1127 }
1128 #endif // __SYCL_DEVICE_ONLY__
1129 
1146 template <typename Ty, __ESIMD_EDNS::lsc_atomic_op Op,
1148  uint16_t AddressScale, int ImmOffset, __ESIMD_ENS::lsc_data_size DS,
1150  __ESIMD_EDNS::lsc_data_order _Transposed, int N>
1151 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()>
1152 __esimd_lsc_xatomic_slm_1(
1153  __ESIMD_DNS::simd_mask_storage_t<N> pred,
1154  __ESIMD_DNS::vector_type_t<uint32_t, N> offsets,
1155  __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()> src0)
1156 #ifdef __SYCL_DEVICE_ONLY__
1157  ;
1158 #else // __SYCL_DEVICE_ONLY__
1159 {
1160  __ESIMD_UNSUPPORTED_ON_HOST;
1161  return 0;
1162 }
1163 #endif // __SYCL_DEVICE_ONLY__
1164 
1182 template <typename Ty, __ESIMD_EDNS::lsc_atomic_op Op,
1184  uint16_t AddressScale, int ImmOffset, __ESIMD_ENS::lsc_data_size DS,
1186  __ESIMD_EDNS::lsc_data_order _Transposed, int N>
1187 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()>
1188 __esimd_lsc_xatomic_slm_2(
1189  __ESIMD_DNS::simd_mask_storage_t<N> pred,
1190  __ESIMD_DNS::vector_type_t<uint32_t, N> offsets,
1191  __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()> src0,
1192  __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()> src1)
1193 #ifdef __SYCL_DEVICE_ONLY__
1194  ;
1195 #else // __SYCL_DEVICE_ONLY__
1196 {
1197  __ESIMD_UNSUPPORTED_ON_HOST;
1198  return 0;
1199 }
1200 #endif // __SYCL_DEVICE_ONLY__
1201 
1219 template <
1221  __ESIMD_ENS::cache_hint L3H, uint16_t AddressScale, int ImmOffset,
1223  __ESIMD_EDNS::lsc_data_order _Transposed, int N, typename SurfIndAliasTy>
1224 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()>
1225 __esimd_lsc_xatomic_bti_0(__ESIMD_DNS::simd_mask_storage_t<N> pred,
1226  __ESIMD_DNS::vector_type_t<uint32_t, N> offsets,
1227  SurfIndAliasTy surf_ind)
1228 #ifdef __SYCL_DEVICE_ONLY__
1229  ;
1230 #else // __SYCL_DEVICE_ONLY__
1231 {
1232  __ESIMD_UNSUPPORTED_ON_HOST;
1233  return 0;
1234 }
1235 #endif // __SYCL_DEVICE_ONLY__
1236 
1255 template <
1257  __ESIMD_ENS::cache_hint L3H, uint16_t AddressScale, int ImmOffset,
1259  __ESIMD_EDNS::lsc_data_order _Transposed, int N, typename SurfIndAliasTy>
1260 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()>
1261 __esimd_lsc_xatomic_bti_1(
1262  __ESIMD_DNS::simd_mask_storage_t<N> pred,
1263  __ESIMD_DNS::vector_type_t<uint32_t, N> offsets,
1264  __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()> src0,
1265  SurfIndAliasTy surf_ind)
1266 #ifdef __SYCL_DEVICE_ONLY__
1267  ;
1268 #else // __SYCL_DEVICE_ONLY__
1269 {
1270  __ESIMD_UNSUPPORTED_ON_HOST;
1271  return 0;
1272 }
1273 #endif // __SYCL_DEVICE_ONLY__
1274 
1294 template <
1296  __ESIMD_ENS::cache_hint L3H, uint16_t AddressScale, int ImmOffset,
1298  __ESIMD_EDNS::lsc_data_order _Transposed, int N, typename SurfIndAliasTy>
1299 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()>
1300 __esimd_lsc_xatomic_bti_2(
1301  __ESIMD_DNS::simd_mask_storage_t<N> pred,
1302  __ESIMD_DNS::vector_type_t<uint32_t, N> offsets,
1303  __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()> src0,
1304  __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()> src1,
1305  SurfIndAliasTy surf_ind)
1306 #ifdef __SYCL_DEVICE_ONLY__
1307  ;
1308 #else // __SYCL_DEVICE_ONLY__
1309 {
1310  __ESIMD_UNSUPPORTED_ON_HOST;
1311  return 0;
1312 }
1313 #endif // __SYCL_DEVICE_ONLY__
1314 
1330 template <typename Ty, __ESIMD_EDNS::lsc_atomic_op Op,
1332  uint16_t AddressScale, int ImmOffset, __ESIMD_ENS::lsc_data_size DS,
1334  __ESIMD_EDNS::lsc_data_order _Transposed, int N>
1335 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()>
1336 __esimd_lsc_xatomic_stateless_0(__ESIMD_DNS::simd_mask_storage_t<N> pred,
1337  __ESIMD_DNS::vector_type_t<uintptr_t, N> addrs)
1338 #ifdef __SYCL_DEVICE_ONLY__
1339  ;
1340 #else // __SYCL_DEVICE_ONLY__
1341 {
1342  __ESIMD_UNSUPPORTED_ON_HOST;
1343  return 0;
1344 }
1345 #endif // __SYCL_DEVICE_ONLY__
1346 
1360 
1364 template <typename Ty, __ESIMD_EDNS::lsc_atomic_op Op,
1366  uint16_t AddressScale, int ImmOffset, __ESIMD_ENS::lsc_data_size DS,
1368  __ESIMD_EDNS::lsc_data_order _Transposed, int N>
1369 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()>
1370 __esimd_lsc_xatomic_stateless_1(
1371  __ESIMD_DNS::simd_mask_storage_t<N> pred,
1372  __ESIMD_DNS::vector_type_t<uintptr_t, N> addrs,
1373  __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()> src0)
1374 #ifdef __SYCL_DEVICE_ONLY__
1375  ;
1376 #else // __SYCL_DEVICE_ONLY__
1377 {
1378  __ESIMD_UNSUPPORTED_ON_HOST;
1379  return 0;
1380 }
1381 #endif // __SYCL_DEVICE_ONLY__
1382 
1400 template <typename Ty, __ESIMD_EDNS::lsc_atomic_op Op,
1402  uint16_t AddressScale, int ImmOffset, __ESIMD_ENS::lsc_data_size DS,
1404  __ESIMD_EDNS::lsc_data_order _Transposed, int N>
1405 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()>
1406 __esimd_lsc_xatomic_stateless_2(
1407  __ESIMD_DNS::simd_mask_storage_t<N> pred,
1408  __ESIMD_DNS::vector_type_t<uintptr_t, N> addrs,
1409  __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()> src0,
1410  __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_EDNS::to_int<VS>()> src1)
1411 #ifdef __SYCL_DEVICE_ONLY__
1412  ;
1413 #else // __SYCL_DEVICE_ONLY__
1414 {
1415  __ESIMD_UNSUPPORTED_ON_HOST;
1416  return 0;
1417 }
1418 #endif // __SYCL_DEVICE_ONLY__
1419 
1429  __ESIMD_ENS::lsc_scope Scope, int N>
1430 __ESIMD_INTRIN void __esimd_lsc_fence(__ESIMD_DNS::simd_mask_storage_t<N> pred)
1431 #ifdef __SYCL_DEVICE_ONLY__
1432  ;
1433 #else // __SYCL_DEVICE_ONLY__
1434 {
1435  __ESIMD_UNSUPPORTED_ON_HOST;
1436 }
1437 #endif // __SYCL_DEVICE_ONLY__
1438 
cl::sycl::ext::intel::experimental::esimd::lsc_scope
lsc_scope
The scope that lsc_fence operation should apply to Supported platforms: DG2, PVC.
Definition: common.hpp:45
cl::sycl::detail::getESIMDDeviceInterface
ESIMDDeviceInterface * getESIMDDeviceInterface()
Definition: esimd_emulator_device_interface.cpp:25
cl::sycl::detail::conditional_t
typename std::conditional< B, T, F >::type conditional_t
Definition: stl_type_traits.hpp:27
cl::sycl::ext::intel::experimental::esimd::detail::lsc_vector_size
lsc_vector_size
Definition: common.hpp:112
cl::sycl::ext::intel::experimental::esimd::lsc_fence_op
lsc_fence_op
The lsc_fence operation to apply to caches Supported platforms: DG2, PVC.
Definition: common.hpp:57
cl::sycl::detail::getNextPowerOfTwo
constexpr size_t getNextPowerOfTwo(size_t Var)
Definition: common.hpp:330
cl::sycl::ext::intel::experimental::esimd::lsc_memory_kind
lsc_memory_kind
The specific LSC shared function to fence with lsc_fence Supported platforms: DG2,...
Definition: common.hpp:69
cl::sycl::ext::intel::experimental::esimd::cache_hint
cache_hint
L1 or L3 cache hint kinds.
Definition: common.hpp:338
cl::sycl::ext::intel::experimental::esimd::detail::lsc_atomic_op
lsc_atomic_op
LSC atomic operations op codes.
Definition: common.hpp:90
cl::sycl::ext::intel::experimental::esimd::lsc_data_size
lsc_data_size
Data size or format to read or store.
Definition: common.hpp:77
cl::sycl::access::mode
mode
Definition: access.hpp:28
memory_intrin.hpp
cl::sycl::ext::intel::experimental::esimd::split_barrier_action
split_barrier_action
Represents a split barrier action.
Definition: common.hpp:416
cl::sycl::uint
unsigned int uint
Definition: aliases.hpp:73
cl::sycl::ext::intel::experimental::esimd::detail::lsc_data_order
lsc_data_order
Definition: common.hpp:123