Exercises — CUDA cores and execution model

A warp contains 32 threads. "In lockstep" means all 32 execute the same instruction on the same cycle — like 32 dancers who are only ever allowed to make the same move at the same time. They differ only in the data each one holds in its own registers, never in which instruction they run at a given tick.

False. A CUDA core is just an ALU (arithmetic logic unit) — a small calculator that does one FP32 or integer operation per cycle. It has no private decoder, branch predictor, or independent instruction stream. It shares all that control logic with the other 31 cores in its warp. That sharing is exactly why GPUs are cheap-per-ALU but bad at branchy code.

1. CUDA core (a scalar ALU, 1 op/cycle)
2. Warp (32 threads in lockstep)
3. Streaming Multiprocessor (SM) (holds many cores + schedulers + registers + shared memory)
4. GPU chip (holds many SMs)

Use the indexing formula $\text{Global thread ID}=\text{blockIdx.x}\times \text{blockDim.x}+\text{threadIdx.x}.$ What each piece means: blockIdx.x is which block (which platoon), blockDim.x is how many threads per block (platoon size), and threadIdx.x is your seat inside your own block. Multiplying block index by block size skips over all the earlier platoons; adding your seat lands you exactly. $=7\times 256+40=1792+40=\boxed{1832}.$

The formula rounds up (that is the floor bracket $\lfloor \cdot \rfloor$ applied to a number nudged past the next whole value) so no element is left uncovered: $\text{numBlocks}=\lfloor \frac{N+\text{threadsPerBlock}-1}{\text{threadsPerBlock}}\rfloor .$ Why the + threadsPerBlock - 1? Plain integer division rounds down, which would drop the last, partly-full block. Adding one-less-than-a-block first forces a round-up whenever there is any remainder — this is exactly a ceiling in disguise. $=\lfloor \frac{1,000,000+255}{256}\rfloor =\lfloor \frac{1,000,255}{256}\rfloor =3907.$ That launches $3907\times 256=1,000,192$ threads — the extra $192$ are switched off by the if (i < N) guard in the kernel.

Warps are minted in chunks of 32, so we round the thread count up with the ceiling bracket: $\text{warps}=\lceil \frac{\text{threads}}{32}\rceil .$

• $256/32=8$ warps exactly. Clean.
• $250/32=7.8\to \lceil 7.8\rceil =8$ warps. The 8th warp is only $250-7\times 32=26$ threads full; the remaining $6$ lanes are inactive but still occupy the warp slot. You pay for a whole warp either way.

The correct minimum is 1 cycle, not a quarter cycle.

Why the student is wrong: a warp instruction is atomic per issue — the smallest slice of time the scheduler deals in is one cycle. Having 128 cores does not chop that cycle into pieces.

What to observe in the figure: the grid shows all 128 cores of the SM on a single tick. The 32 magenta squares are the lanes running this one warp — they light up and finish together on that tick. The 96 pale squares are idle for this warp, but the orange label reminds you they are free to run a different warp on the same tick. So the picture makes the point visual: extra cores widen the SM sideways (more warps at once), they never make this one warp's single tick shorter.

The real benefit of extra cores is throughput, not lower latency: with 128 cores and multiple schedulers, the SM can issue several warp-instructions in the same cycle (from different warps), keeping all 128 cores busy — but each individual warp still takes at least one cycle.

Efficiency is bytes you needed ÷ bytes hardware moved: $\eta =\frac{\text{useful bytes}}{\text{transactions}\times 128}.$ Useful bytes in both cases $=32\times 4=128$.

Case A (coalesced): all 32 consecutive floats span exactly $128$ bytes = one aligned transaction. ${\eta }_A=\frac{128}{1\times 128}=1=100\%.$ Case B (stride 32): consecutive threads land $32\times 4=128$ bytes apart, so each thread's float sits in its own 128-byte line → 32 separate transactions. ${\eta }_B=\frac{128}{32\times 128}=\frac{1}{32}=3.125\%.$ Case B moves $32\times$ more bytes for the same useful data — that is the "$32\times$ slower" you were warned about.

What to observe in the figure: the top row (magenta bracket) is the coalesced case — the 32 requested floats sit side by side and the single magenta box shows they all fit in one 128-byte transaction, so no byte is wasted. The bottom row (orange) is the strided case — only every 8th cell is a float we actually need (orange cells), and each one is wrapped in its own transaction box, so the hardware drags 32 full lines to deliver the same 128 useful bytes. The eye sees instantly why the top wastes nothing and the bottom wastes almost everything.

