Common operations, declarations, call flows

A from-first-principles tour of the CUDA mechanical model: how the host queues work, how the device executes it, and the exact vocabulary an interviewer expects you to use when you talk about it.

1. The split-execution model

CUDA is two computers cooperating. The host is the CPU and its DRAM. The device is the GPU and its on-board memory (HBM/GDDR). They have separate physical address spaces: a raw pointer the CPU dereferences names host DRAM; a raw pointer the GPU dereferences names device memory. A host pointer dereferenced on the device (or vice versa) is undefined behavior, not a copy.

The relationship is producer/consumer through a queue. The host does not run GPU code; it enqueues work onto a stream, and the device's hardware scheduler consumes that queue asynchronously. This is the single most important fact in the whole API:

Because execution is asynchronous and the address spaces are separate, synchronization is always explicit. There is no automatic "wait for the GPU and copy the answer back." You issue a copy, and you issue a wait. Nothing is implicit except the ordering within a single stream (covered below).

2. A minimal complete program: SAXPY

SAXPY is the "hello world" of dense compute: out = a*x + y for vectors x and y and a scalar a. It is embarrassingly parallel — every output element is independent — so it shows the full host/device call flow with no algorithmic distraction.

First the error-checking macro. Every runtime call returns a cudaError_t; ignoring it is the most common cause of silent corruption. Wrap them.

#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>

// Wrap every runtime call. __FILE__/__LINE__ pinpoint the failing site.
#define CHECK_CUDA(call)                                                  \
    do {                                                                  \
        cudaError_t err_ = (call);                                        \
        if (err_ != cudaSuccess) {                                        \
            fprintf(stderr, "CUDA error %s at %s:%d: %s\n",               \
                    cudaGetErrorName(err_), __FILE__, __LINE__,           \
                    cudaGetErrorString(err_));                            \
            exit(EXIT_FAILURE);                                           \
        }                                                                 \
    } while (0)

The kernel. Note the three linearization answers are visible right here: the owner is i, the bytes touched are x[i], y[i], out[i] (contiguous, one element per thread), and nothing needs synchronizing because threads never touch each other's elements.

__global__ void saxpy(int n, float a,
                      const float* __restrict__ x,
                      const float* __restrict__ y,
                      float* __restrict__ out)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;   // this thread's global index
    if (i < n)                                       // boundary check: grid may overshoot n
        out[i] = a * x[i] + y[i];
}

The host driver. Read it as: allocate on both sides, push inputs to the device, launch, check, wait, pull the result back, free.

int main(void)
{
    const int    n     = 1 << 20;          // ~1M elements
    const size_t bytes = (size_t)n * sizeof(float);
    const float  a     = 2.0f;

    // 1. Host allocation + init.
    float* h_x   = (float*)malloc(bytes);
    float* h_y   = (float*)malloc(bytes);
    float* h_out = (float*)malloc(bytes);
    for (int i = 0; i < n; ++i) { h_x[i] = 1.0f; h_y[i] = 2.0f; }

    // 2. Device allocation.
    float *d_x, *d_y, *d_out;
    CHECK_CUDA(cudaMalloc(&d_x,   bytes));
    CHECK_CUDA(cudaMalloc(&d_y,   bytes));
    CHECK_CUDA(cudaMalloc(&d_out, bytes));

    // 3. Host -> Device copies (inputs).
    CHECK_CUDA(cudaMemcpy(d_x, h_x, bytes, cudaMemcpyHostToDevice));
    CHECK_CUDA(cudaMemcpy(d_y, h_y, bytes, cudaMemcpyHostToDevice));

    // 4. Launch. Cover n with ceil(n / blockSize) blocks of 256 threads.
    int block = 256;
    int grid  = (n + block - 1) / block;
    saxpy<<<grid, block>>>(n, a, d_x, d_y, d_out);

    // 5. Catch launch-config errors immediately (synchronous).
    CHECK_CUDA(cudaGetLastError());
    // 6. Wait for the kernel and surface any runtime (async) error.
    CHECK_CUDA(cudaDeviceSynchronize());

    // 7. Device -> Host copy (result). cudaMemcpy also blocks until done.
    CHECK_CUDA(cudaMemcpy(h_out, d_out, bytes, cudaMemcpyDeviceToHost));

    // 8. Verify + clean up.
    printf("out[0] = %f (expect 4.0)\n", h_out[0]);
    cudaFree(d_x); cudaFree(d_y); cudaFree(d_out);
    free(h_x); free(h_y); free(h_out);
    return 0;
}

3. The call flow, drawn out

4. The execution hierarchy

When you launch a kernel you spawn a grid of threads, organized into a strict three-level hierarchy. Knowing the level each builtin lives at is core interview vocabulary.

