Possible driver bug on passing NULL pointers

Accelerated Computing CUDA CUDA Programming and Performance
1

Hey,

I am stumbling across some weird SIGSEGV occuring within libcuda.so, when passing NULL memory buffers to a kernel. Here is the exemplary kernel:

__kernel void test(

   __global unsigned int * data, 

   __global unsigned int * flag) {

   int gid = get_global_id(0);

   if (flag)

      data[gid] = 4;

   else

      data[gid] = 2;

}

Now, if I run the following code:

// allocate buffer object

  // allocate kernel object

  clSetKernelArg(kernel, 0, sizeof(cl_mem), &buffer);

  clSetKernelArg(kernel, 1, sizeof(cl_mem), &buffer);

  // schedule kernel

  // wait on queue

}

everything works fine. However, when I try to pass a NULL cl_mem object:

// allocate buffer object

  // allocate kernel object

  clSetKernelArg(kernel, 0, sizeof(cl_mem), &buffer);

  cl_mem null_buffer = NULL;

  clSetKernelArg(kernel, 1, sizeof(cl_mem), &null_buffer);

  // schedule kernel

  // wait on queue

}

I am getting a SIGSEGV from libcuda.so on the call to clFinish(…).

I am using the latest 280.13 driver on a 64bit Scientific Linux, running on a GTX 460. This example should run according to the OpenCL 1.1 Specification - and I have run it successfully on prior driver versions. Can someone reproduce this, or am I mising something here?

Thanks

Max

2

Which part of the spec makes you think this should run without a SIGSEGV? If I look at “5.7.2 Setting Kernel Arguments” it says "If the argument is a buffer object, the arg_value pointer can be

NULL or point to a NULL value in which case a NULL value will be used as the value for the argument declared as a pointer to __global or __constant memory in the kernel". So the argument can indeed be NULL, but you’ll have to catch that in the kernel code. If it worked before, that was probably due to some sanity checks that were removed now for performance and conformance reasons.

3

If you look at the example kernel I posted, you will notice that I do indeed check for the flag pointer to be NULL. Furthermore, I don’t even access the memory pointed to by the flag pointer - it is only used to distinguish code paths. So, there should be nothing that causes a segfault from within the kernel, if flags is NULL.

PS:

This is also a case, where the NVIDIA implementation diverges from the specificaction. If i try to pass a NULL pointer, I’m getting an error code returned by clSetKernelArg.

4

You’re right, sorry. For some reason I was only looking for a check for the “data” pointer.

5

Just as a follow-up: Nvidia engineers could reproduce the issue and are currently investigating.