cuda-c-optimization

CUDA C 性能优化指南

1. 性能优化策略

1.1 块大小选择策略

• 基础: 使用 2 的幂（128, 256, 512, 1024）
• 推荐: 256 或 512 线程每块
• 限制: 每块最多 1024 线程（大多数 GPU）
• 调优: 平衡并行度与资源占用，避免过大或过小

算子类型	推荐块大小	网格配置
Element-wise	256 / 512	一维
Reduce	256	一维 + 共享内存
MatMul	dim3(16,16) 或 dim3(32,32)	二维
图像处理	dim3(16,16)	二维

1.2 内存访问优化

合并访问 (Coalesced Access)

连续线程访问连续内存地址，GPU 将多次请求合并为少量内存事务。

// ✅ 合并访问（连续线程访问连续地址）
__global__ void coalesced(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        data[idx] = data[idx] * 2.0f;  // 连续访问
    }
}

// ❌ 非合并访问（跳跃访问）
__global__ void strided(float* data, int n, int stride) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx * stride < n) {
        data[idx * stride] = data[idx * stride] * 2.0f;  // 跳跃访问
    }
}

对齐访问

数据按 128 字节边界对齐，提高内存带宽利用率。

避免 Bank 冲突

共享内存由 32 个 bank 组成，避免同一 warp 内多个线程访问同一 bank。

// ✅ 无 bank 冲突
__shared__ float s[256];
s[threadIdx.x] = input[idx];  // 连续线程访问连续 bank

// ❌ bank 冲突
s[threadIdx.x * 32] = input[idx];  // 所有线程访问同一 bank

1.3 计算优化

避免分支发散

同一 warp 内的 32 个线程应执行相同的控制路径。

// ❌ 分支发散：同一 warp 内线程走不同路径
if (threadIdx.x % 2 == 0) {
    // 偶数线程路径
} else {
    // 奇数线程路径
}

// ✅ 使用条件赋值替代分支
float result = (threadIdx.x % 2 == 0) ? value_a : value_b;

使用内置快速数学函数

// 标准精度
float r = expf(x);

// ✅ 快速版本（精度略低但速度更快）
float r = __expf(x);
float r = __logf(x);
float r = __sinf(x);

减少原子操作

尽量使用块内归约代替全局原子操作。

// ❌ 大量原子操作
atomicAdd(&global_sum, local_val);

// ✅ 先块内归约，再原子写回
__shared__ float sdata[256];
sdata[tid] = local_val;
__syncthreads();

// 块内归约
for (int s = blockDim.x / 2; s > 0; s >>= 1) {
    if (tid < s) sdata[tid] += sdata[tid + s];
    __syncthreads();
}

// 只有一次原子操作
if (tid == 0) atomicAdd(&global_sum, sdata[0]);

1.4 Occupancy 优化

• 寄存器使用: 减少每个线程的寄存器使用量，增加并发 block 数
• 共享内存: 合理使用共享内存，不超过硬件限制
• Block 大小: 选择能整除 SM 最大线程数的 block 大小