Seg fault inside kernel with more than constant args

Accelerated Computing CUDA CUDA Programming and Performance
1

I have kernel which is defined as

Function(__constant uchar *buf, __constant ushort buflen, __constant uint m, __constant uchar *table, __constant uint *search_B2G, __global uint *mts, __global uint *o)

If I call this kernel I get a Segfault from within the libOpenCL. If I modify the kernel definition to

Function(__constant uchar *buf, __constant ushort buflen, __constant uint m, __global uchar *table, __global uint *search, __global uint *mts, __global uint *o)

where I have changed table and search to __global instead of __constant and it works fine now without any seg fault. I checked the max no of __constant kernel args that can be supplied to a kernel on my device is 9, so exceeding the limit is ruled out. Also none of my buffers cross 64K, which is the minimum for a constant buffer. Both table and search are buffers defined using CL_MEM_READ_ONLY. I am not sure why I am getting the seg fault while changing the modifier to any of these 2 args to __constant from __global.

P.S. Sorry for the topic. Never noticed the errors. A more apt title would be “Seg fault inside kernel with varying constant args”

2

I also tried putting the __constant args after the __global args, and I still seem to be getting the seg faults

"__kernel void Search(__global uint *mts, __global uint *o, __constant uchar *table, __constant uint *search_B2G, __constant uchar *buf, ushort buflen, uint m)\n",

^

Gives me a Seg fault

declaring search_B2G and table as __global gives no seg fault

"__kernel void Search(__global uint *mts, __global uint *o, __global uchar *table, __global uint *search_B2G, __constant uchar *buf, ushort buflen, uint m)\n",

Table is defined as -

ocl_table = SCOclCreateOpenCLBuffer(ctx->ocl_context,

																	 (CL_MEM_READ_ONLY |

																	  CL_MEM_COPY_HOST_PTR),

																	 sizeof(table), table);

//where table is defined as

uint8_t table[256];

So table buffer size would be 256 bytes.

3

For other’s reference this is a bug in the 190.29 driver. 195.30 seems to solve it. Thanks all