CUDA 中多线程内核中不必要的写操作是否效率低下？

Question

我的 CUDA 代码中有一个内核，我希望一堆线程在一些共享内存上做一堆计算（因为它比在全局内存上这样做快得多），然后写入结果到全局内存（这样我就可以在以后的内核中使用它）。内核看起来像这样：

__global__ void calc(float * globalmem)
{
    __shared__ float sharemem; //initialize shared memory
    sharemem = 0; //set it to initial value
    __syncthreads();

   //do various calculations on the shared memory
   //for example I use atomicAdd() to add each thread's
   //result to sharedmem...

   __syncthreads();
   *globalmem = sharedmem;//write shared memory to global memory
}

事实上，每个线程都将数据从共享内存写出到全局内存，而我真的只需要将其写出一次，这让我觉得很可疑。我也从每个线程在代码开始时将共享内存初始化为零的事实中得到同样的感觉。有没有比我当前的实现更快的方法？

Answer 1

在 warp 级别，进行冗余读取或写入与使用单个线程进行读取或写入之间可能没有太大的性能差异。

但是，我希望通过在线程块中使用多个 warp 进行冗余读取或写入（与单个线程相比），可能会出现可测量的性能差异。

通过让单个线程执行读取或写入操作而不是冗余操作应该足以解决这些问题：

__global__ void calc(float * globalmem)
{
    __shared__ float sharemem; //initialize shared memory
    if (!threadIdx.x) sharemem = 0; //set it to initial value
    __syncthreads();

   //do various calculations on the shared memory
   //for example I use atomicAdd() to add each thread's
   //result to sharedmem...

   __syncthreads();
   if (!threadIdx.x) *globalmem = sharemem;//write shared memory to global memory
}

尽管您没有询问，但在共享内存上的线程块中使用原子可能可以通过共享内存缩减方法替换（以获得更好的性能）。

CUDA 中多线程内核中不必要的写操作是否效率低下？

Are needless write operations in multi-thread kernels in CUDA inefficient?

cuda

gpu

gpgpu

shared-memory