出现 CUDA 错误“声明与之前的“variable_name”不兼容

Question

我正在尝试编译一个包含 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"

信息。什么可能导致这种情况？

Answer 1

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

可以通过类进行模板专门化的形式提供解决方法。您可以选择 NVIDIA 的实现，或者下面提到的更好方便的实现。

NVIDIA 实施

看：

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_floatextern __shared__ int* shared_mem_int}

更好的实现

在我自己的cuda-kat图书馆中，有一个用于此目的的设施。你只需写：

auto foo = kat::shared_memory::dynamic::proxy<T>();

是你们共同的记忆foo。T*你也可以写：

auto n = kat::shared_memory::dynamic::size<T>();

它可以获取适合分配的动态共享内存的 T 类型元素的数量。

当然，我偏爱我自己的解决方案，所以 - 选择适合你的解决方案。

_{（*） - 并不真地。在 NVidia 提供的头文件中，它们专门用于一些基本类型，仅此而已。}