PY-NUMB0002 · conditional_block_barrier

Count Numba block barriers reached through divergent control flow.

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

Definition

Report cuda.syncthreads() inside a branch or loop in a Numba CUDA kernel. A block barrier is valid only when every thread in the block reaches it. Thread-dependent control flow can leave part of the block waiting forever or produce undefined behavior.

Evidence

Each finding identifies the barrier and its kernel. The value is the number of block barriers nested in conditional or iterative control flow.

Exceptions

Uniform conditions proven from block-invariant values may be safe, but the syntax alone cannot prove that uniformity. Keep such a barrier only with an explicit project waiver and evidence.

Examples

Bad

if cuda.threadIdx.x < active:
    cuda.syncthreads()

Good

value = tile[cuda.threadIdx.x] if cuda.threadIdx.x < active else 0
cuda.syncthreads()

References

• Cites “Numba CUDA documentation”, CUDA Kernel API and synchronization. Open reference
• Cites “CUDA C++ Programming Guide”, synchronization functions. Open reference