Chapter 1: Thinking in Threads: The CUDA Model with Numba
Write your first real CUDA kernels in Python with Numba. Understand the thread/block/grid launch hierarchy, how each thread finds its own global index, why every kernel needs a bounds check, and how grid-stride loops make a kernel work for any input size.
Chapter Overview
CUDA is the language GPUs speak, but you do not need C++ to start speaking it. Numba compiles a subset of Python decorated with @cuda.jit straight to GPU machine code (PTX), so you can write genuine CUDA kernels — with threadIdx, blockIdx, shared memory, and atomics — in plain Python.
The mental shift this chapter asks of you is the one that matters most: on a CPU you write a loop that visits each element one at a time; on a GPU you write the body of the loop once and launch thousands of threads that each run that body on a different element, all at the same time.
This chapter builds the foundation everything else rests on:
- Why Numba: real CUDA kernels without leaving Python
- Threads, Blocks & Grids: the launch hierarchy you configure on every kernel
- The Global Index: how a thread figures out which element it owns
- Bounds Checks & Grid-Stride Loops: making a kernel correct for any size, with any launch configuration
Chapter Roadmap
Click any topic to jump in
Why Numba for CUDA
@cuda.jit compiles a Python subset to PTX — real CUDA kernels without a C++ toolchain.
Threads, Blocks & Grids
kernel[blocks, threads](...) configures the launch hierarchy; threadIdx/blockIdx/blockDim/gridDim locate a thread.
The Global Index
i = blockIdx.x*blockDim.x + threadIdx.x (or cuda.grid(1)) maps each thread to one element.
Bounds Checks & Grid-Stride Loops
if i < n guards the ragged tail; grid-stride loops decouple launch size from data size.
Programming an NVIDIA GPU has traditionally meant writing CUDA C/C++, installing the nvcc toolchain, and leaving Python behind to manage kernels, memory, and build steps by hand. That barrier keeps many numerical Python programmers away from the hardware their workloads would benefit from most. Numba removes it: it is a just-in-time compiler that turns a decorated subset of Python directly into GPU machine code, so you can write, launch, and debug real kernels without ever opening a C++ file. This first topic sets up why that matters before you write any parallel code. You will see that the CUDA programming model — threads, blocks, shared memory, atomics — is identical whether you express it in C++ or in Python, so nothing you learn here is throwaway. What changes is only the surface syntax and the speed of iteration. The rest of the chapter builds on this foundation: once you trust that @cuda.jit emits genuine device code, the launch hierarchy, the global index, and bounds checking all follow as the mechanics of telling thousands of threads what to do.
Definition
Numba is a just-in-time compiler that translates a restricted subset of Python and NumPy into optimized machine code. Its cuda submodule targets NVIDIA GPUs, compiling functions decorated with @cuda.jit into PTX device assembly so genuine CUDA kernels can be authored, launched, and debugged entirely from Python.
@cuda.jit: Python In, PTX Out
Numba is a just-in-time compiler. When you decorate a function with @cuda.jit, Numba does not run it as Python — the first time you launch it, Numba compiles that function to PTX (NVIDIA's GPU assembly) and runs it on the device. Inside the kernel you get the real CUDA vocabulary: cuda.threadIdx.x, cuda.blockIdx.x, cuda.shared.array, cuda.syncthreads(), cuda.atomic.add.
The trade-offs versus CUDA C++ are worth knowing. You give up some low-level control and the absolute peak of hand-tuned C++, but you gain the ability to write, launch, and debug a kernel without leaving Python or installing nvcc. The programming model is identical, so everything you learn here transfers directly to CUDA C++.
A kernel is a function executed times in parallel, once per thread. If launching the kernel costs a fixed overhead and each thread does work , the wall-clock time is roughly where is the number of threads running concurrently. Numba's one-time compile adds to on the first call only; amortized over many launches it disappears.
What actually happens the first time you call a @cuda.jit kernel, and why is the second call faster?