Skip to content

[SYCL][E2E] Add cross-SYCLBIN function pointer test case #19503

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Draft
wants to merge 2 commits into
base: sycl
Choose a base branch
from
Draft
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
119 changes: 119 additions & 0 deletions sycl/test-e2e/SYCLBIN/function_pointer_case.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,119 @@

// REQUIRES: aspect-usm_device_allocations, aspect-usm_shared_allocations

// 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
//
// TODO: Add the following options to the object case once linking is supported
// for AOT binaries:
// -fgpu-rdc -fsycl-targets=... --offload-arch=...

#include <sycl/detail/core.hpp>
#include <sycl/ext/oneapi/experimental/syclbin_kernel_bundle.hpp>
#include <sycl/ext/oneapi/free_function_queries.hpp>
#include <sycl/kernel_bundle.hpp>
#include <sycl/usm.hpp>

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<sycl::bundle_state::input>(
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<sycl::bundle_state::object>(
Q.get_context(), std::string{argv[2]});

std::cout << "Link objects." << std::endl;
auto KBExe = sycl::link({SYCLBINInputObj, SYCLBINObj});

ArgsT *Args = sycl::malloc_shared<ArgsT>(N, Q);
Args->Ptr = sycl::malloc_shared<size_t>(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
Loading