1

Consider the following code:

constexpr __host__   void foo() {  }

__global__ void baz()
{
    if constexpr(1==2) { foo(); }
}

this fails to compile with CUDA 11.3.1's NVCC. However, if I remove the constexpr - it does compile.

How come? Is this a bug?

einpoklum
  • 118,144
  • 57
  • 340
  • 684
  • 2
    it's a [stated limitation](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#constexpr-functions). The reason things work when you remove the first of your two `constexpr` keywords, is because another aspect of C++ allows it to. But since execution space decorators are not part of the C++ language, the compiler may enforce those rules as it sees fit, and in the order it sees fit. If you think its a bug, feel free to [file a bug](https://forums.developer.nvidia.com/t/how-to-report-a-bug/67911). – Robert Crovella Jun 17 '21 at 01:10
  • @RobertCrovella: But: my code doesn't break the limitation: It doesn't call foo from the `__device__` version of the function. If it were considered to be calling it, then the second version should not have compiled either. – einpoklum Jun 17 '21 at 11:31

1 Answers1

1

This is a bug. I've filed it with NVIDIA:

https://developer.nvidia.com/nvidia_bug/3328502

(it's difficult to access their bug tracking system, try going through https://developer.nvidia.com first to login, then look at your bugs, then replace the bug number.)

A fix is apparently now in the works!

einpoklum
  • 118,144
  • 57
  • 340
  • 684