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 #include <CL/__spirv/spirv_types.hpp>
11 #include <cstddef>
12 #include <cstdint>
13 #include <sycl/detail/defines.hpp>
14 #include <sycl/detail/export.hpp>
15 #include <sycl/detail/stl_type_traits.hpp>
16 
17 // Convergent attribute
18 #ifdef __SYCL_DEVICE_ONLY__
19 #define __SYCL_CONVERGENT__ __attribute__((convergent))
20 #else
21 #define __SYCL_CONVERGENT__
22 #endif
23 
24 #ifdef __SYCL_DEVICE_ONLY__
25 
26 #if (SYCL_EXT_ONEAPI_MATRIX_VERSION > 1)
27 template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
28  __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
29  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
30 extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *
31 __spirv_JointMatrixLoadINTEL(T *Ptr, std::size_t Stride,
32  __spv::MatrixLayout Layout = L,
33  __spv::Scope::Flag Sc = S, int MemOperand = 0);
34 
35 template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
36  __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
37  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
38 extern SYCL_EXTERNAL void __spirv_JointMatrixStoreINTEL(
39  T *Ptr, __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *Object,
40  std::size_t Stride, __spv::MatrixLayout Layout = L,
41  __spv::Scope::Flag Sc = S, int MemOperand = 0);
42 
43 template <typename T1, typename T2, std::size_t M, std::size_t K, std::size_t N,
44  __spv::MatrixUse UA, __spv::MatrixUse UB, __spv::MatrixUse UC,
45  __spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
46  __spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
47  __spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
48  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
49 extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T2, M, N, LC, S, UC> *
50 __spirv_JointMatrixMadINTEL(
51  __spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S, UA> *A,
52  __spv::__spirv_JointMatrixINTEL<T1, K, N, LB, S, UB> *B,
53  __spv::__spirv_JointMatrixINTEL<T2, M, N, LC, S, UC> *C,
54  __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
55 
56 template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
57  std::size_t N, __spv::MatrixUse UA, __spv::MatrixUse UB,
58  __spv::MatrixUse UC,
59  __spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
60  __spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
61  __spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
62  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
63 extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *
64 __spirv_JointMatrixUUMadINTEL(
65  __spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S, UA> *A,
66  __spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S, UB> *B,
67  __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *C,
68  __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
69 
70 template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
71  std::size_t N, __spv::MatrixUse UA, __spv::MatrixUse UB,
72  __spv::MatrixUse UC,
73  __spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
74  __spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
75  __spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
76  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
77 extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *
78 __spirv_JointMatrixUSMadINTEL(
79  __spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S, UA> *A,
80  __spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S, UB> *B,
81  __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *C,
82  __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
83 
84 template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
85  std::size_t N, __spv::MatrixUse UA, __spv::MatrixUse UB,
86  __spv::MatrixUse UC,
87  __spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
88  __spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
89  __spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
90  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
91 extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *
92 __spirv_JointMatrixSUMadINTEL(
93  __spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S, UA> *A,
94  __spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S, UB> *B,
95  __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *C,
96  __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
97 
98 template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
99  __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
100  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
101 extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *
102 __spirv_CompositeConstruct(const T v);
103 
104 template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
105  __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
106  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
107 extern SYCL_EXTERNAL size_t __spirv_JointMatrixWorkItemLengthINTEL(
108  __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *);
109 
110 template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
111  __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
112  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
113 extern SYCL_EXTERNAL T __spirv_VectorExtractDynamic(
114  __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *, size_t i);
115 
116 template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
117  __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
118  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
119 extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *
120 __spirv_VectorInsertDynamic(__spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *,
121  T val, size_t i);
122 #else
123 template <typename T, std::size_t R, std::size_t C,
124  __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
125  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
126 extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *
127 __spirv_JointMatrixLoadINTEL(T *Ptr, std::size_t Stride,
128  __spv::MatrixLayout Layout = L,
129  __spv::Scope::Flag Sc = S, int MemOperand = 0);
130 
131 template <typename T, std::size_t R, std::size_t C,
132  __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
133  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
134 extern SYCL_EXTERNAL void __spirv_JointMatrixStoreINTEL(
135  T *Ptr, __spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *Object,
136  std::size_t Stride, __spv::MatrixLayout Layout = L,
137  __spv::Scope::Flag Sc = S, int MemOperand = 0);
138 
139 template <typename T1, typename T2, std::size_t M, std::size_t K, std::size_t N,
140  __spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
141  __spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
142  __spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
143  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
144 extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T2, M, N, LC, S> *
145 __spirv_JointMatrixMadINTEL(
146  __spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S> *A,
147  __spv::__spirv_JointMatrixINTEL<T1, K, N, LB, S> *B,
148  __spv::__spirv_JointMatrixINTEL<T2, M, N, LC, S> *C,
149  __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
150 
151 template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
152  std::size_t N, __spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
153  __spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
154  __spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
155  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
156 extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S> *
157 __spirv_JointMatrixUUMadINTEL(
158  __spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S> *A,
159  __spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S> *B,
160  __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S> *C,
161  __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
162 
163 template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
164  std::size_t N, __spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
165  __spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
166  __spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
167  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
168 extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S> *
169 __spirv_JointMatrixUSMadINTEL(
170  __spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S> *A,
171  __spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S> *B,
172  __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S> *C,
173  __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
174 
175 template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
176  std::size_t N, __spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
177  __spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
178  __spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
179  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
180 extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S> *
181 __spirv_JointMatrixSUMadINTEL(
182  __spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S> *A,
183  __spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S> *B,
184  __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S> *C,
185  __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
186 
187 template <typename T, std::size_t R, std::size_t C,
188  __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
189  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
190 extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *
191 __spirv_CompositeConstruct(const T v);
192 
193 template <typename T, std::size_t R, std::size_t C,
194  __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
195  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
196 extern SYCL_EXTERNAL size_t __spirv_JointMatrixWorkItemLengthINTEL(
197  __spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *);
198 
199 template <typename T, std::size_t R, std::size_t C,
200  __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
201  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
202 extern SYCL_EXTERNAL T __spirv_VectorExtractDynamic(
203  __spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *, size_t i);
204 
205 template <typename T, std::size_t R, std::size_t C,
206  __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
207  __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
208 extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *
209 __spirv_VectorInsertDynamic(__spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *,
210  T val, size_t i);
211 #endif // SYCL_EXT_ONEAPI_MATRIX_VERSION
212 
213 #ifndef __SPIRV_BUILTIN_DECLARATIONS__
214 #error \
215  "SPIR-V built-ins are not available. Please set -fdeclare-spirv-builtins flag."
216 #endif
217 
218 template <typename RetT, typename ImageT>
219 extern SYCL_EXTERNAL RetT __spirv_ImageQueryFormat(ImageT);
220 
221 template <typename RetT, typename ImageT>
222 extern SYCL_EXTERNAL RetT __spirv_ImageQueryOrder(ImageT);
223 
224 template <typename RetT, typename ImageT>
225 extern SYCL_EXTERNAL RetT __spirv_ImageQuerySize(ImageT);
226 
227 template <typename ImageT, typename CoordT, typename ValT>
228 extern SYCL_EXTERNAL void __spirv_ImageWrite(ImageT, CoordT, ValT);
229 
230 template <class RetT, typename ImageT, typename TempArgT>
231 extern SYCL_EXTERNAL RetT __spirv_ImageRead(ImageT, TempArgT);
232 
233 template <typename ImageT, typename SampledType>
234 extern SYCL_EXTERNAL SampledType __spirv_SampledImage(ImageT, __ocl_sampler_t);
235 
236 template <typename SampledType, typename TempRetT, typename TempArgT>
237 extern SYCL_EXTERNAL TempRetT __spirv_ImageSampleExplicitLod(SampledType,
238  TempArgT, int,
239  float);
240 
241 #define __SYCL_OpGroupAsyncCopyGlobalToLocal __spirv_GroupAsyncCopy
242 #define __SYCL_OpGroupAsyncCopyLocalToGlobal __spirv_GroupAsyncCopy
243 
244 // Atomic SPIR-V builtins
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, \
252  Type V);
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, \
256  Type V);
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, \
264  Type V);
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, \
268  Type V);
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, \
272  Type V);
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, \
276  Type V);
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, \
280  Type V);
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, \
284  Type V);
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, \
288  Type V);
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, \
292  Type V);
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, \
296  Type V);
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, \
300  Type V);
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, \
304  Type V);
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, \
308  Type V);
309 
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)
317 
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)
326 
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)
331 
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)
336 
337 // Helper atomic operations which select correct signed/unsigned version
338 // of atomic min/max based on the 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, \
345  T Value) { \
346  return __spirv_AtomicS##Op(Ptr, Memory, Semantics, Value); \
347  } \
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, \
353  T Value) { \
354  return __spirv_AtomicU##Op(Ptr, Memory, Semantics, Value); \
355  } \
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, \
360  T Value) { \
361  return __spirv_AtomicF##Op##EXT(Ptr, Memory, Semantics, Value); \
362  }
363 
364 #define __SPIRV_ATOMICS(macro, Arg) \
365  macro(__attribute__((opencl_global)), Arg) \
366  macro(__attribute__((opencl_local)), Arg) macro(, Arg)
367 
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)
378 
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
401 
402 template <typename dataT>
403 extern __attribute__((opencl_global)) dataT *
404 __SYCL_GenericCastToPtrExplicit_ToGlobal(void *Ptr) noexcept {
405  return (__attribute__((opencl_global)) dataT *)
406  __spirv_GenericCastToPtrExplicit_ToGlobal(
407  Ptr, __spv::StorageClass::CrossWorkgroup);
408 }
409 
410 template <typename dataT>
411 extern const __attribute__((opencl_global)) dataT *
412 __SYCL_GenericCastToPtrExplicit_ToGlobal(const void *Ptr) noexcept {
413  return (const __attribute__((opencl_global)) dataT *)
414  __spirv_GenericCastToPtrExplicit_ToGlobal(
415  Ptr, __spv::StorageClass::CrossWorkgroup);
416 }
417 
418 template <typename dataT>
419 extern volatile __attribute__((opencl_global)) dataT *
420 __SYCL_GenericCastToPtrExplicit_ToGlobal(volatile void *Ptr) noexcept {
421  return (volatile __attribute__((opencl_global)) dataT *)
422  __spirv_GenericCastToPtrExplicit_ToGlobal(
423  Ptr, __spv::StorageClass::CrossWorkgroup);
424 }
425 
426 template <typename dataT>
427 extern const volatile __attribute__((opencl_global)) dataT *
428 __SYCL_GenericCastToPtrExplicit_ToGlobal(const volatile void *Ptr) noexcept {
429  return (const volatile __attribute__((opencl_global)) dataT *)
430  __spirv_GenericCastToPtrExplicit_ToGlobal(
431  Ptr, __spv::StorageClass::CrossWorkgroup);
432 }
433 
434 template <typename dataT>
435 extern __attribute__((opencl_local)) dataT *
436 __SYCL_GenericCastToPtrExplicit_ToLocal(void *Ptr) noexcept {
437  return (__attribute__((opencl_local)) dataT *)
438  __spirv_GenericCastToPtrExplicit_ToLocal(Ptr,
439  __spv::StorageClass::Workgroup);
440 }
441 
442 template <typename dataT>
443 extern const __attribute__((opencl_local)) dataT *
444 __SYCL_GenericCastToPtrExplicit_ToLocal(const void *Ptr) noexcept {
445  return (const __attribute__((opencl_local)) dataT *)
446  __spirv_GenericCastToPtrExplicit_ToLocal(Ptr,
447  __spv::StorageClass::Workgroup);
448 }
449 
450 template <typename dataT>
451 extern volatile __attribute__((opencl_local)) dataT *
452 __SYCL_GenericCastToPtrExplicit_ToLocal(volatile void *Ptr) noexcept {
453  return (volatile __attribute__((opencl_local)) dataT *)
454  __spirv_GenericCastToPtrExplicit_ToLocal(Ptr,
455  __spv::StorageClass::Workgroup);
456 }
457 
458 template <typename dataT>
459 extern const volatile __attribute__((opencl_local)) dataT *
460 __SYCL_GenericCastToPtrExplicit_ToLocal(const volatile void *Ptr) noexcept {
461  return (const volatile __attribute__((opencl_local)) dataT *)
462  __spirv_GenericCastToPtrExplicit_ToLocal(Ptr,
463  __spv::StorageClass::Workgroup);
464 }
465 
466 template <typename dataT>
467 extern __attribute__((opencl_private)) dataT *
468 __SYCL_GenericCastToPtrExplicit_ToPrivate(void *Ptr) noexcept {
469  return (__attribute__((opencl_private)) dataT *)
470  __spirv_GenericCastToPtrExplicit_ToPrivate(Ptr,
471  __spv::StorageClass::Function);
472 }
473 
474 template <typename dataT>
475 extern const __attribute__((opencl_private)) dataT *
476 __SYCL_GenericCastToPtrExplicit_ToPrivate(const void *Ptr) noexcept {
477  return (const __attribute__((opencl_private)) dataT *)
478  __spirv_GenericCastToPtrExplicit_ToPrivate(Ptr,
479  __spv::StorageClass::Function);
480 }
481 
482 template <typename dataT>
483 extern volatile __attribute__((opencl_private)) dataT *
484 __SYCL_GenericCastToPtrExplicit_ToPrivate(volatile void *Ptr) noexcept {
485  return (volatile __attribute__((opencl_private)) dataT *)
486  __spirv_GenericCastToPtrExplicit_ToPrivate(Ptr,
487  __spv::StorageClass::Function);
488 }
489 
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,
495  __spv::StorageClass::Function);
496 }
497 
498 template <typename dataT>
499 __SYCL_CONVERGENT__ extern SYCL_EXTERNAL dataT
500 __spirv_SubgroupShuffleINTEL(dataT Data, uint32_t InvocationId) noexcept;
501 template <typename dataT>
502 __SYCL_CONVERGENT__ extern SYCL_EXTERNAL dataT __spirv_SubgroupShuffleDownINTEL(
503  dataT Current, dataT Next, uint32_t Delta) noexcept;
504 template <typename dataT>
505 __SYCL_CONVERGENT__ extern SYCL_EXTERNAL dataT __spirv_SubgroupShuffleUpINTEL(
506  dataT Previous, dataT Current, uint32_t Delta) noexcept;
507 template <typename dataT>
508 __SYCL_CONVERGENT__ extern SYCL_EXTERNAL dataT
509 __spirv_SubgroupShuffleXorINTEL(dataT Data, uint32_t Value) noexcept;
510 
511 template <typename dataT>
512 __SYCL_CONVERGENT__ extern SYCL_EXTERNAL dataT __spirv_SubgroupBlockReadINTEL(
513  const __attribute__((opencl_global)) uint8_t *Ptr) noexcept;
514 
515 template <typename dataT>
516 __SYCL_CONVERGENT__ extern SYCL_EXTERNAL void
517 __spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_global)) uint8_t *Ptr,
518  dataT Data) noexcept;
519 
520 template <typename dataT>
521 __SYCL_CONVERGENT__ extern SYCL_EXTERNAL dataT __spirv_SubgroupBlockReadINTEL(
522  const __attribute__((opencl_global)) uint16_t *Ptr) noexcept;
523 
524 template <typename dataT>
525 __SYCL_CONVERGENT__ extern SYCL_EXTERNAL void
526 __spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_global)) uint16_t *Ptr,
527  dataT Data) noexcept;
528 
529 template <typename dataT>
530 __SYCL_CONVERGENT__ extern SYCL_EXTERNAL dataT __spirv_SubgroupBlockReadINTEL(
531  const __attribute__((opencl_global)) uint32_t *Ptr) noexcept;
532 
533 template <typename dataT>
534 __SYCL_CONVERGENT__ extern SYCL_EXTERNAL void
535 __spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_global)) uint32_t *Ptr,
536  dataT Data) noexcept;
537 
538 template <typename dataT>
539 __SYCL_CONVERGENT__ extern SYCL_EXTERNAL dataT __spirv_SubgroupBlockReadINTEL(
540  const __attribute__((opencl_global)) uint64_t *Ptr) noexcept;
541 
542 template <typename dataT>
543 __SYCL_CONVERGENT__ extern SYCL_EXTERNAL void
544 __spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_global)) uint64_t *Ptr,
545  dataT Data) noexcept;
546 template <int W, int rW>
547 extern SYCL_EXTERNAL sycl::detail::ap_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>
551 extern SYCL_EXTERNAL sycl::detail::ap_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>
556 extern SYCL_EXTERNAL sycl::detail::ap_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>
561 extern SYCL_EXTERNAL sycl::detail::ap_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>
565 extern SYCL_EXTERNAL sycl::detail::ap_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>
569 extern SYCL_EXTERNAL sycl::detail::ap_int<2 * 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>
574 extern SYCL_EXTERNAL sycl::detail::ap_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>
579 extern SYCL_EXTERNAL sycl::detail::ap_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>
584 extern SYCL_EXTERNAL sycl::detail::ap_int<2 * 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>
589 extern SYCL_EXTERNAL sycl::detail::ap_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>
593 extern SYCL_EXTERNAL sycl::detail::ap_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;
596 
597 // In the following built-ins width of arbitrary precision integer type for
598 // a floating point variable should be equal to sum of corresponding
599 // exponent width E, mantissa width M and 1 for sign bit. I.e. WA = EA + MA + 1.
600 template <int WA, int Wout>
601 extern SYCL_EXTERNAL sycl::detail::ap_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;
606 
607 template <int WA, int Wout>
608 extern SYCL_EXTERNAL sycl::detail::ap_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;
614 
615 template <int WA, int Wout>
616 extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
617 __spirv_ArbitraryFloatCastToIntINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
618  bool ToSign = false,
619  int32_t EnableSubnormals = 0,
620  int32_t RoundingMode = 0,
621  int32_t RoundingAccuracy = 0) noexcept;
622 
623 template <int WA, int WB, int Wout>
624 extern SYCL_EXTERNAL sycl::detail::ap_int<Wout> __spirv_ArbitraryFloatAddINTEL(
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;
628 
629 template <int WA, int WB, int Wout>
630 extern SYCL_EXTERNAL sycl::detail::ap_int<Wout> __spirv_ArbitraryFloatSubINTEL(
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;
634 
635 template <int WA, int WB, int Wout>
636 extern SYCL_EXTERNAL sycl::detail::ap_int<Wout> __spirv_ArbitraryFloatMulINTEL(
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;
640 
641 template <int WA, int WB, int Wout>
642 extern SYCL_EXTERNAL sycl::detail::ap_int<Wout> __spirv_ArbitraryFloatDivINTEL(
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;
646 
647 // Comparison built-ins don't use Subnormal Support, Rounding Mode and
648 // Rounding Accuracy.
649 template <int WA, int WB>
650 extern SYCL_EXTERNAL bool
651 __spirv_ArbitraryFloatGTINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
652  sycl::detail::ap_int<WB> B, int32_t MB) noexcept;
653 
654 template <int WA, int WB>
655 extern SYCL_EXTERNAL bool
656 __spirv_ArbitraryFloatGEINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
657  sycl::detail::ap_int<WB> B, int32_t MB) noexcept;
658 
659 template <int WA, int WB>
660 extern SYCL_EXTERNAL bool
661 __spirv_ArbitraryFloatLTINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
662  sycl::detail::ap_int<WB> B, int32_t MB) noexcept;
663 
664 template <int WA, int WB>
665 extern SYCL_EXTERNAL bool
666 __spirv_ArbitraryFloatLEINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
667  sycl::detail::ap_int<WB> B, int32_t MB) noexcept;
668 
669 template <int WA, int WB>
670 extern SYCL_EXTERNAL bool
671 __spirv_ArbitraryFloatEQINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
672  sycl::detail::ap_int<WB> B, int32_t MB) noexcept;
673 
674 template <int WA, int Wout>
675 extern SYCL_EXTERNAL sycl::detail::ap_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;
680 
681 template <int WA, int Wout>
682 extern SYCL_EXTERNAL sycl::detail::ap_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;
687 
688 template <int WA, int Wout>
689 extern SYCL_EXTERNAL sycl::detail::ap_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;
694 
695 template <int WA, int WB, int Wout>
696 extern SYCL_EXTERNAL sycl::detail::ap_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;
702 
703 template <int WA, int Wout>
704 extern SYCL_EXTERNAL sycl::detail::ap_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;
709 
710 template <int WA, int Wout>
711 extern SYCL_EXTERNAL sycl::detail::ap_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;
716 
717 template <int WA, int Wout>
718 extern SYCL_EXTERNAL sycl::detail::ap_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;
723 
724 template <int WA, int Wout>
725 extern SYCL_EXTERNAL sycl::detail::ap_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;
730 
731 template <int WA, int Wout>
732 extern SYCL_EXTERNAL sycl::detail::ap_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;
737 
738 template <int WA, int Wout>
739 extern SYCL_EXTERNAL sycl::detail::ap_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;
744 
745 template <int WA, int Wout>
746 extern SYCL_EXTERNAL sycl::detail::ap_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;
751 
752 template <int WA, int Wout>
753 extern SYCL_EXTERNAL sycl::detail::ap_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;
758 
759 template <int WA, int Wout>
760 extern SYCL_EXTERNAL sycl::detail::ap_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;
765 
766 template <int WA, int Wout>
767 extern SYCL_EXTERNAL sycl::detail::ap_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;
772 
773 template <int WA, int Wout>
774 extern SYCL_EXTERNAL sycl::detail::ap_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;
779 
780 // Result value contains both values of sine and cosine and so has the size of
781 // 2 * Wout where Wout is equal to (1 + Eout + Mout).
782 template <int WA, int Wout>
783 extern SYCL_EXTERNAL sycl::detail::ap_int<2 * 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;
788 
789 template <int WA, int Wout>
790 extern SYCL_EXTERNAL sycl::detail::ap_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;
795 
796 template <int WA, int Wout>
797 extern SYCL_EXTERNAL sycl::detail::ap_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;
802 
803 // Result value contains both values of sine(A*pi) and cosine(A*pi) and so has
804 // the size of 2 * Wout where Wout is equal to (1 + Eout + Mout).
805 template <int WA, int Wout>
806 extern SYCL_EXTERNAL sycl::detail::ap_int<2 * 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;
811 
812 template <int WA, int Wout>
813 extern SYCL_EXTERNAL sycl::detail::ap_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;
818 
819 template <int WA, int Wout>
820 extern SYCL_EXTERNAL sycl::detail::ap_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;
825 
826 template <int WA, int Wout>
827 extern SYCL_EXTERNAL sycl::detail::ap_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;
832 
833 template <int WA, int Wout>
834 extern SYCL_EXTERNAL sycl::detail::ap_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;
839 
840 template <int WA, int Wout>
841 extern SYCL_EXTERNAL sycl::detail::ap_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;
846 
847 template <int WA, int Wout>
848 extern SYCL_EXTERNAL sycl::detail::ap_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;
853 
854 template <int WA, int WB, int Wout>
855 extern SYCL_EXTERNAL sycl::detail::ap_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;
861 
862 template <int WA, int WB, int Wout>
863 extern SYCL_EXTERNAL sycl::detail::ap_int<Wout> __spirv_ArbitraryFloatPowINTEL(
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;
867 
868 template <int WA, int WB, int Wout>
869 extern SYCL_EXTERNAL sycl::detail::ap_int<Wout> __spirv_ArbitraryFloatPowRINTEL(
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;
873 
874 // PowN built-in calculates `A^B` where `A` is arbitrary precision floating
875 // point number and `B` is signed or unsigned arbitrary precision integer,
876 // i.e. its width doesn't depend on sum of exponent and mantissa.
877 template <int WA, int WB, int Wout>
878 extern SYCL_EXTERNAL sycl::detail::ap_int<Wout> __spirv_ArbitraryFloatPowNINTEL(
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;
882 
883 template <typename dataT>
884 extern SYCL_EXTERNAL int32_t __spirv_ReadPipe(__ocl_RPipeTy<dataT> Pipe,
885  dataT *Data, int32_t Size,
886  int32_t Alignment) noexcept;
887 template <typename dataT>
888 extern SYCL_EXTERNAL int32_t __spirv_WritePipe(__ocl_WPipeTy<dataT> Pipe,
889  const dataT *Data, int32_t Size,
890  int32_t Alignment) noexcept;
891 template <typename dataT>
892 extern SYCL_EXTERNAL void
893 __spirv_ReadPipeBlockingINTEL(__ocl_RPipeTy<dataT> Pipe, dataT *Data,
894  int32_t Size, int32_t Alignment) noexcept;
895 template <typename dataT>
896 extern SYCL_EXTERNAL void
897 __spirv_WritePipeBlockingINTEL(__ocl_WPipeTy<dataT> Pipe, const dataT *Data,
898  int32_t Size, int32_t Alignment) noexcept;
899 template <typename dataT>
900 extern SYCL_EXTERNAL __ocl_RPipeTy<dataT>
901 __spirv_CreatePipeFromPipeStorage_read(
902  const ConstantPipeStorage *Storage) noexcept;
903 template <typename dataT>
904 extern SYCL_EXTERNAL __ocl_WPipeTy<dataT>
905 __spirv_CreatePipeFromPipeStorage_write(
906  const ConstantPipeStorage *Storage) noexcept;
907 
908 extern SYCL_EXTERNAL void
909 __spirv_ocl_prefetch(const __attribute__((opencl_global)) char *Ptr,
910  size_t NumBytes) noexcept;
911 
912 extern SYCL_EXTERNAL uint16_t __spirv_ConvertFToBF16INTEL(float) noexcept;
913 extern SYCL_EXTERNAL float __spirv_ConvertBF16ToFINTEL(uint16_t) noexcept;
914 
915 __SYCL_CONVERGENT__ extern SYCL_EXTERNAL __SYCL_EXPORT __ocl_vec_t<uint32_t, 4>
916 __spirv_GroupNonUniformBallot(uint32_t Execution, bool Predicate) noexcept;
917 
918 extern SYCL_EXTERNAL __SYCL_EXPORT void
919 __clc_BarrierInitialize(int64_t *state, int32_t expected_count) noexcept;
920 
921 extern SYCL_EXTERNAL __SYCL_EXPORT void
922 __clc_BarrierInvalidate(int64_t *state) noexcept;
923 
924 extern SYCL_EXTERNAL __SYCL_EXPORT int64_t
925 __clc_BarrierArrive(int64_t *state) noexcept;
926 
927 extern SYCL_EXTERNAL __SYCL_EXPORT int64_t
928 __clc_BarrierArriveAndDrop(int64_t *state) noexcept;
929 
930 extern SYCL_EXTERNAL __SYCL_EXPORT int64_t
931 __clc_BarrierArriveNoComplete(int64_t *state, int32_t count) noexcept;
932 
933 extern SYCL_EXTERNAL __SYCL_EXPORT int64_t
934 __clc_BarrierArriveAndDropNoComplete(int64_t *state, int32_t count) noexcept;
935 
936 extern SYCL_EXTERNAL __SYCL_EXPORT void
937 __clc_BarrierCopyAsyncArrive(int64_t *state) noexcept;
938 
939 extern SYCL_EXTERNAL __SYCL_EXPORT void
940 __clc_BarrierCopyAsyncArriveNoInc(int64_t *state) noexcept;
941 
942 __SYCL_CONVERGENT__ extern SYCL_EXTERNAL __SYCL_EXPORT void
943 __clc_BarrierWait(int64_t *state, int64_t arrival) noexcept;
944 
945 extern SYCL_EXTERNAL __SYCL_EXPORT bool
946 __clc_BarrierTestWait(int64_t *state, int64_t arrival) noexcept;
947 
948 __SYCL_CONVERGENT__ extern SYCL_EXTERNAL __SYCL_EXPORT void
949 __clc_BarrierArriveAndWait(int64_t *state) noexcept;
950 
951 #ifdef __SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__
952 template <typename... Args>
953 extern SYCL_EXTERNAL int
954 __spirv_ocl_printf(const __attribute__((opencl_constant)) char *Format,
955  Args... args);
956 template <typename... Args>
957 extern SYCL_EXTERNAL int __spirv_ocl_printf(const char *Format, Args... args);
958 #else
959 extern SYCL_EXTERNAL int
960 __spirv_ocl_printf(const __attribute__((opencl_constant)) char *Format, ...);
961 extern SYCL_EXTERNAL int __spirv_ocl_printf(const char *Format, ...);
962 #endif
963 
964 // Native builtin extension
965 
966 extern SYCL_EXTERNAL float __clc_native_tanh(float);
967 extern SYCL_EXTERNAL __ocl_vec_t<float, 2>
968  __clc_native_tanh(__ocl_vec_t<float, 2>);
969 extern SYCL_EXTERNAL __ocl_vec_t<float, 3>
970  __clc_native_tanh(__ocl_vec_t<float, 3>);
971 extern SYCL_EXTERNAL __ocl_vec_t<float, 4>
972  __clc_native_tanh(__ocl_vec_t<float, 4>);
973 extern SYCL_EXTERNAL __ocl_vec_t<float, 8>
974  __clc_native_tanh(__ocl_vec_t<float, 8>);
975 extern SYCL_EXTERNAL __ocl_vec_t<float, 16>
976  __clc_native_tanh(__ocl_vec_t<float, 16>);
977 
978 extern SYCL_EXTERNAL _Float16 __clc_native_tanh(_Float16);
979 extern SYCL_EXTERNAL __ocl_vec_t<_Float16, 2>
980  __clc_native_tanh(__ocl_vec_t<_Float16, 2>);
981 extern SYCL_EXTERNAL __ocl_vec_t<_Float16, 3>
982  __clc_native_tanh(__ocl_vec_t<_Float16, 3>);
983 extern SYCL_EXTERNAL __ocl_vec_t<_Float16, 4>
984  __clc_native_tanh(__ocl_vec_t<_Float16, 4>);
985 extern SYCL_EXTERNAL __ocl_vec_t<_Float16, 8>
986  __clc_native_tanh(__ocl_vec_t<_Float16, 8>);
987 extern SYCL_EXTERNAL __ocl_vec_t<_Float16, 16>
988  __clc_native_tanh(__ocl_vec_t<_Float16, 16>);
989 
990 extern SYCL_EXTERNAL _Float16 __clc_native_exp2(_Float16);
991 extern SYCL_EXTERNAL __ocl_vec_t<_Float16, 2>
992  __clc_native_exp2(__ocl_vec_t<_Float16, 2>);
993 extern SYCL_EXTERNAL __ocl_vec_t<_Float16, 3>
994  __clc_native_exp2(__ocl_vec_t<_Float16, 3>);
995 extern SYCL_EXTERNAL __ocl_vec_t<_Float16, 4>
996  __clc_native_exp2(__ocl_vec_t<_Float16, 4>);
997 extern SYCL_EXTERNAL __ocl_vec_t<_Float16, 8>
998  __clc_native_exp2(__ocl_vec_t<_Float16, 8>);
999 extern SYCL_EXTERNAL __ocl_vec_t<_Float16, 16>
1000  __clc_native_exp2(__ocl_vec_t<_Float16, 16>);
1001 
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;
1011 
1012 #define __CLC_BF16_SCAL_VEC(TYPE) \
1013  __CLC_BF16(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>)
1019 
1020 __CLC_BF16_SCAL_VEC(uint16_t)
1021 __CLC_BF16_SCAL_VEC(uint32_t)
1022 
1023 #undef __CLC_BF16_SCAL_VEC
1024 #undef __CLC_BF16
1025 
1026 #else // if !__SYCL_DEVICE_ONLY__
1027 
1028 template <typename dataT>
1029 __SYCL_CONVERGENT__ extern __ocl_event_t
1030 __SYCL_OpGroupAsyncCopyGlobalToLocal(__spv::Scope::Flag, dataT *Dest,
1031  dataT *Src, size_t NumElements,
1032  size_t Stride, __ocl_event_t) noexcept {
1033  for (size_t i = 0; i < NumElements; i++) {
1034  Dest[i] = Src[i * Stride];
1035  }
1036  // A real instance of the class is not needed, return dummy pointer.
1037  return nullptr;
1038 }
1039 
1040 template <typename dataT>
1041 __SYCL_CONVERGENT__ extern __ocl_event_t
1042 __SYCL_OpGroupAsyncCopyLocalToGlobal(__spv::Scope::Flag, dataT *Dest,
1043  dataT *Src, size_t NumElements,
1044  size_t Stride, __ocl_event_t) noexcept {
1045  for (size_t i = 0; i < NumElements; i++) {
1046  Dest[i * Stride] = Src[i];
1047  }
1048  // A real instance of the class is not needed, return dummy pointer.
1049  return nullptr;
1050 }
1051 
1052 extern __SYCL_EXPORT void __spirv_ocl_prefetch(const char *Ptr,
1053  size_t NumBytes) noexcept;
1054 
1055 __SYCL_CONVERGENT__ extern SYCL_EXTERNAL __SYCL_EXPORT void
1056 __spirv_ControlBarrier(__spv::Scope Execution, __spv::Scope Memory,
1057  uint32_t Semantics) noexcept;
1058 
1059 __SYCL_CONVERGENT__ extern SYCL_EXTERNAL __SYCL_EXPORT void
1060 __spirv_MemoryBarrier(__spv::Scope Memory, uint32_t Semantics) noexcept;
1061 
1062 __SYCL_CONVERGENT__ extern SYCL_EXTERNAL __SYCL_EXPORT void
1063 __spirv_GroupWaitEvents(__spv::Scope Execution, uint32_t NumEvents,
1064  __ocl_event_t *WaitEvents) noexcept;
1065 
1066 #endif // !__SYCL_DEVICE_ONLY__
__spv::StorageClass::Function
@ Function
Definition: spirv_types.hpp:52
__spirv_GroupWaitEvents
__SYCL_CONVERGENT__ SYCL_EXTERNAL void __spirv_GroupWaitEvents(__spv::Scope Execution, uint32_t NumEvents, __ocl_event_t *WaitEvents) noexcept
Definition: spirv_ops.cpp:18
__spirv_ControlBarrier
__SYCL_CONVERGENT__ SYCL_EXTERNAL void __spirv_ControlBarrier(__spv::Scope Execution, __spv::Scope Memory, uint32_t Semantics) noexcept
Definition: spirv_ops.cpp:26
__spv::MatrixLayout::RowMajor
@ RowMajor
__SYCL_OpGroupAsyncCopyGlobalToLocal
__SYCL_CONVERGENT__ __ocl_event_t __SYCL_OpGroupAsyncCopyGlobalToLocal(__spv::Scope::Flag, dataT *Dest, dataT *Src, size_t NumElements, size_t Stride, __ocl_event_t) noexcept
Definition: spirv_ops.hpp:1030
sycl::_V1::detail::int64_t
conditional< sizeof(long)==8, long, long long >::type int64_t
Definition: kernel_desc.hpp:34
__spirv_ocl_prefetch
void __spirv_ocl_prefetch(const char *Ptr, size_t NumBytes) noexcept
Definition: spirv_ops.cpp:47
__spv::__spirv_JointMatrixINTEL
Definition: spirv_types.hpp:135
sycl
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14
__ocl_event_t
void * __ocl_event_t
Definition: spirv_types.hpp:178
__spv::MatrixUse
MatrixUse
Definition: spirv_types.hpp:129
stl_type_traits.hpp
__spv::Scope::Flag
Flag
Definition: spirv_types.hpp:28
export.hpp
defines.hpp
SYCL_EXTERNAL
#define SYCL_EXTERNAL
Definition: defines_elementary.hpp:30
__SYCL_CONVERGENT__
#define __SYCL_CONVERGENT__
Definition: spirv_ops.hpp:21
sycl::_V1::ext::oneapi::experimental::detail::Alignment
@ Alignment
Definition: property.hpp:189
__spirv_MemoryBarrier
__SYCL_CONVERGENT__ SYCL_EXTERNAL void __spirv_MemoryBarrier(__spv::Scope Memory, uint32_t Semantics) noexcept
Definition: spirv_ops.cpp:36
__SYCL_OpGroupAsyncCopyLocalToGlobal
__SYCL_CONVERGENT__ __ocl_event_t __SYCL_OpGroupAsyncCopyLocalToGlobal(__spv::Scope::Flag, dataT *Dest, dataT *Src, size_t NumElements, size_t Stride, __ocl_event_t) noexcept
Definition: spirv_ops.hpp:1042
__spv::StorageClass::CrossWorkgroup
@ CrossWorkgroup
Definition: spirv_types.hpp:50
__spv::StorageClass::Workgroup
@ Workgroup
Definition: spirv_types.hpp:49
__spv::MatrixLayout
MatrixLayout
Definition: spirv_types.hpp:113
__spv::Scope
Definition: spirv_types.hpp:26
sycl::_V1::ext::oneapi::experimental::__attribute__
__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: invoke_simd.hpp:332
__ocl_sampler_t
void * __ocl_sampler_t
Definition: spirv_types.hpp:179
spirv_types.hpp