The ceiling is the total budget split evenly: $\frac{M}{P}=\frac{65,536}{2048}=32\text{ registers/thread}.$ If the kernel truly needs 64 registers/thread, you can only keep $P^{\prime }=\frac{M}{64}=\frac{65,536}{64}=1024\text{ threads}$ resident — half the threads. Fewer resident threads means fewer warps to hide latency behind, so occupancy drops. (This is the exact hinge into 6.2.04-occupancy-and-performance.)

Build the bookkeeping. During the $L=200$ stall cycles of one warp, the SM's $S=4$ schedulers can issue $S\cdot L=4\times 200=800$ warp-instructions of other work. Each warp only supplies $C=10$ instructions before it too stalls, so the number of warps needed to keep the pipeline fed is $T_{\text{required}}=\frac{S\cdot L}{C}=\frac{4\times 200}{10}=80\text{ warps}.$ The SM caps at 64 resident warps. Since $80>64$, full latency hiding is impossible for this kernel on this SM. The ALUs will sit idle part of the time — the kernel is memory-bound.

Set $T_{\text{required}}\le 64$ and solve for $C$: $\frac{S\cdot L}{C}\le 64⟹C\ge \frac{S\cdot L}{64}=\frac{4\times 200}{64}=\frac{800}{64}=\mathrm{12.5.}$ Since $C$ counts whole instructions, you need $C\ge 13$ cycles of compute per stall. Design meaning: by raising compute per memory op — fusing work, reusing values from registers/shared memory (see 6.2.02-GPU-memory-hierarchy), or the tiling patterns in 7.3.01-parallel-programming-patterns — you shrink $T_{\text{required}}$ until 64 warps suffice.

Choose threadsPerBlock = 256 (a multiple of 32) so no warp is partly empty: $256/32=8\text{ full warps per block, no wasted lanes.}$ Blocks needed (ceiling, because a partial block still counts): $\text{numBlocks}=\lceil \frac{1,000,000}{256}\rceil =3907.$ Why not 1000 threads/block? $1000$ is not a multiple of 32: $1000/32=31.25\to \lceil 31.25\rceil =32$ warps, and the last warp runs only $1000-31\times 32=8$ active lanes out of 32 — 24 lanes wasted in every block. Also 1000 is close to the 1024 hard cap, leaving no headroom for register pressure. A round multiple of 32 (128, 256, 512) keeps warps full and occupancy tuning flexible.

Steel-man (the strongest honest version): more ALUs usually means more work per unit time, and $128/32=4$, so it feels like a $4\times$ speedup is on the table.

Refutation: the $4\times$ is real but it applies to throughput, not to a single warp's latency. A warp instruction is atomic per issue — its floor is one cycle regardless of whether 32 or 128 cores exist. On the 32-core SM, one warp/cycle fully occupies the cores. On the 128-core SM, one warp still finishes in one cycle, but now up to 4 warps (via multiple schedulers) can be issued per cycle. So the 128-core SM completes up to $4\times$ more warps per cycle — but any one warp is exactly as fast as before.

Verdict: aggregate warp throughput scales ~$4\times$; per-warp latency is unchanged. This is the difference between finishing one race faster (latency) and running four races side by side (throughput).

(a) Effective bandwidth is efficiency times peak: $\text{eff BW}=\eta \times \text{peak}=0.25\times 1000=250\text{ GB/s}.$ (b) At $\eta =100\%$: $\text{eff BW}=1.00\times 1000=1000\text{ GB/s},\text{factor}=\frac{1000}{250}=4\times .$ The compute per element is unchanged — coalescing only fixes how the bytes move, not how many arithmetic ops each thread does. You made the memory $4\times$ better; the ALUs do the exact same math. This is why memory-bound kernels can leap in speed from a pure access-pattern rewrite, no algorithm change. (Directly relevant to 9.1.02-neural-network-training-on-GPUs, where huge tensor reads dominate.)

Start from $T_{\text{required}}(C)=\frac{S\cdot L}{C}.$ Here $S$ (schedulers) and $L$ (stall cycles) are fixed positive numbers, and $C>0$. As a function of $C$, this is $\text{constant}/C$, which is strictly decreasing for $C>0$: doubling $C$ halves $T_{\text{required}}$. So more compute per memory op always reduces the warps needed. $■$

Degenerate case: $C=0$. If a warp does zero useful compute between stalls, then $T_{\text{required}}=S\cdot L/0$ is undefined (division by zero) — physically, a warp that never computes can never help hide anything; you cannot cover latency with pure stalls no matter how many warps you launch. The formula assumes $C\ge 1$ (at least one real instruction between memory ops).