Skip to content

Commit 4222148

Browse files
azabaznoigcbot
azabazno
authored and
igcbot
committed
Split SPIR-V VC BiF module into multiple files
1 parent d2c9fb2 commit 4222148

File tree

4 files changed

+333
-309
lines changed

4 files changed

+333
-309
lines changed

‎IGC/VectorCompiler/lib/BiF/CMakeLists.txt

Lines changed: 9 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -129,8 +129,10 @@ else()
129129
set(PRINTF_OCL_SRC_PATH ${CMAKE_CURRENT_SOURCE_DIR}/printf_ocl_genx.cpp)
130130
set(PRINTF_ZE_SRC_PATH ${CMAKE_CURRENT_SOURCE_DIR}/printf_ze_genx.cpp)
131131
set(EMU_DIVREM_SRC_PATH ${CMAKE_CURRENT_SOURCE_DIR}/emulation_divrem.cpp)
132-
set(SPIRV_BUILTINS_SRC_PATH ${CMAKE_CURRENT_SOURCE_DIR}/spirv_builtins_genx.cpp)
133132
set(FP2UI_SRC_PATH ${CMAKE_CURRENT_SOURCE_DIR}/fp2ui_conversion.cpp)
133+
set(SPIRV_MATH_BUILTINS_SRC_PATH ${CMAKE_CURRENT_SOURCE_DIR}/spirv_math_builtins_genx.cpp)
134+
set(SPIRV_ATOMIC_BUILTINS_SRC_PATH ${CMAKE_CURRENT_SOURCE_DIR}/spirv_atomic_builtins_genx.cpp)
135+
set(SPIRV_EXEC_BUILTINS_SRC_PATH ${CMAKE_CURRENT_SOURCE_DIR}/spirv_exec_builtins_genx.cpp)
134136

135137
vc_embed_bif(PRINTF_OCL_32_CPP_PATH ${PRINTF_OCL_SRC_PATH} VCBiFPrintfOCL 32
136138
DEPENDS ${PRINTF_NOT_CM_COMMON_H_PATH})
@@ -146,8 +148,9 @@ else()
146148
DEPENDS "${DIVREM_EXTRA_SOURCES}")
147149

148150
get_target_include_opt_list(IGC_SPIRV_HEADERS_INCLUDES IGCSPIRVHeaders)
149-
vc_embed_bif(SPIRV_BUILTINS_CPP_PATH ${SPIRV_BUILTINS_SRC_PATH} VCSPIRVBuiltins 64
150-
CLANG_INCLUDES ${IGC_SPIRV_HEADERS_INCLUDES})
151+
vc_embed_bif(SPIRV_BUILTINS_CPP_PATH
152+
"${SPIRV_MATH_BUILTINS_SRC_PATH};${SPIRV_ATOMIC_BUILTINS_SRC_PATH};${SPIRV_EXEC_BUILTINS_SRC_PATH}"
153+
VCSPIRVBuiltins 64 CLANG_INCLUDES ${IGC_SPIRV_HEADERS_INCLUDES})
151154

152155
add_custom_target(VCBiFPreparation
153156
DEPENDS ${PRINTF_OCL_32_CPP_PATH}
@@ -158,7 +161,9 @@ else()
158161
${SPIRV_BUILTINS_CPP_PATH}
159162
SOURCES ${PRINTF_OCL_SRC_PATH}
160163
${EMU_DIVREM_SRC_PATH}
161-
${SPIRV_BUILTINS_SRC_PATH})
164+
${SPIRV_MATH_BUILTINS_SRC_PATH}
165+
${SPIRV_EXEC_BUILTINS_SRC_PATH}
166+
${SPIRV_ATOMIC_BUILTINS_SRC_PATH})
162167

