c++ - 具有动态共享内存的模板化 CUDA 内核

Question

我想在一个程序中调用具有动态分配的共享内存的模板化 CUDA 内核的不同实例化。我第一个天真的方法是写:

template<typename T>
__global__ void kernel(T* ptr)
{
  extern __shared__ T smem[];
  // calculations here ...                                                                                                                                          
}

template<typename T>
void call_kernel( T* ptr, const int n )
{
  dim3 dimBlock(n), dimGrid;
  kernel<<<dimGrid, dimBlock, n*sizeof(T)>>>(ptr);
}

int main(int argc, char *argv[])
{
  const int n = 32;
  float *float_ptr;
  double *double_ptr;
  cudaMalloc( (void**)&float_ptr, n*sizeof(float) );
  cudaMalloc( (void**)&double_ptr, n*sizeof(double) );

  call_kernel( float_ptr, n );
  call_kernel( double_ptr, n ); // problem, 2nd instantiation

  cudaFree( (void*)float_ptr );
  cudaFree( (void*)double_ptr );
  return 0;
}

但是，这段代码无法编译。 nvcc 给我以下错误消息:

main.cu(4): error: declaration is incompatible with previous "smem"
(4): here
          detected during:
            instantiation of "void kernel(T *) [with T=double]"
(12): here
            instantiation of "void call_kernel(T *, int) [with T=double]"
(24): here

我知道我遇到了名称冲突，因为共享内存被声明为 extern。然而，据我所知，如果我想在运行时定义它的大小，就没有办法解决这个问题。

所以，我的问题是:是否有任何优雅的方式来获得所需的行为？优雅的意思是没有代码重复等。

Answer 1

动态分配的共享内存实际上只是一个大小(以字节为单位)和一个为内核设置的指针。所以这样的事情应该有效:

替换这个:

extern __shared__ T smem[];

用这个:

extern __shared__ __align__(sizeof(T)) unsigned char my_smem[];
T *smem = reinterpret_cast<T *>(my_smem);

您可以在 programming guide 中看到重新转换动态分配的共享内存指针的其他示例可以满足其他需求。

编辑:更新了我的答案以反射(reflect)@njuffa 的评论。