2 — Inside the silicon

Requirements: for the reading, nothing. For the experiments: Linux, an NVIDIA GPU, the CUDA Toolkit, CMake 3.20+, and Nsight Compute. Any modern card works for the method. The specific numbers below are for H100, so yours will differ.

A CPU and a GPU compared

The last row is the key. A CPU tries to avoid waiting for memory. A GPU accepts the wait and runs other work. SMs issue in order, with no branch prediction or speculative execution (CUDA guide).

The layout

Here’s how an H100 is organized, from the outside in:

+-----------------------------------------------------------------------+
|                         H100 die (SXM5 figures)                       |
|  +-----------------------------------------------------------------+  |
|  | GPC (General Processing Cluster) x 8                            |  |
|  |  +-----------------------------------------------------------+  |  |
|  |  | TPC (Thread Processing Cluster), 66 in total              |  |  |
|  |  |  +--------------------------+ +--------------------------+|  |  |
|  |  |  | SM, 132 in total         | | SM                       ||  |  |
|  |  |  |  +--------------------+  | |                          ||  |  |
|  |  |  |  | SMSP x 4           |  | |                          ||  |  |
|  |  |  |  | - Warp scheduling  |  | |                          ||  |  |
|  |  |  |  | - Dispatch         |  | |                          ||  |  |
|  |  |  |  | - INT32/FP32 cores |  | |                          ||  |  |
|  |  |  |  | - Tensor Core      |  | |                          ||  |  |
|  |  |  |  +--------------------+  | |                          ||  |  |
|  |  |  |  256 KB register file    | |                          ||  |  |
|  |  |  |  256 KB unified L1 and   | |                          ||  |  |
|  |  |  |  shared memory           | |                          ||  |  |
|  |  |  +--------------------------+ +--------------------------+|  |  |
|  |  +-----------------------------------------------------------+  |  |
|  +-----------------------------------------------------------------+  |
|  +-----------------------------------------------------------------+  |
|  |                     L2 cache, 50 MB, partitioned                |  |
|  +-----------------------------------------------------------------+  |
+-----------------------------------^-----------------------------------+
                                    | 5120-bit HBM3 bus, 3.35 TB/s
+-----------------------------------v-----------------------------------+
|                          HBM3 stacks, 80 GB                           |
+-----------------------------------------------------------------------+

An H100 SXM5 has 8 GPCs, 66 TPCs, 132 SMs, 16,896 FP32 CUDA cores and 528 Tensor Cores (Hopper architecture, Profiling Guide). A full GH100 die has 144 SMs and 60 MB of L2, while H100 PCIe has 114 SMs.

Old material often gets these wrong: GPC stands for General Processing Cluster, and current NVIDIA docs expand TPC as Thread Processing Cluster.

Inside an SM

Each SM splits into four sub-partitions (SMSPs). The important parts:

• Warp scheduling and dispatch, which choose work and issue it to the pipelines.
• A 256 KB register file, 64K 32-bit registers per SM, partitioned among the warps (CUDA guide).
• CUDA cores: scalar FP32 and INT32 units.
• Tensor Cores: matrix multiply-accumulate units. They work on tiles such as m16n8k16, or warp-group shapes such as m64nNk16 on Hopper (PTX ISA). They’re pipelined, so sustained throughput is high, but one instruction takes tens to over a hundred cycles (Luo et al.).
• Other execution units, covering sin, cos, reciprocal square root and conversions (Profiling Guide).
• Load/store units.

The register file explains why a GPU can cost so little to switch: the state of a resident warp is already in physical registers, so there’s nothing to save to a stack. Compare that with what a Linux context switch does in the scheduler chapter.

The memory hierarchy

Bandwidths for shared memory and L2 come from FlashAttention-3, and HBM from the H100 datasheet. The latency column was measured on H800 PCIe, which uses HBM2e, so an H100 SXM5 will differ (Luo et al.).

Read the bandwidth column from top to bottom. Most GPU performance engineering is keeping data as high in this table as you can, for as long as you can. Chapter 6 shows that this is the whole story of attention kernels.

The 256 and 228 figures aren’t a contradiction. The SM has a 256 KB unified data cache, and up to 228 KB of it can be configured as shared memory. A single block can use up to 227 KB, and going above 48 KB needs an explicit opt-in (Compute Capabilities).

Two ways to waste it

Bank conflicts. Shared memory is divided into 32 banks, and successive 32-bit words map to successive banks:

$\text{Bank Index} = (\text{Byte Address} \div 4) \pmod{32}$

If threads hit different words in the same bank at once, the hardware serializes those requests. If they read the same 32-bit word, the read is broadcast with no conflict (Profiling Guide). NVIDIA doesn’t publish the cycle penalty, so the effect has to be measured.

Scattered loads. Memory is read in 32-byte sectors, and four sectors make a 128-byte line. Thirty-two threads reading neighbouring floats fit in one line. If each thread reads from a different sector instead, only one 4-byte float in each 32-byte sector is useful, which is 12.5% of what you fetched (Profiling Guide).

The same code can differ by that much depending on array layout, and nothing in the source tells you which one you wrote.

What changed in Blackwell

Keep vendor statements separate from outside measurements here.

NVIDIA states: 208 billion transistors, two dies joined at 10 TB/s, 186 GB of memory at 8 TB/s on GB200 NVL72 or 180 GB at 7.7 TB/s on HGX B200, NVLink at 1.8 TB/s per GPU, a Decompression Engine and a RAS engine (datasheet, marked preliminary). Blackwell adds tcgen05 tensor instructions and FP4 and FP6 support, while keeping the 256 KB and 228 KB figures (Compute Capabilities).

