218 lines · c
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-exception6 *7 *===-----------------------------------------------------------------------===8 */9 10#ifndef __SPIRV_BUILTIN_VARS_H11#define __SPIRV_BUILTIN_VARS_H12 13#if __cplusplus >= 201103L14#define __SPIRV_NOEXCEPT noexcept15#else16#define __SPIRV_NOEXCEPT17#endif18 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 __generic35#else36#define __generic __attribute__((opencl_generic))37#endif38 39// Check if SPIR-V builtins are supported.40// As the translator doesn't use the LLVM intrinsics (which would be emitted if41// we use the SPIR-V builtins) we can't rely on the SPIRV32/SPIRV64 etc macros42// to establish if we can use the builtin alias. We disable builtin altogether43// if we do not intent to use the backend. So instead of use target macros, rely44// 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#else49#define __SPIRV_BUILTIN_ALIAS(builtin)50#endif51 52// Builtin IDs and sizes53 54extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_num_workgroups) __size_t55 __spirv_BuiltInNumWorkgroups(int);56extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_workgroup_size) __size_t57 __spirv_BuiltInWorkgroupSize(int);58extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_workgroup_id) __size_t59 __spirv_BuiltInWorkgroupId(int);60extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_local_invocation_id) __size_t61 __spirv_BuiltInLocalInvocationId(int);62extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_global_invocation_id) __size_t63 __spirv_BuiltInGlobalInvocationId(int);64 65extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_global_size) __size_t66 __spirv_BuiltInGlobalSize(int);67extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_global_offset) __size_t68 __spirv_BuiltInGlobalOffset(int);69extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_size) __uint32_t70 __spirv_BuiltInSubgroupSize();71extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_max_size) __uint32_t72 __spirv_BuiltInSubgroupMaxSize();73extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_num_subgroups) __uint32_t74 __spirv_BuiltInNumSubgroups();75extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_id) __uint32_t76 __spirv_BuiltInSubgroupId();77extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_local_invocation_id)78 __uint32_t __spirv_BuiltInSubgroupLocalInvocationId();79 80// OpGenericCastToPtrExplicit81 82extern __SPIRV_overloadable83__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)84__global void *__spirv_GenericCastToPtrExplicit_ToGlobal(__generic void *,85 int) __SPIRV_NOEXCEPT;86extern __SPIRV_overloadable87__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_overloadable92__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_overloadable97__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_overloadable102__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)103__local void *__spirv_GenericCastToPtrExplicit_ToLocal(__generic void *,104 int) __SPIRV_NOEXCEPT;105extern __SPIRV_overloadable106__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_overloadable111__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_overloadable116__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_overloadable121__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_overloadable126__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_overloadable131__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_overloadable136__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// OpGenericCastToPtr142 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_overloadable206#undef __SPIRV_convergent207#undef __SPIRV_inline208 209#undef __global210#undef __local211#undef __constant212#undef __generic213 214#undef __SPIRV_BUILTIN_ALIAS215#undef __SPIRV_NOEXCEPT216 217#endif /* __SPIRV_BUILTIN_VARS_H */218