第4章-优化技巧

第4章：Kernel 优化技巧

掌握内存访问优化和计算优化，让你的 CUDA 程序飞起来

本章定位

本章负责把优化技巧串成可执行的判断流程：先测量，再判断 memory-bound、compute-bound 或资源受限，然后只改一个因素并复测。具体指标解释以 Nsight Compute NCU 分析方法与优化思路 为准。

配套主文档：

• CUDA Shared Memory 与 Bank Conflict
• CUDA 线程配置与占用率
• Nsight Compute NCU 分析方法与优化思路
• GPU 初学者术语表

学习目标

• 掌握内存访问优化（合并访问、共享内存、Bank Conflict）
• 掌握计算优化技术（循环展开、Warp 级操作）
• 理解流水线和指令级优化
• 学会使用性能分析工具定位瓶颈

---

4.1 内存访问优化

4.1.1 合并内存访问

什么是合并访问？

当 Warp 中的 32 个线程访问连续的内存地址时，GPU 可以一次性读取所有数据，这就是合并访问。

合并访问（高效）:
线程:  0   1   2   3   4   5   6   7
地址:  0   4   8   12  16  20  24  28  ✓ 连续访问

跨步访问（低效）:
线程:  0   1   2   3   4   5   6   7
地址:  0   32  64  96  128 160 192 224 ✗ 间隔访问

代码示例：

// ❌ 不好的访问模式：跨步访问
__global__ void bad_access(float* data, int stride) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    data[idx * stride] = 0;  // 跨步访问，效率低
}
 
// ✅ 优化后的访问模式：合并访问
__global__ void good_access(float* data) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    data[idx] = 0;  // 连续访问，效率高
}

4.1.2 共享内存优化

共享内存的特点：

• 位于 SM 上，访问速度比全局内存快 10-20 倍
• 容量有限（通常 48-164KB per SM）
• 需要显式管理（手动加载和存储）

Bank Conflict 问题：

共享内存被分成 32 个 Bank，每个 Bank 每个时钟周期只能服务一个请求。如果多个线程同时访问同一个 Bank，就会发生 Bank Conflict，导致串行访问。

// ❌ 产生 Bank Conflict
__shared__ float data[32];
// 所有线程访问同一列，导致 32-way bank conflict
float val = data[threadIdx.x * 32];  // 32个线程访问同一 bank
 
// ✅ 无 Bank Conflict
float val = data[threadIdx.x];  // 每个线程访问不同 bank

Padding 技术解决 Bank Conflict：

// ❌ 可能产生 bank conflict（当访问列时）
__shared__ float matrix[32][32];
float val = matrix[threadIdx.x][threadIdx.y];  // 列访问冲突
 
// ✅ 使用 padding 避免 bank conflict
__shared__ float matrix[32][33];  // 多一列用于 padding
// 现在相邻列在不同 bank 中
float val = matrix[threadIdx.x][threadIdx.y];  // 无冲突

4.1.3 内存访问模式对比

模式	效率	示例	适用场景
合并访问	最高	data[tid]	顺序数据处理
跨步访问	较低	data[tid * stride]	矩阵行/列访问
随机访问	最低	data[index[tid]]	查找表、稀疏数据

4.1.4 内存预取

__global__ void prefetch_example(float* data, float* output, int n) {
    __shared__ float cache[256];
 
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int cache_idx = threadIdx.x;
 
    // 预取下一块数据到共享内存
    if (idx + blockDim.x < n) {
        cache[cache_idx] = data[idx + blockDim.x];
    }
    __syncthreads();
 
    // 处理当前数据（同时硬件在加载下一批）
    if (idx < n) {
        output[idx] = cache[cache_idx] * 2;
    }
}

---

4.2 计算优化

4.2.1 减少寄存器使用

寄存器是 GPU 上最快的内存，但数量有限。过多的寄存器使用会降低占用率（Occupancy）。

// ❌ 过多寄存器使用
__global__ void high_reg_count(float* data, int n) {
    float a[32], b[32], c[32];  // 太多寄存器
    // ... 复杂计算
}
 
// ✅ 减少寄存器使用
__global__ void low_reg_count(float* data, int n) {
    float a, b, c;  // 复用变量，减少寄存器压力
    // ... 逐元素处理
}

4.2.2 循环展开

循环展开可以减少循环开销，增加指令级并行。

// ❌ 不展开，每次循环都有判断开销
for (int i = 0; i < 4; i++) {
    sum += data[i];
}
 
// ✅ 手动展开，无循环开销
sum += data[0];
sum += data[1];
sum += data[2];
sum += data[3];
 
// ✅ 编译器自动展开（推荐）
#pragma unroll
for (int i = 0; i < 4; i++) {
    sum += data[i];
}
 
// ✅ 条件展开（根据模板参数）
#pragma unroll 4
for (int i = 0; i < TILE_SIZE; i++) {
    sum += As[ty][i] * Bs[i][tx];
}

4.2.3 Warp 级原语操作

Warp 级操作可以在线程之间快速交换数据，无需共享内存。

