C++, Cuda and template using constexpr parameter -> linking error

Accelerated Computing CUDA CUDA NVCC Compiler
#1

Hi
I have a linking error, wihout Cuda or without constexpr it works (this is a minimized proof of concept here) :

struct CParams {
  constexpr CParams(int _n) : n(_n) {};

  const int n;
};

constexpr CParams cparams0(3);

template<const CParams &cparams> __global__ void f()
{
    int t=2;
}

int main()
{
    f< cparams0 ><<<1,32>>>();
    return 0;
}

The error is :

/usr/local/cuda-11/bin/nvcc --device-debug --debug -gencode arch=compute_86,code=sm_86 -gencode arch=compute_86,code=compute_86 -ccbin g++ -c -o "main.o" "../main.cu"

In file included from tmpxft_00012d25_00000000-6_main.cudafe1.stub.c:1:
/tmp/tmpxft_00012d25_00000000-6_main.cudafe1.stub.c:14:41: error: template-id ‘__wrapper__device_stub_f<(& cparams0)>’ for ‘void __wrapper__device_stub_f()’ does not match any template declaration
   14 | template<> __specialization_static void __wrapper__device_stub_f<(&cparams0)>(void){__device_stub__Z1fIL_ZN57_INTERNAL_39_tmpxft_00012d25_00000000_7_main_cpp1_ii_main8cparams0EEEvv();}
      |                                         ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
../main.cu:10:47: note: candidate is: ‘template void __wrapper__device_stub_f()’
   10 | template __global__ void f()
      |                                               ^~~~~~

The cost to not use a constexpr parameter would be very high. I don’t know if I’m wrong or if it’s a compilator limitation ? It works well without cuda.

Versions & OS :

nvcc: Built on Mon_Oct_11_21:27:02_PDT_2021 Cuda compilation tools, release 11.4, V11.4.152 Build cuda_11.4.r11.4/compiler.30521435_0
gcc : (Debian 10.2.1-6) 10.2.1 20210110

Thanks