CUDA Programming Model

TL;DR: A kernel runs many threads organized into blocks and a grid. Host code manages allocations and submits work, often asynchronously. Check launch errors and execution errors, and establish completion before consuming results. Operations in different streams need explicit dependencies when they share data.

The split in one picture

A stream is an ordered queue of GPU work: the host adds copies and kernels to it, and the GPU runs them in that order while the host carries on. In the explicit device-allocation example below, the host (CPU) manages data movement rather than dereferencing a device pointer. 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 key distinction is submission versus completion. A kernel launch is asynchronous with respect to the host; the kernel may or may not have started when the call returns. Check cudaGetLastError() for launch errors and a completion operation for asynchronous execution errors. Consuming a result without the required dependency can read stale or undefined data.

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. This is a residency ceiling, not a minimum grid size for good performance. Fewer blocks may still keep every SM busy; whether more resident warps help depends on the kernel’s latency and resource use.

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

Start with a correct global index and a guard for the final partial block. If the result is wrong, check allocation, initialization, copy direction, launch status and completion status before changing the algorithm. Once correctness is established, use a timeline to determine whether kernel work, transfers or launch overhead dominate.

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 host must establish completion before consuming the result.
• 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: account for the SM count, resource-limited residency and the final partial wave.
• 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.