CUDA cores and execution model

What is a CUDA core?

Key distinction: A CPU core has its own instruction decoder, branch predictor, L1/L2 cache, and can run independent instruction streams. A CUDA core shares control logic with 31 other cores in its warp (a group of 32 threads that execute in lockstep).

Architecture hierarchy (from bottom up)

CUDA core (scalar ALU): 1 FP32/integer op per cycle
Warp (execution unit): 32 threads executing the same instruction
Streaming Multiprocessor (SM): Contains 64-128 CUDA cores (architecture dependent), shared memory, warp schedulers, registers
GPU chip: 10-100+ SMs (e.g., RTX 4090 has 128 SMs × 128 cores/SM = 16,384 CUDA cores)

SIMT execution model

Derivation: Why 32 threads per warp?

Design constraint: The SM issues one instruction to one warp per cycle. The warp scheduler needs to:

Hide memory latency (100-400 cycles for global memory)
Keep ALUs fed with work
Fit register file constraints

How a warp maps onto cores: As long as the SM has at least as many ALUs as threads in a warp (i.e. ≥ 32), a single warp instruction completes in one cycle; any extra cores simply sit idle for that instruction (they get used by other warps). So:

An SM with 32 cores → one warp instruction per cycle, all cores busy.
An SM with 64 or 128 cores → one warp instruction still completes in one cycle, but the SM can issue multiple warp instructions per cycle (from different warps, via multiple schedulers) to keep the extra cores busy.

The warp size of 32 is a hardware convention (a power of 2 that simplifies address masking and indexing, and matches the SIMD-width lineage of 8/16/32). It is not derived from any single SM's core count.

Register pressure (why warp size interacts with occupancy): The register file per SM is architecture-dependent (commonly 65,536 32-bit registers on many recent NVIDIA SMs, but this varies). Occupancy is limited by total registers: if the SM caps at $M$ registers and you want $P$ resident threads, each thread may use at most about $\frac{M}{P}$ registers. For example, with $M = 65{,}536$ and $P = 2048$ resident threads, that is roughly $\frac{65536}{2048} = 32$ registers/thread. This is a rough ceiling — real allocation is quantized per warp and per architecture, so treat 32 as an illustrative figure, not a universal law.

Thread hierarchy and scheduling

Why blocks? They provide a scalability abstraction. A kernel with 1,000,000 threads organized into 1000 blocks of 1000 threads can run on a GPU with 10 SMs (100 blocks/SM) or 80 SMs (~12 blocks/SM) without code changes.

Warp scheduling mechanism

Each SM has 1-4 warp schedulers. Each cycle, a scheduler:

