Skip to content
Gains Summary
Main Navigation 首页 / Home
C++ 编程 / C++ Programming
系统与高性能 / Systems & Performance
Web 开发 / Web Development
人工智能 / Artificial Intelligence
工业软件 / Industrial Software
其他内容 / Other Topics
C++ 编程 / C++系统与性能 / SystemsWeb 开发 / Web人工智能 / AI工业软件 / Industrial

外观

本页目录

CUDA 内存管理——百亿参数如何装进显存 / CUDA Memory Management for Large Models ​

📅 创建时间：2026-06-02 🏷️ 标签：#CUDA #显存 #GlobalMemory #SharedMemory #UnifiedMemory #OOM 📚 前置知识：[[04-cuda-programming-model]]（CUDA 编程模型） [[03-gpu-architecture]]（GPU 架构） 📚 相关知识：[[06-cuda-optimization]]（CUDA 优化） [[10-dl-training-optimization]]（训练优化实战）

先抓住直觉 ​

GPU 计算常像“厨房很快，仓库取料很慢”。优化内存的目标，是让相邻线程一起搬连续数据，并把会重复使用的小块数据暂存在离计算单元更近的位置。

  • 必须理解：Global Memory 慢但大；Shared Memory 快但小；合并访问；同步的作用。
  • 用到再查：Bank 数量、各代 GPU 容量限制和 Unified Memory API。
  • 分清两类问题：容量不足问“装不装得下”，性能不足问“送得够不够快”。

场景：175B 参数的模型，能装进一张 A100 吗？ ​

┌─────────────────────────────────────────────────────────────┐
│                                                             │
│  175B 参数，每个参数占多少字节？                          │
│                                                             │
│  FP32（32位浮点）：175B × 4 字节 = 700 GB               │
│  FP16（16位浮点）：175B × 2 字节 = 350 GB               │
│  BF16（16位浮点）：175B × 2 字节 = 350 GB               │
│                                                             │
│  A100 显存：40 GB / 80 GB（取决于型号）                   │
│                                                             │
│  结论：单卡 A100 根本装不下 175B 参数的模型！              │
│  仅存 FP16 参数也至少要 5 张 A100（80GB 版本），           │
│  训练时还要为梯度、优化器状态和激活值预留更多空间。        │
│                                                             │
│  显存不够怎么办？→ 这就是本章要解决的问题。               │
│                                                             │
└─────────────────────────────────────────────────────────────┘
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17

第1节：CUDA 内存类型全景图 ​

五种内存的特点 ​

┌─────────────────────────────────────────────────────────────┐
│                    CUDA 内存类型速查表                         │
├─────────────────────────────────────────────────────────────┤
│                                                             │
│  类型           │ 位置    │ 容量       │ 延迟    │ 访问方式  │
│  ──────────────┼─────────┼───────────┼────────┼─────────  │
│  Register      │ SM 内   │ 256 KB/SM │ ~1ns  │ 线程私有  │
│  Local Memory  │ 全局显存│ ~8 GB     │ ~500ns│ 线程私有  │
│  Shared Memory │ SM 内   │ 128 KB/SM │ ~1ns  │ block 内  │
│  L1/L2 Cache  │ GPU 内  │ 128KB/40MB│ ~30ns │ 自动      │
│  Global Memory │ GPU HBM │ 40-80 GB │ ~500ns│ 所有线程  │
│  Constant Mem │ 全局显存│ 64 KB    │ ~30ns │ 所有只读  │
│  Texture Mem  │ 全局显存│ 依赖显存  │ ~500ns│ 只读      │
│                                                             │
│  带宽（GB/s）：                                             │
│  Shared/L1: ~10,000 GB/s   ← 接近寄存器速度！             │
│  L2 Cache:   ~3,500 GB/s                                      │
│  Global Mem: ~2,000 GB/s                                      │
│                                                             │
└─────────────────────────────────────────────────────────────┘
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20

从 CPU 到 GPU 的内存层次对比 ​

