CUDA 是否将共享内存广播到块中的所有线程而不会发生内存冲突？

Question

在 CUDA 编程指南中，在共享内存中 section，它声明 warp 的共享内存访问不是序列化的，而是广播读取的。

但是它没有说明如果整个块请求相同的内存地址会发生什么。 warp 之间的访问是序列化的还是 CUDA 可以广播到整个块。

我的案例的演示代码

// Assume 1024 sized int array
__global__ add_from_shared(int* i, int* j, int* out)
{
    __shared__ int shmem[1024];
    shmem[threadIdx.x] = i[threadIdx.x];
    ...
    Do some stuff
    ...
    // Is the shared memory call here serilized between warps or is it a broadcast over the entire block?
    j[threadIdx.x] += shmem[0];
}

谢谢

Answer 1

共享内存库冲突仅与特定 instruction/cycle 上 warp 内的线程相关。 GPU中的所有指令都发出warp-wide。它们不会在同一周期内从单个 warp 调度程序发布到 threadblock 中的所有 warp。

不存在不同线程间共享内存库冲突的概念，也不存在执行不同指令的线程间共享内存库冲突的概念。

warp 调度程序将单独向每个 warp 发出共享读取指令 (LDS)。根据该 warp 中线程之间明显的访问模式，对于该发出的指令，可能会或可能不会发生库冲突。一个 warp 的线程与另一个 warp 的线程之间不可能存在库冲突。

同样没有超出 warp 的广播机制。

GPU 中的所有指令都是按 warp 发出的。

如果一个块中的所有线程都读取相同的地址，则 warp 调度程序将向一个 warp 发出该指令，并且对于该 warp 中的线程，将应用广播。在同一时间或不同时间，从同一个 warp 调度程序或不同的 warp 调度程序，相同的指令（即来自指令流中的同一点）将被发布到另一个 warp。广播将在该请求内发生。在线程块中重复尽可能多的扭曲。

您的代码不包含原子，或共享内存写入同一位置，而且我在这里所说的几乎所有内容都与原子无关。原子是 warp-aggregated 或由原子处理机制序列化，并且多次 (non-atomic) 写入同一位置会导致未定义的行为。您可以预期其中一个写入将出现在该位置，但哪个是未定义的。从性能的角度来看，我不知道有任何关于 same-location-shared-write 性能的陈述。从性能的角度来看，原子是完全不同的动物。

CUDA 是否将共享内存广播到块中的所有线程而不会发生内存冲突？

Does CUDA broadcast shared memory to all threads in a block without a bank conflict?

cuda

gpu

nvidia

shared-memory

gpgpu