CUDA Cores and Streaming Multiprocessors: Inside the GPU, from Warps to Register Files

What a CUDA core is, and what the count means

NVIDIA's CUDA core count is the number of FP32 arithmetic lanes on the chip. The H100 SXM has 132 SMs with 128 FP32 lanes each, and 132 × 128 = 16,896. Each lane can perform one fused multiply-add per clock, which counts as two floating-point operations, so the FP32 peak is simply lanes × 2 × clock. The datasheet's 67 TFLOP/s implies a clock close to 1.98 GHz; that is a back-calculation, not a quoted specification, and sustained clocks under load vary with power and temperature.

The count is useful for FP32 peak arithmetic and almost nothing else. It says nothing about memory bandwidth, cache sizes, tensor-core throughput or how many threads the chip can keep resident, and those determine the speed of real kernels. It also does not compare cleanly across generations: the A100 has 108 SMs with 64 FP32 lanes each, so the H100's lane count grew by more than the SM count because each SM doubled its FP32 lanes. Treat a CUDA core as a unit of FP32 throughput, and treat the SM as the unit of execution.

Anatomy of an H100 SM

Each H100 SM contains 128 FP32 lanes, 64 INT32 lanes, 64 FP64 lanes and 4 tensor cores, along with load/store units and special-function units for operations such as reciprocal square root and exponentials. These are organized into four processing blocks, each with a warp scheduler and a slice of the register file. The SM holds a 256 KB register file, up to 2,048 resident threads (64 warps), and 256 KB of combined L1 data cache and shared memory, of which up to 228 KB can be configured as shared memory. Outside the SMs sit a 50 MB L2 cache shared by the whole chip and 80 GB of HBM3.

• Per SM: 128 FP32, 64 INT32, 64 FP64, 4 tensor cores; 65,536 32-bit registers; 64 warps or 2,048 threads resident; 256 KB L1 and shared memory.
• Per thread: at most 255 registers. Per block: at most 1,024 threads and 65,536 registers.
• Per chip (SXM): 132 SMs, 16,896 FP32 lanes, 528 tensor cores, 50 MB L2, 3.35 TB/s HBM3.

The PCIe variant of the H100 has fewer SMs enabled, so always read the SM count from the device rather than from a table. For how these levels relate as a memory system, see The GPU Memory Hierarchy.

How a kernel lands on SMs

A kernel launch specifies a grid of thread blocks. The hardware block scheduler assigns each block to an SM with enough free resources, and the block stays on that SM until every one of its threads has finished; blocks never migrate. A block's threads are split into warps of 32 consecutive threads, and each warp is assigned to one of the SM's four processing blocks. Blocks keep being placed until some per-SM limit is reached: warp slots, registers, shared memory or the maximum number of resident blocks. Whatever runs out first sets how many blocks are resident at once.

That is why launch configuration is expressed in terms of the machine. The example below sizes a grid-stride kernel so that every SM gets as many blocks as it can hold, using the occupancy API instead of a hard-coded constant.

// y = a * x + y over n elements. One thread per element is the simplest
// mapping; the grid-stride loop lets a fixed grid cover any n.
__global__ void saxpy(int n, float a, const float* x, float* y) {
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n;
         i += gridDim.x * blockDim.x) {
        y[i] = a * x[i] + y[i];          // one FP32 FMA per element
    }
}

int main() {
    int dev = 0, sms = 0, blocksPerSm = 0;
    cudaGetDevice(&dev);
    cudaDeviceGetAttribute(&sms, cudaDevAttrMultiProcessorCount, dev);

    const int threads = 256;             // 8 warps per block
    cudaOccupancyMaxActiveBlocksPerMultiprocessor(&blocksPerSm, saxpy, threads, 0);

    // Size the grid to the machine, not to n: every SM gets full residency.
    int blocks = sms * blocksPerSm;
    saxpy<<<blocks, threads>>>(n, 2.0f, d_x, d_y);
}

Two consequences follow. A grid smaller than the SM count leaves SMs idle no matter how fast each block runs: a normalization kernel that assigns one block per row and runs on eight rows uses eight of 132 SMs. And a grid that is not a multiple of the chip's resident-block capacity finishes with a partial wave in which many SMs are idle, the staircase effect covered under wave quantization in GPU Tensor Cores.

Warps and the issue slot

Every clock, each processing block's warp scheduler may issue one instruction from one of its eligible warps. An eligible warp is one whose next instruction has its operands ready and its execution unit free. With 32 FP32 lanes per processing block, an FP32 multiply-add for a full warp occupies the lanes for one issue, so a processing block reaches FP32 peak only if it issues an FP32 instruction on every clock, and the whole SM only if all four blocks do.

All 32 threads of a warp execute the same instruction. When threads in a warp take different branches, the warp executes each path in turn with the non-participating threads masked off, so a divergent warp does less useful work per issue. Latency is hidden not by any single warp running fast but by having other eligible warps to issue from while one waits on memory. The detailed taxonomy of why warps become ineligible, and how to read it from the profiler, is in Warp Scheduling.

