[PATCH] D88786: [CUDA] Don't call __cudaRegisterVariable on C++17 inline variables

[Fangrui Song via Phabricator via cfe-commits]

MaskRay marked an inline comment as done.
MaskRay added inline comments.


================
Comment at: clang/lib/CodeGen/CodeGenModule.cpp:4137
+        // TODO: Reject __device__ constexpr and __device__ inline in Sema.
+        if (!D->hasExternalStorage() && !D->isInline())
           getCUDARuntime().registerDeviceVar(D, *GV, !D->hasDefinition(),
----------------
tra wrote:
> > For rdc mode, device vars are merged. Host shadow vars should also be in comdat and merged. 
> 
> Right. I think we need to add `|| (getLangOpts().HIP && getLangOpts().GPURelocatableDeviceCode)`. Maybe even for both CUDA and HIP as `rdc` should work similarly in CUDA, too.
I don't know `-rdc=true`. Hope @tra and @yaxunl can make the change with a description.

I confirm that `__device__ inline int` works under nvcc with -rdc=true but I don't know its implication with `__cudaRegisterVariable`.
constexpr is still forbidden.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D88786/new/

https://reviews.llvm.org/D88786