Support zero-length arrays in numba.cuda

[eric-wieser]

As an example,

@numba.cuda.jit
def foo():
    x = numba.cuda.local.array(shape=(2, 0), dtype=numba.int64)

foo()

gives ValueError: array length <= 0.

Currently these lines contain a workaround because of this issue:

numba/numba/cuda/cudadrv/devicearray.py

Lines 432 to 455 in 3399d61

	if ndim == 0:
	# the (2, ndim) allocation below is not yet supported, so avoid it
	@cuda.jit
	def kernel(lhs, rhs):
	lhs[()] = rhs[()]
	return kernel
	@cuda.jit
	def kernel(lhs, rhs):
	location = cuda.grid(1)
	n_elements = 1
	for i in range(lhs.ndim):
	n_elements *= lhs.shape[i]
	if location >= n_elements:
	# bake n_elements into the kernel, better than passing it in
	# as another argument.
	return
	# [0, :] is the to-index (into `lhs`)
	# [1, :] is the from-index (into `rhs`)
	idx = cuda.local.array(
	shape=(2, ndim),
	dtype=types.int64)

[gmarkall]

Thanks for the report / request!

[gmarkall]

See also #5234

[eric-wieser]

#5234 is about arrays with shape () (len(shape) == 0), this issue is about arrays with shape (0,), (0, n) etc (0 in shape).

[gmarkall]

Yes, I don't think they're duplicates, I was just trying to gather together a few shape-related issues in case someone is looking at fixing up things in these areas.

[eric-wieser]

I think I already fixed the other one :)

[JuanFML]

Hello everyone! I am new here and I wanted to see if I could work on this as my first issue?

[AruparnaMaity]

Hi, @eric-wieser , @gmarkall is this issue still open? Can I contribute in any way?

[gmarkall]

@AruparnaMaity If you can see a way to fix it, you could open a PR with the fix, if you'd like - note that as per https://numba.discourse.group/t/rfc-moving-the-cuda-target-to-a-new-package-maintained-by-nvidia/2628 new PRs should go to https://github.com/nvidia/numba-cuda

Many thanks in advance!