163168
add_library(VCEmbeddedBiF
164169
${PRINTF_OCL_32_CPP_PATH}

‎IGC/VectorCompiler/lib/BiF/spirv_atomic_builtins_genx.cpp

Lines changed: 290 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,290 @@
1+
/*========================== begin_copyright_notice ============================
2+
3+
Copyright (C) 2021 Intel Corporation
4+
5+
SPDX-License-Identifier: MIT
6+
7+
============================= end_copyright_notice ===========================*/
8+
9+
#include <cm-cl/atomic.h>
10+
#include <cm-cl/vector.h>
11+
12+
using namespace cm;
13+
14+
extern "C" {
15+
#include "spirv_atomics_common.h"
16+
}
17+
18+
namespace {
19+
20+
template <typename PtrT, atomic::operation op, memory_order OCLSemantics,
21+
memory_scope OCLScope, typename... OpT>
22+
constexpr auto invokeConcreteAtomic(PtrT *ptr, OpT... operands) {
23+
return atomic::execute<op, OCLSemantics, OCLScope, PtrT>(ptr, operands...);
24+
}
25+
26+
template <typename PtrT, atomic::operation operation, memory_scope OCLScope,
27+
typename... OpT>
28+
static CM_NODEBUG CM_INLINE auto
29+
spirvAtomicHelperWithKnownScope(PtrT *ptr, int Semantics, OpT... operands) {
30+
switch (Semantics) {
31+
default:
32+
case SequentiallyConsistent:
33+
return invokeConcreteAtomic<PtrT, operation, memory_order_seq_cst,
34+
OCLScope>(ptr, operands...);
35+
case Relaxed:
36+
return invokeConcreteAtomic<PtrT, operation, memory_order_relaxed,
37+
OCLScope>(ptr, operands...);
38+
case Acquire:
39+
return invokeConcreteAtomic<PtrT, operation, memory_order_acquire,
40+
OCLScope>(ptr, operands...);
41+
case Release:
42+
return invokeConcreteAtomic<PtrT, operation, memory_order_release,
43+
OCLScope>(ptr, operands...);
44+
case AcquireRelease:
45+
return invokeConcreteAtomic<PtrT, operation, memory_order_acq_rel,
46+
OCLScope>(ptr, operands...);
47+
}
48+
}
49+
50+
// Iterate through all possible values of non-constant semantics and scope.
51+
// Use the strict possible values if unknown: sequential-consistency semantics
52+
// and cross-device scope.
53+
template <atomic::operation operation, typename PtrT, typename... OpT>
54+
static CM_NODEBUG CM_INLINE auto spirvAtomicHelper(PtrT *ptr, int Semantics,
55+
int Scope, OpT... operands) {
56+
switch (Scope) {
57+
default:
58+
case CrossDevice:
59+
return spirvAtomicHelperWithKnownScope<PtrT, operation,
60+
memory_scope_all_devices>(
61+
ptr, Semantics, operands...);
62+
case Device:
63+
return spirvAtomicHelperWithKnownScope<PtrT, operation,
64+
memory_scope_device>(ptr, Semantics,
65+
operands...);
66+
case Workgroup:
67+
return spirvAtomicHelperWithKnownScope<PtrT, operation,
68+
memory_scope_work_group>(
69+
ptr, Semantics, operands...);
70+
case Subgroup:
71+
return spirvAtomicHelperWithKnownScope<PtrT, operation,
72+
memory_scope_sub_group>(
73+
ptr, Semantics, operands...);
74+
case Invocation:
75+
return spirvAtomicHelperWithKnownScope<PtrT, operation,
76+
memory_scope_work_item>(
77+
ptr, Semantics, operands...);
78+
}
79+
}
80+
81+
} // namespace
82+
83+
#define SPIRV_ATOMIC_BUILTIN_LOAD(ADDRESS_SPACE, TYPE) \
84+
CM_NODEBUG CM_INLINE TYPE __spirv_AtomicLoad(ADDRESS_SPACE TYPE *ptr, \
85+
int Scope, int Semantics) { \
86+
return spirvAtomicHelper<atomic::operation::load, ADDRESS_SPACE TYPE>( \
87+
ptr, Semantics, Scope); \
88+
}
89+
90+
// FIXME: strict aliasing violation.
91+
#define SPIRV_ATOMIC_BUILTIN_FLOATS(ADDRESS_SPACE) \
92+
CM_NODEBUG CM_INLINE float __spirv_AtomicLoad(ADDRESS_SPACE float *ptr, \
93+
int Scope, int Semantics) { \
94+
ADDRESS_SPACE int *int_ptr = reinterpret_cast<ADDRESS_SPACE int *>(ptr); \
95+
return as_float(__spirv_AtomicLoad(int_ptr, Scope, Semantics)); \
96+
} \
97+
CM_NODEBUG CM_INLINE double __spirv_AtomicLoad(ADDRESS_SPACE double *ptr, \
98+
int Scope, int Semantics) { \
99+
ADDRESS_SPACE long *long_ptr = \
100+
reinterpret_cast<ADDRESS_SPACE long *>(ptr); \
101+
return as_double(__spirv_AtomicLoad(long_ptr, Scope, Semantics)); \
102+
}
103+
104+
SPIRV_ATOMIC_BUILTIN_LOAD(__global, int)
105+
SPIRV_ATOMIC_BUILTIN_LOAD(__global, long)
106+
SPIRV_ATOMIC_BUILTIN_FLOATS(__global)
107+
SPIRV_ATOMIC_BUILTIN_LOAD(__local, int)
108+
SPIRV_ATOMIC_BUILTIN_LOAD(__local, long)
109+
SPIRV_ATOMIC_BUILTIN_FLOATS(__local)
110+
SPIRV_ATOMIC_BUILTIN_LOAD(__generic, int)
111+
SPIRV_ATOMIC_BUILTIN_LOAD(__generic, long)
112+
SPIRV_ATOMIC_BUILTIN_FLOATS(__generic)
113+
114+
#define SPIRV_ATOMIC_BUILTIN_STORE(ADDRESS_SPACE, TYPE) \
115+
CM_NODEBUG CM_INLINE void __spirv_AtomicStore( \
116+
ADDRESS_SPACE TYPE *ptr, int Scope, int Semantics, TYPE Value) { \
117+
spirvAtomicHelper<atomic::operation::store, ADDRESS_SPACE TYPE>( \
118+
ptr, Semantics, Scope, Value); \
119+
}
120+
121+
SPIRV_ATOMIC_BUILTIN_STORE(__global, int)
122+
SPIRV_ATOMIC_BUILTIN_STORE(__global, long)
123+
SPIRV_ATOMIC_BUILTIN_STORE(__global, float)
124+
SPIRV_ATOMIC_BUILTIN_STORE(__global, double)
125+
SPIRV_ATOMIC_BUILTIN_STORE(__local, int)
126+
SPIRV_ATOMIC_BUILTIN_STORE(__local, long)
127+
SPIRV_ATOMIC_BUILTIN_STORE(__local, float)
128+
SPIRV_ATOMIC_BUILTIN_STORE(__local, double)
129+
SPIRV_ATOMIC_BUILTIN_STORE(__generic, int)
130+
SPIRV_ATOMIC_BUILTIN_STORE(__generic, long)
131+
SPIRV_ATOMIC_BUILTIN_STORE(__generic, float)
132+
SPIRV_ATOMIC_BUILTIN_STORE(__generic, double)
133+
134+
#define SPIRV_ATOMIC_BUILTIN_BINARY(SPIRV_ATOMIC_OP, GEN_ATOMIC_OP, \
135+
ADDRESS_SPACE, TYPE) \
136+
CM_NODEBUG CM_INLINE TYPE __spirv_Atomic##SPIRV_ATOMIC_OP( \
137+
ADDRESS_SPACE TYPE *ptr, int Scope, int Semantics, TYPE Value) { \
138+
return spirvAtomicHelper<GEN_ATOMIC_OP, ADDRESS_SPACE TYPE>( \
139+
ptr, Semantics, Scope, Value); \
140+
}
141+
142+
SPIRV_ATOMIC_BUILTIN_BINARY(SMin, atomic::operation::minsint, __global, int)
143+
SPIRV_ATOMIC_BUILTIN_BINARY(SMin, atomic::operation::minsint, __global, long)
144+
SPIRV_ATOMIC_BUILTIN_BINARY(SMin, atomic::operation::minsint, __local, int)
145+
SPIRV_ATOMIC_BUILTIN_BINARY(SMin, atomic::operation::minsint, __local, long)
146+
SPIRV_ATOMIC_BUILTIN_BINARY(SMin, atomic::operation::minsint, __generic, int)
147+
SPIRV_ATOMIC_BUILTIN_BINARY(SMin, atomic::operation::minsint, __generic, long)
148+
149+
SPIRV_ATOMIC_BUILTIN_BINARY(SMax, atomic::operation::maxsint, __global, int)
150+
SPIRV_ATOMIC_BUILTIN_BINARY(SMax, atomic::operation::maxsint, __global, long)
151+
SPIRV_ATOMIC_BUILTIN_BINARY(SMax, atomic::operation::maxsint, __local, int)
152+
SPIRV_ATOMIC_BUILTIN_BINARY(SMax, atomic::operation::maxsint, __local, long)
153+
SPIRV_ATOMIC_BUILTIN_BINARY(SMax, atomic::operation::maxsint, __generic, int)
154+
SPIRV_ATOMIC_BUILTIN_BINARY(SMax, atomic::operation::maxsint, __generic, long)
155+
156+
SPIRV_ATOMIC_BUILTIN_BINARY(UMin, atomic::operation::min, __global, uint)
157+
SPIRV_ATOMIC_BUILTIN_BINARY(UMin, atomic::operation::min, __global, ulong)
158+
SPIRV_ATOMIC_BUILTIN_BINARY(UMin, atomic::operation::min, __local, uint)
159+
SPIRV_ATOMIC_BUILTIN_BINARY(UMin, atomic::operation::min, __local, ulong)
160+
SPIRV_ATOMIC_BUILTIN_BINARY(UMin, atomic::operation::min, __generic, uint)
161+
SPIRV_ATOMIC_BUILTIN_BINARY(UMin, atomic::operation::min, __generic, ulong)
162+
163+
SPIRV_ATOMIC_BUILTIN_BINARY(UMax, atomic::operation::max, __global, uint)
164+
SPIRV_ATOMIC_BUILTIN_BINARY(UMax, atomic::operation::max, __global, ulong)
165+
SPIRV_ATOMIC_BUILTIN_BINARY(UMax, atomic::operation::max, __local, uint)
166+
SPIRV_ATOMIC_BUILTIN_BINARY(UMax, atomic::operation::max, __local, ulong)
167+
SPIRV_ATOMIC_BUILTIN_BINARY(UMax, atomic::operation::max, __generic, uint)
168+
SPIRV_ATOMIC_BUILTIN_BINARY(UMax, atomic::operation::max, __generic, ulong)
169+
170+
SPIRV_ATOMIC_BUILTIN_BINARY(IAdd, atomic::operation::add, __global, int)
171+
SPIRV_ATOMIC_BUILTIN_BINARY(IAdd, atomic::operation::add, __global, long)
172+
SPIRV_ATOMIC_BUILTIN_BINARY(IAdd, atomic::operation::add, __local, int)
173+
SPIRV_ATOMIC_BUILTIN_BINARY(IAdd, atomic::operation::add, __local, long)
174+
SPIRV_ATOMIC_BUILTIN_BINARY(IAdd, atomic::operation::add, __generic, int)
175+
SPIRV_ATOMIC_BUILTIN_BINARY(IAdd, atomic::operation::add, __generic, long)
176+
177+
SPIRV_ATOMIC_BUILTIN_BINARY(ISub, atomic::operation::sub, __global, int)
178+
SPIRV_ATOMIC_BUILTIN_BINARY(ISub, atomic::operation::sub, __global, long)
179+
SPIRV_ATOMIC_BUILTIN_BINARY(ISub, atomic::operation::sub, __local, int)
180+
SPIRV_ATOMIC_BUILTIN_BINARY(ISub, atomic::operation::sub, __local, long)
181+
SPIRV_ATOMIC_BUILTIN_BINARY(ISub, atomic::operation::sub, __generic, int)
182+
SPIRV_ATOMIC_BUILTIN_BINARY(ISub, atomic::operation::sub, __generic, long)
183+
184+
SPIRV_ATOMIC_BUILTIN_BINARY(Or, atomic::operation::orl, __global, int)
185+
SPIRV_ATOMIC_BUILTIN_BINARY(Or, atomic::operation::orl, __global, long)
186+
SPIRV_ATOMIC_BUILTIN_BINARY(Or, atomic::operation::orl, __local, int)
187+
SPIRV_ATOMIC_BUILTIN_BINARY(Or, atomic::operation::orl, __local, long)
188+
SPIRV_ATOMIC_BUILTIN_BINARY(Or, atomic::operation::orl, __generic, int)
189+
SPIRV_ATOMIC_BUILTIN_BINARY(Or, atomic::operation::orl, __generic, long)
190+
191+
SPIRV_ATOMIC_BUILTIN_BINARY(Xor, atomic::operation::xorl, __global, int)
192+
SPIRV_ATOMIC_BUILTIN_BINARY(Xor, atomic::operation::xorl, __global, long)
193+
SPIRV_ATOMIC_BUILTIN_BINARY(Xor, atomic::operation::xorl, __local, int)
194+
SPIRV_ATOMIC_BUILTIN_BINARY(Xor, atomic::operation::xorl, __local, long)
195+
SPIRV_ATOMIC_BUILTIN_BINARY(Xor, atomic::operation::xorl, __generic, int)
196+
SPIRV_ATOMIC_BUILTIN_BINARY(Xor, atomic::operation::xorl, __generic, long)
197+
198+
SPIRV_ATOMIC_BUILTIN_BINARY(And, atomic::operation::andl, __global, int)
199+
SPIRV_ATOMIC_BUILTIN_BINARY(And, atomic::operation::andl, __global, long)
200+
SPIRV_ATOMIC_BUILTIN_BINARY(And, atomic::operation::andl, __local, int)
201+
SPIRV_ATOMIC_BUILTIN_BINARY(And, atomic::operation::andl, __local, long)
202+
SPIRV_ATOMIC_BUILTIN_BINARY(And, atomic::operation::andl, __generic, int)
203+
SPIRV_ATOMIC_BUILTIN_BINARY(And, atomic::operation::andl, __generic, long)
204+
205+
SPIRV_ATOMIC_BUILTIN_BINARY(Exchange, atomic::operation::xchg, __global, int)
206+
SPIRV_ATOMIC_BUILTIN_BINARY(Exchange, atomic::operation::xchg, __global, long)
207+
SPIRV_ATOMIC_BUILTIN_BINARY(Exchange, atomic::operation::xchg, __global, float)
208+
SPIRV_ATOMIC_BUILTIN_BINARY(Exchange, atomic::operation::xchg, __global, double)
209+
SPIRV_ATOMIC_BUILTIN_BINARY(Exchange, atomic::operation::xchg, __local, int)
210+
SPIRV_ATOMIC_BUILTIN_BINARY(Exchange, atomic::operation::xchg, __local, long)
211+
SPIRV_ATOMIC_BUILTIN_BINARY(Exchange, atomic::operation::xchg, __local, float)
212+
SPIRV_ATOMIC_BUILTIN_BINARY(Exchange, atomic::operation::xchg, __local, double)
213+
SPIRV_ATOMIC_BUILTIN_BINARY(Exchange, atomic::operation::xchg, __generic, int)
214+
SPIRV_ATOMIC_BUILTIN_BINARY(Exchange, atomic::operation::xchg, __generic, long)
215+
SPIRV_ATOMIC_BUILTIN_BINARY(Exchange, atomic::operation::xchg, __generic, float)
216+
SPIRV_ATOMIC_BUILTIN_BINARY(Exchange, atomic::operation::xchg, __generic,
217+
double)
218+
219+
#define SPIRV_ATOMIC_BUILTIN_BINARY_WITH_OPVALUE( \
220+
SPIRV_ATOMIC_OP, GEN_ATOMIC_OP, OPVALUE, ADDRESS_SPACE, TYPE) \
221+
CM_NODEBUG CM_INLINE TYPE __spirv_Atomic##SPIRV_ATOMIC_OP( \
222+
ADDRESS_SPACE TYPE *ptr, int Scope, int Semantics) { \
223+
TYPE Value = OPVALUE; \
224+
return spirvAtomicHelper<GEN_ATOMIC_OP, ADDRESS_SPACE TYPE>( \
225+
ptr, Semantics, Scope, Value); \
226+
}
227+
228+
SPIRV_ATOMIC_BUILTIN_BINARY_WITH_OPVALUE(IIncrement, atomic::operation::add, 1,
229+
__global, int)
230+
SPIRV_ATOMIC_BUILTIN_BINARY_WITH_OPVALUE(IIncrement, atomic::operation::add, 1,
231+
__global, long)
232+
SPIRV_ATOMIC_BUILTIN_BINARY_WITH_OPVALUE(IIncrement, atomic::operation::add, 1,
233+
__local, int)
234+
SPIRV_ATOMIC_BUILTIN_BINARY_WITH_OPVALUE(IIncrement, atomic::operation::add, 1,
235+
__local, long)
236+
SPIRV_ATOMIC_BUILTIN_BINARY_WITH_OPVALUE(IIncrement, atomic::operation::add, 1,
237+
__generic, int)
238+
SPIRV_ATOMIC_BUILTIN_BINARY_WITH_OPVALUE(IIncrement, atomic::operation::add, 1,
239+
__generic, long)
240+
241+
SPIRV_ATOMIC_BUILTIN_BINARY_WITH_OPVALUE(IDecrement, atomic::operation::sub, 1,
242+
__global, int)
243+
SPIRV_ATOMIC_BUILTIN_BINARY_WITH_OPVALUE(IDecrement, atomic::operation::sub, 1,
244+
__global, long)
245+
SPIRV_ATOMIC_BUILTIN_BINARY_WITH_OPVALUE(IDecrement, atomic::operation::sub, 1,
246+
__local, int)
247+
SPIRV_ATOMIC_BUILTIN_BINARY_WITH_OPVALUE(IDecrement, atomic::operation::sub, 1,
248+
__local, long)
249+
SPIRV_ATOMIC_BUILTIN_BINARY_WITH_OPVALUE(IDecrement, atomic::operation::sub, 1,
250+
__generic, int)
251+
SPIRV_ATOMIC_BUILTIN_BINARY_WITH_OPVALUE(IDecrement, atomic::operation::sub, 1,
252+
__generic, long)
253+
254+
// FIXME: Unequal semantics will be eventually merged with
255+
// Equal semantics which will result in Equal semantics,
256+
// so we're skipping here for now.
257+
#define SPIRV_ATOMIC_BUILTIN_TERNARY(SPIRV_ATOMIC_OP, GEN_ATOMIC_OP, \
258+
ADDRESS_SPACE, TYPE) \
259+
CM_NODEBUG CM_INLINE TYPE __spirv_Atomic##SPIRV_ATOMIC_OP( \
260+
ADDRESS_SPACE TYPE *ptr, int Scope, int Equal, int Unequal, TYPE Value1, \
261+
TYPE Value2) { \
262+
return spirvAtomicHelper<GEN_ATOMIC_OP, ADDRESS_SPACE TYPE>( \
263+
ptr, Equal, Scope, Value1, Value2); \
264+
}
265+
266+
SPIRV_ATOMIC_BUILTIN_TERNARY(CompareExchange, atomic::operation::cmpxchg,
267+
__global, int)
268+
SPIRV_ATOMIC_BUILTIN_TERNARY(CompareExchange, atomic::operation::cmpxchg,
269+
__global, long)
270+
SPIRV_ATOMIC_BUILTIN_TERNARY(CompareExchange, atomic::operation::cmpxchg,
271+
__local, int)
272+
SPIRV_ATOMIC_BUILTIN_TERNARY(CompareExchange, atomic::operation::cmpxchg,
273+
__local, long)
274+
SPIRV_ATOMIC_BUILTIN_TERNARY(CompareExchange, atomic::operation::cmpxchg,
275+
__generic, int)
276+
SPIRV_ATOMIC_BUILTIN_TERNARY(CompareExchange, atomic::operation::cmpxchg,
277+
__generic, long)
278+
279+
SPIRV_ATOMIC_BUILTIN_TERNARY(CompareExchangeWeak, atomic::operation::cmpxchg,
280+
__global, int)
281+
SPIRV_ATOMIC_BUILTIN_TERNARY(CompareExchangeWeak, atomic::operation::cmpxchg,
282+
__global, long)
283+
SPIRV_ATOMIC_BUILTIN_TERNARY(CompareExchangeWeak, atomic::operation::cmpxchg,
284+
__local, int)
285+
SPIRV_ATOMIC_BUILTIN_TERNARY(CompareExchangeWeak, atomic::operation::cmpxchg,
286+
__local, long)
287+
SPIRV_ATOMIC_BUILTIN_TERNARY(CompareExchangeWeak, atomic::operation::cmpxchg,
288+
__generic, int)
289+
SPIRV_ATOMIC_BUILTIN_TERNARY(CompareExchangeWeak, atomic::operation::cmpxchg,
290+
__generic, long)

