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