// Warp 级归约（求和）
__device__ float warp_reduce_sum(float val) {
    for (int offset = 16; offset > 0; offset >>= 1) {
        val += __shfl_down_sync(0xffffffff, val, offset);
    }
    return val;
}
 
// Warp 级最大值归约
__device__ float warp_reduce_max(float val) {
    for (int offset = 16; offset > 0; offset >>= 1) {
        val = fmaxf(val, __shfl_down_sync(0xffffffff, val, offset));
    }
    return val;
}
 
// 使用示例：向量点积
__global__ void dot_product(float* a, float* b, float* result, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float sum = 0.0f;
 
    for (int i = idx; i < n; i += blockDim.x * gridDim.x) {
        sum += a[i] * b[i];
    }
 
    // Warp 内归约
    sum = warp_reduce_sum(sum);
 
    // Warp 0 的 lane 0 写回结果
    if (threadIdx.x % 32 == 0) {
        atomicAdd(result, sum);
    }
}

4.2.4 使用 Tensor Core 加速

Tensor Core 是专门用于矩阵运算的硬件单元，可以大幅加速 GEMM 运算。

#include <mma.h>
using namespace nvcuda::wmma;
 
#define WMMA_M 16
#define WMMA_N 16
#define WMMA_K 16
 
__global__ void matmul_tensor_core(half* A, half* B, float* C,
                                    int M, int N, int K) {
    // 声明 Tensor Core 片段
    fragment<matrix_a, WMMA_M, WMMA_N, WMMA_K, half, row_major> a_frag;
    fragment<matrix_b, WMMA_M, WMMA_N, WMMA_K, half, col_major> b_frag;
    fragment<accumulator, WMMA_M, WMMA_N, WMMA_K, float> c_frag;
 
    // 初始化累加器
    fill_fragment(c_frag, 0.0f);
 
    // 遍历 K 维度
    for (int k = 0; k < K; k += WMMA_K) {
        // 加载矩阵片段
        load_matrix_sync(a_frag, A + blockIdx.y * WMMA_M * K + k, K);
        load_matrix_sync(b_frag, B + k * N + blockIdx.x * WMMA_N, N);
 
        // 执行矩阵乘加
        mma_sync(c_frag, a_frag, b_frag, c_frag);
    }
 
    // 存储结果
    store_matrix_sync(C + blockIdx.y * WMMA_M * N + blockIdx.x * WMMA_N,
                      c_frag, N, mem_row_major);
}

---

4.3 流水线优化

4.3.1 双缓冲（Double Buffering）

双缓冲可以让计算和内存传输重叠，隐藏内存延迟。

__global__ void double_buffer(float* in, float* out, int n) {
    __shared__ float buffer[2][256];
    int cur = 0, next = 1;
 
    // 预加载第一块
    buffer[cur][threadIdx.x] = in[threadIdx.x];
    __syncthreads();
 
    for (int i = 0; i < n / blockDim.x - 1; i++) {
        int idx = (i + 1) * blockDim.x + threadIdx.x;
 
        // 异步加载下一块（加载和计算重叠）
        if (idx < n) {
            buffer[next][threadIdx.x] = in[idx];
        }
 
        // 处理当前数据
        float result = buffer[cur][threadIdx.x] * 2;
 
        // 等待加载完成
        __syncthreads();
 
        // 写回结果
        out[i * blockDim.x + threadIdx.x] = result;
 
        // 交换缓冲区
        cur = next;
        next = 1 - next;
    }
 
    // 处理最后一块
    int last_idx = (n / blockDim.x - 1) * blockDim.x + threadIdx.x;
    if (last_idx < n) {
        out[last_idx] = buffer[cur][threadIdx.x] * 2;
    }
}

4.3.2 异步拷贝（CUDA 11+）

#include <cuda/pipeline>
#include <cooperative_groups.h>
#include <cooperative_groups/memcpy_async.h>
 
namespace cg = cooperative_groups;
 
__global__ void async_copy(float* src, float* dst, int n) {
    __shared__ float buffer[256];
 
    // 单 stage、单线程级 pipeline 的标准 libcu++ 写法
    auto block = cg::this_thread_block();
    __shared__ cuda::pipeline_shared_state<cuda::thread_scope_block, 2> pss;
    auto pipe = cuda::make_pipeline(block, &pss);
 
    // 异步加载到共享内存（producer 端）
    pipe.producer_acquire();
    cuda::memcpy_async(block, buffer,
                       src + blockIdx.x * blockDim.x,
                       blockDim.x * sizeof(float), pipe);
    pipe.producer_commit();
 
    // 消费端：等待最久未释放的那个 stage 完成
    pipe.consumer_wait();
 
    // 处理数据
    buffer[threadIdx.x] *= 2;
    __syncthreads();
 
    // 释放该 stage，供 producer 复用
    pipe.consumer_release();
 
    // 写回全局内存
    dst[blockIdx.x * blockDim.x + threadIdx.x] = buffer[threadIdx.x];
}

