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节：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 │
│ │
└─────────────────────────────────────────────────────────────┘从 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 对程序员透明 │
│ │
└─────────────────────────────────────────────────────────────┘第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内存合并访问——性能的关键
┌─────────────────────────────────────────────────────────────┐
│ 合并访问（Coalesced Access） │
├─────────────────────────────────────────────────────────────┤
│ │
│ 理想情况：连续线程访问连续内存 │
│ │
│ Thread 0 → 读取 addr+0 │
│ Thread 1 → 读取 addr+4 │
│ Thread 2 → 读取 addr+8 │
│ Thread 3 → 读取 addr+12 ... │
│ ↓ │
│ GPU 一次事务读取 128 字节（32 线程 × 4 字节） │
│ = 一次内存事务搞定，效率 100%！ │
│ │
└─────────────────────────────────────────────────────────────┘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];
}带宽计算
┌─────────────────────────────────────────────────────────────┐
│ 全局内存带宽计算 │
├─────────────────────────────────────────────────────────────┤
│ │
│ A100 HBM2e 带宽：2.0 TB/s = 2000 GB/s │
│ │
│ 如果每次内存访问只读取 4 字节（float）： │
│ 最大带宽利用率时：2000 / 4 = 500 亿次读取/秒 │
│ 如果每个线程做 1 FLOP：最大算力 = 500 GFLOPS（实际 19.5T）│
│ → 严重受限于内存带宽！ │
│ │
│ 解决：用向量化加载，一次读取多个字节 │
│ - float4：一次读取 16 字节，效率提升 4 倍 │
│ │
└─────────────────────────────────────────────────────────────┘第3节：Shared Memory——可编程的片上存储
为什么需要共享内存？
┌─────────────────────────────────────────────────────────────┐
│ 共享内存 vs 全局内存 │
├─────────────────────────────────────────────────────────────┤
│ │
│ 全局内存访问： │
│ Thread 0 ──┐ │
│ Thread 1 ──┼──→ 全局内存（500ns，HBM） │
│ Thread 2 ──┘ │
│ 每次访问 500ns，多个线程排队等待 │
│ │
│ 共享内存访问： │
│ Thread 0 ──┐ │
│ Thread 1 ──┼──→ 共享内存（1ns，片上 SRAM） │
│ Thread 2 ──┘ │
│ 所有线程同时读取，极快 │
│ │
│ 策略： │
│ 1. 先从全局内存一次性加载数据到共享内存（所有线程协同） │
│ 2. 然后所有线程从共享内存高速读取 │
│ 3. 共享内存中的结果写回全局内存 │
│ │
└─────────────────────────────────────────────────────────────┘共享内存 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 │
│ │
└─────────────────────────────────────────────────────────────┘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 不是整数，列之间自动错开，无冲突！__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() 可能导致读到未写入的数据第4节：Unified Memory——统一内存管理
什么是 Unified Memory？
┌─────────────────────────────────────────────────────────────┐
│ Unified Memory 原理 │
├─────────────────────────────────────────────────────────────┤
│ │
│ 传统 CUDA： │
│ CPU 内存 ──cudaMemcpy──→ GPU 显存 │
│ 必须手动管理数据迁移 │
│ │
│ Unified Memory（CUDA 6.0+）： │
│ CPU 内存 ←─── 统一地址空间 ───→ GPU 显存 │
│ 操作系统自动迁移数据 │
│ │
│ cudaMallocManaged() 分配统一内存 │
│ CPU 或 GPU 访问同一指针，OS 自动迁移 │
│ │
└─────────────────────────────────────────────────────────────┘使用示例
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);预取和亲和性
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]);第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. 混合精度训练（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 │
│ 可以用单卡训练超大模型（速度很慢） │
│ │
└─────────────────────────────────────────────────────────────┘升华：内存优化的核心原则
┌─────────────────────────────────────────────────────────────┐
│ CUDA 内存优化 3 原则 │
├─────────────────────────────────────────────────────────────┤
│ │
│ 原则1：能用共享内存就不用全局内存 │
│ 共享内存带宽是全局内存的 5 倍，延迟是 1/500 │
│ │
│ 原则2：合并访问是全局内存优化的关键 │
│ 连续线程访问连续内存 → 一次事务搞定 │
│ 跳跃访问 → 多次事务 → 带宽利用率暴跌 │
│ │
│ 原则3：减少全局内存访问量 │
│ - 共享内存复用数据 │
│ - 向量化加载（float4 一次 16 字节） │
│ - 避免不必要的数据传输 │
│ │
└─────────────────────────────────────────────────────────────┘"AI 可查 vs 必须理解"清单
AI 可查：
✅ cudaMalloc/cudaMemcpy 的具体参数
✅ 共享内存 bank 冲突的具体计算公式
✅ DeepSpeed ZeRO 的具体配置参数
必须理解：
🔴 全局内存延迟 500ns vs 共享内存 1ns → 为什么共享内存如此重要
🔴 合并访问：连续线程访问连续内存，效率 100%
🔴 Bank 冲突：同 Bank 访问串行化，效率降低
🔴 __syncthreads() 的必要性：防止先读完的线程读到未写入的数据
🔴 Unified Memory：统一地址空间，OS 自动迁移
🔴 175B 模型为什么单卡装不下：FP32 训练需要 ~2500 GB 显存学习状态：🟡 开始学习