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