Skip to content

Commit 67df4b9

Browse files
authored
[SYCL][E2E] Add free function kernel queries tests (#18904)
Add 2 tests for the free function kernel extension as per the [test plan](https://github.com/intel/llvm/blob/6af08fe9c6fd4dc8e433555646606898c71d92fb/sycl/test-e2e/FreeFunctionKernels/test-plan.md). More specifically, - Check that the 3 overloads of `get_kernel_info` for free function kernels as defined [here](https://github.com/intel/llvm/blob/6af08fe9c6fd4dc8e433555646606898c71d92fb/sycl/doc/extensions/proposed/sycl_ext_oneapi_free_function_kernels.asciidoc#new-free-functions-to-query-kernel-information-descriptors) return the same values as the regular `get_info` called on the kernel object. - Check that `get_kernel_ids` on a kernel bundle as defined [here](https://github.com/intel/llvm/blob/6af08fe9c6fd4dc8e433555646606898c71d92fb/sycl/doc/extensions/proposed/sycl_ext_oneapi_free_function_kernels.asciidoc#behavior-with-kernel-bundle-functions-in-the-core-sycl-specification) returns all kernels defined in the source be they free function kernels, lambdas or function objects.
1 parent 8cec943 commit 67df4b9

File tree

2 files changed

+91
-0
lines changed

2 files changed

+91
-0
lines changed

sycl/test-e2e/KernelAndProgram/free_function_apis.cpp

Lines changed: 36 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,7 @@
66
#include <sycl/detail/core.hpp>
77
#include <sycl/ext/oneapi/experimental/free_function_traits.hpp>
88
#include <sycl/ext/oneapi/free_function_queries.hpp>
9+
#include <sycl/ext/oneapi/get_kernel_info.hpp>
910
#include <sycl/kernel_bundle.hpp>
1011
#include <sycl/usm.hpp>
1112

@@ -207,6 +208,41 @@ bool test_bundle_apis(queue Queue) {
207208
std::cout << "PassR=" << PassR << std::endl;
208209
Pass &= PassR;
209210

211+
kernel_bundle Bundle_ff2 = ext::oneapi::experimental::get_kernel_bundle<
212+
ff_2, bundle_state::executable>(Context);
213+
214+
bool PassS =
215+
ext::oneapi::experimental::get_kernel_info<ff_2,
216+
info::kernel::function_name>(
217+
Context) == Bundle_ff2.ext_oneapi_get_kernel<ff_2>()
218+
.get_info<info::kernel::function_name>();
219+
std::cout << "Test retrieving function_name using context: " << PassS
220+
<< std::endl;
221+
Pass &= PassS;
222+
223+
kernel_bundle Bundle_ff3 = ext::oneapi::experimental::get_kernel_bundle<
224+
ff_3<int>, bundle_state::executable>(Context);
225+
226+
bool PassT =
227+
ext::oneapi::experimental::get_kernel_info<
228+
ff_3<int>, info::kernel_device_specific::work_group_size>(Context,
229+
Device) ==
230+
Bundle_ff3.ext_oneapi_get_kernel<ff_3<int>>()
231+
.get_info<info::kernel_device_specific::work_group_size>(Device);
232+
std::cout << "Test retrieving work_group_size using context and device: "
233+
<< PassT << std::endl;
234+
Pass &= PassT;
235+
236+
bool PassU =
237+
ext::oneapi::experimental::get_kernel_info<
238+
ff_3<int>, info::kernel_device_specific::work_group_size>(Queue) ==
239+
Bundle_ff3.ext_oneapi_get_kernel<ff_3<int>>()
240+
.get_info<info::kernel_device_specific::work_group_size>(
241+
Queue.get_device());
242+
std::cout << "Test retrieving work_group_size using queue: " << PassU
243+
<< std::endl;
244+
Pass &= PassU;
245+
210246
std::cout << "Test bundle APIs: " << (Pass ? "PASS" : "FAIL") << std::endl;
211247
return Pass;
212248
}
Lines changed: 55 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,55 @@
1+
// RUN: %{build} -o %t.out
2+
// RUN: %{run} %t.out
3+
4+
// This test checks that get_kernel_ids returns all the kernels defined
5+
// in the source regardless of whether they are expressed as lambdas,
6+
// function objects or free functions.
7+
8+
#include <sycl/detail/core.hpp>
9+
#include <sycl/ext/oneapi/experimental/free_function_traits.hpp>
10+
#include <sycl/kernel_bundle.hpp>
11+
12+
class FunctionObjectKernel {
13+
public:
14+
void operator()() const {}
15+
};
16+
17+
class LambdaKernel;
18+
19+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
20+
(sycl::ext::oneapi::experimental::nd_range_kernel<1>))
21+
void FreeFunctionKernel() {}
22+
23+
int main() {
24+
sycl::queue Queue;
25+
sycl::device Dev = Queue.get_device();
26+
sycl::context Context = Queue.get_context();
27+
28+
sycl::kernel_id FunctionObjectKernelId =
29+
sycl::get_kernel_id<FunctionObjectKernel>();
30+
sycl::kernel_id LambdaKernelId = sycl::get_kernel_id<LambdaKernel>();
31+
sycl::kernel_id FreeFunctionKernelId =
32+
sycl::ext::oneapi::experimental::get_kernel_id<FreeFunctionKernel>();
33+
34+
sycl::kernel_bundle KernelBundle =
35+
sycl::get_kernel_bundle<sycl::bundle_state::executable>(Context, {Dev});
36+
sycl::kernel FreeFunction = KernelBundle.get_kernel(FreeFunctionKernelId);
37+
Queue.submit(
38+
[&](sycl::handler &CGH) { CGH.single_task<LambdaKernel>([=]() {}); });
39+
Queue.submit([&](sycl::handler &CGH) {
40+
FunctionObjectKernel kernel;
41+
CGH.single_task(kernel);
42+
});
43+
Queue.submit([&](sycl::handler &CGH) {
44+
CGH.parallel_for(sycl::nd_range{{1}, {1}}, FreeFunction);
45+
});
46+
47+
std::vector<sycl::kernel_id> KernelIds = KernelBundle.get_kernel_ids();
48+
assert(KernelIds.size() == 3);
49+
for (const sycl::kernel_id &KernelId :
50+
{FunctionObjectKernelId, LambdaKernelId, FreeFunctionKernelId}) {
51+
auto FoundId = std::find(KernelIds.begin(), KernelIds.end(), KernelId);
52+
assert(FoundId != KernelIds.end());
53+
}
54+
return 0;
55+
}

0 commit comments

Comments
 (0)