CUDA programming model basics: Question bank — Hardware deep dive

Every term used here (host, device, kernel, thread, block, grid, warp, coalescing, boundary check) is built in the parent note — if a word feels unfamiliar, define it there first. Related depth: Thread Warps and SIMT, Memory Hierarchy, Streaming Multiprocessors.

Pictures to lean on

Three traps below are geometric — keep these open while you answer.

The three-level hierarchy and how a 2D global index is built (used in the "2D indexing" traps):

Why warps interleave on one SM — the latency-hiding timeline (used in the "warp scheduling" traps):

The full memory-space map — registers, shared, constant, texture, local, global (used in the "which memory" traps):

True or false — justify

A kernel launch runs the code once, and the GPU loops it internally.

False. A launch spawns one thread per index that all run the kernel body simultaneously; there is no internal loop — you replaced the loop with thousands of parallel copies.

Every thread in a grid can share data through shared memory.

False. Shared memory is per-block; only threads inside the same block see it. Cross-block sharing must go through slow global memory.

threadIdx.x is unique across the whole grid.

False. It is only unique within a block (0 to blockDim.x-1); every block restarts it at 0. The globally unique id is blockIdx.x * blockDim.x + threadIdx.x.

For a 2D grid, threadIdx.x alone tells me my row in the image.

False. In 2D you need both axes: row = blockIdx.y*blockDim.y + threadIdx.y and col = blockIdx.x*blockDim.x + threadIdx.x — see the left panel of figure s01.

A 2D block is a genuinely different hardware thing from a 1D block.

False. dim3 block/grid shapes are purely a labelling convenience; the hardware still flattens them into warps of 32. 2D/3D just makes image and volume indexing readable.

A block with 256 threads always runs all 256 at literally the same instant.

False. Threads execute in warps of 32; a 256-thread block is 8 warps that the SM schedules, possibly interleaved (figure s02), not one giant simultaneous step.

Interleaving warps means the GPU is slower because it keeps switching.

False. Switching is free on an SM and is exactly how it hides memory latency: while one warp waits ~300 cycles for global memory, another warp computes — figure s02 shows the stalls being filled.

If I launch fewer threads than elements, the leftover elements just get skipped silently and safely.

True but dangerous. They are skipped, but "safely" only if you meant to skip them; usually it means you under-launched and your output array has stale/garbage entries — unless you use a grid-stride loop (see below).

cudaMalloc gives you a pointer the CPU can dereference directly.

False. It returns a device pointer into VRAM; dereferencing it on the host is undefined behaviour — data must cross the PCIe bus via cudaMemcpy first.

Registers are the smallest and slowest memory in CUDA.

False. Registers are the fastest (~1 cycle) — smallest in scope (per-thread), but speed and scope are different axes. The slowest reachable memory is host RAM across PCIe.

Constant memory is faster than global memory for every access pattern.

False. Constant memory is only fast when all threads in a warp read the same address (it broadcasts from a small cache); divergent addresses serialize and lose the benefit — figure s03.

"Local memory" lives close to the thread like a register.

False. Despite the name, local memory is off-chip global memory used for register spills and thread-private arrays; it is as slow as global. The name refers to scope, not location — figure s03.

Using more threads per block is always faster because more parallelism.

False. Blocks cap at 1024 threads and each block competes for the SM's registers/shared memory; too-large blocks reduce how many blocks fit per SM, hurting occupancy.

Spot the error

int idx = threadIdx.x; used as the global index in a multi-block launch.

Wrong — this ignores blockIdx.x. Every block would write indices 0..blockDim-1, so all blocks stomp the same elements and most of the array is never touched.

For a 2D image kernel, int idx = blockIdx.x*blockDim.x + threadIdx.x; used as the pixel index.

Wrong — that only covers columns of one row. The flattened pixel index is row*width + col, using both the x and y global indices (figure s01, right panel).

cudaMemcpy(h_C, d_C, bytes, cudaMemcpyHostToDevice); after the kernel.

Wrong direction. Copying results back is GPU→CPU, so it must be cudaMemcpyDeviceToHost; as written it overwrites the GPU result with host garbage.

numBlocks = N / threadsPerBlock; for N = 1000, threadsPerBlock = 256.

Wrong — integer division gives 3 blocks = 768 threads, so elements 768–999 are never computed. Use ceiling division (N + tpb - 1) / tpb.

Kernel body: A[idx] = idx; with no if (idx < N).

Wrong — the last (over-launched) threads write past the end of A, causing out-of-bounds memory corruption or a crash.

A grid-stride loop written as for (int i = idx; i < N; i++) to cover a huge array.

Wrong — every thread would walk to the end and re-process the same tail. The stride must be the total number of threads: for (int i = idx; i < N; i += blockDim.x*gridDim.x).

