Skip to content

Commit 7960931

Browse files
kbobrovskbobrovs
committed
[SYCL] Make item's and group's kernel initialization code reusable.
Device code in the SYCL runtime library may want to reuse SPIRV-based initialization of current group and item id, size and offset - factor them out into a separate header. Signed-off-by: Konstantin S Bobrovsky <[email protected]>
1 parent c05692e commit 7960931

File tree

2 files changed

+54
-50
lines changed

2 files changed

+54
-50
lines changed

‎sycl/include/CL/__spirv/spirv_vars.hpp

Lines changed: 45 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -24,13 +24,17 @@ extern "C" const __constant size_t_vec __spirv_BuiltInGlobalOffset;
2424
template <> size_t get##POSTFIX<1>() { return __spirv_BuiltIn##POSTFIX.y; } \
2525
template <> size_t get##POSTFIX<2>() { return __spirv_BuiltIn##POSTFIX.z; }
2626

27+
namespace __spirv {
28+
2729
DEFINE_INT_ID_TO_XYZ_CONVERTER(GlobalSize);
2830
DEFINE_INT_ID_TO_XYZ_CONVERTER(GlobalInvocationId)
2931
DEFINE_INT_ID_TO_XYZ_CONVERTER(WorkgroupSize)
3032
DEFINE_INT_ID_TO_XYZ_CONVERTER(LocalInvocationId)
3133
DEFINE_INT_ID_TO_XYZ_CONVERTER(WorkgroupId)
3234
DEFINE_INT_ID_TO_XYZ_CONVERTER(GlobalOffset)
3335

36+
} // namespace __spirv
37+
3438
#undef DEFINE_INT_ID_TO_XYZ_CONVERTER
3539

3640
extern "C" const __constant uint32_t __spirv_BuiltInSubgroupSize;
@@ -40,4 +44,45 @@ extern "C" const __constant uint32_t __spirv_BuiltInNumEnqueuedSubgroups;
4044
extern "C" const __constant uint32_t __spirv_BuiltInSubgroupId;
4145
extern "C" const __constant uint32_t __spirv_BuiltInSubgroupLocalInvocationId;
4246

47+
#define DEFINE_INIT_SIZES(POSTFIX) \
48+
\
49+
template <int Dim, class DstT> struct InitSizesST##POSTFIX; \
50+
\
51+
template <class DstT> struct InitSizesST##POSTFIX<1, DstT> { \
52+
static void initSize(DstT &Dst) { \
53+
Dst[0] = get##POSTFIX<0>(); \
54+
} \
55+
}; \
56+
\
57+
template <class DstT> struct InitSizesST##POSTFIX<2, DstT> { \
58+
static void initSize(DstT &Dst) { \
59+
Dst[1] = get##POSTFIX<1>(); \
60+
InitSizesST##POSTFIX<1, DstT>::initSize(Dst); \
61+
} \
62+
}; \
63+
\
64+
template <class DstT> struct InitSizesST##POSTFIX<3, DstT> { \
65+
static void initSize(DstT &Dst) { \
66+
Dst[2] = get##POSTFIX<2>(); \
67+
InitSizesST##POSTFIX<2, DstT>::initSize(Dst); \
68+
} \
69+
}; \
70+
\
71+
template <int Dims, class DstT> static void init##POSTFIX(DstT &Dst) { \
72+
InitSizesST##POSTFIX<Dims, DstT>::initSize(Dst); \
73+
}
74+
75+
namespace __spirv {
76+
77+
DEFINE_INIT_SIZES(GlobalSize);
78+
DEFINE_INIT_SIZES(GlobalInvocationId)
79+
DEFINE_INIT_SIZES(WorkgroupSize)
80+
DEFINE_INIT_SIZES(LocalInvocationId)
81+
DEFINE_INIT_SIZES(WorkgroupId)
82+
DEFINE_INIT_SIZES(GlobalOffset)
83+
84+
} // namespace __spirv
85+
86+
#undef DEFINE_INIT_SIZES
87+
4388
#endif // __SYCL_DEVICE_ONLY__

‎sycl/include/CL/sycl/handler.hpp