• Thread — one instance of the kernel function. It has its own registers and program counter. Identified within its block by threadIdx.
• Warp — 32 consecutive threads that execute in lockstep under SIMT (Single Instruction, Multiple Thread). The hardware issues one instruction to all 32 lanes at once. The warp is the GPU's true unit of scheduling and execution — not the thread.
• Block — a group of threads (up to 1024) that are guaranteed to run on the same Streaming Multiprocessor (SM). Only threads in the same block can share shared memory and synchronize via __syncthreads(). Identified by blockIdx.
• Grid — all the blocks of one launch. Blocks are independent and may run in any order, concurrently or serially; you may not assume an ordering between blocks.

The built-in variables, all of type dim3 (with .x/.y/.z fields):

5. The launch configuration

The triple-angle-bracket syntax takes up to four parameters:

kernel<<< grid, block, sharedBytes, stream >>>(args...);

An int is promoted to a 1D dim3, so saxpy<<<grid, block>>> means dim3(grid,1,1) blocks of dim3(block,1,1) threads. For a 2D launch you build the dim3 explicitly: dim3 b(16,16); dim3 g((W+15)/16,(H+15)/16);

6. Function and variable declarations

CUDA extends C++ with execution-space qualifiers (where a function may be called from and where it runs) and memory-space qualifiers (where a variable lives).

7. Indexing patterns

1D global index with boundary check

The grid usually covers more elements than exist (because block size rarely divides n), so the boundary check is not optional — without it the tail threads write out of bounds.

int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) out[i] = /* ... */;

2D row-major: make the column the fastest-varying axis

For a row-major matrix A[row][col] stored as A[row*width + col], you want consecutive threads (consecutive threadIdx.x, i.e. one warp) to hit consecutive col values so their global addresses are contiguous and coalesce. Therefore map threadIdx.x to the column:

int col = blockIdx.x * blockDim.x + threadIdx.x;   // fastest-varying -> contiguous
int row = blockIdx.y * blockDim.y + threadIdx.y;
if (row < height && col < width)
    out[row * width + col] = in[row * width + col];

Grid-stride loop

Instead of sizing the grid to n, launch a fixed, hardware-sized grid and have each thread stride through the data by the total thread count:

__global__ void saxpy_gs(int n, float a, const float* x, const float* y, float* out)
{
    int stride = blockDim.x * gridDim.x;          // total threads launched
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n; i += stride)
        out[i] = a * x[i] + y[i];
}

Why it exists: it decouples the launch size from the problem size. You pick a grid that just saturates the GPU (e.g. a few blocks per SM) once, and the same binary handles n = 1000 or n = 10^9 without relaunching or overflowing gridDim limits. It keeps every SM busy, makes the kernel trivially correct for any n, and improves data reuse since each thread processes multiple elements. It is the idiomatic default for elementwise kernels.

8. Memory spaces

The performance of a CUDA kernel is dominated by which memory space each access hits. Memorize this hierarchy — it is the most-asked table in interviews.

9. Host/device memory operations

10. Streams

A stream is an ordered queue of device work. The ordering rule is the whole point:

• Within one stream, operations execute in issue order; each finishes before the next begins. This implicit ordering is the only automatic synchronization CUDA gives you.
• Across different streams, operations may overlap and have no ordering guarantee. To impose a cross-stream dependency you must use events (next section).

Putting independent copy/compute/copy chains on multiple streams lets the GPU overlap a transfer on one stream with a kernel on another:

cudaStream_t s;
CHECK_CUDA(cudaStreamCreate(&s));

float* h_pinned;
CHECK_CUDA(cudaMallocHost(&h_pinned, bytes));    // pinned -> async is real
// ... fill h_pinned ...

CHECK_CUDA(cudaMemcpyAsync(d_x, h_pinned, bytes, cudaMemcpyHostToDevice, s));
saxpy<<<grid, block, 0, s>>>(n, a, d_x, d_y, d_out);   // 4th launch arg = stream
CHECK_CUDA(cudaGetLastError());
CHECK_CUDA(cudaMemcpyAsync(h_pinned, d_out, bytes, cudaMemcpyDeviceToHost, s));

CHECK_CUDA(cudaStreamSynchronize(s));            // block host until s is drained
CHECK_CUDA(cudaStreamDestroy(s));
CHECK_CUDA(cudaFreeHost(h_pinned));

The three operations above run in order because they share stream s: the copy-in finishes before the kernel starts, the kernel finishes before the copy-out starts. Split them across two streams and the kernel could read d_x before the copy-in completed — a classic race.

11. Events

An event is a marker you record into a stream. It serves two jobs: expressing cross-stream dependencies, and timing GPU work accurately.

cudaEvent_t start, stop;
CHECK_CUDA(cudaEventCreate(&start));
CHECK_CUDA(cudaEventCreate(&stop));

CHECK_CUDA(cudaEventRecord(start, s));               // mark "now" on stream s
saxpy<<<grid, block, 0, s>>>(n, a, d_x, d_y, d_out);
CHECK_CUDA(cudaEventRecord(stop, s));

CHECK_CUDA(cudaEventSynchronize(stop));              // wait until 'stop' is reached
float ms = 0.0f;
CHECK_CUDA(cudaEventElapsedTime(&ms, start, stop));  // GPU time between the two marks
printf("kernel took %.3f ms\n", ms);

