DPC++ Runtime
Runtime libraries for oneAPI DPC++
spirv_ops.hpp
Go to the documentation of this file.
1 //==----------- spirv_ops.hpp --- SPIRV operations -------------------------==//
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/__spirv/spirv_types.hpp> // for Scope, __ocl_event_t
12 #include <sycl/detail/defines_elementary.hpp> // for __DPCPP_SYCL_EXTERNAL
13 #include <sycl/detail/export.hpp> // for __SYCL_EXPORT
14 
15 #include <stddef.h> // for size_t
16 #include <stdint.h> // for uint32_t
17 #include <type_traits>
18 
19 // Convergent attribute
20 #ifdef __SYCL_DEVICE_ONLY__
21 #define __SYCL_CONVERGENT__ __attribute__((convergent))
22 #else
23 #define __SYCL_CONVERGENT__
24 #endif
25 
26 #ifdef __SYCL_DEVICE_ONLY__
27 
28 extern __DPCPP_SYCL_EXTERNAL float __spirv_RoundFToTF32INTEL(float a);
29 
30 template <typename T, typename Tp, std::size_t R, std::size_t C,
33  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
36  __spirv_JointMatrixLoadINTEL(T *Ptr, std::size_t Stride,
37  __spv::MatrixLayout Layout = L,
38  __spv::Scope::Flag Sc = S, int MemOperand = 0);
39 
40 template <typename T, typename Tp, std::size_t R, std::size_t C,
43  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
44 extern __DPCPP_SYCL_EXTERNAL void __spirv_JointMatrixStoreINTEL(
46  std::size_t Stride, __spv::MatrixLayout Layout = L,
47  __spv::Scope::Flag Sc = S, int MemOperand = 0);
48 
49 template <typename T, typename Tp, std::size_t R, std::size_t C,
52  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
55  __spirv_CooperativeMatrixConstructCheckedINTEL(int32_t CoordX,
56  int32_t CoordY,
57  uint32_t Height,
58  uint32_t Width,
59  const T Value);
60 
61 template <typename T, typename Tp, std::size_t R, std::size_t C,
64  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
67  __spirv_CooperativeMatrixLoadCheckedINTEL(
68  T *Ptr, int32_t CoordX, int32_t CoordY, __spv::MatrixLayout Layout = L,
69  uint32_t Height = 0, uint32_t Width = 0, std::size_t Stride = 0,
70  int MemOperand = 0);
71 
72 template <typename T, typename Tp, std::size_t R, std::size_t C,
75  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
76 extern __DPCPP_SYCL_EXTERNAL void __spirv_CooperativeMatrixStoreCheckedINTEL(
77  T *Ptr, int32_t CoordX, int32_t CoordY,
79  __spv::MatrixLayout Layout = L, uint32_t Height = 0, uint32_t Width = 0,
80  std::size_t Stride = 0, int MemOperand = 0);
81 
82 template <typename TA, typename TB, typename TC, std::size_t M, std::size_t K,
83  std::size_t N, __spv::MatrixUse UA, __spv::MatrixUse UB,
88  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
91  __spirv_JointMatrixMadINTEL(
95  __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
96 
97 template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
98  std::size_t N, __spv::MatrixUse UA, __spv::MatrixUse UB,
103  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
106  __spirv_JointMatrixUUMadINTEL(
110  __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
111 
112 template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
113  std::size_t N, __spv::MatrixUse UA, __spv::MatrixUse UB,
114  __spv::MatrixUse UC,
118  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
121  __spirv_JointMatrixUSMadINTEL(
125  __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
126 
127 template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
128  std::size_t N, __spv::MatrixUse UA, __spv::MatrixUse UB,
129  __spv::MatrixUse UC,
133  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
136  __spirv_JointMatrixSUMadINTEL(
140  __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
141 
142 template <typename T, typename Tp, std::size_t R, std::size_t C,
145  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
148  __spirv_CompositeConstruct(const T v);
149 
150 template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
152  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
153 extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<uint32_t, 2>
154 __spirv_JointMatrixGetElementCoordINTEL(
156 
157 template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
159  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
160 extern __DPCPP_SYCL_EXTERNAL size_t __spirv_JointMatrixWorkItemLengthINTEL(
162 
163 template <typename Ts, typename T, std::size_t R, std::size_t C,
166  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
167 extern __DPCPP_SYCL_EXTERNAL Ts __spirv_VectorExtractDynamic(
169 
170 template <typename Ts, typename T, std::size_t R, std::size_t C,
173  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
175 __spirv_VectorInsertDynamic(__spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *,
176  Ts val, size_t i);
177 
178 template <typename T>
179 extern __DPCPP_SYCL_EXTERNAL void __spirv_CooperativeMatrixPrefetchINTEL(
180  T *Ptr, uint32_t NumRows, uint32_t NumCols, unsigned int CacheLevel,
181  __spv::MatrixLayout Layout, size_t Stride);
182 
183 #ifndef __SPIRV_BUILTIN_DECLARATIONS__
184 #error \
185  "SPIR-V built-ins are not available. Please set -fdeclare-spirv-builtins flag."
186 #endif
187 
188 template <typename RetT, typename ImageT>
189 extern __DPCPP_SYCL_EXTERNAL RetT __spirv_ImageQueryFormat(ImageT);
190 
191 template <typename RetT, typename ImageT>
192 extern __DPCPP_SYCL_EXTERNAL RetT __spirv_ImageQueryOrder(ImageT);
193 
194 template <typename RetT, typename ImageT>
195 extern __DPCPP_SYCL_EXTERNAL RetT __spirv_ImageQuerySize(ImageT);
196 
197 template <typename ImageT, typename CoordT, typename ValT>
198 extern __DPCPP_SYCL_EXTERNAL void __spirv_ImageWrite(ImageT, CoordT, ValT);
199 
200 template <class RetT, typename ImageT, typename TempArgT>
201 extern __DPCPP_SYCL_EXTERNAL RetT __spirv_ImageRead(ImageT, TempArgT);
202 
203 template <class RetT, typename ImageT, typename TempArgT>
204 extern __DPCPP_SYCL_EXTERNAL RetT __spirv_ImageFetch(ImageT, TempArgT);
205 
206 template <class RetT, typename ImageT, typename TempArgT>
207 extern __DPCPP_SYCL_EXTERNAL RetT __spirv_SampledImageFetch(ImageT, TempArgT);
208 
209 template <class RetT, typename ImageT, typename TempArgT>
210 extern __DPCPP_SYCL_EXTERNAL RetT __spirv_ImageArrayFetch(ImageT, TempArgT,
211  int);
212 
213 template <typename ImageT, typename CoordT, typename ValT>
214 extern __DPCPP_SYCL_EXTERNAL void __spirv_ImageArrayWrite(ImageT, CoordT, int,
215  ValT);
216 
217 template <typename ImageT, typename SampledType>
218 extern __DPCPP_SYCL_EXTERNAL SampledType __spirv_SampledImage(ImageT,
220 
221 template <typename SampledType, typename TempRetT, typename TempArgT>
222 extern __DPCPP_SYCL_EXTERNAL TempRetT
223 __spirv_ImageSampleExplicitLod(SampledType, TempArgT, int, float);
224 
225 template <typename SampledType, typename TempRetT, typename TempArgT>
226 extern __DPCPP_SYCL_EXTERNAL TempRetT
227 __spirv_ImageSampleExplicitLod(SampledType, TempArgT, int, TempArgT, TempArgT);
228 
229 template <typename SampledType, typename TempRetT, typename TempArgT>
230 extern __DPCPP_SYCL_EXTERNAL TempRetT __spirv_ImageSampleCubemap(SampledType,
231  TempArgT);
232 
233 #define __SYCL_OpGroupAsyncCopyGlobalToLocal __spirv_GroupAsyncCopy
234 #define __SYCL_OpGroupAsyncCopyLocalToGlobal __spirv_GroupAsyncCopy
235 
236 // Atomic SPIR-V builtins
237 #define __SPIRV_ATOMIC_LOAD(AS, Type) \
238  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicLoad( \
239  AS const Type *P, __spv::Scope::Flag S, \
240  __spv::MemorySemanticsMask::Flag O);
241 #define __SPIRV_ATOMIC_STORE(AS, Type) \
242  extern __DPCPP_SYCL_EXTERNAL void __spirv_AtomicStore( \
243  AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
244  Type V);
245 #define __SPIRV_ATOMIC_EXCHANGE(AS, Type) \
246  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicExchange( \
247  AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
248  Type V);
249 #define __SPIRV_ATOMIC_CMP_EXCHANGE(AS, Type) \
250  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicCompareExchange( \
251  AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag E, \
252  __spv::MemorySemanticsMask::Flag U, Type V, Type C);
253 #define __SPIRV_ATOMIC_IADD(AS, Type) \
254  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicIAdd( \
255  AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
256  Type V);
257 #define __SPIRV_ATOMIC_ISUB(AS, Type) \
258  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicISub( \
259  AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
260  Type V);
261 #define __SPIRV_ATOMIC_FADD(AS, Type) \
262  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicFAddEXT( \
263  AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
264  Type V);
265 #define __SPIRV_ATOMIC_SMIN(AS, Type) \
266  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicSMin( \
267  AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
268  Type V);
269 #define __SPIRV_ATOMIC_UMIN(AS, Type) \
270  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicUMin( \
271  AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
272  Type V);
273 #define __SPIRV_ATOMIC_FMIN(AS, Type) \
274  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicFMinEXT( \
275  AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
276  Type V);
277 #define __SPIRV_ATOMIC_SMAX(AS, Type) \
278  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicSMax( \
279  AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
280  Type V);
281 #define __SPIRV_ATOMIC_UMAX(AS, Type) \
282  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicUMax( \
283  AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
284  Type V);
285 #define __SPIRV_ATOMIC_FMAX(AS, Type) \
286  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicFMaxEXT( \
287  AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
288  Type V);
289 #define __SPIRV_ATOMIC_AND(AS, Type) \
290  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicAnd( \
291  AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
292  Type V);
293 #define __SPIRV_ATOMIC_OR(AS, Type) \
294  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicOr( \
295  AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
296  Type V);
297 #define __SPIRV_ATOMIC_XOR(AS, Type) \
298  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicXor( \
299  AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
300  Type V);
301 
302 #define __SPIRV_ATOMIC_FLOAT(AS, Type) \
303  __SPIRV_ATOMIC_FADD(AS, Type) \
304  __SPIRV_ATOMIC_FMIN(AS, Type) \
305  __SPIRV_ATOMIC_FMAX(AS, Type) \
306  __SPIRV_ATOMIC_LOAD(AS, Type) \
307  __SPIRV_ATOMIC_STORE(AS, Type) \
308  __SPIRV_ATOMIC_EXCHANGE(AS, Type)
309 
310 #define __SPIRV_ATOMIC_BASE(AS, Type) \
311  __SPIRV_ATOMIC_FLOAT(AS, Type) \
312  __SPIRV_ATOMIC_CMP_EXCHANGE(AS, Type) \
313  __SPIRV_ATOMIC_IADD(AS, Type) \
314  __SPIRV_ATOMIC_ISUB(AS, Type) \
315  __SPIRV_ATOMIC_AND(AS, Type) \
316  __SPIRV_ATOMIC_OR(AS, Type) \
317  __SPIRV_ATOMIC_XOR(AS, Type)
318 
319 #define __SPIRV_ATOMIC_SIGNED(AS, Type) \
320  __SPIRV_ATOMIC_BASE(AS, Type) \
321  __SPIRV_ATOMIC_SMIN(AS, Type) \
322  __SPIRV_ATOMIC_SMAX(AS, Type)
323 
324 #define __SPIRV_ATOMIC_UNSIGNED(AS, Type) \
325  __SPIRV_ATOMIC_BASE(AS, Type) \
326  __SPIRV_ATOMIC_UMIN(AS, Type) \
327  __SPIRV_ATOMIC_UMAX(AS, Type)
328 
329 // Helper atomic operations which select correct signed/unsigned version
330 // of atomic min/max based on the type
331 #define __SPIRV_ATOMIC_MINMAX(AS, Op) \
332  template <typename T> \
333  typename std::enable_if_t< \
334  std::is_integral<T>::value && std::is_signed<T>::value, T> \
335  __spirv_Atomic##Op(AS T *Ptr, __spv::Scope::Flag Memory, \
336  __spv::MemorySemanticsMask::Flag Semantics, \
337  T Value) { \
338  return __spirv_AtomicS##Op(Ptr, Memory, Semantics, Value); \
339  } \
340  template <typename T> \
341  typename std::enable_if_t< \
342  std::is_integral<T>::value && !std::is_signed<T>::value, T> \
343  __spirv_Atomic##Op(AS T *Ptr, __spv::Scope::Flag Memory, \
344  __spv::MemorySemanticsMask::Flag Semantics, \
345  T Value) { \
346  return __spirv_AtomicU##Op(Ptr, Memory, Semantics, Value); \
347  } \
348  template <typename T> \
349  typename std::enable_if_t<std::is_floating_point<T>::value, T> \
350  __spirv_Atomic##Op(AS T *Ptr, __spv::Scope::Flag Memory, \
351  __spv::MemorySemanticsMask::Flag Semantics, \
352  T Value) { \
353  return __spirv_AtomicF##Op##EXT(Ptr, Memory, Semantics, Value); \
354  }
355 
356 #define __SPIRV_ATOMICS(macro, Arg) \
357  macro(__attribute__((opencl_global)), Arg) \
358  macro(__attribute__((opencl_local)), Arg) macro(, Arg)
359 
360 __SPIRV_ATOMICS(__SPIRV_ATOMIC_FLOAT, float)
361 __SPIRV_ATOMICS(__SPIRV_ATOMIC_FLOAT, double)
362 __SPIRV_ATOMICS(__SPIRV_ATOMIC_SIGNED, int)
363 __SPIRV_ATOMICS(__SPIRV_ATOMIC_SIGNED, long)
364 __SPIRV_ATOMICS(__SPIRV_ATOMIC_SIGNED, long long)
365 __SPIRV_ATOMICS(__SPIRV_ATOMIC_UNSIGNED, unsigned int)
366 __SPIRV_ATOMICS(__SPIRV_ATOMIC_UNSIGNED, unsigned long)
367 __SPIRV_ATOMICS(__SPIRV_ATOMIC_UNSIGNED, unsigned long long)
368 __SPIRV_ATOMICS(__SPIRV_ATOMIC_MINMAX, Min)
369 __SPIRV_ATOMICS(__SPIRV_ATOMIC_MINMAX, Max)
370 
371 #undef __SPIRV_ATOMICS
372 #undef __SPIRV_ATOMIC_AND
373 #undef __SPIRV_ATOMIC_BASE
374 #undef __SPIRV_ATOMIC_CMP_EXCHANGE
375 #undef __SPIRV_ATOMIC_EXCHANGE
376 #undef __SPIRV_ATOMIC_FADD
377 #undef __SPIRV_ATOMIC_FLOAT
378 #undef __SPIRV_ATOMIC_FMAX
379 #undef __SPIRV_ATOMIC_FMIN
380 #undef __SPIRV_ATOMIC_IADD
381 #undef __SPIRV_ATOMIC_ISUB
382 #undef __SPIRV_ATOMIC_LOAD
383 #undef __SPIRV_ATOMIC_MINMAX
384 #undef __SPIRV_ATOMIC_OR
385 #undef __SPIRV_ATOMIC_SIGNED
386 #undef __SPIRV_ATOMIC_SMAX
387 #undef __SPIRV_ATOMIC_SMIN
388 #undef __SPIRV_ATOMIC_STORE
389 #undef __SPIRV_ATOMIC_UMAX
390 #undef __SPIRV_ATOMIC_UMIN
391 #undef __SPIRV_ATOMIC_UNSIGNED
392 #undef __SPIRV_ATOMIC_XOR
393 
394 template <typename dataT>
395 extern __attribute__((opencl_global)) dataT *
396 __SYCL_GenericCastToPtrExplicit_ToGlobal(void *Ptr) noexcept {
397  return (__attribute__((opencl_global)) dataT *)
398  __spirv_GenericCastToPtrExplicit_ToGlobal(
400 }
401 
402 template <typename dataT>
403 extern const __attribute__((opencl_global)) dataT *
404 __SYCL_GenericCastToPtrExplicit_ToGlobal(const void *Ptr) noexcept {
405  return (const __attribute__((opencl_global)) dataT *)
406  __spirv_GenericCastToPtrExplicit_ToGlobal(
408 }
409 
410 template <typename dataT>
411 extern volatile __attribute__((opencl_global)) dataT *
412 __SYCL_GenericCastToPtrExplicit_ToGlobal(volatile void *Ptr) noexcept {
413  return (volatile __attribute__((opencl_global)) dataT *)
414  __spirv_GenericCastToPtrExplicit_ToGlobal(
416 }
417 
418 template <typename dataT>
419 extern const volatile __attribute__((opencl_global)) dataT *
420 __SYCL_GenericCastToPtrExplicit_ToGlobal(const volatile void *Ptr) noexcept {
421  return (const volatile __attribute__((opencl_global)) dataT *)
422  __spirv_GenericCastToPtrExplicit_ToGlobal(
424 }
425 
426 template <typename dataT>
427 extern __attribute__((opencl_local)) dataT *
428 __SYCL_GenericCastToPtrExplicit_ToLocal(void *Ptr) noexcept {
429  return (__attribute__((opencl_local)) dataT *)
430  __spirv_GenericCastToPtrExplicit_ToLocal(Ptr,
432 }
433 
434 template <typename dataT>
435 extern const __attribute__((opencl_local)) dataT *
436 __SYCL_GenericCastToPtrExplicit_ToLocal(const void *Ptr) noexcept {
437  return (const __attribute__((opencl_local)) dataT *)
438  __spirv_GenericCastToPtrExplicit_ToLocal(Ptr,
440 }
441 
442 template <typename dataT>
443 extern volatile __attribute__((opencl_local)) dataT *
444 __SYCL_GenericCastToPtrExplicit_ToLocal(volatile void *Ptr) noexcept {
445  return (volatile __attribute__((opencl_local)) dataT *)
446  __spirv_GenericCastToPtrExplicit_ToLocal(Ptr,
448 }
449 
450 template <typename dataT>
451 extern const volatile __attribute__((opencl_local)) dataT *
452 __SYCL_GenericCastToPtrExplicit_ToLocal(const volatile void *Ptr) noexcept {
453  return (const volatile __attribute__((opencl_local)) dataT *)
454  __spirv_GenericCastToPtrExplicit_ToLocal(Ptr,
456 }
457 
458 template <typename dataT>
459 extern __attribute__((opencl_private)) dataT *
460 __SYCL_GenericCastToPtrExplicit_ToPrivate(void *Ptr) noexcept {
461  return (__attribute__((opencl_private)) dataT *)
462  __spirv_GenericCastToPtrExplicit_ToPrivate(Ptr,
464 }
465 
466 template <typename dataT>
467 extern const __attribute__((opencl_private)) dataT *
468 __SYCL_GenericCastToPtrExplicit_ToPrivate(const void *Ptr) noexcept {
469  return (const __attribute__((opencl_private)) dataT *)
470  __spirv_GenericCastToPtrExplicit_ToPrivate(Ptr,
472 }
473 
474 template <typename dataT>
475 extern volatile __attribute__((opencl_private)) dataT *
476 __SYCL_GenericCastToPtrExplicit_ToPrivate(volatile void *Ptr) noexcept {
477  return (volatile __attribute__((opencl_private)) dataT *)
478  __spirv_GenericCastToPtrExplicit_ToPrivate(Ptr,
480 }
481 
482 template <typename dataT>
483 extern const volatile __attribute__((opencl_private)) dataT *
484 __SYCL_GenericCastToPtrExplicit_ToPrivate(const volatile void *Ptr) noexcept {
485  return (const volatile __attribute__((opencl_private)) dataT *)
486  __spirv_GenericCastToPtrExplicit_ToPrivate(Ptr,
488 }
489 
490 template <typename dataT>
491 extern __attribute__((opencl_global)) dataT *
492 __SYCL_GenericCastToPtr_ToGlobal(void *Ptr) noexcept {
493  return (__attribute__((opencl_global)) dataT *)
494  __spirv_GenericCastToPtr_ToGlobal(Ptr,
496 }
497 
498 template <typename dataT>
499 extern const __attribute__((opencl_global)) dataT *
500 __SYCL_GenericCastToPtr_ToGlobal(const void *Ptr) noexcept {
501  return (const __attribute__((opencl_global)) dataT *)
502  __spirv_GenericCastToPtr_ToGlobal(Ptr,
504 }
505 
506 template <typename dataT>
507 extern volatile __attribute__((opencl_global)) dataT *
508 __SYCL_GenericCastToPtr_ToGlobal(volatile void *Ptr) noexcept {
509  return (volatile __attribute__((opencl_global)) dataT *)
510  __spirv_GenericCastToPtr_ToGlobal(Ptr,
512 }
513 
514 template <typename dataT>
515 extern const volatile __attribute__((opencl_global)) dataT *
516 __SYCL_GenericCastToPtr_ToGlobal(const volatile void *Ptr) noexcept {
517  return (const volatile __attribute__((opencl_global)) dataT *)
518  __spirv_GenericCastToPtr_ToGlobal(Ptr,
520 }
521 
522 template <typename dataT>
523 extern __attribute__((opencl_local)) dataT *
524 __SYCL_GenericCastToPtr_ToLocal(void *Ptr) noexcept {
525  return (__attribute__((opencl_local)) dataT *)
526  __spirv_GenericCastToPtr_ToLocal(Ptr, __spv::StorageClass::Workgroup);
527 }
528 
529 template <typename dataT>
530 extern const __attribute__((opencl_local)) dataT *
531 __SYCL_GenericCastToPtr_ToLocal(const void *Ptr) noexcept {
532  return (const __attribute__((opencl_local)) dataT *)
533  __spirv_GenericCastToPtr_ToLocal(Ptr, __spv::StorageClass::Workgroup);
534 }
535 
536 template <typename dataT>
537 extern volatile __attribute__((opencl_local)) dataT *
538 __SYCL_GenericCastToPtr_ToLocal(volatile void *Ptr) noexcept {
539  return (volatile __attribute__((opencl_local)) dataT *)
540  __spirv_GenericCastToPtr_ToLocal(Ptr, __spv::StorageClass::Workgroup);
541 }
542 
543 template <typename dataT>
544 extern const volatile __attribute__((opencl_local)) dataT *
545 __SYCL_GenericCastToPtr_ToLocal(const volatile void *Ptr) noexcept {
546  return (const volatile __attribute__((opencl_local)) dataT *)
547  __spirv_GenericCastToPtr_ToLocal(Ptr, __spv::StorageClass::Workgroup);
548 }
549 
550 template <typename dataT>
551 extern __attribute__((opencl_private)) dataT *
552 __SYCL_GenericCastToPtr_ToPrivate(void *Ptr) noexcept {
553  return (__attribute__((opencl_private)) dataT *)
554  __spirv_GenericCastToPtr_ToPrivate(Ptr, __spv::StorageClass::Function);
555 }
556 
557 template <typename dataT>
558 extern const __attribute__((opencl_private)) dataT *
559 __SYCL_GenericCastToPtr_ToPrivate(const void *Ptr) noexcept {
560  return (const __attribute__((opencl_private)) dataT *)
561  __spirv_GenericCastToPtr_ToPrivate(Ptr, __spv::StorageClass::Function);
562 }
563 
564 template <typename dataT>
565 extern volatile __attribute__((opencl_private)) dataT *
566 __SYCL_GenericCastToPtr_ToPrivate(volatile void *Ptr) noexcept {
567  return (volatile __attribute__((opencl_private)) dataT *)
568  __spirv_GenericCastToPtr_ToPrivate(Ptr, __spv::StorageClass::Function);
569 }
570 
571 template <typename dataT>
572 extern const volatile __attribute__((opencl_private)) dataT *
573 __SYCL_GenericCastToPtr_ToPrivate(const volatile void *Ptr) noexcept {
574  return (const volatile __attribute__((opencl_private)) dataT *)
575  __spirv_GenericCastToPtr_ToPrivate(Ptr, __spv::StorageClass::Function);
576 }
577 
578 template <typename dataT>
580 __spirv_SubgroupShuffleINTEL(dataT Data, uint32_t InvocationId) noexcept;
581 template <typename dataT>
583 __spirv_SubgroupShuffleDownINTEL(dataT Current, dataT Next,
584  uint32_t Delta) noexcept;
585 template <typename dataT>
587 __spirv_SubgroupShuffleUpINTEL(dataT Previous, dataT Current,
588  uint32_t Delta) noexcept;
589 template <typename dataT>
591 __spirv_SubgroupShuffleXorINTEL(dataT Data, uint32_t Value) noexcept;
592 
593 template <typename dataT>
595 __spirv_SubgroupBlockReadINTEL(const __attribute__((opencl_global))
596  uint8_t *Ptr) noexcept;
597 
598 template <typename dataT>
600 __spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_global)) uint8_t *Ptr,
601  dataT Data) noexcept;
602 
603 template <typename dataT>
605 __spirv_SubgroupBlockReadINTEL(const __attribute__((opencl_global))
606  uint16_t *Ptr) noexcept;
607 
608 template <typename dataT>
610 __spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_global)) uint16_t *Ptr,
611  dataT Data) noexcept;
612 
613 template <typename dataT>
615 __spirv_SubgroupBlockReadINTEL(const __attribute__((opencl_global))
616  uint32_t *Ptr) noexcept;
617 
618 template <typename dataT>
620 __spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_global)) uint32_t *Ptr,
621  dataT Data) noexcept;
622 
623 template <typename dataT>
625 __spirv_SubgroupBlockReadINTEL(const __attribute__((opencl_global))
626  uint64_t *Ptr) noexcept;
627 
628 template <typename dataT>
630 __spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_global)) uint64_t *Ptr,
631  dataT Data) noexcept;
632 template <int W, int rW>
633 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<rW>
634 __spirv_FixedSqrtINTEL(sycl::detail::ap_int<W> a, bool S, int32_t I, int32_t rI,
635  int32_t Quantization = 0, int32_t Overflow = 0) noexcept;
636 template <int W, int rW>
637 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<rW>
638 __spirv_FixedRecipINTEL(sycl::detail::ap_int<W> a, bool S, int32_t I,
639  int32_t rI, int32_t Quantization = 0,
640  int32_t Overflow = 0) noexcept;
641 template <int W, int rW>
642 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<rW>
643 __spirv_FixedRsqrtINTEL(sycl::detail::ap_int<W> a, bool S, int32_t I,
644  int32_t rI, int32_t Quantization = 0,
645  int32_t Overflow = 0) noexcept;
646 template <int W, int rW>
647 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<rW>
648 __spirv_FixedSinINTEL(sycl::detail::ap_int<W> a, bool S, int32_t I, int32_t rI,
649  int32_t Quantization = 0, int32_t Overflow = 0) noexcept;
650 template <int W, int rW>
651 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<rW>
652 __spirv_FixedCosINTEL(sycl::detail::ap_int<W> a, bool S, int32_t I, int32_t rI,
653  int32_t Quantization = 0, int32_t Overflow = 0) noexcept;
654 template <int W, int rW>
655 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<2 * rW>
656 __spirv_FixedSinCosINTEL(sycl::detail::ap_int<W> a, bool S, int32_t I,
657  int32_t rI, int32_t Quantization = 0,
658  int32_t Overflow = 0) noexcept;
659 template <int W, int rW>
660 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<rW>
661 __spirv_FixedSinPiINTEL(sycl::detail::ap_int<W> a, bool S, int32_t I,
662  int32_t rI, int32_t Quantization = 0,
663  int32_t Overflow = 0) noexcept;
664 template <int W, int rW>
665 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<rW>
666 __spirv_FixedCosPiINTEL(sycl::detail::ap_int<W> a, bool S, int32_t I,
667  int32_t rI, int32_t Quantization = 0,
668  int32_t Overflow = 0) noexcept;
669 template <int W, int rW>
670 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<2 * rW>
671 __spirv_FixedSinCosPiINTEL(sycl::detail::ap_int<W> a, bool S, int32_t I,
672  int32_t rI, int32_t Quantization = 0,
673  int32_t Overflow = 0) noexcept;
674 template <int W, int rW>
675 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<rW>
676 __spirv_FixedLogINTEL(sycl::detail::ap_int<W> a, bool S, int32_t I, int32_t rI,
677  int32_t Quantization = 0, int32_t Overflow = 0) noexcept;
678 template <int W, int rW>
679 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<rW>
680 __spirv_FixedExpINTEL(sycl::detail::ap_int<W> a, bool S, int32_t I, int32_t rI,
681  int32_t Quantization = 0, int32_t Overflow = 0) noexcept;
682 
683 // In the following built-ins width of arbitrary precision integer type for
684 // a floating point variable should be equal to sum of corresponding
685 // exponent width E, mantissa width M and 1 for sign bit. I.e. WA = EA + MA + 1.
686 template <int WA, int Wout>
687 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
688 __spirv_ArbitraryFloatCastINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
689  int32_t Mout, int32_t EnableSubnormals = 0,
690  int32_t RoundingMode = 0,
691  int32_t RoundingAccuracy = 0) noexcept;
692 
693 template <int WA, int Wout>
694 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
695 __spirv_ArbitraryFloatCastFromIntINTEL(sycl::detail::ap_int<WA> A, int32_t Mout,
696  bool FromSign = false,
697  int32_t EnableSubnormals = 0,
698  int32_t RoundingMode = 0,
699  int32_t RoundingAccuracy = 0) noexcept;
700 
701 template <int WA, int Wout>
702 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
703 __spirv_ArbitraryFloatCastToIntINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
704  bool ToSign = false,
705  int32_t EnableSubnormals = 0,
706  int32_t RoundingMode = 0,
707  int32_t RoundingAccuracy = 0) noexcept;
708 
709 template <int WA, int WB, int Wout>
710 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
711 __spirv_ArbitraryFloatAddINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
712  sycl::detail::ap_int<WB> B, int32_t MB,
713  int32_t Mout, int32_t EnableSubnormals = 0,
714  int32_t RoundingMode = 0,
715  int32_t RoundingAccuracy = 0) noexcept;
716 
717 template <int WA, int WB, int Wout>
718 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
719 __spirv_ArbitraryFloatSubINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
720  sycl::detail::ap_int<WB> B, int32_t MB,
721  int32_t Mout, int32_t EnableSubnormals = 0,
722  int32_t RoundingMode = 0,
723  int32_t RoundingAccuracy = 0) noexcept;
724 
725 template <int WA, int WB, int Wout>
726 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
727 __spirv_ArbitraryFloatMulINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
728  sycl::detail::ap_int<WB> B, int32_t MB,
729  int32_t Mout, int32_t EnableSubnormals = 0,
730  int32_t RoundingMode = 0,
731  int32_t RoundingAccuracy = 0) noexcept;
732 
733 template <int WA, int WB, int Wout>
734 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
735 __spirv_ArbitraryFloatDivINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
736  sycl::detail::ap_int<WB> B, int32_t MB,
737  int32_t Mout, int32_t EnableSubnormals = 0,
738  int32_t RoundingMode = 0,
739  int32_t RoundingAccuracy = 0) noexcept;
740 
741 // Comparison built-ins don't use Subnormal Support, Rounding Mode and
742 // Rounding Accuracy.
743 template <int WA, int WB>
744 extern __DPCPP_SYCL_EXTERNAL bool
745 __spirv_ArbitraryFloatGTINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
746  sycl::detail::ap_int<WB> B, int32_t MB) noexcept;
747 
748 template <int WA, int WB>
749 extern __DPCPP_SYCL_EXTERNAL bool
750 __spirv_ArbitraryFloatGEINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
751  sycl::detail::ap_int<WB> B, int32_t MB) noexcept;
752 
753 template <int WA, int WB>
754 extern __DPCPP_SYCL_EXTERNAL bool
755 __spirv_ArbitraryFloatLTINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
756  sycl::detail::ap_int<WB> B, int32_t MB) noexcept;
757 
758 template <int WA, int WB>
759 extern __DPCPP_SYCL_EXTERNAL bool
760 __spirv_ArbitraryFloatLEINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
761  sycl::detail::ap_int<WB> B, int32_t MB) noexcept;
762 
763 template <int WA, int WB>
764 extern __DPCPP_SYCL_EXTERNAL bool
765 __spirv_ArbitraryFloatEQINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
766  sycl::detail::ap_int<WB> B, int32_t MB) noexcept;
767 
768 template <int WA, int Wout>
769 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
770 __spirv_ArbitraryFloatRecipINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
771  int32_t Mout, int32_t EnableSubnormals = 0,
772  int32_t RoundingMode = 0,
773  int32_t RoundingAccuracy = 0) noexcept;
774 
775 template <int WA, int Wout>
776 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
777 __spirv_ArbitraryFloatRSqrtINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
778  int32_t Mout, int32_t EnableSubnormals = 0,
779  int32_t RoundingMode = 0,
780  int32_t RoundingAccuracy = 0) noexcept;
781 
782 template <int WA, int Wout>
783 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
784 __spirv_ArbitraryFloatCbrtINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
785  int32_t Mout, int32_t EnableSubnormals = 0,
786  int32_t RoundingMode = 0,
787  int32_t RoundingAccuracy = 0) noexcept;
788 
789 template <int WA, int WB, int Wout>
790 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
791 __spirv_ArbitraryFloatHypotINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
792  sycl::detail::ap_int<WB> B, int32_t MB,
793  int32_t Mout, int32_t EnableSubnormals = 0,
794  int32_t RoundingMode = 0,
795  int32_t RoundingAccuracy = 0) noexcept;
796 
797 template <int WA, int Wout>
798 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
799 __spirv_ArbitraryFloatSqrtINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
800  int32_t Mout, int32_t EnableSubnormals = 0,
801  int32_t RoundingMode = 0,
802  int32_t RoundingAccuracy = 0) noexcept;
803 
804 template <int WA, int Wout>
805 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
806 __spirv_ArbitraryFloatLogINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
807  int32_t Mout, int32_t EnableSubnormals = 0,
808  int32_t RoundingMode = 0,
809  int32_t RoundingAccuracy = 0) noexcept;
810 
811 template <int WA, int Wout>
812 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
813 __spirv_ArbitraryFloatLog2INTEL(sycl::detail::ap_int<WA> A, int32_t MA,
814  int32_t Mout, int32_t EnableSubnormals = 0,
815  int32_t RoundingMode = 0,
816  int32_t RoundingAccuracy = 0) noexcept;
817 
818 template <int WA, int Wout>
819 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
820 __spirv_ArbitraryFloatLog10INTEL(sycl::detail::ap_int<WA> A, int32_t MA,
821  int32_t Mout, int32_t EnableSubnormals = 0,
822  int32_t RoundingMode = 0,
823  int32_t RoundingAccuracy = 0) noexcept;
824 
825 template <int WA, int Wout>
826 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
827 __spirv_ArbitraryFloatLog1pINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
828  int32_t Mout, int32_t EnableSubnormals = 0,
829  int32_t RoundingMode = 0,
830  int32_t RoundingAccuracy = 0) noexcept;
831 
832 template <int WA, int Wout>
833 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
834 __spirv_ArbitraryFloatExpINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
835  int32_t Mout, int32_t EnableSubnormals = 0,
836  int32_t RoundingMode = 0,
837  int32_t RoundingAccuracy = 0) noexcept;
838 
839 template <int WA, int Wout>
840 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
841 __spirv_ArbitraryFloatExp2INTEL(sycl::detail::ap_int<WA> A, int32_t MA,
842  int32_t Mout, int32_t EnableSubnormals = 0,
843  int32_t RoundingMode = 0,
844  int32_t RoundingAccuracy = 0) noexcept;
845 
846 template <int WA, int Wout>
847 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
848 __spirv_ArbitraryFloatExp10INTEL(sycl::detail::ap_int<WA> A, int32_t MA,
849  int32_t Mout, int32_t EnableSubnormals = 0,
850  int32_t RoundingMode = 0,
851  int32_t RoundingAccuracy = 0) noexcept;
852 
853 template <int WA, int Wout>
854 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
855 __spirv_ArbitraryFloatExpm1INTEL(sycl::detail::ap_int<WA> A, int32_t MA,
856  int32_t Mout, int32_t EnableSubnormals = 0,
857  int32_t RoundingMode = 0,
858  int32_t RoundingAccuracy = 0) noexcept;
859 
860 template <int WA, int Wout>
861 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
862 __spirv_ArbitraryFloatSinINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
863  int32_t Mout, int32_t EnableSubnormals = 0,
864  int32_t RoundingMode = 0,
865  int32_t RoundingAccuracy = 0) noexcept;
866 
867 template <int WA, int Wout>
868 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
869 __spirv_ArbitraryFloatCosINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
870  int32_t Mout, int32_t EnableSubnormals = 0,
871  int32_t RoundingMode = 0,
872  int32_t RoundingAccuracy = 0) noexcept;
873 
874 // Result value contains both values of sine and cosine and so has the size of
875 // 2 * Wout where Wout is equal to (1 + Eout + Mout).
876 template <int WA, int Wout>
877 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<2 * Wout>
878 __spirv_ArbitraryFloatSinCosINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
879  int32_t Mout, int32_t EnableSubnormals = 0,
880  int32_t RoundingMode = 0,
881  int32_t RoundingAccuracy = 0) noexcept;
882 
883 template <int WA, int Wout>
884 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
885 __spirv_ArbitraryFloatSinPiINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
886  int32_t Mout, int32_t EnableSubnormals = 0,
887  int32_t RoundingMode = 0,
888  int32_t RoundingAccuracy = 0) noexcept;
889 
890 template <int WA, int Wout>
891 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
892 __spirv_ArbitraryFloatCosPiINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
893  int32_t Mout, int32_t EnableSubnormals = 0,
894  int32_t RoundingMode = 0,
895  int32_t RoundingAccuracy = 0) noexcept;
896 
897 // Result value contains both values of sine(A*pi) and cosine(A*pi) and so has
898 // the size of 2 * Wout where Wout is equal to (1 + Eout + Mout).
899 template <int WA, int Wout>
900 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<2 * Wout>
901 __spirv_ArbitraryFloatSinCosPiINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
902  int32_t Mout, int32_t EnableSubnormals = 0,
903  int32_t RoundingMode = 0,
904  int32_t RoundingAccuracy = 0) noexcept;
905 
906 template <int WA, int Wout>
907 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
908 __spirv_ArbitraryFloatASinINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
909  int32_t Mout, int32_t EnableSubnormals = 0,
910  int32_t RoundingMode = 0,
911  int32_t RoundingAccuracy = 0) noexcept;
912 
913 template <int WA, int Wout>
914 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
915 __spirv_ArbitraryFloatASinPiINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
916  int32_t Mout, int32_t EnableSubnormals = 0,
917  int32_t RoundingMode = 0,
918  int32_t RoundingAccuracy = 0) noexcept;
919 
920 template <int WA, int Wout>
921 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
922 __spirv_ArbitraryFloatACosINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
923  int32_t Mout, int32_t EnableSubnormals = 0,
924  int32_t RoundingMode = 0,
925  int32_t RoundingAccuracy = 0) noexcept;
926 
927 template <int WA, int Wout>
928 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
929 __spirv_ArbitraryFloatACosPiINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
930  int32_t Mout, int32_t EnableSubnormals = 0,
931  int32_t RoundingMode = 0,
932  int32_t RoundingAccuracy = 0) noexcept;
933 
934 template <int WA, int Wout>
935 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
936 __spirv_ArbitraryFloatATanINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
937  int32_t Mout, int32_t EnableSubnormals = 0,
938  int32_t RoundingMode = 0,
939  int32_t RoundingAccuracy = 0) noexcept;
940 
941 template <int WA, int Wout>
942 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
943 __spirv_ArbitraryFloatATanPiINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
944  int32_t Mout, int32_t EnableSubnormals = 0,
945  int32_t RoundingMode = 0,
946  int32_t RoundingAccuracy = 0) noexcept;
947 
948 template <int WA, int WB, int Wout>
949 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
950 __spirv_ArbitraryFloatATan2INTEL(sycl::detail::ap_int<WA> A, int32_t MA,
951  sycl::detail::ap_int<WB> B, int32_t MB,
952  int32_t Mout, int32_t EnableSubnormals = 0,
953  int32_t RoundingMode = 0,
954  int32_t RoundingAccuracy = 0) noexcept;
955 
956 template <int WA, int WB, int Wout>
957 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
958 __spirv_ArbitraryFloatPowINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
959  sycl::detail::ap_int<WB> B, int32_t MB,
960  int32_t Mout, int32_t EnableSubnormals = 0,
961  int32_t RoundingMode = 0,
962  int32_t RoundingAccuracy = 0) noexcept;
963 
964 template <int WA, int WB, int Wout>
965 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
966 __spirv_ArbitraryFloatPowRINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
967  sycl::detail::ap_int<WB> B, int32_t MB,
968  int32_t Mout, int32_t EnableSubnormals = 0,
969  int32_t RoundingMode = 0,
970  int32_t RoundingAccuracy = 0) noexcept;
971 
972 // PowN built-in calculates `A^B` where `A` is arbitrary precision floating
973 // point number and `B` is signed or unsigned arbitrary precision integer,
974 // i.e. its width doesn't depend on sum of exponent and mantissa.
975 template <int WA, int WB, int Wout>
976 extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
977 __spirv_ArbitraryFloatPowNINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
978  sycl::detail::ap_int<WB> B, bool SignOfB,
979  int32_t Mout, int32_t EnableSubnormals = 0,
980  int32_t RoundingMode = 0,
981  int32_t RoundingAccuracy = 0) noexcept;
982 
983 template <typename dataT>
984 extern __DPCPP_SYCL_EXTERNAL int32_t
985 __spirv_ReadPipe(__ocl_RPipeTy<dataT> Pipe, dataT *Data, int32_t Size,
986  int32_t Alignment) noexcept;
987 template <typename dataT>
988 extern __DPCPP_SYCL_EXTERNAL int32_t
989 __spirv_WritePipe(__ocl_WPipeTy<dataT> Pipe, const dataT *Data, int32_t Size,
990  int32_t Alignment) noexcept;
991 template <typename dataT>
992 extern __DPCPP_SYCL_EXTERNAL void
993 __spirv_ReadPipeBlockingINTEL(__ocl_RPipeTy<dataT> Pipe, dataT *Data,
994  int32_t Size, int32_t Alignment) noexcept;
995 template <typename dataT>
996 extern __DPCPP_SYCL_EXTERNAL void
997 __spirv_WritePipeBlockingINTEL(__ocl_WPipeTy<dataT> Pipe, const dataT *Data,
998  int32_t Size, int32_t Alignment) noexcept;
999 template <typename dataT>
1000 extern __DPCPP_SYCL_EXTERNAL __ocl_RPipeTy<dataT>
1001 __spirv_CreatePipeFromPipeStorage_read(
1002  const ConstantPipeStorage *Storage) noexcept;
1003 template <typename dataT>
1004 extern __DPCPP_SYCL_EXTERNAL __ocl_WPipeTy<dataT>
1005 __spirv_CreatePipeFromPipeStorage_write(
1006  const ConstantPipeStorage *Storage) noexcept;
1007 
1008 extern __DPCPP_SYCL_EXTERNAL void
1009 __spirv_ocl_prefetch(const __attribute__((opencl_global)) char *Ptr,
1010  size_t NumBytes) noexcept;
1011 
1012 extern __DPCPP_SYCL_EXTERNAL uint16_t
1013 __spirv_ConvertFToBF16INTEL(float) noexcept;
1014 extern __DPCPP_SYCL_EXTERNAL float
1015  __spirv_ConvertBF16ToFINTEL(uint16_t) noexcept;
1016 
1018  __SYCL_EXPORT __ocl_vec_t<uint32_t, 4>
1019  __spirv_GroupNonUniformBallot(uint32_t Execution, bool Predicate) noexcept;
1020 
1021 // TODO: I'm not 100% sure that these NonUniform instructions should be
1022 // convergent Following precedent set for GroupNonUniformBallot above
1023 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT uint32_t
1024 __spirv_GroupNonUniformBallotBitCount(__spv::Scope::Flag, int,
1025  __ocl_vec_t<uint32_t, 4>) noexcept;
1026 
1027 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT int
1028  __spirv_GroupNonUniformBallotFindLSB(__spv::Scope::Flag,
1029  __ocl_vec_t<uint32_t, 4>) noexcept;
1030 
1031 template <typename ValueT, typename IdT>
1032 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1033  __spirv_GroupNonUniformBroadcast(__spv::Scope::Flag, ValueT, IdT);
1034 
1035 template <typename ValueT, typename IdT>
1036 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1037  __spirv_GroupNonUniformShuffle(__spv::Scope::Flag, ValueT, IdT) noexcept;
1038 
1039 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT bool
1040 __spirv_GroupNonUniformAll(__spv::Scope::Flag, bool);
1041 
1042 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT bool
1043 __spirv_GroupNonUniformAny(__spv::Scope::Flag, bool);
1044 
1045 template <typename ValueT>
1046 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1047 __spirv_GroupNonUniformSMin(__spv::Scope::Flag, unsigned int, ValueT);
1048 
1049 template <typename ValueT>
1050 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1051 __spirv_GroupNonUniformUMin(__spv::Scope::Flag, unsigned int, ValueT);
1052 
1053 template <typename ValueT>
1054 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1055 __spirv_GroupNonUniformFMin(__spv::Scope::Flag, unsigned int, ValueT);
1056 
1057 template <typename ValueT>
1058 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1059 __spirv_GroupNonUniformSMax(__spv::Scope::Flag, unsigned int, ValueT);
1060 
1061 template <typename ValueT>
1062 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1063 __spirv_GroupNonUniformUMax(__spv::Scope::Flag, unsigned int, ValueT);
1064 
1065 template <typename ValueT>
1066 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1067 __spirv_GroupNonUniformFMax(__spv::Scope::Flag, unsigned int, ValueT);
1068 
1069 template <typename ValueT>
1070 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1071 __spirv_GroupNonUniformIAdd(__spv::Scope::Flag, unsigned int, ValueT);
1072 
1073 template <typename ValueT>
1074 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1075 __spirv_GroupNonUniformFAdd(__spv::Scope::Flag, unsigned int, ValueT);
1076 
1077 template <typename ValueT>
1078 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1079 __spirv_GroupNonUniformIMul(__spv::Scope::Flag, unsigned int, ValueT);
1080 
1081 template <typename ValueT>
1082 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1083 __spirv_GroupNonUniformFMul(__spv::Scope::Flag, unsigned int, ValueT);
1084 
1085 template <typename ValueT>
1086 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1087 __spirv_GroupNonUniformBitwiseOr(__spv::Scope::Flag, unsigned int, ValueT);
1088 
1089 template <typename ValueT>
1090 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1091 __spirv_GroupNonUniformBitwiseXor(__spv::Scope::Flag, unsigned int, ValueT);
1092 
1093 template <typename ValueT>
1094 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1095 __spirv_GroupNonUniformBitwiseAnd(__spv::Scope::Flag, unsigned int, ValueT);
1096 
1097 template <typename ValueT>
1098 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1099 __spirv_GroupNonUniformSMin(__spv::Scope::Flag, unsigned int, ValueT,
1100  unsigned int);
1101 
1102 template <typename ValueT>
1103 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1104 __spirv_GroupNonUniformUMin(__spv::Scope::Flag, unsigned int, ValueT,
1105  unsigned int);
1106 
1107 template <typename ValueT>
1108 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1109 __spirv_GroupNonUniformFMin(__spv::Scope::Flag, unsigned int, ValueT,
1110  unsigned int);
1111 
1112 template <typename ValueT>
1113 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1114 __spirv_GroupNonUniformSMax(__spv::Scope::Flag, unsigned int, ValueT,
1115  unsigned int);
1116 
1117 template <typename ValueT>
1118 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1119 __spirv_GroupNonUniformUMax(__spv::Scope::Flag, unsigned int, ValueT,
1120  unsigned int);
1121 
1122 template <typename ValueT>
1123 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1124 __spirv_GroupNonUniformFMax(__spv::Scope::Flag, unsigned int, ValueT,
1125  unsigned int);
1126 
1127 template <typename ValueT>
1128 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1129 __spirv_GroupNonUniformIAdd(__spv::Scope::Flag, unsigned int, ValueT,
1130  unsigned int);
1131 
1132 template <typename ValueT>
1133 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1134 __spirv_GroupNonUniformFAdd(__spv::Scope::Flag, unsigned int, ValueT,
1135  unsigned int);
1136 
1137 template <typename ValueT>
1138 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1139 __spirv_GroupNonUniformIMul(__spv::Scope::Flag, unsigned int, ValueT,
1140  unsigned int);
1141 
1142 template <typename ValueT>
1143 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1144 __spirv_GroupNonUniformFMul(__spv::Scope::Flag, unsigned int, ValueT,
1145  unsigned int);
1146 
1147 template <typename ValueT>
1148 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1149 __spirv_GroupNonUniformBitwiseOr(__spv::Scope::Flag, unsigned int, ValueT,
1150  unsigned int);
1151 
1152 template <typename ValueT>
1153 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1154 __spirv_GroupNonUniformBitwiseXor(__spv::Scope::Flag, unsigned int, ValueT,
1155  unsigned int);
1156 
1157 template <typename ValueT>
1158 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1159 __spirv_GroupNonUniformBitwiseAnd(__spv::Scope::Flag, unsigned int, ValueT,
1160  unsigned int);
1161 
1162 extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void
1163 __clc_BarrierInitialize(int64_t *state, int32_t expected_count) noexcept;
1164 
1165 extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void
1166 __clc_BarrierInvalidate(int64_t *state) noexcept;
1167 
1168 extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT int64_t
1169 __clc_BarrierArrive(int64_t *state) noexcept;
1170 
1171 extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT int64_t
1172 __clc_BarrierArriveAndDrop(int64_t *state) noexcept;
1173 
1174 extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT int64_t
1175 __clc_BarrierArriveNoComplete(int64_t *state, int32_t count) noexcept;
1176 
1177 extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT int64_t
1178 __clc_BarrierArriveAndDropNoComplete(int64_t *state, int32_t count) noexcept;
1179 
1180 extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void
1181 __clc_BarrierCopyAsyncArrive(int64_t *state) noexcept;
1182 
1183 extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void
1184 __clc_BarrierCopyAsyncArriveNoInc(int64_t *state) noexcept;
1185 
1186 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void
1187 __clc_BarrierWait(int64_t *state, int64_t arrival) noexcept;
1188 
1189 extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT bool
1190 __clc_BarrierTestWait(int64_t *state, int64_t arrival) noexcept;
1191 
1192 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void
1193 __clc_BarrierArriveAndWait(int64_t *state) noexcept;
1194 
1195 #ifdef __SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__
1196 extern __DPCPP_SYCL_EXTERNAL int
1197 __spirv_ocl_printf(const __attribute__((opencl_constant)) char *Format, ...);
1198 extern __DPCPP_SYCL_EXTERNAL int __spirv_ocl_printf(const char *Format, ...);
1199 #else
1200 template <typename... Args>
1201 extern __DPCPP_SYCL_EXTERNAL int
1202 __spirv_ocl_printf(const __attribute__((opencl_constant)) char *Format,
1203  Args... args);
1204 template <typename... Args>
1205 extern __DPCPP_SYCL_EXTERNAL int __spirv_ocl_printf(const char *Format,
1206  Args... args);
1207 #endif
1208 
1209 // Native builtin extension
1210 
1211 extern __DPCPP_SYCL_EXTERNAL float __clc_native_tanh(float);
1212 extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<float, 2>
1213  __clc_native_tanh(__ocl_vec_t<float, 2>);
1214 extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<float, 3>
1215  __clc_native_tanh(__ocl_vec_t<float, 3>);
1216 extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<float, 4>
1217  __clc_native_tanh(__ocl_vec_t<float, 4>);
1218 extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<float, 8>
1219  __clc_native_tanh(__ocl_vec_t<float, 8>);
1220 extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<float, 16>
1221  __clc_native_tanh(__ocl_vec_t<float, 16>);
1222 
1223 extern __DPCPP_SYCL_EXTERNAL _Float16 __clc_native_tanh(_Float16);
1224 extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<_Float16, 2>
1225  __clc_native_tanh(__ocl_vec_t<_Float16, 2>);
1226 extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<_Float16, 3>
1227  __clc_native_tanh(__ocl_vec_t<_Float16, 3>);
1228 extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<_Float16, 4>
1229  __clc_native_tanh(__ocl_vec_t<_Float16, 4>);
1230 extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<_Float16, 8>
1231  __clc_native_tanh(__ocl_vec_t<_Float16, 8>);
1232 extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<_Float16, 16>
1233  __clc_native_tanh(__ocl_vec_t<_Float16, 16>);
1234 
1235 extern __DPCPP_SYCL_EXTERNAL _Float16 __clc_native_exp2(_Float16);
1236 extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<_Float16, 2>
1237  __clc_native_exp2(__ocl_vec_t<_Float16, 2>);
1238 extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<_Float16, 3>
1239  __clc_native_exp2(__ocl_vec_t<_Float16, 3>);
1240 extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<_Float16, 4>
1241  __clc_native_exp2(__ocl_vec_t<_Float16, 4>);
1242 extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<_Float16, 8>
1243  __clc_native_exp2(__ocl_vec_t<_Float16, 8>);
1244 extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<_Float16, 16>
1245  __clc_native_exp2(__ocl_vec_t<_Float16, 16>);
1246 
1247 #define __CLC_BF16(...) \
1248  extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT __VA_ARGS__ __clc_fabs( \
1249  __VA_ARGS__) noexcept; \
1250  extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT __VA_ARGS__ __clc_fmin( \
1251  __VA_ARGS__, __VA_ARGS__) noexcept; \
1252  extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT __VA_ARGS__ __clc_fmax( \
1253  __VA_ARGS__, __VA_ARGS__) noexcept; \
1254  extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT __VA_ARGS__ __clc_fma( \
1255  __VA_ARGS__, __VA_ARGS__, __VA_ARGS__) noexcept;
1256 
1257 #define __CLC_BF16_SCAL_VEC(TYPE) \
1258  __CLC_BF16(TYPE) \
1259  __CLC_BF16(__ocl_vec_t<TYPE, 2>) \
1260  __CLC_BF16(__ocl_vec_t<TYPE, 3>) \
1261  __CLC_BF16(__ocl_vec_t<TYPE, 4>) \
1262  __CLC_BF16(__ocl_vec_t<TYPE, 8>) \
1263  __CLC_BF16(__ocl_vec_t<TYPE, 16>)
1264 
1265 __CLC_BF16_SCAL_VEC(uint16_t)
1266 __CLC_BF16_SCAL_VEC(uint32_t)
1267 
1268 #undef __CLC_BF16_SCAL_VEC
1269 #undef __CLC_BF16
1270 
1271 extern __DPCPP_SYCL_EXTERNAL int32_t __spirv_BuiltInGlobalHWThreadIDINTEL();
1272 extern __DPCPP_SYCL_EXTERNAL int32_t __spirv_BuiltInSubDeviceIDINTEL();
1273 
1274 template <typename from, typename to>
1275 extern __DPCPP_SYCL_EXTERNAL
1276  std::enable_if_t<std::is_integral_v<to> && std::is_unsigned_v<to>, to>
1277  __spirv_ConvertPtrToU(from val) noexcept;
1278 
1279 template <typename RetT, typename... ArgsT>
1280 extern __DPCPP_SYCL_EXTERNAL __spv::__spirv_TaskSequenceINTEL *
1281 __spirv_TaskSequenceCreateINTEL(RetT (*f)(ArgsT...), int Pipelined = -1,
1282  int ClusterMode = -1,
1283  unsigned int ResponseCapacity = 0,
1284  unsigned int InvocationCapacity = 0) noexcept;
1285 
1286 template <typename... ArgsT>
1287 extern __DPCPP_SYCL_EXTERNAL void
1288 __spirv_TaskSequenceAsyncINTEL(__spv::__spirv_TaskSequenceINTEL *TaskSequence,
1289  ArgsT... Args) noexcept;
1290 
1291 template <typename RetT>
1292 extern __DPCPP_SYCL_EXTERNAL RetT __spirv_TaskSequenceGetINTEL(
1293  __spv::__spirv_TaskSequenceINTEL *TaskSequence) noexcept;
1294 
1295 extern __DPCPP_SYCL_EXTERNAL void __spirv_TaskSequenceReleaseINTEL(
1296  __spv::__spirv_TaskSequenceINTEL *TaskSequence) noexcept;
1297 
1298 #else // if !__SYCL_DEVICE_ONLY__
1299 
1300 template <typename dataT>
1303  const dataT *Src, size_t NumElements,
1304  size_t Stride, __ocl_event_t) noexcept {
1305  for (size_t i = 0; i < NumElements; i++) {
1306  Dest[i] = Src[i * Stride];
1307  }
1308  // A real instance of the class is not needed, return dummy pointer.
1309  return nullptr;
1310 }
1311 
1312 template <typename dataT>
1315  const dataT *Src, size_t NumElements,
1316  size_t Stride, __ocl_event_t) noexcept {
1317  for (size_t i = 0; i < NumElements; i++) {
1318  Dest[i * Stride] = Src[i];
1319  }
1320  // A real instance of the class is not needed, return dummy pointer.
1321  return nullptr;
1322 }
1323 
1324 extern __SYCL_EXPORT void __spirv_ocl_prefetch(const char *Ptr,
1325  size_t NumBytes) noexcept;
1326 
1327 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void
1329  uint32_t Semantics) noexcept;
1330 
1331 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void
1332 __spirv_MemoryBarrier(__spv::Scope Memory, uint32_t Semantics) noexcept;
1333 
1334 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void
1335 __spirv_GroupWaitEvents(__spv::Scope Execution, uint32_t NumEvents,
1336  __ocl_event_t *WaitEvents) noexcept;
1337 #endif // !__SYCL_DEVICE_ONLY__
#define __DPCPP_SYCL_EXTERNAL
conditional< sizeof(long)==8, long, long long >::type int64_t
Definition: kernel_desc.hpp:35
__attribute__((always_inline)) auto invoke_simd(sycl
The invoke_simd free function invokes a SIMD function using all work-items in a sub_group.
Definition: access.hpp:18
__SYCL_CONVERGENT__ __ocl_event_t __SYCL_OpGroupAsyncCopyGlobalToLocal(__spv::Scope::Flag, dataT *Dest, const dataT *Src, size_t NumElements, size_t Stride, __ocl_event_t) noexcept
Definition: spirv_ops.hpp:1302
__SYCL_CONVERGENT__ __DPCPP_SYCL_EXTERNAL void __spirv_ControlBarrier(__spv::Scope Execution, __spv::Scope Memory, uint32_t Semantics) noexcept
Definition: spirv_ops.cpp:26
__SYCL_CONVERGENT__ __DPCPP_SYCL_EXTERNAL void __spirv_MemoryBarrier(__spv::Scope Memory, uint32_t Semantics) noexcept
Definition: spirv_ops.cpp:36
__SYCL_CONVERGENT__ __DPCPP_SYCL_EXTERNAL void __spirv_GroupWaitEvents(__spv::Scope Execution, uint32_t NumEvents, __ocl_event_t *WaitEvents) noexcept
Definition: spirv_ops.cpp:18
#define __SYCL_CONVERGENT__
Definition: spirv_ops.hpp:23
void __spirv_ocl_prefetch(const char *Ptr, size_t NumBytes) noexcept
Definition: spirv_ops.cpp:47
__SYCL_CONVERGENT__ __ocl_event_t __SYCL_OpGroupAsyncCopyLocalToGlobal(__spv::Scope::Flag, dataT *Dest, const dataT *Src, size_t NumElements, size_t Stride, __ocl_event_t) noexcept
Definition: spirv_ops.hpp:1314
void * __ocl_event_t
void * __ocl_sampler_t
_Abi const simd< _Tp, _Abi > & noexcept
Definition: simd.hpp:1324