c – 带有动态共享内存的模板化CUDA内核

我想在一个程序中调用带有动态分配的共享内存的模板化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给我以下错误信息：

@H_734_2@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

我知道我遇到了@L_607_7@冲突,因为共享内存被声明为extern.然而,据我所知,如果我想在运行时定义它的大小,那就无法解决这个问题.

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

extern __shared__ T smem[];

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