Selects a ready warp (not stalled on memory/dependency)
Issues its next instruction to the CUDA cores
Moves to next warp next cycle (zero-overhead context switch, since all warps' registers are already loaded)

Latency hiding derivation (done carefully):

Set up the bookkeeping. Suppose:

$S$ = number of warp schedulers on the SM (each issues one warp-instruction per cycle),
a warp does $C$ cycles of useful compute and then issues a long-latency memory op that stalls it for $L$ cycles,
$T$ = number of resident (schedulable) warps.

While one warp is stalled for $L$ cycles, we want the schedulers to have other ready warps to issue. Each cycle the SM can issue $S$ warp-instructions, so over $L$ stall cycles the SM has capacity to issue $S \cdot L$ warp-instructions. To keep the ALUs busy during that window, we need enough independent warps supplying that many instructions. Since each warp supplies about $C$ compute instructions before it too stalls, the number of warps needed to fully cover the latency is:

$T_{\text{required}} \;=\; \frac{S \cdot L}{C}.$

Worked numbers: with $S = 4$ , $L = 200$ , $C = 10$ : $T_{\text{required}} = \frac{4 \times 200}{10} = 80 \text{ warps}.$

But a typical SM caps at ~64 resident warps (2048 threads ÷ 32). Since $80 > 64$ , full latency hiding is impossible here — which is exactly why memory-efficient kernels (coalescing, caching, higher compute intensity $C$ ) matter: raising $C$ lowers $T_{\text{required}}$ .

Memory hierarchy and CUDA cores

CUDA cores access data through a hierarchy:

Registers (per-thread, ~1 cycle latency): Fastest, limited to ~32-64 per thread
Shared memory (per-block, ~30 cycle latency): Explicitly managed, 48-96 KB/SM
L1 cache (per-SM, ~30 cycles): Implicit, combined with shared mem in recent arches
L2 cache (global, ~200 cycles): Shared across all SMs, 6-40 MB
Global memory (DRAM, ~400 cycles): 8-24 GB, high bandwidth (1 TB/s) but high latency

Coalesced memory access: When threads in a warp access consecutive memory addresses (e.g., thread 0 → A[0], thread 1 → A[1], ..., thread 31 → A[31]), the hardware combines them into a single 128-byte transaction. Uncoalesced access (random addresses) → up to 32 separate transactions → up to 32× slower.

Formula Memory transaction efficiency (a fraction, not a bandwidth)

Define transaction efficiency as the ratio of bytes you actually needed to bytes the hardware had to move: $\eta \;=\; \frac{\text{useful bytes requested}}{\text{transactions} \times 128 \text{ bytes/transaction}} \quad (0 < \eta \le 1)$ This is dimensionless (bytes ÷ bytes). Effective bandwidth is then $\eta$ times the peak DRAM bandwidth. Best case (sequential): 32 threads × 4 bytes/float = 128 useful bytes in 1 transaction → $\eta = \frac{128}{1 \times 128} = 100\%$ . Worst case (stride 32): thread $i$ reads A[i*32] → 32 separate 128-byte transactions for 128 useful bytes → $\eta = \frac{128}{32 \times 128} = 3.125\%$ .

Worked example Simple vector addition kernel

__global__ void vectorAdd(float *A, float *B, float *C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}
 
// Launch: <<<numBlocks, threadsPerBlock>>>
int threadsPerBlock = 256;
int numBlocks = (N + threadsPerBlock - 1) / threadsPerBlock;
vectorAdd<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C, N);

What happens under the hood:

Grid launch: GPU distributes blocks across SMs (e.g., N=1,000,000 → 3907 blocks @ 256 threads)
Block → SM mapping: Each SM gets multiple blocks (limited by registers, shared mem)
Warp creation: Each block's 256 threads → 8 warps (256 ÷ 32)
Execution: All 32 threads in warp 0 compute i = blockIdx.x * 256 + [0...31], load A[i] and B[i] in parallel (if memory access is coalesced), add, store C[i]
Warp 1 executes next with threadIdx.x = [32...63], and so on Why this step? The if (i < N) is needed because N might not be divisible by 256. Within a warp, some threads (i ≥ N) will be inactive, but the warp still executes the full instruction path (predication, not branching).

Worked example Warp divergence (the performance killer)

__global__ void divergentKernel(int *data) {
    int i = threadIdx.x;
    if (i % 2 == 0) {
        data[i] = data[i] * 2;      // Even threads
    } else {
        data[i] = data[i] + 100;    // Odd threads
    }
}

What happens:

Threads 0, 2, 4, ..., 30 take the if branch
Threads 1, 3, 5, ..., 31 take the else branch
The warp cannot execute both paths simultaneously (single instruction stream)
Serialization: Warp executes data[i] * 2 with threads 1,3,5,... masked off (idle), then executes data[i] + 100 with threads 0,2,4,... masked off
Effective throughput: Half (or worse with more branches)

Why this step? Hardware implements predication: both paths are fetched, but each thread only commits results for its active path. This doubles execution time for this warp.

Common mistakes

Common mistake "More CUDA cores = faster"

Why it feels right: Spec sheets advertise core counts (RTX 4090: 16,384 cores!). More cores = more parallelism = faster.

Why it's wrong: Performance is limited by:

