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

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl added a comment.

In D79237#2031870 <https://reviews.llvm.org/D79237#2031870>, @tra wrote:

> > 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/55dEx-


It seems adding an implicit `__constant__` attribute to constexpr variables on device side may be a possible solution. I will give it a try.


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

https://reviews.llvm.org/D79237