Skip to content

Commit 0df930d

Browse files
PietroGhgPietroGhg
committed
Add tests for jit RTC
1 parent 94c50ec commit 0df930d

File tree

4 files changed

+246
-1
lines changed

4 files changed

+246
-1
lines changed

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

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

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

Lines changed: 82 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,82 @@
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: %{run} %t.out 1
17+
// RUN: %{l0_leak_check} %{run} %t.out 1
18+
19+
#include <sycl/detail/core.hpp>
20+
#include <sycl/kernel_bundle.hpp>
21+
#include <sycl/usm.hpp>
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 two kernels defined in different namespaces
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+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
37+
void iota(int start, int *ptr) {
38+
size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
39+
ptr[id] = start + static_cast<int>(id);
40+
}
41+
42+
namespace mykernels {
43+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
44+
void iota(int start, int *ptr) {
45+
size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
46+
ptr[id] = start + static_cast<int>(id);
47+
}
48+
} // namespace mykernels
49+
)""";
50+
51+
// Create a kernel bundle in "source" state.
52+
sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source> kb_src =
53+
syclexp::create_kernel_bundle_from_source(
54+
q.get_context(), syclexp::source_language::sycl_jit, source);
55+
56+
// Compile the kernel. Select kernel from the mykernels namespace
57+
std::string iota_name{"mykernels::iota"};
58+
sycl::kernel_bundle<sycl::bundle_state::executable> kb_exe = syclexp::build(
59+
kb_src, syclexp::properties{syclexp::registered_kernel_names{iota_name}});
60+
61+
sycl::kernel iota = kb_exe.ext_oneapi_get_kernel(iota_name);
62+
63+
int *ptr = sycl::malloc_shared<int>(NUM, q);
64+
q.submit([&](sycl::handler &cgh) {
65+
// Set the values of the kernel arguments.
66+
cgh.set_args(3, ptr);
67+
68+
// Launch the kernel according to its type, in this case an nd-range
69+
// kernel.
70+
sycl::nd_range ndr{{NUM}, {WGSIZE}};
71+
cgh.parallel_for(ndr, iota);
72+
}).wait();
73+
74+
for (int i = 0; i < NUM; i++) {
75+
if (ptr[i] != i + 3) {
76+
std::cout << "Result: " << ptr[i] << " expected " << i << "\n";
77+
sycl::free(ptr, q);
78+
exit(1);
79+
}
80+
}
81+
sycl::free(ptr, q);
82+
}

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

Lines changed: 84 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,84 @@
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: %{run} %t.out 1
17+
// RUN: %{l0_leak_check} %{run} %t.out 1
18+
19+
#include <sycl/detail/core.hpp>
20+
#include <sycl/kernel_bundle.hpp>
21+
#include <sycl/usm.hpp>
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 two kernels defined as overloaded functions.
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+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
37+
void iota(float start, float *ptr) {
38+
size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
39+
ptr[id] = start + static_cast<float>(id);
40+
}
41+
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+
)""";
48+
49+
// Create a kernel bundle in "source" state.
50+
sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source> kb_src =
51+
syclexp::create_kernel_bundle_from_source(
52+
q.get_context(), syclexp::source_language::sycl_jit, source);
53+
54+
// Compile the kernel. Because there are two overloads of "iota", we need to
55+
// use a C++ cast to disambiguate between them. Here, we are selecting the
56+
// "int" overload.
57+
std::string iota_name{"(void(*)(int, int*))iota"};
58+
sycl::kernel_bundle<sycl::bundle_state::executable> kb_exe = syclexp::build(
59+
kb_src, syclexp::properties{syclexp::registered_kernel_names{iota_name}});
60+
61+
// Get the kernel by passing the same string we used to construct the
62+
// "registered_kernel_names" property.
63+
sycl::kernel iota = kb_exe.ext_oneapi_get_kernel(iota_name);
64+
65+
int *ptr = sycl::malloc_shared<int>(NUM, q);
66+
q.submit([&](sycl::handler &cgh) {
67+
// Set the values of the kernel arguments.
68+
cgh.set_args(3, ptr);
69+
70+
// Launch the kernel according to its type, in this case an nd-range
71+
// kernel.
72+
sycl::nd_range ndr{{NUM}, {WGSIZE}};
73+
cgh.parallel_for(ndr, iota);
74+
}).wait();
75+
76+
for (int i = 0; i < NUM; i++) {
77+
if (ptr[i] != i + 3) {
78+
std::cout << "Result: " << ptr[i] << " expected " << i << "\n";
79+
sycl::free(ptr, q);
80+
exit(1);
81+
}
82+
}
83+
sycl::free(ptr, q);
84+
}

‎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)