Skip to content

Commit ef28856

Browse files
PietroGhgPietroGhg
authored
[SYCL][RTC] Add tests for JIT RTC (#17182)
This PR is based on #17032. It adds runtime tests that match the examples from the spec changes in #11985, showing that the test run successfully using the JIT approach for RTC. This PR also contains a few minor fixes in the example code in the RTC docs.
1 parent 0e155f0 commit ef28856

File tree

5 files changed

+249
-5
lines changed

5 files changed

+249
-5
lines changed

‎sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc

Lines changed: 6 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -816,7 +816,7 @@ int main() {
816816
extern "C"
817817
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
818818
void iota(float start, float *ptr) {
819-
size_t id = syclext::this_work_item::get_nd_item().get_global_linear_id();
819+
size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
820820
ptr[id] = start + static_cast<float>(id);
821821
}
822822
)""";
@@ -845,6 +845,7 @@ int main() {
845845
sycl::nd_range ndr{{NUM}, {WGSIZE}};
846846
cgh.parallel_for(ndr, iota);
847847
}).wait();
848+
sycl::free(ptr, q);
848849
}
849850
----
850851

@@ -872,13 +873,13 @@ int main() {
872873
873874
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
874875
void iota(float start, float *ptr) {
875-
size_t id = syclext::this_work_item::get_nd_item().get_global_linear_id();
876+
size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
876877
ptr[id] = start + static_cast<float>(id);
877878
}
878879
879-
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::range_kernel<1>))
880+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
880881
void iota(int start, int *ptr) {
881-
size_t id = syclext::this_work_item::get_nd_item().get_global_linear_id();
882+
size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
882883
ptr[id] = start + static_cast<int>(id);
883884
}
884885
)""";
@@ -910,6 +911,7 @@ int main() {
910911
sycl::nd_range ndr{{NUM}, {WGSIZE}};
911912
cgh.parallel_for(ndr, iota);
912913
}).wait();
914+
sycl::free(ptr, q);
913915
}
914916
----
915917

‎sycl/test-e2e/KernelCompiler/kernel_compiler_basic.cpp

Lines changed: 78 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,78 @@
1+
//==- kernel_compiler_sycl_jit.cpp --- kernel_compiler extension tests -----==//
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+
// REQUIRES: (opencl || level_zero)
10+
// REQUIRES: aspect-usm_device_allocations
11+
12+
// UNSUPPORTED: accelerator
13+
// UNSUPPORTED-INTENDED: while accelerator is AoT only, this cannot run there.
14+
15+
// RUN: %{build} -o %t.out
16+
// RUN: %{l0_leak_check} %{run} %t.out
17+
18+
#include <sycl/detail/core.hpp>
19+
#include <sycl/kernel_bundle.hpp>
20+
#include <sycl/usm.hpp>
21+
22+
namespace syclexp = sycl::ext::oneapi::experimental;
23+
24+
static constexpr size_t NUM = 1024;
25+
static constexpr size_t WGSIZE = 16;
26+
27+
int main() {
28+
sycl::queue q;
29+
30+
// The source code for a kernel, defined as a SYCL "free function kernel".
31+
std::string source = R"""(
32+
#include <sycl/sycl.hpp>
33+
namespace syclext = sycl::ext::oneapi;
34+
namespace syclexp = sycl::ext::oneapi::experimental;
35+
36+
extern "C"
37+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
38+
void iota(float start, float *ptr) {
39+
size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
40+
ptr[id] = start + static_cast<float>(id);
41+
}
42+
)""";
43+
44+
// Create a kernel bundle in "source" state.
45+
sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source> kb_src =
46+
syclexp::create_kernel_bundle_from_source(
47+
q.get_context(), syclexp::source_language::sycl_jit, source);
48+
49+
// Compile the kernel. There is no need to use the "registered_names"
50+
// property because the kernel is declared extern "C".
51+
sycl::kernel_bundle<sycl::bundle_state::executable> kb_exe =
52+
syclexp::build(kb_src);
53+
54+
// Get the kernel via its compiler-generated name.
55+
sycl::kernel iota = kb_exe.ext_oneapi_get_kernel("iota");
56+
57+
float *ptr = sycl::malloc_shared<float>(NUM, q);
58+
q.submit([&](sycl::handler &cgh) {
59+
// Set the values of the kernel arguments.
60+
cgh.set_args(3.14f, ptr);
61+
62+
// Launch the kernel according to its type, in this case an nd-range
63+
// kernel.
64+
sycl::nd_range ndr{{NUM}, {WGSIZE}};
65+
cgh.parallel_for(ndr, iota);
66+
}).wait();
67+
68+
constexpr float eps = 0.001;
69+
for (int i = 0; i < NUM; i++) {
70+
const float truth = 3.14f + static_cast<float>(i);
71+
if (std::abs(ptr[i] - truth) > eps) {
72+
std::cout << "Result: " << ptr[i] << " expected " << i << "\n";
73+
sycl::free(ptr, q);
74+
exit(1);
75+
}
76+
}
77+
sycl::free(ptr, q);
78+
}