┌─────────────────────────────────────────────────────────────┐
│                    GPU vs CPU 内存层次对比                     │
├─────────────────────────────────────────────────────────────┤
│                                                             │
│  CPU：                                                    │
│  Register → L1 → L2 → L3 → DDR Memory                    │
│  (~1ns)   (~3ns) (~10ns) (~100ns)                       │
│                                                             │
│  GPU：                                                    │
│  Register → Shared Memory → L2 → Global Memory (HBM)       │
│  (~1ns)    (~1ns)          (~30ns)  (~500ns)            │
│                                                             │
│  关键区别：                                                 │
│  - GPU Shared Memory 比 CPU Cache 快 10 倍                  │
│  - Shared Memory 可编程（用户显式管理）                     │
│  - CPU L1/L2 Cache 对程序员透明                            │
│                                                             │
└─────────────────────────────────────────────────────────────┘
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18

第2节：Global Memory——容量最大、速度最慢 ​

基本操作 ​

cpp
// 全局内存分配
float* d_array;
cudaMalloc(&d_array, N * sizeof(float));

// 全局内存释放
cudaFree(d_array);

// 全局内存拷贝
cudaMemcpy(d_dst, d_src, size, cudaMemcpyDeviceToDevice);  // GPU→GPU
cudaMemcpy(h_dst, d_src, size, cudaMemcpyDeviceToHost);    // GPU→CPU
cudaMemcpy(d_dst, h_src, size, cudaMemcpyHostToDevice);    // CPU→GPU
1
2
3
4
5
6
7
8
9
10
11

内存合并访问——性能的关键 ​

┌─────────────────────────────────────────────────────────────┐
│                    合并访问（Coalesced Access）               │
├─────────────────────────────────────────────────────────────┤
│                                                             │
│  理想情况：连续线程访问连续内存                            │
│                                                             │
│  Thread 0 → 读取 addr+0                                  │
│  Thread 1 → 读取 addr+4                                  │
│  Thread 2 → 读取 addr+8                                  │
│  Thread 3 → 读取 addr+12  ...                            │
│  ↓                                                         │
│  GPU 一次事务读取 128 字节（32 线程 × 4 字节）             │
│  = 一次内存事务搞定，效率 100%！                          │
│                                                             │
└─────────────────────────────────────────────────────────────┘
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
cpp
// ✅ 好：合并访问（行优先存储，连续线程访问连续内存）
__global__
void copy_row_major(float* src, float* dst, int rows, int cols) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    if (row < rows && col < cols) {
        // src[row * cols + col] 在内存中是连续的
        dst[row * cols + col] = src[row * cols + col];
    }
}

// ❌ 差：交叉访问（列优先存储，线程访问跳跃内存）
__global__
void copy_col_major_bad(float* src, float* dst, int rows, int cols) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    if (row < rows && col < cols) {
        // src[col * rows + row] 访问跳跃地址
        dst[col * rows + row] = src[col * rows + row];
    }
}

