From d3522888f727a3e38b50b9f10b385213300ebba0 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Thu, 17 Jul 2025 09:10:52 -0700 Subject: [PATCH 1/2] [SYCL][E2E] Add cross-SYCLBIN function pointer test case This commit adds an E2E test case for passing and calling a function pointer across images from different SYCLBIN binaries that have been linked together. Signed-off-by: Larsen, Steffen --- .../SYCLBIN/function_pointer_case.cpp | 113 ++++++++++++++++++ 1 file changed, 113 insertions(+) create mode 100644 sycl/test-e2e/SYCLBIN/function_pointer_case.cpp diff --git a/sycl/test-e2e/SYCLBIN/function_pointer_case.cpp b/sycl/test-e2e/SYCLBIN/function_pointer_case.cpp new file mode 100644 index 0000000000000..79566d2df5d00 --- /dev/null +++ b/sycl/test-e2e/SYCLBIN/function_pointer_case.cpp @@ -0,0 +1,113 @@ + +// REQUIRES: aspect-usm_device_allocations, aspect-usm_shared_allocations + +// RUN: %clangxx --offload-new-driver -fsyclbin=input %{sycl_target_opts} -fsycl-allow-device-image-dependencies -DSYCLBIN_INPUT %s -o %t.input.syclbin +// RUN: %clangxx --offload-new-driver -fsyclbin=object %{sycl_target_opts} -fsycl-allow-device-image-dependencies -DSYCLBIN_OBJECT -Xclang -fsycl-allow-func-ptr %s -o %t.object.syclbin +// RUN: %{build} -o %t.out +// +// RUN: %{l0_leak_check} %{run} %t.out %t.input.syclbin %t.object.syclbin +// +// TODO: Add the following options to the object case once linking is supported +// for AOT binaries: +// -fgpu-rdc -fsycl-targets=... --offload-arch=... + +#include +#include +#include +#include +#include + +namespace syclext = sycl::ext::oneapi; +namespace syclexp = sycl::ext::oneapi::experimental; + +typedef void (*FuncPtrT)(size_t *); + +struct ArgsT { + size_t *Ptr; + FuncPtrT FuncPtr; +}; + +#if defined(SYCLBIN_INPUT) + +SYCL_EXTERNAL size_t GetID(); + +SYCL_EXTERNAL void Func(size_t *Ptr) { + size_t GlobalID = GetID(); + Ptr[GlobalID] = GlobalID; +} + +extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( + (syclexp::single_task_kernel)) void GetFuncPtr(ArgsT *Args) { + Args->FuncPtr = Func; +} + +#elif defined(SYCLBIN_OBJECT) + +SYCL_EXTERNAL size_t GetID() { + return syclext::this_work_item::get_nd_item<1>().get_global_id(); +} + +extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( + (syclexp::nd_range_kernel<1>)) void Kernel(ArgsT *Args) { + (*Args->FuncPtr)(Args->Ptr); +} + +#else + +constexpr size_t N = 32; + +int main(int argc, char *argv[]) { + assert(argc == 3); + + sycl::queue Q; + + std::cout << "Load input SYCLBIN and compile it to object state." + << std::endl; + auto SYCLBINInput = syclexp::get_kernel_bundle( + Q.get_context(), std::string{argv[1]}); + auto SYCLBINInputObj = sycl::compile(SYCLBINInput); + + std::cout << "Load object SYCLBIN." << std::endl; + auto SYCLBINObj = syclexp::get_kernel_bundle( + Q.get_context(), std::string{argv[2]}); + + std::cout << "Link objects." << std::endl; + auto KBExe = sycl::link({SYCLBINInputObj, SYCLBINObj}); + + ArgsT *Args = sycl::malloc_shared(N, Q); + Args->Ptr = sycl::malloc_shared(N, Q); + + // Get function pointer through kernel. This deviates from the original. + sycl::kernel GetFuncPtrKern = KBExe.ext_oneapi_get_kernel("GetFuncPtr"); + std::cout << "Launching GetFuncPtr" << std::endl; + Q.submit([&](sycl::handler &CGH) { + CGH.set_args(Args); + CGH.single_task(GetFuncPtrKern); + }).wait(); + + // Launch kernel. + sycl::kernel Kern = KBExe.ext_oneapi_get_kernel("Kernel"); + std::cout << "Launching Kernel" << std::endl; + Q.submit([&](sycl::handler &CGH) { + CGH.set_args(Args); + CGH.parallel_for(sycl::nd_range{{N}, {N}}, Kern); + }).wait(); + + int Failed = 0; + for (size_t I = 0; I < N; ++I) { + if (Args->Ptr[I] != I) { + std::cout << Args->Ptr[I] << " != " << I << std::endl; + ++Failed; + } + } + + if (!Failed) + std::cout << "Results are a-okay!" << std::endl; + + sycl::free(Args->Ptr, Q); + sycl::free(Args, Q); + + return Failed; +} + +#endif From 8f0fada1d4e3dfc54011b1ac60e9a3f6cc279fd4 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Thu, 17 Jul 2025 21:15:55 -0700 Subject: [PATCH 2/2] Limit to target-spir Signed-off-by: Larsen, Steffen --- sycl/test-e2e/SYCLBIN/function_pointer_case.cpp | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/SYCLBIN/function_pointer_case.cpp b/sycl/test-e2e/SYCLBIN/function_pointer_case.cpp index 79566d2df5d00..26d1246c1ec5d 100644 --- a/sycl/test-e2e/SYCLBIN/function_pointer_case.cpp +++ b/sycl/test-e2e/SYCLBIN/function_pointer_case.cpp @@ -1,8 +1,14 @@ // REQUIRES: aspect-usm_device_allocations, aspect-usm_shared_allocations -// RUN: %clangxx --offload-new-driver -fsyclbin=input %{sycl_target_opts} -fsycl-allow-device-image-dependencies -DSYCLBIN_INPUT %s -o %t.input.syclbin -// RUN: %clangxx --offload-new-driver -fsyclbin=object %{sycl_target_opts} -fsycl-allow-device-image-dependencies -DSYCLBIN_OBJECT -Xclang -fsycl-allow-func-ptr %s -o %t.object.syclbin +// ptxas currently fails to compile images with unresolved symbols. Disable for +// other targets than SPIR-V until this has been resolved. (CMPLRLLVM-68810) +// Note: %{sycl_target_opts} should be added to the SYCLBIN compilation lines +// once fixed. +// REQUIRES: target-spir + +// RUN: %clangxx --offload-new-driver -fsyclbin=input -fsycl-allow-device-image-dependencies -DSYCLBIN_INPUT %s -o %t.input.syclbin +// RUN: %clangxx --offload-new-driver -fsyclbin=object -fsycl-allow-device-image-dependencies -DSYCLBIN_OBJECT -Xclang -fsycl-allow-func-ptr %s -o %t.object.syclbin // RUN: %{build} -o %t.out // // RUN: %{l0_leak_check} %{run} %t.out %t.input.syclbin %t.object.syclbin