‎sycl/test-e2e/KernelCompiler/kernel_compiler_namespaces.cpp

Lines changed: 81 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,81 @@
1+
//==- kernel_compiler_namespaces.cpp --- kernel_compiler extension tests ---==//
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+
// REQUIRES: (opencl || level_zero)
10+
// REQUIRES: aspect-usm_device_allocations
11+
12+
// UNSUPPORTED: accelerator
13+
// UNSUPPORTED-INTENDED: while accelerator is AoT only, this cannot run there.
14+
15+
// RUN: %{build} -o %t.out
16+
// RUN: %{l0_leak_check} %{run} %t.out
17+
18+
#include <sycl/detail/core.hpp>
19+
#include <sycl/kernel_bundle.hpp>
20+
#include <sycl/usm.hpp>
21+
namespace syclexp = sycl::ext::oneapi::experimental;
22+
23+
static constexpr size_t NUM = 1024;
24+
static constexpr size_t WGSIZE = 16;
25+
26+
int main() {
27+
sycl::queue q;
28+
29+
// The source code for two kernels defined in different namespaces
30+
std::string source = R"""(
31+
#include <sycl/sycl.hpp>
32+
namespace syclext = sycl::ext::oneapi;
33+
namespace syclexp = sycl::ext::oneapi::experimental;
34+
35+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
36+
void iota(int start, int *ptr) {
37+
size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
38+
ptr[id] = start + static_cast<int>(id);
39+
}
40+
41+
namespace mykernels {
42+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
43+
void iota(int start, int *ptr) {
44+
size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
45+
ptr[id] = start + static_cast<int>(id);
46+
}
47+
} // namespace mykernels
48+
)""";
49+
50+
// Create a kernel bundle in "source" state.
51+
sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source> kb_src =
52+
syclexp::create_kernel_bundle_from_source(
53+
q.get_context(), syclexp::source_language::sycl_jit, source);
54+
55+
// Compile the kernel. Select kernel from the mykernels namespace
56+
std::string iota_name{"mykernels::iota"};
57+
sycl::kernel_bundle<sycl::bundle_state::executable> kb_exe = syclexp::build(
58+
kb_src, syclexp::properties{syclexp::registered_names{iota_name}});
59+
60+
sycl::kernel iota = kb_exe.ext_oneapi_get_kernel(iota_name);
61+
62+
int *ptr = sycl::malloc_shared<int>(NUM, q);
63+
q.submit([&](sycl::handler &cgh) {
64+
// Set the values of the kernel arguments.
65+
cgh.set_args(3, ptr);
66+
67+
// Launch the kernel according to its type, in this case an nd-range
68+
// kernel.
69+
sycl::nd_range ndr{{NUM}, {WGSIZE}};
70+
cgh.parallel_for(ndr, iota);
71+
}).wait();
72+
73+
for (int i = 0; i < NUM; i++) {
74+
if (ptr[i] != i + 3) {
75+
std::cout << "Result: " << ptr[i] << " expected " << i << "\n";
76+
sycl::free(ptr, q);
77+
exit(1);
78+
}
79+
}
80+
sycl::free(ptr, q);
81+
}

