D17779: host-side shadow variables of external declarations of device-side
global variables have internal linkage and are referenced by
nvcc from CUDA 11 does not allow __device__ inline or __device__ constexpr
(C++17 inline variables) but clang has incorrectly supported them for a while:
error: A __device__ variable cannot be marked constexpr error: An inline __device__/__constant__/__managed__ variable must have internal linkage when the program is compiled in whole program mode (-rdc=false)
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:
A symbol table entry with STB_LOCAL binding that is defined relative to one of a group's sections, and that is contained in a symbol table section that is not part of the group, must be discarded if the group members are discarded. References to this symbol table entry from outside the group are not allowed.
As a workaround, don't register such inline variables for now.
(If we register the variables in all TUs, we will keep multiple instances of the shadow and break the C++ semantics for inline variables).
We should reject such variables in Sema but our internal users need some time to migrate.