性能优化

性能优化是 CUDA 编程的核心目标。本章将介绍内存访问优化、计算优化和占用率优化等关键技术。

内存访问优化

合并访问

合并访问是指 Warp 中的线程访问连续的内存地址，这是内存优化的基础。

__global__ void coalescedAccess(float *data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (idx < n) {
        data[idx] = data[idx] * 2.0f;
    }
}

非合并访问示例：

__global__ void stridedAccess(float *data, int n, int stride) {
    int idx = (blockIdx.x * blockDim.x + threadIdx.x) * stride;
    
    if (idx < n) {
        data[idx] = data[idx] * 2.0f;
    }
}

共享内存优化

使用共享内存减少全局内存访问：

#define TILE_SIZE 16

__global__ void matrixMulShared(float *A, float *B, float *C, int width) {
    __shared__ float sA[TILE_SIZE][TILE_SIZE];
    __shared__ float sB[TILE_SIZE][TILE_SIZE];
    
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    
    float sum = 0.0f;
    
    for (int t = 0; t < width / TILE_SIZE; t++) {
        sA[threadIdx.y][threadIdx.x] = A[row * width + t * TILE_SIZE + threadIdx.x];
        sB[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * width + col];
        
        __syncthreads();
        
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += sA[threadIdx.y][k] * sB[k][threadIdx.x];
        }
        
        __syncthreads();
    }
    
    C[row * width + col] = sum;
}

避免 Bank 冲突

共享内存被划分为 32 个 Bank，不当的访问模式会导致 Bank 冲突：

#define TILE_SIZE 32

__global__ void avoidBankConflict(float *data) {
    __shared__ float sData[TILE_SIZE][TILE_SIZE + 1];
    
    int idx = threadIdx.x;
    sData[idx][idx] = data[idx];
    
    __syncthreads();
    
    data[idx] = sData[idx][idx];
}

添加 +1 填充可以避免 Bank 冲突。

使用只读缓存

对于只读数据，使用 __ldg 或纹理内存：

__global__ void readOnlyCache(const float * __restrict__ input, 
                              float *output, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (idx < n) {
        output[idx] = __ldg(&input[idx]) * 2.0f;
    }
}

计算优化

减少分支分化

Warp 内的分支分化会降低性能：

__global__ void badBranch(int *data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (idx < n) {
        if (data[idx] > 0) {
            data[idx] = data[idx] * 2;
        } else {
            data[idx] = data[idx] + 1;
        }
    }
}

__global__ void goodBranch(int *data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (idx < n) {
        int cond = (data[idx] > 0);
        data[idx] = cond ? data[idx] * 2 : data[idx] + 1;
    }
}

循环展开

使用 #pragma unroll 展开循环：

__global__ void loopUnroll(float *data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (idx < n) {
        float sum = 0.0f;
        
        #pragma unroll 8
        for (int i = 0; i < 32; i++) {
            sum += data[idx + i * blockDim.x];
        }
        
        data[idx] = sum;
    }
}

使用内联函数

__device__ __forceinline__ float fastSqrt(float x) {
    return sqrtf(x);
}

__global__ void kernel(float *data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (idx < n) {
        data[idx] = fastSqrt(data[idx]);
    }
}

使用快速数学函数

CUDA 提供了快速但精度较低的数学函数：

标准函数	快速函数	说明
`sin(x)`	`__sinf(x)`	快速正弦
`cos(x)`	`__cosf(x)`	快速余弦
`exp(x)`	`__expf(x)`	快速指数
`log(x)`	`__logf(x)`	快速对数
`sqrt(x)`	`__sqrtf(x)`	快速平方根

__global__ void fastMath(float *data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (idx < n) {
        data[idx] = __sinf(data[idx]) + __cosf(data[idx]);
    }
}

编译选项：-use_fast_math

占用率优化

什么是占用率？

占用率 = 活跃 Warp 数 / 最大 Warp 数

高占用率可以隐藏内存延迟，但不是越高越好。

计算占用率

cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);

int maxThreadsPerSM = prop.maxThreadsPerMultiProcessor;
int maxBlocksPerSM = prop.maxBlocksPerMultiProcessor;

int threadsPerBlock = 256;
int blocksPerSM = min(maxBlocksPerSM, maxThreadsPerSM / threadsPerBlock);

float occupancy = (float)(blocksPerSM * threadsPerBlock) / maxThreadsPerSM;
printf("Occupancy: %.2f%%\n", occupancy * 100);

使用 CUDA Occupancy Calculator

NVIDIA 提供 CUDA Occupancy Calculator 工具，帮助选择最优的 Block 大小。

使用 cudaOccupancyMaxActiveBlocksPerMultiprocessor

int blockSize = 256;
int numBlocks;

cudaOccupancyMaxActiveBlocksPerMultiprocessor(
    &numBlocks, kernel, blockSize, 0);

