GPU 内核性能优化进阶：占用率、访存与指令级调优

一个「能跑」的 GPU 内核，和「把硬件吃到饱」的内核，性能差距可以高达 10-100 倍。现代 NVIDIA GPU 单 SM 内含数千个执行单元，需要足够多的并发 warp 来隐藏访存延迟、足够规整的数据流来喂饱吞吐单元、足够低的指令依赖来填满流水线。本文从占用率开始，系统拆解 GPU 内核调优的四大维度，并覆盖 cuBLAS/cuDNN/NCCL 这类库的黑盒调优路径。

1. Occupancy：占用率计算

占用率（Occupancy） = 单个 SM 上活跃的 warp 数 / SM 最大可容纳 warp 数。占用率决定了有多少并发 warp 可用于隐藏内存延迟和指令延迟。A100/H100 的 SM 最大 64 warp（2048 线程），但寄存器数与共享内存容量会共同压低实际占用率。

SM 资源预算（以 A100 为例）：
  max threads/SM   = 2048
  registers/SM     = 65536  （每线程最多 255）
  shared mem/SM    = 164 KB
  max blocks/SM    = 32

约束链：blockSize=256  → 最多 8 block（2048 线程）
        每线程 40 寄存器 → 65536/40 = 1638 线程 → 仍 6 block（1536 线程）
        每 block 用 32KB 共享内存 → 164/32 = 5 block → 占用率降为 5×256/2048=62.5%

三个约束取最小，这才是实际占用率。计算并优化的正确姿势是交给 CUDA 工具：

#include <cuda_runtime.h>
// 查询给定 blockSize 与资源下最多活跃 block 数
int maxBlocks;
cudaOccupancyMaxActiveBlocksPerMultiprocessor(
    &maxBlocks, kernel, blockSize, dynamicSMemSize);

// 或让库自动选最优 block 数
int blockSize;
cudaOccupancyMaxPotentialBlockSize(&minGrid, &blockSize, kernel,
                                   dynamicSMemSize, 0);

也可以直接用 NVIDIA 的占用率计算器（Spreadsheet）手工推演：给定每线程寄存器数、blockSize、共享内存用量，得出占用率表格。

关键权衡：占用率不是越高越好。寄存器换占用率的经典例子：

高占用率 + 低寄存器命中率，不如中占用率 + 高指令级并行。用 -maxrregcount 或 __launch_bounds__ 限制寄存器时，务必实测 IPC 而非只看占用率数字。

手工推演一张占用率表（A100，blockSize=256，共享内存每 block 8KB）：

资源数：maxThreads=2048, regs=65536, smem=164KB, maxBlocks=32
推演：
  threads 约束 → 2048/256  = 8 blocks
  regs   约束 → 65536/40  = 1638 线程 → 6 blocks（1536）
  smem   约束 → 164/8     = 20 blocks
  取最小 → 6 blocks → 占用率 6×256/2048 = 75%

推演结论：当前内核的瓶颈是每线程 40 个寄存器；把它压到 32 个即可让寄存器约束达到 8 blocks，占用率升到 100%。这就是手工推演的意义——知道该往哪调。

一句话：占用率是「用并发换延迟隐藏」的资源博弈，由寄存器、共享内存、block 数三个约束取最小决定，且要与 ILP 权衡。

2. 访存合并与共享内存

GPU 全局内存的带宽只有被**合并访问（Coalesced Access）**时才发挥得出来：同一个 warp 的 32 个线程访问同一 128 字节 cache line 内连续地址，硬件一次事务搬完；若地址分散在不同 cache line，则退化为 32 次事务，带宽断崖式下降。

合并访问（理想）：             非合并访问（灾难）：
线程 t0..t31 → 连续 128B       线程 t0..t31 → 随机 128B×N
   +------------------------+      +----+----+----+----+
   | 128-byte aligned line  |      | 散 | 乱 | 地 | 址 |
   +------------------------+      +----+----+----+----+
   1 次内存事务 → 32×4B 全利用        32 次事务 → 每次只用一个字节

正确写法：让连续线程访问连续地址。以下面的矩阵转置为例：

// 错误示范：线程 i 按列跳着读，非合并
__global__ void transpose_bad(const float* in, float* out, int n) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    out[col * n + row] = in[row * n + col];  // 读合并、写非合并
}

优化：用共享内存做 tile 转置，让读写两侧都合并：

