Remove special-casing of 0d arrays

An experiment to answer gh-5358

Ok seems this type of experiment doesn't work, I know for a fact from the discussion over in gh-5358 that there are tests looking for this behavior. Presumably there are no CI tests for CUDA code?

There is an internal CI system we can trigger to run PRs on platforms that aren't available on Azure (ARM, CUDA, etc).

I have triggered numba_smoketest_cuda_31 to test this PR.

Thanks!

OK, generates these failing tests:
======================================================================
ERROR: test_broadcast (numba.cuda.tests.cudadrv.test_cuda_array_slicing.CudaArraySetting)
----------------------------------------------------------------------
Traceback (most recent call last):
File "/opt/conda/envs/testenv_8f41183b-f264-4d62-a619-bce1be92cf39/lib/python3.6/site-packages/numba/cuda/tests/cudadrv/test_cuda_array_slicing.py", line 232, in test_broadcast
darr[:, 2] = 500
File "/opt/conda/envs/testenv_8f41183b-f264-4d62-a619-bce1be92cf39/lib/python3.6/site-packages/numba/cuda/cudadrv/devices.py", line 224, in _require_cuda_context
return fn(*args, **kws)
File "/opt/conda/envs/testenv_8f41183b-f264-4d62-a619-bce1be92cf39/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py", line 574, in __setitem__
return self._do_setitem(key, value)
File "/opt/conda/envs/testenv_8f41183b-f264-4d62-a619-bce1be92cf39/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py", line 613, in _do_setitem
rhs_shape[-rhs.ndim:] = rhs.shape
ValueError: cannot copy sequence with size 0 to array axis with dimension 1

======================================================================
ERROR: test_rank (numba.cuda.tests.cudadrv.test_cuda_array_slicing.CudaArraySetting)
----------------------------------------------------------------------
Traceback (most recent call last):
File "/opt/conda/envs/testenv_8f41183b-f264-4d62-a619-bce1be92cf39/lib/python3.6/site-packages/numba/cuda/tests/cudadrv/test_cuda_array_slicing.py", line 225, in test_rank
darr[2] = 500
File "/opt/conda/envs/testenv_8f41183b-f264-4d62-a619-bce1be92cf39/lib/python3.6/site-packages/numba/cuda/cudadrv/devices.py", line 224, in _require_cuda_context
return fn(*args, **kws)
File "/opt/conda/envs/testenv_8f41183b-f264-4d62-a619-bce1be92cf39/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py", line 574, in __setitem__
return self._do_setitem(key, value)
File "/opt/conda/envs/testenv_8f41183b-f264-4d62-a619-bce1be92cf39/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py", line 613, in _do_setitem
rhs_shape[-rhs.ndim:] = rhs.shape
ValueError: cannot copy sequence with size 0 to array axis with dimension 1

======================================================================
ERROR: test_scalar (numba.cuda.tests.cudadrv.test_cuda_array_slicing.CudaArraySetting)
----------------------------------------------------------------------
Traceback (most recent call last):
File "/opt/conda/envs/testenv_8f41183b-f264-4d62-a619-bce1be92cf39/lib/python3.6/site-packages/numba/cuda/tests/cudadrv/test_cuda_array_slicing.py", line 218, in test_scalar
darr[2, 2] = 500
File "/opt/conda/envs/testenv_8f41183b-f264-4d62-a619-bce1be92cf39/lib/python3.6/site-packages/numba/cuda/cudadrv/devices.py", line 224, in _require_cuda_context
return fn(*args, **kws)
File "/opt/conda/envs/testenv_8f41183b-f264-4d62-a619-bce1be92cf39/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py", line 574, in __setitem__
return self._do_setitem(key, value)
File "/opt/conda/envs/testenv_8f41183b-f264-4d62-a619-bce1be92cf39/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py", line 613, in _do_setitem
rhs_shape[-rhs.ndim:] = rhs.shape
ValueError: cannot copy sequence with size 0 to array axis with dimension 1

======================================================================
ERROR: test_strides (numba.cuda.tests.cudadrv.test_cuda_array_slicing.CudaArraySetting)
----------------------------------------------------------------------
Traceback (most recent call last):
File "/opt/conda/envs/testenv_8f41183b-f264-4d62-a619-bce1be92cf39/lib/python3.6/site-packages/numba/cuda/tests/cudadrv/test_cuda_array_slicing.py", line 279, in test_strides
darr[::2] = 500
File "/opt/conda/envs/testenv_8f41183b-f264-4d62-a619-bce1be92cf39/lib/python3.6/site-packages/numba/cuda/cudadrv/devices.py", line 224, in _require_cuda_context
return fn(*args, **kws)
File "/opt/conda/envs/testenv_8f41183b-f264-4d62-a619-bce1be92cf39/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py", line 574, in __setitem__
return self._do_setitem(key, value)
File "/opt/conda/envs/testenv_8f41183b-f264-4d62-a619-bce1be92cf39/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py", line 613, in _do_setitem
rhs_shape[-rhs.ndim:] = rhs.shape
ValueError: cannot copy sequence with size 0 to array axis with dimension 1

======================================================================
ERROR: test_array_views (numba.cuda.tests.cudapy.test_cuda_array_interface.TestCudaArrayInterface)
Views created via array interface support:
----------------------------------------------------------------------
Traceback (most recent call last):
File "/opt/conda/envs/testenv_8f41183b-f264-4d62-a619-bce1be92cf39/lib/python3.6/site-packages/numba/cuda/tests/cudapy/test_cuda_array_interface.py", line 147, in test_array_views
arr[:5] = np.pi
File "/opt/conda/envs/testenv_8f41183b-f264-4d62-a619-bce1be92cf39/lib/python3.6/site-packages/numba/cuda/cudadrv/devices.py", line 224, in _require_cuda_context
return fn(*args, **kws)
File "/opt/conda/envs/testenv_8f41183b-f264-4d62-a619-bce1be92cf39/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py", line 574, in __setitem__
return self._do_setitem(key, value)
File "/opt/conda/envs/testenv_8f41183b-f264-4d62-a619-bce1be92cf39/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py", line 613, in _do_setitem
rhs_shape[-rhs.ndim:] = rhs.shape
ValueError: cannot copy sequence with size 0 to array axis with dimension 1

======================================================================
FAIL: test_device_array_interface (numba.cuda.tests.cudadrv.test_cuda_ndarray.TestCudaNDArray)
----------------------------------------------------------------------
Traceback (most recent call last):
File "/opt/conda/envs/testenv_8f41183b-f264-4d62-a619-bce1be92cf39/lib/python3.6/site-packages/numba/cuda/tests/cudadrv/test_cuda_ndarray.py", line 19, in test_device_array_interface
self.assertEquals(dary.ndim, 1)
AssertionError: 0 != 1

----------------------------------------------------------------------
Ran 564 tests in 118.800s

FAILED (failures=1, errors=5, skipped=8)

Updated, hopefully in a way that will cause all of those tests to fail differently.

@seibert, would you mind re-running the private CI?

running numba_smoketest_cuda_33

======================================================================
ERROR: test_broadcast (numba.cuda.tests.cudadrv.test_cuda_array_slicing.CudaArraySetting)
----------------------------------------------------------------------
Traceback (most recent call last):
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/tests/cudadrv/test_cuda_array_slicing.py", line 232, in test_broadcast
darr[:, 2] = 500
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/cudadrv/devices.py", line 224, in _require_cuda_context
return fn(*args, **kws)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py", line 574, in __setitem__
return self._do_setitem(key, value)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py", line 613, in _do_setitem
rhs_shape[-rhs.ndim:] = rhs.shape
ValueError: cannot copy sequence with size 0 to array axis with dimension 1

======================================================================
ERROR: test_rank (numba.cuda.tests.cudadrv.test_cuda_array_slicing.CudaArraySetting)
----------------------------------------------------------------------
Traceback (most recent call last):
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/tests/cudadrv/test_cuda_array_slicing.py", line 225, in test_rank
darr[2] = 500
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/cudadrv/devices.py", line 224, in _require_cuda_context
return fn(*args, **kws)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py", line 574, in __setitem__
return self._do_setitem(key, value)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py", line 613, in _do_setitem
rhs_shape[-rhs.ndim:] = rhs.shape
ValueError: cannot copy sequence with size 0 to array axis with dimension 1

======================================================================
ERROR: test_scalar (numba.cuda.tests.cudadrv.test_cuda_array_slicing.CudaArraySetting)
----------------------------------------------------------------------
Traceback (most recent call last):
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/errors.py", line 720, in new_error_context
yield
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/lowering.py", line 266, in lower_block
self.lower_inst(inst)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/lowering.py", line 338, in lower_inst
val = self.lower_assign(ty, inst)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/lowering.py", line 512, in lower_assign
return self.lower_expr(ty, value)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/lowering.py", line 977, in lower_expr
res = self.lower_call(resty, expr)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/lowering.py", line 769, in lower_call
res = self._lower_call_normal(fnty, expr, signature)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/lowering.py", line 948, in _lower_call_normal
res = impl(self.builder, argvals, self.loc)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/base.py", line 1154, in __call__
res = self._imp(self._context, builder, self._sig, args, loc=loc)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/base.py", line 1184, in wrapper
return fn(*args, **kwargs)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/cudaimpl.py", line 187, in ptx_lmem_alloc_array
can_dynsized=False)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/cudaimpl.py", line 626, in _generic_array
raise ValueError("array length <= 0")
ValueError: array length <= 0

During handling of the above exception, another exception occurred:

Traceback (most recent call last):
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/tests/cudadrv/test_cuda_array_slicing.py", line 218, in test_scalar
darr[2, 2] = 500
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/cudadrv/devices.py", line 224, in _require_cuda_context
return fn(*args, **kws)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py", line 574, in __setitem__
return self._do_setitem(key, value)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py", line 625, in _do_setitem
_assign_kernel(lhs.ndim).forall(n_elements, stream=stream)(lhs, rhs)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/compiler.py", line 266, in __call__
kernel = self.kernel.specialize(*args)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/compiler.py", line 768, in specialize
kernel = self.compile(argtypes)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/compiler.py", line 784, in compile
**self.targetoptions)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
return func(*args, **kwargs)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/compiler.py", line 56, in compile_kernel
cres = compile_cuda(pyfunc, types.void, args, debug=debug, inline=inline)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
return func(*args, **kwargs)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/compiler.py", line 45, in compile_cuda
locals={})
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/compiler.py", line 549, in compile_extra
return pipeline.compile_extra(func)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/compiler.py", line 329, in compile_extra
return self._compile_bytecode()
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/compiler.py", line 391, in _compile_bytecode
return self._compile_core()
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/compiler.py", line 371, in _compile_core
raise e
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/compiler.py", line 362, in _compile_core
pm.run(self.state)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/compiler_machinery.py", line 341, in run
raise patched_exception
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/compiler_machinery.py", line 332, in run
self._runPass(idx, pass_inst, state)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
return func(*args, **kwargs)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/compiler_machinery.py", line 291, in _runPass
mutated |= check(pss.run_pass, internal_state)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/compiler_machinery.py", line 264, in check
mangled = func(compiler_state)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/typed_passes.py", line 419, in run_pass
NativeLowering().run_pass(state)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/typed_passes.py", line 361, in run_pass
lower.lower()
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/lowering.py", line 178, in lower
self.lower_normal_function(self.fndesc)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/lowering.py", line 226, in lower_normal_function
entry_block_tail = self.lower_function_body()
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/lowering.py", line 251, in lower_function_body
self.lower_block(block)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/lowering.py", line 266, in lower_block
self.lower_inst(inst)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/contextlib.py", line 99, in __exit__
self.gen.throw(type, value, traceback)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/errors.py", line 727, in new_error_context
reraise(type(newerr), newerr, tb)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/core/utils.py", line 78, in reraise
raise value
numba.core.errors.LoweringError: Failed in nopython mode pipeline (step: nopython mode backend)
array length <= 0

File "opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py", line 452:
def kernel(lhs, rhs):
<source elided>
shape=(2, ndim),
dtype=types.int64)
^

[1] During: lowering "$62.10 = call ptx.lmem.alloc(func=ptx.lmem.alloc, args=[], kws=[('shape', Var($62.6, devicearray.py:451)), ('dtype', Var($62.8, devicearray.py:452))], vararg=None)" at /opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py (452)

======================================================================
ERROR: test_strides (numba.cuda.tests.cudadrv.test_cuda_array_slicing.CudaArraySetting)
----------------------------------------------------------------------
Traceback (most recent call last):
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/tests/cudadrv/test_cuda_array_slicing.py", line 279, in test_strides
darr[::2] = 500
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/cudadrv/devices.py", line 224, in _require_cuda_context
return fn(*args, **kws)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py", line 574, in __setitem__
return self._do_setitem(key, value)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py", line 613, in _do_setitem
rhs_shape[-rhs.ndim:] = rhs.shape
ValueError: cannot copy sequence with size 0 to array axis with dimension 1

======================================================================
ERROR: test_array_views (numba.cuda.tests.cudapy.test_cuda_array_interface.TestCudaArrayInterface)
Views created via array interface support:
----------------------------------------------------------------------
Traceback (most recent call last):
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/tests/cudapy/test_cuda_array_interface.py", line 147, in test_array_views
arr[:5] = np.pi
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/cudadrv/devices.py", line 224, in _require_cuda_context
return fn(*args, **kws)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py", line 574, in __setitem__
return self._do_setitem(key, value)
File "/opt/conda/envs/testenv_dd1c91cc-3d96-4dd2-a069-deec58eb57b3/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py", line 613, in _do_setitem
rhs_shape[-rhs.ndim:] = rhs.shape
ValueError: cannot copy sequence with size 0 to array axis with dimension 1

----------------------------------------------------------------------
Ran 564 tests in 116.557s

FAILED (errors=5, skipped=8)

Thanks, that's enough for me to continue

Yes, now there is only one error:
======================================================================
ERROR: test_scalar (numba.cuda.tests.cudadrv.test_cuda_array_slicing.CudaArraySetting)
----------------------------------------------------------------------
Traceback (most recent call last):
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/errors.py", line 720, in new_error_context
yield
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/lowering.py", line 266, in lower_block
self.lower_inst(inst)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/lowering.py", line 338, in lower_inst
val = self.lower_assign(ty, inst)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/lowering.py", line 512, in lower_assign
return self.lower_expr(ty, value)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/lowering.py", line 977, in lower_expr
res = self.lower_call(resty, expr)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/lowering.py", line 769, in lower_call
res = self._lower_call_normal(fnty, expr, signature)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/lowering.py", line 948, in _lower_call_normal
res = impl(self.builder, argvals, self.loc)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/base.py", line 1154, in __call__
res = self._imp(self._context, builder, self._sig, args, loc=loc)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/base.py", line 1184, in wrapper
return fn(*args, **kwargs)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/cuda/cudaimpl.py", line 187, in ptx_lmem_alloc_array
can_dynsized=False)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/cuda/cudaimpl.py", line 626, in _generic_array
raise ValueError("array length <= 0")
ValueError: array length <= 0

During handling of the above exception, another exception occurred:

Traceback (most recent call last):
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/cuda/tests/cudadrv/test_cuda_array_slicing.py", line 218, in test_scalar
darr[2, 2] = 500
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/cuda/cudadrv/devices.py", line 224, in _require_cuda_context
return fn(*args, **kws)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py", line 574, in __setitem__
return self._do_setitem(key, value)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py", line 626, in _do_setitem
_assign_kernel(lhs.ndim).forall(n_elements, stream=stream)(lhs, rhs)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/cuda/compiler.py", line 266, in __call__
kernel = self.kernel.specialize(*args)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/cuda/compiler.py", line 768, in specialize
kernel = self.compile(argtypes)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/cuda/compiler.py", line 784, in compile
**self.targetoptions)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
return func(*args, **kwargs)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/cuda/compiler.py", line 56, in compile_kernel
cres = compile_cuda(pyfunc, types.void, args, debug=debug, inline=inline)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
return func(*args, **kwargs)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/cuda/compiler.py", line 45, in compile_cuda
locals={})
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/compiler.py", line 549, in compile_extra
return pipeline.compile_extra(func)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/compiler.py", line 329, in compile_extra
return self._compile_bytecode()
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/compiler.py", line 391, in _compile_bytecode
return self._compile_core()
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/compiler.py", line 371, in _compile_core
raise e
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/compiler.py", line 362, in _compile_core
pm.run(self.state)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/compiler_machinery.py", line 341, in run
raise patched_exception
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/compiler_machinery.py", line 332, in run
self._runPass(idx, pass_inst, state)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
return func(*args, **kwargs)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/compiler_machinery.py", line 291, in _runPass
mutated |= check(pss.run_pass, internal_state)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/compiler_machinery.py", line 264, in check
mangled = func(compiler_state)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/typed_passes.py", line 419, in run_pass
NativeLowering().run_pass(state)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/typed_passes.py", line 361, in run_pass
lower.lower()
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/lowering.py", line 178, in lower
self.lower_normal_function(self.fndesc)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/lowering.py", line 226, in lower_normal_function
entry_block_tail = self.lower_function_body()
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/lowering.py", line 251, in lower_function_body
self.lower_block(block)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/lowering.py", line 266, in lower_block
self.lower_inst(inst)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/contextlib.py", line 99, in __exit__
self.gen.throw(type, value, traceback)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/errors.py", line 727, in new_error_context
reraise(type(newerr), newerr, tb)
File "/opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/core/utils.py", line 78, in reraise
raise value
numba.core.errors.LoweringError: Failed in nopython mode pipeline (step: nopython mode backend)
array length <= 0

File "opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py", line 452:
def kernel(lhs, rhs):
<source elided>
shape=(2, ndim),
dtype=types.int64)
^

[1] During: lowering "$62.10 = call ptx.lmem.alloc(func=ptx.lmem.alloc, args=[], kws=[('shape', Var($62.6, devicearray.py:451)), ('dtype', Var($62.8, devicearray.py:452))], vararg=None)" at /opt/conda/envs/testenv_56cb10ce-5669-47b8-9f0d-62da54429a32/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py (452)

----------------------------------------------------------------------
Ran 564 tests in 119.809s

FAILED (errors=1, skipped=8)

@seibert : Any chance you could paste that without stripping the leading whitespace?

