A GPU kernel is an entry point called by the CPU to run the same code across many GPU threads. In an elementwise example, each thread updates one array element; a Numba kernel uses `cuda.jit` and receives its output through a passed array rather than a return value. section 9.1.4 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
This one-dimensional kernel computes each thread’s array position and returns when that position is outside the array, making oversized launch configurations safe. section 9.2.2
Array doubling and profiling
A 1,000-element array can be doubled with 50 blocks of 20 threads and checked with an assertion. In a Mandelbrot comparison, kernel-only timings were 180 ms for Numba and 5.8 ms for CUDA C; profiling also showed that data transfers can dominate total time. section 9.2.2 section 9.3.6