// ✅ 好：转置时用共享内存合并写入
__global__
void transpose_optimized(float* src, float* dst, int rows, int cols) {
    __shared__ float tile[32][32];  // 共享内存缓冲

    int x = blockIdx.x * 32 + threadIdx.x;
    int y = blockIdx.y * 32 + threadIdx.y;

    // 合并读取到共享内存
    if (y < rows && x < cols)
        tile[threadIdx.y][threadIdx.x] = src[y * cols + x];

    __syncthreads();

    // 转置后，合并写入全局内存
    x = blockIdx.y * 32 + threadIdx.x;  // 交换 blockIdx
    y = blockIdx.x * 32 + threadIdx.y;
    if (y < cols && x < rows)
        dst[y * rows + x] = tile[threadIdx.x][threadIdx.y];
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44

带宽计算 ​

┌─────────────────────────────────────────────────────────────┐
│                    全局内存带宽计算                           │
├─────────────────────────────────────────────────────────────┤
│                                                             │
│  A100 HBM2e 带宽：2.0 TB/s = 2000 GB/s                    │
│                                                             │
│  如果每次内存访问只读取 4 字节（float）：                    │
│  最大带宽利用率时：2000 / 4 = 500 亿次读取/秒              │
│  如果每个线程做 1 FLOP：最大算力 = 500 GFLOPS（实际 19.5T）│
│  → 严重受限于内存带宽！                                    │
│                                                             │
│  解决：用向量化加载，一次读取多个字节                       │
│  - float4：一次读取 16 字节，效率提升 4 倍                  │
│                                                             │
└─────────────────────────────────────────────────────────────┘
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15

第3节：Shared Memory——可编程的片上存储 ​

为什么需要共享内存？ ​

┌─────────────────────────────────────────────────────────────┐
│                    共享内存 vs 全局内存                      │
├─────────────────────────────────────────────────────────────┤
│                                                             │
│  全局内存访问：                                            │
│  Thread 0 ──┐                                             │
│  Thread 1 ──┼──→ 全局内存（500ns，HBM）                │
│  Thread 2 ──┘                                             │
│  每次访问 500ns，多个线程排队等待                          │
│                                                             │
│  共享内存访问：                                            │
│  Thread 0 ──┐                                             │
│  Thread 1 ──┼──→ 共享内存（1ns，片上 SRAM）             │
│  Thread 2 ──┘                                             │
│  所有线程同时读取，极快                                     │
│                                                             │
│  策略：                                                   │
│  1. 先从全局内存一次性加载数据到共享内存（所有线程协同）    │
│  2. 然后所有线程从共享内存高速读取                          │
│  3. 共享内存中的结果写回全局内存                           │
│                                                             │
└─────────────────────────────────────────────────────────────┘
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22

共享内存 Bank 冲突 ​

┌─────────────────────────────────────────────────────────────┐
│                    Bank 冲突（Bank Conflict）                  │
├─────────────────────────────────────────────────────────────┤
│                                                             │
│  共享内存被分成 32 个 Bank（A100）                         │
│  每个 Bank 每次只能服务一次访问                             │
│                                                             │
│  无冲突：32 个线程访问 32 个不同 Bank → 同时完成           │
│                                                             │
│  2-way Bank 冲突：                                         │
│  线程 0 和线程 16 访问同一个 Bank → 需要 2 次访问          │
│  效率降低为 50%                                            │
│                                                             │
│  8-way Bank 冲突：8 个线程访问同一 Bank → 需要 8 次访问    │
│  效率降低为 12.5%                                          │
│                                                             │
│  A100 Bank 大小：8 字节（64 位）                           │
│  Bank 索引 = (地址 / 8) % 32                             │
│                                                             │
└─────────────────────────────────────────────────────────────┘
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
cpp
// ✅ 好：避免 Bank 冲突（行优先访问）
__global__
void smem_row_access(float* data, int N) {
    __shared__ float tile[32][32];  // 行优先排列

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

    // 线程 threadIdx.x 读取不同列 → 不同 Bank，无冲突
    tile[threadIdx.y][threadIdx.x] = data[row * N + col];
    // tile[row][col]，col 由 threadIdx.x 决定
    // 每个线程访问不同列 → 不同 Bank
}

// ❌ 差：Bank 冲突（列优先访问）
__global__
void smem_col_access_bad(float* data, int N) {
    __shared__ float tile[32][32];  // 仍是行优先排列

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

    // 线程 threadIdx.y 读取不同行 → 同一 Bank，冲突！
    tile[threadIdx.x][threadIdx.y] = data[row * N + col];
    // tile[row][col]，col 由 threadIdx.y 决定
    // threadIdx.y=0 和 threadIdx.y=1 → 同列 → 同 Bank！
}

// ✅ 好：列优先排列，匹配列优先访问
__shared__ float tile[32][33];  // +1 padding 避免 Bank 冲突
// 33 = 32 + 1，每行多一个元素
// Bank 索引计算：(col * 33 / 8) % 32
// 33/8 不是整数，列之间自动错开，无冲突！
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33

__syncthreads()——线程同步 ​

cpp
// __syncthreads() 的两个用途：

// 用途1：共享内存写入后，等待所有线程完成，再读取
__global__
void kernel_with_sync(float* global_data, float* result, int N) {
    __shared__ float shared_data[256];

    int tid = threadIdx.x;

    // Step 1: 每个线程加载一个元素到共享内存
    shared_data[tid] = global_data[tid];
    __syncthreads();  // ⚠️ 必须！确保所有线程都写完了

    // Step 2: 所有线程都读完了，可以安全使用
    float sum = 0.0f;
    for (int i = 0; i < tid; i++) {
        sum += shared_data[i];
    }

    if (tid == 0) {
        *result = sum;
    }
}

// 用途2：防止先完成的线程读取后完成线程的数据（race condition）
// ⚠️ 没有 __syncthreads() 可能导致读到未写入的数据
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26

第4节：Unified Memory——统一内存管理 ​

什么是 Unified Memory？ ​

┌─────────────────────────────────────────────────────────────┐
│                    Unified Memory 原理                       │
├─────────────────────────────────────────────────────────────┤
│                                                             │
│  传统 CUDA：                                               │
│  CPU 内存 ──cudaMemcpy──→ GPU 显存                       │
│  必须手动管理数据迁移                                        │
│                                                             │
│  Unified Memory（CUDA 6.0+）：                             │
│  CPU 内存 ←─── 统一地址空间 ───→ GPU 显存               │
│            操作系统自动迁移数据                             │
│                                                             │
│  cudaMallocManaged() 分配统一内存                          │
│  CPU 或 GPU 访问同一指针，OS 自动迁移                       │
│                                                             │
└─────────────────────────────────────────────────────────────┘
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16

使用示例 ​

cpp
// 传统方式：手动管理
float *d_data, *h_data;
h_data = (float*)malloc(N * sizeof(float));
cudaMalloc(&d_data, N * sizeof(float));

// CPU 处理
for (int i = 0; i < N; i++) h_data[i] = i;

// CPU→GPU 拷贝
cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice);