Memory bandwidth: If your kernel is memory-bound (spending cycles waiting for data), adding cores doesn't help. The bottleneck is DRAM speed (~1 TB/s), not compute (100 TFLOPS).
Warp divergence: Branchy code serializes, wasting cores.
Occupancy: If your kernel uses too many registers, the SM can't fit enough warps to hide latency.

Correct intuition: Core count sets the ceiling for compute-bound workloads (dense matrix multiply, convolutions). For real applications, memory optimization (coalescing, tiling, shared memory) matters more.

Fix: Profile with nvprof or Nsight Compute. If "memory throughput" is near 100%, optimize memory patterns, not just adding cores.

Recall Explain to a 12-year-old

Imagine a factory where you need to paint 10,000 toy cars, each a different color based on its car number.

CPU approach: You hire 8 really smart painters. Each painter can read the instruction sheet, mix the exact paint color, paint one car perfectly, and move to the next. Fast for each car, but only 8 at a time.

GPU (CUDA) approach: You hire 16,384 simple painters, but they work in teams of 32 called "warps." Each warp gets ONE instruction at a time: "Everyone paint your car red!" All 32 painters red simultaneously. Next instruction: "Add blue stripes!" All 32 add blue stripes together.

Problem: What if car #1 needs red and car #2 needs blue? The warp can't split. First, painter #1 paints red while #2 waits (doing nothing). Then #2 paints blue while #1 waits. This is "warp divergence"—half the team is idle, so you're really only using 8 painters effectively, not 16.

The trick: Organize work so all 32 painters in a warp do the same thing. If cars 0-31 all get the same paint pattern (or very similar), the GPU crushes the CPU. That's why GPUs are amazing at graphics (every pixel gets similar shading math) and AI (every neuron does the same multiply-add) but bad at sorting a messy list where each item needs custom logic.

Connections

6.2.01-GPU-vs-CPU-architecture: Why GPUs chose SIMT over MIMD
6.2.02-GPU-memory-hierarchy: How shared memory and coalescing affect CUDA core efficiency
6.2.04-occupancy-and-performance: Maximizing warp occupancy to hide latency
7.3.01-parallel-programming-patterns: Map/reduce patterns that exploit warp execution
9.1.02-neural-network-training-on-GPUs: Why matrix multiply is perfectly suited to CUDA
6.1.03-SIMD-vs-SIMT: Comparing GPU SIMT with CPU SIMD vectorization

#flashcards/hardware

What is a CUDA core and how does it differ from a CPU core?

A CUDA core is a general-purpose ALU (NVIDIA marketing term) that executes one FP32 or integer operation per clock — most notably an FMA. Unlike a CPU core, it lacks independent instruction decode, branch prediction, and control logic—these are shared across 32 cores in a warp. A CPU core runs independent instruction streams; CUDA cores execute in lockstep within their warp.

Does adding more CUDA cores per SM reduce the latency of a single warp instruction?

No. As long as the SM has at least 32 ALUs (one per thread in a warp), a single warp instruction completes in one cycle. Extra cores add throughput (more warps issued per cycle via multiple schedulers), not lower latency for one warp.

What is a warp in CUDA execution model?

A warp is a group of 32 threads that execute the same instruction simultaneously. Threads in a warp share control logic and execute in lockstep (SIMT). If threads diverge (different branch paths), the warp serializes those paths, reducing efficiency.

Derive the global thread ID for a 1D CUDA grid

globalID = blockIdx.x * blockDim.x + threadIdx.x. This maps a thread's position in its block (threadIdx) to a unique global position by offsetting by how many complete blocks precede it (blockIdx * blockDim).

Why is warp divergence a performance problem?

When threads in a warp take different execution paths (if/else), the warp serializes them—executing one path with non-participating threads masked, then the other. This can cut effective throughput to 1/N for N divergent paths, since cores sit idle during paths they don't take.

What is the formula for warps needed to fully hide memory latency, and what does each symbol mean?

T_{\text{required}} = \frac{S \cdot L}{C}

, where

S

= number of warp schedulers,

L

