CUDA Kernel、内存与性能

CUDA 优化的主线不是背参数，而是追踪：每个线程做什么、读取哪些字节、在哪里等待、哪种资源限制并发。

1. Grid-stride loop

cpp

__global__ void scale(float* x, int n, float k) {
    for (int i = blockIdx.x * blockDim.x + threadIdx.x;
         i < n;
         i += blockDim.x * gridDim.x) {
        x[i] *= k;
    }
}

它让有限 Grid 覆盖任意长度输入，并保持连续线程首次访问连续元素。仍需根据设备和 Kernel 特征选择 Grid/Block，而不是固定迷信 256。

2. 合并访存

text

good: lane 0->x[i], lane 1->x[i+1], lane 2->x[i+2] ...
bad:  lane 0->x[0], lane 1->x[stride], lane 2->x[2*stride] ...

前者通常能合并为较少的内存事务。后者可能读取许多 Cache line 却只使用少量字节。真实事务规则依架构、类型、对齐和缓存路径而定，应由 profiler 验证。

3. Shared Memory：显式片上复用

矩阵分块、Stencil 等算法可让一个 Block 先把数据协作加载到 Shared Memory，再重复使用：

text

global memory -> cooperative tile load -> shared memory
                                      -> many arithmetic uses
                                      -> write result

收益来自减少 Global Memory 流量；代价包括加载、barrier、容量占用和 bank conflict。只搬一次又只用一次通常没有意义。

4. 寄存器、Shared Memory 与 Occupancy

每线程寄存器越多、每 Block Shared Memory 越多，一个 SM 能同时驻留的 Block/Warp 可能越少。强行限制寄存器有时会造成 spill 到本地内存，反而更慢。

优化判断链：

text

correctness
  -> enough parallel work?
  -> transfer or kernel bound?
  -> memory bandwidth / latency / compute / synchronization?
  -> which resource limits residency?
  -> change one variable and measure

5. Stream 与重叠

Stream 是有序命令队列：同一 Stream 内按序，不同 Stream 间可能并发。使用 pinned host memory、设备能力允许且依赖正确时，传输可能与计算重叠。

text

chunk 0: H2D | kernel | D2H
chunk 1:       H2D | kernel | D2H

“可能并发”不等于保证重叠；还受拷贝引擎、默认 Stream 语义、依赖和资源竞争影响。

6. 常见反优化

小数据也搬到 GPU，传输和启动成本超过计算；
每个小阶段都同步，破坏流水线；
只看 Kernel 时间，不算 H2D/D2H；
为提高 Occupancy 牺牲过多寄存器；
用 Shared Memory 却没有数据复用；
未先调用成熟库就手写通用矩阵乘。

深入原理与工程实践

前面的内容负责建立统一心智模型；下面把同一主题继续拆到执行过程、代码、性能代价与工程判断。

完整迁入：原 CUDA Kernel 与内存全文

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 显存

学习状态：🟡 开始学习

完整迁入：原 CUDA 性能优化全文

CUDA 性能优化——从 30 天缩短到 10 天 / CUDA Performance Optimization

📅 创建时间：2026-06-02 🏷️ 标签：#CUDA #优化 #Occupancy #WarpDivergence #TensorCore #CUDAStream 📚 前置知识：[[05-cuda-kernel-and-memory]]（内存管理） [[04-cuda-programming-model]]（编程模型） 📚 相关知识：[[11-performance-analysis-tools]]（性能分析工具） [[10-dl-training-optimization]]（训练优化）

先抓住直觉

性能优化不是把所有技巧都加上，而是先找到最窄的那段管道。算力没吃满可能是数据供不上、并发线程不够、分支浪费，也可能是 CPU 没有及时提交工作。

必须理解：先测量再优化；计算瓶颈与带宽瓶颈；Occupancy 只是手段，不是最终目标。
用到再查：具体 profiler 指标、Shuffle 写法和 Stream 优先级 API。
建议顺序：先读第 11 章的测量流程，再回本章选择优化方法。

场景：你的 CUDA 代码跑得比预期慢 10 倍

┌─────────────────────────────────────────────────────────────┐
│                                                             │
│  你写了一个矩阵乘法的 CUDA kernel，自测了一下：           │
│                                                             │
│  理论算力（A100）：156 TFLOPS（Tensor Core FP16）         │
│  你的实际性能：15 GFLOPS                                   │
│                                                             │
│  差距：10,000 倍！                                         │
│                                                             │
│  问题出在哪里？                                            │
│  → 内存带宽瓶颈？                                         │
│  → Warp 分支分化？                                       │
│  → 共享内存 bank 冲突？                                   │
│  → occupancy 太低？                                       │
│                                                             │
│  本章就是来解决这些问题的。                               │
│                                                             │
└─────────────────────────────────────────────────────────────┘

第1节：性能分析——先定位瓶颈

屋顶线模型（Roofline Model）

┌─────────────────────────────────────────────────────────────┐
│                    Roofline 分析                           │
├─────────────────────────────────────────────────────────────┤
│                                                             │
│  性能上限 = min(算力上限, 带宽上限 × 算术强度)           │
│                                                             │
│       GFLOPS                                              │
│  156k ┤                                          ★ GPU   │
│       │                                        /           │
│   20k ┤─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─           │
│       │                              /│                  │
│    5k ┤─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─  │
│       │                         /│                       │
│    1k ┤─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─  │
│       │              /│                                   │
│  312G ┤─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─  │
│       │       /│                                          │
│       └───────┼─────────────────────────────→             │
│              500      1000    2000  算术强度(FLOP/Byte)   │
│                                                             │
│  你的 kernel：算术强度 = 5 FLOP / 8 Byte = 0.6             │
│  → 落在带宽受限区，性能上限 = 带宽 × 0.6                   │
│  → 不是 GPU 算力不够，是内存访问拖了后腿                   │
│                                                             │
└─────────────────────────────────────────────────────────────┘

CUDA 性能分析工具

bash

### NVIDIA Nsight Compute（kernel 级别）
ncu --set full ./matrix_mul
### 输出：SM 利用率、内存带宽、warp 执行效率等

### NVIDIA Nsight Systems（系统级别）
nsys profile ./my_cuda_app
### 输出：CPU-GPU 传输时间、kernel 执行 timeline

### 简单计时
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

cudaEventRecord(start);
my_kernel<<<blocks, threads>>>(...);
cudaEventRecord(stop);
cudaEventSynchronize(stop);

float ms;
cudaEventElapsedTime(&ms, start, stop);
printf("Kernel time: %f ms\n", ms);

第2节：Occupancy——GPU 利用率的根本

什么是 Occupancy？

┌─────────────────────────────────────────────────────────────┐
│                    Occupancy（占用率）                        │
├─────────────────────────────────────────────────────────────┤
│                                                             │
│  Occupancy = 实际运行线程数 / 最大可运行线程数              │
│                                                             │
│  A100 每 SM 最大：2048 线程 = 64 Warp × 32 线程           │
│                                                             │
│  如果只运行了 1024 线程：Occupancy = 1024/2048 = 50%      │
│                                                             │
│  低 Occupancy 的后果：                                      │
│  → Warp Scheduler 选择的余地小                            │
│  → 等待内存时没有足够的 Warp 可切换                        │
│  → GPU 闲置                                                │
│                                                             │
└─────────────────────────────────────────────────────────────┘

计算 Occupancy

┌─────────────────────────────────────────────────────────────┐
│                    A100 Occupancy 计算                       │
├─────────────────────────────────────────────────────────────┤
│                                                             │
│  每 SM 资源：                                              │
│  - 最大线程数：2048                                        │
│  - 最大 Block 数：16                                       │
│  - 寄存器数：65536 × 32-bit = 256 KB                     │
│  - 共享内存：128 KB                                        │
│                                                             │
│  限制因素：                                                 │
│  1. 寄存器限制：每个线程用的寄存器越多，能开的线程越少      │
│     Max threads = min(2048, 65536 / registers_per_thread) │
│                                                             │
│  2. 共享内存限制：每 block 用的共享内存越多，能开的 block越少│
│     Max blocks = 128KB / shared_mem_per_block             │
│                                                             │
│  3. Block 数限制：最多 16 个 block/SM                      │
│                                                             │
│  示例：                                                    │
│  每线程用 32 个寄存器，每 block 16KB 共享内存              │
│  → 寄存器限制：65536/32 = 2048 线程 → 64 Warp            │
│  → 共享内存限制：128KB/16KB = 8 个 block                 │
│  → Block 数限制：min(8, 16) = 8                          │
│  → 实际线程：8 block × 256 threads/block = 2048          │
│  → Occupancy = 2048/2048 = 100%                          │
│                                                             │
└─────────────────────────────────────────────────────────────┘

提高 Occupancy

cpp

// ❌ 差：每线程寄存器太多，Occupancy 低
__global__
void heavy_register_kernel(float* data, int N) {
    float temp1, temp2, temp3, temp4, temp5;  // 5 个临时变量
    float temp6, temp7, temp8, temp9, temp10;
    float temp11, temp12, temp13, temp14, temp15;

    // 大量使用寄存器
    temp1 = data[threadIdx.x];
    temp2 = temp1 * 2.0f;
    // ...
    data[threadIdx.x] = temp15;
}

// ✅ 好：减少临时变量，使用共享内存
__global__
void light_register_kernel(float* data, int N) {
    __shared__ float temp[256];  // 共享内存代替寄存器

    temp[threadIdx.x] = data[threadIdx.x];
    temp[threadIdx.x] *= 2.0f;
    __syncthreads();

    data[threadIdx.x] = temp[threadIdx.x];
}

第3节：向量化加载——充分利用内存带宽

什么是向量化加载？

┌─────────────────────────────────────────────────────────────┐
│                    向量化加载原理                           │
├─────────────────────────────────────────────────────────────┤
│                                                             │
│  普通加载：                                                │
│  Thread 0: 读取 addr+0   (4 字节)                        │
│  Thread 1: 读取 addr+4   (4 字节)                        │
│  Thread 2: 读取 addr+8   (4 字节)                        │
│  Thread 3: 读取 addr+12  (4 字节)                        │
│  → 4 次内存事务                                          │
│                                                             │
│  向量化加载（float4）：                                     │
│  Thread 0-3: 一次读取 addr+0~15  (16 字节)              │
│  → 1 次内存事务，效率提升 4 倍！                          │
│                                                             │
│  A100 支持的最大向量化：float4（128 位）                   │
│                                                             │
└─────────────────────────────────────────────────────────────┘

向量化代码示例

cpp

// 普通版本
__global__
void vector_add_basic(float* a, float* b, float* c, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        c[i] = a[i] + b[i];  // 每个线程读取 3×4=12 字节
    }
}

// 向量化版本（效率提升 ~2-3 倍）
__global__
void vector_add_vectorized(float4* a, float4* b, float4* c, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N / 4) {  // 数据量变为 1/4（每次处理 4 个 float）
        float4 va = a[i];
        float4 vb = b[i];
        c[i] = make_float4(
            va.x + vb.x,
            va.y + vb.y,
            va.z + vb.z,
            va.w + vb.w
        );
    }
}

// 调用
int N_aligned = (N + 3) / 4;  // 对齐到 4
vector_add_vectorized<<<blocks, threads>>>(
    (float4*)d_a, (float4*)d_b, (float4*)d_c, N_aligned);

第4节：Warp 级优化——分支和洗牌

减少分支分化

cpp

// ❌ 差：Warp 内分支分化
__global__
void update_if_else(float* data, float threshold, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        if (data[i] > threshold) {
            data[i] = sqrt(data[i]);
        } else {
            data[i] = data[i] * data[i];
        }
    }
}

// ✅ 好：用 step 函数替代分支
__global__
void update_branchless(float* data, float threshold, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        float val = data[i];
        float greater = (val > threshold);  // 0.0 或 1.0

        // 选择：greater * sqrt + (1-greater) * square
        float res = greater * sqrtf(val) + (1.0f - greater) * val * val;
        data[i] = res;
    }
}

