Volta架构中规约计算的同步问题

  • 5 replies
  • 192 views
Volta架构中规约计算的同步问题
« 于: 八月 22, 2019, 07:31:32 am »
CUDA C 编程指南说到从Volta起,线程束内将不再有默认的同步。这篇博文(https://devblogs.nvidia.com/using-cuda-warp-level-primitives/)建议用如下代码进行线程束内部的规约:

int tid = threadIdx.x;
int v = 0;
v += shmem[tid+16]; __syncwarp();
shmem[tid] = v;     __syncwarp();
v += shmem[tid+8];  __syncwarp();
shmem[tid] = v;     __syncwarp();
v += shmem[tid+4];  __syncwarp();
shmem[tid] = v;     __syncwarp();
v += shmem[tid+2];  __syncwarp();
shmem[tid] = v;     __syncwarp();
v += shmem[tid+1];  __syncwarp();
shmem[tid] = v;

也就是说,认为这样是不对的:
int tid = threadIdx.x;
shmem[tid] += shmem[tid+16]; __syncwarp();
shmem[tid] += shmem[tid+8];  __syncwarp();
shmem[tid] += shmem[tid+4];  __syncwarp();
shmem[tid] += shmem[tid+2];  __syncwarp();
shmem[tid] += shmem[tid+1];  __syncwarp();

我不太懂是什么意思。在第二段代码中的第二行,展开的话就是
shmem[tid] = shmem[tid] + shmem[tid+16]; __syncwarp();
等号右边要读取shmem[tid],左边要写入shmem[tid]。意思是这个读取和写入的次序不能保证?

另外,线程束之前的规约会不会也有类似的要求呢?也就是说,写成如下形式是否有错?
int tid = threadIdx.x;
__syncthreads();
for (int offset = blockDim.x / 2; offset > 0; offset /= 2)
{
    if (tid < offset)
    {
        shmem[tid] += shmem[tid + offset]; // 有问题吗?为什么?
    }
    __syncthreads();
}


Re: Volta架构中规约计算的同步问题
« 回复 #1 于: 八月 22, 2019, 07:36:56 am »
我的最后一段代码的循环指标有误,应该改为:

for (int offset = blockDim.x / 2; offset >= 32; offset /= 2)

Re: Volta架构中规约计算的同步问题
« 回复 #2 于: 八月 23, 2019, 03:08:04 pm »
brucefan:

关于是否一个多个线程参与的,在某区域内的有读有写的过程是安全的,则要看:
(1)编译器能否正确的产生代码(软件读写barrier),和,
(2)硬件在执行的时候,是否能保证一定的次序。
这两点缺一不可。

我们都知道从计算能力7.0+开始,warp内不在严格的locked-step的执行了,因此需要额外引入的__syncwarp()。那么在插入了__syncwarp()的代码中,例如你给出的例子:

引用
int tid = threadIdx.x;
shmem[tid] += shmem[tid+16]; __syncwarp();
shmem[tid] += shmem[tid+8];  __syncwarp();
shmem[tid] += shmem[tid+4];  __syncwarp();
shmem[tid] += shmem[tid+2];  __syncwarp();
shmem[tid] += shmem[tid+1];  __syncwarp();

也就是, 这种常见处理1个warp最后的32个数据的规约求和的代码中,我们从warp中具体的几个线程来看来看,会出现:
线程8:[8] = [8] + [8 + 4];
线程4:[4] = [4] + [4 + 4];
然后一起__syncwarp();
则,在sync warp之前,线程8和线程4的硬件上的执行次序是无法确定的(上面说的第2点),很可能线程8已经写入改变[8]元素的值,而线程4才读取[8];(或者线程8没有改变[8]元素的值,线程[4]已经读取了[8]),

这样同一个线程4,读取到的[8]的值,存在多种可能,因此程序无法得到一个确定的结果的。(所谓的竞态,race condition, 和具体的硬件执行的时序)有关,因此是不安全的。

所以这也是为何你的下一个例子,将写入读取隔离开,额外的插入__syncwarp(), 却可以保证次序,从而结果正确的原因。

Re: Volta架构中规约计算的同步问题
« 回复 #3 于: 八月 23, 2019, 03:11:13 pm »
接上:

你的原始理解,同一个线程内部的:
引用
我不太懂是什么意思。在第二段代码中的第二行,展开的话就是
shmem[tid] = shmem[tid] + shmem[tid+16]; __syncwarp();
等号右边要读取shmem[tid],左边要写入shmem[tid]。意思是这个读取和写入的次序不能保证?

问题并不在于同一个线程内部的读取,然后写入(这个是安全的)。而在于多个线程间的读写次序。因此你之前的解释看似正确,实际上不对。这点额外强调一下。

Re: Volta架构中规约计算的同步问题
« 回复 #4 于: 八月 23, 2019, 03:18:33 pm »
引用
for (int offset = blockDim.x / 2; offset > 0; offset /= 2)
{
    if (tid < offset)
    {
        shmem[tid] += shmem[tid + offset]; // 有问题吗?为什么?
    }
    __syncthreads();
}

同时让我们来看看你的第三个代码,这代码其实和之前的并不同,这里多了这么一行:
if (tid < offset)   <---- 注意这里的约束条件
{[tid] = [tid] + [tid + offset];}
__syncthreads();

此代码在多了这个约束条件, 在此条件下,对于同1个block内部的线程来说(假设block是1D的),
不可能在该约束条件下,存在任意2个线程,使得读取的[tid]或者[tid + offset]和另外一个线程的写入的[tid']元素相同。

这个容易证明:
如果相同,
要么(1)下标tid = tid' (tid来自于threadIdx.x, 刚才强调了是1D的,因此不可能),这是不可能的。
要么(2)下标tid + offset = tid', 即存在2个线程,他们的下标差为offset。但是刚才约束了所有的下表范围均为[0, offset), 在此范围内是找不到任意两个坐标,能坐标差为offset的(显然!)

结合(1)(2),这种是安全的。

所以你看,我们分析的时候只需要线程间的读写。而无视了同一个线程对同一个下标位置的读取,然后写入(上面说了,同一个线程这样做是安全的)。

你以后可以用这种方式分析任意类似代码。

Re: Volta架构中规约计算的同步问题
« 回复 #5 于: 八月 23, 2019, 03:37:16 pm »
brucefan:

关于是否一个多个线程参与的,在某区域内的有读有写的过程是安全的,则要看:
(1)编译器能否正确的产生代码(软件读写barrier),和,
(2)硬件在执行的时候,是否能保证一定的次序。
这两点缺一不可。

我们都知道从计算能力7.0+开始,warp内不在严格的locked-step的执行了,因此需要额外引入的__syncwarp()。那么在插入了__syncwarp()的代码中,例如你给出的例子:

也就是, 这种常见处理1个warp最后的32个数据的规约求和的代码中,我们从warp中具体的几个线程来看来看,会出现:
线程8:[8] = [8] + [8 + 4];
线程4:[4] = [4] + [4 + 4];
然后一起__syncwarp();
则,在sync warp之前,线程8和线程4的硬件上的执行次序是无法确定的(上面说的第2点),很可能线程8已经写入改变[8]元素的值,而线程4才读取[8];(或者线程8没有改变[8]元素的值,线程[4]已经读取了[8]),

这样同一个线程4,读取到的[8]的值,存在多种可能,因此程序无法得到一个确定的结果的。(所谓的竞态,race condition, 和具体的硬件执行的时序)有关,因此是不安全的。

所以这也是为何你的下一个例子,将写入读取隔离开,额外的插入__syncwarp(), 却可以保证次序,从而结果正确的原因。

我之前的理解确实不对。你分析得非常到位。感谢解惑!