= memory stall cycles,

C

= useful compute cycles per warp between stalls. Raising compute intensity

C

lowers the number of warps needed.

What is memory transaction efficiency, and is it a bandwidth?

It is

\eta = \frac{\text{useful bytes}}{\text{transactions} \times 128}

, a dimensionless fraction (bytes ÷ bytes), not a bandwidth. Effective bandwidth =

\eta \times

peak DRAM bandwidth. Coalesced access →

\eta \approx 1

; strided access can drop

\eta

to ~3%.

How many warps are in a block of 256 threads?

8 warps (256 ÷ 32).

Why is the warp size 32 threads?

It's a hardware convention: a power of 2 (simplifies masking/indexing) matching the SIMD-width lineage of 8/16/32. It is NOT derived from an SM's core count. Warp size also interacts with register-file occupancy limits.

What is SIMT and how does it differ from SIMD?

SIMT (Single Instruction Multiple Thread): one instruction operates on multiple threads, each with its own registers, which can theoretically diverge. SIMD: one instruction on packed data in a single register, no per-lane divergence. SIMT gives the illusion of independent threads while executing them in lockstep.

Concept Map

Hinglish (regional understanding)

Intuition Hinglish mein samjho

Hinglish (regional understanding)

Intuition Hinglish mein samjho

Dekho, sabse pehle ye samajhna zaroori hai ki CUDA core aur CPU core bilkul alag cheezein hain. CPU core ek pura powerful processor hota hai jiska apna brain hota hai — decoder, cache, branch predictor sab kuch. Lekin CUDA core sirf ek chhota sa ALU hai, matlab ek tiny calculator jo per clock cycle mein ek arithmetic operation kar sakta hai (jaise FMA: $c = a \times b + c$ ). Ab GPU mein aise thousands cores hote hain, par ye akele akele kaam nahi karte — ye 32 threads ke groups mein chalte hain jinhe hum warp kehte hain, aur ek hi instruction pure warp par ek saath run hoti hai. Isi ko SIMT model kehte hain, aur isi wajah se GPU data-parallel kaam (jaise matrix multiply, image processing) mein bahut fast hai, par branchy sequential code mein struggle karta hai.

Ab ek important trap jisme students phans jaate hain: log sochte hain "128 cores hain aur warp mein sirf 32 threads, toh warp toh quarter cycle mein khatam ho jayega." Yeh galat hai! Ek warp instruction atomic hoti hai — minimum ek pura cycle lagega, chahe kitne bhi extra cores ho. Extra cores latency kam nahi karte, balki throughput badhate hain, matlab SM ek saath alag-alag warps ki multiple instructions issue kar sakta hai. Isliye 32 ka warp size koi core-count se derived number nahi hai — yeh bas ek hardware convention hai (power of 2 jo addressing aur indexing ko simple banata hai).

Yeh cheez matter kyun karti hai? Kyunki jab tum GPU code likhoge ya optimize karoge, toh tumhe occupancy samajhni padegi — matlab SM par kitne threads ek saath rakh sakte ho. Ismein register pressure ka bada role hai: har SM ke paas limited registers hote hain (jaise 65,536), aur agar tum 2048 threads resident chahte ho, toh har thread ko roughly $\frac{65536}{2048} = 32$ registers milenge. Agar tumhara code zyada registers use karta hai, toh kam threads fit honge, aur GPU apni full parallel power use nahi kar payega. Toh yeh architecture ki samajh seedha tumhare code ki performance decide karti hai — isliye yeh foundational concept hai jo har GPU programmer ko clear hona chahiye.

What is a CUDA core?

Architecture hierarchy (from bottom up)

SIMT execution model

Derivation: Why 32 threads per warp?

Thread hierarchy and scheduling

Warp scheduling mechanism

Memory hierarchy and CUDA cores

Common mistakes

Connections

Concept Map

Hinglish (regional understanding)

Hinglish (regional understanding)

Go deeper — visual, from zero

Test yourself — GPU Architecture

Connections