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