‎sycl/test-e2e/KernelCompiler/kernel_compiler_overload.cpp

Lines changed: 83 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,83 @@
1+
//==- kernel_compiler_overload.cpp --- kernel_compiler extension tests -----==//
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+
// REQUIRES: (opencl || level_zero)
10+
// REQUIRES: aspect-usm_device_allocations
11+
12+
// UNSUPPORTED: accelerator
13+
// UNSUPPORTED-INTENDED: while accelerator is AoT only, this cannot run there.
14+
15+
// RUN: %{build} -o %t.out
16+
// RUN: %{l0_leak_check} %{run} %t.out
17+
18+
#include <sycl/detail/core.hpp>
19+
#include <sycl/kernel_bundle.hpp>
20+
#include <sycl/usm.hpp>
21+
namespace syclexp = sycl::ext::oneapi::experimental;
22+
23+
static constexpr size_t NUM = 1024;
24+
static constexpr size_t WGSIZE = 16;
25+
26+
int main() {
27+
sycl::queue q;
28+
29+
// The source code for two kernels defined as overloaded functions.
30+
std::string source = R"""(
31+
#include <sycl/sycl.hpp>
32+
namespace syclext = sycl::ext::oneapi;
33+
namespace syclexp = sycl::ext::oneapi::experimental;
34+
35+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
36+
void iota(float start, float *ptr) {
37+
size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
38+
ptr[id] = start + static_cast<float>(id);
39+
}
40+
41+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
42+
void iota(int start, int *ptr) {
43+
size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
44+
ptr[id] = start + static_cast<int>(id);
45+
}
46+
)""";
47+
48+
// Create a kernel bundle in "source" state.
49+
sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source> kb_src =
50+
syclexp::create_kernel_bundle_from_source(
51+
q.get_context(), syclexp::source_language::sycl_jit, source);
52+
53+
// Compile the kernel. Because there are two overloads of "iota", we need to
54+
// use a C++ cast to disambiguate between them. Here, we are selecting the
55+
// "int" overload.
56+
std::string iota_name{"(void(*)(int, int*))iota"};
57+
sycl::kernel_bundle<sycl::bundle_state::executable> kb_exe = syclexp::build(
58+
kb_src, syclexp::properties{syclexp::registered_names{iota_name}});
59+
60+
// Get the kernel by passing the same string we used to construct the
61+
// "registered_names" property.
62+
sycl::kernel iota = kb_exe.ext_oneapi_get_kernel(iota_name);
63+
64+
int *ptr = sycl::malloc_shared<int>(NUM, q);
65+
q.submit([&](sycl::handler &cgh) {
66+
// Set the values of the kernel arguments.
67+
cgh.set_args(3, ptr);
68+
69+
// Launch the kernel according to its type, in this case an nd-range
70+
// kernel.
71+
sycl::nd_range ndr{{NUM}, {WGSIZE}};
72+
cgh.parallel_for(ndr, iota);
73+
}).wait();
74+
75+
for (int i = 0; i < NUM; i++) {
76+
if (ptr[i] != i + 3) {
77+
std::cout << "Result: " << ptr[i] << " expected " << i << "\n";
78+
sycl::free(ptr, q);
79+
exit(1);
80+
}
81+
}
82+
sycl::free(ptr, q);
83+
}

‎sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -6,7 +6,7 @@
66
// CHECK-DAG: README.md
77
// CHECK-DAG: lit.cfg.py
88
//
9-
// CHECK-NUM-MATCHES: 7
9+
// CHECK-NUM-MATCHES: 10
1010
//
1111
// This test verifies that `<sycl/sycl.hpp>` isn't used in E2E tests. Instead,
1212
// fine-grained includes should used, see

0 commit comments

Comments
 (0)