diff --git a/SYCL/Plugin/interop-opencl-make-kernel-bundle.cpp b/SYCL/Plugin/interop-opencl-make-kernel-bundle.cpp new file mode 100644 index 0000000000..399d1d8d08 --- /dev/null +++ b/SYCL/Plugin/interop-opencl-make-kernel-bundle.cpp @@ -0,0 +1,62 @@ +// REQUIRES: opencl, opencl_icd + +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out %opencl_lib +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out + +#include + +#include +#include +#include +#include +#include + +char const *source = R"OCL(kernel void do_nothing() {})OCL"; + +using namespace sycl; + +int select_opencl(const sycl::device &dev) { + if (dev.get_platform().get_backend() == backend::opencl) { + return 1; + } + return -1; +} + +int main() { + auto q = queue{select_opencl}; + + std::cerr << "Get native context." << std::endl; + auto native_context = get_native(q.get_context()); + std::cerr << "Create native program." << std::endl; + cl_program p = + clCreateProgramWithSource(native_context, 1, &source, nullptr, nullptr); + std::cerr << "Build native program." << std::endl; + clBuildProgram(p, 0, nullptr, nullptr, nullptr, nullptr); + std::cerr << "Release native context." << std::endl; + clReleaseContext(native_context); + + std::cerr << "Make kernel bundle." << std::endl; + auto bundle = make_kernel_bundle( + p, q.get_context()); + std::cerr << "Release native program." << std::endl; + // cl_program must have been retained by the above call. + clReleaseProgram(p); + + std::cerr << "Get native program." << std::endl; + std::vector device_image = + get_native(bundle); + assert(device_image.size() == 1); + std::cerr << "Create native kernel." << std::endl; + cl_kernel k = clCreateKernel(device_image.front(), "do_nothing", nullptr); + // get_native must have retained cl_program as well. + clReleaseProgram(device_image.front()); + + std::cerr << "Make kernel." << std::endl; + make_kernel(k, q.get_context()); + std::cerr << "Release native kernel." << std::endl; + clReleaseKernel(k); + + return 0; +}