// ✅ 好：把相同分支的线程放到一起（Block 分组）
__global__
void update_grouped(float* data, float threshold, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        // Block 0: 大于阈值，Block 1: 小于等于阈值
        if (blockIdx.x == 0) {
            if (data[i] > threshold)
                data[i] = sqrt(data[i]);
        } else {
            if (data[i] <= threshold)
                data[i] = data[i] * data[i];
        }
    }
}

Warp 洗牌（Shuffle）指令

cpp

// Warp Shuffle：同 Warp 内线程直接交换数据，不需要共享内存！

// __shfl_sync：同一 Warp 内广播某个线程的值
__global__
void warp_reduce_sum(float* data, float* result, int N) {
    float val = data[threadIdx.x];

    // Warp 内归约
    for (int offset = 16; offset > 0; offset >>= 1) {
        // 把 offset 线程的值加到当前线程
        val += __shfl_down_sync(0xffffffff, val, offset);
    }

    // thread 0 保存结果
    if (threadIdx.x == 0) {
        *result = val;
    }
}

// __shfl_down_sync 的效果（8 线程示例）：
// 初始：  [t0=1, t1=2, t2=3, t3=4, t4=5, t5=6, t6=7, t7=8]
// offset=4: [t0=5, t1=6, t2=7, t3=8, t4=5, t5=6, t6=7, t7=8]
//          当前线程从 thread+4 获取值并相加
//          t0 = 1+5=6, t1=2+6=8, ...

// 对比用共享内存的版本：
__global__
void shared_reduce_sum(float* data, float* result, int N) {
    __shared__ float sdata[32];
    sdata[threadIdx.x] = data[threadIdx.x];
    __syncthreads();

    for (int s = 16; s > 0; s >>= 1) {
        if (threadIdx.x < s) {
            sdata[threadIdx.x] += sdata[threadIdx.x + s];
        }
        __syncthreads();
    }

    if (threadIdx.x == 0) *result = sdata[0];
}

// Shuffle 版本不需要 __syncthreads（因为同一 Warp 同步执行）
// 速度更快！

第5节：CUDA Stream——异步并行

什么是 CUDA Stream？

┌─────────────────────────────────────────────────────────────┐
│                    CUDA Stream 概念                         │
├─────────────────────────────────────────────────────────────┤
│                                                             │
│  默认 Stream（NULL）：                                      │
│  所有操作顺序执行                                           │
│  H2D ──→ Kernel ──→ D2H                                  │
│                                                             │
│  多 Stream：                                               │
│  Stream 0: H2D ──→ Kernel ──→ D2H                        │
│  Stream 1:      H2D ──→ Kernel ──→ D2H                   │
│  ↑ 不同 Stream 的操作可以重叠！                            │
│                                                             │
│  内存拷贝和 Kernel 执行可以同时进行                         │
│                                                             │
└─────────────────────────────────────────────────────────────┘

Stream 使用示例

cpp

// 单 Stream 版本（顺序执行）
cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
kernel<<<blocks, threads>>>(d_a, d_b, d_c);
cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);
// 总时间 = H2D + Kernel + D2H

// ================================================================
// 多 Stream 版本（重叠执行）
// ================================================================
cudaStream_t stream0, stream1;
cudaStreamCreate(&stream0);
cudaStreamCreate(&stream1);

const int chunk = size / 2;

// Stream 0: 前半部分
cudaMemcpyAsync(d_a, h_a, chunk, cudaMemcpyHostToDevice, stream0);
kernel<<<blocks, threads, 0, stream0>>>(d_a, d_b, d_c, chunk);
cudaMemcpyAsync(h_c, d_c, chunk, cudaMemcpyDeviceToHost, stream0);

// Stream 1: 后半部分（与 Stream 0 并行！）
cudaMemcpyAsync(d_a + chunk/4, h_a + chunk/4, chunk,
                 cudaMemcpyHostToDevice, stream1);
kernel<<<blocks, threads, 0, stream1>>>(d_a + chunk/4, d_b + chunk/4,
                                         d_c + chunk/4, chunk);
