breakpoints in cl files #included by a another cl file are not hit

[Volker_Bruns]

Hi,

my main cl file does not have any kernels, but instead #includes other cl files. This works just fine, except that the debugger will not hit the breakpoints in the #included file. It seems they will only be hit if they are in the top-level cl file, i.e. the file that I provided to clBuildWithSource and as "-s" value to clBuildProgram's build options.

Stepping into a function that resides in a #included file works fine, but in my case the top-level cl file does not have a kernel to begin with. Can you think of a workaround?

cheers

[Yuri_K_Intel]

Hi, Yes, I was able to reproduce such behavior. Now I tend to think that this is design limitation, but I will try to analyze it further and I will submit either a bug or feature request. Regarding the workaround. I think you already mentioned it - since the 'step into' works fine, create a kernel function at the top level file and call the one from included file. Thanks, Yuri

[Volker_Bruns]

Hi,

thanks for looking into this.
The reason I stumbled accross this limitation is that I ported CUDA code to OpenCL.
In contrast to OpenCL, CUDA supports template functions.
I am guessing I am not the only one facing this task, so let me make my case:

There is design pattern on how to emulate templates in C, as described here:
http://stackoverflow.com/questions/4465309/how-to-use-c-templates-in-opencl-kernels

For a single templated kernel function, you end up with three files:

somekrnl.h:
//DST and SRC are template parameters, i.e. DST = float, SRC = int
#define someKrnl(DST, SRC) someKrnl_##DST##_##SRC  // becomes someKrnl_float_int

somekrnl_krnl.cl:
#include "somekrnl.h"
//invoke macro. Compiling somekrnl_krnl.cl by itself only works if DST and SRC are defined with -D, e.g. "-D DST=float -D SRC=int"
__kernel void someKrnl(DST, SRC)(__global DST* dst,  __global DST* src){
    //[...]
}

somekrnl.cl:
// instantiate template with DST=float and SRC=int
#define DST float
#define SRC int
#include "somekrnl_krnl.cl" // produces kernel someKrnl_float_int
#undef DST
#undef SRC

// instantiate template with DST=float2 and SRC=int2
#define DST float2
#define SRC int2
#include "somekrnl_krnl.cl" // produces kernel someKrnl_float2_int2
#undef DST
#undef SRC

The file that I would compile with the offline-compiler or pass to clCreateProgram is somekrnl.cl. Now it would be great if I could also add "-g -s somekrnl.cl" to the build options and still be able to put my breakpoints into the template kernel.
However, since that's not possible I have to explicitely compile the *_krnl.cl file for each template instantiation and define all template parameters with -D. It's doable, but a lot more complicated.

Thanks