在代码中过度使用__syncthread

Question

我理解的目的__syncthreads(),但我有时会发现它在某些代码中被过度使用.

例如,在下面的代码中,从NVIDIA注释中,每个线程主要计算s_data[tx]-s_data[tx-1].每个线程需要从全局内存中读取的数据以及由其相邻线程读取的数据.两个线程将处于相同的warp中,因此应完成从全局内存中检索其数据并安排同时执行.

我相信代码仍然无法使用__syncthread(),但显然NVIDIA的说明不是这样.有什么评论吗？

// Example – shared variables
// optimized version of adjacent difference
__global__ void adj_diff(int *result, int *input)
{
    // shorthand for threadIdx.x
    int tx = threadIdx.x;
    // allocate a __shared__ array, one element per thread
    __shared__ int s_data[BLOCK_SIZE];
    // each thread reads one element to s_data
    unsigned int i = blockDim.x * blockIdx.x + tx;
    s_data[tx] = input[i];
    // avoid race condition: ensure all loads
    // complete before continuing
    __syncthreads();

    if(tx > 0)
        result[i] = s_data[tx] – s_data[tx–1];
    else if(i > 0)
    {
        // handle thread block boundary
        result[i] = s_data[tx] – input[i-1];
    }
}

Answer 1

如果你在"Nvidia笔记"中出现了一个链接,那就好了.

两个线程将处于相同的warp中

不,他们不会,至少在所有情况下都不会.当tx= 32 时会发生什么？然后对应的线程tx属于块中的warp 1,并且对应的线程tx-1属于块中的warp 0.

不能保证warp 0在warp 1之前执行,因此代码可能会在没有调用的情况下失败__synchtreads()(因为没有它,值s_data[tx-1]可能无效,因为warp 0没有运行,因此尚未加载它. )