Inside a modern NVIDIA GPU: SMs, schedulers, CUDA cores, and tensor cores

Prerequisites

This article assumes you have read CUDA thread hierarchy: grids, blocks, warps, and how blockIdx/threadIdx map threads to data. That article described the programming model. This one describes the hardware underneath it: what a Streaming Multiprocessor (SM) actually contains, how it decides which instruction to run next, and why a block size that is a multiple of 32 is not a stylistic preference but a direct consequence of the silicon.

From the whole chip down to one SM

A GPU die is organized as a hierarchy of clusters, though the exact naming and count of levels has shifted across NVIDIA architectures. At the top, the whole GPU contains several GPCs (Graphics Processing Clusters). Each GPC contains several TPCs (Texture Processing Clusters), and on most recent architectures each TPC contains a pair of SMs (Streaming Multiprocessors). This GPC/TPC framing shows up in NVIDIA’s architecture whitepapers and matters for understanding overall die layout, but as a CUDA programmer you almost never reason about GPCs or TPCs directly. The unit you actually care about, because it is the unit the CUDA programming model maps onto, is the SM.

graph TD
  GPU["GPU die"] --> GPC1["GPC"]
  GPU --> GPC2["GPC"]
  GPU --> GPCn["... more GPCs"]
  GPC1 --> TPC1["TPC"]
  GPC1 --> TPC2["TPC"]
  TPC1 --> SM1["SM"]
  TPC1 --> SM2["SM"]
  SM1 --> L2["Shared L2 cache (device-wide)"]
  SM2 --> L2
  L2 --> DRAM["GDDR/HBM DRAM"]
  style GPU fill:#4a90d9,color:#fff
  style SM1 fill:#e8744f,color:#fff
  style SM2 fill:#e8744f,color:#fff
  style L2 fill:#9b59b6,color:#fff

Every kernel you launch is distributed across SMs at the granularity of a thread block: the hardware scheduler assigns whole blocks to SMs, never splitting a block across two SMs. A GPU with more SMs runs more blocks concurrently, not faster blocks; this is the core reason “more SMs” and “more parallelism” are the same statement in CUDA.

What lives inside one SM

An SM is not one big execution unit. It is itself a small parallel machine with several kinds of functional units, a chunk of fast on-chip memory, and its own instruction schedulers.

