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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