I was just looking at this too - here is my version:
======================================================================
ERROR: test_scalar (numba.cuda.tests.cudadrv.test_cuda_array_slicing.CudaArraySetting)
----------------------------------------------------------------------
Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 720, in new_error_context
    yield
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 338, in lower_inst
    val = self.lower_assign(ty, inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 512, in lower_assign
    return self.lower_expr(ty, value)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 977, in lower_expr
    res = self.lower_call(resty, expr)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 769, in lower_call
    res = self._lower_call_normal(fnty, expr, signature)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 948, in _lower_call_normal
    res = impl(self.builder, argvals, self.loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1154, in __call__
    res = self._imp(self._context, builder, self._sig, args, loc=loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1184, in wrapper
    return fn(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 187, in ptx_lmem_alloc_array
    can_dynsized=False)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 626, in _generic_array
    raise ValueError("array length <= 0")
ValueError: array length <= 0

During handling of the above exception, another exception occurred:

Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudadrv/test_cuda_array_slicing.py", line 218, in test_scalar
    darr[2, 2] = 500
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudadrv/devices.py", line 224, in _require_cuda_context
    return fn(*args, **kws)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudadrv/devicearray.py", line 574, in __setitem__
    return self._do_setitem(key, value)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudadrv/devicearray.py", line 626, in _do_setitem
    _assign_kernel(lhs.ndim).forall(n_elements, stream=stream)(lhs, rhs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 266, in __call__
    kernel = self.kernel.specialize(*args)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 768, in specialize
    kernel = self.compile(argtypes)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 784, in compile
    **self.targetoptions)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 56, in compile_kernel
    cres = compile_cuda(pyfunc, types.void, args, debug=debug, inline=inline)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 45, in compile_cuda
    locals={})
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 549, in compile_extra
    return pipeline.compile_extra(func)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 329, in compile_extra
    return self._compile_bytecode()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 391, in _compile_bytecode
    return self._compile_core()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 371, in _compile_core
    raise e
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 362, in _compile_core
    pm.run(self.state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 341, in run
    raise patched_exception
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 332, in run
    self._runPass(idx, pass_inst, state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 291, in _runPass
    mutated |= check(pss.run_pass, internal_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 264, in check
    mangled = func(compiler_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 419, in run_pass
    NativeLowering().run_pass(state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 361, in run_pass
    lower.lower()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 178, in lower
    self.lower_normal_function(self.fndesc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 226, in lower_normal_function
    entry_block_tail = self.lower_function_body()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 251, in lower_function_body
    self.lower_block(block)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/miniconda3/envs/numbaenv/lib/python3.7/contextlib.py", line 130, in __exit__
    self.gen.throw(type, value, traceback)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 727, in new_error_context
    reraise(type(newerr), newerr, tb)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/utils.py", line 78, in reraise
    raise value
numba.core.errors.LoweringError: Failed in nopython mode pipeline (step: nopython mode backend)
array length <= 0

File "numba/cuda/cudadrv/devicearray.py", line 452:
    def kernel(lhs, rhs):
        <source elided>
            shape=(2, ndim),
            dtype=types.int64)
            ^

[1] During: lowering "$80call_function_kw.9 = call ptx.lmem.alloc(func=ptx.lmem.alloc, args=[], kws=[('shape', Var($72build_tuple.5, devicearray.py:451)), ('dtype', Var($76load_attr.7, devicearray.py:452))], vararg=None)" at /home/nfs/gmarkall/numbadev/numba/numba/cuda/cudadrv/devicearray.py (452)

----------------------------------------------------------------------
Ran 564 tests in 150.298s

Doesn't look too hard to fix if you're willing to add a special case for ndim == 0 in _assign_kernel:
--- C:\Users\wiese\Repos\ga\numba\numba\cuda\cudadrv\devicearray.py	Mon Mar  9 14:24:39 2020
+++ C:\Users\wiese\Repos\ga\numba\numba\cuda\cudadrv\devicearray.py	Mon Mar  9 16:07:40 2020
@@ -433,6 +433,12 @@
     """
     from numba import cuda  # circular!

+    if ndim == 0:
+        @cuda.jit
+        def kernel(lhs, rhs):
+            lhs[()] = rhs[()]
+        return kernel
+
     @cuda.jit
     def kernel(lhs, rhs):
         location = cuda.grid(1)

I wondered if this would also enable declaration of local arrays with 0 dimensions, so I tried:
diff --git a/numba/cuda/tests/cudapy/test_localmem.py b/numba/cuda/tests/cudapy/test_localmem.py
index b35a456..8151f02 100644
--- a/numba/cuda/tests/cudapy/test_localmem.py
+++ b/numba/cuda/tests/cudapy/test_localmem.py
@@ -28,6 +28,13 @@ def culocal1tuple(A, B):
         B[i] = C[i]

+def culocal0tuple(A, B):
+    C = cuda.local.array((), dtype=int32)
+    i = cuda.grid(1)
+    C = A[i]
+    B[i] = C
+
+
 class TestCudaLocalMem(SerialMixin, unittest.TestCase):
     def test_local_array(self):
         jculocal = cuda.jit('void(int32[:], int32[:])')(culocal)
@@ -37,6 +44,14 @@ class TestCudaLocalMem(SerialMixin, unittest.TestCase):
         jculocal[1, 1](A, B)
         self.assertTrue(np.all(A == B))

+    def test_local_array_0(self):
+        jculocal = cuda.jit('void(int32[:], int32[:])')(culocal0tuple)
+        self.assertTrue('.local' in jculocal.ptx)
+        A = np.arange(256, dtype='int32')
+        B = np.zeros_like(A)
+        jculocal[1, 256](A, B)
+        self.assertTrue(np.all(A == B))
+
     def test_local_array_1_tuple(self):
         """Ensure that the macro can be use with 1-tuple
         """

and also with
    C[0] = A[i]
    B[i] = C[0]

but that still doesn't compile:
======================================================================
ERROR: test_local_array_0 (numba.cuda.tests.cudapy.test_localmem.TestCudaLocalMem)
----------------------------------------------------------------------
Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 720, in new_error_context
    yield
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 338, in lower_inst
    val = self.lower_assign(ty, inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 512, in lower_assign
    return self.lower_expr(ty, value)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 977, in lower_expr
    res = self.lower_call(resty, expr)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 769, in lower_call
    res = self._lower_call_normal(fnty, expr, signature)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 948, in _lower_call_normal
    res = impl(self.builder, argvals, self.loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1154, in __call__
    res = self._imp(self._context, builder, self._sig, args, loc=loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1184, in wrapper
    return fn(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 187, in ptx_lmem_alloc_array
    can_dynsized=False)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 704, in _generic_array
    shape=cgutils.pack_array(builder, kshape),
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/cgutils.py", line 629, in pack_array
    ty = values[0].type
IndexError: list index out of range

During handling of the above exception, another exception occurred:

Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_localmem.py", line 48, in test_local_array_0
    jculocal = cuda.jit('void(int32[:], int32[:])')(culocal0tuple)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/decorators.py", line 96, in kernel_jit
    inline=inline, fastmath=fastmath)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 56, in compile_kernel
    cres = compile_cuda(pyfunc, types.void, args, debug=debug, inline=inline)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 45, in compile_cuda
    locals={})
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 549, in compile_extra
    return pipeline.compile_extra(func)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 329, in compile_extra
    return self._compile_bytecode()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 391, in _compile_bytecode
    return self._compile_core()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 371, in _compile_core
    raise e
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 362, in _compile_core
    pm.run(self.state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 341, in run
    raise patched_exception
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 332, in run
    self._runPass(idx, pass_inst, state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 291, in _runPass
    mutated |= check(pss.run_pass, internal_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 264, in check
    mangled = func(compiler_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 419, in run_pass
    NativeLowering().run_pass(state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 361, in run_pass
    lower.lower()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 178, in lower
    self.lower_normal_function(self.fndesc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 226, in lower_normal_function
    entry_block_tail = self.lower_function_body()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 251, in lower_function_body
    self.lower_block(block)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/miniconda3/envs/numbaenv/lib/python3.7/contextlib.py", line 130, in __exit__
    self.gen.throw(type, value, traceback)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 727, in new_error_context
    reraise(type(newerr), newerr, tb)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/utils.py", line 78, in reraise
    raise value
numba.core.errors.LoweringError: Failed in nopython mode pipeline (step: nopython mode backend)
list index out of range

File "numba/cuda/tests/cudapy/test_localmem.py", line 32:
def culocal0tuple(A, B):
    C = cuda.local.array((), dtype=int32)
    ^

[1] During: lowering "$14call_function_kw.6 = call ptx.lmem.alloc($const8.3, func=ptx.lmem.alloc, args=[Var($const8.3, test_localmem.py:32)], kws=[('dtype', Var($10load_global.4, test_localmem.py:32))], vararg=None)" at /home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_localmem.py (32)

I don't think this matters though - I was mainly worried that this would start to compile, but do something weird / unexpected instead of giving the user some form of traceback.

... I should have been testing with [()], but that also doesn't work - I still don't think that's an issue though.

Can you show the output for when you tried ()?

File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 704, in _generic_array
shape=cgutils.pack_array(builder, kshape),
File "/home/nfs/gmarkall/numbadev/numba/numba/core/cgutils.py", line 629, in pack_array
ty = values[0].type
IndexError: list index out of range

Updated with a fix for this error, and the patch I suggested above

failed due to a typo, repushed

Can you show the output for when you tried ()?

With this test:
diff --git a/numba/cuda/tests/cudapy/test_localmem.py b/numba/cuda/tests/cudapy/test_localmem.py
index b35a456..8346d91 100644
--- a/numba/cuda/tests/cudapy/test_localmem.py
+++ b/numba/cuda/tests/cudapy/test_localmem.py
@@ -28,6 +28,13 @@ def culocal1tuple(A, B):
         B[i] = C[i]

+def culocal0tuple(A, B):
+    C = cuda.local.array((), dtype=int32)
+    i = cuda.grid(1)
+    C[()] = A[i]
+    B[i] = C[()]
+
+
 class TestCudaLocalMem(SerialMixin, unittest.TestCase):
     def test_local_array(self):
         jculocal = cuda.jit('void(int32[:], int32[:])')(culocal)
@@ -37,6 +44,14 @@ class TestCudaLocalMem(SerialMixin, unittest.TestCase):
         jculocal[1, 1](A, B)
         self.assertTrue(np.all(A == B))

+    def test_local_array_0(self):
+        jculocal = cuda.jit('void(int32[:], int32[:])')(culocal0tuple)
+        self.assertTrue('.local' in jculocal.ptx)
+        A = np.arange(256, dtype='int32')
+        B = np.zeros_like(A)
+        jculocal[1, 256](A, B)
+        self.assertTrue(np.all(A == B))
+
     def test_local_array_1_tuple(self):
         """Ensure that the macro can be use with 1-tuple
         """

the test output is:
$ python -m numba.runtests numba.cuda.tests.cudapy.test_localmem
.E..
======================================================================
ERROR: test_local_array_0 (numba.cuda.tests.cudapy.test_localmem.TestCudaLocalMem)
----------------------------------------------------------------------
Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 720, in new_error_context
    yield
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 338, in lower_inst
    val = self.lower_assign(ty, inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 512, in lower_assign
    return self.lower_expr(ty, value)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 977, in lower_expr
    res = self.lower_call(resty, expr)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 769, in lower_call
    res = self._lower_call_normal(fnty, expr, signature)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 948, in _lower_call_normal
    res = impl(self.builder, argvals, self.loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1154, in __call__
    res = self._imp(self._context, builder, self._sig, args, loc=loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1184, in wrapper
    return fn(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 187, in ptx_lmem_alloc_array
    can_dynsized=False)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 704, in _generic_array
    shape=cgutils.pack_array(builder, kshape),
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/cgutils.py", line 629, in pack_array
    ty = values[0].type
IndexError: list index out of range

During handling of the above exception, another exception occurred:

Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_localmem.py", line 48, in test_local_array_0
    jculocal = cuda.jit('void(int32[:], int32[:])')(culocal0tuple)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/decorators.py", line 96, in kernel_jit
    inline=inline, fastmath=fastmath)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 56, in compile_kernel
    cres = compile_cuda(pyfunc, types.void, args, debug=debug, inline=inline)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 45, in compile_cuda
    locals={})
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 549, in compile_extra
    return pipeline.compile_extra(func)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 329, in compile_extra
    return self._compile_bytecode()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 391, in _compile_bytecode
    return self._compile_core()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 371, in _compile_core
    raise e
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 362, in _compile_core
    pm.run(self.state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 341, in run
    raise patched_exception
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 332, in run
    self._runPass(idx, pass_inst, state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 291, in _runPass
    mutated |= check(pss.run_pass, internal_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 264, in check
    mangled = func(compiler_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 419, in run_pass
    NativeLowering().run_pass(state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 361, in run_pass
    lower.lower()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 178, in lower
    self.lower_normal_function(self.fndesc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 226, in lower_normal_function
    entry_block_tail = self.lower_function_body()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 251, in lower_function_body
    self.lower_block(block)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/miniconda3/envs/numbaenv/lib/python3.7/contextlib.py", line 130, in __exit__
    self.gen.throw(type, value, traceback)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 727, in new_error_context
    reraise(type(newerr), newerr, tb)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/utils.py", line 78, in reraise
    raise value
numba.core.errors.LoweringError: Failed in nopython mode pipeline (step: nopython mode backend)
list index out of range

File "numba/cuda/tests/cudapy/test_localmem.py", line 32:
def culocal0tuple(A, B):
    C = cuda.local.array((), dtype=int32)
    ^

[1] During: lowering "$14call_function_kw.6 = call ptx.lmem.alloc($const8.3, func=ptx.lmem.alloc, args=[Var($const8.3, test_localmem.py:32)], kws=[('dtype', Var($10load_global.4, test_localmem.py:32))], vararg=None)" at /home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_localmem.py (32)

----------------------------------------------------------------------
Ran 4 tests in 1.778s

FAILED (errors=1)

Can you try again with my latest update? I fixed the offending line.

Apologies - with the latest changes, all local mem tests fail:
EEEE
======================================================================
ERROR: test_local_array (numba.cuda.tests.cudapy.test_localmem.TestCudaLocalMem)
----------------------------------------------------------------------
Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 720, in new_error_context
    yield
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 338, in lower_inst
    val = self.lower_assign(ty, inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 512, in lower_assign
    return self.lower_expr(ty, value)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 977, in lower_expr
    res = self.lower_call(resty, expr)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 769, in lower_call
    res = self._lower_call_normal(fnty, expr, signature)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 948, in _lower_call_normal
    res = impl(self.builder, argvals, self.loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1154, in __call__
    res = self._imp(self._context, builder, self._sig, args, loc=loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1184, in wrapper
    return fn(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 187, in ptx_lmem_alloc_array
    can_dynsized=False)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 704, in _generic_array
    shape=cgutils.pack_array(builder, kshape, types.intp),
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/cgutils.py", line 632, in pack_array
    ary = builder.insert_value(ary, v, i)
  File "/home/nfs/gmarkall/numbadev/llvmlite/llvmlite/ir/builder.py", line 944, in insert_value
    instr = instructions.InsertValue(self.block, agg, value, idx, name=name)
  File "/home/nfs/gmarkall/numbadev/llvmlite/llvmlite/ir/instructions.py", line 652, in __init__
    % (typ, list(indices), agg.type, elem.type))
TypeError: Can only insert int64 at [0] in [1 x int64]: got i64

During handling of the above exception, another exception occurred:

Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_localmem.py", line 40, in test_local_array
    jculocal = cuda.jit('void(int32[:], int32[:])')(culocal)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/decorators.py", line 96, in kernel_jit
    inline=inline, fastmath=fastmath)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 56, in compile_kernel
    cres = compile_cuda(pyfunc, types.void, args, debug=debug, inline=inline)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 45, in compile_cuda
    locals={})
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 549, in compile_extra
    return pipeline.compile_extra(func)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 329, in compile_extra
    return self._compile_bytecode()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 391, in _compile_bytecode
    return self._compile_core()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 371, in _compile_core
    raise e
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 362, in _compile_core
    pm.run(self.state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 341, in run
    raise patched_exception
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 332, in run
    self._runPass(idx, pass_inst, state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 291, in _runPass
    mutated |= check(pss.run_pass, internal_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 264, in check
    mangled = func(compiler_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 419, in run_pass
    NativeLowering().run_pass(state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 361, in run_pass
    lower.lower()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 178, in lower
    self.lower_normal_function(self.fndesc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 226, in lower_normal_function
    entry_block_tail = self.lower_function_body()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 251, in lower_function_body
    self.lower_block(block)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/miniconda3/envs/numbaenv/lib/python3.7/contextlib.py", line 130, in __exit__
    self.gen.throw(type, value, traceback)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 727, in new_error_context
    reraise(type(newerr), newerr, tb)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/utils.py", line 78, in reraise
    raise value
numba.core.errors.LoweringError: Failed in nopython mode pipeline (step: nopython mode backend)
Can only insert int64 at [0] in [1 x int64]: got i64

File "numba/cuda/tests/cudapy/test_localmem.py", line 8:
def culocal(A, B):
    C = cuda.local.array(1000, dtype=int32)
    ^

[1] During: lowering "$14call_function_kw.6 = call ptx.lmem.alloc($const8.3, func=ptx.lmem.alloc, args=[Var($const8.3, test_localmem.py:8)], kws=[('dtype', Var($10load_global.4, test_localmem.py:8))], vararg=None)" at /home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_localmem.py (8)

======================================================================
ERROR: test_local_array_0 (numba.cuda.tests.cudapy.test_localmem.TestCudaLocalMem)
----------------------------------------------------------------------
Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 720, in new_error_context
    yield
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 338, in lower_inst
    val = self.lower_assign(ty, inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 512, in lower_assign
    return self.lower_expr(ty, value)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 977, in lower_expr
    res = self.lower_call(resty, expr)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 769, in lower_call
    res = self._lower_call_normal(fnty, expr, signature)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 948, in _lower_call_normal
    res = impl(self.builder, argvals, self.loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1154, in __call__
    res = self._imp(self._context, builder, self._sig, args, loc=loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1184, in wrapper
    return fn(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 187, in ptx_lmem_alloc_array
    can_dynsized=False)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 705, in _generic_array
    strides=cgutils.pack_array(builder, kstrides, types.intp),
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/cgutils.py", line 632, in pack_array
    ary = builder.insert_value(ary, v, i)
  File "/home/nfs/gmarkall/numbadev/llvmlite/llvmlite/ir/builder.py", line 944, in insert_value
    instr = instructions.InsertValue(self.block, agg, value, idx, name=name)
  File "/home/nfs/gmarkall/numbadev/llvmlite/llvmlite/ir/instructions.py", line 652, in __init__
    % (typ, list(indices), agg.type, elem.type))
TypeError: Can only insert int64 at [0] in [1 x int64]: got i64

During handling of the above exception, another exception occurred:

Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_localmem.py", line 48, in test_local_array_0
    jculocal = cuda.jit('void(int32[:], int32[:])')(culocal0tuple)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/decorators.py", line 96, in kernel_jit
    inline=inline, fastmath=fastmath)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 56, in compile_kernel
    cres = compile_cuda(pyfunc, types.void, args, debug=debug, inline=inline)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 45, in compile_cuda
    locals={})
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 549, in compile_extra
    return pipeline.compile_extra(func)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 329, in compile_extra
    return self._compile_bytecode()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 391, in _compile_bytecode
    return self._compile_core()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 371, in _compile_core
    raise e
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 362, in _compile_core
    pm.run(self.state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 341, in run
    raise patched_exception
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 332, in run
    self._runPass(idx, pass_inst, state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 291, in _runPass
    mutated |= check(pss.run_pass, internal_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 264, in check
    mangled = func(compiler_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 419, in run_pass
    NativeLowering().run_pass(state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 361, in run_pass
    lower.lower()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 178, in lower
    self.lower_normal_function(self.fndesc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 226, in lower_normal_function
    entry_block_tail = self.lower_function_body()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 251, in lower_function_body
    self.lower_block(block)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/miniconda3/envs/numbaenv/lib/python3.7/contextlib.py", line 130, in __exit__
    self.gen.throw(type, value, traceback)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 727, in new_error_context
    reraise(type(newerr), newerr, tb)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/utils.py", line 78, in reraise
    raise value
numba.core.errors.LoweringError: Failed in nopython mode pipeline (step: nopython mode backend)
Can only insert int64 at [0] in [1 x int64]: got i64

File "numba/cuda/tests/cudapy/test_localmem.py", line 32:
def culocal0tuple(A, B):
    C = cuda.local.array((), dtype=int32)
    ^

[1] During: lowering "$14call_function_kw.6 = call ptx.lmem.alloc($const8.3, func=ptx.lmem.alloc, args=[Var($const8.3, test_localmem.py:32)], kws=[('dtype', Var($10load_global.4, test_localmem.py:32))], vararg=None)" at /home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_localmem.py (32)

======================================================================
ERROR: test_local_array_1_tuple (numba.cuda.tests.cudapy.test_localmem.TestCudaLocalMem)
Ensure that the macro can be use with 1-tuple
----------------------------------------------------------------------
Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 720, in new_error_context
    yield
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 338, in lower_inst
    val = self.lower_assign(ty, inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 512, in lower_assign
    return self.lower_expr(ty, value)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 977, in lower_expr
    res = self.lower_call(resty, expr)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 769, in lower_call
    res = self._lower_call_normal(fnty, expr, signature)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 948, in _lower_call_normal
    res = impl(self.builder, argvals, self.loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1154, in __call__
    res = self._imp(self._context, builder, self._sig, args, loc=loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1184, in wrapper
    return fn(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 187, in ptx_lmem_alloc_array
    can_dynsized=False)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 704, in _generic_array
    shape=cgutils.pack_array(builder, kshape, types.intp),
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/cgutils.py", line 632, in pack_array
    ary = builder.insert_value(ary, v, i)
  File "/home/nfs/gmarkall/numbadev/llvmlite/llvmlite/ir/builder.py", line 944, in insert_value
    instr = instructions.InsertValue(self.block, agg, value, idx, name=name)
  File "/home/nfs/gmarkall/numbadev/llvmlite/llvmlite/ir/instructions.py", line 652, in __init__
    % (typ, list(indices), agg.type, elem.type))
TypeError: Can only insert int64 at [0] in [1 x int64]: got i64

During handling of the above exception, another exception occurred:

Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_localmem.py", line 58, in test_local_array_1_tuple
    jculocal = cuda.jit('void(int32[:], int32[:])')(culocal1tuple)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/decorators.py", line 96, in kernel_jit
    inline=inline, fastmath=fastmath)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 56, in compile_kernel
    cres = compile_cuda(pyfunc, types.void, args, debug=debug, inline=inline)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 45, in compile_cuda
    locals={})
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 549, in compile_extra
    return pipeline.compile_extra(func)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 329, in compile_extra
    return self._compile_bytecode()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 391, in _compile_bytecode
    return self._compile_core()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 371, in _compile_core
    raise e
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 362, in _compile_core
    pm.run(self.state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 341, in run
    raise patched_exception
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 332, in run
    self._runPass(idx, pass_inst, state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 291, in _runPass
    mutated |= check(pss.run_pass, internal_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 264, in check
    mangled = func(compiler_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 419, in run_pass
    NativeLowering().run_pass(state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 361, in run_pass
    lower.lower()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 178, in lower
    self.lower_normal_function(self.fndesc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 226, in lower_normal_function
    entry_block_tail = self.lower_function_body()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 251, in lower_function_body
    self.lower_block(block)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/miniconda3/envs/numbaenv/lib/python3.7/contextlib.py", line 130, in __exit__
    self.gen.throw(type, value, traceback)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 727, in new_error_context
    reraise(type(newerr), newerr, tb)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/utils.py", line 78, in reraise
    raise value
numba.core.errors.LoweringError: Failed in nopython mode pipeline (step: nopython mode backend)
Can only insert int64 at [0] in [1 x int64]: got i64

File "numba/cuda/tests/cudapy/test_localmem.py", line 24:
def culocal1tuple(A, B):
    C = cuda.local.array((5,), dtype=int32)
    ^

[1] During: lowering "$14call_function_kw.6 = call ptx.lmem.alloc($const8.3, func=ptx.lmem.alloc, args=[Var($const8.3, test_localmem.py:24)], kws=[('dtype', Var($10load_global.4, test_localmem.py:24))], vararg=None)" at /home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_localmem.py (24)

======================================================================
ERROR: test_local_array_complex (numba.cuda.tests.cudapy.test_localmem.TestCudaLocalMem)
----------------------------------------------------------------------
Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 720, in new_error_context
    yield
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 338, in lower_inst
    val = self.lower_assign(ty, inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 512, in lower_assign
    return self.lower_expr(ty, value)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 977, in lower_expr
    res = self.lower_call(resty, expr)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 769, in lower_call
    res = self._lower_call_normal(fnty, expr, signature)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 948, in _lower_call_normal
    res = impl(self.builder, argvals, self.loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1154, in __call__
    res = self._imp(self._context, builder, self._sig, args, loc=loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1184, in wrapper
    return fn(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 187, in ptx_lmem_alloc_array
    can_dynsized=False)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 704, in _generic_array
    shape=cgutils.pack_array(builder, kshape, types.intp),
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/cgutils.py", line 632, in pack_array
    ary = builder.insert_value(ary, v, i)
  File "/home/nfs/gmarkall/numbadev/llvmlite/llvmlite/ir/builder.py", line 944, in insert_value
    instr = instructions.InsertValue(self.block, agg, value, idx, name=name)
  File "/home/nfs/gmarkall/numbadev/llvmlite/llvmlite/ir/instructions.py", line 652, in __init__
    % (typ, list(indices), agg.type, elem.type))
TypeError: Can only insert int64 at [0] in [1 x int64]: got i64

During handling of the above exception, another exception occurred:

Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_localmem.py", line 68, in test_local_array_complex
    jculocalcomplex = cuda.jit(sig)(culocalcomplex)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/decorators.py", line 96, in kernel_jit
    inline=inline, fastmath=fastmath)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 56, in compile_kernel
    cres = compile_cuda(pyfunc, types.void, args, debug=debug, inline=inline)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 45, in compile_cuda
    locals={})
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 549, in compile_extra
    return pipeline.compile_extra(func)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 329, in compile_extra
    return self._compile_bytecode()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 391, in _compile_bytecode
    return self._compile_core()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 371, in _compile_core
    raise e
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 362, in _compile_core
    pm.run(self.state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 341, in run
    raise patched_exception
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 332, in run
    self._runPass(idx, pass_inst, state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 291, in _runPass
    mutated |= check(pss.run_pass, internal_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 264, in check
    mangled = func(compiler_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 419, in run_pass
    NativeLowering().run_pass(state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 361, in run_pass
    lower.lower()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 178, in lower
    self.lower_normal_function(self.fndesc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 226, in lower_normal_function
    entry_block_tail = self.lower_function_body()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 251, in lower_function_body
    self.lower_block(block)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/miniconda3/envs/numbaenv/lib/python3.7/contextlib.py", line 130, in __exit__
    self.gen.throw(type, value, traceback)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 727, in new_error_context
    reraise(type(newerr), newerr, tb)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/utils.py", line 78, in reraise
    raise value
numba.core.errors.LoweringError: Failed in nopython mode pipeline (step: nopython mode backend)
Can only insert int64 at [0] in [1 x int64]: got i64

File "numba/cuda/tests/cudapy/test_localmem.py", line 16:
def culocalcomplex(A, B):
    C = cuda.local.array(100, dtype=complex128)
    ^

[1] During: lowering "$14call_function_kw.6 = call ptx.lmem.alloc($const8.3, func=ptx.lmem.alloc, args=[Var($const8.3, test_localmem.py:16)], kws=[('dtype', Var($10load_global.4, test_localmem.py:16))], vararg=None)" at /home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_localmem.py (16)

----------------------------------------------------------------------
Ran 4 tests in 0.161s

FAILED (errors=4)

TypeError: Can only insert int64 at [0] in [1 x int64]: got i64

What's the difference between int64 and i64?

Ah, think I've got it - numba types vs llvm types. Fixed by eliminating the pack_array calls altogether I think.

Thanks for all the testing @gmarkall!

That fixes three fails - I'm still seeing:
$ python -m numba.runtests numba.cuda.tests.cudapy.test_localmem
.E..
======================================================================
ERROR: test_local_array_0 (numba.cuda.tests.cudapy.test_localmem.TestCudaLocalMem)
----------------------------------------------------------------------
Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 720, in new_error_context
    yield
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 338, in lower_inst
    val = self.lower_assign(ty, inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 512, in lower_assign
    return self.lower_expr(ty, value)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 977, in lower_expr
    res = self.lower_call(resty, expr)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 769, in lower_call
    res = self._lower_call_normal(fnty, expr, signature)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 948, in _lower_call_normal
    res = impl(self.builder, argvals, self.loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1154, in __call__
    res = self._imp(self._context, builder, self._sig, args, loc=loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1184, in wrapper
    return fn(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 187, in ptx_lmem_alloc_array
    can_dynsized=False)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 707, in _generic_array
    meminfo=None)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 996, in populate_array
    return arrayobj.populate_array(arr, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/np/arrayobj.py", line 231, in populate_array
    setattr(array, k, v)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/cgutils.py", line 164, in __setattr__
    self[self._datamodel.get_field_position(field)] = value
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/cgutils.py", line 193, in __setitem__
    index=index))
TypeError: Invalid store of [1 x i64] to [0 x i64] in <numba.core.datamodel.models.ArrayModel object at 0x7fac7e0929d0> (trying to write member #6)

During handling of the above exception, another exception occurred:

Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_localmem.py", line 48, in test_local_array_0
    jculocal = cuda.jit('void(int32[:], int32[:])')(culocal0tuple)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/decorators.py", line 96, in kernel_jit
    inline=inline, fastmath=fastmath)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 56, in compile_kernel
    cres = compile_cuda(pyfunc, types.void, args, debug=debug, inline=inline)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 45, in compile_cuda
    locals={})
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 549, in compile_extra
    return pipeline.compile_extra(func)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 329, in compile_extra
    return self._compile_bytecode()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 391, in _compile_bytecode
    return self._compile_core()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 371, in _compile_core
    raise e
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 362, in _compile_core
    pm.run(self.state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 341, in run
    raise patched_exception
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 332, in run
    self._runPass(idx, pass_inst, state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 291, in _runPass
    mutated |= check(pss.run_pass, internal_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 264, in check
    mangled = func(compiler_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 419, in run_pass
    NativeLowering().run_pass(state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 361, in run_pass
    lower.lower()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 178, in lower
    self.lower_normal_function(self.fndesc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 226, in lower_normal_function
    entry_block_tail = self.lower_function_body()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 251, in lower_function_body
    self.lower_block(block)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/miniconda3/envs/numbaenv/lib/python3.7/contextlib.py", line 130, in __exit__
    self.gen.throw(type, value, traceback)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 727, in new_error_context
    reraise(type(newerr), newerr, tb)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/utils.py", line 78, in reraise
    raise value
numba.core.errors.LoweringError: Failed in nopython mode pipeline (step: nopython mode backend)
Invalid store of [1 x i64] to [0 x i64] in <numba.core.datamodel.models.ArrayModel object at 0x7fac7e0929d0> (trying to write member #6)

File "numba/cuda/tests/cudapy/test_localmem.py", line 32:
def culocal0tuple(A, B):
    C = cuda.local.array((), dtype=int32)
    ^

[1] During: lowering "$14call_function_kw.6 = call ptx.lmem.alloc($const8.3, func=ptx.lmem.alloc, args=[Var($const8.3, test_localmem.py:32)], kws=[('dtype', Var($10load_global.4, test_localmem.py:32))], vararg=None)" at /home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_localmem.py (32)

----------------------------------------------------------------------
Ran 4 tests in 1.735s

FAILED (errors=1)

Thanks for all the testing @gmarkall!

No problem - glad to help out a bit!

It occurred to me that maybe shared memory should also be checked. So, with:
diff --git a/numba/cuda/tests/cudapy/test_sm.py b/numba/cuda/tests/cudapy/test_sm.py
index 9265baf..6f56f64 100644
--- a/numba/cuda/tests/cudapy/test_sm.py
+++ b/numba/cuda/tests/cudapy/test_sm.py
@@ -375,6 +375,19 @@ class TestSharedMemory(SerialMixin, unittest.TestCase):
         host_result = d_result.copy_to_host()
         np.testing.assert_array_equal(arr, host_result)

+    def test_zero_dims(self):
+        @cuda.jit
+        def shared0tuple(A, B):
+            C = cuda.shared.array((), dtype=int32)
+            i = cuda.grid(1)
+            C[()] = A[i]
+            B[i] = C[()]
+
+        A = np.arange(256, dtype='int32')
+        B = np.zeros_like(A)
+        shared0tuple[256, 1](A, B)
+        self.assertTrue(np.all(A == B))
+

 if __name__ == '__main__':
     unittest.main()

there is a similar problem:
$ python -m numba.runtests numba.cuda.tests.cudapy.test_sm
................E
======================================================================
ERROR: test_zero_dims (numba.cuda.tests.cudapy.test_sm.TestSharedMemory)
----------------------------------------------------------------------
Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 720, in new_error_context
    yield
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 338, in lower_inst
    val = self.lower_assign(ty, inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 512, in lower_assign
    return self.lower_expr(ty, value)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 977, in lower_expr
    res = self.lower_call(resty, expr)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 769, in lower_call
    res = self._lower_call_normal(fnty, expr, signature)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 948, in _lower_call_normal
    res = impl(self.builder, argvals, self.loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1154, in __call__
    res = self._imp(self._context, builder, self._sig, args, loc=loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1184, in wrapper
    return fn(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 169, in ptx_smem_alloc_array
    can_dynsized=True)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 707, in _generic_array
    meminfo=None)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 996, in populate_array
    return arrayobj.populate_array(arr, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/np/arrayobj.py", line 231, in populate_array
    setattr(array, k, v)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/cgutils.py", line 164, in __setattr__
    self[self._datamodel.get_field_position(field)] = value
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/cgutils.py", line 193, in __setitem__
    index=index))
TypeError: Invalid store of [1 x i64] to [0 x i64] in <numba.core.datamodel.models.ArrayModel object at 0x7f5a9c364090> (trying to write member #6)

During handling of the above exception, another exception occurred:

Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_sm.py", line 388, in test_zero_dims
    shared0tuple[256, 1](A, B)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 757, in __call__
    kernel = self.specialize(*args)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 768, in specialize
    kernel = self.compile(argtypes)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 784, in compile
    **self.targetoptions)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 56, in compile_kernel
    cres = compile_cuda(pyfunc, types.void, args, debug=debug, inline=inline)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 45, in compile_cuda
    locals={})
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 549, in compile_extra
    return pipeline.compile_extra(func)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 329, in compile_extra
    return self._compile_bytecode()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 391, in _compile_bytecode
    return self._compile_core()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 371, in _compile_core
    raise e
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 362, in _compile_core
    pm.run(self.state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 341, in run
    raise patched_exception
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 332, in run
    self._runPass(idx, pass_inst, state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 291, in _runPass
    mutated |= check(pss.run_pass, internal_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 264, in check
    mangled = func(compiler_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 419, in run_pass
    NativeLowering().run_pass(state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 361, in run_pass
    lower.lower()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 178, in lower
    self.lower_normal_function(self.fndesc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 226, in lower_normal_function
    entry_block_tail = self.lower_function_body()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 251, in lower_function_body
    self.lower_block(block)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/miniconda3/envs/numbaenv/lib/python3.7/contextlib.py", line 130, in __exit__
    self.gen.throw(type, value, traceback)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 727, in new_error_context
    reraise(type(newerr), newerr, tb)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/utils.py", line 78, in reraise
    raise value
numba.core.errors.LoweringError: Failed in nopython mode pipeline (step: nopython mode backend)
Invalid store of [1 x i64] to [0 x i64] in <numba.core.datamodel.models.ArrayModel object at 0x7f5a9c364090> (trying to write member #6)

File "numba/cuda/tests/cudapy/test_sm.py", line 381:
        def shared0tuple(A, B):
            C = cuda.shared.array((), dtype=int32)
            ^

[1] During: lowering "$14call_function_kw.6 = call ptx.smem.alloc($const8.3, func=ptx.smem.alloc, args=[Var($const8.3, test_sm.py:381)], kws=[('dtype', Var($10load_global.4, test_sm.py:381))], vararg=None)" at /home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_sm.py (381)

----------------------------------------------------------------------
Ran 17 tests in 3.175s

FAILED (errors=1)

I'm also not sure it's hugely useful to have 0d local and shared arrays, so if it's a lot of hassle to sort out, I think it would be acceptable to raise an error for a 0d shared or local array to inform the user that they're not supported.

So, your traceback goes through ptx_smem_alloc_array, which calls _generic_array with can_dynsized=True. From the source for that situation:
https://github.com/numba/numba/blob/master/numba/cuda/cudaimpl.py#L679-L693
The crash here is because len(shape) != len(kshape), since in that path len(kshape) == 1 is always true. I think it would indeed make sense to raise an exception if an attempt to use dynamic memory on a non-1d array is used.

Updated once again to not use the dynamic path for non-1d arrays.

@gmarkall: Mind re-testing?

@eric-wieser There's something about this particular set of changes that manages to trigger https://github.com/numba/numba/issues/4954 reliably, so I'm having trouble running the whole test suite - I don't think this is directly related to your changes, though - I suspect someone with a different GPU / driver might not run into it.
However, for the shared and local memory tests:
$ python -m numba.runtests numba.cuda.tests.cudapy.test_localmem
.E..
======================================================================
ERROR: test_local_array_0 (numba.cuda.tests.cudapy.test_localmem.TestCudaLocalMem)
----------------------------------------------------------------------
Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 720, in new_error_context
    yield
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 338, in lower_inst
    val = self.lower_assign(ty, inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 512, in lower_assign
    return self.lower_expr(ty, value)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 977, in lower_expr
    res = self.lower_call(resty, expr)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 769, in lower_call
    res = self._lower_call_normal(fnty, expr, signature)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 948, in _lower_call_normal
    res = impl(self.builder, argvals, self.loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1154, in __call__
    res = self._imp(self._context, builder, self._sig, args, loc=loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1184, in wrapper
    return fn(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 187, in ptx_lmem_alloc_array
    can_dynsized=False)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 708, in _generic_array
    meminfo=None)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 996, in populate_array
    return arrayobj.populate_array(arr, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/np/arrayobj.py", line 231, in populate_array
    setattr(array, k, v)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/cgutils.py", line 164, in __setattr__
    self[self._datamodel.get_field_position(field)] = value
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/cgutils.py", line 193, in __setitem__
    index=index))
TypeError: Invalid store of [1 x i64] to [0 x i64] in <numba.core.datamodel.models.ArrayModel object at 0x7fb8dfebbd10> (trying to write member #6)

During handling of the above exception, another exception occurred:

Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_localmem.py", line 48, in test_local_array_0
    jculocal = cuda.jit('void(int32[:], int32[:])')(culocal0tuple)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/decorators.py", line 96, in kernel_jit
    inline=inline, fastmath=fastmath)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 56, in compile_kernel
    cres = compile_cuda(pyfunc, types.void, args, debug=debug, inline=inline)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 45, in compile_cuda
    locals={})
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 549, in compile_extra
    return pipeline.compile_extra(func)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 329, in compile_extra
    return self._compile_bytecode()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 391, in _compile_bytecode
    return self._compile_core()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 371, in _compile_core
    raise e
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 362, in _compile_core
    pm.run(self.state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 341, in run
    raise patched_exception
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 332, in run
    self._runPass(idx, pass_inst, state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 291, in _runPass
    mutated |= check(pss.run_pass, internal_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 264, in check
    mangled = func(compiler_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 419, in run_pass
    NativeLowering().run_pass(state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 361, in run_pass
    lower.lower()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 178, in lower
    self.lower_normal_function(self.fndesc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 226, in lower_normal_function
    entry_block_tail = self.lower_function_body()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 251, in lower_function_body
    self.lower_block(block)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/miniconda3/envs/numbaenv/lib/python3.7/contextlib.py", line 130, in __exit__
    self.gen.throw(type, value, traceback)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 727, in new_error_context
    reraise(type(newerr), newerr, tb)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/utils.py", line 78, in reraise
    raise value
numba.core.errors.LoweringError: Failed in nopython mode pipeline (step: nopython mode backend)
Invalid store of [1 x i64] to [0 x i64] in <numba.core.datamodel.models.ArrayModel object at 0x7fb8dfebbd10> (trying to write member #6)

File "numba/cuda/tests/cudapy/test_localmem.py", line 32:
def culocal0tuple(A, B):
    C = cuda.local.array((), dtype=int32)
    ^

[1] During: lowering "$14call_function_kw.6 = call ptx.lmem.alloc($const8.3, func=ptx.lmem.alloc, args=[Var($const8.3, test_localmem.py:32)], kws=[('dtype', Var($10load_global.4, test_localmem.py:32))], vararg=None)" at /home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_localmem.py (32)

----------------------------------------------------------------------
Ran 4 tests in 1.482s

FAILED (errors=1)

$ python -m numba.runtests numba.cuda.tests.cudapy.test_sm
................E
======================================================================
ERROR: test_zero_dims (numba.cuda.tests.cudapy.test_sm.TestSharedMemory)
----------------------------------------------------------------------
Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 720, in new_error_context
    yield
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 338, in lower_inst
    val = self.lower_assign(ty, inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 512, in lower_assign
    return self.lower_expr(ty, value)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 977, in lower_expr
    res = self.lower_call(resty, expr)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 769, in lower_call
    res = self._lower_call_normal(fnty, expr, signature)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 948, in _lower_call_normal
    res = impl(self.builder, argvals, self.loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1154, in __call__
    res = self._imp(self._context, builder, self._sig, args, loc=loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1184, in wrapper
    return fn(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 169, in ptx_smem_alloc_array
    can_dynsized=True)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 708, in _generic_array
    meminfo=None)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 996, in populate_array
    return arrayobj.populate_array(arr, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/np/arrayobj.py", line 231, in populate_array
    setattr(array, k, v)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/cgutils.py", line 164, in __setattr__
    self[self._datamodel.get_field_position(field)] = value
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/cgutils.py", line 193, in __setitem__
    index=index))
TypeError: Invalid store of [1 x i64] to [0 x i64] in <numba.core.datamodel.models.ArrayModel object at 0x7fe7e6d1ed10> (trying to write member #6)

During handling of the above exception, another exception occurred:

Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_sm.py", line 388, in test_zero_dims
    shared0tuple[256, 1](A, B)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 757, in __call__
    kernel = self.specialize(*args)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 768, in specialize
    kernel = self.compile(argtypes)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 784, in compile
    **self.targetoptions)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 56, in compile_kernel
    cres = compile_cuda(pyfunc, types.void, args, debug=debug, inline=inline)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 45, in compile_cuda
    locals={})
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 549, in compile_extra
    return pipeline.compile_extra(func)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 329, in compile_extra
    return self._compile_bytecode()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 391, in _compile_bytecode
    return self._compile_core()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 371, in _compile_core
    raise e
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 362, in _compile_core
    pm.run(self.state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 341, in run
    raise patched_exception
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 332, in run
    self._runPass(idx, pass_inst, state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 291, in _runPass
    mutated |= check(pss.run_pass, internal_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 264, in check
    mangled = func(compiler_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 419, in run_pass
    NativeLowering().run_pass(state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 361, in run_pass
    lower.lower()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 178, in lower
    self.lower_normal_function(self.fndesc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 226, in lower_normal_function
    entry_block_tail = self.lower_function_body()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 251, in lower_function_body
    self.lower_block(block)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/miniconda3/envs/numbaenv/lib/python3.7/contextlib.py", line 130, in __exit__
    self.gen.throw(type, value, traceback)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 727, in new_error_context
    reraise(type(newerr), newerr, tb)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/utils.py", line 78, in reraise
    raise value
numba.core.errors.LoweringError: Failed in nopython mode pipeline (step: nopython mode backend)
Invalid store of [1 x i64] to [0 x i64] in <numba.core.datamodel.models.ArrayModel object at 0x7fe7e6d1ed10> (trying to write member #6)

File "numba/cuda/tests/cudapy/test_sm.py", line 381:
        def shared0tuple(A, B):
            C = cuda.shared.array((), dtype=int32)
            ^

[1] During: lowering "$14call_function_kw.6 = call ptx.smem.alloc($const8.3, func=ptx.smem.alloc, args=[Var($const8.3, test_sm.py:381)], kws=[('dtype', Var($10load_global.4, test_sm.py:381))], vararg=None)" at /home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_sm.py (381)

----------------------------------------------------------------------
Ran 17 tests in 3.389s

FAILED (errors=1)

Spotted the next issue, thanks - it would be handy if the Invalid store of ... error could include the field name.

With your latest commit, all tests run for me, and there are 5 fails:
======================================================================
ERROR: test_result_on_device (numba.cuda.tests.cudapy.test_reduction.TestReduction)
----------------------------------------------------------------------
Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_reduction.py", line 70, in test_result_on_device
    res = sum_reduce(A, res=got)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/kernels/reduction.py", line 251, in __call__
    not full_blockct)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 759, in __call__
    cfg(*args)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 523, in __call__
    sharedmem=self.sharedmem)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 597, in _kernel_call
    cu_func(*kernelargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudadrv/driver.py", line 1584, in __call__
    self.sharedmem, streamhandle, args)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudadrv/driver.py", line 1628, in launch_kernel
    None)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudadrv/driver.py", line 293, in safe_cuda_api_call
    self._check_error(fname, retcode)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudadrv/driver.py", line 328, in _check_error
    raise CudaAPIError(retcode, msg)
numba.cuda.cudadrv.driver.CudaAPIError: [701] Call to cuLaunchKernel results in CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES

======================================================================
ERROR: test_sum_reduce (numba.cuda.tests.cudapy.test_reduction.TestReduction)
----------------------------------------------------------------------
Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_reduction.py", line 30, in test_sum_reduce
    self._sum_reduce(n)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_reduction.py", line 15, in _sum_reduce
    got = sum_reduce(A)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/kernels/reduction.py", line 251, in __call__
    not full_blockct)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 759, in __call__
    cfg(*args)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 523, in __call__
    sharedmem=self.sharedmem)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 597, in _kernel_call
    cu_func(*kernelargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudadrv/driver.py", line 1584, in __call__
    self.sharedmem, streamhandle, args)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudadrv/driver.py", line 1628, in launch_kernel
    None)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudadrv/driver.py", line 293, in safe_cuda_api_call
    self._check_error(fname, retcode)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudadrv/driver.py", line 328, in _check_error
    raise CudaAPIError(retcode, msg)
numba.cuda.cudadrv.driver.CudaAPIError: [701] Call to cuLaunchKernel results in CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES

======================================================================
ERROR: test_uniform_float32 (numba.cuda.tests.cudapy.test_random.TestCudaRandomXoroshiro128p)
----------------------------------------------------------------------
Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_random.py", line 77, in test_uniform_float32
    self.check_uniform(rng_kernel_float32, np.float32)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_random.py", line 70, in check_uniform
    kernel_func[2, 32](states, out, 32, UNIFORM)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 759, in __call__
    cfg(*args)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 523, in __call__
    sharedmem=self.sharedmem)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 597, in _kernel_call
    cu_func(*kernelargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudadrv/driver.py", line 1584, in __call__
    self.sharedmem, streamhandle, args)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudadrv/driver.py", line 1628, in launch_kernel
    None)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudadrv/driver.py", line 293, in safe_cuda_api_call
    self._check_error(fname, retcode)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudadrv/driver.py", line 328, in _check_error
    raise CudaAPIError(retcode, msg)
numba.cuda.cudadrv.driver.CudaAPIError: [201] Call to cuLaunchKernel results in CUDA_ERROR_INVALID_CONTEXT

======================================================================
ERROR: test_local_array_0 (numba.cuda.tests.cudapy.test_localmem.TestCudaLocalMem)
----------------------------------------------------------------------
Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 720, in new_error_context
    yield
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 338, in lower_inst
    val = self.lower_assign(ty, inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 512, in lower_assign
    return self.lower_expr(ty, value)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 977, in lower_expr
    res = self.lower_call(resty, expr)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 769, in lower_call
    res = self._lower_call_normal(fnty, expr, signature)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 948, in _lower_call_normal
    res = impl(self.builder, argvals, self.loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1154, in __call__
    res = self._imp(self._context, builder, self._sig, args, loc=loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1184, in wrapper
    return fn(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 187, in ptx_lmem_alloc_array
    can_dynsized=False)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 708, in _generic_array
    meminfo=None)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 996, in populate_array
    return arrayobj.populate_array(arr, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/np/arrayobj.py", line 231, in populate_array
    setattr(array, k, v)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/cgutils.py", line 164, in __setattr__
    self[self._datamodel.get_field_position(field)] = value
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/cgutils.py", line 193, in __setitem__
    index=index))
TypeError: Invalid store of [1 x i64] to [0 x i64] in <numba.core.datamodel.models.ArrayModel object at 0x7fa3d45e21d0> (trying to write member #6)

During handling of the above exception, another exception occurred:

Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_localmem.py", line 48, in test_local_array_0
    jculocal = cuda.jit('void(int32[:], int32[:])')(culocal0tuple)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/decorators.py", line 96, in kernel_jit
    inline=inline, fastmath=fastmath)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 56, in compile_kernel
    cres = compile_cuda(pyfunc, types.void, args, debug=debug, inline=inline)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 45, in compile_cuda
    locals={})
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 549, in compile_extra
    return pipeline.compile_extra(func)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 329, in compile_extra
    return self._compile_bytecode()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 391, in _compile_bytecode
    return self._compile_core()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 371, in _compile_core
    raise e
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 362, in _compile_core
    pm.run(self.state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 341, in run
    raise patched_exception
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 332, in run
    self._runPass(idx, pass_inst, state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 291, in _runPass
    mutated |= check(pss.run_pass, internal_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 264, in check
    mangled = func(compiler_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 419, in run_pass
    NativeLowering().run_pass(state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 361, in run_pass
    lower.lower()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 178, in lower
    self.lower_normal_function(self.fndesc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 226, in lower_normal_function
    entry_block_tail = self.lower_function_body()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 251, in lower_function_body
    self.lower_block(block)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/miniconda3/envs/numbaenv/lib/python3.7/contextlib.py", line 130, in __exit__
    self.gen.throw(type, value, traceback)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 727, in new_error_context
    reraise(type(newerr), newerr, tb)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/utils.py", line 78, in reraise
    raise value
numba.core.errors.LoweringError: Failed in nopython mode pipeline (step: nopython mode backend)
Invalid store of [1 x i64] to [0 x i64] in <numba.core.datamodel.models.ArrayModel object at 0x7fa3d45e21d0> (trying to write member #6)

File "numba/cuda/tests/cudapy/test_localmem.py", line 32:
def culocal0tuple(A, B):
    C = cuda.local.array((), dtype=int32)
    ^

[1] During: lowering "$14call_function_kw.6 = call ptx.lmem.alloc($const8.3, func=ptx.lmem.alloc, args=[Var($const8.3, test_localmem.py:32)], kws=[('dtype', Var($10load_global.4, test_localmem.py:32))], vararg=None)" at /home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_localmem.py (32)

======================================================================
ERROR: test_zero_dims (numba.cuda.tests.cudapy.test_sm.TestSharedMemory)
----------------------------------------------------------------------
Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 720, in new_error_context
    yield
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 338, in lower_inst
    val = self.lower_assign(ty, inst)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 512, in lower_assign
    return self.lower_expr(ty, value)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 977, in lower_expr
    res = self.lower_call(resty, expr)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 769, in lower_call
    res = self._lower_call_normal(fnty, expr, signature)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 948, in _lower_call_normal
    res = impl(self.builder, argvals, self.loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1154, in __call__
    res = self._imp(self._context, builder, self._sig, args, loc=loc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 1184, in wrapper
    return fn(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 169, in ptx_smem_alloc_array
    can_dynsized=True)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/cudaimpl.py", line 708, in _generic_array
    meminfo=None)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/base.py", line 996, in populate_array
    return arrayobj.populate_array(arr, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/np/arrayobj.py", line 231, in populate_array
    setattr(array, k, v)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/cgutils.py", line 164, in __setattr__
    self[self._datamodel.get_field_position(field)] = value
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/cgutils.py", line 193, in __setitem__
    index=index))
TypeError: Invalid store of [1 x i64] to [0 x i64] in <numba.core.datamodel.models.ArrayModel object at 0x7fa3d45e21d0> (trying to write member #6)

During handling of the above exception, another exception occurred:

Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_sm.py", line 388, in test_zero_dims
    shared0tuple[256, 1](A, B)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 757, in __call__
    kernel = self.specialize(*args)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 768, in specialize
    kernel = self.compile(argtypes)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 784, in compile
    **self.targetoptions)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 56, in compile_kernel
    cres = compile_cuda(pyfunc, types.void, args, debug=debug, inline=inline)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 45, in compile_cuda
    locals={})
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 549, in compile_extra
    return pipeline.compile_extra(func)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 329, in compile_extra
    return self._compile_bytecode()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 391, in _compile_bytecode
    return self._compile_core()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 371, in _compile_core
    raise e
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler.py", line 362, in _compile_core
    pm.run(self.state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 341, in run
    raise patched_exception
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 332, in run
    self._runPass(idx, pass_inst, state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 291, in _runPass
    mutated |= check(pss.run_pass, internal_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 264, in check
    mangled = func(compiler_state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 419, in run_pass
    NativeLowering().run_pass(state)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 361, in run_pass
    lower.lower()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 178, in lower
    self.lower_normal_function(self.fndesc)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 226, in lower_normal_function
    entry_block_tail = self.lower_function_body()
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 251, in lower_function_body
    self.lower_block(block)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/lowering.py", line 266, in lower_block
    self.lower_inst(inst)
  File "/home/nfs/gmarkall/miniconda3/envs/numbaenv/lib/python3.7/contextlib.py", line 130, in __exit__
    self.gen.throw(type, value, traceback)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/errors.py", line 727, in new_error_context
    reraise(type(newerr), newerr, tb)
  File "/home/nfs/gmarkall/numbadev/numba/numba/core/utils.py", line 78, in reraise
    raise value
numba.core.errors.LoweringError: Failed in nopython mode pipeline (step: nopython mode backend)
Invalid store of [1 x i64] to [0 x i64] in <numba.core.datamodel.models.ArrayModel object at 0x7fa3d45e21d0> (trying to write member #6)

File "numba/cuda/tests/cudapy/test_sm.py", line 381:
        def shared0tuple(A, B):
            C = cuda.shared.array((), dtype=int32)
            ^

[1] During: lowering "$14call_function_kw.6 = call ptx.smem.alloc($const8.3, func=ptx.smem.alloc, args=[Var($const8.3, test_sm.py:381)], kws=[('dtype', Var($10load_global.4, test_sm.py:381))], vararg=None)" at /home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_sm.py (381)

----------------------------------------------------------------------
Ran 566 tests in 156.689s

FAILED (errors=5, skipped=5)

That doesn't look like you actually have my change. From what I deduced this error comes from the strides, which is what I fixed in the last diff.

TypeError: Invalid store of [1 x i64] to [0 x i64] in <numba.core.datamodel.models.ArrayModel object at 0x7fa3d45e21d0> (trying to write member #6)

Would you prefer me not to force push in future?

Oops, sorry... I did git fetch but not git reset ... - I don't think force pushing is the problem, because I'd probably just forget to do git merge instead (I never do git pull because I prefer to check what's in the remote before merging).

Were you able to retest?

@eric-wieser apologies for the delay in a re-test - it seems that #4954 is breaking the running of tests quite consistently right now so it's hard to run the test suite as a whole. I'll try and run again now and report back what I can.

@eric-wieser As far as I can tell, there is one relevant failure:
FAIL: test_local_array_0 (numba.cuda.tests.cudapy.test_localmem.TestCudaLocalMem)
----------------------------------------------------------------------
Traceback (most recent call last):
  File "/home/nfs/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_localmem.py", line 49, in test_local_array_0
    self.assertTrue('.local' in jculocal.ptx)
AssertionError: False is not true

----------------------------------------------------------------------

... from the test I suggested earlier. The PTX is:
test_local_array (numba.cuda.tests.cudapy.test_localmem.TestCudaLocalMem) ... ok
test_local_array_0 (numba.cuda.tests.cudapy.test_localmem.TestCudaLocalMem) ... //
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-24330188
// Cuda compilation tools, release 9.2, V9.2.148
// Based on LLVM 3.4svn
//

.version 6.2
.target sm_70
.address_size 64

	// .globl	_ZN6cudapy5numba4cuda5tests6cudapy13test_localmem17culocal0tuple$242E5ArrayIiLi1E1A7mutable7alignedE5ArrayIiLi1E1A7mutable7alignedE
.visible .global .align 4 .u32 _ZN6cudapy5numba4cuda5tests6cudapy13test_localmem17culocal0tuple$242E5ArrayIiLi1E1A7mutable7alignedE5ArrayIiLi1E1A7mutable7alignedE__errcode__;
.visible .global .align 4 .u32 _ZN6cudapy5numba4cuda5tests6cudapy13test_localmem17culocal0tuple$242E5ArrayIiLi1E1A7mutable7alignedE5ArrayIiLi1E1A7mutable7alignedE__tidx__;
.visible .global .align 4 .u32 _ZN6cudapy5numba4cuda5tests6cudapy13test_localmem17culocal0tuple$242E5ArrayIiLi1E1A7mutable7alignedE5ArrayIiLi1E1A7mutable7alignedE__ctaidx__;
.visible .global .align 4 .u32 _ZN6cudapy5numba4cuda5tests6cudapy13test_localmem17culocal0tuple$242E5ArrayIiLi1E1A7mutable7alignedE5ArrayIiLi1E1A7mutable7alignedE__tidy__;
.visible .global .align 4 .u32 _ZN6cudapy5numba4cuda5tests6cudapy13test_localmem17culocal0tuple$242E5ArrayIiLi1E1A7mutable7alignedE5ArrayIiLi1E1A7mutable7alignedE__ctaidy__;
.visible .global .align 4 .u32 _ZN6cudapy5numba4cuda5tests6cudapy13test_localmem17culocal0tuple$242E5ArrayIiLi1E1A7mutable7alignedE5ArrayIiLi1E1A7mutable7alignedE__tidz__;
.visible .global .align 4 .u32 _ZN6cudapy5numba4cuda5tests6cudapy13test_localmem17culocal0tuple$242E5ArrayIiLi1E1A7mutable7alignedE5ArrayIiLi1E1A7mutable7alignedE__ctaidz__;
.common .global .align 8 .u64 _ZN08NumbaEnv5numba4cuda5tests6cudapy13test_localmem17culocal0tuple$242E5ArrayIiLi1E1A7mutable7alignedE5ArrayIiLi1E1A7mutable7alignedE;

.visible .entry _ZN6cudapy5numba4cuda5tests6cudapy13test_localmem17culocal0tuple$242E5ArrayIiLi1E1A7mutable7alignedE5ArrayIiLi1E1A7mutable7alignedE(
	.param .u64 _ZN6cudapy5numba4cuda5tests6cudapy13test_localmem17culocal0tuple$242E5ArrayIiLi1E1A7mutable7alignedE5ArrayIiLi1E1A7mutable7alignedE_param_0,
	.param .u64 _ZN6cudapy5numba4cuda5tests6cudapy13test_localmem17culocal0tuple$242E5ArrayIiLi1E1A7mutable7alignedE5ArrayIiLi1E1A7mutable7alignedE_param_1,
	.param .u64 _ZN6cudapy5numba4cuda5tests6cudapy13test_localmem17culocal0tuple$242E5ArrayIiLi1E1A7mutable7alignedE5ArrayIiLi1E1A7mutable7alignedE_param_2,
	.param .u64 _ZN6cudapy5numba4cuda5tests6cudapy13test_localmem17culocal0tuple$242E5ArrayIiLi1E1A7mutable7alignedE5ArrayIiLi1E1A7mutable7alignedE_param_3,
	.param .u64 _ZN6cudapy5numba4cuda5tests6cudapy13test_localmem17culocal0tuple$242E5ArrayIiLi1E1A7mutable7alignedE5ArrayIiLi1E1A7mutable7alignedE_param_4,
	.param .u64 _ZN6cudapy5numba4cuda5tests6cudapy13test_localmem17culocal0tuple$242E5ArrayIiLi1E1A7mutable7alignedE5ArrayIiLi1E1A7mutable7alignedE_param_5,