Why can't I create a kernel (CL_INVALID_PROGRAM_EXECUTABLE) after successfully compiling an OpenCL program?

Question

In the following program, I compile a kernel for the first device on the first platform:

const char* kernel_source_code = R"(
__kernel void vectorAdd(
 __global float * __restrict C,
 __global float const * __restrict A,
 __global float const * __restrict B,
 unsigned long length)
{
 int i = get_global_id(0);
 if (i < length)
 C[i] = A[i] + B[i];
}
)";
cl_int status;
cl_platform platform_id;
status = clGetPlatformIDs(1, &platform_id, nullptr);
ensure_success(status);
cl_device device_id;
status = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, nullptr);
ensure_success(status);
cl_context context_id = clCreateContext(nullptr, 1, &device_id, nullptr, nullptr, &status);
ensure_success(status);
size_t length = strlen(kernel_source_code);
cl_program program_id = clCreateProgramWithSource(context_id, 1, &kernel_source_code, &length, &status);
ensure_success(status);
status = clCompileProgram(
 program_id,
 1, &device_id, // single device
 "", // no special options
 0, nullptr, nullptr, // num headers, header sources, header names
 nullptr, nullptr // no callback
 );
ensure_success(status);
cl_kernel kernel = clCreateKernel(program_id, "vectorAdd", &status);
ensure_success(status);

but - the last API call, the kernel creation, fails with CL_INVALID_PROGRAM_EXECUTABLE, which means "there is no successfully built executable for [the] program".

Why am I getting this failure, if my compilation has succeeded?

Answer 1

The OpenCL documentation doesn't clarify this, but - it's not enough that your compiled program has "binaries", you need it to have linked "binaries", i.e. you have to link your program before you can get any kernels. After your clCompileProgram() call, add:

cl_program linked = clLinkProgram(context_id,
 1, &device_id, // devices,
 "", // No link options
 1, &program_id,
 nullptr, nullptr, // no callback
 &status);
ensure_success(status);

Now you can create your kernel (that is, create it using the linked variable, not the program_id.

Alternatively you could use clBuildKernel().