Skip to content

Commit ca5cc18

Browse files
againullagainull
authored
[SYCL] Implement sycl_ext_oneapi_get_kernel_info extension (#15650)
Extension: [sycl_ext_oneapi_get_kernel_info](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_get_kernel_info.asciidoc) Outlined kernel::attributes test to a separate file with existing XFAIL markings as it is the only descriptor which has a problem (GSD-8971) and reenabled kernel_info.cpp for other descriptors.
1 parent 66f2868 commit ca5cc18

File tree

6 files changed

+208
-64
lines changed

6 files changed

+208
-64
lines changed

‎sycl/doc/extensions/proposed/sycl_ext_oneapi_get_kernel_info.asciidoc‎ renamed to ‎sycl/doc/extensions/supported/sycl_ext_oneapi_get_kernel_info.asciidoc‎

Lines changed: 1 addition & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -43,11 +43,7 @@ SYCL specification refer to that revision.
4343

4444
== Status
4545

46-
This is a proposed extension specification, intended to gather community
47-
feedback. Interfaces defined in this specification may not be implemented yet
48-
or may be in a preliminary state. The specification itself may also change in
49-
incompatible ways before it is finalized. *Shipping software products should
50-
not rely on APIs defined in this specification.*
46+
This extension is implemented and fully supported by DPC++.
5147

5248

5349
== Overview

‎sycl/include/sycl/ext/oneapi/get_kernel_info.hpp‎

Lines changed: 48 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,48 @@
1+
//==----- get_kernel_info.hpp --- SYCL get_kernel_info extension -------==//
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+
#include <sycl/context.hpp>
11+
#include <sycl/detail/export.hpp>
12+
#include <sycl/detail/info_desc_helpers.hpp>
13+
#include <sycl/device.hpp>
14+
#include <sycl/queue.hpp>
15+
16+
namespace sycl {
17+
inline namespace _V1 {
18+
namespace ext::oneapi {
19+
20+
template <typename KernelName, typename Param>
21+
typename sycl::detail::is_kernel_info_desc<Param>::return_type
22+
get_kernel_info(const context &Ctx) {
23+
auto Bundle =
24+
sycl::get_kernel_bundle<KernelName, sycl::bundle_state::executable>(Ctx);
25+
return Bundle.template get_kernel<KernelName>().template get_info<Param>();
26+
}
27+
28+
template <typename KernelName, typename Param>
29+
typename sycl::detail::is_kernel_device_specific_info_desc<Param>::return_type
30+
get_kernel_info(const context &Ctx, const device &Dev) {
31+
auto Bundle =
32+
sycl::get_kernel_bundle<KernelName, sycl::bundle_state::executable>(Ctx);
33+
return Bundle.template get_kernel<KernelName>().template get_info<Param>(Dev);
34+
}
35+
36+
template <typename KernelName, typename Param>
37+
typename sycl::detail::is_kernel_device_specific_info_desc<Param>::return_type
38+
get_kernel_info(const queue &Q) {
39+
auto Bundle =
40+
sycl::get_kernel_bundle<KernelName, sycl::bundle_state::executable>(
41+
Q.get_context());
42+
return Bundle.template get_kernel<KernelName>().template get_info<Param>(
43+
Q.get_device());
44+
}
45+
46+
} // namespace ext::oneapi
47+
} // namespace _V1
48+
} // namespace sycl

‎sycl/include/sycl/sycl.hpp‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -103,6 +103,7 @@
103103
#include <sycl/ext/oneapi/filter_selector.hpp>
104104
#include <sycl/ext/oneapi/free_function_queries.hpp>
105105
#include <sycl/ext/oneapi/functional.hpp>
106+
#include <sycl/ext/oneapi/get_kernel_info.hpp>
106107
#include <sycl/ext/oneapi/group_local_memory.hpp>
107108
#include <sycl/ext/oneapi/kernel_properties/properties.hpp>
108109
#include <sycl/ext/oneapi/matrix/matrix.hpp>

‎sycl/source/feature_test.hpp.in‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -108,6 +108,7 @@ inline namespace _V1 {
108108
#define SYCL_EXT_ONEAPI_RAW_KERNEL_ARG 1
109109
#define SYCL_EXT_ONEAPI_PROFILING_TAG 1
110110
#define SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND 1
111+
#define SYCL_EXT_ONEAPI_GET_KERNEL_INFO 1
111112
// In progress yet
112113
#define SYCL_EXT_ONEAPI_ATOMIC16 0
113114

‎sycl/test-e2e/Basic/kernel_info.cpp‎

Lines changed: 112 additions & 59 deletions
Original file line numberDiff line numberDiff line change
@@ -1,12 +1,6 @@
11
// RUN: %{build} -o %t.out
22
// RUN: %{run} %t.out
33
//
4-
// Fail is flaky for level_zero, enable when fixed.
5-
// UNSUPPORTED: level_zero
6-
//
7-
// Consistently fails with opencl gpu, enable when fixed.
8-
// XFAIL: opencl && gpu
9-
104
//==--- kernel_info.cpp - SYCL kernel info test ----------------------------==//
115
//
126
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
@@ -17,17 +11,36 @@
1711

1812
#include <cassert>
1913
#include <sycl/detail/core.hpp>
14+
#include <sycl/ext/oneapi/get_kernel_info.hpp>
2015

2116
using namespace sycl;
17+
namespace syclex = sycl::ext::oneapi;
18+
19+
auto checkExceptionIsThrown = [](auto &getInfoFunc,
20+
const std::string &refErrMsg,
21+
std::error_code refErrc) {
22+
std::string errMsg = "";
23+
std::error_code errc;
24+
bool exceptionWasThrown = false;
25+
try {
26+
std::ignore = getInfoFunc();
27+
} catch (exception &e) {
28+
errMsg = e.what();
29+
errc = e.code();
30+
exceptionWasThrown = true;
31+
}
32+
assert(exceptionWasThrown);
33+
assert(errMsg == refErrMsg);
34+
assert(errc == refErrc);
35+
};
2236

2337
int main() {
2438
queue q;
25-
39+
auto ctx = q.get_context();
2640
buffer<int, 1> buf(range<1>(1));
27-
auto KernelID = sycl::get_kernel_id<class SingleTask>();
28-
auto KB =
29-
get_kernel_bundle<bundle_state::executable>(q.get_context(), {KernelID});
30-
kernel krn = KB.get_kernel(KernelID);
41+
auto kernelID = sycl::get_kernel_id<class SingleTask>();
42+
auto kb = get_kernel_bundle<bundle_state::executable>(ctx, {kernelID});
43+
kernel krn = kb.get_kernel(kernelID);
3144

3245
q.submit([&](handler &cgh) {
3346
auto acc = buf.get_access<access::mode::read_write>(cgh);
@@ -37,30 +50,34 @@ int main() {
3750
const std::string krnName = krn.get_info<info::kernel::function_name>();
3851
assert(!krnName.empty());
3952

40-
std::string ErrMsg = "";
41-
std::error_code Errc;
42-
bool ExceptionWasThrown = false;
43-
try {
44-
const cl_uint krnArgCount = krn.get_info<info::kernel::num_args>();
45-
} catch (exception &e) {
46-
ErrMsg = e.what();
47-
Errc = e.code();
48-
ExceptionWasThrown = true;
49-
}
50-
assert(ExceptionWasThrown && "Invalid using of \"info::kernel::num_args\" "
51-
"query should throw an exception.");
52-
assert(ErrMsg ==
53-
"info::kernel::num_args descriptor may only be used to query a kernel "
54-
"that resides in a kernel bundle constructed using a backend specific"
55-
"interoperability function or to query a device built-in kernel");
56-
assert(Errc == errc::invalid);
53+
auto refErrMsg =
54+
"info::kernel::num_args descriptor may only be used to query a kernel "
55+
"that resides in a kernel bundle constructed using a backend specific"
56+
"interoperability function or to query a device built-in kernel";
57+
auto refErrc = errc::invalid;
58+
auto getInfoNumArgsFunc = [&]() -> cl_uint {
59+
return krn.get_info<info::kernel::num_args>();
60+
};
61+
checkExceptionIsThrown(getInfoNumArgsFunc, refErrMsg, refErrc);
62+
auto getInfoNumArgsFuncExt = [&]() {
63+
return syclex::get_kernel_info<SingleTask, info::kernel::num_args>(ctx);
64+
};
65+
checkExceptionIsThrown(getInfoNumArgsFuncExt, refErrMsg, refErrc);
5766

5867
const context krnCtx = krn.get_info<info::kernel::context>();
5968
assert(krnCtx == q.get_context());
6069
const cl_uint krnRefCount = krn.get_info<info::kernel::reference_count>();
6170
assert(krnRefCount > 0);
62-
const std::string krnAttr = krn.get_info<info::kernel::attributes>();
63-
assert(krnAttr.empty());
71+
72+
// Use ext_oneapi_get_kernel_info extension and check that answers match.
73+
const context krnCtxExt =
74+
syclex::get_kernel_info<SingleTask, info::kernel::context>(ctx);
75+
assert(krnCtxExt == krnCtx);
76+
// Reference count might be different because we have to retain the kernel
77+
// handle first to fetch the info. So just check that it is not 0.
78+
const cl_uint krnRefCountExt =
79+
syclex::get_kernel_info<SingleTask, info::kernel::reference_count>(ctx);
80+
assert(krnRefCountExt > 0);
6481

6582
device dev = q.get_device();
6683
const size_t wgSize =
@@ -82,34 +99,70 @@ int main() {
8299
krn.get_info<info::kernel_device_specific::compile_num_sub_groups>(dev);
83100
assert(compileNumSg <= maxNumSg);
84101

85-
{
86-
std::error_code Errc;
87-
std::string ErrMsg = "";
88-
bool IsExceptionThrown = false;
89-
try {
90-
krn.get_info<sycl::info::kernel_device_specific::global_work_size>(dev);
91-
auto BuiltInIds = dev.get_info<info::device::built_in_kernel_ids>();
92-
bool isBuiltInKernel = std::find(BuiltInIds.begin(), BuiltInIds.end(),
93-
KernelID) != BuiltInIds.end();
94-
bool isCustomDevice = dev.get_info<sycl::info::device::device_type>() ==
95-
sycl::info::device_type::custom;
96-
assert((isCustomDevice || isBuiltInKernel) &&
97-
"info::kernel_device_specific::global_work_size descriptor can "
98-
"only be used with custom device "
99-
"or built-in kernel.");
102+
// Use ext_oneapi_get_kernel_info extension and check that answers match.
103+
const size_t wgSizeExt = syclex::get_kernel_info<
104+
SingleTask, info::kernel_device_specific::work_group_size>(ctx, dev);
105+
assert(wgSizeExt == wgSize);
106+
const size_t prefWGSizeMultExt = syclex::get_kernel_info<
107+
SingleTask,
108+
info::kernel_device_specific::preferred_work_group_size_multiple>(ctx,
109+
dev);
110+
assert(prefWGSizeMultExt == prefWGSizeMult);
111+
const cl_uint maxSgSizeExt = syclex::get_kernel_info<
112+
SingleTask, info::kernel_device_specific::max_sub_group_size>(ctx, dev);
113+
assert(maxSgSizeExt == maxSgSize);
114+
const cl_uint compileSgSizeExt = syclex::get_kernel_info<
115+
SingleTask, info::kernel_device_specific::compile_sub_group_size>(ctx,
116+
dev);
117+
assert(compileSgSizeExt == compileSgSize);
118+
const cl_uint maxNumSgExt = syclex::get_kernel_info<
119+
SingleTask, info::kernel_device_specific::max_num_sub_groups>(ctx, dev);
120+
assert(maxNumSgExt == maxNumSg);
121+
const cl_uint compileNumSgExt = syclex::get_kernel_info<
122+
SingleTask, info::kernel_device_specific::compile_num_sub_groups>(ctx,
123+
dev);
124+
assert(compileNumSgExt == compileNumSg);
100125

101-
} catch (sycl::exception &e) {
102-
IsExceptionThrown = true;
103-
Errc = e.code();
104-
ErrMsg = e.what();
105-
}
106-
assert(IsExceptionThrown &&
107-
"Invalid using of info::kernel_device_specific::global_work_size "
108-
"query should throw an exception.");
109-
assert(Errc == errc::invalid);
110-
assert(ErrMsg ==
111-
"info::kernel_device_specific::global_work_size descriptor may only "
112-
"be used if the device type is device_type::custom or if the "
113-
"kernel is a built-in kernel.");
114-
}
126+
// Use ext_oneapi_get_kernel_info extension with queue parameter and check the
127+
// result.
128+
const size_t wgSizeExtQ =
129+
syclex::get_kernel_info<SingleTask,
130+
info::kernel_device_specific::work_group_size>(q);
131+
assert(wgSizeExtQ == wgSize);
132+
const size_t prefWGSizeMultExtQ = syclex::get_kernel_info<
133+
SingleTask,
134+
info::kernel_device_specific::preferred_work_group_size_multiple>(q);
135+
assert(prefWGSizeMultExtQ == prefWGSizeMult);
136+
const cl_uint maxSgSizeExtQ = syclex::get_kernel_info<
137+
SingleTask, info::kernel_device_specific::max_sub_group_size>(q);
138+
assert(maxSgSizeExtQ == maxSgSize);
139+
const cl_uint compileSgSizeExtQ = syclex::get_kernel_info<
140+
SingleTask, info::kernel_device_specific::compile_sub_group_size>(q);
141+
assert(compileSgSizeExtQ == compileSgSize);
142+
const cl_uint maxNumSgExtQ = syclex::get_kernel_info<
143+
SingleTask, info::kernel_device_specific::max_num_sub_groups>(q);
144+
assert(maxNumSgExtQ == maxNumSg);
145+
const cl_uint compileNumSgExtQ = syclex::get_kernel_info<
146+
SingleTask, info::kernel_device_specific::compile_num_sub_groups>(q);
147+
assert(compileNumSgExtQ == compileNumSg);
148+
149+
refErrMsg =
150+
"info::kernel_device_specific::global_work_size descriptor may only "
151+
"be used if the device type is device_type::custom or if the "
152+
"kernel is a built-in kernel.";
153+
auto getInfoGWSFunc = [&]() {
154+
return krn.get_info<sycl::info::kernel_device_specific::global_work_size>(
155+
dev);
156+
};
157+
checkExceptionIsThrown(getInfoGWSFunc, refErrMsg, refErrc);
158+
auto getInfoGWSFuncExt = [&]() {
159+
return syclex::get_kernel_info<
160+
SingleTask, info::kernel_device_specific::global_work_size>(ctx, dev);
161+
};
162+
checkExceptionIsThrown(getInfoGWSFuncExt, refErrMsg, refErrc);
163+
auto getInfoGWSFuncExtQ = [&]() {
164+
return syclex::get_kernel_info<
165+
SingleTask, info::kernel_device_specific::global_work_size>(q);
166+
};
167+
checkExceptionIsThrown(getInfoGWSFuncExtQ, refErrMsg, refErrc);
115168
}

‎sycl/test-e2e/Basic/kernel_info_attr.cpp‎

Lines changed: 45 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,45 @@
1+
// RUN: %{build} -o %t.out
2+
// RUN: %{run} %t.out
3+
//
4+
// Fail is flaky for level_zero, enable when fixed.
5+
// UNSUPPORTED: level_zero
6+
//
7+
// Consistently fails with opencl gpu, enable when fixed.
8+
// XFAIL: opencl && gpu
9+
// XFAIL-TRACKER: GSD-8971
10+
11+
//==--- kernel_info_attr.cpp - SYCL info::kernel::attributes test ---==//
12+
//
13+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
14+
// See https://llvm.org/LICENSE.txt for license information.
15+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
16+
//
17+
//===---------------------------------------------------------------===//
18+
19+
#include <cassert>
20+
#include <sycl/detail/core.hpp>
21+
#include <sycl/ext/oneapi/get_kernel_info.hpp>
22+
23+
using namespace sycl;
24+
namespace syclex = sycl::ext::oneapi;
25+
26+
int main() {
27+
queue q;
28+
auto ctx = q.get_context();
29+
buffer<int, 1> buf(range<1>(1));
30+
auto KernelID = sycl::get_kernel_id<class SingleTask>();
31+
auto KB = get_kernel_bundle<bundle_state::executable>(ctx, {KernelID});
32+
kernel krn = KB.get_kernel(KernelID);
33+
34+
q.submit([&](handler &cgh) {
35+
auto acc = buf.get_access<access::mode::read_write>(cgh);
36+
cgh.single_task<class SingleTask>(krn, [=]() { acc[0] = acc[0] + 1; });
37+
});
38+
39+
const std::string krnAttr = krn.get_info<info::kernel::attributes>();
40+
assert(krnAttr.empty());
41+
const std::string krnAttrExt =
42+
syclex::get_kernel_info<SingleTask, info::kernel::attributes>(ctx);
43+
assert(krnAttr == krnAttrExt);
44+
return 0;
45+
}

0 commit comments

Comments
 (0)