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