printf("Max blocks per SM: %d\n", numBlocks);
printf("Max threads per SM: %d\n", numBlocks * blockSize);

优化寄存器使用

过多的寄存器使用会降低占用率：

__global__ void __launch_bounds__(256, 4) kernel() {
}

优化共享内存使用

__global__ void __launch_bounds__(256) kernel() {
    extern __shared__ float sharedData[];
}

内存带宽优化

使用对齐内存访问

size_t pitch;
cudaMallocPitch(&d_data, &pitch, width * sizeof(float), height);

cudaMemcpy2D(d_data, pitch, h_data, width * sizeof(float),
             width * sizeof(float), height, cudaMemcpyHostToDevice);

使用合并访问模式

__global__ void coalescedCopy(float *dst, float *src, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (idx < n) {
        dst[idx] = src[idx];
    }
}

使用向量加载

__global__ void vectorLoad(float4 *dst, float4 *src, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (idx < n) {
        dst[idx] = src[idx];
    }
}

指令级优化

使用 FMA（融合乘加）

__global__ void fmaExample(float *a, float *b, float *c, float *d, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (idx < n) {
        d[idx] = a[idx] * b[idx] + c[idx];
    }
}

编译器会自动使用 FMA 指令。

减少类型转换

__global__ void avoidConversion(float *data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (idx < n) {
        float val = data[idx];
        int intVal = __float2int_rn(val);
        data[idx] = __int2float_rn(intVal);
    }
}

性能分析工具

nvprof

nvprof ./my_program

nvprof --print-gpu-trace ./my_program

Nsight Compute

ncu ./my_program

ncu --set full ./my_program

Nsight Systems

nsys profile ./my_program

nsys profile --stats=true ./my_program

优化示例：矩阵乘法

基础版本

__global__ void matrixMulNaive(float *A, float *B, float *C, int width) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    
    float sum = 0.0f;
    
    for (int k = 0; k < width; k++) {
        sum += A[row * width + k] * B[k * width + col];
    }
    
    C[row * width + col] = sum;
}

共享内存版本

#define TILE_SIZE 16

__global__ void matrixMulShared(float *A, float *B, float *C, int width) {
    __shared__ float sA[TILE_SIZE][TILE_SIZE];
    __shared__ float sB[TILE_SIZE][TILE_SIZE];
    
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    
    float sum = 0.0f;
    
    for (int t = 0; t < width / TILE_SIZE; t++) {
        sA[threadIdx.y][threadIdx.x] = A[row * width + t * TILE_SIZE + threadIdx.x];
        sB[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * width + col];
        
        __syncthreads();
        
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += sA[threadIdx.y][k] * sB[k][threadIdx.x];
        }
        
        __syncthreads();
    }
    
    C[row * width + col] = sum;
}

优化版本（避免 Bank 冲突）

#define TILE_SIZE 32

__global__ void matrixMulOptimized(float *A, float *B, float *C, int width) {
    __shared__ float sA[TILE_SIZE][TILE_SIZE + 1];
    __shared__ float sB[TILE_SIZE][TILE_SIZE + 1];
    
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    
    float sum = 0.0f;
    
    for (int t = 0; t < width / TILE_SIZE; t++) {
        sA[threadIdx.y][threadIdx.x] = A[row * width + t * TILE_SIZE + threadIdx.x];
        sB[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * width + col];
        
        __syncthreads();
        
        #pragma unroll
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += sA[threadIdx.y][k] * sB[k][threadIdx.x];
        }
        
        __syncthreads();
    }
    
    C[row * width + col] = sum;
}

小结

本章介绍了 CUDA 性能优化的关键技术：

内存访问优化：合并访问、共享内存、避免 Bank 冲突
计算优化：减少分支分化、循环展开、快速数学函数
占用率优化：平衡寄存器和共享内存使用
内存带宽优化：对齐访问、向量加载
性能分析工具：nvprof、Nsight Compute、Nsight Systems

性能优化是一个迭代过程，需要不断分析和调整。下一章将介绍高级特性，学习动态并行、统一内存等高级功能。

内存访问优化​

合并访问​

共享内存优化​

避免 Bank 冲突​

使用只读缓存​

计算优化​

减少分支分化​

循环展开​

使用内联函数​

使用快速数学函数​

占用率优化​

什么是占用率？​

计算占用率​

使用 CUDA Occupancy Calculator​

使用 cudaOccupancyMaxActiveBlocksPerMultiprocessor​

优化寄存器使用​

优化共享内存使用​

内存带宽优化​

使用对齐内存访问​

使用合并访问模式​

使用向量加载​

指令级优化​

使用 FMA（融合乘加）​

减少类型转换​

性能分析工具​

nvprof​

Nsight Compute​

Nsight Systems​

优化示例：矩阵乘法​

基础版本​

共享内存版本​

优化版本（避免 Bank 冲突）​

小结​