同步线程和全局内存

Question

我写了一个小内核，使用并行归约对2^k个元素求和。这里没有什么新鲜事。。。。我的向量存储在全局内存中，我将向量的每个部分分配给不同的块，并将每个块减少到一个位置。剩下的我在CPU中完成。

__global__ void sum(real *v, long int s){
long int ix     =  threadIdx.x;
long int shift = blockIdx.x*blockDim.x;
long int h = blockDim.x/2;
while (h >= 1){
if (ix < h){
v[ix +  shift] = v[2*ix + shift] + v[2*ix + 1 + shift];
}
__syncthreads(); 
h = h / 2;
}
}

代码有效。然而，经过仔细检查，我意识到也许它不应该起作用。所以我很困惑。。。。在thread_id＝0能够读取元素0和1之前，对元素2和3求和的线程_id＝1可以将其和写入位置1。从而使结果无效。

我本以为，为了安全起见，代码必须是

__global__ void sumsafe(real *v, long int s){
long int ix     =  threadIdx.x;
long int shift = blockIdx.x*blockDim.x;
real x = 0;
long int h = blockDim.x/2;
while (h >= 1){
if (ix < h){
x = v[2*ix + shift] + v[2*ix + 1 + shift];
}
__syncthreads(); 
if (ix < h){
v[ix +  shift] = x;
}
__syncthreads();
h = h / 2;
}
}

因此，我保证所有线程在开始更改它们之前都读取了它们的值。但正如我所说。。。两个代码都有效。。。他们的时间实际上也差不多。

为什么会这样？

我知道GPU不能保证一个线程写入全局内存的内容对其他线程不可见。但它也不能保证这种情况永远不会发生。

有什么想法吗！？我正在开发GTX1080。

Answer 1

您确实很"幸运"，因为CUDA无法保证扭曲的执行顺序。以下描述(这是推测)不应被解释为你所展示的是一个好主意的陈述。任何人都不应该这样做。

但对于一个小的测试用例(没有其他代码，只对一个数据块进行操作)，我希望它能起作用。

从全局内存读取通常是高延迟的。当执行遇到这行代码时：

v[ix +  shift] = v[2*ix + shift] + v[2*ix + 1 + shift];

它将转化为SASS指令，类似于以下内容：

LD  R0, v[2*ix + shift]        (let's call this LD0)
LD  R1, v[2*ix + 1 + shift];   (let's call this LD1)
ADD R3, R0, R1
ST  v[ix + shift], R3

现在，前两个LD操作不会导致失速。但是，如果R1和R0还无效，ADD操作将导致暂停(无法发出)。

失速的结果将是SM中的曲速调度引擎将寻找其他可用的工作。其他可用的工作可能会构成其他翘曲的上述代码。

由于ADD指令在读取完成之前不能发出，并且由于扭曲调度程序对暂停的响应，读取(跨扭曲)都有效地背靠背发出，因此在ADD指令完成发出时，读取操作将趋向于所有完成，这意味着在发出所有ADD操作时所有读取都已完成(并且在其对应的ADD完成之前不能发出ST)。ADD也有管道延迟，因此ADD操作可能也会按顺序发布(但这里的短管道延迟可能会增加危险的概率)，并且在相应的ADD操作完成之前，无法发布给定的ST操作。净效应可能是：

LD0 W0
LD1 W0
LD0 W1
LD1 W1
...   (all LD0 and LD1 get issued across all warps W0..WN)
<read latency stall --  eventually the first 2 LD0 and LD1 complete>
ADD W0
<read pipeline latency - 1 cycle>
ADD W1
<read pipeline latency - 1 cycle>
ADD W2
...
<add pipeline latency>
ST W0
<add pipeline latency>
ST W1
...

延迟的结果是，在任何ADD操作开始之前，所有读取都很可能被发送到全局内存。由于管道效应，在任何ST操作开始之前，所有读取操作也有可能完成，因此在这种有限的测试情况下，可能不会发生实际的危险错误。

我预计，即使数据在二级缓存中，从二级缓存读取的延迟可能仍然足以使上述功能发挥作用。我怀疑，如果数据在一级缓存中，从一级缓存读取的延迟(并假设扭曲的最大补码)可能不足以使上述描述成立，但我没有仔细研究算术。由于ADD管道延迟是固定的，但与ADD管道延迟相比，LD到ST操作的危险是由ADD操作的数量决定的，因此实际危险概率会随着在线程块中加载更多扭曲而增加。

请注意，上面的所有描述都试图解开while循环的单个迭代的行为。__syncthreads()的存储器屏障效应应保证迭代i+1的读取不会被迭代i的写入(未能见证)破坏。