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