Skip to content

Commit 24ae95b

Browse files
author
Alexander Batashev
authored
[SYCL] Implement SYCL_INTEL_device_specific_kernel_queries (#2549)
Also adjust subgroups implementation
1 parent da8929e commit 24ae95b

File tree

11 files changed

+291
-69
lines changed

11 files changed

+291
-69
lines changed

‎sycl/include/CL/sycl/info/info_desc.hpp‎

Lines changed: 34 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -214,6 +214,19 @@ enum class kernel_sub_group : cl_kernel_sub_group_info {
214214
compile_sub_group_size = CL_KERNEL_COMPILE_SUB_GROUP_SIZE_INTEL
215215
};
216216

217+
enum class kernel_device_specific : cl_kernel_work_group_info {
218+
global_work_size = CL_KERNEL_GLOBAL_WORK_SIZE,
219+
work_group_size = CL_KERNEL_WORK_GROUP_SIZE,
220+
compile_work_group_size = CL_KERNEL_COMPILE_WORK_GROUP_SIZE,
221+
preferred_work_group_size_multiple =
222+
CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE,
223+
private_mem_size = CL_KERNEL_PRIVATE_MEM_SIZE,
224+
max_sub_group_size = CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE,
225+
max_num_sub_groups = CL_KERNEL_MAX_NUM_SUB_GROUPS,
226+
compile_num_sub_groups = CL_KERNEL_COMPILE_NUM_SUB_GROUPS,
227+
compile_sub_group_size = CL_KERNEL_COMPILE_SUB_GROUP_SIZE_INTEL
228+
};
229+
217230
// A.6 Program information desctiptors
218231
enum class program : cl_program_info {
219232
context = CL_PROGRAM_CONTEXT,
@@ -242,6 +255,8 @@ enum class event_profiling : cl_profiling_info {
242255
// Provide an alias to the return type for each of the info parameters
243256
template <typename T, T param> class param_traits {};
244257

258+
template <typename T, T param> struct compatibility_param_traits {};
259+
245260
#define PARAM_TRAITS_SPEC(param_type, param, ret_type) \
246261
template <> class param_traits<param_type, param_type::param> { \
247262
public: \
@@ -263,6 +278,7 @@ template <typename T, T param> class param_traits {};
263278

264279
#include <CL/sycl/info/event_profiling_traits.def>
265280

281+
#include <CL/sycl/info/kernel_device_specific_traits.def>
266282
#include <CL/sycl/info/kernel_sub_group_traits.def>
267283
#include <CL/sycl/info/kernel_traits.def>
268284
#include <CL/sycl/info/kernel_work_group_traits.def>
@@ -276,6 +292,24 @@ template <typename T, T param> class param_traits {};
276292
#undef PARAM_TRAITS_SPEC
277293
#undef PARAM_TRAITS_SPEC_WITH_INPUT
278294

295+
#define PARAM_TRAITS_SPEC(param_type, param, ret_type) \
296+
template <> \
297+
struct compatibility_param_traits<param_type, param_type::param> { \
298+
static constexpr auto value = kernel_device_specific::param; \
299+
};
300+
301+
#define PARAM_TRAITS_SPEC_WITH_INPUT(param_type, param, ret_type, in_type) \
302+
template <> \
303+
struct compatibility_param_traits<param_type, param_type::param> { \
304+
static constexpr auto value = kernel_device_specific::param; \
305+
};
306+
307+
#include <CL/sycl/info/kernel_sub_group_traits.def>
308+
#include <CL/sycl/info/kernel_work_group_traits.def>
309+
310+
#undef PARAM_TRAITS_SPEC
311+
#undef PARAM_TRAITS_SPEC_WITH_INPUT
312+
279313
} // namespace info
280314
} // namespace sycl
281315
} // __SYCL_INLINE_NAMESPACE(cl)

‎sycl/include/CL/sycl/info/kernel_device_specific_traits.def‎

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,12 @@
1+
PARAM_TRAITS_SPEC(kernel_device_specific, compile_work_group_size,
2+
cl::sycl::range<3>)
3+
PARAM_TRAITS_SPEC(kernel_device_specific, global_work_size, cl::sycl::range<3>)
4+
PARAM_TRAITS_SPEC(kernel_device_specific,
5+
preferred_work_group_size_multiple, size_t)
6+
PARAM_TRAITS_SPEC(kernel_device_specific, private_mem_size, cl_ulong)
7+
PARAM_TRAITS_SPEC(kernel_device_specific, work_group_size, size_t)
8+
PARAM_TRAITS_SPEC_WITH_INPUT(kernel_device_specific, max_sub_group_size,
9+
uint32_t, cl::sycl::range<3>)
10+
PARAM_TRAITS_SPEC(kernel_device_specific, max_num_sub_groups, uint32_t)
11+
PARAM_TRAITS_SPEC(kernel_device_specific, compile_num_sub_groups, uint32_t)
12+
PARAM_TRAITS_SPEC(kernel_device_specific, compile_sub_group_size, uint32_t)

‎sycl/include/CL/sycl/kernel.hpp‎

Lines changed: 30 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -92,6 +92,27 @@ class __SYCL_EXPORT kernel {
9292
typename info::param_traits<info::kernel, param>::return_type
9393
get_info() const;
9494

95+
/// Query device-specific information from the kernel object using the
96+
/// info::kernel_device_specific descriptor.
97+
///
98+
/// \param Device is a valid SYCL device to query info for.
99+
/// \return depends on information being queried.
100+
template <info::kernel_device_specific param>
101+
typename info::param_traits<info::kernel_device_specific, param>::return_type
102+
get_info(const device &Device) const;
103+
104+
/// Query device-specific information from a kernel using the
105+
/// info::kernel_device_specific descriptor for a specific device and value.
106+
///
107+
/// \param Device is a valid SYCL device.
108+
/// \param Value depends on information being queried.
109+
/// \return depends on information being queried.
110+
template <info::kernel_device_specific param>
111+
typename info::param_traits<info::kernel_device_specific, param>::return_type
112+
get_info(const device &Device,
113+
typename info::param_traits<info::kernel_device_specific,
114+
param>::input_type Value) const;
115+
95116
/// Query work-group information from a kernel using the
96117
/// info::kernel_work_group descriptor for a specific device.
97118
///
@@ -107,8 +128,11 @@ class __SYCL_EXPORT kernel {
107128
/// \param Device is a valid SYCL device.
108129
/// \return depends on information being queried.
109130
template <info::kernel_sub_group param>
131+
// clang-format off
110132
typename info::param_traits<info::kernel_sub_group, param>::return_type
133+
__SYCL_DEPRECATED("Use get_info with info::kernel_device_specific instead.")
111134
get_sub_group_info(const device &Device) const;
135+
// clang-format on
112136

113137
/// Query sub-group information from a kernel using the
114138
/// info::kernel_sub_group descriptor for a specific device and value.
@@ -117,11 +141,13 @@ class __SYCL_EXPORT kernel {
117141
/// \param Value depends on information being queried.
118142
/// \return depends on information being queried.
119143
template <info::kernel_sub_group param>
144+
// clang-format off
120145
typename info::param_traits<info::kernel_sub_group, param>::return_type
121-
get_sub_group_info(
122-
const device &Device,
123-
typename info::param_traits<info::kernel_sub_group, param>::input_type
124-
Value) const;
146+
__SYCL_DEPRECATED("Use get_info with info::kernel_device_specific instead.")
147+
get_sub_group_info(const device &Device,
148+
typename info::param_traits<info::kernel_sub_group,
149+
param>::input_type Value) const;
150+
// clang-format on
125151

126152
private:
127153
/// Constructs a SYCL kernel object from a valid kernel_impl instance.

‎sycl/source/detail/kernel_impl.cpp‎

Lines changed: 46 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -82,43 +82,59 @@ template <> program kernel_impl::get_info<info::kernel::program>() const {
8282
return createSyclObjFromImpl<program>(MProgramImpl);
8383
}
8484

85-
template <info::kernel_work_group param>
86-
typename info::param_traits<info::kernel_work_group, param>::return_type
87-
kernel_impl::get_work_group_info(const device &Device) const {
85+
template <info::kernel_device_specific param>
86+
typename info::param_traits<info::kernel_device_specific, param>::return_type
87+
kernel_impl::get_info(const device &Device) const {
8888
if (is_host()) {
89-
return get_kernel_work_group_info_host<param>(Device);
89+
return get_kernel_device_specific_info_host<param>(Device);
9090
}
91-
return get_kernel_work_group_info<
92-
typename info::param_traits<info::kernel_work_group, param>::return_type,
91+
return get_kernel_device_specific_info<
92+
typename info::param_traits<info::kernel_device_specific,
93+
param>::return_type,
9394
param>::get(this->getHandleRef(), getSyclObjImpl(Device)->getHandleRef(),
9495
getPlugin());
9596
}
9697

97-
template <info::kernel_sub_group param>
98-
typename info::param_traits<info::kernel_sub_group, param>::return_type
99-
kernel_impl::get_sub_group_info(const device &Device) const {
98+
template <info::kernel_device_specific param>
99+
typename info::param_traits<info::kernel_device_specific, param>::return_type
100+
kernel_impl::get_info(
101+
const device &Device,
102+
typename info::param_traits<info::kernel_device_specific, param>::input_type
103+
Value) const {
100104
if (is_host()) {
101105
throw runtime_error("Sub-group feature is not supported on HOST device.",
102106
PI_INVALID_DEVICE);
103107
}
104-
return get_kernel_sub_group_info<param>::get(
105-
this->getHandleRef(), getSyclObjImpl(Device)->getHandleRef(),
108+
return get_kernel_device_specific_info_with_input<param>::get(
109+
this->getHandleRef(), getSyclObjImpl(Device)->getHandleRef(), Value,
106110
getPlugin());
107111
}
108112

113+
template <info::kernel_work_group param>
114+
typename info::param_traits<info::kernel_work_group, param>::return_type
115+
kernel_impl::get_work_group_info(const device &Device) const {
116+
return get_info<
117+
info::compatibility_param_traits<info::kernel_work_group, param>::value>(
118+
Device);
119+
}
120+
121+
template <info::kernel_sub_group param>
122+
typename info::param_traits<info::kernel_sub_group, param>::return_type
123+
kernel_impl::get_sub_group_info(const device &Device) const {
124+
return get_info<
125+
info::compatibility_param_traits<info::kernel_sub_group, param>::value>(
126+
Device);
127+
}
128+
109129
template <info::kernel_sub_group param>
110130
typename info::param_traits<info::kernel_sub_group, param>::return_type
111131
kernel_impl::get_sub_group_info(
112132
const device &Device,
113133
typename info::param_traits<info::kernel_sub_group, param>::input_type
114134
Value) const {
115-
if (is_host()) {
116-
throw runtime_error("Sub-group feature is not supported on HOST device.",
117-
PI_INVALID_DEVICE);
118-
}
119-
return get_kernel_sub_group_info_with_input<param>::get(
120-
this->getHandleRef(), getSyclObjImpl(Device)->getHandleRef(), Value,
121-
getPlugin());
135+
return get_info<
136+
info::compatibility_param_traits<info::kernel_sub_group, param>::value>(
137+
Device, Value);
122138
}
123139

124140
#define PARAM_TRAITS_SPEC(param_type, param, ret_type) \
@@ -128,6 +144,18 @@ kernel_impl::get_sub_group_info(
128144

129145
#undef PARAM_TRAITS_SPEC
130146

147+
#define PARAM_TRAITS_SPEC(param_type, param, ret_type) \
148+
template ret_type kernel_impl::get_info<info::param_type::param>( \
149+
const device &) const;
150+
#define PARAM_TRAITS_SPEC_WITH_INPUT(param_type, param, ret_type, in_type) \
151+
template ret_type kernel_impl::get_info<info::param_type::param>( \
152+
const device &, in_type) const;
153+
154+
#include <CL/sycl/info/kernel_device_specific_traits.def>
155+
156+
#undef PARAM_TRAITS_SPEC
157+
#undef PARAM_TRAITS_SPEC_WITH_INPUT
158+
131159
#define PARAM_TRAITS_SPEC(param_type, param, ret_type) \
132160
template ret_type kernel_impl::get_work_group_info<info::param_type::param>( \
133161
const device &) const;

‎sycl/source/detail/kernel_impl.hpp‎

Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -103,6 +103,27 @@ class kernel_impl {
103103
typename info::param_traits<info::kernel, param>::return_type
104104
get_info() const;
105105

106+
/// Query device-specific information from a kernel object using the
107+
/// info::kernel_device_specific descriptor.
108+
///
109+
/// \param Device is a valid SYCL device to query info for.
110+
/// \return depends on information being queried.
111+
template <info::kernel_device_specific param>
112+
typename info::param_traits<info::kernel_device_specific, param>::return_type
113+
get_info(const device &Device) const;
114+
115+
/// Query device-specific information from a kernel using the
116+
/// info::kernel_device_specific descriptor for a specific device and value.
117+
///
118+
/// \param Device is a valid SYCL device.
119+
/// \param Value depends on information being queried.
120+
/// \return depends on information being queried.
121+
template <info::kernel_device_specific param>
122+
typename info::param_traits<info::kernel_device_specific, param>::return_type
123+
get_info(const device &Device,
124+
typename info::param_traits<info::kernel_device_specific,
125+
param>::input_type Value) const;
126+
106127
/// Query work-group information from a kernel using the
107128
/// info::kernel_work_group descriptor for a specific device.
108129
///

0 commit comments

Comments
 (0)