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

概述我想在一个程序中调用带有动态分配的共享内存的模板化CUDA内核的不同实例.我的第一个天真的方法是写： template<typename T>__global__ void kernel(T* ptr){ extern __shared__ T smem[]; // calculations here ... 我想在一个程序中调用带有动态分配的共享内存的模板化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.然而,据我所知,如果我想在运行时定义它的大小,那就无法解决这个问题.

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

解决方法 动态分配的共享内存实际上只是一个大小(以字节为单位)和为内核设置的指针.所以像这样的东西应该工作：

替换这个：

extern __shared__ T smem[];

有了这个：

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

您可以在programming guide中看到重新构建动态分配的共享内存指针的其他示例,这些示例可以满足其他需求.

编辑：更新我的答案,以反映@njuffa的评论.

总结

以上是内存溢出为你收集整理的c – 带有动态共享内存的模板化CUDA内核全部内容，希望文章能够帮你解决c – 带有动态共享内存的模板化CUDA内核所遇到的程序开发问题。

如果觉得内存溢出网站内容还不错，欢迎将内存溢出网站推荐给程序员好友。