CL_INVALID_BINARY from clCreateProgramWithSource/clBuildProgram

I have a very simple kernel that I’m trying to compile (using my own modified version of clcc). This is the kernel:

__kernel void init(__global float * v, int n)

{

    int i = get_local_id(0)

    + get_local_size(0) * get_group_id(0)

    + get_local_size(0) * get_global_size(0) * get_local_size(1) * get_group_id(1)

    + get_local_size(0) * get_global_size(0) * get_local_id(1)

    ;

if (i > n)

        return;

if (i < 0)

        return;

float w = 1.0 / n;

    v[i] = i * w;

}

The first statement in the kernel computes the mapping of the work item and group into an array index. This statements seems to cause some problems with the compiler.

I compiled this using calls to clCreateProgramWithSource then clBuildProgram. However, clBuildProgram returned -42 (= CL_INVALID_BINARY).

The following rewritten version is accepted.

__kernel void init(__global float * v, int n)

{

    int j = get_local_id(1);

    int i = get_local_id(0)

    + get_local_size(0) * get_group_id(0)

    + get_local_size(0) * get_global_size(0) * get_local_size(1) * get_group_id(1)

//    + get_local_size(0) * get_global_size(0) * get_local_id(1)

    + get_local_size(0) * get_global_size(0) * j

    ;

if (i > n)

        return;

if (i < 0)

        return;

float w = 1.0 / n;

    v[i] = i * w;

}

This seems like compiler bug.

In addition, the return value CL_INVALID_BINARY a little weird because according to the OpenCL Specification “CL_INVALID_BINARY if program is created with clCreateWithProgramBinary and devices listed in device_list do not have a valid program binary loaded.”. This is the only way it can return CL_INVALID_BINARY. However, I never called clCreateWithProgramBinary. This seems like the OpenCL driver is not compliant with the spec.

Ken

When I compile the kernel with -cl-opt-disable, the compilation (via call to clBuildProgram) works, and the compiled kernel executes fine. This definitely seems to be a compiler bug (OpenCL 1.0 CUDA 4.0.1, GeForce GTX 470, driver version = 275.33).

Ken

Hi Ken,

Did you manage to solve your problem? It happened to me too today.

Yes, I got it to compile and run. But, I had to disable optimization for the compiler, using -cl-opt-disable, e.g., *err = clBuildProgram(program, 1, &device, “-cl-opt-disable”, NULL, NULL);

Kind of disappointing, but NVIDIA is supposed to be releasing v1.1 of OpenCL “soon”.

Ken