1/*===---- spirv_builtin_vars.h - SPIR-V built-in ---------------------------===
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
10#ifndef __SPIRV_BUILTIN_VARS_H
11#define __SPIRV_BUILTIN_VARS_H
12
13#if __cplusplus >= 201103L
14#define __SPIRV_NOEXCEPT noexcept
15#else
16#define __SPIRV_NOEXCEPT
17#endif
18
19#pragma push_macro("__size_t")
20#pragma push_macro("__uint32_t")
21#pragma push_macro("__uint64_t")
22#define __size_t __SIZE_TYPE__
23#define __uint32_t __UINT32_TYPE__
24
25#define __SPIRV_overloadable __attribute__((overloadable))
26#define __SPIRV_convergent __attribute__((convergent))
27#define __SPIRV_inline __attribute__((always_inline))
28
29#define __global __attribute__((opencl_global))
30#define __local __attribute__((opencl_local))
31#define __private __attribute__((opencl_private))
32#define __constant __attribute__((opencl_constant))
33#ifdef __SYCL_DEVICE_ONLY__
34#define __generic
35#else
36#define __generic __attribute__((opencl_generic))
37#endif
38
39// Check if SPIR-V builtins are supported.
40// As the translator doesn't use the LLVM intrinsics (which would be emitted if
41// we use the SPIR-V builtins) we can't rely on the SPIRV32/SPIRV64 etc macros
42// to establish if we can use the builtin alias. We disable builtin altogether
43// if we do not intent to use the backend. So instead of use target macros, rely
44// on a __has_builtin test.
45#if (__has_builtin(__builtin_spirv_num_workgroups))
46#define __SPIRV_BUILTIN_ALIAS(builtin) \
47 __attribute__((clang_builtin_alias(builtin)))
48#else
49#define __SPIRV_BUILTIN_ALIAS(builtin)
50#endif
51
52// Builtin IDs and sizes
53
54extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_num_workgroups) __size_t
55 __spirv_BuiltInNumWorkgroups(int);
56extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_workgroup_size) __size_t
57 __spirv_BuiltInWorkgroupSize(int);
58extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_workgroup_id) __size_t
59 __spirv_BuiltInWorkgroupId(int);
60extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_local_invocation_id) __size_t
61 __spirv_BuiltInLocalInvocationId(int);
62extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_global_invocation_id) __size_t
63 __spirv_BuiltInGlobalInvocationId(int);
64
65extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_global_size) __size_t
66 __spirv_BuiltInGlobalSize(int);
67extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_global_offset) __size_t
68 __spirv_BuiltInGlobalOffset(int);
69extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_size) __uint32_t
70 __spirv_BuiltInSubgroupSize();
71extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_max_size) __uint32_t
72 __spirv_BuiltInSubgroupMaxSize();
73extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_num_subgroups) __uint32_t
74 __spirv_BuiltInNumSubgroups();
75extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_id) __uint32_t
76 __spirv_BuiltInSubgroupId();
77extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_local_invocation_id)
78 __uint32_t __spirv_BuiltInSubgroupLocalInvocationId();
79
80// OpGenericCastToPtrExplicit
81
82extern __SPIRV_overloadable
83__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)
84__global void *__spirv_GenericCastToPtrExplicit_ToGlobal(__generic void *,
85 int) __SPIRV_NOEXCEPT;
86extern __SPIRV_overloadable
87__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)
88__global const void *
89__spirv_GenericCastToPtrExplicit_ToGlobal(__generic const void *,
90 int) __SPIRV_NOEXCEPT;
91extern __SPIRV_overloadable
92__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)
93__global volatile void *
94__spirv_GenericCastToPtrExplicit_ToGlobal(__generic volatile void *,
95 int) __SPIRV_NOEXCEPT;
96extern __SPIRV_overloadable
97__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)
98__global const volatile void *
99__spirv_GenericCastToPtrExplicit_ToGlobal(__generic const volatile void *,
100 int) __SPIRV_NOEXCEPT;
101extern __SPIRV_overloadable
102__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)
103__local void *__spirv_GenericCastToPtrExplicit_ToLocal(__generic void *,
104 int) __SPIRV_NOEXCEPT;
105extern __SPIRV_overloadable
106__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)
107__local const void *
108__spirv_GenericCastToPtrExplicit_ToLocal(__generic const void *,
109 int) __SPIRV_NOEXCEPT;
110extern __SPIRV_overloadable
111__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)
112__local volatile void *
113__spirv_GenericCastToPtrExplicit_ToLocal(__generic volatile void *,
114 int) __SPIRV_NOEXCEPT;
115extern __SPIRV_overloadable
116__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)
117__local const volatile void *
118__spirv_GenericCastToPtrExplicit_ToLocal(__generic const volatile void *,
119 int) __SPIRV_NOEXCEPT;
120extern __SPIRV_overloadable
121__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)
122__private void *
123__spirv_GenericCastToPtrExplicit_ToPrivate(__generic void *,
124 int) __SPIRV_NOEXCEPT;
125extern __SPIRV_overloadable
126__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)
127__private const void *
128__spirv_GenericCastToPtrExplicit_ToPrivate(__generic const void *,
129 int) __SPIRV_NOEXCEPT;
130extern __SPIRV_overloadable
131__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)
132__private volatile void *
133__spirv_GenericCastToPtrExplicit_ToPrivate(__generic volatile void *,
134 int) __SPIRV_NOEXCEPT;
135extern __SPIRV_overloadable
136__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)
137__private const volatile void *
138__spirv_GenericCastToPtrExplicit_ToPrivate(__generic const volatile void *,
139 int) __SPIRV_NOEXCEPT;
140
141// OpGenericCastToPtr
142
143static __SPIRV_overloadable __SPIRV_inline __global void *
144__spirv_GenericCastToPtr_ToGlobal(__generic void *p, int) __SPIRV_NOEXCEPT {
145 return (__global void *)p;
146}
147static __SPIRV_overloadable __SPIRV_inline __global const void *
148__spirv_GenericCastToPtr_ToGlobal(__generic const void *p,
149 int) __SPIRV_NOEXCEPT {
150 return (__global const void *)p;
151}
152static __SPIRV_overloadable __SPIRV_inline __global volatile void *
153__spirv_GenericCastToPtr_ToGlobal(__generic volatile void *p,
154 int) __SPIRV_NOEXCEPT {
155 return (__global volatile void *)p;
156}
157static __SPIRV_overloadable __SPIRV_inline __global const volatile void *
158__spirv_GenericCastToPtr_ToGlobal(__generic const volatile void *p,
159 int) __SPIRV_NOEXCEPT {
160 return (__global const volatile void *)p;
161}
162static __SPIRV_overloadable __SPIRV_inline __local void *
163__spirv_GenericCastToPtr_ToLocal(__generic void *p, int) __SPIRV_NOEXCEPT {
164 return (__local void *)p;
165}
166static __SPIRV_overloadable __SPIRV_inline __local const void *
167__spirv_GenericCastToPtr_ToLocal(__generic const void *p,
168 int) __SPIRV_NOEXCEPT {
169 return (__local const void *)p;
170}
171static __SPIRV_overloadable __SPIRV_inline __local volatile void *
172__spirv_GenericCastToPtr_ToLocal(__generic volatile void *p,
173 int) __SPIRV_NOEXCEPT {
174 return (__local volatile void *)p;
175}
176static __SPIRV_overloadable __SPIRV_inline __local const volatile void *
177__spirv_GenericCastToPtr_ToLocal(__generic const volatile void *p,
178 int) __SPIRV_NOEXCEPT {
179 return (__local const volatile void *)p;
180}
181static __SPIRV_overloadable __SPIRV_inline __private void *
182__spirv_GenericCastToPtr_ToPrivate(__generic void *p, int) __SPIRV_NOEXCEPT {
183 return (__private void *)p;
184}
185static __SPIRV_overloadable __SPIRV_inline __private const void *
186__spirv_GenericCastToPtr_ToPrivate(__generic const void *p,
187 int) __SPIRV_NOEXCEPT {
188 return (__private const void *)p;
189}
190static __SPIRV_overloadable __SPIRV_inline __private volatile void *
191__spirv_GenericCastToPtr_ToPrivate(__generic volatile void *p,
192 int) __SPIRV_NOEXCEPT {
193 return (__private volatile void *)p;
194}
195static __SPIRV_overloadable __SPIRV_inline __private const volatile void *
196__spirv_GenericCastToPtr_ToPrivate(__generic const volatile void *p,
197 int) __SPIRV_NOEXCEPT {
198 return (__private const volatile void *)p;
199}
200
201#pragma pop_macro("__size_t")
202#pragma pop_macro("__uint32_t")
203#pragma pop_macro("__uint64_t")
204
205#undef __SPIRV_overloadable
206#undef __SPIRV_convergent
207#undef __SPIRV_inline
208
209#undef __global
210#undef __local
211#undef __constant
212#undef __generic
213
214#undef __SPIRV_BUILTIN_ALIAS
215#undef __SPIRV_NOEXCEPT
216
217#endif /* __SPIRV_BUILTIN_VARS_H */