CUDA代替__syncthreads而不是__threadfence()的区别

Question

我从NVIDIA手册中复制了以下代码,例如:for __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"记忆围栏功能"