PY-NUMB0004 · dynamic_kernel_array_shape

Count local or shared arrays whose shape comes from a kernel parameter.

This is a deterministic rule for python. Read its implementation.

Definition

Report cuda.local.array and cuda.shared.array calls whose shape argument is a parameter of the surrounding Numba CUDA kernel. These arrays require a simple compile-time constant shape. A launch-time parameter cannot satisfy that contract.

Evidence

Each finding identifies the allocation call, kernel, and dynamic shape parameter. The value is the number of kernel arrays with parameter-dependent shapes.

Exceptions

Literal shapes and names bound to module constants are accepted. Dynamic shared memory passed through the launch configuration is a different mechanism and does not call shared.array with a parameter shape.

Examples

Bad

@cuda.jit
def reduce(values, width):
    tile = cuda.shared.array(width, float32)

Good

TILE_WIDTH = 256

@cuda.jit
def reduce(values):
    tile = cuda.shared.array(TILE_WIDTH, float32)

References

• Cites “Numba CUDA documentation”, memory management and local memory. Open reference
• Cites “Numba CUDA documentation”, CUDA Kernel API and memory management. Open reference