我正在尝试用MSVS 2012和CUDA编译一个包含内核的程序。我使用共享内存,但与这个关于同一问题的问题不同,我只为这个内核的共享内存使用了一次变量名,所以不存在重新定义的问题。代码如下:
template<typename T>
__global__ void mykernel(
const T* __restrict__ data,
T* __restrict__ results)
{
extern __shared__ T warp_partial_results[];
/* ... */
warp_partial_results[lane_id] = something;
/* ... */
results[something_else] = warp_partial_results[something_else];
/* ... */
}
它被实例化为几种类型(例如float、int、unsigned int),我得到了可怕的
declaration is incompatible with previous "warp_partial_results"
消息。是什么原因造成的?
CUDA不会立即"支持"模板化函数中动态分配的共享内存数组,因为它(显然)会生成这些外部的实际定义。如果为多个类型实例化一个模板化函数,则定义会发生冲突。
可以通过类以模板专门化的形式提供一种变通方法。您可以选择NVIDIA的实现,也可以选择下面提到的更方便的实现。
NVIDIA实现
参见:
http://www.ecse.rpi.edu/~wrf/wiki/PallelComputingSpring2015/cuda/nvidia/samples/0_Simple/simpleTemplates/sharedmem.cuh
您使用的解决方法如下:
template<class T> __global__ void foo( T* g_idata, T* g_odata)
{
// shared memory
// the size is determined by the host application
SharedMem<T> shared;
T* sdata = shared.getPointer();
// .. the rest of the code remains unchanged!
}
getPointer()
对每种类型都有一个专门的实现,它返回不同的指针,例如extern __shared__ float* shared_mem_float
或extern __shared__ int* shared_mem_int
等。
更好的实现
在我自己的库达卡特图书馆里,有一个设施。你只需写:
auto foo = kat::shared_memory::dynamic::proxy<T>();
并且CCD_ 4是对共享存储器的CCD_。你也可以写:
auto n = kat::shared_memory::dynamic::size<T>();
它为您提供了适合于分配的动态共享内存的T类型元素的数量。
当然,我偏爱自己的解决方案,所以——选择适合你的方案。
(*)-不是。在NVidia提供的头文件中,它们专门用于一些基本类型,仅此而已