[FEA] Allow passing of NULL pointer in CUDA function

[gmarkall]

From numba/numba#9655:

There is currently not a way to pass a NULL pointer to a C function from a compiled python function in CUDA. The value cffi.NULL should be valid in such functions.

For example:

from numba import cuda
import cffi

ffi = cffi.FFI()

func = cuda.declare_device('func', 'int32(CPointer(int32))')

@cuda.jit(link=['/content/func.cu'])
def kernel(values) :
  i_thread = cuda.threadIdx.x

  values[i_thread] = func(ffi.NULL)

cc @ed-o-saurus

[jakirkham]

In Cython code, we have had some success using uintptr_t (for example)

Cython treats uintptr_t as coercible to/from a Python int object. So users are then able to pass 0 then have that treated as NULL at the Cython/C/C++ layer. There may be some explicit casting between uintptr_t and C/C++ pointer types like void* (for example), but this is pretty trivial/fast to do

Not sure exactly how this maps to numba-cuda's model of working with C/C++, but hopefully this is a helpful way of thinking about the problem