__constant args cause Segfault in the OpenCL library

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.

The address qualifiers (__global, __constant, __local etc.) are used to declare address regions of memory. When used with kernel arguments, these qualifiers can only be used for arguments that are declared to be a pointer.

The __constant ushort buflen, and __constant uint m declaration is incorrect. If you want to mark that these arguments are read-only then use the const qualifier instead.

Other than this, you kernel declaration looks good. And __constant uchar *table and __constant uint *search_B2G should work. Can you provide details on where you are seeing this problem i.e. which device and what OS?

Actually I pasted the wrong definition. I had declared the non-pointers without any qualifiers since I wanted them to be private.

"__kernel void Function(__global uint *offsets, __constant uint *search, __constant uchar *table, __constant uchar *buf, ushort arg_buflen, uint m)

Changing all the constant to global results in no seg fault. And all the mem objects except “offsets” are specified as READ_ONLY.

Using __constant should not cause a segfault. When does the segfault occur? Is it on the call to clEnqueueNDRangeKernel for this kernel? I would suggest working with the vendor where this problem is occuring to determine the cause of this problem.

Make sure that the total size of your constant buffers does not exceed what the hardware supports. The size returned for CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE is the total size of all constant buffers. You should get an error from clEnqueueNDRangeKernel if you exceed this on a friendly implementation.

Also make sure you are at no point accidentally writing to a buffer which you have declared as constant. These are stored in separate physical memory areas on most GPUs so this will cause a real problem.

Yes, it ocurs at clEnqueuenNDRangeKernel. I think I will contact the vendor. The driver is a beta release from nvidia. Thanks for replying affie

I haven’t exceeded the limit. I am not crossing 1 kb of constant buffer size, let alone the min limit of 64 kb. Running it through gdb at times, makes clEnqueueNDRangeKernel throw errors. Otherwise it’s a seg fault. Changing the constant to global makes it run fine without any segvs

I have defined all the buffers as read only. I have double checked my kernel and I am not performing any writes on the constant buffers.

Also, I have also noticed really bad performance. I think poor data transfer bandwidth. I should profile it more closely with actual values to back up my claim. Want to see how it performs against CUDA which is stable and allegedly faster(as claimed by others). These bugs from the driver are a nuisance. Can’t blame the driver, since it’s still beta.

This is a bug in the 190.29 nvidia driver. The 195.30 beta solves it. Thanks all