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
}

虽然您没有问过这个问题，但在共享内存上的线程块中使用原子可能可以通过共享内存减少方法来替换（以获得更好的性能）。