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 并行）

学习状态：🟡 开始学习

原体系结构与硬件 / Original Architecture

原并行计算 / Original Parallel Computing

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 性能优化——从 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 性能优化——从 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 必须理解"清单