Foundations — CUDA cores and execution model

This page assumes you have seen none of the notation in the parent note. We build every term, one at a time, and each new idea leans only on the ones before it. When you finish, re-read CUDA cores and execution model and every symbol will already be familiar.

1. The clock cycle — the heartbeat of everything

Before any core, warp, or memory word, there is time, chopped into equal ticks.

Picture a drummer keeping a steady beat. On each beat, every worker in the factory is allowed to move their hands exactly once.

Why the topic needs it: the parent note counts everything in cycles — "1 op per cycle", "200 cycles for global memory", "one warp instruction per cycle". Cycles are the unit of time for the whole page. You cannot read the latency-hiding formula without this.

Look at the picture: each vertical grid line is one beat. A fast operation fits in one beat; a slow one (a trip to far-away memory) spans hundreds of beats. That gap is the whole reason GPUs are built the way they are.

2. The ALU — the tiny calculator

Picture a single pocket calculator with exactly one button pressed per beat.

The parent note says a CUDA core is just an ALU. This is the single most important de-mystification on the whole page: a "CUDA core" is not a computer, it is a calculator. Once you believe that, "16,384 CUDA cores" stops sounding magical — it just means 16,384 calculators.

Reading $c=a\times b+c$: take the old value of $c$, add to it the product $a\times b$, store the result back in $c$. The $\times$ means multiply; the $+$ means add; the $=$ means "becomes".

3. Thread — one worker following one recipe

Picture a person at a workstation. The recipe (the program) is shared; the notepad (their registers, section 6) is theirs alone.

Why the topic needs it: the whole GPU model is "launch a million threads, let the hardware sort out when each one runs." A thread is the atom of that story.

4. The warp — 32 workers marching in lockstep

Here is the idea that makes a GPU a GPU, not just a big pile of calculators.

Look at the picture: one instruction arrow at the top fans out to 32 ALUs. Each ALU works on a different number (thread 0 on A[0], thread 1 on A[1], …) but they all perform the identical operation.

The number 32 is a hardware convention (a power of 2, easy to index), not something you derive. The parent note is explicit about that.

The formal name for "one instruction, many threads that could in principle diverge" is SIMT (Single Instruction, Multiple Thread). Compare it with plain lockstep SIMD in SIMD vs SIMT — SIMT is SIMD that allows (but penalises) divergence.

5. Blocks, grids, and the indexing formula

A warp is 32 threads; but you launch millions. How are they organised?

Now the parent's key formula. It answers: "I am one specific worker — which slot of the array is mine?"

Let us earn every symbol:

• $\text{threadIdx.x}$ — your seat number inside your own block (0, 1, 2, …).
• $\text{blockDim.x}$ — how many threads are in each block (the block's size).
• $\text{blockIdx.x}$ — which block you belong to (block 0, block 1, …).

Why this exact formula? Imagine seats numbered per row. To get your seat number in the whole theatre, you take (your row number) × (seats per row), then add your seat within the row. That is precisely $\text{blockIdx.x}\times \text{blockDim.x}+\text{threadIdx.x}$.

Look at the picture: block 2, thread 3, with block size 4, gives global ID $2\times 4+3=11$ — the 12th slot counting from zero. The .x just means "the x-direction"; blocks can also be 2D or 3D (adding .y, .z), but 1D is enough to see the idea.

6. Registers — the private notepad

Picture the notepad at each worker's elbow: instant to read, but there are only a few pages.

Why the topic needs it: registers are scarce and shared out among all resident threads. The parent's occupancy estimate $\frac{M}{P}$ is just "total register slots $M$ divided by number of threads $P$ you want present at once". With $M=65,536$ slots and $P=2048$ threads, each thread may use at most $\frac{65536}{2048}=32$ registers. Run out, and fewer threads fit — the topic of occupancy and performance.

7. Latency, throughput, and why they are different

Two words the parent uses constantly, and people constantly confuse.

Picture a highway. Latency = how long your car takes to drive the road. Throughput = how many cars cross the finish line per minute. Adding lanes raises throughput but does not make your single car arrive sooner.

Latency hiding is the trick of filling the wait: while warp A is stalled 400 cycles waiting on memory, the scheduler runs warps B, C, D… so no ALU sits idle. The parent's formula $T_{\text{required}}=\frac{S\cdot L}{C}$ is just "how many spare warps do I need to fill the gap", where $S$ = schedulers, $L$ = stall length in cycles, $C$ = compute cycles a warp does before stalling. This connects straight to the memory hierarchy, which is why the stalls exist.

8. Reading the efficiency symbol $\eta$

The parent introduces $\eta$ (Greek letter "eta") for memory transaction efficiency.

Why a fraction and not a speed? Because it separates how well you used the road from how fast the road is — you multiply $\eta$ by the peak bandwidth to get your real effective bandwidth. This matters for the coalescing story and connects to parallel programming patterns.

Prerequisite map

Everything on the left feeds the parent topic on the right. Notice the clock cycle and the thread are the two true roots — every other idea grows from those two.

Equipment checklist

Test yourself: cover the right side and answer out loud.

When every line above comes easily, you are ready for CUDA cores and execution model.