OpenCL clBuildProgram caches source, and does not recompile if #include'd source changes

Question

I have implemented a project with opencl. I have a file which contains the kernel function and the functions which are used by the kernel are included in a seperate header file but when I change the file which is included, sometimes the changes are applied and sometimes they are not and it makes me confused if the application has bug or not.

I checked the other posts in stackoverflow and see nvidia has serious problem with passing -I{include directory}, so I changed it and give the header files address explicitly, but still the opencl compiler is not able to find the errors in the header file which is included in the kernel file name.

Also, I am using nvidia gtx 980 and I have intalled CUDA 7.0 on my computer.

Anyone has the same experience? how can I fix it?

So, Assume I have a kernel like this:

#include "../../src/cl/test_kernel_include.cl"

void __kernel test_kernel(
  __global int* result,
  int n
  )
{
  int thread_idx = get_global_id(0);
  result[thread_idx] = test_func();
}

which the test_kernel_include.cl is as follows:

int test_func()
{
  return 1;
}

Then I run the code and I get an array which all the members are equal to 1 as we expect. Now, I change the test_kernel_include.cl to:

int test_func()
{
  return 2;
}

but the result is still an array which all the members are equal to 1 which should change to 2 but they are not.

Answer 1

Do this before platform initialization:

setenv("CUDA_CACHE_DISABLE", "1", 1);

It will disable caching mechanism for the build. It also works for the OpenCL platform, even though it says CUDA.

Answer 2

In order to improve kernel compilation times, NVIDIA implement a caching scheme, whereby a compiled kernel binary is stored to disk and loaded next time the same kernel is compiled. Some hash is computed on the kernel source code which is then used as the index into the compiled kernel cache.

Unfortunately, these hashes do not include any header files that are included by the main kernel source. This means that when you change something in an included header file, the driver will essentially ignore the change and reload the previous kernel binary from disk (unless something changed in the main kernel source as well).

On Linux systems, the kernel cache can be found in ~/.nv/ComputeCache. If you delete this directory after making a change to one of your include files, then it should force the driver to actually recompile the OpenCL kernel.