‎IGC/VectorCompiler/lib/BiF/spirv_exec_builtins_genx.cpp

Lines changed: 34 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,34 @@
1+
/*========================== begin_copyright_notice ============================
2+
3+
Copyright (C) 2021 Intel Corporation
4+
5+
SPDX-License-Identifier: MIT
6+
7+
============================= end_copyright_notice ===========================*/
8+
9+
#include <cm-cl/exec.h>
10+
11+
using namespace cm;
12+
13+
CM_NODEBUG CM_INLINE ulong __spirv_BuiltInGlobalInvocationId(int dim) {
14+
return cm::exec::get_local_id(dim) +
15+
static_cast<ulong>(cm::exec::get_group_id(dim)) *
16+
cm::exec::get_local_size(dim);
17+
}
18+
19+
CM_NODEBUG CM_INLINE ulong __spirv_BuiltInWorkgroupSize(int dim) {
20+
return cm::exec::get_local_size(dim);
21+
}
22+
23+
CM_NODEBUG CM_INLINE ulong __spirv_BuiltInLocalInvocationId(int dim) {
24+
return cm::exec::get_local_id(dim);
25+
}
26+
27+
CM_NODEBUG CM_INLINE ulong __spirv_BuiltInWorkgroupId(int dim) {
28+
return cm::exec::get_group_id(dim);
29+
}
30+
31+
CM_NODEBUG CM_INLINE ulong __spirv_BuiltInGlobalSize(int dim) {
32+
return static_cast<ulong>(cm::exec::get_local_size(dim)) *
33+
cm::exec::get_group_count(dim);
34+
}

0 commit comments

Comments
 (0)