graph TD
  SM["Streaming Multiprocessor"] --> WS1["Warp scheduler 0
+ dispatch unit"]
  SM --> WS2["Warp scheduler 1
+ dispatch unit"]
  SM --> WSn["... more scheduler partitions"]
  WS1 --> RF["Register file
(partitioned per scheduler)"]
  WS1 --> CC["CUDA cores
(FP32/INT32 ALUs)"]
  WS1 --> TC["Tensor cores
(matrix multiply-accumulate)"]
  WS1 --> SFU["Special function units
(sin, cos, exp, rsqrt, ...)"]
  WS1 --> LSU["Load/store units"]
  RF --> SHM["Shared memory / L1 cache
(configurable split, on-chip)"]
  SHM --> L2["L2 cache (device-wide)"]
  style SM fill:#4a90d9,color:#fff
  style TC fill:#9b59b6,color:#fff
  style CC fill:#e8744f,color:#fff

• Warp schedulers. An SM is partitioned into a small number of independent scheduler blocks (the exact count and how work is subdivided has changed across generations). Each partition owns a slice of the SM’s warps, registers, and execution units, and issues one (or occasionally more) instruction per warp per cycle from among its resident, ready warps.
• CUDA cores. The basic FP32/INT32 arithmetic lanes. A “CUDA core” is really a single-precision ALU capable of one fused multiply-add per cycle per lane; the marketing number you see on a spec sheet (“N,000 CUDA cores”) is simply the total count of these lanes across every SM on the chip.
• Tensor cores. Dedicated matrix multiply-accumulate units introduced with the Volta architecture and extended in every generation since. A tensor core performs a small matrix multiply-accumulate (the exact operand shapes and supported precisions differ by generation) in far fewer cycles than the equivalent sequence of scalar FMAs on CUDA cores. They are the hardware floating-point performance later in this series calls out specifically, because their reduced-precision inputs with higher-precision accumulation change both throughput and numerical error characteristics.
• Special function units (SFUs). Hardware implementations of transcendental functions (sin, cos, exp2, rsqrt, reciprocal) that would otherwise require many CUDA-core cycles to approximate in software. There are far fewer SFUs than CUDA cores per SM, so a kernel that calls sinf on every element in a hot loop can become SFU-throughput-bound even though it looks “compute-light.”
• Load/store units (LSUs). Compute the memory addresses for a warp’s load or store instruction and issue the resulting memory transaction. This is the unit that determines whether a warp’s accesses coalesce into one wide transaction or fragment into many, the subject of memory coalescing.
• Register file. Fast, per-thread storage, physically partitioned across the SM’s scheduler partitions. Every thread in a resident block draws its registers from this shared pool; a kernel that uses more registers per thread leaves fewer registers available for other concurrently resident blocks.
• Shared memory / L1 cache. A fast, on-chip memory local to the SM, split (in a configurable ratio on most architectures) between programmer-managed shared memory and hardware-managed L1 cache. Covered in depth in CUDA memory hierarchy.
• L2 cache and DRAM. Shared across the whole device, not per-SM. Every SM’s L1/shared memory backs onto the same L2, and every L2 miss goes to device DRAM (GDDR or HBM depending on the product line).

How a warp actually gets issued

A thread block resident on an SM is broken into warps of 32 threads each, exactly as described in the thread hierarchy article. What that article did not cover is what happens after a warp is resident:

1. Every cycle, each warp scheduler partition looks at its resident warps and identifies which ones are eligible: not waiting on a memory operation, not waiting on a __syncthreads() barrier held up by other warps in the block, not waiting on a dependent instruction’s result.
2. The scheduler picks one (or, on architectures with dual-issue capability, up to two independent) eligible warp and issues its next instruction to the appropriate execution unit (CUDA cores, tensor cores, SFU, or LSU).
3. If a warp is not eligible, it simply is not considered that cycle. There is no penalty for a stalled warp beyond the fact that its own progress is delayed; other resident warps are unaffected and can be issued instead.

sequenceDiagram
  participant Sched as Warp scheduler
  participant W0 as Warp 0
  participant W1 as Warp 1
  participant W2 as Warp 2
  participant Mem as Global memory

  Sched->>W0: issue instruction (cycle N)
  W0->>Mem: load request (long latency)
  Sched->>W1: issue instruction (cycle N+1), W0 stalled
  Sched->>W2: issue instruction (cycle N+2), W1 now stalled
  Mem-->>W0: data returns
  Sched->>W0: issue next instruction, W0 eligible again

This is latency hiding through massive multithreading, and it is the single most important idea in this article. A memory load from global memory can take hundreds of cycles to return. A CPU core spends enormous transistor and power budget (out-of-order execution, deep pipelines, branch prediction, huge caches) trying to avoid ever stalling on that latency. An SM does the opposite: it accepts that individual warps will stall constantly, and instead keeps enough other warps resident that the scheduler always has something eligible to issue. The SM does not run faster than a stalled warp’s memory system allows; it simply keeps busy with other work while that warp waits.

This is why occupancy (the fraction of an SM’s maximum warp capacity that is actually resident) matters, and it is why a kernel that keeps every thread constantly busy with independent memory requests in flight tends to hide latency far better than one with only a few warps resident. The full mechanics of tuning for this are the subject of CUDA occupancy and register pressure later in the series; here, the point is only to establish why occupancy is the lever it is.

Block resource allocation: a worked example

An SM has a fixed budget of resources: a maximum number of resident threads, a maximum number of resident blocks, a fixed-size register file, and a fixed amount of shared memory. When you launch a kernel, the hardware computes how many blocks can be simultaneously resident on one SM by checking every one of these limits and taking the most restrictive.

Consider a hypothetical SM (numbers below are illustrative, not a specific chip’s datasheet; always check cudaGetDeviceProperties for your actual hardware) with:

• Maximum 2048 resident threads per SM
• Maximum 32 resident blocks per SM
• 65,536 registers per SM
• 100 KB of shared memory available per SM for this configuration

Suppose your kernel launches blocks of 256 threads, each thread uses 40 registers, and each block requests 12 KB of shared memory. Check every limit independently:

Thread limit:    2048 threads / 256 threads-per-block   = 8 blocks
Register limit:  65536 registers / (256 * 40 registers)  = 65536 / 10240 = 6 blocks
Shared mem limit: 100 KB / 12 KB per block                = 8 blocks (floor)
Block-count limit: 32 (hardware maximum blocks per SM)     = 32 blocks

The binding constraint is the register limit: only 6 blocks (1536 threads, 48 warps) can be resident simultaneously, even though the thread and shared-memory budgets would have allowed 8. Reducing register usage per thread to 32 would relax the register limit to 65536 / (256*32) = 8 blocks, matching the thread and shared-memory ceilings and letting the SM host 2048 resident threads (100% of its thread capacity) instead of 1536 (75%).

This example intentionally stops at “how many blocks fit”; it does not yet ask “how does that residency percentage affect throughput on a real kernel,” because that question depends on the kernel’s memory access pattern and is exactly what occupancy and register pressure covers in depth once you have more tools (profiling, roofline analysis) to answer it precisely.

The programming model is an abstraction, not a promise

It is worth being explicit about what the CUDA programming model does and does not guarantee, because the mapping from “grid of blocks” to “SMs and warps” is intentionally underspecified:

• You do not choose which SM a block runs on. The hardware scheduler decides, and it can differ between runs, between GPUs, and between architecture generations.
• You do not choose the order in which blocks execute, or whether two blocks run concurrently or sequentially on the same SM. Your kernel’s correctness must never depend on block execution order (if it does, you need a redesign, not a workaround).
• Warp size has been 32 threads across every architecture generation this series covers, but treating that as an eternal hardware constant baked into your code (rather than querying warpSize or using cuda::std constants) is fragile, because it is a hardware property, not a language guarantee.
• The specific number of CUDA cores, tensor cores, schedulers, and their exact issue/dual-issue behavior differ by architecture and even by SKU within the same architecture. Code that hardcodes “32 CUDA cores per scheduler” or similar assumptions from one whitepaper will silently become inaccurate documentation (though not incorrect code, since none of these numbers are compiled into your source) on the next GPU generation.

The abstraction exists precisely so that code written against the thread/block/grid model keeps working, and keeps getting faster, as NVIDIA changes the underlying SM count, core count, and scheduling details generation after generation. This is the same reason a CPU program compiled once keeps running (and often keeps getting faster) on newer CPUs with more cores and bigger caches without being rewritten.

How the SM has evolved, at a conceptual level

Rather than list specific core counts per generation (which vary by SKU and go stale quickly), it is more useful to track what capability was introduced or changed at the SM level across major architecture generations:

Every one of these changes preserved the same programming model: __global__ kernels, <<<grid, block>>> launches, threadIdx/blockIdx. What changed underneath was capacity (more of everything) and new optional hardware features (tensor cores, clusters, the Transformer Engine) that a kernel can choose to target explicitly, without breaking kernels that do not.

Concurrent kernels on one GPU

A single SM at any moment executes warps belonging to one or more resident thread blocks, and those blocks can come from different kernel launches, not just the same one, when the launches happen on different CUDA streams and the device has spare resident capacity. This is what makes concurrent kernel execution possible: it is not a separate hardware mode, it is a direct consequence of the same block-scheduling mechanism described above being fed blocks from more than one kernel’s grid at a time. The advanced stream patterns article covers how to structure your host code to make this actually happen (independent streams, no implicit synchronization between them); the point here is that the SM itself has no special “concurrent mode” to enable, because concurrency at this level is just normal block scheduling with more than one grid’s worth of blocks available to schedule.

Connecting this to Nsight

Everything in this article is directly observable, not theoretical. Nsight Compute reports achieved occupancy (resident warps as a fraction of the SM’s maximum), warp state statistics (the fraction of cycles warps spent stalled on each specific cause: memory dependency, execution dependency, barrier wait, and others), and per-unit throughput for CUDA cores, tensor cores, and the LSU, all as direct measurements of the concepts described here. Nsight Systems shows concurrent kernel execution across streams directly on its timeline as overlapping bars. The full profiling workflow, including how to read those reports and act on them, is covered in debugging and profiling CUDA programs; this article’s job was only to make sure the hardware concepts those reports are describing are not a black box when you first open them.

What comes next

You now have a hardware-level model to go with the programming-level model from the previous article: grids and blocks map onto SMs, warps map onto scheduler partitions, and latency hiding through many resident warps is what makes the whole system fast despite individual memory operations being slow.

The next article, warp divergence, uses this exact model to explain what happens when threads within a warp take different branches: since a warp scheduler issues one instruction to all 32 lanes of a warp at once, divergent control flow is not a minor inefficiency but a direct consequence of the SIMT execution unit you just learned about.