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

Mon Oct 5 11:31:05 PDT 2020

MaskRay added a comment.

In D88786#2312329 <https://reviews.llvm.org/D88786#2312329>, @tra wrote:

>> If such a variable (which has a comdat group) is discarded (a copy from another
>> translation unit is prevailing and selected), accessing the variable from
>> outside the section group (__cuda_register_globals) is a violation of the ELF
>> specification and will be rejected by linkers:
>
> Every TU  is the whole program on the GPU side, provided we compile w/o `-frdc`, so there's no other TU to prevail.
> I don't have a good idea yet what's the best way to handle this in CUDA, but not registering the variables will likely to create other issues, only visible at runtime. E.g. some host-side code will attempt to use cudaMemcpy() on the symbol and will fail, because it's not been registered, even though we do have all other glue in place.
>
> 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.

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

Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D88786