Compile and dynamically link generated c-code in SYCL

Question

I have "*.c" files generated during runtime with function implementation int foo(int, int):

extern "C"{
int foo(int a, int b)
{
return a + b;
}
}

I want to use these implementations inside SYCL-kernels:

q.submit([&](sycl::handler& h){
auto accessor_read = ...
auto accessor_write = ...
h.parallel_for(N_tasks, [=](sycl::id id){
accessor_write[id] = foo(accessor_read[2*id], accessor_read[2*id+1]);
});
});

SYCL previously had a method build_with_source in the sycl::program class, which is deprecated, but which allowed strings to be compiled to kernels at runtime.

How would I proceed to do the same with the current standard?

I have tried to find alternative method definitions in the updated SYCL standard, but was not able to find anything.

I do not have control of the code-generator, which makes it harder to adapt the code to fit with OpenCL. (Kernel-bundles can accept OpenCL kernels)

shs_sf · Accepted Answer · 2023-05-15 15:30:38Z

0

I'm also interesting in this question. I'm not sure but I think sycl::make_kernel_bundle() should be used.

Also, I see sycl::compile() and sycl::build() methods exists but I see no examples how to use them for LevelZero backend.

Perhaps this link could shed some lingh on your question. https://www.intel.com/content/www/us/en/developer/articles/technical/sycl-interoperability-study-opencl-kernel-in-dpc.html

answered May 15, 2023 at 15:30

shs_sf

1992 silver badges11 bronze badges

Sign up to request clarification or add additional context in comments.

Comments

Anon232 · Accepted Answer · 2023-05-16 12:47:51Z

According to Intel, dynamic linkage of device code is a work in progress for DPCPP. No immediate solution was found on the Khronos forum either.

The link provided by shs_sf contains an example where a const char* OpenCL kernel source is specified and compiled to run under SYCL. Interoperating this way requires linking to OpenCL, which makes it difficult to run the kernels on Nvidia devices.

#include<CL/sycl.hpp>
#include<iostream>

int main() {
constexpr size_t size = 16;
std::array<int, size> data;

for (int i = 0; i < size; i++) { data[i] = i; }

sycl::device dev(sycl::cpu_selector{});
sycl::context ctx=sycl::context(dev);

auto ocl_dev=sycl::get_native<cl::sycl::backend::opencl,sycl::device>(dev);
auto ocl_ctx=sycl::get_native<cl::sycl::backend::opencl,sycl::context>(ctx);

cl_int err = CL_SUCCESS;
cl_command_queue ocl_queue = clCreateCommandQueueWithProperties(ocl_ctx, ocl_dev,0,&err);
sycl::queue q=sycl::make_queue<sycl::backend::opencl>(ocl_queue,ctx); 

cl_mem ocl_buf = clCreateBuffer(ocl_ctx,CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, size * sizeof(int), &data[0],&err);
sycl::buffer<int, 1> buffer =sycl::make_buffer<sycl::backend::opencl, int>(ocl_buf, ctx);

const char* kernelSource =
R"CLC(
kernel void add(global int* data) {
int index = get_global_id(0);
data[index] = data[index] + 1;
}
)CLC"; 

cl_program ocl_program = clCreateProgramWithSource(ocl_ctx,1,&kernelSource, nullptr, &err);
clBuildProgram(ocl_program, 1, &ocl_dev, nullptr, nullptr, nullptr);

cl_kernel ocl_kernel = clCreateKernel(ocl_program, "add", nullptr);
sycl::kernel add_kernel = sycl::make_kernel<sycl::backend::opencl>(ocl_kernel, ctx);

q.submit([&](sycl::handler& h){
                auto data_acc =buffer.get_access<sycl::access_mode::read_write, sycl::target::device>(h);
                h.set_args(data_acc);
                h.parallel_for(size,add_kernel);
                }).wait();

clEnqueueReadBuffer(ocl_queue, ocl_buf, CL_TRUE, 0, size*sizeof(int), &data[0], 0, NULL, NULL); 

for (int i = 0; i < size; i++) {
    if (data[i] != i + 1) { std::cout << "Results did not validate at index " << i << "!\n"; return -1; }
}
std::cout << "Success!\n";
return 0;
}

Since SYCL has been able to circumvent these issues by transpiling to nvptx, it would be nice to find a way within its standard to compile the kernels.

Collectives™ on Stack Overflow

Compile and dynamically link generated c-code in SYCL

2 Answers 2

Comments

Comments

Your Answer

Hot Network Questions

Collectives™ on Stack Overflow

2 Answers 2

Comments

Comments

Your Answer

Sign up or log in

Post as a guest

Related