我想在一个程序中调用模板化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。然而，据我所知，如果我想在运行时定义它的大小，就没有办法了。
所以我的问题是**是否有任何优雅的方式来获得所需的行为？**优雅的意思是没有代码重复等。
2条答案
按热度按时间gab6jxml1#
动态分配的共享内存实际上只是一个大小（以字节为单位）和一个为内核设置的指针。这样的东西应该可以工作：
替换此：
用这个：
你可以在编程指南中看到重新转换动态分配的共享内存指针的其他例子，它们可以满足其他需求。
**编辑：**更新了我的回答，以反映@njuffa的评论。
qq24tv8q2#
（@RobertCrovella的answer的变体）
NVCC不愿意接受两个名称相同但类型不同的
extern __shared__数组--即使它们从来不在对方的作用域中。为了满足NVCC的要求，我们需要让模板示例都在底层使用相同的共享内存类型，同时让使用它们的内核代码看到它喜欢的类型。所以我们替换这个指令：
用这个
其中：
在一些设备端代码包括文件中。
优点：
extern，或对齐说明符，或重新解释转换等。注意事项：
shared_memory.cuh（在这里它被命名为shared_memory::dynamic::proxy()）。