cudaMemcpyAsync(h_c + chunk/4, d_c + chunk/4, chunk,
                 cudaMemcpyDeviceToHost, stream1);

// 等待所有 Stream 完成
cudaDeviceSynchronize();

// 总时间 = max(H2D, Kernel) + max(Kernel, D2H) < H2D + Kernel + D2H

Stream 优先级

cpp

// 设置 Stream 优先级
int priority_low, priority_high;
cudaDeviceGetStreamPriorityRange(&priority_low, &priority_high);
// priority_high > priority_low = 更高优先级

cudaStream_t stream_high, stream_low;
cudaStreamCreateWithPriority(&stream_high, cudaStreamNonBlocking, priority_high);
cudaStreamCreateWithPriority(&stream_low, cudaStreamNonBlocking, priority_low);

// 高优先级 Stream 的 kernel 更早被调度

第6节：Tensor Core 加速矩阵乘法

使用 cuBLAS 库

cpp

// 使用 cuBLAS 自动利用 Tensor Core
#include <cublas_v2.h>

cublasHandle_t cublas_handle;
cublasCreate(&cublas_handle);

// 使用 Tensor Core（FP16 矩阵乘法，自动选择最快实现）
cublasGemmEx(cublas_handle,
    CUBLAS_OP_N, CUBLAS_OP_N,  // 不转置
    N, M, K,                     // C[M×N] = A[M×K] × B[K×N]
    &alpha,                       // = 1.0
    d_A, CUDA_R_16F, K,          // A 矩阵
    d_B, CUDA_R_16F, N,          // B 矩阵
    &beta,                        // = 0.0（不累加）
    d_C, CUDA_R_16F, M,          // C 矩阵
    CUDA_R_16F,                   // 计算精度
    CUBLAS_GEMM_DEFAULT_TENSOR_OP  // 使用 Tensor Core
);

// cuBLAS 自动处理：
// 1. 选择最优的 kernel
// 2. 自动使用 Tensor Core（如果可用）
// 3. 内存布局优化

第7节：优化检查清单

┌─────────────────────────────────────────────────────────────┐
│                    CUDA 优化检查清单                          │
├─────────────────────────────────────────────────────────────┤
│                                                             │
│  [ ] 1. 内存访问                                           │
│      [ ] 合并访问：连续线程访问连续内存                      │
│      [ ] 共享内存复用：减少全局内存访问                     │
│      [ ] 向量化加载：float4/float2 一次读多个值             │
│      [ ] 避免 Bank 冲突                                    │
│                                                             │
│  [ ] 2. 计算优化                                           │
│      [ ] 使用 Tensor Core（cuBLAS/cuDNN）                  │
│      [ ] 减少分支分化（Warp 内 if-else）                   │
│      [ ] 使用 Warp Shuffle 代替共享内存做归约               │
│                                                             │
│  [ ] 3. 调度优化                                           │
│      [ ] Occupancy >= 50%（通常足够好）                   │
│      [ ] Block 维度 32 的倍数                              │
│      [ ] 使用 CUDA Stream 重叠数据传输和计算                │
│                                                             │
│  [ ] 4. 测量验证                                           │
│      [ ] 用 Nsight Compute/Nsight Systems profiling        │
│      [ ] 对比优化前后的实际 FLOPS                          │
│      [ ] 确认达到了 Roofline 上限                          │
│                                                             │
└─────────────────────────────────────────────────────────────┘

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

AI 可查：
✅ Tensor Core 的具体 API（cuBLAS/cuDNN 参数）
✅ CUDA Stream 的具体函数
✅ Nsight Compute 的具体命令选项

必须理解：
🔴 Roofline 模型：带宽受限 vs 算力受限
🔴 Occupancy = 实际线程数 / 最大线程数，影响延迟隐藏能力
🔴 合并访问：连续线程访问连续内存
🔴 Warp 分支分化导致 50% 效率损失
🔴 Warp Shuffle：同 Warp 内直接交换数据，无需共享内存
🔴 CUDA Stream：不同 Stream 的操作可以重叠（H2D 和 Kernel 并行）

学习状态：🟡 开始学习

面试速答

