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

Mon Oct 5 11:23:14 PDT 2020

tra added a comment.

> 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?

Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D88786