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