#define TILE 32

__global__ void transpose_tile(const float* in, float* out, int n) {
    __shared__ float tile[TILE][TILE + 1];   // +1 padding 消除 bank conflict
    int x = blockIdx.x * TILE + threadIdx.x;
    int y = blockIdx.y * TILE + threadIdx.y;

    if (x < n && y < n)
        tile[threadIdx.y][threadIdx.x] = in[y * n + x];  // 合并读

    __syncthreads();

    x = blockIdx.y * TILE + threadIdx.x;
    y = blockIdx.x * TILE + threadIdx.y;
    if (x < n && y < n)
        out[y * n + x] = tile[threadIdx.x][threadIdx.y]; // 合并写
}

共享内存把「跨层级的随机访问」转成「片内转置 + 合并搬移」，是访存优化最常用的手法。更系统的访存分析可用 Nsight Compute 的 Memory Workload Analysis 视图确认合并率。

向量化访存（vectorized access）：当每个线程需要连续多个元素时，用 float4 让单条指令搬 16 字节，既减少指令数又保证合并：

// 每线程处理 4 个连续元素（合并 + 少指令）
const float4* a4 = reinterpret_cast<const float4*>(a);
const float4* b4 = reinterpret_cast<const float4*>(b);
float4* c4 = reinterpret_cast<float4*>(c);

int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n / 4) {
    float4 va = a4[i], vb = b4[i];
    float4 vc = { va.x + vb.x, va.y + vb.y, va.z + vb.z, va.w + vb.w };
    c4[i] = vc;
}

要求 n 是 4 的倍数且指针 16 字节对齐；尾部剩余元素用普通循环补。向量化与合并是「减少事务数」的同一目标的两面：合并保证每次事务不浪费，向量化保证每线程指令更少。

3. Bank Conflict：共享内存的性能陷阱

共享内存被划分为 32 个 bank（每 bank 4 字节，A100 后为 32-bit 粒度）。当同一个 warp 内两个线程同时访问同一 bank 的不同地址时，硬件被迫将访问串行化为多次事务——这就是 bank conflict。

bank:     0  1  2  3  ...  30  31
地址:    [0][4][8][12]...          ← 线程 t 访问地址 4t（无冲突）
地址:    [0][4][8][12]...          ← 线程 t 访问地址 8t（冲突：t 和 t+2 撞同 bank）

消除手法：

1. 列方向 padding：二维数组把列宽从 TILE 改为 TILE + 1，让相邻行的同列错开 bank。
2. 交错布局：把行优先改为斜对角映射（如 XOR 散列）。
3. 广播优化：所有线程读同一个地址时，硬件广播为一次事务——这不是冲突，是免费的。

// 用 padding 消除 bank conflict 的标准写法（TILE=32，pad=1）
__shared__ float sdata[32][33];
int idx = threadIdx.x;
float v = sdata[idx][threadIdx.y];   // 行间错位，无冲突

判断冲突：Nsight Compute 的 Shared Memory 指标会显示 bank conflicts per access；实测中，冲突严重的共享内存读写比全局未合并访问更难发现——因为代码看起来很正常，必须靠工具确认。

4. 指令级并行与 ILP

GPU 单线程的算力上限由指令级并行（ILP）决定：一个 warp 内的指令依赖链越长，流水线越空。提升 ILP 的核心手段是多重累加器（multiple accumulators）与循环展开（unrolling），让相互独立的指令穿插在依赖链之间。

// 低 ILP：归约链串行，每次累加依赖上一次
float sum = 0;
for (int i = tid; i < n; i += blockDim.x) {
    sum += a[i];           // sum 的依赖链 = n 次串行等待
}

// 高 ILP：4 路累加器打散依赖链，4 条链并行前进
float sum0 = 0, sum1 = 0, sum2 = 0, sum3 = 0;
for (int i = tid; i < n; i += 4 * blockDim.x) {
    sum0 += a[i];
    sum1 += a[i + blockDim.x];
    sum2 += a[i + 2 * blockDim.x];
    sum3 += a[i + 3 * blockDim.x];
}
float sum = (sum0 + sum1) + (sum2 + sum3);  // 尾部合并，仅 log 级依赖

