将数据从全局内存移动到共享内存是否会使线程停滞

Question

__shared__ float smem[2];
smem[0] = global_memory[0];
smem[1] = global_memory[1];
/*process smem[0]...*/
/*process smem[1]...*/

我的问题是，smem[1] = global_memory[1];在smem[0]上阻塞计算吗？在Cuda线程调度-延迟隐藏和Cuda全局内存加载和存储中，他们说内存读取不会使线程停滞，直到读取的数据被使用为止。将其存储到共享存储器是否计数为"0"；使用数据"；？我应该这样做吗：

__shared__ float smem[2];
float a = global_memory[0];
float b = global_memory[1];
smem[0] = a;
/* process smem[0]*/
smem[1] = b;
/* process smem[1]*/

或者可能是编译器帮我做的？但是它使用额外的寄存器吗？

Answer 1

是的，在一般情况下，这会阻塞CUDA线程：

smem[0] = global_memory[0];

原因是这种操作将分为两个步骤：

LDG  Rx, [Ry]
STS  [Rz], Rx

第一条SASS指令从全局内存加载。此操作不会阻塞CUDA线程。它可以被发送到LD/ST单元，并且线程可以继续。但是，会跟踪该操作(Rx(的寄存器目标，并且如果任何指令需要使用Rx中的值，则CUDA线程将在该点停止。

当然，下一条指令是STS(存储共享(指令，它将使用Rx中的值，因此CUDA线程将在这一点上停滞(直到满足全局负载(。

当然，编译器可能会对指令进行重新排序，以便稍后出现STS指令，但这并不能保证。不管怎样，每当编译器对STS指令进行排序时，CUDA线程都会在该点停止，直到全局加载完成。对于您给出的例子，我认为编译器很可能会创建如下代码：

LDG  Rx, [Ry]
LDG  Rw, [Ry+1]
STS  [Rz], Rx
STS  [Rz+1], Rw

换句话说，我认为编译器很可能会组织这些加载，以便在可能的暂停发生之前发出两个全局加载。然而，这并不能保证，您的代码的具体行为只能通过研究实际的SASS来推断，但在一般情况下，我们应该假设线程停滞的可能性。

是的，如果你可以像你在代码中显示的那样分解负载和存储，那么这个操作：

float b = global_memory[1];

不应阻止此操作：

smem[0] = a;
/* process smem[0]*/

话虽如此，CUDA在CUDA 11中引入了一种新的机制来解决这种情况，该机制由计算能力8.0及更高的设备(因此，目前所有的Ampere GPU(支持。这一新功能被称为从全局内存到共享内存的数据异步复制。它允许在不停止CUDA线程的情况下继续执行这些复制操作。然而，此功能需要正确使用屏障，以确保当您需要实际使用共享内存中的数据时，它是存在的。