PY-NUMB0003 · unguarded_grid_index

Count Numba grid indices never bounded by a branch or grid-stride loop.

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

Definition

Report a value assigned from cuda.grid() when no branch or loop in the kernel reads that value. Launch grids normally round up to a whole block, so the extra threads need a bounds check or a grid-stride loop before indexing an array.

Evidence

Each finding identifies the grid index assignment and kernel. The value is the number of grid indices that never participate in bounded control flow.

Exceptions

A kernel intentionally launched over an exact multiple may be safe, but that launch contract is not local to the kernel. A project can retain it with an explicit waiver. A grid index read by any branch or loop is accepted without guessing the comparison operator.

Examples

Bad

position = cuda.grid(1)
output[position] = input[position]

Good

position = cuda.grid(1)
if position < output.size:
    output[position] = input[position]

References

• Cites “Numba CUDA documentation”, writing CUDA kernels and absolute positions. Open reference
• Cites “Numba CUDA documentation”, CUDA Kernel API and forall. Open reference