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

问题描述

__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 hidingCuda 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线程。但是,此功能需要正确使用屏障,以确保当您需要实际使用共享内存中的数据时,该数据就会存在。