访问 __device__ 函数模板的专门实例中的 CUDA 库函数答案

【问题标题】：Access to CUDA library functions inside specialized instantiations of __device__ function templates访问 __device__ 函数模板的专门实例中的 CUDA 库函数
【发布时间】：2013-06-20 22:43:38
【问题描述】：

我在 CUDA 中有以下模板 __device__ 函数：

template<typename T>
__device__ void MyatomicAdd(T *address, T val){
atomicAdd(address , val);
}

如果将 T 实例化为浮点数，则编译和运行都很好，即

__global__ void myKernel(float *a, float b){
MyatomicAdd<float>(a,b);
}

将毫无问题地运行。

我想专门化这个函数，因为没有 atomicAdd() 用于双精度，所以我可以手动编写双精度实现。暂时忽略双精度特化，单精度特化和模板如下所示：

template<typename T>
__device__ void MyatomicAdd(T *address, T val){
};


template<>
__device__ void MyatomicAdd<float>(float *address, float val){
atomicAdd(address , val);
}

现在编译器抱怨 atomicAdd() 在我的专业中未定义，当我尝试在专业中使用任何 CUDA 函数（如 __syncthreads() ）时也是如此。有任何想法吗？谢谢。

【问题讨论】：

听起来您正在尝试使用错误的编译器来编译专业化。您真的在使用 nvcc，正在编译的文件是否具有 .cu 扩展名？
第一个非专业例子和第二个专业例子都存在于同一个文件中。在写完第二个示例后，我刚刚注释掉了第一个示例。这似乎是专业化的具体问题。
问题是我可以拿你贴的代码，用内核把它粘在一个文件里，用nvcc编译它，没有错误。因此，您忽略了告诉我们一些不同的事情..
我会尝试将它放在不同的文件中。同一个 cu 文件中还有一些其他内核和设备函数，所以可能某处存在一些干扰。
它确实用另一个内核在不同的 cu 文件中编译，搜索仍在继续..

标签： cuda template-specialization

【解决方案1】：

它最终成为与同事开发的一些 OpenGL 代码的链接问题。强制内联专业化解决了问题，尽管显然不是根本原因。不过，在我懒得去挖掘其他人的代码之前，它现在还是可以的。

【讨论】：