编译器 -O3 会自动做部分展开，但循环内访存依赖（每次 a[i] 都是新地址）无法消除时，4 路累加器的收益依然可观。关键度量：用 Nsight Compute 看 Issue Slots（每周期发射槽）与 ILP 统计，若指令发射率低而依赖停顿高，就继续加累加器；若发射率已接近上限，再展开只会浪费寄存器。

寄存器与 ILP 的联动：多重累加器消耗寄存器，可能在占用率上打折扣。实践上先测「8 block 低 ILP」与「4 block 高 ILP」两个配置的吞吐，选快的那个——这就是第 1 节强调的权衡。

5. Warp Divergence：分支的代价

GPU 是 SIMT（单指令多线程）：一个 warp 的 32 个线程共享一条指令流。if/else 造成两路线程走不同分支时，硬件先执行一路、屏蔽另一路，再执行第二路——串行化 + 屏蔽开销。

// 分歧：偶数/奇数线程走不同路径 → 两轮执行
if (tid % 2 == 0) { costly_even(); } else { costly_odd(); }

消除分歧的策略：

1. 按 warp 对齐数据：把数据按奇偶重排，使每个 warp 内同质（见下例）。对「每线程处理一个格子、某些格子需特殊处理」的问题尤其有效。
2. 掩码函数（predication）：短分支用三元表达式或 ?: 让编译器生成谓词指令（两条都执行，用谓词选择写回）——短分支下谓词比分支快。
3. 分歧上移/下移：把分歧从「每个线程内」移到「block 级别」——用 if (threadIdx.x < threshold) 在入口处一次性分流。

// 消除分歧的经典做法：先压缩，后统一处理
// 把「需特殊处理的元素」复制到 warp 内连续位置（stream compaction），
// 使每个 warp 要么全特殊、要么全普通，然后分别整块处理
bool flag = need_special(idx);
__shared__ int bucket;
if (flag) special_list[atomicAdd(&bucket, 1)] = idx;

分歧的度量工具：perf/ncu 的 Branch 指标，或编译器报告中的 divergent branches。注意：分歧只影响同一 warp 内部——不同 warp 之间的分支各自独立，完全并行。

一句话：GPU 性能密码 = 合并访存（数据流规整）+ 消除冲突（地址错位）+ 提升 ILP（打断依赖链）+ 消除分歧（warp 内同质）。

6. cuBLAS/cuDNN 调优参数

使用 cuBLAS/cuDNN 这类库时，性能关键不在代码而在算法选择（heuristics）与数学模式。这些库为同一运算提供多个算法内核，选错可以差 5-10 倍。

cuBLAS（GEMM 示例）：

cublasHandle_t h;
cublasCreate(&h);

// 1. 数学模式：允许 TF32 张量核心（A100/H100 上关键加速）
cublasSetMathMode(h, CUBLAS_TF32_TENSOR_OP_MATH);

// 2. 用 cublasGemmEx 指定数据类型与算法
cublasGemmAlgo_t algo = CUBLAS_GEMM_DEFAULT_TENSOR_OP;  // 自动选
cublasGemmEx(h,
    CUBLAS_OP_N, CUBLAS_OP_N,
    m, n, k,
    &alpha, A, CUDA_R_16F, lda,
             B, CUDA_R_16F, ldb,
    &beta,  C, CUDA_R_16F, ldc,
    CUDA_R_16F, algo);

// 3. 对固定形状做 autotune：暴力尝试所有算法
for (int a = CUBLAS_GEMM_DEFAULT; a <= CUBLAS_GEMM_ALGO14_TENSOR_OP; ++a) {
    // 逐个计时，选最快的保存
}

cuDNN（卷积示例）：

// 1. 允许 Tensor Core
cudnnSetConvolutionMathType(convDesc, CUDNN_TENSOR_OP_MATH);

// 2. 自动挑最快的 forward 算法（heuristic）
int nAlgo = 0;
cudnnGetConvolutionForwardAlgorithm_v7(handle, srcDesc, filtDesc,
    convDesc, dstDesc, CUDNN_CONVOLUTION_FWD_ALGO_COUNT, &nAlgo, perfResults);
// perfResults[0].algo 即为推荐，algo 之间性能差异可达数倍

// 3. 处理 workspace 大小
size_t wsSize;
cudnnGetConvolutionForwardWorkspaceSize(handle, ..., perfResults[0].algo, &wsSize);
cudaMalloc(&workspace, wsSize);

