CUDA 优化

Mini-ImagePipe 中使用的 GPU 优化技术。

内存层次结构

内存类型

内存	位置	速度	作用域
全局内存	DRAM	~800 GB/s	所有线程
共享内存	片上	~2 TB/s	线程块
常量内存	缓存	~1 TB/s	所有线程
寄存器	片上	~10 TB/s	线程

共享内存分块

对于卷积操作，我们将分块和边缘区域加载到共享内存：

分块配置

// For 5×5 kernel with 16×16 tiles
constexpr int TILE = 16;
constexpr int HALO = 2;  // (kernel_size - 1) / 2

__shared__ float sharedMem[TILE + 2*HALO][TILE + 2*HALO];

边缘区域加载

+------------------+
|  Halo (top)      |
+--+------------+--+
|H |   Tile     | H|
|a |   Data     | a|
|l |            | l|
|o |            | o|
+--+------------+--+
|  Halo (bottom)   |
+------------------+

常量内存

Sobel 内核存储在常量内存中用于广播：

__constant__ float c_sobelX[9] = {-1, 0, 1, -2, 0, 2, -1, 0, 1};
__constant__ float c_sobelY[9] = {-1, -2, -1, 0, 0, 0, 1, 2, 1};

// All threads in a warp access the same address → single memory fetch

流并发

独立执行

// Stream 0: Operator A
kernelA<<<grid, block, 0, stream0>>>(...);
cudaEventRecord(eventA, stream0);

// Stream 1: Operator B (independent)
kernelB<<<grid, block, 0, stream1>>>(...);

// Stream 0: Operator C (depends on A)
cudaStreamWaitEvent(stream0, eventA);
kernelC<<<grid, block, 0, stream0>>>(...);

占用率优化

线程块大小

// Target: maximize occupancy
// Consider: registers per thread, shared memory per block

dim3 blockSize(16, 16);  // 256 threads per block
dim3 gridSize((width + 15) / 16, (height + 15) / 16);

寄存器压力

// Reduce register usage for higher occupancy
__launch_bounds__(256, 4)  // 256 threads, min 4 blocks per SM
__global__ void myKernel(...) {
    // ...
}

性能指南

1. 合并访问：确保线程访问连续的内存地址
2. 共享内存 bank 冲突：填充数组以避免 bank 冲突
3. 分支发散：最小化 warp 内的分支
4. 内存传输：使用固定内存进行异步传输

参考资料

• CUDA Best Practices Guide
• Harris, M. (2007). "Optimizing Parallel Reduction in CUDA."