@@ -82,6 +82,18 @@ void vector_add_esimd(float *A, float *B, float *C) {
8282 }
8383)===" ;
8484
85+ auto constexpr DeviceCodeSplitSource = R"===(
86+ #include <sycl/sycl.hpp>
87+
88+ template<typename T, unsigned WG = 16> SYCL_EXTERNAL
89+ SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(sycl::ext::oneapi::experimental::nd_range_kernel<1>)
90+ [[sycl::reqd_work_group_size(WG)]]
91+ void vec_add(T* in1, T* in2, T* out){
92+ size_t id = sycl::ext::oneapi::this_work_item::get_nd_item<1>().get_global_linear_id();
93+ out[id] = in1[id] + in2[id];
94+ }
95+ )===" ;
96+
8597auto constexpr BadSource = R"===(
8698#include <sycl/sycl.hpp>
8799
@@ -206,12 +218,7 @@ int test_build_and_run() {
206218 ctx, syclex::source_language::sycl_jit, SYCLSource,
207219 syclex::properties{incFiles2});
208220
209- exe_kb kbExe3 = syclex::build (
210- kbSrc2, syclex::properties{
211- syclex::build_options{" -fsycl-device-code-split=per_kernel" },
212- syclex::registered_kernel_names{" ff_templated<int>" }});
213- assert (std::distance (kbExe3.begin (), kbExe3.end ()) == 2 &&
214- " Expected 2 device images" );
221+ exe_kb kbExe3 = syclex::build (kbSrc2);
215222 sycl::kernel k3 = kbExe3.ext_oneapi_get_kernel (" ff_cp" );
216223 test_1 (q, k3, 37 + 7 );
217224
@@ -222,6 +229,58 @@ int test_build_and_run() {
222229 return 0 ;
223230}
224231
232+ int test_device_code_split () {
233+ namespace syclex = sycl::ext::oneapi::experimental;
234+ using source_kb = sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source>;
235+ using exe_kb = sycl::kernel_bundle<sycl::bundle_state::executable>;
236+
237+ sycl::queue q;
238+ sycl::context ctx = q.get_context ();
239+
240+ bool ok =
241+ q.get_device ().ext_oneapi_can_compile (syclex::source_language::sycl_jit);
242+ if (!ok) {
243+ std::cout << " Apparently this device does not support `sycl_jit` source "
244+ " kernel bundle extension: "
245+ << q.get_device ().get_info <sycl::info::device::name>()
246+ << std::endl;
247+ return -1 ;
248+ }
249+
250+ source_kb kbSrc = syclex::create_kernel_bundle_from_source (
251+ ctx, syclex::source_language::sycl_jit, DeviceCodeSplitSource);
252+
253+ // Test explicit device code split
254+ std::vector<std::string> names{" vec_add<float>" , " vec_add<int>" ,
255+ " vec_add<short>" };
256+ auto build = [&](const std::string &mode) -> size_t {
257+ exe_kb kbExe = syclex::build (
258+ kbSrc, syclex::properties{
259+ syclex::registered_kernel_names{names},
260+ syclex::build_options{" -fsycl-device-code-split=" + mode}});
261+ return std::distance (kbExe.begin (), kbExe.end ());
262+ };
263+
264+ size_t perKernelNImg = build (" per_kernel" );
265+ size_t perSourceNImg = build (" per_source" );
266+ size_t offNImg = build (" off" );
267+ size_t autoNImg = build (" auto" );
268+
269+ assert (perKernelNImg == 3 );
270+ assert (perSourceNImg == 1 );
271+ assert (offNImg == 1 );
272+ assert (autoNImg >= offNImg && autoNImg <= perKernelNImg);
273+
274+ // Test implicit device code split
275+ names = {" vec_add<float, 8>" , " vec_add<float, 16>" };
276+ exe_kb kbDiffWorkGroupSizes = syclex::build (
277+ kbSrc, syclex::properties{syclex::registered_kernel_names{names}});
278+ assert (std::distance (kbDiffWorkGroupSizes.begin (),
279+ kbDiffWorkGroupSizes.end ()) == 2 );
280+
281+ return 0 ;
282+ }
283+
225284int test_esimd () {
226285 namespace syclex = sycl::ext::oneapi::experimental;
227286 using source_kb = sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source>;
@@ -393,8 +452,8 @@ int test_warning() {
393452int main (int argc, char **) {
394453#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER
395454 int optional_tests = (argc > 1 ) ? test_warning () : 0 ;
396- return test_build_and_run () || test_esimd () || test_unsupported_options () ||
397- test_error () || optional_tests;
455+ return test_build_and_run () || test_device_code_split () || test_esimd () ||
456+ test_unsupported_options () || test_error () || optional_tests;
398457#else
399458 static_assert (false , " Kernel Compiler feature test macro undefined" );
400459#endif
0 commit comments