Instantiating a CUDA kernel with a type predicated on `__CUDA_ARCH__` causes `cudaErrorSymbolNotFound`

[OgnianM]

Example:

#include <iostream>
#include <typeinfo>
#include <cxxabi.h>

template<typename T>
__global__ void SomeKernel(T) {}

template<int>
struct SomeType { };

int main() {
#ifdef __CUDA_ARCH__
    using T = SomeType<0>;
#else
    using T = SomeType<1>;
#endif
    SomeKernel<T><<<1,1>>>({});

    if (cudaGetLastError() == cudaErrorSymbolNotFound) {
        std::cout << abi::__cxa_demangle(typeid(SomeKernel<T>).name(),0,0,0) << " not found!" << std::endl;
    }
}

[Artem-B]

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#cuda-arch

The type signature of the following entities shall not depend on whether CUDA_ARCH is defined or not, or on a particular value of CUDA_ARCH:

• global functions and function templates
• device and constant variables textures and surfaces

Host and GPU-side compilation must have exactly the same idea about the kernel name.
Your code defines SomeKernel<SomeType<0>> on the GPU side, but tries to call SomeKernel<SomeType<1>> from the host. Predictably, that's not going to work, as we've never instantiated SomeKernel<SomeType<1>> on the GPU.