学习笔记

GPU 与 CUDA 编程基础 — 从硬件到核函数 ​

#GPU · #CUDA · #TensorCore · #显存层次 · #并行计算 · #性能优化

LLM 推理和训练的性能瓶颈几乎都在 GPU 上。理解 GPU 架构和 CUDA 编程模型，是理解 Flash Attention、量化、TensorRT 等优化的前提。本文从 GPU 硬件架构入手，逐步深入到 CUDA 编程实战。

GPU 硬件架构 ​

CPU vs GPU 架构差异 ​

mermaid
graph TD
    subgraph "CPU (Intel i9)"
        CTRL_CPU["复杂控制单元<br/>分支预测/乱序执行"] --> ALU_CPU["少量 ALU<br/>大核心 × 16-24 核"]
        CACHE_CPU["大容量缓存<br/>L1 64KB | L2 2MB | L3 36MB"]
    end

    subgraph "GPU (NVIDIA H100)"
        CTRL_GPU["简单控制单元<br/>Warp Scheduler"] --> ALU_GPU["海量 ALU<br/>小核心 × 18432 核"]
        CACHE_GPU["小容量缓存<br/>L1 256KB | L2 50MB"]
    end

    NOTE["CPU 擅长: 复杂控制流、低延迟单线程<br/>GPU 擅长: 数据并行、高吞吐量计算"]
特性CPUGPU
核心数16-64数千~上万
单核性能极高 (高频 + 超标量)较低 (简化核心)
吞吐量低极高
缓存大 (三级缓存 ~100MB)小 (L2 ~50MB)
控制逻辑复杂 (分支预测/乱序执行)简单 (SIMT)
适用操作系统、Web 服务、数据库矩阵运算、图像处理、深度学习

NVIDIA GPU 演进 ​

GPU年份显存FP16 TFLOPSTensor Core关键特性
V100201732GB HBM2125第 1 代Volta 架构
A100202080GB HBM2e312第 3 代稀疏化 2x
H100202380GB HBM3989第 4 代FP8, Transformer Engine
H2002024141GB HBM3e989第 4 代更大显存
B2002025192GB HBM3e~2250第 5 代Blackwell, FP4

GPU 内存层次 ​

mermaid
graph TD
    subgraph "GPU 内存层次 (延迟从小到大)"
        REG["寄存器 (Register)<br/>每个线程私有<br/>最快 (~0 cycles)<br/>容量: 每 SM 65536 个 32-bit"]
        SMEM["共享内存 (Shared Memory)<br/>线程块内共享<br/>极快 (~20 cycles)<br/>容量: 每 SM 最大 228KB (H100)"]
        L1["L1 缓存<br/>每 SM 私有<br/>快 (~30 cycles)<br/>容量: 256KB/SM"]
        L2["L2 缓存<br/>所有 SM 共享<br/>较慢 (~200 cycles)<br/>容量: 50MB (H100)"]
        GMEM["全局显存 (Global Memory / HBM)<br/>所有 SM 和 CPU 可访问<br/>慢 (~400-800 cycles)<br/>容量: 80-192GB (HBM3e)"]
    end

    REG --> SMEM --> L1 --> L2 --> GMEM

关键洞察：共享内存比全局显存快 20-40 倍。Flash Attention 的核心优化——分块计算中所有中间结果都放在 SRAM（共享内存+L1） 中而非写回 HBM。

内存带宽 ​

GPU带宽带宽利用率 80%
A100 80GB2.0 TB/s1.6 TB/s
H100 80GB3.35 TB/s2.68 TB/s
H200 141GB4.8 TB/s3.84 TB/s
RTX 40901.0 TB/s0.8 TB/s

LLM 推理的 Decode 阶段是 Memory-bound：kernels 等数据的时间远多于计算时间。因此带宽是推理吞吐的天花板。

CUDA 编程模型 ​

线程层级 ​

mermaid
graph TD
    GRID["Grid (网格)<br/>一次 kernel 启动"] --> BLOCK0["Block 0<br/>(0,0)"]
    GRID --> BLOCK1["Block 1<br/>(0,1)"]
    GRID --> BLOCK2["Block 2<br/>(1,0)"]
    GRID --> BLOCK3["Block 3<br/>(1,1)"]

    BLOCK0 --> THR00["Thread (0,0)"]
    BLOCK0 --> THR01["Thread (0,1)"]
    BLOCK0 --> THR_MORE["..."]
    BLOCK0 --> THR15["Thread (3,3)"]

    WARP["32 个线程 = 1 个 Warp<br/>最小调度单元"]
层级标识数量限制通信方式
ThreadthreadIdx每 Block 最多 1024寄存器
Warp—固定 32 线程__shfl_sync
BlockblockIdx1D/2D/3Dshared 内存 + __syncthreads
Grid—1D/2D/3D全局内存 (需 kernel 结束)