The register file: 256 KB per SM, allocated per thread

It is easy to overlook, but registers are the fastest storage the GPU has, and there is a lot of it. Each H100 SM has 256 KB of registers, as much as its combined L1 and shared memory, and 132 SMs together hold about 33 MB, roughly two-thirds the size of the 50 MB L2. Registers are private to a thread and are allocated at launch: the compiler fixes a per-thread register count, and the SM reserves that many for every resident thread.

This allocation is the most common link between code and occupancy. At the full 2,048 resident threads, 65,536 registers allow 32 per thread. A kernel compiled to 64 registers per thread can keep at most 1,024 threads, 32 warps, resident per SM; one using 128 registers, typical of tensor-core GEMM kernels holding accumulator tiles, can keep 16 warps. Fewer resident warps mean fewer candidates for each scheduler, so each warp needs more independent work in flight. Going the other way, forcing the register count down can spill values to local memory, which lives in device memory, and make the kernel slower despite higher occupancy. Register Pressure on GPUs walks through the compiler output and the levers.

Why training FLOPs run on tensor cores

The H100 datasheet lists 67 TFLOP/s for FP32 on CUDA cores and about 989 TFLOP/s for dense BF16 on tensor cores, a ratio of roughly 15. Tensor cores execute small matrix multiply-accumulate operations on tiles in a single instruction, feeding many multiplications from each operand fetch, and that is how they reach throughput the scalar lanes cannot.

In a training step, the matmuls in attention and feed-forward layers, forward and backward, are dispatched by cuBLAS, cuDNN and framework kernels to tensor cores, provided the data type allows it. That leaves the CUDA cores with everything else: normalization, activation functions, softmax, residual adds, the optimizer update, loss computation, and the address and index arithmetic inside the GEMM kernels themselves. Those operations are mostly memory-bound, so the CUDA cores rarely limit training throughput. Two practical consequences follow. Running matmuls in plain FP32 leaves most of the chip's arithmetic unused, which is why frameworks offer TF32 and mixed precision. And speeding up the non-matmul portion is about fusion and memory traffic, not more CUDA cores. Tensor Cores: Matmul Accelerators covers the matrix units themselves.

Reading your device

Kernel and launch decisions should use values read from the device at runtime. From PyTorch, the device properties give the name, compute capability, SM count and memory; CUDA C exposes the per-SM limits directly.

import torch

p = torch.cuda.get_device_properties(0)
print(p.name, f"compute capability {p.major}.{p.minor}")
print("SMs:", p.multi_processor_count)
print("memory GB:", round(p.total_memory / 1e9, 1))

# CUDA C exposes more through cudaGetDeviceProperties / cudaDeviceGetAttribute:
#   multiProcessorCount, regsPerMultiprocessor, sharedMemPerMultiprocessor,
#   maxThreadsPerMultiProcessor, warpSize, l2CacheSize

Use the SM count to size persistent and grid-stride kernels, the per-SM register and shared-memory limits when choosing tile sizes, and the compute capability to select code paths, since features such as thread block clusters and newer matrix instructions are architecture-specific. On a GPU partitioned with Multi-Instance GPU, a process sees only the SMs of its slice, which is another reason not to hard-code 132.

Failure modes and misreadings

• Comparing GPUs by CUDA core count. For training throughput the tensor-core rate and memory bandwidth matter far more.
• Hard-coding the SM count. PCIe parts, other SKUs and MIG slices have different counts, and a grid sized for one machine under-fills or over-subscribes another.
• Grids smaller than the machine. Per-row kernels on small batches leave most SMs idle; fuse, batch or split rows across blocks.
• Chasing occupancy with register caps. Spills to local memory cost more than the extra warps gain.
• Divergent branches inside warps. Data-dependent branches on per-thread values serialize the paths; reorganize data so neighboring threads take the same branch.
• FP32 GEMMs by default. Without TF32 or mixed precision, matmuls run on the slower path.
• Trusting nominal clocks. Power and thermal limits lower sustained clocks, so measured peaks sit below datasheet peaks.

Operational guidance and trade-offs

When a kernel is slow, start from the SM rather than from the code. Check whether the grid fills the machine, what limits resident blocks, whether the schedulers have eligible warps, and whether the kernel is limited by an arithmetic pipe or by memory. Nsight Compute reports each of these in its launch statistics, occupancy and scheduler sections, and the answers map directly onto the resources described here.

The design trade-offs are about how to spend per-SM resources. More registers per thread allow larger per-thread tiles and more independent work, at the cost of fewer resident warps. More shared memory per block enables bigger tiles and more data reuse, at the cost of fewer resident blocks. Larger blocks share data more widely but are harder to place and leave coarser tails. None of these has a universal answer; the best kernels choose deliberately, measure, and accept low occupancy when instruction-level parallelism hides latency instead.