https://github.com/pvelesko updated
https://github.com/llvm/llvm-project/pull/136412
>From fe6a426fc135c7232650b5ebac465ceaa66d7a20 Mon Sep 17 00:00:00 2001
From: Paulius Velesko <pvele...@pglc.io>
Date: Sat, 19 Apr 2025 10:02:59 +0300
Subject: [PATCH 1/3] HIPSPV: Unbundle SDL
This fixes the issue of rdc linking static libraries with device code
---
clang/lib/Driver/ToolChains/HIPSPV.cpp | 9 +++++++++
1 file changed, 9 insertions(+)
diff --git a/clang/lib/Driver/ToolChains/HIPSPV.cpp
b/clang/lib/Driver/ToolChains/HIPSPV.cpp
index bbde5c047c3f9..eaf20f34edb72 100644
--- a/clang/lib/Driver/ToolChains/HIPSPV.cpp
+++ b/clang/lib/Driver/ToolChains/HIPSPV.cpp
@@ -70,6 +70,15 @@ void HIPSPV::Linker::constructLinkAndEmitSpirvCommand(
// Link LLVM bitcode.
ArgStringList LinkArgs{};
+
+ // Add static device libraries using the common helper function.
+ // This handles unbundling archives (.a) containing bitcode bundles.
+ const HIPSPVToolChain &TC = static_cast<const HIPSPVToolChain
&>(getToolChain());
+ StringRef Arch = TC.getTriple().getArchName();
+ StringRef Target = "generic"; // SPIR-V is generic, no specific target ID
like -mcpu
+ tools::AddStaticDeviceLibsLinking(C, *this, JA, Inputs, Args, LinkArgs, Arch,
+ Target, /*IsBitCodeSDL=*/true);
+
for (auto Input : Inputs)
LinkArgs.push_back(Input.getFilename());
LinkArgs.append({"-o", TempFile});
>From 116d6894516cff02c11b69380f4f86c3208d0b06 Mon Sep 17 00:00:00 2001
From: Paulius Velesko <pvele...@pglc.io>
Date: Tue, 17 Jun 2025 15:23:54 +0300
Subject: [PATCH 2/3] Update clang/lib/Driver/ToolChains/HIPSPV.cpp
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit
Co-authored-by: Henry Linjamäki <henry.linjam...@gmail.com>
---
clang/lib/Driver/ToolChains/HIPSPV.cpp | 3 +--
1 file changed, 1 insertion(+), 2 deletions(-)
diff --git a/clang/lib/Driver/ToolChains/HIPSPV.cpp
b/clang/lib/Driver/ToolChains/HIPSPV.cpp
index eaf20f34edb72..ae485a2688b9a 100644
--- a/clang/lib/Driver/ToolChains/HIPSPV.cpp
+++ b/clang/lib/Driver/ToolChains/HIPSPV.cpp
@@ -73,8 +73,7 @@ void HIPSPV::Linker::constructLinkAndEmitSpirvCommand(
// Add static device libraries using the common helper function.
// This handles unbundling archives (.a) containing bitcode bundles.
- const HIPSPVToolChain &TC = static_cast<const HIPSPVToolChain
&>(getToolChain());
- StringRef Arch = TC.getTriple().getArchName();
+ StringRef Arch = getToolChain().getTriple().getArchName();
StringRef Target = "generic"; // SPIR-V is generic, no specific target ID
like -mcpu
tools::AddStaticDeviceLibsLinking(C, *this, JA, Inputs, Args, LinkArgs, Arch,
Target, /*IsBitCodeSDL=*/true);
>From 9e4e9f3b04c748e30039ff384a456a140db74198 Mon Sep 17 00:00:00 2001
From: Paulius Velesko <pvele...@pglc.io>
Date: Tue, 17 Jun 2025 16:27:22 +0300
Subject: [PATCH 3/3] Address PR comments: Fix HIPSPV SDL linking
* Also ordering and add --no-offloadlib support
---
clang/lib/Driver/ToolChains/CommonArgs.cpp | 5 +++
clang/lib/Driver/ToolChains/HIPSPV.cpp | 6 +--
.../Driver/hipspv-link-static-library.hip | 38 +++++++++++++++++++
3 files changed, 46 insertions(+), 3 deletions(-)
create mode 100644 clang/test/Driver/hipspv-link-static-library.hip
diff --git a/clang/lib/Driver/ToolChains/CommonArgs.cpp
b/clang/lib/Driver/ToolChains/CommonArgs.cpp
index 109316d0a27e7..2f0235cec4c31 100644
--- a/clang/lib/Driver/ToolChains/CommonArgs.cpp
+++ b/clang/lib/Driver/ToolChains/CommonArgs.cpp
@@ -2622,6 +2622,11 @@ void tools::AddStaticDeviceLibs(Compilation *C, const
Tool *T,
llvm::opt::ArgStringList &CC1Args,
StringRef Arch, StringRef Target,
bool isBitCodeSDL) {
+
+ // Check if offload libraries are disabled
+ if (!DriverArgs.hasFlag(options::OPT_offloadlib, options::OPT_no_offloadlib,
+ true))
+ return;
SmallVector<std::string, 8> LibraryPaths;
// Add search directories from LIBRARY_PATH env variable
diff --git a/clang/lib/Driver/ToolChains/HIPSPV.cpp
b/clang/lib/Driver/ToolChains/HIPSPV.cpp
index ae485a2688b9a..15e7c2c9e9418 100644
--- a/clang/lib/Driver/ToolChains/HIPSPV.cpp
+++ b/clang/lib/Driver/ToolChains/HIPSPV.cpp
@@ -71,15 +71,15 @@ void HIPSPV::Linker::constructLinkAndEmitSpirvCommand(
// Link LLVM bitcode.
ArgStringList LinkArgs{};
+ for (auto Input : Inputs)
+ LinkArgs.push_back(Input.getFilename());
+
// Add static device libraries using the common helper function.
// This handles unbundling archives (.a) containing bitcode bundles.
StringRef Arch = getToolChain().getTriple().getArchName();
StringRef Target = "generic"; // SPIR-V is generic, no specific target ID
like -mcpu
tools::AddStaticDeviceLibsLinking(C, *this, JA, Inputs, Args, LinkArgs, Arch,
Target, /*IsBitCodeSDL=*/true);
-
- for (auto Input : Inputs)
- LinkArgs.push_back(Input.getFilename());
LinkArgs.append({"-o", TempFile});
const char *LlvmLink =
Args.MakeArgString(getToolChain().GetProgramPath("llvm-link"));
diff --git a/clang/test/Driver/hipspv-link-static-library.hip
b/clang/test/Driver/hipspv-link-static-library.hip
new file mode 100644
index 0000000000000..5447d35014bbf
--- /dev/null
+++ b/clang/test/Driver/hipspv-link-static-library.hip
@@ -0,0 +1,38 @@
+// Test HIPSPV static device library linking
+// REQUIRES: system-linux
+// UNSUPPORTED: system-windows
+
+// Create a dummy archive to test SDL linking
+// RUN: rm -rf %t && mkdir %t
+// RUN: touch %t/dummy.bc
+// RUN: llvm-ar cr %t/libSDL.a %t/dummy.bc
+
+// Test that -l options are passed to llvm-link for --offload=spirv64
+// RUN: %clang -### --target=x86_64-linux-gnu --offload=spirv64 \
+// RUN: --hip-path=%S/Inputs/hipspv -nohipwrapperinc %s \
+// RUN: -L%t -lSDL \
+// RUN: 2>&1 | FileCheck -check-prefix=SDL-LINK %s
+
+// Test that .a files are properly unbundled and passed to llvm-link
+// RUN: %clang -### --target=x86_64-linux-gnu --offload=spirv64 \
+// RUN: --hip-path=%S/Inputs/hipspv -nohipwrapperinc %s \
+// RUN: %t/libSDL.a \
+// RUN: 2>&1 | FileCheck -check-prefix=SDL-ARCHIVE %s
+
+// Verify that the input files are added before the SDL files in llvm-link
command
+// This tests the ordering fix to match HIPAMD behavior
+// SDL-LINK: "{{.*}}clang-offload-bundler" "-unbundle" "-type=a"
"-input={{.*}}libSDL.a" "-targets=hip-spirv64-unknown-unknown-unknown-generic"
"-output=[[SDL_A:.*\.a]]" "-allow-missing-bundles"
+// SDL-LINK: "{{.*}}llvm-link" "{{.*}}.bc" "[[SDL_A]]" "-o"
+
+// SDL-ARCHIVE: "{{.*}}clang-offload-bundler" "-unbundle" "-type=a"
"-input={{.*}}libSDL.a" "-targets=hip-spirv64-unknown-unknown-unknown-generic"
"-output=[[SDL_A:.*\.a]]" "-allow-missing-bundles"
+// SDL-ARCHIVE: "{{.*}}llvm-link" "{{.*}}.bc" "[[SDL_A]]" "-o"
+
+// Test that no SDL linking occurs when --no-offloadlib is used
+// RUN: %clang -### --target=x86_64-linux-gnu --offload=spirv64 \
+// RUN: --hip-path=%S/Inputs/hipspv -nohipwrapperinc --no-offloadlib %s \
+// RUN: -L%t -lSDL \
+// RUN: 2>&1 | FileCheck -check-prefix=NO-SDL %s
+
+// NO-SDL-NOT: "{{.*}}clang-offload-bundler" "-unbundle" "-type=a"
"-input={{.*}}libSDL.a"
+
+__global__ void kernel() {}
\ No newline at end of file
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
