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

Question

在 CUDA 编程指南的共享内存部分中，它指出 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

共享内存组冲突仅与特定指令/周期上的线程束内的线程相关。GPU 中的所有指令均在扭曲范围内发出。它们不会在同一周期中从单个 warp 调度程序发布到线程块中的所有 warp。

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

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

同样，不存在超越扭曲的广播机制。

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

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

您的代码不包含原子，也不包含共享内存写入同一位置，并且我在这里所说的几乎所有内容都与原子有关。原子通过原子处理机制进行扭曲聚合或序列化，并且对同一位置的多个（非原子）写入会导致未定义的行为。您可以预期其中一项写入将显示在该位置，但具体是哪一项尚未定义。从性能角度来看，我不知道有任何关于同一位置共享写入性能的说法。从性能角度来看，原子是一种完全不同的动物。