将数据从全局内存移动到共享内存会使线程停止吗？
Does moving data from global memory to shared memory stall the thread?
__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] 的计算吗？
在 and Cuda global memory load and store 他们说内存读取不会停止线程，直到使用读取的数据。将其存储到共享内存算作“使用数据”吗？我应该这样做吗：
__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]*/
或者也许是编译器帮我做了？但是它会使用额外的寄存器吗？
是的，在一般情况下这会阻塞 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）都支持该机制。此新功能称为 asynchronous copy of data from global to shared memory。它允许这些复制操作在不停止 CUDA 线程的情况下继续进行。然而，此功能需要正确使用屏障，以确保当您需要实际使用共享内存中的数据时，它就存在。
__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] 的计算吗？
在
__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]*/
或者也许是编译器帮我做了？但是它会使用额外的寄存器吗？
是的，在一般情况下这会阻塞 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）都支持该机制。此新功能称为 asynchronous copy of data from global to shared memory。它允许这些复制操作在不停止 CUDA 线程的情况下继续进行。然而，此功能需要正确使用屏障，以确保当您需要实际使用共享内存中的数据时，它就存在。