[PATCH] D88786: [CUDA] Don't call __cudaRegisterVariable on C++17 inline variables

[Artem Belevich via Phabricator via cfe-commits]

tra added a comment.

In D88786#2312365 <https://reviews.llvm.org/D88786#2312365>, @MaskRay wrote:

>> Could you provide an example where this is causing an issue?
>
> If the C++17 inline variable appears in two TUs. They have the same comdat group. The first comdat group is prevailing and the second one is disarded. `__cudaRegisterVar(...)` in the second TU references a local symbol in a discarded section.

So, if I understand you correctly, it's the *host* side which ends up dropping it in one of TUs. It is a bit of a problem, considering that both of those TUs will need their own register call for their own GPU-side counterpart of the variable.

  a.h:
    __device__ inline int foo;
  a1.cu: #inlcude "a.h"
    a1.o/host : inline int foo; // 'shadow' variable. 
                register(foo, gpu-side-foo) // tell runtime that when we use host-side foo we want to access device-side foo.
    a1/GPU: int foo; // the only device-side instance. It's always there.
  a2.cu: #inlcude "a.h"
    a2.o/host : inline int foo; // 'shadow' variable. 
                register(foo, gpu-side-foo) // tell runtime that when we use host-side foo we want to access device-side foo.
    a2/GPU: int foo; // the only device-side instance. It's always there.
  
  host_exe: a1.o, a2.o
    only one instance of inline int foo survives and we lose ability to tell which GPU-side `foo` we want to access when we use host-side foo shadow.

Not allowing inline/constexpr variables seems to be the only choice here. Otherwise we's have to keep multiple instances of the shadow and that would break the C++ semantics for `inline` and `constexpr`

> The previous revision (https://reviews.llvm.org/D88786?id=295997 ) drops the comdat, but I think it is inferior to this one.

Silently dropping variable registration shifts the problem from link time to runtime. It may be OK as a temporary workaround for the build issues and only fraction of those will run into it at runtime, so it's technically an improvement, but we will need to catch it in Sema ASAP.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D88786