To make stream B wait for work in stream A, record an event in A and call cudaStreamWaitEvent(B, evt, 0). This is the correct, fine-grained alternative to a full cudaDeviceSynchronize(), which would stall everything.

12. Shared memory and __syncthreads()

A neighbor/stencil computation is the canonical reason to use shared memory: each output reads several adjacent inputs, so cooperatively staging a tile into shared memory lets the block read each global element once instead of repeatedly. Here each thread averages itself with its two neighbors, using dynamic shared memory plus a halo of one element on each side.

// Launch with sharedBytes = (blockDim.x + 2) * sizeof(float)
__global__ void stencil3(const float* __restrict__ in, float* __restrict__ out, int n)
{
    extern __shared__ float tile[];          // size set by the 3rd launch param
    int g = blockIdx.x * blockDim.x + threadIdx.x;
    int t = threadIdx.x + 1;                 // +1 leaves room for the left halo

    // Each thread loads its own element into the interior of the tile.
    tile[t] = (g < n) ? in[g] : 0.0f;

    // The two edge threads also load the halo cells.
    if (threadIdx.x == 0) {
        tile[0]                = (g > 0)     ? in[g - 1] : 0.0f;
        int r = g + blockDim.x;
        tile[blockDim.x + 1]   = (r < n)     ? in[r]     : 0.0f;
    }

    __syncthreads();                         // ALL loads must finish before any read

    if (g < n)
        out[g] = (tile[t - 1] + tile[t] + tile[t + 1]) / 3.0f;
}

__syncthreads() is a block-wide barrier: no thread proceeds past it until every thread in the block has reached it. It is what makes "thread A writes tile[...], thread B reads it" safe. Without the barrier, B could read tile before A's store landed.

if (threadIdx.x < 16) {
    __syncthreads();    // WRONG: threads 16..blockDim never reach it -> deadlock
}

13. Warp-level primitives (intro)

Threads in a warp can exchange register values directly, with no shared memory and no barrier, via shuffle and vote intrinsics. The leading argument is an active mask — a 32-bit bitset naming which lanes participate (use 0xffffffff when the whole warp is active). Every lane in the mask must execute the same intrinsic.

• __shfl_down_sync(mask, val, d) — lane i receives lane i+d's val. The workhorse of warp reductions (sum a warp in log2(32)=5 steps).
• __shfl_xor_sync(mask, val, m) — lane i exchanges with lane i ^ m; gives a butterfly pattern for an all-reduce within a warp.
• __ballot_sync(mask, pred) — returns a 32-bit value whose bit i is lane i's pred; the basis for warp-wide counting and stream compaction.

// Sum the 32 values held by one warp; result lands in lane 0.
__device__ float warp_sum(float v) {
    for (int d = 16; d > 0; d >>= 1)
        v += __shfl_down_sync(0xffffffff, v, d);
    return v;
}

These are register-to-register, so they are dramatically faster than the shared-memory equivalent. Lesson 02 builds full reductions and scans on top of them.

14. Atomics

When many threads must update the same address — e.g. building a histogram — an ordinary ++ is a read-modify-write race. atomicAdd makes the update indivisible:

__global__ void histogram(const unsigned char* data, int n, unsigned int* bins)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n)
        atomicAdd(&bins[data[i]], 1u);    // indivisible RMW on bins[...]
}

15. Error handling done right

CUDA errors come in two flavors, and they surface at different times. Getting this wrong is the most common interview red flag.

• Synchronous (launch-configuration) errors are known the instant the launch is issued: invalid block dimensions, requesting more shared memory than the SM has, an absent kernel. They are returned by cudaGetLastError() called immediately after the launch.
• Asynchronous (runtime) errors occur while the kernel executes on the device — an illegal memory access, a misaligned read. They cannot be reported at launch time because the launch already returned. They surface on the next blocking call (cudaDeviceSynchronize(), a synchronous cudaMemcpy, or cudaStreamSynchronize).

auto t0 = now();
kernel<<<g, b>>>(...);     // returns immediately, kernel still queued
auto t1 = now();          // measures ENQUEUE latency, not kernel runtime!

Interview checklist: the 5-step linearization

When asked to read or write a kernel, walk these in order — out loud. It is the linearized thinking lens turned into a procedure.

1. Owner. Compute the global index (blockIdx*blockDim + threadIdx) and state which output element this thread owns. Add the boundary check.
2. Bytes. Name the memory space of every access (register / shared / global / constant) and check that consecutive lanes hit consecutive addresses (coalescing); fix the index mapping if not.
3. Synchronization. For each cross-thread data dependency, place the right barrier: __syncthreads() within a block, stream order or events across streams, cudaDeviceSynchronize() before the host reads results.
4. Launch config. Choose block size (multiple of 32), grid to cover n (or a fixed grid-stride grid), and the dynamic shared-memory bytes; pick the stream.
5. Errors & timing. cudaGetLastError() after the launch for config errors, a sync for sticky runtime errors, and CUDA events — never the CPU clock — for timing.