CUDA9.2及以上版本中模板默认参数存在无法解决的外部函数错误

Unresolved extern function error with template default parameter in CUDA9.2 and above

本文关键字：存在 解决 错误 函数 外部 参数 默认 版本 CUDA9      更新时间：2023-10-16

我正在使用一些c++/CUDA代码，这些代码大量使用类和函数的模板。我们主要使用CUDA 9.0和9.1，在这两个版本中，一切都可以很好地编译和运行。但是，CUDA的较新版本(特别是9.2和10)的编译失败。

经过进一步调查，试图用CUDA 9.2.88及以上版本编译完全相同的代码似乎会失败，而用CUDA 8至9.1.85版本编译代码并正确运行。

问题代码的一个最小示例可以写如下：

#include <iostream>
template<typename Pt>
using Link_force = void(Pt* x, Pt* y);
template<typename Pt>
__device__ void linear_force(Pt* x, Pt* y)
{
*x += *y;
}
template<typename Pt, Link_force<Pt> force>
__global__ void link(Pt* x, Pt* y)
{
force(x, y);
}
template<typename Pt = float, Link_force<Pt> force = linear_force<Pt>>
void apply_forces(Pt* x, Pt* y)
{
link<Pt, force><<<1, 1, 0>>>(x, y);
}
int main(int argc, const char* argv[])    
{
float *x, *y;
cudaMallocManaged(&x, sizeof(float));
cudaMallocManaged(&y, sizeof(float));
*x = 0.0f;
*y = 42.0f;
std::cout << "Pre :: x = " << *x << ", y = " << *y << 'n';
apply_forces(x, y);
cudaDeviceSynchronize();
std::cout << "Post :: x = " << *x << ", y = " << *y << 'n';
return 0;
}

如果我使用nvcc编译，如下所示，最终结果是ptxas:的错误

$ nvcc --verbose -std=c++11 -arch=sm_61 minimal_example.cu 
#$ _SPACE_= 
#$ _CUDART_=cudart
#$ _HERE_=/usr/local/cuda-9.2/bin
#$ _THERE_=/usr/local/cuda-9.2/bin
#$ _TARGET_SIZE_=
#$ _TARGET_DIR_=
#$ _TARGET_SIZE_=64
#$ TOP=/usr/local/cuda-9.2/bin/..
#$ NVVMIR_LIBRARY_DIR=/usr/local/cuda-9.2/bin/../nvvm/libdevice
#$ LD_LIBRARY_PATH=/usr/local/cuda-9.2/bin/../lib:/usr/local/cuda-9.2/lib64:
#$ PATH=/usr/local/cuda-9.2/bin/../nvvm/bin:/usr/local/cuda-9.2/bin:/usr/local/cuda-9.2/bin:/usr/local/sbin:/usr/local/bin:/usr/sbin:/usr/bin:/sbin:/bin:/usr/games:/usr/local/games:/snap/bin
#$ INCLUDES="-I/usr/local/cuda-9.2/bin/..//include"  
#$ LIBRARIES=  "-L/usr/local/cuda-9.2/bin/..//lib64/stubs" "-L/usr/local/cuda-9.2/bin/..//lib64"
#$ CUDAFE_FLAGS=
#$ PTXAS_FLAGS=
#$ gcc -std=c++11 -D__CUDA_ARCH__=610 -E -x c++  -DCUDA_DOUBLE_MATH_FUNCTIONS -D__CUDACC__ -D__NVCC__  "-I/usr/local/cuda-9.2/bin/..//include"   -D"__CUDACC_VER_BUILD__=148" -D"__CUDACC_VER_MINOR__=2" -D"__CUDACC_VER_MAJOR__=9" -include "cuda_runtime.h" -m64 "minimal_example.cu" > "/tmp/tmpxft_0000119e_00000000-8_minimal_example.cpp1.ii" 
#$ cicc --c++11 --gnu_version=70300 --allow_managed   -arch compute_61 -m64 -ftz=0 -prec_div=1 -prec_sqrt=1 -fmad=1 --include_file_name "tmpxft_0000119e_00000000-2_minimal_example.fatbin.c" -tused -nvvmir-library "/usr/local/cuda-9.2/bin/../nvvm/libdevice/libdevice.10.bc" --gen_module_id_file --module_id_file_name "/tmp/tmpxft_0000119e_00000000-3_minimal_example.module_id" --orig_src_file_name "minimal_example.cu" --gen_c_file_name "/tmp/tmpxft_0000119e_00000000-5_minimal_example.cudafe1.c" --stub_file_name "/tmp/tmpxft_0000119e_00000000-5_minimal_example.cudafe1.stub.c" --gen_device_file_name "/tmp/tmpxft_0000119e_00000000-5_minimal_example.cudafe1.gpu"  "/tmp/tmpxft_0000119e_00000000-8_minimal_example.cpp1.ii" -o "/tmp/tmpxft_0000119e_00000000-5_minimal_example.ptx"
#$ ptxas -arch=sm_61 -m64  "/tmp/tmpxft_0000119e_00000000-5_minimal_example.ptx"  -o "/tmp/tmpxft_0000119e_00000000-9_minimal_example.sm_61.cubin" 
ptxas fatal   : Unresolved extern function '_Z12linear_forceIfEvPT_S1_'
# --error 0xff --

据我所知，只有在apply_forces的模板定义中使用默认模板参数Link_force<Pt> force = linear_force<Pt>时才会出现错误。例如，在主中显式指定模板参数

apply_forces<float, linear_force>(x, y);

我们调用apply_forces的位置将导致所有内容正确编译和运行，以任何其他方式显式定义模板参数也是如此。

这可能是nvcc工具链的问题吗？我没有发现CUDA发布说明中有任何可能是罪魁祸首的变化，所以我有点困惑。

由于这适用于旧版本的nvcc，而现在不是，我不明白这是否真的是对模板默认参数的非法使用？(可能特别是与CUDA功能结合使用时？)

这是CUDA 9.2和10.0中的一个错误，正在进行修复。感谢您指出。

正如你已经指出的，一个可能的解决方案是恢复到CUDA 9.1

另一种可能的解决方法是在函数体中重复有问题的模板实例化(例如，在丢弃的语句中)。这对性能没有影响，它只是迫使编译器发出该函数的代码：

template<typename Pt = float, Link_force<Pt> force = linear_force<Pt>>
void apply_forces(Pt* x, Pt* y)
{
(void)linear_force<Pt>;  // add this
link<Pt, force><<<1, 1, 0>>>(x, y);
}

我没有关于何时可以修复的进一步信息，但它将在未来的CUDA版本中发布。

相关文章：