vectorAdd<<<numBlocks, threadsPerBlock>>>(h_A, h_B, h_C, N); passing host pointers.

Wrong — the kernel runs on the device and can only touch device memory; you must pass d_A, d_B, d_C, not the host arrays.

Freeing device memory with free(d_A);.

Wrong — free is for host allocations from malloc; device memory from cudaMalloc must be released with cudaFree.

Why questions

Why does CUDA split work into blocks instead of one flat pool of threads?

Because blocks map to SMs and give a cooperation boundary (shared memory + barriers) while letting the same code scale on GPUs with different SM counts — the hardware chooses how many blocks run at once.

Why use a grid-stride loop when a plain one-thread-per-element launch already works?

It decouples the kernel from N: a fixed grid of, say, 4096 threads can process an array of any size by each thread hopping forward by blockDim.x*gridDim.x, so the same launch handles arrays bigger than the GPU's thread count.

Why do warps interleave on one SM instead of finishing one before starting the next?

To hide latency (figure s02): a warp stalled ~300 cycles on global memory would idle the compute units, so the scheduler swaps in a ready warp — many resident warps keep the ALUs busy.

Why offer constant and texture memory at all when global memory can hold anything?

They are specialized caches: constant memory broadcasts one value to a whole warp cheaply, and texture memory is optimized for 2D/3D spatially-local reads with hardware interpolation — both beat global memory for their pattern (figure s03).

Why is the boundary check if (idx < N) needed if I computed numBlocks correctly?

Ceiling division over-launches: numBlocks * threadsPerBlock is usually larger than N, so a few extra threads exist and must be told to do nothing.

Why copy data to the device at all — why not let the GPU read system RAM?

The GPU has a physically separate memory space (VRAM); reaching host RAM crosses the slow PCIe bus, so CUDA requires an explicit cudaMemcpy to stage data in fast device memory.

Why prefer 256 threads per block rather than, say, 30?

256 is a multiple of the warp size 32, so no warp runs half-empty; a block of 30 wastes 2 of 32 lanes in its single partial warp on every step.

Why does moving data cost so much more than doing math?

Global memory sits off-chip (~cm of wire + a memory controller), so an access is ~200–400 cycles versus ~1 for an on-chip arithmetic op — the whole memory hierarchy exists to hide this gap.

Why does thread cooperation only exist within a block, not across blocks?

Different blocks may run at different times on different SMs (or not concurrently at all), so there is no guaranteed shared clock or memory to synchronize on — only intra-block barriers are safe.

Edge cases

What happens when N is smaller than threadsPerBlock (e.g. N = 10, tpb = 256)?

numBlocks becomes 1, launching 256 threads; the boundary check lets only threads 0–9 do work and the other 246 sit idle — correct, just underutilized.

What if N is exactly divisible by threadsPerBlock?

No threads are over-launched, so the boundary check never fires — but you still keep it, because relying on perfect divisibility breaks the moment N changes.

What if N = 0?

Ceiling division gives numBlocks = 0, so the kernel launches with an empty grid and simply does nothing — a valid no-op, not an error.

What if a grid-stride loop is launched with more threads than N?

Threads whose starting idx >= N never enter the loop body — the i < N condition is the built-in boundary check, so grid-stride loops are safe at any launch size.

What if a 2D block is 32×32 = 1024 threads?

That is exactly the per-block cap; legal, but leaves no room for even one more thread, and register/shared-memory pressure may still force fewer such blocks per SM.

What if two threads write to the same global address without coordination?

You get a race condition; the final value is whichever write lands last (undefined), which is why such patterns need atomics or a redesigned access map — see Parallel Algorithm Design.

What if a warp's 32 threads take different branches of an if?

The warp executes both paths in sequence with inactive lanes masked off — this branch divergence serializes the branches and costs performance (detail in Thread Warps and SIMT).

If I forget to check the return code of cudaMalloc and it fails, what do I see?

d_A stays an invalid pointer and the kernel silently corrupts memory or crashes later — the failure surfaces far from its cause, which is why every CUDA call should be error-checked.

Recall Self-test before you leave

Global unique thread id formula (1D)? ::: blockIdx.x * blockDim.x + threadIdx.x. 2D pixel index into a width-wide image? ::: (blockIdx.y*blockDim.y+threadIdx.y)*width + (blockIdx.x*blockDim.x+threadIdx.x). Grid-stride loop stride value? ::: blockDim.x * gridDim.x — the total number of threads launched. Which copy direction loads inputs onto the GPU? ::: cudaMemcpyHostToDevice. Why does the boundary check exist at all? ::: Ceiling division over-launches threads; the check idles the extras so they don't write out of bounds.