第16章 §04 · GPU 架构与 CUDA · 讲解稿¶

篇幅: ~4000 字 / 阅读 10 分钟

一 · 一句话总览¶

GPU = 数据并行的"屠龙刀"。NVIDIA H100 1979 TFLOPS FP16, 是 AI 训练 / 推理的事实标准。CUDA 是 GPU 编程的"通用语"。

二 · GPU vs CPU 哲学¶

2.1 核心差异¶

维度	CPU	GPU
核心数	4-64 大核	1000-10000 小核
延迟	极低	中高
吞吐	中	极高
适用	串行 / 分支	数据并行
能耗	100W	300-700W

2.2 SIMT 模型¶

SIMT = Single Instruction Multiple Threads: - 32 个线程组成 warp - 同 warp 同步执行同指令 - 分支时 warp 串行 (divergence)

三 · NVIDIA GPU 架构¶

3.1 历代架构¶

架构	年份	代表	CUDA 核心	关键改进
Volta	2017	V100	5120	Tensor Core
Turing	2018	T4	2560	RT Core
Ampere	2020	A100	6912	TF32 + MIG
Hopper	2022	H100	16896	FP8 Transformer
Blackwell	2024	B200	~20000	4-bit + Dual GPU

3.2 SM 内部结构¶

SM (Streaming Multiprocessor) = GPU 核心计算单元: - 128 个 FP32 核心 (Hopper) - 64 个 INT32 核心 - 4 个 Tensor Core (矩阵) - 256 KB 寄存器 - 256 KB 共享内存 + L1

3.3 内存层次¶

Register (per thread)        1 cycle,  256 KB/SM
Shared Memory (per block)    ~30 cycle, 256 KB/SM
L1 Cache                     ~30 cycle, 256 KB/SM (含 Shared)
L2 Cache                     ~200 cycle, 50 MB
HBM (global)                 ~500 cycle, 80 GB (H100)
Host Memory (PCIe)           ~μs, TB

3.4 H100 vs A100 vs B200¶

指标	A100	H100	B200
FP16 TFLOPS	312	1979	4500
FP8 TFLOPS	-	3958	9000
HBM	80 GB	80 GB	192 GB
带宽	2 TB/s	3.35 TB/s	8 TB/s
功耗	400W	700W	1000W

四 · CUDA 编程模型¶

4.1 核心概念¶

Thread: 最小执行单位
Block: 多个 thread (最多 1024), 共享内存
Grid: 多个 block, 跨 SM 调度

Grid (N blocks)
├── Block 0 (256 threads)
│   ├── Thread 0
│   ├── Thread 1
│   └── ...
├── Block 1
└── ...

4.2 索引与并行¶

__global__ void vec_add(float *a, float *b, float *c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) c[idx] = a[idx] + b[idx];
}

// 启动: 1M 个元素, 256 threads/block
int threads = 256;
int blocks = (n + threads - 1) / threads;
vec_add<<<blocks, threads>>>(a, b, c, n);

4.3 内存合并访问¶

线程 0-31 访问 addr[0..127] → 1 个 128B 事务 ✓ (合并)
线程 0-31 访问 addr[0, 32, 64, ..., 992] → 32 个事务 ✗ (不合并)

4.4 共享内存 + 同步¶

__global__ void reduce_sum(const float *in, float *out, int n) {
    __shared__ float smem[256];
    int tid = threadIdx.x;
    float sum = 0;
    for (int i = tid; i < n; i += blockDim.x) sum += in[i];
    smem[tid] = sum;
    __syncthreads();

    // 树形归约
    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) smem[tid] += smem[tid + s];
        __syncthreads();
    }
    if (tid == 0) out[blockIdx.x] = smem[0];
}

五 · Tensor Core¶

5.1 矩阵乘加速¶

Tensor Core = 专用矩阵乘硬件: - 单指令完成 4×4×4 矩阵乘 - FP16 / BF16 / TF32 / FP8 / INT8

5.2 WMMA 编程¶

#include <mma.h>
using namespace nvcuda::wmma;

__global__ void matmul_tensor_core(const half *A, const half *B, float *C, int M, int N, int K) {
    fragment<matrix_a, 16, 16, 16, half, row_major> a_frag;
    fragment<matrix_b, 16, 16, 16, half, col_major> b_frag;
    fragment<accumulator, 16, 16, 16, float> c_frag;

    fill_fragment(c_frag, 0.0f);

    for (int k = 0; k < K; k += 16) {
        load_matrix_sync(a_frag, A + ..., 16);
        load_matrix_sync(b_frag, B + ..., 16);
        mma_sync(c_frag, a_frag, b_frag, c_frag);
    }

    store_matrix_sync(C + ..., c_frag, 16, mem_row_major);
}

5.3 FP8 Transformer Engine (H100+)¶

动态 FP8 转换
比 FP16 快 2×, 内存减半
Transformer 训练专用

六 · 实战案例¶

6.1 GEMM (矩阵乘) 优化¶

// 128x128 tile, 16x16 warp
template<int BM, int BN, int BK, int TM, int TN>
__global__ void gemm_kernel(const float *A, const float *B, float *C, int M, int N, int K) {
    __shared__ float As[BM][BK], Bs[BK][BN];

    int bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;

    float c[TM][TN] = {0};

    for (int k = 0; k < K; k += BK) {
        // 加载到 shared memory (合并)
        As[ty][tx] = A[(by * BM + ty) * K + k + tx];
        Bs[ty][tx] = B[(k + ty) * N + bx * BN + tx];
        __syncthreads();

        // 计算
        for (int kk = 0; kk < BK; kk++) {
            for (int i = 0; i < TM; i++)
                for (int j = 0; j < TN; j++)
                    c[i][j] += As[ty * TM + i][kk] * Bs[kk][tx * TN + j];
        }
        __syncthreads();
    }

    for (int i = 0; i < TM; i++)
        for (int j = 0; j < TN; j++)
            C[(by * BM + ty * TM + i) * N + bx * BN + tx * TN + j] = c[i][j];
}

6.2 vLLM 的 Flash Attention + PagedAttention¶

Flash Attention: 利用 shared memory 减少 HBM 访问
PagedAttention: 类比 OS 虚拟内存, block table 映射 KV cache
A100 跑 Llama 2 7B: 100+ tokens/s (单 GPU)

七 · 一句话总结¶

CUDA + Tensor Core = GPU 编程的事实标准。H100 1979 TFLOPS FP16, FP8 Transformer Engine 翻倍, 让千亿模型训练从月降到周。未来 Blackwell + FP4 将再翻倍。