Question bank — Thread blocks and grids

True. Every thread of a block shares that SM's on-chip shared memory and registers, so the block cannot be split across two SMs — otherwise the shared memory promise would break.

False. An SM holds as many blocks as its budget allows (threads, registers, shared memory), often 4–8 at once. Interleaving them is precisely how the SM hides memory latency.

False. The runtime schedules blocks in any order and may run them concurrently; block 5 can finish before block 1 even starts. The numbering is just an identity, not a timeline.

False. It is a barrier for threads within one block only. There is no built-in whole-grid barrier inside a single ordinary kernel launch (see Thread Synchronization).

False. Larger blocks consume more of the SM's fixed budget, so fewer of them fit; occupancy can drop and hide less latency. Throughput often peaks around 128–256 threads.

False. A warp issues one instruction for all 32 lanes; on divergence the hardware masks off lanes and runs both paths serially, so divergence costs cycles.

False. The barrier is block-scoped. Cross-block waiting inside one kernel risks deadlock because the blocks you're waiting on may not even be resident yet.

True. The dimensionality only shapes the index arithmetic; hardware ultimately flattens everything. A 2D grid just makes image/matrix mapping read naturally.

False. It's blockIdx * blockDim + threadIdx. You must skip over all the complete blocks before you (each of size blockDim) before adding your local offset.

True. Warps are always 32 lanes; a 40-thread block becomes 2 warps (64 lanes) with 24 idle lanes that still occupy scheduling slots.

Error: this only gives the local position inside the block, so every block would re-index elements 0…blockDim-1 and clobber each other. Correct: blockIdx.x * blockDim.x + threadIdx.x.

Error: integer division gives 3 blocks (768 threads) and drops elements 768–999. Use ceiling division (N + B - 1) / B → 4 blocks.

Error: the last block usually has extra threads whose global ID exceeds N; without the guard they read/write out-of-bounds memory. The guard is mandatory whenever N isn't an exact multiple of block size.

Error: blocks aren't guaranteed to co-reside, so the spinning blocks may occupy every SM slot while block 0 never gets scheduled → deadlock. Use separate kernel launches or cooperative groups instead.

Error: plain division truncates the ragged edge tiles. Use (N+15)/16 and (M+15)/16 so partial border tiles still get a block (with in-kernel bounds checks).

Error (by convention): x is the horizontal axis → column, y is vertical → row. Swapping them breaks coalesced access because adjacent threadIdx.x should map to adjacent columns.

Error in judgement: 1000 isn't a multiple of 32 (wastes lanes) and the huge block crowds out others, hurting occupancy. Pick a 32-multiple like 256.

So the runtime can distribute them across any number of SMs in any order — this is what lets identical code scale from a 4-SM laptop chip to a 100-SM datacenter GPU (the core promise of the CUDA Programming Model).

Only threads that are physically co-located on one SM can cheaply share an on-chip scratchpad; a grid spans many SMs, so grid-wide shared memory would require slow global-memory traffic instead.

More independent blocks give the SM more work to swap in whenever one block stalls on memory, so latency stays hidden and the compute units stay busy.

The hardware issues instructions per 32-lane warp, so a block that isn't a multiple of 32 leaves dead lanes in its final warp, wasting scheduling capacity.

Memory is physically 1D and stored row-major, so the flattened index reconstructs the linear address; getting width (not height) into the formula is what keeps a row contiguous.

A block is hardware-capped (e.g., 1024 threads) and must fit on a single SM, so one block could never span a million-element array — the grid exists precisely to tile beyond one SM's reach.

A kernel launch doesn't complete until all its blocks finish, so the next launch is guaranteed to see all results — an implicit whole-grid synchronization without any in-kernel spinning.

An SM must satisfy every resource simultaneously; if each block hogs shared memory, the SM runs out of that budget before it runs out of thread slots, so the tighter constraint wins (see GPU Occupancy).

You still launch 1 block; the if (i < N) guard idles the surplus 206 threads. Wasteful but correct — the guard is what keeps it safe.

Not strictly for correctness in that specific launch, but keep it anyway: it costs almost nothing and protects you the moment N changes to a non-multiple.

(G-1)*B + (B-1) = G*B - 1, i.e. total threads minus one. If that exceeds N-1, the extra threads are exactly the ones the guard must reject.

Both branches execute serially for the single warp: the taken lanes run while the others are masked, then they swap. You pay for both paths even though only one lane needed the else.

No — extra blocks simply queue and are scheduled as earlier blocks retire and free resources. Resource limits cap concurrent residency, not the total number of blocks you may launch.

The kernel does nothing (no thread runs), which is usually a bug from N = 0 slipping through — always ensure ceiling division and a nonzero problem size before launching.