Skip to content

Commit 7ea6b96

Browse files
oleksandr-pavlykoleksandr-pavlyk
committed
Added all but one kernel_device_specific properties
max_sub_group_size property is currently on hold due to an issue in DPC++ runtime
1 parent bcf1a14 commit 7ea6b96

File tree

3 files changed

+434
-28
lines changed

3 files changed

+434
-28
lines changed

‎libsyclinterface/include/dpctl_sycl_kernel_interface.h

Lines changed: 104 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -39,25 +39,124 @@ DPCTL_C_EXTERN_C_BEGIN
3939
*/
4040

4141
/*!
42-
* @brief Returns the number of arguments for the OpenCL kernel.
42+
* @brief Returns the number of arguments for the sycl
43+
* interoperability kernel.
4344
*
44-
* @param KRef DPCTLSyclKernelRef pointer to an OpenCL
45+
* @param KRef DPCTLSyclKernelRef pointer to an SYCL
4546
* interoperability kernel.
46-
* @return Returns the number of arguments for the OpenCL interoperability
47+
* @return Returns the number of arguments for the interoperability
4748
* kernel.
4849
* @ingroup KernelInterface
4950
*/
5051
DPCTL_API
5152
size_t DPCTLKernel_GetNumArgs(__dpctl_keep const DPCTLSyclKernelRef KRef);
5253

5354
/*!
54-
* @brief Deletes the DPCTLSyclKernelRef after casting it to a sycl::kernel.
55+
* @brief Deletes the DPCTLSyclKernelRef after casting it to a
56+
* ``sycl::kernel``.
5557
*
56-
* @param KRef DPCTLSyclKernelRef pointer to an OpenCL
58+
* @param KRef DPCTLSyclKernelRef pointer to an SYCL
5759
* interoperability kernel.
5860
* @ingroup KernelInterface
5961
*/
6062
DPCTL_API
6163
void DPCTLKernel_Delete(__dpctl_take DPCTLSyclKernelRef KRef);
6264

65+
/*!
66+
* !brief Wrapper around
67+
* `kernel::get_info<info::kernel_device_specific::work_group_size>()`.
68+
*
69+
* @param KRef DPCTLSyclKernelRef pointer to an SYCL
70+
* interoperability kernel.
71+
* @return Returns the maximum number of work-items in a work-group
72+
* that can be used to execute a kernel on a specific device.
73+
* @ingroup KernelInterface
74+
*/
75+
DPCTL_API
76+
size_t DPCTLKernel_GetWorkGroupSize(__dpctl_keep const DPCTLSyclKernelRef KRef);
77+
78+
/*!
79+
* !brief Wrapper around
80+
* `kernel::get_info<info::kernel_device_specific::preferred_work_group_size_multiple>()`.
81+
*
82+
* @param KRef DPCTLSyclKernelRef pointer to an SYCL
83+
* interoperability kernel.
84+
* @return Returns a value, of which work-group size is preferred to be a
85+
* multiple, for executing a kernel on a specific device.
86+
* @ingroup KernelInterface
87+
*/
88+
DPCTL_API
89+
size_t DPCTLKernel_GetPreferredWorkGroupSizeMultiple(
90+
__dpctl_keep const DPCTLSyclKernelRef KRef);
91+
92+
/*!
93+
* !brief Wrapper around
94+
* `kernel::get_info<info::kernel_device_specific::private_mem_size>()`.
95+
*
96+
* @param KRef DPCTLSyclKernelRef pointer to an SYCL
97+
* interoperability kernel.
98+
* @return Returns the minimum amount of private memory, in bytes,
99+
* used by each work-item in the kernel.
100+
* @ingroup KernelInterface
101+
*/
102+
DPCTL_API
103+
size_t
104+
DPCTLKernel_GetPrivateMemSize(__dpctl_keep const DPCTLSyclKernelRef KRef);
105+
106+
/*!
107+
* !brief Wrapper around
108+
* `kernel::get_info<info::kernel_device_specific::max_num_sub_groups>()`.
109+
*
110+
* @param KRef DPCTLSyclKernelRef pointer to an SYCL
111+
* interoperability kernel.
112+
* @return Returns the maximum number of sub-groups for this kernel.
113+
* @ingroup KernelInterface
114+
*/
115+
DPCTL_API
116+
uint32_t
117+
DPCTLKernel_GetMaxNumSubGroups(__dpctl_keep const DPCTLSyclKernelRef KRef);
118+
119+
#if 0
120+
/*!
121+
* !brief Wrapper around
122+
* `kernel::get_info<info::kernel_device_specific::max_sub_group_size>()`.
123+
*
124+
* @param KRef DPCTLSyclKernelRef pointer to an SYCL
125+
* interoperability kernel.
126+
* @return Returns the maximum sub-group size for this kernel.
127+
* @ingroup KernelInterface
128+
*/
129+
DPCTL_API
130+
uint32_t
131+
DPCTLKernel_GetMaxSubGroupSize(__dpctl_keep const DPCTLSyclKernelRef KRef);
132+
#endif
133+
134+
/*!
135+
* !brief Wrapper around
136+
* `kernel::get_info<info::kernel_device_specific::compile_num_sub_groups>()`.
137+
*
138+
* @param KRef DPCTLSyclKernelRef pointer to an SYCL
139+
* interoperability kernel.
140+
* @return Returns the number of sub-groups specified by the kernel,
141+
* or 0 (if not specified).
142+
* @ingroup KernelInterface
143+
*/
144+
DPCTL_API
145+
uint32_t
146+
DPCTLKernel_GetCompileNumSubGroups(__dpctl_keep const DPCTLSyclKernelRef KRef);
147+
148+
/*!
149+
* !brief Wrapper around
150+
* `kernel::get_info<info::kernel_device_specific::compile_sub_group_size>()`.
151+
*
152+
* @param KRef DPCTLSyclKernelRef pointer to an SYCL
153+
* interoperability kernel.
154+
* @return Returns the required sub-group size specified by this kernel,
155+
* or 0 (if not specified).
156+
* @ingroup KernelInterface
157+
*/
158+
DPCTL_API
159+
uint32_t
160+
DPCTLKernel_GetCompileSubGroupSize(__dpctl_keep const DPCTLSyclKernelRef KRef);
161+
63162
DPCTL_C_EXTERN_C_END