Third parties report: 148 active SMs (Chips and Cheese). Still disputed: L2 size (126 MB measured there, while other sites list 50 or 60 MB) and GPC count (8 or 10). Experiment 1 below can help settle these if you have a B200.

Hands-on

Experiment 1: read the GPU’s inventory

Needs: any NVIDIA GPU.

1. Build NVIDIA’s sample. Set the architecture number for your GPU, such as 90 for Hopper:

   git clone https://github.com/NVIDIA/cuda-samples.git
   cd cuda-samples/cpp/1_Utilities/deviceQuery
   mkdir -p build && cd build
   cmake -DCMAKE_CUDA_ARCHITECTURES=90 ..
   make
   ./deviceQuery

2. Look for: Multiprocessors, CUDA Cores/MP, GPU Max Clock, Memory Bus Width and L2 Cache Size.

3. Expected: an H100 SXM5 should report 132 SMs, 128 cores per SM, a 5120-bit bus and about 50 MB of L2. Multiply the SMs by the cores: you should get 16,896.

If you find older tutorials using bandwidthTest, they’re stale. It’s no longer in cpp/1_Utilities (directory).

Experiment 2: read clocks and power

Needs: any NVIDIA GPU. Only read settings here. Don’t change clocks, ECC or MIG on a shared machine.

1. See which fields your driver supports, then read them:

   nvidia-smi --help-query-gpu
   nvidia-smi -q -d MEMORY,CLOCK,POWER,ECC
   nvidia-smi dmon -s c -c 10
   nvidia-smi topo -m

2. Look for: the gap between maximum and current clocks under load, and any clock event reasons such as a power cap.

3. Run the next experiment while watching clocks. When you convert cycles to time, use the clock you observed, not the rated one, because frequency can fall under load (Luo et al.).

Experiment 3: measure bandwidth and latency

Needs: any NVIDIA GPU.

1. Build and run nvbandwidth:

   git clone https://github.com/NVIDIA/nvbandwidth.git
   cd nvbandwidth && cmake . && make
   ./nvbandwidth -l
   ./nvbandwidth -t host_to_device_memcpy_ce device_to_host_memcpy_ce

2. Don’t be fooled: these tests measure the host link, so a low number doesn’t contradict 3.35 TB/s.

3. For on-chip tiers, use gpu-latency from gpu-benches, a pointer chase.

4. Look for: steps in latency as the working set passes L1 and then L2.

Experiment 4: provoke a bank conflict

Needs: a GPU supported by Nsight Compute.

1. Create bank.cu. Each thread reads shared memory at a stride you choose:

   #include <cstdio>
   #include <chrono>
   #include <cuda_runtime.h>
   
   __global__ void stride(unsigned *out, int s, int rounds) {
       __shared__ unsigned buf[1024];
       int lane = threadIdx.x;
       for (int i = lane; i < 1024; i += 32) buf[i] = i;
       __syncthreads();
       unsigned sum = 0;
       for (int r = 0; r < rounds; ++r)
           sum += buf[(lane * s + r) & 1023];
       out[blockIdx.x * blockDim.x + lane] = sum;
   }
   
   int main() {
       const int blocks = 1024, threads = 32, rounds = 4096;
       unsigned *out;
       cudaMalloc(&out, sizeof(unsigned) * blocks * threads);
       for (int s : {0, 1, 32}) {
           auto t0 = std::chrono::steady_clock::now();
           stride<<<blocks, threads>>>(out, s, rounds);
           cudaDeviceSynchronize();
           std::chrono::duration<double, std::milli> ms =
               std::chrono::steady_clock::now() - t0;
           printf("stride %d: %.3f ms\n", s, ms.count());
       }
       cudaFree(out);
   }

2. Stride 0 reads one word, so it should broadcast. Stride 1 puts each lane in a different bank. Stride 32 puts every lane in the same bank but on different words.

3. Build and run it, then confirm the metric exists on your chip:

   nvcc -O3 bank.cu -o bank
   ./bank
   ncu --query-metrics | grep -i bank
   ncu -o bank --metrics l1tex__data_bank_conflicts_pipe_lsu.sum --section SourceCounters ./bank

4. Expected: stride 32 should show more conflicts than stride 1. Trust the metric over timing, because other limits can hide the effect. If the result surprises you, read the machine code with cuobjdump -sass bank, since the compiler may have changed the loop. Reports show Nsight’s Details and Source pages can disagree on H100 (forum thread), so record your Nsight version.

Experiment 5: profile a strided read

Needs: a GPU supported by Nsight Compute.

1. Write two kernels: one reads a 128-byte-aligned array contiguously, and one reads it with a large stride.

2. Profile each:

   ncu --section SpeedOfLight --section MemoryWorkloadAnalysis -o report ./your_app
   ncu -i report.ncu-rep --page=details

3. Look for: more sectors requested in the strided case. Metric names vary by chip, so find them with ncu --query-metrics.

Experiment 6: check the peak rating by hand

Needs: Experiments 1 and 2.

1. Compute SMs × 4096 × clock, using your SM count and observed clock.
2. Expected: with 132 SMs at 1.83 GHz you get about 989 TFLOPS, the dense FP16 peak. The datasheet’s 1,979 figure includes sparsity, so don’t use that as your target.