Tag Archives: cpu

OpenCL Cookbook: Hello World using C host binding

In our OpenCL Cookbook series so far we’ve looked at some preliminary data structures in OpenCL host programming using the C language. This time – we finally arrive at a complete end-to-end example – the customary Hello World!

What this example does is simple. The host program in C passes a character array to the GPU into which the GPU writes the characters of the phrase: “Hello, World!”. The host program then reads the contents of the character array back and prints them on the screen. The output should be “Hello, World!”.

The code is annotated using brief comments. There are some aspects of OpenCL that are new that I have not yet been through in previous articles but don’t worry I’ll go through a full dissection after presenting the complete code.

Note that error handling has been taken out completely to keep the program short for easy viewing. The most important aspect of error handling in the program below is checking the build error, status and log for a failed program build which you can find further details of in my previous post.

Host source

#include
#include
#ifdef __APPLE__
#include <OpenCL/opencl.h>
#else
#include <CL/cl.h>
#endif

int main() {

    cl_platform_id platform; cl_device_id device; cl_context context;
    cl_program program; cl_kernel kernel; cl_command_queue queue;
    cl_mem kernelBuffer;

    FILE* programHandle; char *programBuffer; char *programLog;
    size_t programSize; char hostBuffer[32];

    // get first available sdk and gpu and create context
    clGetPlatformIDs(1, &platform, NULL);
    clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
    context = clCreateContext(NULL, 1, &device, NULL, NULL, NULL);

    // get size of kernel source
    programHandle = fopen("helloWorld.cl", "r");
    fseek(programHandle, 0, SEEK_END);
    programSize = ftell(programHandle);
    rewind(programHandle);

    // read kernel source into buffer
    programBuffer = (char*) malloc(programSize + 1);
    programBuffer[programSize] = '\0';
    fread(programBuffer, sizeof(char), programSize, programHandle);
    fclose(programHandle);

    // create and build program
    program = clCreateProgramWithSource(context, 1,
            (const char**) &programBuffer, &programSize, NULL);
    free(programBuffer);
    clBuildProgram(program, 1, &device, "-Werror -cl-std=CL1.1", NULL, NULL);

    // create kernel and command queue
    kernel = clCreateKernel(program, "hello", NULL);
    queue = clCreateCommandQueue(context, device, 0, NULL);

    // create kernel argument buffer and set it into kernel
    kernelBuffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY,
            32 * sizeof(char), NULL, NULL);
    clSetKernelArg(kernel, 0, sizeof(cl_mem), &kernelBuffer);

    // execute kernel, read back the output and print to screen
    clEnqueueTask(queue, kernel, 0, NULL, NULL);
    clEnqueueReadBuffer(queue, kernelBuffer, CL_TRUE, 0,
            32 * sizeof(char), hostBuffer, 0, NULL, NULL);
    puts(hostBuffer);

    clFlush(queue);
    clFinish(queue);
    clReleaseKernel(kernel);
    clReleaseProgram(program);
    clReleaseMemObject(kernelBuffer);
    clReleaseCommandQueue(queue);
    clReleaseContext(context);
    return 0;

}

The host source runs on the CPU and is written in C in this case though you could also write it in C++, Python or Java whereas the kernel source runs on a device which could be one or more CPUs, GPUs or accelerators. The host source must be written in a host language whereas the kernel source must be written in OpenCL.

Host source by dissection

Here I describe what the host source is doing by dissecting it. A hello world example should ideally be entirely self contained and not rely on other articles to complement the reader’s understanding. With the exception of error handling and particularly how to debug a failed program build which I address elsewhere this example is self contained.

Below I present one snippet of code at a time followed by its dissection.

Creating platforms, devices and contexts

// get first available sdk and gpu and create context
clGetPlatformIDs(1, &platform, NULL);
clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
context = clCreateContext(NULL, 1, &device, NULL, NULL, NULL);

Here I first get a platform (an OpenCL SDK/framework). As I know I only have the Apple OpenCL framework installed on my Mac it will always be the one selected. However, if you have multiple SDKs installed such as AMD, Nvidia and Intel then you may want to select one explicitly. Next I ask for a GPU device. Once again, my machine only has one GPU so it will always be the one that’s selected but if you have multiple GPUs installed you may want to choose one in particular. Finally I create a context which is an incredibly important OpenCL data structure as it is required for the creation of numerous other structures such as programs, command queues and kernel buffers.

Loading kernel sources

// get size of kernel source
programHandle = fopen("helloWorld.cl", "r");
fseek(programHandle, 0, SEEK_END);
programSize = ftell(programHandle);
rewind(programHandle);

