Skip to content

Commit 9e4e9f3

Browse files
committed
Address PR comments: Fix HIPSPV SDL linking
* Also ordering and add --no-offloadlib support
1 parent 116d689 commit 9e4e9f3

File tree

3 files changed

+46
-3
lines changed

3 files changed

+46
-3
lines changed

clang/lib/Driver/ToolChains/CommonArgs.cpp

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2622,6 +2622,11 @@ void tools::AddStaticDeviceLibs(Compilation *C, const Tool *T,
26222622
llvm::opt::ArgStringList &CC1Args,
26232623
StringRef Arch, StringRef Target,
26242624
bool isBitCodeSDL) {
2625+
2626+
// Check if offload libraries are disabled
2627+
if (!DriverArgs.hasFlag(options::OPT_offloadlib, options::OPT_no_offloadlib,
2628+
true))
2629+
return;
26252630

26262631
SmallVector<std::string, 8> LibraryPaths;
26272632
// Add search directories from LIBRARY_PATH env variable

clang/lib/Driver/ToolChains/HIPSPV.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -71,15 +71,15 @@ void HIPSPV::Linker::constructLinkAndEmitSpirvCommand(
7171
// Link LLVM bitcode.
7272
ArgStringList LinkArgs{};
7373

74+
for (auto Input : Inputs)
75+
LinkArgs.push_back(Input.getFilename());
76+
7477
// Add static device libraries using the common helper function.
7578
// This handles unbundling archives (.a) containing bitcode bundles.
7679
StringRef Arch = getToolChain().getTriple().getArchName();
7780
StringRef Target = "generic"; // SPIR-V is generic, no specific target ID like -mcpu
7881
tools::AddStaticDeviceLibsLinking(C, *this, JA, Inputs, Args, LinkArgs, Arch,
7982
Target, /*IsBitCodeSDL=*/true);
80-
81-
for (auto Input : Inputs)
82-
LinkArgs.push_back(Input.getFilename());
8383
LinkArgs.append({"-o", TempFile});
8484
const char *LlvmLink =
8585
Args.MakeArgString(getToolChain().GetProgramPath("llvm-link"));
Lines changed: 38 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,38 @@
1+
// Test HIPSPV static device library linking
2+
// REQUIRES: system-linux
3+
// UNSUPPORTED: system-windows
4+
5+
// Create a dummy archive to test SDL linking
6+
// RUN: rm -rf %t && mkdir %t
7+
// RUN: touch %t/dummy.bc
8+
// RUN: llvm-ar cr %t/libSDL.a %t/dummy.bc
9+
10+
// Test that -l options are passed to llvm-link for --offload=spirv64
11+
// RUN: %clang -### --target=x86_64-linux-gnu --offload=spirv64 \
12+
// RUN: --hip-path=%S/Inputs/hipspv -nohipwrapperinc %s \
13+
// RUN: -L%t -lSDL \
14+
// RUN: 2>&1 | FileCheck -check-prefix=SDL-LINK %s
15+
16+
// Test that .a files are properly unbundled and passed to llvm-link
17+
// RUN: %clang -### --target=x86_64-linux-gnu --offload=spirv64 \
18+
// RUN: --hip-path=%S/Inputs/hipspv -nohipwrapperinc %s \
19+
// RUN: %t/libSDL.a \
20+
// RUN: 2>&1 | FileCheck -check-prefix=SDL-ARCHIVE %s
21+
22+
// Verify that the input files are added before the SDL files in llvm-link command
23+
// This tests the ordering fix to match HIPAMD behavior
24+
// 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"
25+
// SDL-LINK: "{{.*}}llvm-link" "{{.*}}.bc" "[[SDL_A]]" "-o"
26+
27+
// 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"
28+
// SDL-ARCHIVE: "{{.*}}llvm-link" "{{.*}}.bc" "[[SDL_A]]" "-o"
29+
30+
// Test that no SDL linking occurs when --no-offloadlib is used
31+
// RUN: %clang -### --target=x86_64-linux-gnu --offload=spirv64 \
32+
// RUN: --hip-path=%S/Inputs/hipspv -nohipwrapperinc --no-offloadlib %s \
33+
// RUN: -L%t -lSDL \
34+
// RUN: 2>&1 | FileCheck -check-prefix=NO-SDL %s
35+
36+
// NO-SDL-NOT: "{{.*}}clang-offload-bundler" "-unbundle" "-type=a" "-input={{.*}}libSDL.a"
37+
38+
__global__ void kernel() {}

0 commit comments

Comments
 (0)