第一个 CUDA 程序 ​

cuda
// vector_add.cu — 向量加法
#include <cuda_runtime.h>
#include <stdio.h>

// GPU 核函数 (Kernel)
// __global__ 表示 CPU 调用、GPU 执行
__global__ void vector_add(float *a, float *b, float *c, int n) {
    // 计算全局线程 ID
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    // 边界检查
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    int n = 1 << 20;  // 1M 元素
    size_t bytes = n * sizeof(float);

    // 1. 分配主机内存 (CPU)
    float *h_a = (float*)malloc(bytes);
    float *h_b = (float*)malloc(bytes);
    float *h_c = (float*)malloc(bytes);

    // 初始化数据
    for (int i = 0; i < n; i++) {
        h_a[i] = 1.0f;
        h_b[i] = 2.0f;
    }

    // 2. 分配设备内存 (GPU)
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);

    // 3. 拷贝数据到 GPU
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);

    // 4. 启动 Kernel
    int threads_per_block = 256;
    int blocks_per_grid = (n + threads_per_block - 1) / threads_per_block;
    vector_add<<<blocks_per_grid, threads_per_block>>>(d_a, d_b, d_c, n);

    // 5. 拷贝结果回 CPU
    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);

    // 6. 验证
    float max_error = 0.0f;
    for (int i = 0; i < n; i++) {
        max_error = fmax(max_error, fabs(h_c[i] - 3.0f));
    }
    printf("Max error: %f\n", max_error);

    // 7. 释放内存
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    free(h_a); free(h_b); free(h_c);

    return 0;
}
bash
# 编译运行
nvcc -o vector_add vector_add.cu
./vector_add

线程索引计算 ​

cuda
// 1D Grid + 1D Block
int tid = blockIdx.x * blockDim.x + threadIdx.x;

// 2D Grid + 2D Block (图像处理常用)
int col = blockIdx.x * blockDim.x + threadIdx.x;  // 图像宽度方向
int row = blockIdx.y * blockDim.y + threadIdx.y;  // 图像高度方向

Tensor Core — 矩阵乘法的加速器 ​

工作原理 ​

mermaid
graph TD
    subgraph "Tensor Core: D = A × B + C"
        A["矩阵 A<br/>FP16 (4×4)"] --> TC["Tensor Core<br/>单周期完成<br/>D = A×B + C"]
        B["矩阵 B<br/>FP16 (4×4)"] --> TC
        C["矩阵 C<br/>FP32 (4×4)"] --> TC
        TC --> D["矩阵 D<br/>FP32 (4×4)"]
    end

    NOTE["H100 Tensor Core:<br/>FP16: 989 TFLOPS<br/>FP8:  1979 TFLOPS<br/>FP4:  ~4000 TFLOPS<br/>(稀疏化 2x 后)"]
精度相对速度适用阶段
FP321× (基准)训练中的梯度累积
TF328×训练前向/反向的主要精度
FP16/BF1616×推理 + 混合精度训练
FP832×H100+ 推理加速
INT832×推理量化
INT464×极端推理量化

用 CUDA 调用 Tensor Core ​

cuda
// 通过 warp-level matrix multiply (wmma) 使用 Tensor Core
#include <cuda_fp16.h>
#include <mma.h>
using namespace nvcuda;

__global__ void tensor_core_matmul(
    half *A, half *B, float *C, int M, int N, int K
) {
    // 声明 warp-level 操作的 fragment
    wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::row_major> a_frag;
    wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::col_major> b_frag;
    wmma::fragment<wmma::accumulator, 16, 16, 16, float> c_frag;

    // 初始化累加器
    wmma::fill_fragment(c_frag, 0.0f);

    // 加载并计算
    wmma::load_matrix_sync(a_frag, A, K);
    wmma::load_matrix_sync(b_frag, B, K);
    wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);

    // 存储结果
    wmma::store_matrix_sync(C, c_frag, N, wmma::mem_row_major);
}

共享内存优化实战 ​

Bank Conflict ​

共享内存被组织为 32 个 bank（对应 32 线程的 warp）。同一 warp 内多个线程访问同一 bank 的不同地址时，会产生 bank conflict → 串行化。

cuda
// ❌ 坏: bank conflict — 线程 i 访问 bank[i*2]
__shared__ float smem[128];
int idx = threadIdx.x;
float val = smem[idx * 2];  // threads 0,1 都访问 bank 0 → 冲突!

// ✅ 好: 无 bank conflict — 线程 i 访问 bank[i]
float val = smem[idx];      // threads 0,1 访问 bank 0,1 → 无冲突

矩阵乘法分块 (Tiling) ​

cuda
// 使用共享内存的分块矩阵乘法 (简化版)
#define TILE_SIZE 32

