18 #ifdef __SYCL_DEVICE_ONLY__
19 #define __SYCL_CONVERGENT__ __attribute__((convergent))
21 #define __SYCL_CONVERGENT__
24 #ifdef __SYCL_DEVICE_ONLY__
26 #if (SYCL_EXT_ONEAPI_MATRIX_VERSION > 1)
31 __spirv_JointMatrixLoadINTEL(T *Ptr, std::size_t Stride,
43 template <
typename T1,
typename T2, std::size_t M, std::size_t K, std::size_t N,
50 __spirv_JointMatrixMadINTEL(
56 template <
typename T1,
typename T2,
typename T3, std::size_t M, std::size_t K,
64 __spirv_JointMatrixUUMadINTEL(
70 template <
typename T1,
typename T2,
typename T3, std::size_t M, std::size_t K,
78 __spirv_JointMatrixUSMadINTEL(
84 template <
typename T1,
typename T2,
typename T3, std::size_t M, std::size_t K,
92 __spirv_JointMatrixSUMadINTEL(
102 __spirv_CompositeConstruct(
const T v);
107 extern SYCL_EXTERNAL size_t __spirv_JointMatrixWorkItemLengthINTEL(
123 template <
typename T, std::size_t R, std::size_t C,
127 __spirv_JointMatrixLoadINTEL(T *Ptr, std::size_t Stride,
131 template <
typename T, std::size_t R, std::size_t C,
139 template <
typename T1,
typename T2, std::size_t M, std::size_t K, std::size_t N,
145 __spirv_JointMatrixMadINTEL(
151 template <
typename T1,
typename T2,
typename T3, std::size_t M, std::size_t K,
157 __spirv_JointMatrixUUMadINTEL(
163 template <
typename T1,
typename T2,
typename T3, std::size_t M, std::size_t K,
169 __spirv_JointMatrixUSMadINTEL(
175 template <
typename T1,
typename T2,
typename T3, std::size_t M, std::size_t K,
181 __spirv_JointMatrixSUMadINTEL(
187 template <
typename T, std::size_t R, std::size_t C,
191 __spirv_CompositeConstruct(
const T v);
193 template <
typename T, std::size_t R, std::size_t C,
196 extern SYCL_EXTERNAL size_t __spirv_JointMatrixWorkItemLengthINTEL(
199 template <
typename T, std::size_t R, std::size_t C,
205 template <
typename T, std::size_t R, std::size_t C,
211 #endif // SYCL_EXT_ONEAPI_MATRIX_VERSION
213 #ifndef __SPIRV_BUILTIN_DECLARATIONS__
215 "SPIR-V built-ins are not available. Please set -fdeclare-spirv-builtins flag."
218 template <
typename RetT,
typename ImageT>
221 template <
typename RetT,
typename ImageT>
224 template <
typename RetT,
typename ImageT>
227 template <
typename ImageT,
typename CoordT,
typename ValT>
228 extern SYCL_EXTERNAL void __spirv_ImageWrite(ImageT, CoordT, ValT);
230 template <
class RetT,
typename ImageT,
typename TempArgT>
231 extern SYCL_EXTERNAL RetT __spirv_ImageRead(ImageT, TempArgT);
233 template <
typename ImageT,
typename SampledType>
236 template <
typename SampledType,
typename TempRetT,
typename TempArgT>
237 extern SYCL_EXTERNAL TempRetT __spirv_ImageSampleExplicitLod(SampledType,
241 #define __SYCL_OpGroupAsyncCopyGlobalToLocal __spirv_GroupAsyncCopy
242 #define __SYCL_OpGroupAsyncCopyLocalToGlobal __spirv_GroupAsyncCopy
245 #define __SPIRV_ATOMIC_LOAD(AS, Type) \
246 extern SYCL_EXTERNAL Type __spirv_AtomicLoad( \
247 AS const Type *P, __spv::Scope::Flag S, \
248 __spv::MemorySemanticsMask::Flag O);
249 #define __SPIRV_ATOMIC_STORE(AS, Type) \
250 extern SYCL_EXTERNAL void __spirv_AtomicStore( \
251 AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
253 #define __SPIRV_ATOMIC_EXCHANGE(AS, Type) \
254 extern SYCL_EXTERNAL Type __spirv_AtomicExchange( \
255 AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
257 #define __SPIRV_ATOMIC_CMP_EXCHANGE(AS, Type) \
258 extern SYCL_EXTERNAL Type __spirv_AtomicCompareExchange( \
259 AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag E, \
260 __spv::MemorySemanticsMask::Flag U, Type V, Type C);
261 #define __SPIRV_ATOMIC_IADD(AS, Type) \
262 extern SYCL_EXTERNAL Type __spirv_AtomicIAdd( \
263 AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
265 #define __SPIRV_ATOMIC_ISUB(AS, Type) \
266 extern SYCL_EXTERNAL Type __spirv_AtomicISub( \
267 AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
269 #define __SPIRV_ATOMIC_FADD(AS, Type) \
270 extern SYCL_EXTERNAL Type __spirv_AtomicFAddEXT( \
271 AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
273 #define __SPIRV_ATOMIC_SMIN(AS, Type) \
274 extern SYCL_EXTERNAL Type __spirv_AtomicSMin( \
275 AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
277 #define __SPIRV_ATOMIC_UMIN(AS, Type) \
278 extern SYCL_EXTERNAL Type __spirv_AtomicUMin( \
279 AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
281 #define __SPIRV_ATOMIC_FMIN(AS, Type) \
282 extern SYCL_EXTERNAL Type __spirv_AtomicFMinEXT( \
283 AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
285 #define __SPIRV_ATOMIC_SMAX(AS, Type) \
286 extern SYCL_EXTERNAL Type __spirv_AtomicSMax( \
287 AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
289 #define __SPIRV_ATOMIC_UMAX(AS, Type) \
290 extern SYCL_EXTERNAL Type __spirv_AtomicUMax( \
291 AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
293 #define __SPIRV_ATOMIC_FMAX(AS, Type) \
294 extern SYCL_EXTERNAL Type __spirv_AtomicFMaxEXT( \
295 AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
297 #define __SPIRV_ATOMIC_AND(AS, Type) \
298 extern SYCL_EXTERNAL Type __spirv_AtomicAnd( \
299 AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
301 #define __SPIRV_ATOMIC_OR(AS, Type) \
302 extern SYCL_EXTERNAL Type __spirv_AtomicOr( \
303 AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
305 #define __SPIRV_ATOMIC_XOR(AS, Type) \
306 extern SYCL_EXTERNAL Type __spirv_AtomicXor( \
307 AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O, \
310 #define __SPIRV_ATOMIC_FLOAT(AS, Type) \
311 __SPIRV_ATOMIC_FADD(AS, Type) \
312 __SPIRV_ATOMIC_FMIN(AS, Type) \
313 __SPIRV_ATOMIC_FMAX(AS, Type) \
314 __SPIRV_ATOMIC_LOAD(AS, Type) \
315 __SPIRV_ATOMIC_STORE(AS, Type) \
316 __SPIRV_ATOMIC_EXCHANGE(AS, Type)
318 #define __SPIRV_ATOMIC_BASE(AS, Type) \
319 __SPIRV_ATOMIC_FLOAT(AS, Type) \
320 __SPIRV_ATOMIC_CMP_EXCHANGE(AS, Type) \
321 __SPIRV_ATOMIC_IADD(AS, Type) \
322 __SPIRV_ATOMIC_ISUB(AS, Type) \
323 __SPIRV_ATOMIC_AND(AS, Type) \
324 __SPIRV_ATOMIC_OR(AS, Type) \
325 __SPIRV_ATOMIC_XOR(AS, Type)
327 #define __SPIRV_ATOMIC_SIGNED(AS, Type) \
328 __SPIRV_ATOMIC_BASE(AS, Type) \
329 __SPIRV_ATOMIC_SMIN(AS, Type) \
330 __SPIRV_ATOMIC_SMAX(AS, Type)
332 #define __SPIRV_ATOMIC_UNSIGNED(AS, Type) \
333 __SPIRV_ATOMIC_BASE(AS, Type) \
334 __SPIRV_ATOMIC_UMIN(AS, Type) \
335 __SPIRV_ATOMIC_UMAX(AS, Type)
339 #define __SPIRV_ATOMIC_MINMAX(AS, Op) \
340 template <typename T> \
341 typename sycl::detail::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, \
346 return __spirv_AtomicS##Op(Ptr, Memory, Semantics, Value); \
348 template <typename T> \
349 typename sycl::detail::enable_if_t< \
350 std::is_integral<T>::value && !std::is_signed<T>::value, T> \
351 __spirv_Atomic##Op(AS T *Ptr, __spv::Scope::Flag Memory, \
352 __spv::MemorySemanticsMask::Flag Semantics, \
354 return __spirv_AtomicU##Op(Ptr, Memory, Semantics, Value); \
356 template <typename T> \
357 typename sycl::detail::enable_if_t<std::is_floating_point<T>::value, T> \
358 __spirv_Atomic##Op(AS T *Ptr, __spv::Scope::Flag Memory, \
359 __spv::MemorySemanticsMask::Flag Semantics, \
361 return __spirv_AtomicF##Op##EXT(Ptr, Memory, Semantics, Value); \
364 #define __SPIRV_ATOMICS(macro, Arg) \
365 macro(__attribute__((opencl_global)), Arg) \
366 macro(__attribute__((opencl_local)), Arg) macro(, Arg)
368 __SPIRV_ATOMICS(__SPIRV_ATOMIC_FLOAT,
float)
369 __SPIRV_ATOMICS(__SPIRV_ATOMIC_FLOAT,
double)
370 __SPIRV_ATOMICS(__SPIRV_ATOMIC_SIGNED,
int)
371 __SPIRV_ATOMICS(__SPIRV_ATOMIC_SIGNED,
long)
372 __SPIRV_ATOMICS(__SPIRV_ATOMIC_SIGNED,
long long)
373 __SPIRV_ATOMICS(__SPIRV_ATOMIC_UNSIGNED,
unsigned int)
374 __SPIRV_ATOMICS(__SPIRV_ATOMIC_UNSIGNED,
unsigned long)
375 __SPIRV_ATOMICS(__SPIRV_ATOMIC_UNSIGNED,
unsigned long long)
376 __SPIRV_ATOMICS(__SPIRV_ATOMIC_MINMAX, Min)
377 __SPIRV_ATOMICS(__SPIRV_ATOMIC_MINMAX, Max)
379 #undef __SPIRV_ATOMICS
380 #undef __SPIRV_ATOMIC_AND
381 #undef __SPIRV_ATOMIC_BASE
382 #undef __SPIRV_ATOMIC_CMP_EXCHANGE
383 #undef __SPIRV_ATOMIC_EXCHANGE
384 #undef __SPIRV_ATOMIC_FADD
385 #undef __SPIRV_ATOMIC_FLOAT
386 #undef __SPIRV_ATOMIC_FMAX
387 #undef __SPIRV_ATOMIC_FMIN
388 #undef __SPIRV_ATOMIC_IADD
389 #undef __SPIRV_ATOMIC_ISUB
390 #undef __SPIRV_ATOMIC_LOAD
391 #undef __SPIRV_ATOMIC_MINMAX
392 #undef __SPIRV_ATOMIC_OR
393 #undef __SPIRV_ATOMIC_SIGNED
394 #undef __SPIRV_ATOMIC_SMAX
395 #undef __SPIRV_ATOMIC_SMIN
396 #undef __SPIRV_ATOMIC_STORE
397 #undef __SPIRV_ATOMIC_UMAX
398 #undef __SPIRV_ATOMIC_UMIN
399 #undef __SPIRV_ATOMIC_UNSIGNED
400 #undef __SPIRV_ATOMIC_XOR
402 template <
typename dataT>
404 __SYCL_GenericCastToPtrExplicit_ToGlobal(
void *Ptr) noexcept {
406 __spirv_GenericCastToPtrExplicit_ToGlobal(
410 template <
typename dataT>
412 __SYCL_GenericCastToPtrExplicit_ToGlobal(
const void *Ptr) noexcept {
414 __spirv_GenericCastToPtrExplicit_ToGlobal(
418 template <
typename dataT>
420 __SYCL_GenericCastToPtrExplicit_ToGlobal(
volatile void *Ptr) noexcept {
422 __spirv_GenericCastToPtrExplicit_ToGlobal(
426 template <
typename dataT>
428 __SYCL_GenericCastToPtrExplicit_ToGlobal(
const volatile void *Ptr) noexcept {
429 return (
const volatile __attribute__((opencl_global)) dataT *)
430 __spirv_GenericCastToPtrExplicit_ToGlobal(
434 template <
typename dataT>
436 __SYCL_GenericCastToPtrExplicit_ToLocal(
void *Ptr) noexcept {
438 __spirv_GenericCastToPtrExplicit_ToLocal(Ptr,
442 template <
typename dataT>
444 __SYCL_GenericCastToPtrExplicit_ToLocal(
const void *Ptr) noexcept {
446 __spirv_GenericCastToPtrExplicit_ToLocal(Ptr,
450 template <
typename dataT>
452 __SYCL_GenericCastToPtrExplicit_ToLocal(
volatile void *Ptr) noexcept {
454 __spirv_GenericCastToPtrExplicit_ToLocal(Ptr,
458 template <
typename dataT>
460 __SYCL_GenericCastToPtrExplicit_ToLocal(
const volatile void *Ptr) noexcept {
461 return (
const volatile __attribute__((opencl_local)) dataT *)
462 __spirv_GenericCastToPtrExplicit_ToLocal(Ptr,
466 template <
typename dataT>
468 __SYCL_GenericCastToPtrExplicit_ToPrivate(
void *Ptr) noexcept {
470 __spirv_GenericCastToPtrExplicit_ToPrivate(Ptr,
474 template <
typename dataT>
476 __SYCL_GenericCastToPtrExplicit_ToPrivate(
const void *Ptr) noexcept {
478 __spirv_GenericCastToPtrExplicit_ToPrivate(Ptr,
482 template <
typename dataT>
484 __SYCL_GenericCastToPtrExplicit_ToPrivate(
volatile void *Ptr) noexcept {
486 __spirv_GenericCastToPtrExplicit_ToPrivate(Ptr,
490 template <
typename dataT>
491 extern const volatile __attribute__((opencl_private)) dataT *
492 __SYCL_GenericCastToPtrExplicit_ToPrivate(
const volatile void *Ptr) noexcept {
493 return (
const volatile __attribute__((opencl_private)) dataT *)
494 __spirv_GenericCastToPtrExplicit_ToPrivate(Ptr,
498 template <
typename dataT>
500 __spirv_SubgroupShuffleINTEL(dataT Data, uint32_t InvocationId) noexcept;
501 template <
typename dataT>
503 dataT Current, dataT Next, uint32_t Delta) noexcept;
504 template <
typename dataT>
506 dataT Previous, dataT Current, uint32_t Delta) noexcept;
507 template <
typename dataT>
509 __spirv_SubgroupShuffleXorINTEL(dataT Data, uint32_t Value) noexcept;
511 template <
typename dataT>
515 template <
typename dataT>
517 __spirv_SubgroupBlockWriteINTEL(
__attribute__((opencl_global)) uint8_t *Ptr,
518 dataT Data) noexcept;
520 template <
typename dataT>
522 const __attribute__((opencl_global)) uint16_t *Ptr) noexcept;
524 template <
typename dataT>
526 __spirv_SubgroupBlockWriteINTEL(
__attribute__((opencl_global)) uint16_t *Ptr,
527 dataT Data) noexcept;
529 template <
typename dataT>
531 const __attribute__((opencl_global)) uint32_t *Ptr) noexcept;
533 template <
typename dataT>
535 __spirv_SubgroupBlockWriteINTEL(
__attribute__((opencl_global)) uint32_t *Ptr,
536 dataT Data) noexcept;
538 template <
typename dataT>
540 const __attribute__((opencl_global)) uint64_t *Ptr) noexcept;
542 template <
typename dataT>
544 __spirv_SubgroupBlockWriteINTEL(
__attribute__((opencl_global)) uint64_t *Ptr,
545 dataT Data) noexcept;
546 template <
int W,
int rW>
548 __spirv_FixedSqrtINTEL(sycl::detail::ap_int<W> a,
bool S, int32_t I, int32_t rI,
549 int32_t Quantization = 0, int32_t Overflow = 0) noexcept;
550 template <
int W,
int rW>
552 __spirv_FixedRecipINTEL(
sycl::detail::ap_int<W> a,
bool S, int32_t I,
553 int32_t rI, int32_t Quantization = 0,
554 int32_t Overflow = 0) noexcept;
555 template <
int W,
int rW>
557 __spirv_FixedRsqrtINTEL(
sycl::detail::ap_int<W> a,
bool S, int32_t I,
558 int32_t rI, int32_t Quantization = 0,
559 int32_t Overflow = 0) noexcept;
560 template <
int W,
int rW>
562 __spirv_FixedSinINTEL(
sycl::detail::ap_int<W> a,
bool S, int32_t I, int32_t rI,
563 int32_t Quantization = 0, int32_t Overflow = 0) noexcept;
564 template <
int W,
int rW>
566 __spirv_FixedCosINTEL(
sycl::detail::ap_int<W> a,
bool S, int32_t I, int32_t rI,
567 int32_t Quantization = 0, int32_t Overflow = 0) noexcept;
568 template <
int W,
int rW>
570 __spirv_FixedSinCosINTEL(
sycl::detail::ap_int<W> a,
bool S, int32_t I,
571 int32_t rI, int32_t Quantization = 0,
572 int32_t Overflow = 0) noexcept;
573 template <
int W,
int rW>
575 __spirv_FixedSinPiINTEL(
sycl::detail::ap_int<W> a,
bool S, int32_t I,
576 int32_t rI, int32_t Quantization = 0,
577 int32_t Overflow = 0) noexcept;
578 template <
int W,
int rW>
580 __spirv_FixedCosPiINTEL(
sycl::detail::ap_int<W> a,
bool S, int32_t I,
581 int32_t rI, int32_t Quantization = 0,
582 int32_t Overflow = 0) noexcept;
583 template <
int W,
int rW>
585 __spirv_FixedSinCosPiINTEL(
sycl::detail::ap_int<W> a,
bool S, int32_t I,
586 int32_t rI, int32_t Quantization = 0,
587 int32_t Overflow = 0) noexcept;
588 template <
int W,
int rW>
590 __spirv_FixedLogINTEL(
sycl::detail::ap_int<W> a,
bool S, int32_t I, int32_t rI,
591 int32_t Quantization = 0, int32_t Overflow = 0) noexcept;
592 template <
int W,
int rW>
594 __spirv_FixedExpINTEL(
sycl::detail::ap_int<W> a,
bool S, int32_t I, int32_t rI,
595 int32_t Quantization = 0, int32_t Overflow = 0) noexcept;
600 template <
int WA,
int Wout>
602 __spirv_ArbitraryFloatCastINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
603 int32_t Mout, int32_t EnableSubnormals = 0,
604 int32_t RoundingMode = 0,
605 int32_t RoundingAccuracy = 0) noexcept;
607 template <
int WA,
int Wout>
609 __spirv_ArbitraryFloatCastFromIntINTEL(
sycl::detail::ap_int<WA> A, int32_t Mout,
610 bool FromSign = false,
611 int32_t EnableSubnormals = 0,
612 int32_t RoundingMode = 0,
613 int32_t RoundingAccuracy = 0) noexcept;
615 template <
int WA,
int Wout>
617 __spirv_ArbitraryFloatCastToIntINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
619 int32_t EnableSubnormals = 0,
620 int32_t RoundingMode = 0,
621 int32_t RoundingAccuracy = 0) noexcept;
623 template <
int WA,
int WB,
int Wout>
625 sycl::detail::ap_int<WA> A, int32_t MA,
sycl::detail::ap_int<WB> B,
626 int32_t MB, int32_t Mout, int32_t EnableSubnormals = 0,
627 int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept;
629 template <
int WA,
int WB,
int Wout>
631 sycl::detail::ap_int<WA> A, int32_t MA,
sycl::detail::ap_int<WB> B,
632 int32_t MB, int32_t Mout, int32_t EnableSubnormals = 0,
633 int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept;
635 template <
int WA,
int WB,
int Wout>
637 sycl::detail::ap_int<WA> A, int32_t MA,
sycl::detail::ap_int<WB> B,
638 int32_t MB, int32_t Mout, int32_t EnableSubnormals = 0,
639 int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept;
641 template <
int WA,
int WB,
int Wout>
643 sycl::detail::ap_int<WA> A, int32_t MA,
sycl::detail::ap_int<WB> B,
644 int32_t MB, int32_t Mout, int32_t EnableSubnormals = 0,
645 int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept;
649 template <
int WA,
int WB>
651 __spirv_ArbitraryFloatGTINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
652 sycl::detail::ap_int<WB> B, int32_t MB) noexcept;
654 template <
int WA,
int WB>
656 __spirv_ArbitraryFloatGEINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
657 sycl::detail::ap_int<WB> B, int32_t MB) noexcept;
659 template <
int WA,
int WB>
661 __spirv_ArbitraryFloatLTINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
662 sycl::detail::ap_int<WB> B, int32_t MB) noexcept;
664 template <
int WA,
int WB>
666 __spirv_ArbitraryFloatLEINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
667 sycl::detail::ap_int<WB> B, int32_t MB) noexcept;
669 template <
int WA,
int WB>
671 __spirv_ArbitraryFloatEQINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
672 sycl::detail::ap_int<WB> B, int32_t MB) noexcept;
674 template <
int WA,
int Wout>
676 __spirv_ArbitraryFloatRecipINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
677 int32_t Mout, int32_t EnableSubnormals = 0,
678 int32_t RoundingMode = 0,
679 int32_t RoundingAccuracy = 0) noexcept;
681 template <
int WA,
int Wout>
683 __spirv_ArbitraryFloatRSqrtINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
684 int32_t Mout, int32_t EnableSubnormals = 0,
685 int32_t RoundingMode = 0,
686 int32_t RoundingAccuracy = 0) noexcept;
688 template <
int WA,
int Wout>
690 __spirv_ArbitraryFloatCbrtINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
691 int32_t Mout, int32_t EnableSubnormals = 0,
692 int32_t RoundingMode = 0,
693 int32_t RoundingAccuracy = 0) noexcept;
695 template <
int WA,
int WB,
int Wout>
697 __spirv_ArbitraryFloatHypotINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
698 sycl::detail::ap_int<WB> B, int32_t MB,
699 int32_t Mout, int32_t EnableSubnormals = 0,
700 int32_t RoundingMode = 0,
701 int32_t RoundingAccuracy = 0) noexcept;
703 template <
int WA,
int Wout>
705 __spirv_ArbitraryFloatSqrtINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
706 int32_t Mout, int32_t EnableSubnormals = 0,
707 int32_t RoundingMode = 0,
708 int32_t RoundingAccuracy = 0) noexcept;
710 template <
int WA,
int Wout>
712 __spirv_ArbitraryFloatLogINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
713 int32_t Mout, int32_t EnableSubnormals = 0,
714 int32_t RoundingMode = 0,
715 int32_t RoundingAccuracy = 0) noexcept;
717 template <
int WA,
int Wout>
719 __spirv_ArbitraryFloatLog2INTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
720 int32_t Mout, int32_t EnableSubnormals = 0,
721 int32_t RoundingMode = 0,
722 int32_t RoundingAccuracy = 0) noexcept;
724 template <
int WA,
int Wout>
726 __spirv_ArbitraryFloatLog10INTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
727 int32_t Mout, int32_t EnableSubnormals = 0,
728 int32_t RoundingMode = 0,
729 int32_t RoundingAccuracy = 0) noexcept;
731 template <
int WA,
int Wout>
733 __spirv_ArbitraryFloatLog1pINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
734 int32_t Mout, int32_t EnableSubnormals = 0,
735 int32_t RoundingMode = 0,
736 int32_t RoundingAccuracy = 0) noexcept;
738 template <
int WA,
int Wout>
740 __spirv_ArbitraryFloatExpINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
741 int32_t Mout, int32_t EnableSubnormals = 0,
742 int32_t RoundingMode = 0,
743 int32_t RoundingAccuracy = 0) noexcept;
745 template <
int WA,
int Wout>
747 __spirv_ArbitraryFloatExp2INTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
748 int32_t Mout, int32_t EnableSubnormals = 0,
749 int32_t RoundingMode = 0,
750 int32_t RoundingAccuracy = 0) noexcept;
752 template <
int WA,
int Wout>
754 __spirv_ArbitraryFloatExp10INTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
755 int32_t Mout, int32_t EnableSubnormals = 0,
756 int32_t RoundingMode = 0,
757 int32_t RoundingAccuracy = 0) noexcept;
759 template <
int WA,
int Wout>
761 __spirv_ArbitraryFloatExpm1INTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
762 int32_t Mout, int32_t EnableSubnormals = 0,
763 int32_t RoundingMode = 0,
764 int32_t RoundingAccuracy = 0) noexcept;
766 template <
int WA,
int Wout>
768 __spirv_ArbitraryFloatSinINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
769 int32_t Mout, int32_t EnableSubnormals = 0,
770 int32_t RoundingMode = 0,
771 int32_t RoundingAccuracy = 0) noexcept;
773 template <
int WA,
int Wout>
775 __spirv_ArbitraryFloatCosINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
776 int32_t Mout, int32_t EnableSubnormals = 0,
777 int32_t RoundingMode = 0,
778 int32_t RoundingAccuracy = 0) noexcept;
782 template <
int WA,
int Wout>
784 __spirv_ArbitraryFloatSinCosINTEL(
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;
789 template <
int WA,
int Wout>
791 __spirv_ArbitraryFloatSinPiINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
792 int32_t Mout, int32_t EnableSubnormals = 0,
793 int32_t RoundingMode = 0,
794 int32_t RoundingAccuracy = 0) noexcept;
796 template <
int WA,
int Wout>
798 __spirv_ArbitraryFloatCosPiINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
799 int32_t Mout, int32_t EnableSubnormals = 0,
800 int32_t RoundingMode = 0,
801 int32_t RoundingAccuracy = 0) noexcept;
805 template <
int WA,
int Wout>
807 __spirv_ArbitraryFloatSinCosPiINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
808 int32_t Mout, int32_t EnableSubnormals = 0,
809 int32_t RoundingMode = 0,
810 int32_t RoundingAccuracy = 0) noexcept;
812 template <
int WA,
int Wout>
814 __spirv_ArbitraryFloatASinINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
815 int32_t Mout, int32_t EnableSubnormals = 0,
816 int32_t RoundingMode = 0,
817 int32_t RoundingAccuracy = 0) noexcept;
819 template <
int WA,
int Wout>
821 __spirv_ArbitraryFloatASinPiINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
822 int32_t Mout, int32_t EnableSubnormals = 0,
823 int32_t RoundingMode = 0,
824 int32_t RoundingAccuracy = 0) noexcept;
826 template <
int WA,
int Wout>
828 __spirv_ArbitraryFloatACosINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
829 int32_t Mout, int32_t EnableSubnormals = 0,
830 int32_t RoundingMode = 0,
831 int32_t RoundingAccuracy = 0) noexcept;
833 template <
int WA,
int Wout>
835 __spirv_ArbitraryFloatACosPiINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
836 int32_t Mout, int32_t EnableSubnormals = 0,
837 int32_t RoundingMode = 0,
838 int32_t RoundingAccuracy = 0) noexcept;
840 template <
int WA,
int Wout>
842 __spirv_ArbitraryFloatATanINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
843 int32_t Mout, int32_t EnableSubnormals = 0,
844 int32_t RoundingMode = 0,
845 int32_t RoundingAccuracy = 0) noexcept;
847 template <
int WA,
int Wout>
849 __spirv_ArbitraryFloatATanPiINTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
850 int32_t Mout, int32_t EnableSubnormals = 0,
851 int32_t RoundingMode = 0,
852 int32_t RoundingAccuracy = 0) noexcept;
854 template <
int WA,
int WB,
int Wout>
856 __spirv_ArbitraryFloatATan2INTEL(
sycl::detail::ap_int<WA> A, int32_t MA,
857 sycl::detail::ap_int<WB> B, int32_t MB,
858 int32_t Mout, int32_t EnableSubnormals = 0,
859 int32_t RoundingMode = 0,
860 int32_t RoundingAccuracy = 0) noexcept;
862 template <
int WA,
int WB,
int Wout>
864 sycl::detail::ap_int<WA> A, int32_t MA,
sycl::detail::ap_int<WB> B,
865 int32_t MB, int32_t Mout, int32_t EnableSubnormals = 0,
866 int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept;
868 template <
int WA,
int WB,
int Wout>
870 sycl::detail::ap_int<WA> A, int32_t MA,
sycl::detail::ap_int<WB> B,
871 int32_t MB, int32_t Mout, int32_t EnableSubnormals = 0,
872 int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept;
877 template <
int WA,
int WB,
int Wout>
879 sycl::detail::ap_int<WA> A, int32_t MA,
sycl::detail::ap_int<WB> B,
880 bool SignOfB, int32_t Mout, int32_t EnableSubnormals = 0,
881 int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept;
883 template <typename dataT>
884 extern
SYCL_EXTERNAL int32_t __spirv_ReadPipe(__ocl_RPipeTy<dataT> Pipe,
885 dataT *Data, int32_t Size,
887 template <typename dataT>
888 extern
SYCL_EXTERNAL int32_t __spirv_WritePipe(__ocl_WPipeTy<dataT> Pipe,
889 const dataT *Data, int32_t Size,
891 template <typename dataT>
893 __spirv_ReadPipeBlockingINTEL(__ocl_RPipeTy<dataT> Pipe, dataT *Data,
894 int32_t Size, int32_t
Alignment) noexcept;
895 template <typename dataT>
897 __spirv_WritePipeBlockingINTEL(__ocl_WPipeTy<dataT> Pipe, const dataT *Data,
898 int32_t Size, int32_t
Alignment) noexcept;
899 template <typename dataT>
901 __spirv_CreatePipeFromPipeStorage_read(
902 const ConstantPipeStorage *Storage) noexcept;
903 template <typename dataT>
905 __spirv_CreatePipeFromPipeStorage_write(
906 const ConstantPipeStorage *Storage) noexcept;
910 size_t NumBytes) noexcept;
912 extern
SYCL_EXTERNAL uint16_t __spirv_ConvertFToBF16INTEL(
float) noexcept;
913 extern
SYCL_EXTERNAL float __spirv_ConvertBF16ToFINTEL(uint16_t) noexcept;
916 __spirv_GroupNonUniformBallot(uint32_t Execution,
bool Predicate) noexcept;
919 __clc_BarrierInitialize(
int64_t *state, int32_t expected_count) noexcept;
922 __clc_BarrierInvalidate(
int64_t *state) noexcept;
925 __clc_BarrierArrive(
int64_t *state) noexcept;
928 __clc_BarrierArriveAndDrop(
int64_t *state) noexcept;
931 __clc_BarrierArriveNoComplete(
int64_t *state, int32_t count) noexcept;
934 __clc_BarrierArriveAndDropNoComplete(
int64_t *state, int32_t count) noexcept;
937 __clc_BarrierCopyAsyncArrive(
int64_t *state) noexcept;
940 __clc_BarrierCopyAsyncArriveNoInc(
int64_t *state) noexcept;
949 __clc_BarrierArriveAndWait(
int64_t *state) noexcept;
951 #ifdef __SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__
952 template <
typename... Args>
954 __spirv_ocl_printf(
const __attribute__((opencl_constant))
char *Format,
956 template <
typename... Args>
957 extern SYCL_EXTERNAL int __spirv_ocl_printf(
const char *Format, Args... args);
960 __spirv_ocl_printf(
const __attribute__((opencl_constant))
char *Format, ...);
961 extern SYCL_EXTERNAL int __spirv_ocl_printf(
const char *Format, ...);
968 __clc_native_tanh(__ocl_vec_t<float, 2>);
970 __clc_native_tanh(__ocl_vec_t<float, 3>);
972 __clc_native_tanh(__ocl_vec_t<float, 4>);
974 __clc_native_tanh(__ocl_vec_t<float, 8>);
976 __clc_native_tanh(__ocl_vec_t<float, 16>);
980 __clc_native_tanh(__ocl_vec_t<_Float16, 2>);
982 __clc_native_tanh(__ocl_vec_t<_Float16, 3>);
984 __clc_native_tanh(__ocl_vec_t<_Float16, 4>);
986 __clc_native_tanh(__ocl_vec_t<_Float16, 8>);
988 __clc_native_tanh(__ocl_vec_t<_Float16, 16>);
992 __clc_native_exp2(__ocl_vec_t<_Float16, 2>);
994 __clc_native_exp2(__ocl_vec_t<_Float16, 3>);
996 __clc_native_exp2(__ocl_vec_t<_Float16, 4>);
998 __clc_native_exp2(__ocl_vec_t<_Float16, 8>);
1000 __clc_native_exp2(__ocl_vec_t<_Float16, 16>);
1002 #define __CLC_BF16(...) \
1003 extern SYCL_EXTERNAL __SYCL_EXPORT __VA_ARGS__ __clc_fabs( \
1004 __VA_ARGS__) noexcept; \
1005 extern SYCL_EXTERNAL __SYCL_EXPORT __VA_ARGS__ __clc_fmin( \
1006 __VA_ARGS__, __VA_ARGS__) noexcept; \
1007 extern SYCL_EXTERNAL __SYCL_EXPORT __VA_ARGS__ __clc_fmax( \
1008 __VA_ARGS__, __VA_ARGS__) noexcept; \
1009 extern SYCL_EXTERNAL __SYCL_EXPORT __VA_ARGS__ __clc_fma( \
1010 __VA_ARGS__, __VA_ARGS__, __VA_ARGS__) noexcept;
1012 #define __CLC_BF16_SCAL_VEC(TYPE) \
1014 __CLC_BF16(__ocl_vec_t<TYPE, 2>) \
1015 __CLC_BF16(__ocl_vec_t<TYPE, 3>) \
1016 __CLC_BF16(__ocl_vec_t<TYPE, 4>) \
1017 __CLC_BF16(__ocl_vec_t<TYPE, 8>) \
1018 __CLC_BF16(__ocl_vec_t<TYPE, 16>)
1020 __CLC_BF16_SCAL_VEC(uint16_t)
1021 __CLC_BF16_SCAL_VEC(uint32_t)
1023 #undef __CLC_BF16_SCAL_VEC
1026 #else // if !__SYCL_DEVICE_ONLY__
1028 template <
typename dataT>
1031 dataT *Src,
size_t NumElements,
1033 for (
size_t i = 0; i < NumElements; i++) {
1034 Dest[i] = Src[i * Stride];
1040 template <
typename dataT>
1043 dataT *Src,
size_t NumElements,
1045 for (
size_t i = 0; i < NumElements; i++) {
1046 Dest[i * Stride] = Src[i];
1053 size_t NumBytes) noexcept;
1057 uint32_t Semantics) noexcept;
1066 #endif // !__SYCL_DEVICE_ONLY__