LoweringError while allocating shared memory

[jetxeberria]

Numba version: 0.46.0

Running this code with "python error_minimal_execution.py":

import numpy as np
from numba import cuda, types

@cuda.jit
def foo():
    assert (cuda.blockDim.x, cuda.blockDim.y) == (32, 32)
    a_cache = cuda.shared.array((32, 32), types.int32)
    b_cache = cuda.shared.array((cuda.blockDim.x, cuda.blockDim.y), types.int32)

foo[1, (32, 32)]()

Raises an error in 'b_cache' but not in 'a_cache':

Traceback (most recent call last):
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/errors.py", line 717, in new_error_context
    yield
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/lowering.py", line 260, in lower_block
    self.lower_inst(inst)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/lowering.py", line 303, in lower_inst
    val = self.lower_assign(ty, inst)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/lowering.py", line 465, in lower_assign
    return self.lower_expr(ty, value)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/lowering.py", line 927, in lower_expr
    res = self.lower_call(resty, expr)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/lowering.py", line 719, in lower_call
    res = self._lower_call_normal(fnty, expr, signature)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/lowering.py", line 898, in _lower_call_normal
    res = impl(self.builder, argvals, self.loc)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/targets/base.py", line 1132, in __call__
    res = self._imp(self._context, builder, self._sig, args, loc=loc)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/targets/base.py", line 1157, in wrapper
    return fn(*args, **kwargs)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/cuda/cudaimpl.py", line 174, in ptx_smem_alloc_array
    can_dynsized=True)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/cuda/cudaimpl.py", line 616, in _generic_array
    elemcount = reduce(operator.mul, shape)
TypeError: unsupported operand type(s) for *: 'Macro' and 'Macro'

During handling of the above exception, another exception occurred:

Traceback (most recent call last):
  File "error_minimal_execution.py", line 10, in <module>
    foo[1, (32, 32)]()
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/cuda/compiler.py", line 799, in __call__
    kernel = self.specialize(*args)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/cuda/compiler.py", line 810, in specialize
    kernel = self.compile(argtypes)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/cuda/compiler.py", line 826, in compile
    **self.targetoptions)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/cuda/compiler.py", line 62, in compile_kernel
    cres = compile_cuda(pyfunc, types.void, args, debug=debug, inline=inline)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/cuda/compiler.py", line 51, in compile_cuda
    locals={})
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/compiler.py", line 528, in compile_extra
    return pipeline.compile_extra(func)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/compiler.py", line 326, in compile_extra
    return self._compile_bytecode()
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/compiler.py", line 385, in _compile_bytecode
    return self._compile_core()
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/compiler.py", line 365, in _compile_core
    raise e
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/compiler.py", line 356, in _compile_core
    pm.run(self.state)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/compiler_machinery.py", line 328, in run
    raise patched_exception
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/compiler_machinery.py", line 319, in run
    self._runPass(idx, pass_inst, state)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/compiler_machinery.py", line 281, in _runPass
    mutated |= check(pss.run_pass, internal_state)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/compiler_machinery.py", line 268, in check
    mangled = func(compiler_state)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/typed_passes.py", line 380, in run_pass
    NativeLowering().run_pass(state) # TODO: Pull this out into the pipeline
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/typed_passes.py", line 325, in run_pass
    lower.lower()
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/lowering.py", line 179, in lower
    self.lower_normal_function(self.fndesc)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/lowering.py", line 220, in lower_normal_function
    entry_block_tail = self.lower_function_body()
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/lowering.py", line 245, in lower_function_body
    self.lower_block(block)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/lowering.py", line 260, in lower_block
    self.lower_inst(inst)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/contextlib.py", line 130, in __exit__
    self.gen.throw(type, value, traceback)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/errors.py", line 725, in new_error_context
    six.reraise(type(newerr), newerr, tb)
  File "/home/jetxeberria/documents/programs/anaconda/anaconda3/envs/testing3/lib/python3.7/site-packages/numba/six.py", line 669, in reraise
    raise value
numba.errors.LoweringError: Failed in nopython mode pipeline (step: nopython mode backend)
unsupported operand type(s) for *: 'Macro' and 'Macro'

