CUDA共享内存Bank Conflict问题：原因分析与优化方案求助

问题背景 #

在完成卡内基梅隆大学15-418课程作业2的渲染模块时，基于CUDA 12.2开发高性能渲染器，采用单线程计算单个像素的数据并行模式。为加速计算使用共享内存存储中间结果，但通过Nsight Compute分析发现存在Bank冲突问题。使用GPU型号为NVIDIA RTX 3090，曾尝试修改共享内存数组大小（将32改为其他值）但未解决冲突，后续通过__ldg指令加载常量全局内存消除了冲突，但发现该问题并非性能瓶颈。

核心代码 #

__shared__ float smem_rads[32];
__shared__ float smem_colors[32 * 3];
__shared__ float smem_ps[32 * 3];

// Load values into the shared memory arrays every 32 indices
for (int index = 0; index < cuConstRendererParams.numCircles; ++index) {
    int index3 = 3 * index;
    if (index % 32 == 0) {
        __syncthreads();
        if (threadId < 32) {
            smem_ps[threadId] = cuConstRendererParams.position[index3 + threadId];
            smem_ps[threadId + 32] = cuConstRendererParams.position[index3 + threadId + 32];
            smem_ps[threadId + 64] = cuConstRendererParams.position[index3 + threadId + 64];
            smem_rads[threadId] = cuConstRendererParams.radius[index + threadId];
            smem_colors[threadId] = cuConstRendererParams.color[index3 + threadId];
            smem_colors[threadId + 32] = cuConstRendererParams.color[index3 + threadId + 32];
            smem_colors[threadId + 64] = cuConstRendererParams.color[index3 + threadId + 64];
        }
        __syncthreads();
    }

    // Access the shared memory values later
    float3 p = ((float3*)smem_ps)[index % 32];
    float rad = smem_rads[index % 32];
    float3 rgb = ((float3*)smem_colors)[index % 32];
}

Bank冲突产生原因 #

NVIDIA RTX 3090（Ampere架构）的共享内存默认划分为32个Bank，每个Bank一次只能响应一个线程的访问请求。当warp内多个线程同时访问同一Bank的不同地址时，访问会被序列化，引发Bank冲突，降低内存访问效率。

你的代码中，冲突来自smem_ps和smem_colors的float3类型访问：

• 共享内存数组smem_ps和smem_colors均为float类型（4字节/元素），强制转换为float3指针后，每个float3对应连续3个float元素（共12字节）。
• 当warp内32个线程同时访问((float3*)smem_ps)[index%32]时，线程t会访问smem_ps[3*t]、smem_ps[3*t+1]、smem_ps[3*t+2]三个元素。
• 计算Bank编号：Bank = (元素字节偏移 / 4) % 32。对于smem_ps[3*t]，字节偏移为3*t*4=12t，对应Bank编号为(12t/4)%32=3t%32。
• 随着t从0到31，3t%32会出现重复值（例如t=0和t=11的Bank编号分别为0和33%32=1，t=1和t=12的Bank编号为3和36%32=4，以此类推），导致多个线程同时访问同一Bank，触发冲突。

而smem_rads为单float数组，线程t访问smem_rads[t]，对应Bank编号为t%32，每个线程访问不同Bank，因此无冲突。

解决建议 #

1. 调整共享内存布局（添加填充）

在float3元素之间添加填充，确保每个float3的起始地址对应唯一Bank。例如将共享内存数组定义为：

__shared__ float smem_ps[32 * 4];  // 每个float3占4个float空间，1个float作为填充
__shared__ float smem_colors[32 * 4];

此时每个float3对应smem_ps[4*t]、smem_ps[4*t+1]、smem_ps[4*t+2]，第四个元素为填充。字节偏移为4*t*4=16t，Bank编号为(16t/4)%32=4t%32，每个线程的Bank编号唯一，避免冲突。

2. 拆分分量存储

将float3的三个分量分别存储到独立的共享内存数组中，避免连续访问导致的Bank冲突：

__shared__ float smem_ps_x[32];
__shared__ float smem_ps_y[32];
__shared__ float smem_ps_z[32];
__shared__ float smem_colors_r[32];
__shared__ float smem_colors_g[32];
__shared__ float smem_colors_b[32];

加载和访问时直接操作单个分量：

// 加载时
smem_ps_x[threadId] = cuConstRendererParams.position[index3 + threadId];
smem_ps_y[threadId] = cuConstRendererParams.position[index3 + threadId + 32];
smem_ps_z[threadId] = cuConstRendererParams.position[index3 + threadId + 64];

// 访问时
float3 p = make_float3(smem_ps_x[index%32], smem_ps_y[index%32], smem_ps_z[index%32]);

这种方式下每个数组的访问都是线程t对应索引t，Bank编号t%32唯一，完全消除冲突。

3. 调整共享内存Bank宽度（Ampere架构专属）

Ampere架构支持将共享内存Bank宽度配置为8字节（默认4字节），可通过代码或运行时设置：

• 代码中使用对齐属性：__shared__ __align__(8) float smem_ps[32*3];
• 运行时通过cudaFuncSetAttribute设置cudaFuncAttributeSharedMemoryBankSize为cudaSharedMemBankSizeEightByte

该方式会改变Bank的划分规则，缓解部分冲突，但效果不如前两种方式直接可靠。

补充说明 #

通过__ldg指令加载常量全局内存后，共享内存的Bank冲突消失，这是因为__ldg会将数据缓存到只读缓存（L1/L2），绕过了共享内存的访问路径。后续发现该冲突并非性能主要瓶颈，说明渲染模块的核心性能开销可能在光线追踪计算、全局内存访问延迟或其他计算密集型逻辑上。

内容的提问来源于stack exchange，提问作者Sunjnn