使用易失性共享内存的 CUDA 原子函数

Question

我有一个 CUDA 内核，它需要在 volatile 共享整数内存上使用原子函数。但是，当我尝试将共享内存声明为 volatile 并在原子函数中使用它时，我收到一条错误消息。

下面是一些重现错误的极简代码。请注意，下面的内核什么都不做，并且严重滥用了为什么您想将共享内存声明为易失性（甚至根本不使用共享内存）。但它确实重现了错误。

代码在共享内存上使用原子函数，因此，要运行它，您可能需要使用“arch12”或更高版本进行编译（在 Visual Studio 2010 中，右键单击您的项目并转到“属性 -> 配置属性 - > CUDA C/C++ -> Device”并在“Code Generation”行输入“compute_12,sm_12”）。否则代码应按原样编译。

#include <cstdlib>
#include <cuda_runtime.h>

static int const X_THRDS_PER_BLK = 32;
static int const Y_THRDS_PER_BLK = 8;

__global__ void KernelWithSharedMemoryAndAtomicFunction(int * d_array, int numTotX, int numTotY)
{
              __shared__ int s_blk[Y_THRDS_PER_BLK][X_THRDS_PER_BLK]; // compiles
   //volatile __shared__ int s_blk[Y_THRDS_PER_BLK][X_THRDS_PER_BLK]; // will not compile

   int tx = threadIdx.x;
   int ty = threadIdx.y;

   int mx = blockIdx.x*blockDim.x + threadIdx.x;
   int my = blockIdx.y*blockDim.y + threadIdx.y;
   int mi = my*numTotX + mx;

   if (mx < numTotX && my < numTotY)
   {
      s_blk[ty][tx] = d_array[mi];

      __syncthreads();

      atomicMin(&s_blk[ty][tx], 4); // will compile with volatile shared memory only if this line is commented out

      __syncthreads();

      d_array[mi] = s_blk[ty][tx];
   }
}

int main(void)
{
   // Declare and initialize some array on host
   int const NUM_TOT_X = 4*X_THRDS_PER_BLK;
   int const NUM_TOT_Y = 6*Y_THRDS_PER_BLK;

   int * h_array = (int *)malloc(NUM_TOT_X*NUM_TOT_Y*sizeof(int));

   for (int i = 0; i < NUM_TOT_X*NUM_TOT_Y; ++i) h_array[i] = i;

   // Copy array to device
   int * d_array;
   cudaMalloc((void **)&d_array, NUM_TOT_X*NUM_TOT_Y*sizeof(int));

   cudaMemcpy(d_array, h_array, NUM_TOT_X*NUM_TOT_Y*sizeof(int), cudaMemcpyHostToDevice);

   // Declare block and thread variables
   dim3 thdsPerBlk;
   dim3 blks;

   thdsPerBlk.x = X_THRDS_PER_BLK;
   thdsPerBlk.y = Y_THRDS_PER_BLK;
   thdsPerBlk.z = 1;

   blks.x = (NUM_TOT_X + X_THRDS_PER_BLK - 1)/X_THRDS_PER_BLK;
   blks.y = (NUM_TOT_Y + Y_THRDS_PER_BLK - 1)/Y_THRDS_PER_BLK;
   blks.z = 1;

   // Run kernel
   KernelWithSharedMemoryAndAtomicFunction<<<blks, thdsPerBlk>>>(d_array, NUM_TOT_X, NUM_TOT_Y);

   // Cleanup
   free    (h_array);
   cudaFree(d_array);

   return 0;
}

无论如何，如果您注释掉内核顶部的“s_blk”声明并取消注释紧随其后的注释掉的声明，那么您应该得到以下错误：

error : no instance of overloaded function "atomicMin" matches the argument list

我不明白为什么将共享内存声明为 volatile 会影响其类型，因为（我认为）此错误消息表明，也不明白为什么它不能与原子操作一起使用。

谁能提供任何见解？

谢谢，

亚伦

Answer 1

只需替换
atomicMin(&s_blk[ty][tx], 4);
由
atomicMin((int *)&s_blk[ty][tx], 4);.

它对&s_blk[ty][tx] 进行类型转换，因此它与atomicMin(..) 的参数相匹配。