Tag Archives: framework

OpenCL Cookbook: Building a program and debugging failures

Last time, in the opencl cookbook series, we looked at how to create a program data structure in the C OpenCL host programming API as well as how to read kernel source from a separate file. A program is a container or collection of kernels and a kernel, in turn, is a function in OpenCL that executes on an OpenCL device such as a CPU, GPU or accelerator. This time we look at how to build a program, which in turn builds the kernels within it and also how to debug failures that occur in the program build. For the latter we recreate two kinds of failures to see how the program reacts.

But what actually happens when you build a program? The clBuildProgram() function in C takes a program and a program in turn contains the source of one or more kernels read in from one or more files. However, kernel sources in raw form are of little use. To be functionally useful they must be compiled. This is what happens when you build a program.

Every OpenCL framework/SDK/implementation (whatever you want to call it) is mandated by the specification to make a compiler accessible through the clBuildProgram(). Though, they may provide other interfaces to their OpenCL compiler. AMD provides a compile time compiler command called clc whereas NVidia provides only a runtime one. clBuildProgram() compiles and link a program for devices associated with the platform.

Our example host program below is an extension of the one in the previous article in the series. It reads in a kernel function from a separate file into an OpenCL program and builds the program which would normally succeed. However, here I introduce a couple of mistakes to show you how the API reacts in each case and how to debug such issues.

Host source (buildProgramDebug.c)

#include <stdio.h>
#include <stdlib.h>
#ifdef __APPLE__
#include <OpenCL/opencl.h>
#include <CL/cl.h>

int main() {

    cl_platform_id platform; cl_device_id device; cl_context context;
    cl_program program; cl_int error; cl_build_status status;

    FILE* programHandle;
    char *programBuffer; char *programLog;
    size_t programSize; size_t logSize;

    // get first available platform 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("kernel.cl", "r");
    fseek(programHandle, 0, SEEK_END);
    programSize = ftell(programHandle);

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

    // create program from buffer
    program = clCreateProgramWithSource(context, 1,
            (const char**) &programBuffer, &programSize, NULL);

    // build program
    const char options[] = "-Werror -cl-std=CL1.1";
    error = clBuildProgram(program, 1, &device, options, NULL, NULL);

    // build failed
    if (error != CL_SUCCESS) {

        // check build error and build status first
        clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_STATUS, 
                sizeof(cl_build_status), &status, NULL);

        // check build log
        clGetProgramBuildInfo(program, device, 
                CL_PROGRAM_BUILD_LOG, 0, NULL, &logSize);
        programLog = (char*) calloc (logSize+1, sizeof(char));
        clGetProgramBuildInfo(program, device, 
                CL_PROGRAM_BUILD_LOG, logSize+1, programLog, NULL);
        printf("Build failed; error=%d, status=%d, programLog:nn%s", 
                error, status, programLog);


    return 0;


Kernel source (kernel.cl)

__kernel void hello(__global char* string){

string[0] = 'H';
string[1] = 'e';
string[2] = 'l';
string[3] = 'l';
string[4] = 'o';
string[5] = ',';
string[6] = ' ';
string[7] = 'W';
string[8] = 'o';
string[9] = 'r';
string[10] = 'l';
string[11] = 'd';
string[12] = '!';
string[13] = '';


These two programs, as they stand, have no errors and in that state produce no output. The kernel source is read in to a program and the program built based on the command line options provided on line 41. The error variable on line 45 always equals CL_SUCCESS so nothing is printed. Let’s now introduce two errors in our source code one at a time and see what happens. Specifically, when a problem occurs, we’ll be examining three separate variables – the error code (which is the return value from clBuildProgram), the program build status (which has to be specifically requested) and the program build log (which also has to be requested).

Error 1: Rogue build option

Here I change line 41 to contain a bogus command line option called ‘-foobar’ as below.

const char options[] = "-Werror -cl-std=CL1.1 -foobar";

When the program is built, unsurprisingly it fails, with the output below.

tron:opencl dhruba$ clang -framework OpenCL buildProgramDebug.c -o buildProgramDebug && ./buildProgramDebug
Build failed; error=-43, status=-1, programLog:

Above the error code of -43 corresponds to the constant CL_INVALID_BUILD_OPTIONS and a status of -1 corresponds to the constant CL_BUILD_NONE. The former is self explanatory whereas the latter means that the kernel was not compiled which is to be expected as the build options were wrong.

Error 2: Kernel source syntax error

Here, I revert the error I introduced last time, and instead add an extra underscore as the first character of the kernel source to create an OpenCL syntax error. This time the output is longer as the program log is no longer blank.

tron:opencl dhruba$ clang -framework OpenCL buildProgramDebug.c -o buildProgramDebug && ./buildProgramDebug
Build failed; error=-11, status=-2, programLog:

:1:1: error: unknown type name '___kernel'
___kernel void hello(__global char* string){
:1:11: error: expected identifier or '('
___kernel void hello(__global char* string){

Above the error code of -11 refers to the constant CL_BUILD_PROGRAM_FAILURE and a program build status of -2 refers to the constant CL_BUILD_ERROR which makes sense. This time, however, we have some output in the program log field. The syntax error in the kernel source is being reported by the runtime compiler to the program log.

Error code and build status constants

One final tip: you may be asking how I knew which error codes and build statuses corresponded to which constants in OpenCL. Well – this is a bit of a nightmare to be quite honest. I had to open up the following header file in the Apple OpenCL framework to check which constants matched those integers. I’m sure there will be similar places to look in other SDKs.


As you can see, the above three fields, provide critically important means of debugging the failure to build of your program and can point our errors both in your host source and your kernel source. Did this help you, did you face any issues or do you have feedback for improvement? Let me know in the comments!