Compile and launch simple OpenCL C kernels with Clover


Another Clover milestone was hit today ! During my holidays (from July 4 to July 15), I worked on the OpenCL C compiler based on Clang that is used in Clover to compile OpenCL programs. I also implemented all the infrastructure to get and launch the kernels of a program.

All is in place and can be tested (but there are still some limitations). Clover is now able to fully execute functions like clSetKernelArg, clBuildProgram and the best one : clEnqueueNDRangeKernel.

This means that from now, my OpenCL implementation is very close from being complete API-wise (I’ve still a lot of work to do regarding the OpenCL C language and its standard set of functions).

For information, Clover is able to launch this program (found in tests/test_kernel.cpp) :

float simple_function(float a) {
    return a * 2.5f;

__kernel void kernel1(__global float *a, __global float *b, float f) {
    size_t i = get_global_id(0);

    a[i] = simple_function(f) * b[i];

__kernel void kernel2(__global int *buf) {
    size_t i = get_global_id(0);

    buf[i] = 2 * i;

It recognizes that kernel1 and kernel2 are two different kernels (and that simple_function cannot be called as a kernel). The testcase launches kernel2 as it has only one argument (it’s faster to test a simpler kernel, but I’ll also add testcases for the other kernel), and checks that its result is ok.

int buffer[64];

// Create a cl_mem buffer to give to the kernel
buf = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR,
                     sizeof(buffer), buffer, &result);

// Fetch the list of the kernels in the program
result = clCreateKernelsInProgram(program, 2, kernels, 0);

// Set the argument of kernel2
result = clSetKernelArg(kernels[1], 0, sizeof(cl_mem), &buf);

// Execute the kernel
size_t global_size = sizeof(buffer) / sizeof(buffer[0]);
result = clEnqueueNDRangeKernel(queue, kernels[1], 1, 0, &global_size, 0, 0, 0, &event);
result = clWaitForEvents(1, &event);

The buffer contains after the kernel is run 64 ints : (0, 2, 4, 6, 8, …, 126).

The implementation of all of this is a bit tricky. I invite you to take a look at the CGit, but the commits are fairly difficult to understand. I hope the code isn’t too ugly, it works, I even find it somewhat nice at some places, but it’s too early to be sure it’s good.

To end this blog post, a summary of what I did during my holidays :

55 files changed, 5031 insertions(+), 1941 deletions(-)

The line numbers are not truth, but it show I was able to work many hours a day during my holidays (and I also very enjoyed them). I’m happy 🙂 . I promised to be able to work during my holidays, I kept my promise. I have now a month and a half to work on the builtin functions (only get_global_id is implemented, and there are plenty of other functions) and to finish the API (samplers and images aren’t yet completely implemented).

Then I will have to write tests, tests and more tests.


5 responses to “Compile and launch simple OpenCL C kernels with Clover

  • fabiand

    This is great news 🙂

  • Kai

    Hi Denis,

    when trying to compile the current version of Clover using a an up to date llvm and clang, I get the following error:

    /home/kai/clover/src/core/kernel.cpp: In member function ‘cl_int Coal::Kernel::addFunction(Coal::DeviceInterface*, llvm::Function*, llvm::Module*)’:
    /home/kai/clover/src/core/kernel.cpp:110:40: error: ‘class llvm::Module’ has no member named ‘getTypeName’
    make[2]: *** [src/CMakeFiles/OpenCL.dir/core/kernel.cpp.o] Fehler 1
    make[1]: *** [src/CMakeFiles/OpenCL.dir/all] Fehler 2
    make: *** [all] Fehler 2

    kai@linux:~/clover/build> llvm-config –version

    kai@linux:~/clover/build> clang –version
    clang version 3.0 (trunk 135672)
    Target: x86_64-unknown-linux-gnu
    Thread model: posix

    Any idea what is going on here?


  • steckdenis


    When I came back from my holidays, I updated LLVM and Clover refused to build with an error like yours. I have committed in Git a fix allowing Clover to build on recent LLVM-3.0 versions (LLVM-dev as of July).

  • Kai

    Hi Denis,

    thanks for the fix – I can compile clover now.


  • Won Jeon

    Hello Denis,

    I’ve tested Clover with hellocl.c from, and it failed with the following segmentation fault:

    Program received signal SIGSEGV, Segmentation fault.
    0x00f32480 in Coal::Kernel::Arg::file() const ()
    from /home/won.jeon/work/clover-master/build/lib/

    Could you let me know what in the source code causes the fault?


Leave a Reply

Fill in your details below or click an icon to log in: Logo

You are commenting using your account. Log Out /  Change )

Google+ photo

You are commenting using your Google+ account. Log Out /  Change )

Twitter picture

You are commenting using your Twitter account. Log Out /  Change )

Facebook photo

You are commenting using your Facebook account. Log Out /  Change )


Connecting to %s

%d bloggers like this: