我有一个 CUDA 内核,需要在易失性共享整数内存上使用原子函数。但是,当我尝试将共享内存声明为 volatile 并在原子函数中使用它时,我收到一条错误消息。
下面是一些重现错误的极简代码。请注意,下面的内核什么都不做,并且严重滥用了为什么您想将共享内存声明为易失性(甚至根本不使用共享内存)。但它确实重现了错误。
该代码在共享内存上使用原子函数,因此,要运行它,您可能需要使用“arch12”或更高版本进行编译(在 Visual Studio 2010 中,右键单击您的项目并转到“属性 -> 配置属性 -> CUDA C /C++ -> Device”并在“代码生成”行中输入“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 会影响其类型,因为(我认为)此错误消息表明,也不明白为什么它不能与原子操作一起使用。
任何人都可以提供任何见解吗?
谢谢,
亚伦