// GPU 计算
process<<<blocks, threads>>>(d_data, N);

// GPU→CPU 拷贝
cudaMemcpy(h_data, d_data, N * sizeof(float), cudaMemcpyDeviceToHost);

// ================================================================
// Unified Memory 方式：
// ================================================================
float *data;
cudaMallocManaged(&data, N * sizeof(float));  // 分配统一内存

// CPU 处理
for (int i = 0; i < N; i++) data[i] = i;

// GPU 计算（OS 自动把数据迁移到 GPU）
process<<<blocks, threads>>>(data, N);

// CPU 读取结果（OS 自动把数据迁移回 CPU）
cudaDeviceSynchronize();  // 等待 GPU 完成
for (int i = 0; i < N; i++) printf("%f\n", data[i]);

cudaFree(data);
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34

预取和亲和性 ​

cpp
// Unified Memory 性能优化：手动预取

float *data;
cudaMallocManaged(&data, N * sizeof(float));

// 初始化数据（CPU）
for (int i = 0; i < N; i++) data[i] = i;

// 预取到 GPU
cudaMemPrefetchAsync(data, N * sizeof(float), 0);  // deviceId=0

// GPU 计算
process<<<blocks, threads>>>(data, N);
cudaDeviceSynchronize();

// 预取回 CPU
cudaMemPrefetchAsync(data, N * sizeof(float), cudaCpuDeviceId);

// CPU 读取结果
for (int i = 0; i < N; i++) printf("%f\n", data[i]);
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20

第5节：显存不够的解决方案 ​

显存估算 ​

┌─────────────────────────────────────────────────────────────┐
│                    显存占用估算                              │
├─────────────────────────────────────────────────────────────┤
│                                                             │
│  175B 参数模型显存占用（训练）：                             │
│                                                             │
│  FP32 训练：                                               │
│    参数：175B × 4B = 700 GB                               │
│    梯度：175B × 4B = 700 GB                               │
│    优化器状态（Adam）：                                    │
│      m（动量）：175B × 4B = 700 GB                        │
│      v（方差）：175B × 4B = 700 GB                        │
│    激活值：~200 GB（依赖序列长度和 batch size）             │
│    总计：~2500 GB ≈ 25 张 A100（80GB）                   │
│                                                             │
│  FP16 训练（主流）：                                       │
│    参数：175B × 2B = 350 GB                               │
│    梯度：175B × 2B = 350 GB                               │
│    优化器状态（FP32）：                                    │
│      m：175B × 4B = 700 GB                                │
│      v：175B × 4B = 700 GB                                │
│    激活值：~100 GB                                         │
│    总计：~2200 GB ≈ 28 张 A100（80GB）                   │
│                                                             │
│  FP16 + DeepSpeed ZeRO-3：                                │
│    参数/梯度/优化器状态分片到多卡                           │
│    每卡：参数(350/8) + 梯度(350/8) + 优化器(2200/8)      │
│         ≈ 375 GB → 还是装不下！                            │
│                                                             │
│  结论：175B 模型训练至少需要多卡并行                       │
│                                                             │
└─────────────────────────────────────────────────────────────┘
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32