什么是合并访存？ 同一 Warp 的线程以有利地址模式访问内存，使硬件能用较少事务满足请求，从而提高有效带宽。

Shared Memory 为什么快？ 它是 SM 上的软件管理片上存储，延迟和带宽通常优于 Global Memory；但容量有限且需正确同步。

Occupancy 越高越好吗？ 不一定。达到足够隐藏延迟后，计算吞吐、内存带宽、指令依赖等可能主导。

自测

两个 Stream 是否必然同时执行？
寄存器 spill 为什么危险？
优化 Kernel 前为何要算传输时间？

CUDA Kernel、内存与性能 ​

1. Grid-stride loop ​

2. 合并访存 ​

3. Shared Memory：显式片上复用 ​

4. 寄存器、Shared Memory 与 Occupancy ​

5. Stream 与重叠 ​

6. 常见反优化 ​

深入原理与工程实践 ​

完整迁入：原 CUDA Kernel 与内存全文 ​

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

先抓住直觉 ​

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

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

五种内存的特点 ​

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

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

基本操作 ​

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

带宽计算 ​

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

为什么需要共享内存？ ​

共享内存 Bank 冲突 ​

__syncthreads()——线程同步 ​

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

什么是 Unified Memory？ ​

使用示例 ​

预取和亲和性 ​

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

显存估算 ​

显存优化技术 ​

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

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

完整迁入：原 CUDA 性能优化全文 ​

CUDA 性能优化——从 30 天缩短到 10 天 / CUDA Performance Optimization ​

先抓住直觉 ​

场景：你的 CUDA 代码跑得比预期慢 10 倍 ​

第1节：性能分析——先定位瓶颈 ​

屋顶线模型（Roofline Model） ​

CUDA 性能分析工具 ​

第2节：Occupancy——GPU 利用率的根本 ​

什么是 Occupancy？ ​

计算 Occupancy ​

提高 Occupancy ​

第3节：向量化加载——充分利用内存带宽 ​

什么是向量化加载？ ​

向量化代码示例 ​

第4节：Warp 级优化——分支和洗牌 ​

减少分支分化 ​

Warp 洗牌（Shuffle）指令 ​

第5节：CUDA Stream——异步并行 ​

什么是 CUDA Stream？ ​

Stream 使用示例 ​

Stream 优先级 ​

第6节：Tensor Core 加速矩阵乘法 ​

使用 cuBLAS 库 ​

第7节：优化检查清单 ​

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

面试速答 ​

自测 ​

CUDA Kernel、内存与性能

1. Grid-stride loop

2. 合并访存

3. Shared Memory：显式片上复用

4. 寄存器、Shared Memory 与 Occupancy

5. Stream 与重叠

6. 常见反优化

深入原理与工程实践

完整迁入：原 CUDA Kernel 与内存全文

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

先抓住直觉

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

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

五种内存的特点

从 CPU 到 GPU 的内存层次对比

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

基本操作

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

带宽计算

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

为什么需要共享内存？

共享内存 Bank 冲突

`__syncthreads()`——线程同步

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

什么是 Unified Memory？

使用示例

预取和亲和性

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

显存估算

显存优化技术

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

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

完整迁入：原 CUDA 性能优化全文

CUDA 性能优化——从 30 天缩短到 10 天 / CUDA Performance Optimization

先抓住直觉

场景：你的 CUDA 代码跑得比预期慢 10 倍

第1节：性能分析——先定位瓶颈

屋顶线模型（Roofline Model）

CUDA 性能分析工具

第2节：Occupancy——GPU 利用率的根本

什么是 Occupancy？

计算 Occupancy

提高 Occupancy

第3节：向量化加载——充分利用内存带宽

什么是向量化加载？

向量化代码示例

第4节：Warp 级优化——分支和洗牌

减少分支分化

Warp 洗牌（Shuffle）指令

第5节：CUDA Stream——异步并行

什么是 CUDA Stream？

Stream 使用示例

Stream 优先级

第6节：Tensor Core 加速矩阵乘法

使用 cuBLAS 库

第7节：优化检查清单

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

面试速答

自测