Неразрешенная ошибка внешней функции с параметром шаблона по умолчанию в CUDA9.2 и выше

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.