Results 1 to 8 of 8

Thread: __constant args cause Segfault in the OpenCL library

  1. #1
    Junior Member
    Join Date
    Jul 2009
    Posts
    20

    __constant args cause Segfault in the OpenCL library

    I have kernel which is defined as

    Code :
    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

    Code :
    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.

  2. #2
    Senior Member
    Join Date
    Sep 2002
    Location
    Santa Clara
    Posts
    105

    Re: __constant args cause Segfault in the OpenCL library

    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?

  3. #3
    Junior Member
    Join Date
    Jul 2009
    Posts
    20

    Re: __constant args cause Segfault in the OpenCL library

    Quote Originally Posted by affie
    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)\n",

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

  4. #4
    Senior Member
    Join Date
    Sep 2002
    Location
    Santa Clara
    Posts
    105

    Re: __constant args cause Segfault in the OpenCL library

    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.

  5. #5
    Senior Member
    Join Date
    Jul 2009
    Location
    Northern Europe
    Posts
    311

    Re: __constant args cause Segfault in the OpenCL library

    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.

  6. #6
    Junior Member
    Join Date
    Jul 2009
    Posts
    20

    Re: __constant args cause Segfault in the OpenCL library

    Quote Originally Posted by affie
    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.
    Yes, it ocurs at clEnqueuenNDRangeKernel. I think I will contact the vendor. The driver is a beta release from nvidia. Thanks for replying affie

  7. #7
    Junior Member
    Join Date
    Jul 2009
    Posts
    20

    Re: __constant args cause Segfault in the OpenCL library

    Quote Originally Posted by dbs2
    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.
    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

    Quote Originally Posted by dbs2
    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.
    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.

  8. #8
    Junior Member
    Join Date
    Jul 2009
    Posts
    20

    Re: __constant args cause Segfault in the OpenCL library

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

Similar Threads

  1. SegFault at clGetDeviceIDs?
    By ds789 in forum OpenCL
    Replies: 1
    Last Post: 12-28-2012, 03:48 PM
  2. Wrong cl_mem pointer causes segfault
    By SAdam in forum OpenCL
    Replies: 1
    Last Post: 07-12-2012, 06:01 AM

Posting Permissions

  • You may not post new threads
  • You may not post replies
  • You may not post attachments
  • You may not edit your posts
  •