CUDA替代__syncthreads而不是__threadfence()差异

Question

我从NVIDIA手册Eg中复制了以下代码:__threadfence()。他们为什么有
在以下代码中使用了__threadfence()。我认为使用__syncthreads()而不是__threadfence()将给您相同的结果。

有人可以解释__syncthreads()和__threadfence()调用之间的区别吗？

__device__ unsigned int count = 0;
__shared__ bool isLastBlockDone;

__global__ void sum(const float* array, unsigned int N,float* result)
{
    // Each block sums a subset of the input array
    float partialSum = calculatePartialSum(array, N);

    if (threadIdx.x == 0) {
        // Thread 0 of each block stores the partial sum
        // to global memory
        result[blockIdx.x] = partialSum;

        // Thread 0 makes sure its result is visible to
        // all other threads
        __threadfence();

        // Thread 0 of each block signals that it is done
        unsigned int value = atomicInc(&count, gridDim.x);

        // Thread 0 of each block determines if its block is
        // the last block to be done
        isLastBlockDone = (value == (gridDim.x - 1));
    }

    // Synchronize to make sure that each thread reads
    // the correct value of isLastBlockDone
    __syncthreads();

    if (isLastBlockDone) 
    {
        // The last block sums the partial sums
        // stored in result[0 .. gridDim.x-1]
        float totalSum = calculateTotalSum(result);

        if (threadIdx.x == 0)
        {
            // Thread 0 of last block stores total sum
            // to global memory and resets count so that
            // next kernel call works properly
            result[0] = totalSum;
            count = 0;
        }
    }
}

Answer 1

就共享内存而言，__syncthreads()比__threadfence()更加强大。关于全局内存-这是两件不同的事情。

__threadfence_block()暂停当前线程，直到同一块中的其他线程可以看到对共享内存的所有写入。它通过在寄存器中缓存共享内存写来防止编译器优化。它不同步线程，也不需要所有线程实际到达此指令。

__threadfence()暂停当前线程，直到对所有其他线程可见的对共享和全局内存的所有写操作为止。

块中的所有线程都必须达到

__syncthreads()(例如，没有分歧的if语句)，并确保针对该块中的所有线程，在该指令之前的代码之前执行该指令之前的代码。

在您的特定情况下， __threadfence()指令用于确保每个人都可以看到对全局数组 result的写入。 __syncthreads()仅会仅同步当前块中的线程，而不会为其他块强制执行全局内存写入。而且，到那时，您在 if分支内的代码中，只有一个线程在执行该代码。使用 __syncthreads()将导致GPU出现不确定的行为，最有可能导致内核完全不同步。

查阅《 CUDA C编程指南》中的以下章节:

3.2.2“共享内存”-矩阵乘法

的示例

5.4.3“同步指令”

B.2.5“易失”

B.5“存储栅栏功能”