‎libsyclinterface/source/dpctl_sycl_kernel_interface.cpp

Lines changed: 164 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -29,6 +29,7 @@
2929
#include "dpctl_error_handlers.h"
3030
#include "dpctl_string_utils.hpp"
3131
#include <CL/sycl.hpp> /* Sycl headers */
32+
#include <cstdint>
3233

3334
using namespace cl::sycl;
3435

@@ -39,21 +40,177 @@ DEFINE_SIMPLE_CONVERSION_FUNCTIONS(kernel, DPCTLSyclKernelRef)
3940

4041
} /* end of anonymous namespace */
4142

42-
size_t DPCTLKernel_GetNumArgs(__dpctl_keep const DPCTLSyclKernelRef Kernel)
43+
size_t DPCTLKernel_GetNumArgs(__dpctl_keep const DPCTLSyclKernelRef KRef)
4344
{
44-
if (!Kernel) {
45+
if (!KRef) {
4546
error_handler("Cannot get the number of arguments from "
4647
"DPCTLSyclKernelRef as input is a nullptr.",
4748
__FILE__, __func__, __LINE__);
4849
return -1;
4950
}
5051

51-
auto SyclKernel = unwrap(Kernel);
52-
auto num_args = SyclKernel->get_info<info::kernel::num_args>();
53-
return (size_t)num_args;
52+
auto sycl_kernel = unwrap(KRef);
53+
auto num_args = sycl_kernel->get_info<info::kernel::num_args>();
54+
return static_cast<size_t>(num_args);
5455
}
5556

56-
void DPCTLKernel_Delete(__dpctl_take DPCTLSyclKernelRef Kernel)
57+
void DPCTLKernel_Delete(__dpctl_take DPCTLSyclKernelRef KRef)
5758
{
58-
delete unwrap(Kernel);
59+
delete unwrap(KRef);
60+
}
61+
62+
size_t DPCTLKernel_GetWorkGroupSize(__dpctl_keep const DPCTLSyclKernelRef KRef)
63+
{
64+
if (!KRef) {
65+
error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__,
66+
__func__, __LINE__);
67+
return 0;
68+
}
69+
70+
auto sycl_kern = unwrap(KRef);
71+
auto devs = sycl_kern->get_kernel_bundle().get_devices();
72+
if (devs.empty()) {
73+
error_handler("Input DPCTKSyclKernelRef has no associated device.",
74+
__FILE__, __func__, __LINE__);
75+
return 0;
76+
}
77+
auto v = sycl_kern->get_info<info::kernel_device_specific::work_group_size>(
78+
devs[0]);
79+
return static_cast<size_t>(v);
80+
}
81+
82+
size_t DPCTLKernel_GetPreferredWorkGroupSizeMultiple(
83+
__dpctl_keep const DPCTLSyclKernelRef KRef)
84+
{
85+
if (!KRef) {
86+
error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__,
87+
__func__, __LINE__);
88+
return 0;
89+
}
90+
91+
auto sycl_kern = unwrap(KRef);
92+
auto devs = sycl_kern->get_kernel_bundle().get_devices();
93+
if (devs.empty()) {
94+
error_handler("Input DPCTKSyclKernelRef has no associated device.",
95+
__FILE__, __func__, __LINE__);
96+
return 0;
97+
}
98+
auto v = sycl_kern->get_info<
99+
info::kernel_device_specific::preferred_work_group_size_multiple>(
100+
devs[0]);
101+
return static_cast<size_t>(v);
102+
}
103+
104+
size_t DPCTLKernel_GetPrivateMemSize(__dpctl_keep const DPCTLSyclKernelRef KRef)
105+
{
106+
if (!KRef) {
107+
error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__,
108+
__func__, __LINE__);
109+
return 0;
110+
}
111+
112+
auto sycl_kern = unwrap(KRef);
113+
auto devs = sycl_kern->get_kernel_bundle().get_devices();
114+
if (devs.empty()) {
115+
error_handler("Input DPCTKSyclKernelRef has no associated device.",
116+
__FILE__, __func__, __LINE__);
117+
return 0;
118+
}
119+
auto v =
120+
sycl_kern->get_info<info::kernel_device_specific::private_mem_size>(
121+
devs[0]);
122+
return static_cast<size_t>(v);
123+
}
124+
125+
uint32_t
126+
DPCTLKernel_GetMaxNumSubGroups(__dpctl_keep const DPCTLSyclKernelRef KRef)
127+
{
128+
if (!KRef) {
129+
error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__,
130+
__func__, __LINE__);
131+
return 0;
132+
}
133+
134+
auto sycl_kern = unwrap(KRef);
135+
auto devs = sycl_kern->get_kernel_bundle().get_devices();
136+
if (devs.empty()) {
137+
error_handler("Input DPCTKSyclKernelRef has no associated device.",
138+
__FILE__, __func__, __LINE__);
139+
return 0;
140+
}
141+
auto v =
142+
sycl_kern->get_info<info::kernel_device_specific::max_num_sub_groups>(
143+
devs[0]);
144+
return static_cast<uint32_t>(v);
145+
}
146+
147+
#if 0
148+
// commented out due to bug in DPC++ runtime, get_info for max_sub_group_size
149+
// exported by libsycl has different, not SPEC-compliant signature
150+
uint32_t
151+
DPCTLKernel_GetMaxSubGroupSize(__dpctl_keep const DPCTLSyclKernelRef KRef)
152+
{
153+
if (!KRef) {
154+
error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__,
155+
__func__, __LINE__);
156+
return 0;
157+
}
158+
159+
auto sycl_kern = unwrap(KRef);
160+
auto devs = sycl_kern->get_kernel_bundle().get_devices();
161+
if (devs.empty()) {
162+
error_handler("Input DPCTKSyclKernelRef has no associated device.",
163+
__FILE__, __func__, __LINE__);
164+
return 0;
165+
}
166+
auto v = sycl_kern
167+
->get_info<info::kernel_device_specific::max_sub_group_size>(devs[0]);
168+
return v;
169+
}
170+
#endif
171+
172+
uint32_t
173+
DPCTLKernel_GetCompileNumSubGroups(__dpctl_keep const DPCTLSyclKernelRef KRef)
174+
{
175+
if (!KRef) {
176+
error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__,
177+
__func__, __LINE__);
178+
return 0;
179+
}
180+
181+
auto sycl_kern = unwrap(KRef);
182+
auto devs = sycl_kern->get_kernel_bundle().get_devices();
183+
if (devs.empty()) {
184+
error_handler("Input DPCTKSyclKernelRef has no associated device.",
185+
__FILE__, __func__, __LINE__);
186+
return 0;
187+
}
188+
auto v =
189+
sycl_kern
190+
->get_info<info::kernel_device_specific::compile_num_sub_groups>(
191+
devs[0]);
192+
return static_cast<uint32_t>(v);
193+
}
194+
195+
uint32_t
196+
DPCTLKernel_GetCompileSubGroupSize(__dpctl_keep const DPCTLSyclKernelRef KRef)
197+
{
198+
if (!KRef) {
199+
error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__,
200+
__func__, __LINE__);
201+
return 0;
202+
}
203+
204+
auto sycl_kern = unwrap(KRef);
205+
auto devs = sycl_kern->get_kernel_bundle().get_devices();
206+
if (devs.empty()) {
207+
error_handler("Input DPCTKSyclKernelRef has no associated device.",
208+
__FILE__, __func__, __LINE__);
209+
return 0;
210+
}
211+
auto v =
212+
sycl_kern
213+
->get_info<info::kernel_device_specific::compile_sub_group_size>(
214+
devs[0]);
215+
return static_cast<uint32_t>(v);
59216
}

0 commit comments

Comments
 (0)