[PATCH] D123441: [CUDA][HIP] Fix host used external kernel in archive

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl marked an inline comment as done.
yaxunl added a comment.

In D123441#3446719 <https://reviews.llvm.org/D123441#3446719>, @tra wrote:

> So, a main TU with just `__global__ void kernel();` would emit a reference when it's compiled on the GPU side. That, in turn will tell the linker what it needs to pull from the libraries and things should just work.
> If that's the case, then it would work in my example, too.

Yes, this patch creates artificial references in the device IR originating from the host functions in the same TU, or, in other words, it creates the missing references which should be there but are not there and no more references than those.



================
Comment at: clang/lib/CodeGen/CodeGenModule.cpp:602
+        getModule(), ATy, false, llvm::GlobalValue::AppendingLinkage,
+        llvm::ConstantArray::get(ATy, UsedArray), "hip.used.external");
+    addCompilerUsedGlobal(GV);
----------------
tra wrote:
> This is not HIP-specific and should have a more generic name. `@gpu.used.external` ?
will do when committing.


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

https://reviews.llvm.org/D123441