[all-commits] [llvm/llvm-project] a2cc88: [CUDA] Don't call __cudaRegisterVariable on C++17 ...

[Fangrui Song via All-commits]

  Branch: refs/heads/master
  Home:   https://github.com/llvm/llvm-project
  Commit: a2cc8833683dd124cf2ee96f6d17f7f835da1fc8
      https://github.com/llvm/llvm-project/commit/a2cc8833683dd124cf2ee96f6d17f7f835da1fc8
  Author: Fangrui Song <i at maskray.me>
  Date:   2020-10-05 (Mon, 05 Oct 2020)

  Changed paths:
    M clang/lib/CodeGen/CodeGenModule.cpp
    M clang/test/CodeGenCUDA/device-stub.cu

  Log Message:
  -----------
  [CUDA] Don't call __cudaRegisterVariable on C++17 inline variables

D17779: host-side shadow variables of external declarations of device-side
global variables have internal linkage and are referenced by
`__cuda_register_globals`.

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.

Reviewed By: tra

Differential Revision: https://reviews.llvm.org/D88786