显存优化技术 ​

┌─────────────────────────────────────────────────────────────┐
│                    显存优化技术一览                          │
├─────────────────────────────────────────────────────────────┤
│                                                             │
│  1. 混合精度训练（FP16/BF16）                              │
│     前向/反向用 FP16，参数和优化器用 FP32                  │
│     显存减半，精度几乎不变                                  │
│                                                             │
│  2. 梯度累积（Gradient Accumulation）                       │
│     不增大 batch size，减少显存占用                        │
│                                                             │
│  3. Activation Checkpointing（梯度检查点）                  │
│     不保存所有中间激活值，只保存每 N 层的输出               │
│     反向传播时重新计算被丢弃的激活值                        │
│     显存 O(N) → O(N^(1/2))，时间略增                      │
│                                                             │
│  4. DeepSpeed ZeRO（零冗余优化器）                         │
│     ZeRO-1：优化器状态分片（省 4 倍显存）                 │
│     ZeRO-2：梯度分片（再省 2 倍）                         │
│     ZeRO-3：参数分片（线性扩展到无限大模型）               │
│                                                             │
│  5. CPU Offload                                             │
│     把优化器状态卸载到 CPU 内存                            │
│     速度慢，但可以训练更大的模型                           │
│                                                             │
│  6. 梯度检查点（Activation Checkpointing）+ CPU Offload   │
│     激活值和优化器状态都卸载到 CPU                        │
│     可以用单卡训练超大模型（速度很慢）                      │
│                                                             │
└─────────────────────────────────────────────────────────────┘
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30

升华：内存优化的核心原则 ​

┌─────────────────────────────────────────────────────────────┐
│                    CUDA 内存优化 3 原则                      │
├─────────────────────────────────────────────────────────────┤
│                                                             │
│  原则1：能用共享内存就不用全局内存                         │
│  共享内存带宽是全局内存的 5 倍，延迟是 1/500              │
│                                                             │
│  原则2：合并访问是全局内存优化的关键                       │
│  连续线程访问连续内存 → 一次事务搞定                       │
│  跳跃访问 → 多次事务 → 带宽利用率暴跌                     │
│                                                             │
│  原则3：减少全局内存访问量                                 │
│  - 共享内存复用数据                                        │
│  - 向量化加载（float4 一次 16 字节）                      │
│  - 避免不必要的数据传输                                    │
│                                                             │
└─────────────────────────────────────────────────────────────┘
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17

"AI 可查 vs 必须理解"清单 ​

AI 可查：
✅ cudaMalloc/cudaMemcpy 的具体参数
✅ 共享内存 bank 冲突的具体计算公式
✅ DeepSpeed ZeRO 的具体配置参数

必须理解：
🔴 全局内存延迟 500ns vs 共享内存 1ns → 为什么共享内存如此重要
🔴 合并访问：连续线程访问连续内存，效率 100%
🔴 Bank 冲突：同 Bank 访问串行化，效率降低
🔴 __syncthreads() 的必要性：防止先读完的线程读到未写入的数据
🔴 Unified Memory：统一地址空间，OS 自动迁移
🔴 175B 模型为什么单卡装不下：FP32 训练需要 ~2500 GB 显存
1
2
3
4
5
6
7
8
9
10
11
12

学习状态：🟡 开始学习

最后更新于:

Pager
下一篇← 系统与高性能 / Systems & Performance

持续记录，持续成长

Copyright © Tidenflow