库调优通用流程：固定形状（batch、通道、分辨率）+ 开启 Tensor Core 数学模式 + autotune 选算法 + 分配足量 workspace。对推理服务，建议把 autotune 结果固化（缓存最优算法与 workspace），避免每次启动重跑。

7. NCCL 通信优化：多 GPU 的咽喉

多 GPU 训练/仿真的通信库是 NCCL。通信优化的核心是隐藏延迟与选对算法：

# 常见 NCCL 环境变量调优
export NCCL_IB_DISABLE=0        # 启用 InfiniBand/RoCE RDMA
export NCCL_SOCKET_IFNAME=ib0   # 指定网卡（避免误用以太网口）
export NCCL_P2P_DISABLE=0       # 同节点 NVLink 直连（默认开启）
export NCCL_BUFFSIZE=16777216   # 消息缓冲（大消息调大）
export NCCL_ALGO=Ring           # 显式选 Ring/Tree 算法
export NCCL_PROTO=Simple        # Simple / LL / LL128 协议
export NCCL_DEBUG=INFO          # 排查时开启日志

算法选择的经验法则：

NCCL 性能验证：nccl-tests 的 all_reduce_perf 给出总线带宽与延迟：

mpirun -np 8 all_reduce_perf -b 8 -e 256M -f 2 -g 1
# 观察 "Out of place" 带宽是否接近硬件上限（NVLink ~600 GB/s, IB 400G）

若带宽远低于上限：检查网卡（nvidia-smi topo -m 确认 P2P 可达）、NCCL_DEBUG 日志中的传输选择、以及是否有 CPU 拓扑竞争。

# 确认 GPU 间拓扑：PIX/NV 表示 NVLink 直连，PXB/PXB 表示经 PCIe 交换机
nvidia-smi topo -m
#        GPU0   GPU1   GPU2   GPU3
# GPU0   NV     NV     PIX    PIX
# ...

通信与计算的重叠，配合 Nsight Systems 的 timeline 判断瓶颈在哪一段。DDP/AllReduce 优化序列：先用 nccl-tests 验证单条链路，再用 nsys 看训练步内通信占比，最后通过梯度压缩（fp16/bf16）、通信与反向计算重叠（torch.compile 自动调度）降总耗时。

8. 基准对比方法论：让优化可见

优化的最终裁判是可重复的基准。GPU 内核基准要处理三重陷阱：首次调用初始化、CPU-GPU 同步、波动性。

# Python + PyTorch 风格的内核基准（CUDA C 同理）
import torch, time

def bench(fn, args, warmup=10, iters=100):
    # warmup：触发 CUDA 上下文初始化、kernel 编译缓存
    for _ in range(warmup): fn(*args)
    torch.cuda.synchronize()

    # 正式计时：必须在时间点前后同步，否则算的是「异步提交」的时间
    t0 = time.perf_counter()
    for _ in range(iters): fn(*args)
    torch.cuda.synchronize()
    return (time.perf_counter() - t0) / iters

# 用例：对比两种内核
t_tile   = bench(lambda: tile_kernel[grid, block](x, y, n), ())
t_naive  = bench(lambda: naive_kernel[grid, block](x, y, n), ())
print(f"tile {t_tile:.3f} ms, naive {t_naive:.3f} ms, speedup {t_naive/t_tile:.2f}x")

基准方法论要点：

1. 统一输入与形状：所有对比用同一批数据、同一 shape，避免 cache 效应。
2. 报告几何/多次中位数：取多次运行的中位数而非平均，排除调度抖动。
3. 验证正确性：先核对 torch.allclose(ref, out)，性能对比建立在正确性之上。
4. 记录环境：GPU 型号、驱动、CUDA 版本、时钟频率（nvidia-smi -q -d CLOCK）一并存档。
5. 用 profiler 佐证：基准数字 + Nsight Compute 指标（吞吐、带宽、占用率）交叉验证「为什么快/慢」。

一句话：GPU 优化是「假设 → 改内核 → 基准 → 工具佐证」的迭代循环；没有正确性与可重复性的性能数字毫无意义。

总结

GPU 内核调优没有银弹——它是占用率、访存、指令级并行、通信四者的联合优化。掌握了本文的度量工具与调优手法治，配合 CUDA + MPI 异构并行 的多 GPU 扩展知识，你就能把一块 GPU 乃至一个 GPU 集群的算力潜力系统性地榨干。异构平台的性能分析工具可参考 性能剖析与调优工具链。