[PATCH] D79237: [CUDA][HIP] Fix constexpr variables for C++17

[Artem Belevich via Phabricator via cfe-commits]

tra added a comment.

> constexpr variables are compile time constants and implicitly const, therefore
>  they are safe to emit on both device and host side. Besides, in many cases
>  they are intended for both device and host, therefore it makes sense
>  to emit them on both device and host sides if necessary.
> 
> In most cases constexpr variables are used as rvalue and the variables
>  themselves do not need to be emitted.

Agreed so far. constexpr values are usable on both sides.

> However if their address is taken, then they need to be emitted.

Agreed.

That's where the fun starts. Will this assertion trigger or not?

constexpr int ce = 42;
__global__ void kernel() {

  assert(p == &ce); 

}
void f() {

  kernel<<<1,1>>>(&ce);

}

  NVCC does not allow accessing ce, so this does not compile.
  Clang allows taking a reference, but the variable is not emitted, so the failure will happen in ptxas.
  If compiler does emit ce on device side, the addresses will be different, unless we rely on 
   CUDA runtime magic to translate addresses of shadows into real device-side address.
  If we do that, then we should probably add an implicit __constant__ on the constexpr vars.
  
  > 
  > The following example shows constexpr is available on device side without
  > __device__ or __constant__ attribute
  > 
  > https://godbolt.org/z/Uf7CgK
  
  If you clear output filters, you will see that the reference to constexpr is emitted as `.extern .global .align 4 .u32 _ZN1B1bE;` and the varible's address is not actually available. If you were to actually compile the example, ptxas would complain about ti.
  NVCC, predictably, allows using constexpr values, but not the variables themselves: https://godbolt.org/z/DXy65T


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D79237/new/

https://reviews.llvm.org/D79237