c++ - 获取CUDA错误 “declaration is incompatible with previous ” variable_name“

Question

我正在尝试使用MSVS 2012和CUDA编译包含内核的程序。我使用共享内存，但是与this question中的相同问题不同，我只为该内核的共享内存使用了变量名一次，因此没有重新定义的问题。用这样的代码:

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"

信息。是什么原因造成的？

Answer 1

CUDA不会立即“支持”模板化函数中的动态分配的共享内存阵列，因为它(显然)会生成这些外部函数的实际定义。如果为多个类型实例化模板化函数，则定义将发生冲突。

可以通过类以模板专用化的形式提供解决方法。看到:

http://www.ecse.rpi.edu/~wrf/wiki/ParallelComputingSpring2015/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等。

(*) - 并不是的。在NVidia提供的头文件中，它们专门用于某些基本类型，仅此而已。