CUDA程序中的种族状况

Question

我有两段代码。一个用C写，相应的操作用CUDA写。请帮助我理解__syncthreads()是如何在以下程序的上下文中工作的。根据我的理解，__syncthreads()确保了仅限于一个块的线程的同步。

C程序：

{
    for(i=1;i<10000;i++)
    {
        t=a[i]+b[i];
        a[i-1]=t;
    }
}

`

相应的数据自动化系统计划：

__global__ void kernel0(int *b, int *a, int *t, int N)
{
    int b0=blockIdx.x;
    int t0=threadIdx.x;
    int tid=b0*blockDim.x+t0;
    int private_t;
    if(tid<10000)
    {
        private_t=a[tid]+b[tid];
        if(tid>1)
            a[tid-1]=private_t;
        __syncthreads();
        if(tid==9999)
        *t=private_t;
    }
}

内核维数：

dim3 k0_dimBlock(32);
dim3 k0_dimGrid(313);
kernel0 <<<k0_dimGrid, k0_dimBlock>>>

令人惊讶的是，C和CUDA程序的输出是相同的。由于问题的本质是依赖于a[]，所以ai由thrad-ID i加载，并由同一个线程写入ai-1。现在，线程IDI-1也会发生同样的情况。如果问题的大小小于32，则输出是明显的。但是对于10000大小的313个块和块的问题，依赖是如何被尊重的？

Answer 1

根据我的理解，__syncthreads()确保了仅限于一个块的线程的同步。

你是对的。__syncthreads()是块上下文中的同步障碍。因此，例如，当您必须确保在开始算法的下一阶段之前更新所有数据时，这是非常有用的。

考虑到问题的本质，问题本身依赖于a[]，所以ai由线程ID加载，由同一个线程写入ai-1。

假设线程2到达if语句，因为它与它输入到语句的条件相匹配。现在线程执行以下操作：

private_t=a[2]+b[2];
a[1]=private_t;

巫婆相当于：

a[1]=a[2]+b[2];

正如您所指出的，它是数组a上的数据依赖。由于您无法在某个时候控制warps的执行顺序，所以您将使用a数组的更新版本。在我看来，您需要添加一个额外的__syncthreads()语句：

if( tid > 0 && tid<10000)
{
    private_t=a[tid]+b[tid];
    __syncthreads();
    a[tid-1]=private_t;
    __syncthreads();
    if(tid==9999)
    *t=private_t;
}

通过这种方式，每个线程都使用原始数组private_t获得自己版本的a变量，然后并行更新数组。

关于*t值的：

如果您只查看*t的值，您将不会注意到这种随机调度的影响，这取决于启动参数，这是因为带有tid==9999的线程可能与线程tid==9998一起处于最后一曲。由于创建private_t值所需的两个数组位置已经存在同步障碍，所以答案应该是正确的。