c++ - CUDA9.2及更高版本中模板默认参数 Unresolved external 函数错误

Question

我正在处理一些 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 --

据我所知，错误仅在使用默认模板参数时发生 Link_force<Pt> force = linear_force<Pt>在 apply_forces 的模板定义中.例如在main中显式指定模板参数

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

我们称之为 apply_forces 的地方将导致一切都正确编译和运行，就像以任何其他方式明确定义模板参数一样。

这可能是 nvcc 工具链的问题吗？我没有在 CUDA 发行说明中发现任何可能是罪魁祸首的更改，所以我有点难过。

因为这适用于旧版本的 nvcc，现在不行了，我不明白这是否实际上是对模板默认参数的非法使用？ (也许特别是与 CUDA 函数结合使用时？)

Answer 1

这是 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 版本中提供。