常见误写：cuda::pipeline_commit(pipe) / cuda::pipeline_wait_prior<0>(pipe) 这种”自由函数 + 模板下标”的伪 API 不存在于真实的 libcu++ 头文件中。cuda::pipeline 的正确接口是成员函数 producer_acquire/commit 与 consumer_wait/release，需配合 cuda::pipeline_shared_state 与 cuda::make_pipeline() 使用。更多细节见 libcu++ pipeline 文档。

---

4.4 指令级优化

4.4.1 减少分支分歧

分支分歧会导致 Warp 内线程串行执行不同分支。

// ❌ 严重的分支分歧
if (threadIdx.x % 2 == 0) {
    result = a + b;
} else {
    result = a - b;
}
 
// ✅ 使用条件表达式减少分歧
float add = a + b;
float sub = a - b;
result = (threadIdx.x % 2 == 0) ? add : sub;
 
// ✅ 按 Warp 对齐处理（推荐）
int warp_id = threadIdx.x / 32;
if (warp_id % 2 == 0) {  // 整个 Warp 走同一分支
    result = a + b;
} else {
    result = a - b;
}

4.4.2 使用快速数学函数

// 标准函数（精度高）
float r = sqrtf(x);
float e = expf(x);
float s = sinf(x);
 
// 快速函数（精度略低，速度快）
float r = __fsqrt_rn(x);   // 快速平方根
float e = __expf(x);        // 快速指数
float s = __sinf(x);        // 快速正弦
 
// 内联 PTX 汇编（最高控制）
float r;
asm("sqrt.approx.ftz.f32 %0, %1;" : "=f"(r) : "f"(x));

---

4.5 性能分析工具

4.5.1 Nsight Compute

Nsight Compute 提供详细的 Kernel 级性能分析。

# 详细性能分析
ncu --set full -o report ./my_kernel
 
# 关键指标分析
ncu --metrics \
    smsp__cycles_active.avg.pct_of_peak,\
    dram__throughput.avg.pct_of_peak,\
    l1tex__t_requests.avg.pct_of_peak_rate \
    ./my_kernel
 
# 常用指标说明：
# - achieved_occupancy: 实际占用率
# - gld_efficiency: 全局加载效率
# - gst_efficiency: 全局存储效率
# - shared_efficiency: 共享内存效率

4.5.2 Nsight Systems

Nsight Systems 提供系统级时间线分析。

# 采集时间线
nsys profile -o report ./my_program
 
# 查看时间线（图形界面）
nsys-ui report.nsys-rep
 
# 可以看到：
# - Kernel 执行时间
# - CPU-GPU 数据传输
# - 内存分配/释放
# - 多流并行情况

4.5.3 CUDA Profiler（旧版）

# 使用 nvprof（已弃用，但仍可用）
nvprof ./my_program
 
# 指定指标
nvprof --metrics achieved_occupancy,gld_efficiency ./my_program

---

4.6 优化检查清单

内存优化

• 检查是否为合并访问模式
• 检查共享内存是否存在 Bank Conflict
• 评估共享内存使用量是否合理
• 检查寄存器使用量（--ptxas-options=-v）
• 考虑使用常量内存存储只读数据

计算优化

• 检查占用率是否达到预期（> 50%）
• 检查是否存在分支分歧
• 考虑循环展开（#pragma unroll）
• 考虑使用 Warp 级原语
• 评估是否可以使用 Tensor Core

流水线优化

• 考虑使用双缓冲隐藏延迟
• 考虑使用异步拷贝（CUDA 11+）
• 检查内存延迟是否被隐藏
• 评估多流并行

---

4.7 实战练习

练习 1：矩阵转置优化

实现一个高效的矩阵转置 Kernel，考虑：

1. 使用共享内存避免非合并访问
2. 使用 Padding 避免 Bank Conflict
3. 对比不同分块大小的性能

练习 2：归约优化

实现一个高效的数组求和 Kernel：

1. 使用 Warp 级归约
2. 使用共享内存进行 Block 级归约
3. 达到理论带宽的 80% 以上

练习 3：Softmax 优化

实现一个高性能 Softmax：

1. 使用在线算法避免数值溢出
2. 使用 Warp 级原语进行归约
3. 分析内存瓶颈并优化

---

💡 本章要点

1. 合并访问是内存优化的关键，相邻线程访问相邻地址
2. 共享内存很快但要注意 Bank Conflict，用 Padding 解决
3. 循环展开和 Warp 级原语可以显著提升计算效率
4. Tensor Core 可以带来 3-8 倍的矩阵运算加速
5. 双缓冲和 异步拷贝可以隐藏内存延迟
6. 减少分支分歧，尽量让 Warp 内线程走相同路径
7. 使用 Nsight Compute 和 Nsight Systems 分析性能

---

参考指标

优化目标	优秀	良好	需优化
全局内存效率	> 80%	50-80%	< 50%
共享内存效率	> 90%	70-90%	< 70%
占用率	> 70%	50-70%	< 50%
Warp 执行效率	> 90%	70-90%	< 70%

---

← 上一章：第 3 章 GPU 硬件 | 下一章：第 5 章 基础 Kernel →