CUDA能否无bank冲突将共享内存广播至块内所有线程？

不同线程束（warp）访问同一共享内存地址时，不存在覆盖整个线程块（block）的全局广播机制，跨warp的访问不会被强制串行化，每个warp会独立发起读请求，各自通过warp内部的广播通路完成数据读取，整体性能开销极低。

• 共享内存的读广播是warp级别的硬件原生特性，作用范围严格限定在单个warp内部。同一个warp内无论多少个线程访问同一共享内存地址，硬件仅需发起1次读请求，拿到数据后直接广播给warp内所有发起请求的线程，既不会出现warp内访问串行化，也不会触发bank冲突。
• warp是流多处理器（SM）上的最小独立调度单元，不同warp的指令发射、内存访问流程完全独立调度，硬件没有实现跨整个block的广播通路。当多个warp访问同一个共享内存地址时，每个warp都会单独向共享内存子系统发送读请求，这些请求进入访问队列后由SM内存调度器按既定策略处理，不会被强制串行排队。
• 共享内存bank冲突的判定规则仅对单个warp的访问请求生效，跨warp的访问互不参与bank冲突计算，即使多个warp同时读取同一地址，也不会产生额外的性能惩罚。

对应场景的示例代码如下：

// 假设传入的int数组长度为1024
__global__ add_from_shared(int* i, int* j, int* out)
{
    __shared__ int shmem[1024];
    shmem[threadIdx.x] = i[threadIdx.x];
    ...
    // 其他业务逻辑
    ...
    // 这里的共享内存读是跨warp串行化，还是会在整个block内广播？
    j[threadIdx.x] += shmem[0];
}  

以常见的1024线程block配置为例，整个块会被拆分为32个warp（每个warp固定包含32个线程），执行到读取shmem[0]的逻辑时：

1. 每个warp仅需发起1次共享内存读请求，通过warp内广播即可满足该warp全部32个线程的读取需求
2. 32个warp总共会产生32次独立的共享内存读请求，这些请求不会被强制阻塞，共享内存子系统可以在极短周期内完成所有响应，整体开销几乎可以忽略。
3. 如果追求极致性能，你也可以仅让单个线程将shmem[0]读取到寄存器，通过块同步+warp通信原语把值同步给所有线程，但对绝大多数业务场景来说这类优化没有实际收益，原生访问逻辑的效率已经足够高。

注意：以上结论仅针对共享内存读场景，如果涉及多个warp对同一共享内存地址的写操作，必须通过原子操作或者同步机制保证访问正确性，否则会出现未定义行为。

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