Lines changed: 9 additions & 50 deletions
Original file line numberDiff line numberDiff line change
@@ -67,47 +67,6 @@ namespace csd = cl::sycl::detail;
6767
template <typename T, int Dimensions, typename AllocatorT> class buffer;
6868
namespace detail {
6969

70-
#ifdef __SYCL_DEVICE_ONLY__
71-
72-
#define DEFINE_INIT_SIZES(POSTFIX) \
73-
\
74-
template <int Dim, class DstT> struct InitSizesST##POSTFIX; \
75-
\
76-
template <class DstT> struct InitSizesST##POSTFIX<1, DstT> { \
77-
static void initSize(DstT &Dst) { \
78-
Dst[0] = get##POSTFIX<0>(); \
79-
} \
80-
}; \
81-
\
82-
template <class DstT> struct InitSizesST##POSTFIX<2, DstT> { \
83-
static void initSize(DstT &Dst) { \
84-
Dst[1] = get##POSTFIX<1>(); \
85-
InitSizesST##POSTFIX<1, DstT>::initSize(Dst); \
86-
} \
87-
}; \
88-
\
89-
template <class DstT> struct InitSizesST##POSTFIX<3, DstT> { \
90-
static void initSize(DstT &Dst) { \
91-
Dst[2] = get##POSTFIX<2>(); \
92-
InitSizesST##POSTFIX<2, DstT>::initSize(Dst); \
93-
} \
94-
}; \
95-
\
96-
template <int Dims, class DstT> static void init##POSTFIX(DstT &Dst) { \
97-
InitSizesST##POSTFIX<Dims, DstT>::initSize(Dst); \
98-
}
99-
100-
DEFINE_INIT_SIZES(GlobalSize);
101-
DEFINE_INIT_SIZES(GlobalInvocationId)
102-
DEFINE_INIT_SIZES(WorkgroupSize)
103-
DEFINE_INIT_SIZES(LocalInvocationId)
104-
DEFINE_INIT_SIZES(WorkgroupId)
105-
DEFINE_INIT_SIZES(GlobalOffset)
106-
107-
#undef DEFINE_INIT_SIZES
108-
109-
#endif //__SYCL_DEVICE_ONLY__
110-
11170
class queue_impl;
11271
class stream_impl;
11372
template <typename RetType, typename Func, typename Arg>
@@ -548,7 +507,7 @@ class handler {
548507
KernelType>::type KernelFunc) {
549508
id<dimensions> global_id;
550509

551-
detail::initGlobalInvocationId<dimensions>(global_id);
510+
__spirv::initGlobalInvocationId<dimensions>(global_id);
552511

553512
KernelFunc(global_id);
554513
}
@@ -562,8 +521,8 @@ class handler {
562521
id<dimensions> global_id;
563522
range<dimensions> global_size;
564523

565-
detail::initGlobalInvocationId<dimensions>(global_id);
566-
detail::initGlobalSize<dimensions>(global_size);
524+
__spirv::initGlobalInvocationId<dimensions>(global_id);
525+
__spirv::initGlobalSize<dimensions>(global_size);
567526

568527
item<dimensions, false> Item =
569528
detail::Builder::createItem<dimensions, false>(global_size, global_id);
@@ -583,12 +542,12 @@ class handler {
583542
id<dimensions> local_id;
584543
id<dimensions> global_offset;
585544

586-
detail::initGlobalSize<dimensions>(global_size);
587-
detail::initWorkgroupSize<dimensions>(local_size);
588-
detail::initWorkgroupId<dimensions>(group_id);
589-
detail::initGlobalInvocationId<dimensions>(global_id);
590-
detail::initLocalInvocationId<dimensions>(local_id);
591-
detail::initGlobalOffset<dimensions>(global_offset);
545+
__spirv::initGlobalSize<dimensions>(global_size);
546+
__spirv::initWorkgroupSize<dimensions>(local_size);
547+
__spirv::initWorkgroupId<dimensions>(group_id);
548+
__spirv::initGlobalInvocationId<dimensions>(global_id);
549+
__spirv::initLocalInvocationId<dimensions>(local_id);
550+
__spirv::initGlobalOffset<dimensions>(global_offset);
592551

593552
group<dimensions> Group = detail::Builder::createGroup<dimensions>(
594553
global_size, local_size, group_id);

0 commit comments

Comments
 (0)