Chapter 2: Moving Data: GPU Memory with Numba
The GPU has its own memory, separate from your program's. Learn how to move arrays to the device and back, allocate and mutate device arrays, index 2D grids for matrices, and use fast per-block shared memory with a barrier to let threads cooperate.
- Topics
- 4
- Demo
- 1
- Videos
- 3
- Min read
- 11
Chapter Overview
A kernel can only touch memory that lives on the GPU. Your NumPy arrays live in host (CPU) memory, in a completely separate address space connected to the GPU by the PCIe bus. Before a kernel can run, its inputs must be copied to the device, and after it finishes, the results must be copied back.
Managing that movement deliberately is most of what separates a fast GPU program from a slow one — transfers are expensive, so you want to move data up once, do as much work as possible on the device, and move only the results back.
This chapter covers the memory model and the tools Numba gives you to work with it:
- Host vs Device Memory: two separate worlds, bridged by explicit copies
- to_device & copy_to_host: staging inputs and retrieving results
- 2D Grids for Matrices: indexing rows and columns with cuda.grid(2)
- Shared Memory & __syncthreads(): fast on-chip memory that lets a block's threads cooperate
Chapter Roadmap
Click any topic to jump in
Host vs Device Memory
CPU and GPU have separate address spaces bridged by PCIe; transfers are explicit and costly.
to_device & copy_to_host
Stage inputs with to_device, allocate outputs with device_array, retrieve with copy_to_host.
2D Grids for Matrices
cuda.grid(2) gives each thread a (row, col); 2D blocks tile the matrix; guard both dimensions.
Shared Memory & __syncthreads()
Fast per-block on-chip memory; a barrier coordinates cooperative load-then-read across threads.
Topics
4 topics in this chapter
A GPU is not simply a faster CPU bolted onto your program; it is a separate processor with its own physical memory, sitting across the PCIe bus from the host. That physical separation is the single fact that shapes every GPU program, and this topic makes it concrete before you move a single byte. Data your Python code creates lives in host memory (system RAM); a kernel can only read and write device memory (the GPU's onboard VRAM). Nothing on the GPU can dereference a host pointer, and nothing on the CPU can dereference a device pointer — the two address spaces are genuinely distinct. So the shape of every CUDA workflow is the same three-step rhythm: copy inputs from host to device, launch a kernel that computes on the device, then copy results back from device to host. Those transfers cross the PCIe bus and are often the slowest part of a naive program, which is why understanding the two address spaces up front pays off immediately. The next topics turn this model into concrete Numba calls — to_device, copy_to_host, and the device arrays that live entirely in VRAM.
Definition
The host is the CPU and its system RAM; the device is the GPU and its onboard memory. The two form separate address spaces connected by the PCIe bus: kernels access only device memory, host code accesses only host memory, and data must be explicitly copied across the bus in each direction.
Two Address Spaces, One Bus
Your Python process and its NumPy arrays live in host memory (system RAM). The GPU has its own device memory (VRAM) with far higher bandwidth but no direct access to host RAM. The two are connected by the PCIe bus, which is fast in absolute terms but slow compared to either chip's local memory.
The practical consequence: a kernel's arguments must already be on the device. If you pass a NumPy array, Numba will helpfully copy it up and the result back automatically — convenient, but it hides a transfer on every call. For anything beyond a one-off, you manage the device arrays yourself so you can keep data resident on the GPU across multiple kernels and pay the transfer cost only when you truly must.
If a kernel does FLOPs on bytes and the device runs at FLOP/s while PCIe moves bytes/s, then compute takes and the round-trip transfer takes . The GPU only wins when the time saved elsewhere — for low arithmetic intensity ( small) the transfer term dominates, which is why minimizing transfers matters more than tuning the kernel.
You run a pipeline of five kernels on the same array. Why is it much faster to transfer the array to the device once at the start than to pass a NumPy array to each kernel?
Related Problems on PixelBank
Knowing that host and device memory are separate is the theory; moving data between them is the practice, and Numba keeps it refreshingly simple. This topic covers the two calls that bracket almost every GPU computation. cuda.to_device(array) takes a NumPy array in host RAM and returns a device array — a handle to a copy that now lives in GPU memory and can be passed to a kernel. After the kernel runs, device_array.copy_to_host() pulls the results back into a fresh NumPy array your Python code can use. Between those two calls the data never leaves the GPU, so a chain of kernels can operate on the same device arrays without shuttling intermediates back and forth. That detail matters more than it looks: because each transfer crosses the comparatively slow PCIe bus, the performance difference between copying once and copying repeatedly is often larger than any speedup from the kernel itself. This topic also notes the convenience path — passing a NumPy array straight to a kernel triggers an automatic transfer — and why relying on it silently reintroduces the copies you were trying to avoid.
Definition
cuda.to_device copies a host array into GPU memory and returns a device array; copy_to_host copies a device array back into a new host NumPy array. Explicit device arrays let intermediate results stay on the GPU across successive kernels, avoiding redundant transfers over the PCIe bus.
Staging Inputs and Retrieving Results
cuda.to_device(host_array) allocates device memory and copies the data up, returning a device array you pass to kernels. cuda.device_array(shape, dtype) allocates uninitialized device memory for outputs — no copy, since there is nothing to copy yet. After a kernel runs, device_array.copy_to_host() brings the result back into a NumPy array.
Kernels can also mutate a device array in place (read and write the same buffer), which avoids allocating a separate output. The pattern for a typical kernel is: copy inputs up with to_device, allocate the output with device_array, launch, then copy_to_host the output.
Fill in the host-side code to square an array on the GPU and read the result back.
Related Problems on PixelBank
One-dimensional grids are perfect for flat arrays, but a great deal of GPU work — images, matrices, grids of simulation cells — is naturally two-dimensional, and forcing it into a 1D index is awkward and error-prone. This topic introduces 2D launches, where the block and grid dimensions themselves become tuples and each thread receives a pair of coordinates instead of one. You configure it by passing tuples to the launch, for example kernel[(bx, by), (tx, ty)](...), and inside the kernel row, col = cuda.grid(2) hands each thread its unique position in a 2D space. The mapping is the direct generalization of the 1D case: cuda.grid(2) returns the x-index first, so with row, col = cuda.grid(2) the row comes from the x-dimension of the hierarchy (the first entry of each launch tuple) and the column from the y-dimension, and each still needs its own bounds check against the matrix's height and width. Thinking in two dimensions keeps kernel code readable — A[row, col] instead of hand-rolled offset arithmetic — and matches how the data is actually laid out. It is also the natural setting for the tiling patterns you will meet in the next chapter, where a block cooperates on a square sub-region of a matrix.
Definition
A 2D launch gives the grid and block two-dimensional shapes, so the hierarchy spans rows and columns. Passing tuples such as kernel[(bx, by), (tx, ty)](args) configures it, and cuda.grid(2) returns each thread's (row, col) coordinate — the natural mapping for images and matrices, with a bounds check on each axis.
cuda.grid(2) and 2D Launches
When you launch with 2-tuples — kernel[blocks_per_grid, threads_per_block] where each is (x, y) — the grid becomes two-dimensional. Inside the kernel, i, j = cuda.grid(2) returns the thread's row and column directly. The launch configuration mirrors the 1D case in each dimension: threads_per_block = (16, 16) is a common choice (256 threads), and blocks_per_grid rounds up each matrix dimension.
Every thread still needs its bounds check, now in both dimensions: if i < rows and j < cols. This pattern underlies elementwise matrix ops, transposes, and the per-output-tile structure of matrix multiplication.
A 2D launch tiles the index plane with blocks of size . The grid is blocks, covering threads — the overhang in each dimension is exactly what the two-part bounds check discards.
Launch a kernel that scales an M×N matrix, with 16×16 blocks. What are the grid dimensions and the in-kernel index/guard?
Related Problems on PixelBank
Every thread can reach the GPU's large global memory, but global memory is comparatively slow, and a kernel that re-reads the same values from it again and again leaves most of the hardware's speed on the table. This topic introduces the tool that fixes that: shared memory, a small pool of fast on-chip storage that every thread in a block can see. Because it sits on the Streaming Multiprocessor rather than out in VRAM, shared memory is roughly an order of magnitude faster to access, and a block can use it as a hand-managed cache — load a chunk of data from global memory once, then let all the block's threads read it many times cheaply. That cooperation only works if the threads agree on timing, which is why shared memory comes paired with cuda.syncthreads(), a barrier that makes every thread in the block wait until all of them have arrived. Without the barrier, one thread might read a shared slot before another has written it, a subtle race. Shared memory plus synchronization is the foundation of the high-performance patterns ahead — reductions and tiled matrix multiplication both live or die on using it correctly.
Definition
Shared memory is a small, fast region of on-chip storage private to a single block, declared with cuda.shared.array. It lets a block's threads cooperate by caching global data for repeated reuse. cuda.syncthreads() is the barrier that coordinates them, forcing all threads in the block to reach it before any proceeds.
On-Chip Memory and the Barrier
cuda.shared.array(shape, dtype) allocates a buffer in shared memory, visible to all threads in a block and roughly an order of magnitude faster than global memory. The shape must be a compile-time constant, so you typically define it from a module-level TPB (threads-per-block) constant.
Shared memory only helps if threads cooperate correctly, and that requires a barrier: cuda.syncthreads(). It forces every thread in the block to wait until all of them reach that point, so that values one thread wrote to shared memory are guaranteed visible to the others afterward. The canonical pattern is: each thread loads one element into shared memory → cuda.syncthreads() → every thread now safely reads its neighbors' values. Forgetting the barrier is the most common shared-memory bug — threads race and read stale or uninitialized data.
Shared memory turns repeated global reads into one. If a block of threads each needs neighboring values, naive code does global reads; staging the block in shared memory does global reads plus shared reads. Since shared bandwidth global bandwidth, the effective cost drops from slow accesses to — the reuse factor is the speedup.
Why does reversing a block of data through shared memory need a cuda.syncthreads() between the load and the read?
Theory Exercise
Problem:
A kernel stages a block's data into a shared array, then each thread reads a value its neighbor wrote — but the developer omitted cuda.syncthreads(). The kernel returns correct results on tiny inputs during testing but produces wrong, non-deterministic answers in production. Explain the bug and why testing missed it.
Hints:
- What guarantees that a value written by one thread is visible to another?
- Do all 32 threads in a warp execute in lockstep, and what about threads in different warps of the same block?
- How could a very small block accidentally hide the problem?