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

Tue Apr 12 14:13:57 PDT 2022

tra accepted this revision.
tra added a comment.
This revision is now accepted and ready to land.

In D123441#3446499 <https://reviews.llvm.org/D123441#3446499>, @yaxunl wrote:

> You are talking about a use case that actually needs --whole-archive option. If the main TU does not reference some symbols in the archive but wants all symbols in the archive to be linked in, it is justifiable to use --whole-archive and HIP toolchain can support passing -Wl,--whole-archive specified in the command line to the device linking step.

Hmm. My point was the opposite -- only one object should be linked and I saw no way to do that without conservatively including everything.
I think I've misunderstood what your patch does.

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.

> However, in normal use cases, users only want to link in symbols referenced by the main TU. They do not need to link every symbol in the archive.

Agreed.

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

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

https://reviews.llvm.org/D123441