A boundary check tests whether an index lies within an array’s extent before accessing it. This matters in GPU kernels, where Python’s usual bounds checks are absent and invalid accesses can cause memory errors or silent bugs. section 9.2.2
The safety and cost of a boundary check depend on the implementation:
- A CUDA kernel can launch more threads than there are array elements; checking each thread’s position against the extent prevents excess threads from indexing beyond the array. section 9.2.2
- Cython memoryviews provide bounds checking, but disabling it removes a guard rail and can let invalid indices crash the program. section 5.4.2
Guarding an overprovisioned GPU launch
A 1,000-element array launched with 63 blocks of 16 threads has 1,008 thread positions. Return when a position is greater than or equal to the array length before accessing it. section 9.2.2
@cuda.jit
def double_safe_explicit(my_array):
position = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.x
if position >= my_array.shape[0]:
return
my_array[position] *= 2
See the explicit CUDA position calculation and extent check in a safe doubling kernel. section 9.2.2
