It appears that the nvcc compiler doesn’t detect correctly when someone calls __device__ functions from __host__ functions if both are templates. This should cause a compilation error, but doesn’t.
Instead of a compilation error, the code compiles and crashes when the __device__ function is called.
This was tested on nvcc-12.0
Very minimal code example:
// bug.cu
template <typename T>
__device__ int device_func() { return 0; }
template <typename T>
__host__ int host_func() { return device_func<T>(); }
int func() { return host_func<int>(); }
// main.cc
#include <iostream>
int func();
int main() {
std::cerr << "Before\n";
int res = func();
std::cerr << "After: " << res << "\n";
return 0;
}
compiling and running it with
nvcc bug.cu main.cc
./a.out
the compilation succeeds, and the program prints Before and then exists.
If we remove the template from either of these functions, we get the expected compilation error.
Is this a known issue? Is this fixed in newer nvcc versions?