DPC++ Runtime
Runtime libraries for oneAPI DPC++
spirv_vars.hpp
Go to the documentation of this file.
1 //==----------- spirv_vars.hpp --- SPIRV variables -------------------------==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 // ===-------------------------------------------------------------------=== //
8 
9 #pragma once
10 
11 #ifdef __SYCL_DEVICE_ONLY__
12 
13 #include <sycl/detail/defines_elementary.hpp> // for __DPCPP_SYCL_EXTERNAL
14 
15 #include <cstddef> // for size_t
16 #include <cstdint> // for uint8_t
17 
18 #define __SPIRV_VAR_QUALIFIERS extern "C" const
19 
20 #if defined(__NVPTX__) || defined(__AMDGCN__) || defined(__SYCL_NATIVE_CPU__)
21 
22 __DPCPP_SYCL_EXTERNAL size_t __spirv_GlobalInvocationId_x();
23 __DPCPP_SYCL_EXTERNAL size_t __spirv_GlobalInvocationId_y();
24 __DPCPP_SYCL_EXTERNAL size_t __spirv_GlobalInvocationId_z();
25 
26 __DPCPP_SYCL_EXTERNAL size_t __spirv_GlobalSize_x();
27 __DPCPP_SYCL_EXTERNAL size_t __spirv_GlobalSize_y();
28 __DPCPP_SYCL_EXTERNAL size_t __spirv_GlobalSize_z();
29 
30 __DPCPP_SYCL_EXTERNAL size_t __spirv_GlobalOffset_x();
31 __DPCPP_SYCL_EXTERNAL size_t __spirv_GlobalOffset_y();
32 __DPCPP_SYCL_EXTERNAL size_t __spirv_GlobalOffset_z();
33 
34 __DPCPP_SYCL_EXTERNAL size_t __spirv_NumWorkgroups_x();
35 __DPCPP_SYCL_EXTERNAL size_t __spirv_NumWorkgroups_y();
36 __DPCPP_SYCL_EXTERNAL size_t __spirv_NumWorkgroups_z();
37 
38 __DPCPP_SYCL_EXTERNAL size_t __spirv_WorkgroupSize_x();
39 __DPCPP_SYCL_EXTERNAL size_t __spirv_WorkgroupSize_y();
40 __DPCPP_SYCL_EXTERNAL size_t __spirv_WorkgroupSize_z();
41 
42 __DPCPP_SYCL_EXTERNAL size_t __spirv_WorkgroupId_x();
43 __DPCPP_SYCL_EXTERNAL size_t __spirv_WorkgroupId_y();
44 __DPCPP_SYCL_EXTERNAL size_t __spirv_WorkgroupId_z();
45 
46 __DPCPP_SYCL_EXTERNAL size_t __spirv_LocalInvocationId_x();
47 __DPCPP_SYCL_EXTERNAL size_t __spirv_LocalInvocationId_y();
48 __DPCPP_SYCL_EXTERNAL size_t __spirv_LocalInvocationId_z();
49 
50 __DPCPP_SYCL_EXTERNAL uint32_t __spirv_SubgroupSize();
51 __DPCPP_SYCL_EXTERNAL uint32_t __spirv_SubgroupMaxSize();
52 __DPCPP_SYCL_EXTERNAL uint32_t __spirv_NumSubgroups();
53 __DPCPP_SYCL_EXTERNAL uint32_t __spirv_SubgroupId();
54 __DPCPP_SYCL_EXTERNAL uint32_t __spirv_SubgroupLocalInvocationId();
55 
56 #else // defined(__NVPTX__) || defined(__AMDGCN__)
57 
58 typedef size_t size_t_vec __attribute__((ext_vector_type(3)));
59 __SPIRV_VAR_QUALIFIERS size_t_vec __spirv_BuiltInGlobalSize;
60 __SPIRV_VAR_QUALIFIERS size_t_vec __spirv_BuiltInGlobalInvocationId;
61 __SPIRV_VAR_QUALIFIERS size_t_vec __spirv_BuiltInWorkgroupSize;
62 __SPIRV_VAR_QUALIFIERS size_t_vec __spirv_BuiltInNumWorkgroups;
63 __SPIRV_VAR_QUALIFIERS size_t_vec __spirv_BuiltInLocalInvocationId;
64 __SPIRV_VAR_QUALIFIERS size_t_vec __spirv_BuiltInWorkgroupId;
65 __SPIRV_VAR_QUALIFIERS size_t_vec __spirv_BuiltInGlobalOffset;
66 
67 __SPIRV_VAR_QUALIFIERS uint32_t __spirv_BuiltInSubgroupSize;
68 __SPIRV_VAR_QUALIFIERS uint32_t __spirv_BuiltInSubgroupMaxSize;
69 __SPIRV_VAR_QUALIFIERS uint32_t __spirv_BuiltInNumSubgroups;
70 __SPIRV_VAR_QUALIFIERS uint32_t __spirv_BuiltInSubgroupId;
71 __SPIRV_VAR_QUALIFIERS uint32_t __spirv_BuiltInSubgroupLocalInvocationId;
72 
73 __SPIRV_VAR_QUALIFIERS __ocl_vec_t<uint32_t, 4> __spirv_BuiltInSubgroupEqMask;
74 __SPIRV_VAR_QUALIFIERS __ocl_vec_t<uint32_t, 4> __spirv_BuiltInSubgroupGeMask;
75 __SPIRV_VAR_QUALIFIERS __ocl_vec_t<uint32_t, 4> __spirv_BuiltInSubgroupGtMask;
76 __SPIRV_VAR_QUALIFIERS __ocl_vec_t<uint32_t, 4> __spirv_BuiltInSubgroupLeMask;
77 __SPIRV_VAR_QUALIFIERS __ocl_vec_t<uint32_t, 4> __spirv_BuiltInSubgroupLtMask;
78 
79 __DPCPP_SYCL_EXTERNAL inline size_t __spirv_GlobalInvocationId_x() {
80  return __spirv_BuiltInGlobalInvocationId.x;
81 }
82 __DPCPP_SYCL_EXTERNAL inline size_t __spirv_GlobalInvocationId_y() {
83  return __spirv_BuiltInGlobalInvocationId.y;
84 }
85 __DPCPP_SYCL_EXTERNAL inline size_t __spirv_GlobalInvocationId_z() {
86  return __spirv_BuiltInGlobalInvocationId.z;
87 }
88 
89 __DPCPP_SYCL_EXTERNAL inline size_t __spirv_GlobalSize_x() {
90  return __spirv_BuiltInGlobalSize.x;
91 }
92 __DPCPP_SYCL_EXTERNAL inline size_t __spirv_GlobalSize_y() {
93  return __spirv_BuiltInGlobalSize.y;
94 }
95 __DPCPP_SYCL_EXTERNAL inline size_t __spirv_GlobalSize_z() {
96  return __spirv_BuiltInGlobalSize.z;
97 }
98 
99 __DPCPP_SYCL_EXTERNAL inline size_t __spirv_GlobalOffset_x() {
100  return __spirv_BuiltInGlobalOffset.x;
101 }
102 __DPCPP_SYCL_EXTERNAL inline size_t __spirv_GlobalOffset_y() {
103  return __spirv_BuiltInGlobalOffset.y;
104 }
105 __DPCPP_SYCL_EXTERNAL inline size_t __spirv_GlobalOffset_z() {
106  return __spirv_BuiltInGlobalOffset.z;
107 }
108 
109 __DPCPP_SYCL_EXTERNAL inline size_t __spirv_NumWorkgroups_x() {
110  return __spirv_BuiltInNumWorkgroups.x;
111 }
112 __DPCPP_SYCL_EXTERNAL inline size_t __spirv_NumWorkgroups_y() {
113  return __spirv_BuiltInNumWorkgroups.y;
114 }
115 __DPCPP_SYCL_EXTERNAL inline size_t __spirv_NumWorkgroups_z() {
116  return __spirv_BuiltInNumWorkgroups.z;
117 }
118 
119 __DPCPP_SYCL_EXTERNAL inline size_t __spirv_WorkgroupSize_x() {
120  return __spirv_BuiltInWorkgroupSize.x;
121 }
122 __DPCPP_SYCL_EXTERNAL inline size_t __spirv_WorkgroupSize_y() {
123  return __spirv_BuiltInWorkgroupSize.y;
124 }
125 __DPCPP_SYCL_EXTERNAL inline size_t __spirv_WorkgroupSize_z() {
126  return __spirv_BuiltInWorkgroupSize.z;
127 }
128 
129 __DPCPP_SYCL_EXTERNAL inline size_t __spirv_WorkgroupId_x() {
130  return __spirv_BuiltInWorkgroupId.x;
131 }
132 __DPCPP_SYCL_EXTERNAL inline size_t __spirv_WorkgroupId_y() {
133  return __spirv_BuiltInWorkgroupId.y;
134 }
135 __DPCPP_SYCL_EXTERNAL inline size_t __spirv_WorkgroupId_z() {
136  return __spirv_BuiltInWorkgroupId.z;
137 }
138 
139 __DPCPP_SYCL_EXTERNAL inline size_t __spirv_LocalInvocationId_x() {
140  return __spirv_BuiltInLocalInvocationId.x;
141 }
142 __DPCPP_SYCL_EXTERNAL inline size_t __spirv_LocalInvocationId_y() {
143  return __spirv_BuiltInLocalInvocationId.y;
144 }
145 __DPCPP_SYCL_EXTERNAL inline size_t __spirv_LocalInvocationId_z() {
146  return __spirv_BuiltInLocalInvocationId.z;
147 }
148 
149 __DPCPP_SYCL_EXTERNAL inline uint32_t __spirv_SubgroupSize() {
150  return __spirv_BuiltInSubgroupSize;
151 }
152 __DPCPP_SYCL_EXTERNAL inline uint32_t __spirv_SubgroupMaxSize() {
153  return __spirv_BuiltInSubgroupMaxSize;
154 }
155 __DPCPP_SYCL_EXTERNAL inline uint32_t __spirv_NumSubgroups() {
156  return __spirv_BuiltInNumSubgroups;
157 }
158 __DPCPP_SYCL_EXTERNAL inline uint32_t __spirv_SubgroupId() {
159  return __spirv_BuiltInSubgroupId;
160 }
161 __DPCPP_SYCL_EXTERNAL inline uint32_t __spirv_SubgroupLocalInvocationId() {
162  return __spirv_BuiltInSubgroupLocalInvocationId;
163 }
164 
165 #endif // defined(__NVPTX__) || defined(__AMDGCN__)
166 
167 #undef __SPIRV_VAR_QUALIFIERS
168 
169 namespace __spirv {
170 
171 // Helper function templates to initialize and get vector component from SPIR-V
172 // built-in variables
173 #define __SPIRV_DEFINE_INIT_AND_GET_HELPERS(POSTFIX) \
174  template <int ID> static size_t get##POSTFIX(); \
175  template <> size_t get##POSTFIX<0>() { return __spirv_##POSTFIX##_x(); } \
176  template <> size_t get##POSTFIX<1>() { return __spirv_##POSTFIX##_y(); } \
177  template <> size_t get##POSTFIX<2>() { return __spirv_##POSTFIX##_z(); } \
178  \
179  template <int Dim, class DstT> struct InitSizesST##POSTFIX; \
180  \
181  template <class DstT> struct InitSizesST##POSTFIX<1, DstT> { \
182  static DstT initSize() { return {get##POSTFIX<0>()}; } \
183  }; \
184  \
185  template <class DstT> struct InitSizesST##POSTFIX<2, DstT> { \
186  static DstT initSize() { return {get##POSTFIX<1>(), get##POSTFIX<0>()}; } \
187  }; \
188  \
189  template <class DstT> struct InitSizesST##POSTFIX<3, DstT> { \
190  static DstT initSize() { \
191  return {get##POSTFIX<2>(), get##POSTFIX<1>(), get##POSTFIX<0>()}; \
192  } \
193  }; \
194  \
195  template <int Dims, class DstT> DstT init##POSTFIX() { \
196  return InitSizesST##POSTFIX<Dims, DstT>::initSize(); \
197  }
198 
199 __SPIRV_DEFINE_INIT_AND_GET_HELPERS(GlobalSize);
200 __SPIRV_DEFINE_INIT_AND_GET_HELPERS(GlobalInvocationId)
201 __SPIRV_DEFINE_INIT_AND_GET_HELPERS(WorkgroupSize)
202 __SPIRV_DEFINE_INIT_AND_GET_HELPERS(NumWorkgroups)
203 __SPIRV_DEFINE_INIT_AND_GET_HELPERS(LocalInvocationId)
204 __SPIRV_DEFINE_INIT_AND_GET_HELPERS(WorkgroupId)
205 __SPIRV_DEFINE_INIT_AND_GET_HELPERS(GlobalOffset)
206 
207 #undef __SPIRV_DEFINE_INIT_AND_GET_HELPERS
208 
209 } // namespace __spirv
210 
211 #endif // __SYCL_DEVICE_ONLY__
defines_elementary.hpp
__DPCPP_SYCL_EXTERNAL
#define __DPCPP_SYCL_EXTERNAL
Definition: defines_elementary.hpp:33
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:464