__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 thread scheduling - latency hiding和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]*/或者是编译器为我做的？但它会使用额外的寄存器吗？
1条答案
按热度按时间ugmeyewa1#
是的，在一般情况下，这会阻塞CUDA线程：
原因是该操作将被分成两个步骤：
第一条SASS指令从全局存储器加载。此操作不会阻止CUDA线程。它可以被发送到LD/ST单元，并且线程可以继续。然而，该操作的寄存器目标（Rx）被跟踪，并且如果任何指令需要使用来自
Rx的值，则CUDA线程将在该点处停止。当然，下一条指令是STS（存储共享）指令，它将使用来自
Rx的值，因此CUDA线程将在该点停止（直到满足全局加载）。当然，编译器可能会对指令进行重新排序，以便稍后出现
STS指令，但不能保证这一点。无论如何，每当编译器命令STS指令时，CUDA线程将在该点停止，直到全局加载完成。对于你给出的例子，我认为编译器很可能会创建这样的代码：换句话说，我认为编译器可能会组织这些加载，以便在可能的停顿发生之前，可以发出两个全局加载。然而，这并不能保证，代码的特定行为只能通过研究实际的SASS来推断，但在一般情况下，我们应该假设线程停顿的可能性。
是的，如果你可以像你的代码中所示的那样分解加载和存储，那么这个操作：
不应阻止此操作：
话虽如此，CUDA在CUDA 11中引入了一种新的机制来解决这种情况，由计算能力8.0及更高版本的设备（因此，此时所有Ampere GPU）支持。这个新特性被称为从全局内存到共享内存的异步数据复制。它允许这些复制操作继续进行，而不会停止CUDA线程。但是，这个特性需要正确使用屏障，以确保当您需要实际使用共享内存中的数据时，它存在。