File "error_minimal_execution.py", line 8:
def foo():
    <source elided>
    a_cache = cuda.shared.array((32, 32), types.int32)
    b_cache = cuda.shared.array((cuda.blockDim.x, cuda.blockDim.y), types.int32)
    ^

[1] During: lowering "$24.20 = call ptx.smem.alloc($24.17, $24.19, func=ptx.smem.alloc, args=[Var($24.17, error_minimal_execution.py (8)), Var($24.19, error_minimal_execution.py (8))], kws=(), vararg=None)" at error_minimal_execution.py (8)

However, running the same code with "NUMBA_ENABLE_CUDASIM=1 python error_minimal_execution.py" raises no error and I can debug normally.

[sklam]

Bug at type inference that should have stopped the use of non-constant value as shared memory size.

[stuartarchibald]

Thanks for the report. This should have been caught earlier at type inference as the shared memory size must be a compile time constant. Marking as a bug.

[gmarkall]

This is a TypeError now (probably since the macro implementation of the CUDA target was replaced with typing/lowering):

$ python repro.py 
Traceback (most recent call last):
  File "repro.py", line 10, in <module>
    foo[1, (32, 32)]()
  File "/home/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 769, in __call__
    return self.dispatcher.call(args, self.griddim, self.blockdim,
  File "/home/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 861, in call
    kernel = self.compile(argtypes)
  File "/home/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 930, in compile
    kernel = compile_kernel(self.py_func, argtypes,
  File "/home/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 57, in compile_kernel
    cres = compile_cuda(pyfunc, types.void, args, debug=debug, inline=inline)
  File "/home/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 40, in compile_cuda
    cres = compiler.compile_extra(typingctx=typingctx,
  File "/home/gmarkall/numbadev/numba/numba/core/compiler.py", line 625, in compile_extra
    return pipeline.compile_extra(func)
  File "/home/gmarkall/numbadev/numba/numba/core/compiler.py", line 361, in compile_extra
    return self._compile_bytecode()
  File "/home/gmarkall/numbadev/numba/numba/core/compiler.py", line 423, in _compile_bytecode
    return self._compile_core()
  File "/home/gmarkall/numbadev/numba/numba/core/compiler.py", line 403, in _compile_core
    raise e
  File "/home/gmarkall/numbadev/numba/numba/core/compiler.py", line 394, in _compile_core
    pm.run(self.state)
  File "/home/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 341, in run
    raise patched_exception
  File "/home/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 332, in run
    self._runPass(idx, pass_inst, state)
  File "/home/gmarkall/numbadev/numba/numba/core/compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 291, in _runPass
    mutated |= check(pss.run_pass, internal_state)
  File "/home/gmarkall/numbadev/numba/numba/core/compiler_machinery.py", line 264, in check
    mangled = func(compiler_state)
  File "/home/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 92, in run_pass
    typemap, return_type, calltypes = type_inference_stage(
  File "/home/gmarkall/numbadev/numba/numba/core/typed_passes.py", line 70, in type_inference_stage
    infer.propagate(raise_errors=raise_errors)
  File "/home/gmarkall/numbadev/numba/numba/core/typeinfer.py", line 1071, in propagate
    raise errors[0]
numba.core.errors.TypingError: Failed in nopython mode pipeline (step: nopython frontend)
No implementation of function Function(<function shared.array at 0x7f67a77248b0>) found for signature:
 
 >>> array(UniTuple(int32 x 2), dtype=class(int32))
 
There are 2 candidate implementations:
  - Of which 2 did not match due to:
  Overload of function 'array': File: numba/cuda/cudadecl.py: Line 44.
    With argument(s): '(UniTuple(int32 x 2), dtype=class(int32))':
   No match.

During: resolving callee type: Function(<function shared.array at 0x7f67a77248b0>)
During: typing of call at repro.py (8)


File "repro.py", line 8:
def foo():
    <source elided>
    a_cache = cuda.shared.array((32, 32), dtype=types.int32)
    b_cache = cuda.shared.array((cuda.blockDim.x, cuda.blockDim.y), dtype=types.int32)