// read kernel source into buffer
programBuffer = (char*) malloc(programSize + 1);
programBuffer[programSize] = '\0';
fread(programBuffer, sizeof(char), programSize, programHandle);
fclose(programHandle);

As this is a host source file it has the responsibility of involving the kernel source. Generally speaking the kernel source is usually compiled at runtime as part of the execution of the host source. Therefore, the host source file must pull in the kernel source and compile it. Above I first calculate the size of the kernel source file and then read the source in into a buffer of that calculated size.

Creating a program and compiling kernel sources

// create and build program
program = clCreateProgramWithSource(context, 1,
        (const char**) &programBuffer, &programSize, NULL);
free(programBuffer);
clBuildProgram(program, 1, &device, "-Werror -cl-std=CL1.1", NULL, NULL);

Here I construct a program structure by passing in a context and the buffer containing the kernel source. Then I build the program which essentially compiles the kernel source based on supplied build options. Note that a program can contain numerous kernel sources containing multiple OpenCL functions potentially drawn in from a number of files. This program build steps builds the sum total of all kernels sources read in. At this point the build could fail for a variety of reasons and it’s critically important to be able to narrow the cause easily. Here I’ve skipped this error handling but I address this subject in detail on my previous post in the series.

Creating kernels and command queues

// create kernel and command queue
kernel = clCreateKernel(program, "hello", NULL);
queue = clCreateCommandQueue(context, device, 0, NULL);

Here I create a kernel and a command queue structure. Let’s look at what each one means in turn.

A kernel is an OpenCL function that executes on one or more devices. The program structure above may contain numerous functions so the purpose of creating the kernel structure above is to pinpoint one particular one called ‘hello’. I need a reference to this kernel in order to pass it an argument later on in the process.

A command queue is exactly what the name implies. The host program invokes a device by sending it a command. The sending mechanism for that command is a queue. Commands are by default processed in FIFO order but that can be changed by a configuration option. Sending a command, also known as a task, to a command queue is a way of requesting its execution.

Setting kernel arguments

// create kernel argument buffer and set it into kernel
kernelBuffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY,
        32 * sizeof(char), NULL, NULL);
clSetKernelArg(kernel, 0, sizeof(cl_mem), &kernelBuffer);

Here, I create an OpenCL memory object. There are two types of memory objects – image and buffer. Here I am not dealing with image data so I choose a buffer memory object. Our goal here is to provide the kernel with a character array big enough to hold the phrase ‘Hello, World!’. However I cannot pass a character array into a kernel directly. I must create an OpenCL buffer memory object of a given size and then set it as the first kernel argument and that’s what I’m doing above. You’ll notice that I set the memory object to be write only as the device only needs to write to it.

Executing kernels and reading output data

// execute kernel, read back the output and print to screen
clEnqueueTask(queue, kernel, 0, NULL, NULL);
clEnqueueReadBuffer(queue, kernelBuffer, CL_TRUE, 0,
        32 * sizeof(char), hostBuffer, 0, NULL, NULL);
puts(hostBuffer);

This is the final step. Earlier I created a command queue and a kernel structure for the hello function and passed in a buffer memory object as the first argument. Here I complete the entire process by enqueuing the kernel for execution as a task into the command queue and reading back the output by passing in a character array of the same size as the original kernel buffer memory object. I then print the contents of that array onto the screen to prove that it contains what the GPU originally wrote into it.

Cleaning up

clFlush(queue);
clFinish(queue);
clReleaseKernel(kernel);
clReleaseProgram(program);
clReleaseMemObject(kernelBuffer);
clReleaseCommandQueue(queue);
clReleaseContext(context);

Above I first ensure that all commands have been issed to the device associated with the command queue by calling clFlush(). Then I block until all commands have been issued and completed by calling clFinish(). The rest of the functions above simple deallocate their own respective named structures.

Kernel source

__kernel void hello(__global char* message){
message[0] = 'H';
message[1] = 'e';
message[2] = 'l';
message[3] = 'l';
message[4] = 'o';
message[5] = ',';
message[6] = ' ';
message[7] = 'W';
message[8] = 'o';
message[9] = 'r';
message[10] = 'l';
message[11] = 'd';
message[12] = '!';
message[13] = '';
}

The kernel source is fairly self explanatory. It simply receives a character array called message and writes its message into it. Kernel functions get infinitely more complex than this one but this one has been kept deliberately simple.

Compile and run as follows keeping both source files in the same directory.

clang -framework OpenCL helloWorld.c -o helloWorld && ./helloWorld

As always if you have any feedback or if this helped you let me know in the comments!