CUDA Programming Model

TL;DR: A CUDA program is a CPU that owns memory allocation and control flow, plus a GPU that runs one function body across a grid of thread blocks, each block scheduled to a streaming multiprocessor on its own. Launches, copies and kernels are queued into streams and execute asynchronously, so the host must synchronize before it reads a result, and errors surface only when it does.

The split in one picture

The host (CPU) never computes on device memory directly. It allocates it, copies into it, enqueues kernels that operate on it, and waits. The device runs kernels: one function, written once, executed by every thread with a different index. Everything the host enqueues goes into a stream, and a stream executes its work in the order it was issued.

The line to remember is the note. A kernel launch costs the host a few microseconds and returns before a single thread has started. Reading the output buffer before a synchronize reads whatever was there before, which is usually zeros, and the launch error (a bad grid size, too much shared memory) is reported by the next synchronizing call, not by the launch.

A first kernel that runs

SAXPY, y = a * x + y, is the standard first kernel because it shows every moving part and nothing else. Each thread computes its global index from the block it belongs to and its position inside that block, then handles one element. The if guard matters: the grid is rounded up to whole blocks, so the last block has threads with no element to process.

extern "C" __global__ void saxpy(int n, float a, const float* x, float* y) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) y[i] = a * x[i] + y[i];
}

Host side, in the form from the CUDA Programming Guide:

int n = 1 << 20;
float *dx, *dy;
cudaMalloc(&dx, n * sizeof(float));
cudaMalloc(&dy, n * sizeof(float));
cudaMemcpy(dx, hx, n * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(dy, hy, n * sizeof(float), cudaMemcpyHostToDevice);
int block = 256, grid = (n + block - 1) / block;   // 4096 blocks
saxpy<<<grid, block>>>(n, 2.0f, dx, dy);
cudaError_t e = cudaGetLastError();               // launch-config errors land here
cudaMemcpy(hy, dy, n * sizeof(float), cudaMemcpyDeviceToHost);  // implicit sync

The kernel body was compiled with NVRTC and executed on a consumer Blackwell GPU for this page; it matched a NumPy reference to the last bit. The host block is the standard form and was not compiled here because no nvcc was available in this environment.

Grid, block, SM

A block is the unit of scheduling. When the grid launches, the hardware hands blocks to streaming multiprocessors as capacity frees up, in no promised order. Threads inside a block share a shared-memory allocation and can synchronize with __syncthreads(). Threads in different blocks cannot, because the two blocks may never be resident at the same time. That is the reason a reduction across the whole grid is written as two launches, or as one launch with atomics, or with cooperative groups' grid sync, which requires every block to be co-resident and refuses to launch otherwise.

The numbers make this concrete. An H100 SXM has 132 SMs (NVIDIA H100 architecture whitepaper), and each SM holds at most 2,048 resident threads, which is 64 warps of 32. The SAXPY above at 1M elements and 256 threads per block is 4,096 blocks. With 8 blocks of 256 threads filling an SM, 132 SMs hold 1,056 blocks at once, so the grid runs in roughly four waves. Below about a thousand blocks on this part, some SMs sit idle for part of the launch, which is the "grid too small" failure that shows up as a kernel far under peak bandwidth.

Streams and what actually overlaps

Two consequences drive real code. Overlapping a copy with compute needs pinned host memory (cudaMallocHost) and a non-default stream; a cudaMemcpy from ordinary malloc memory blocks the host and cannot overlap. And PyTorch runs everything on a per-device current stream, so custom kernels launched on a stream you created run concurrently with, and unordered against, the framework's own work unless you record an event to fence them.

What interviewers are listening for

The live exercise is usually "write a kernel that does X, then tell me why it is slow." The screening question inside it is whether you compute the global index correctly and guard the tail. The follow-up held in reserve is "your output is all zeros, what do you check first?" The strong answer is a sequence: cudaGetLastError() right after the launch for configuration errors, then a synchronize and cudaPeekAtLastError() for a fault during execution, then compute-sanitizer for out-of-bounds writes. Candidates who answer "add printf" have not debugged a kernel under time pressure.

A second follow-up separates people who have shipped from people who have read: "how long does an empty launch take, and when does it matter?" A few microseconds per launch is irrelevant for a 40 ms GEMM and dominates a decode step made of hundreds of small kernels, which is exactly why CUDA graphs exist.

Common misconceptions

• "The launch runs the kernel." It enqueues it. The kernel runs when the stream reaches it, and the host has moved on.
• "More blocks always means more parallelism." Beyond the resident capacity of the SMs, extra blocks queue into later waves. Parallelism is bounded by resident warps per SM (see occupancy), and the tail wave is often half empty.
• "Threads in different blocks can wait on each other." Not without grid-wide cooperative launch. A spin-wait on a global flag across blocks deadlocks when the writer block has not been scheduled yet.
• "Unified memory removes the host/device split." It hides the copies. The page migrations still happen, and on a first touch they cost far more than an explicit cudaMemcpyAsync you scheduled yourself.

Key takeaways

• Host allocates and enqueues, device executes; a launch returns in microseconds and the result exists only after a synchronize.
• One block is the scheduling unit and the synchronization boundary. Grid-wide sync needs a second launch or cooperative groups.
• Size the grid in waves against the SM count: on an H100's 132 SMs, fewer than about a thousand 256-thread blocks leaves hardware idle.
• Overlap needs pinned memory, non-default streams and events. Pageable cudaMemcpy blocks the host.
• The first debugging move is cudaGetLastError() after the launch, then a synchronize, then compute-sanitizer.