forked from OSchip/llvm-project
a2cc883368
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 |
||
|---|---|---|
| .. | ||
| Inputs | ||
| address-spaces.cu | ||
| alias.cu | ||
| amdgpu-hip-implicit-kernarg.cu | ||
| amdgpu-kernel-arg-pointer-type.cu | ||
| amdgpu-kernel-attrs.cu | ||
| amdgpu-visibility.cu | ||
| amdgpu-workgroup-size.cu | ||
| builtins-amdgcn.cu | ||
| constexpr-variables.cu | ||
| convergent.cu | ||
| cuda-builtin-vars.cu | ||
| debug-info-address-class.cu | ||
| debug-info-template.cu | ||
| deferred-diag.cu | ||
| dependent-libs.cu | ||
| device-init-fun.cu | ||
| device-stub.cu | ||
| device-var-init.cu | ||
| device-vtable.cu | ||
| dft-func-attr-skip-intrinsic.hip | ||
| filter-decl.cu | ||
| flush-denormals.cu | ||
| fp-contract.cu | ||
| function-overload.cu | ||
| kernel-amdgcn.cu | ||
| kernel-args-alignment.cu | ||
| kernel-args.cu | ||
| kernel-call.cu | ||
| kernel-dbg-info.cu | ||
| kernel-stub-name.cu | ||
| lambda.cu | ||
| launch-bounds.cu | ||
| library-builtin.cu | ||
| link-device-bitcode.cu | ||
| llvm-used.cu | ||
| ms-linker-options.cu | ||
| norecurse.cu | ||
| nothrow.cu | ||
| openmp-target.cu | ||
| printf-aggregate.cu | ||
| printf.cu | ||
| profile-coverage-mapping.cu | ||
| propagate-metadata.cu | ||
| ptx-kernels.cu | ||
| static-device-var-no-rdc.cu | ||
| surface.cu | ||
| texture.cu | ||
| types.cu | ||
| unnamed-types.cu | ||
| usual-deallocators.cu | ||