__global__ void matmul_tiled(
    float *A, float *B, float *C, int M, int N, int K
) {
    __shared__ float As[TILE_SIZE][TILE_SIZE];
    __shared__ float Bs[TILE_SIZE][TILE_SIZE];

    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;

    float sum = 0.0f;

    // 遍历所有 tile
    for (int t = 0; t < (K + TILE_SIZE - 1) / TILE_SIZE; t++) {
        // 协作加载 A 和 B 的 tile 到共享内存
        int a_col = t * TILE_SIZE + threadIdx.x;
        As[threadIdx.y][threadIdx.x] = (row < M && a_col < K) ? A[row * K + a_col] : 0.0f;

        int b_row = t * TILE_SIZE + threadIdx.y;
        Bs[threadIdx.y][threadIdx.x] = (b_row < K && col < N) ? B[b_row * N + col] : 0.0f;

        __syncthreads();  // 确保所有线程加载完毕

        // 在共享内存上计算
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }

        __syncthreads();  // 确保所有线程计算完毕再加载下一 tile
    }

    if (row < M && col < N) {
        C[row * N + col] = sum;
    }
}

性能对比 ​

实现方式相对速度瓶颈
朴素实现1×全局内存带宽
共享内存分块10-20×共享内存延迟
cuBLAS50-100×Tensor Core 理论极限

GPU 性能分析工具 ​

nvidia-smi ​

bash
# 实时监控 GPU 利用率、温度、功耗
nvidia-smi dmon -s pucvmet -d 1

# 输出:
# gpu  pwr  gtemp  mtemp   sm  mem  enc  dec  mclk  pclk
#   0  320     65     65   98   85    0    0  9501  1980
#   sm=98% → 计算利用率饱和
#   mem=85% → 显存带宽利用率

Nsight Systems / Compute ​

bash
# 性能分析: 查看 kernel 耗时和内存传输
nsys profile --stats=true ./my_cuda_program

# 深入分析单个 kernel 的瓶颈
ncu --set full -o profile_report ./my_cuda_program
# 输出: 计算吞吐、内存带宽、寄存器占用、bank conflict 次数

PyTorch 中的 GPU 编程技巧 ​

python
"""
PyTorch 中的 GPU 性能最佳实践
"""
import torch

# 1. 避免 CPU-GPU 数据传输
# ❌ 坏: 频繁 .item() 或 .cpu()
for i in range(1000):
    loss = model(batch).item()  # 1000 次 GPU→CPU 传输!
    print(loss)

# ✅ 好: 批量收集后再传回
losses = []
for i in range(1000):
    losses.append(model(batch).detach())
losses = torch.stack(losses).cpu()  # 1 次传输

# 2. 使用 pin_memory 加速数据传输
dataloader = DataLoader(
    dataset,
    batch_size=64,
    pin_memory=True,       # 锁页内存, 加速 CPU→GPU
    num_workers=4,         # 多进程加载
)

# 3. 异步传输 (Non-blocking)
data = data.cuda(non_blocking=True)  # CPU→GPU 异步, 不阻塞 CPU

# 4. CUDA Stream 并行
stream1 = torch.cuda.Stream()
stream2 = torch.cuda.Stream()

with torch.cuda.stream(stream1):
    output1 = model(batch1)  # Stream 1 推理

with torch.cuda.stream(stream2):
    output2 = model(batch2)  # Stream 2 并行推理

torch.cuda.synchronize()  # 等待所有 Stream 完成

GPU 显存管理 ​

python
# 查看显存使用
print(f"已分配: {torch.cuda.memory_allocated() / 1e9:.2f} GB")
print(f"已缓存:  {torch.cuda.memory_reserved() / 1e9:.2f} GB")
print(f"最大分配: {torch.cuda.max_memory_allocated() / 1e9:.2f} GB")

# 清空缓存（释放碎片，不释放正在使用的显存）
torch.cuda.empty_cache()

# 重置统计
torch.cuda.reset_peak_memory_stats()

# 训练后清理
del model
del optimizer
torch.cuda.empty_cache()

常见性能陷阱 ​

陷阱表现原因解决
数据传输瓶颈CPU 利用率高但 GPU 空闲.item()/.cpu() 过于频繁批量操作，减少 CPU-GPU 传输
小 batch 推理吞吐量极低每个 kernel 启动开销 > 计算量动态批处理 (Continuous Batching)
Bank Conflict共享内存延迟异常同一 warp 访问相同 bankpadding 或重排访问模式
Occupancy 不足大量 SM 空闲Block 太少或寄存器/共享内存超限调整 Block 大小, 减少寄存器使用
分支发散warp 内线程走不同分支if/else 导致部分线程闲置重构为无分支逻辑

参考 ​

批注模式

💬 文章评论

暂无评论，来说点什么吧 👇

编程学习笔记

Copyright © 2026