From efa4d7cc4565389935f26b23a397e178a175fcc6 Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Sun, 19 Jan 2025 23:09:34 +0000 Subject: [PATCH 1/2] [SYCL][RTC] Add E2E test for implicit device code split Signed-off-by: Julian Oppermann --- .../kernel_compiler_sycl_jit.cpp | 75 +++++++++++++++++-- 1 file changed, 67 insertions(+), 8 deletions(-) diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp index d0240bc9b8964..f0cc68dfea37b 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp @@ -79,6 +79,18 @@ void vector_add_esimd(float *A, float *B, float *C) { } )==="; +auto constexpr DeviceCodeSplitSource = R"===( +#include + +template SYCL_EXTERNAL +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(sycl::ext::oneapi::experimental::nd_range_kernel<1>) +[[sycl::reqd_sub_group_size(SG)]] +void vec_add(T* in1, T* in2, T* out){ + size_t id = sycl::ext::oneapi::this_work_item::get_nd_item<1>().get_global_linear_id(); + out[id] = in1[id] + in2[id]; +} +)==="; + auto constexpr BadSource = R"===( #include @@ -203,12 +215,7 @@ int test_build_and_run() { ctx, syclex::source_language::sycl_jit, SYCLSource, syclex::properties{incFiles2}); - exe_kb kbExe3 = syclex::build( - kbSrc2, syclex::properties{ - syclex::build_options{"-fsycl-device-code-split=per_kernel"}, - syclex::registered_kernel_names{"ff_templated"}}); - assert(std::distance(kbExe3.begin(), kbExe3.end()) == 2 && - "Expected 2 device images"); + exe_kb kbExe3 = syclex::build(kbSrc2); sycl::kernel k3 = kbExe3.ext_oneapi_get_kernel("ff_cp"); test_1(q, k3, 37 + 7); @@ -219,6 +226,58 @@ int test_build_and_run() { return 0; } +int test_device_code_split() { + namespace syclex = sycl::ext::oneapi::experimental; + using source_kb = sycl::kernel_bundle; + using exe_kb = sycl::kernel_bundle; + + sycl::queue q; + sycl::context ctx = q.get_context(); + + bool ok = + q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl_jit); + if (!ok) { + std::cout << "Apparently this device does not support `sycl_jit` source " + "kernel bundle extension: " + << q.get_device().get_info() + << std::endl; + return -1; + } + + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::sycl_jit, DeviceCodeSplitSource); + + // Test explicit device code split + std::vector names{"vec_add", "vec_add", + "vec_add"}; + auto build = [&](const std::string &mode) -> size_t { + exe_kb kbExe = syclex::build( + kbSrc, syclex::properties{ + syclex::registered_kernel_names{names}, + syclex::build_options{"-fsycl-device-code-split=" + mode}}); + return std::distance(kbExe.begin(), kbExe.end()); + }; + + size_t perKernelNImg = build("per_kernel"); + size_t perSourceNImg = build("per_source"); + size_t offNImg = build("off"); + size_t autoNImg = build("auto"); + + assert(perKernelNImg == 3); + assert(perSourceNImg == 1); + assert(offNImg == 1); + assert(autoNImg >= offNImg && autoNImg <= perKernelNImg); + + // Test implicit device code split + names = {"vec_add", "vec_add"}; + exe_kb kbDiffSubgroupSizes = syclex::build( + kbSrc, syclex::properties{syclex::registered_kernel_names{names}}); + assert(std::distance(kbDiffSubgroupSizes.begin(), + kbDiffSubgroupSizes.end()) == 2); + + return 0; +} + int test_esimd() { namespace syclex = sycl::ext::oneapi::experimental; using source_kb = sycl::kernel_bundle; @@ -390,8 +449,8 @@ int test_warning() { int main(int argc, char **) { #ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER int optional_tests = (argc > 1) ? test_warning() : 0; - return test_build_and_run() || test_esimd() || test_unsupported_options() || - test_error() || optional_tests; + return test_build_and_run() || test_device_code_split() || test_esimd() || + test_unsupported_options() || test_error() || optional_tests; #else static_assert(false, "Kernel Compiler feature test macro undefined"); #endif From 2559ebafe16dfd4d6c845bb90d0d5308a74ebd6e Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Mon, 20 Jan 2025 08:50:57 +0000 Subject: [PATCH 2/2] Use reqd_work_group_size. Signed-off-by: Julian Oppermann --- .../KernelCompiler/kernel_compiler_sycl_jit.cpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp index f0cc68dfea37b..e7a6118dd3c46 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp @@ -82,9 +82,9 @@ void vector_add_esimd(float *A, float *B, float *C) { auto constexpr DeviceCodeSplitSource = R"===( #include -template SYCL_EXTERNAL +template SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(sycl::ext::oneapi::experimental::nd_range_kernel<1>) -[[sycl::reqd_sub_group_size(SG)]] +[[sycl::reqd_work_group_size(WG)]] void vec_add(T* in1, T* in2, T* out){ size_t id = sycl::ext::oneapi::this_work_item::get_nd_item<1>().get_global_linear_id(); out[id] = in1[id] + in2[id]; @@ -270,10 +270,10 @@ int test_device_code_split() { // Test implicit device code split names = {"vec_add", "vec_add"}; - exe_kb kbDiffSubgroupSizes = syclex::build( + exe_kb kbDiffWorkGroupSizes = syclex::build( kbSrc, syclex::properties{syclex::registered_kernel_names{names}}); - assert(std::distance(kbDiffSubgroupSizes.begin(), - kbDiffSubgroupSizes.end()) == 2); + assert(std::distance(kbDiffWorkGroupSizes.begin(), + kbDiffWorkGroupSizes.end()) == 2); return 0; }