================
@@ -53,19 +53,19 @@ void AMDGCN::Linker::constructLldCommand(Compilation &C,
const JobAction &JA,
const llvm::opt::ArgList &Args) const
{
// Construct lld command.
// The output from ld.lld is an HSA code object file.
- ArgStringList LldArgs{"-flavor",
- "gnu",
- "-m",
- "elf64_amdgpu",
- "--no-undefined",
- "-shared",
- "-plugin-opt=-amdgpu-internalize-symbols"};
+ const ToolChain &TC = getToolChain();
+ LTOKind LTOMode = TC.getLTOMode(Args, Action::OFK_HIP);
+ ArgStringList LldArgs{"-flavor", "gnu", "-m", "elf64_amdgpu",
+ "--no-undefined", "-shared"};
+ // Native object references are invisible when LTO compiles bitcode libraries
+ // in a non-LTO link, so library definitions must remain external.
+ if (LTOMode != LTOK_None)
+ LldArgs.push_back("-plugin-opt=-amdgpu-internalize-symbols");
----------------
yxsamliu wrote:
Thanks, Joseph. You’re right that LTO runs and handles internalization here.
The case I’m trying to fix is device-only offload PGO without `-foffload-lto`.
A reduced test case is:
```cpp
#define __global__ __attribute__((global))
extern "C" __global__ void kernel(int *out) { out[0] = 42; }
```
```sh
clang++ -x hip --offload-device-only --offload-arch=gfx1100 \
-fprofile-generate -nogpuinc -nogpulib -v -save-temps \
kernel.hip -o kernel.co
```
On current main, Clang generates a native AMDGPU object for the kernel, then
links it with the bitcode profile library. The object comes from this normal
driver flow; it isn’t a separate object supplied by the user. LLD’s normal LTO
keeps the profile helpers referenced by that object, but the extra
`-amdgpu-internalize-symbols` pass sees only the bitcode and removes them. The
link fails with undefined symbols. Omitting just that extra pass makes it work.
https://github.com/llvm/llvm-project/pull/225859
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
