Skip to content

[NewOffloadModel] Remove compiler backend option and linker option to be passed as argument for ClangLinkerWrapper #20691

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
YixingZhang007 wants to merge 14 commits into
base: sycl
Choose a base branch
from
+52 −86
Open

[NewOffloadModel] Remove compiler backend option and linker option to be passed as argument for ClangLinkerWrapper #20691

Show file tree
Hide file tree
Changes from 8 commits
Commits
Show all changes
14 commits
Select commit Hold shift + click to select a range
2ee12b0
enable new offloading model by default
YixingZhang007 Nov 18, 2025
1488454
resolve the issue of backend compile option got passed into ld
YixingZhang007 Nov 19, 2025
da1c270
remove linker and backend option to be passed into clang-linker-wrapper
YixingZhang007 Nov 20, 2025
0452eb7
code clean up
YixingZhang007 Nov 20, 2025
71c475b
fix the CI failure
YixingZhang007 Nov 25, 2025
2993152
remove the backend option to be passed into clanglinkerwrapper, and a…
YixingZhang007 Nov 25, 2025
7d0ee38
add passing the link time compiler option for cpu as well
YixingZhang007 Dec 2, 2025
e88a3fb
fix driver test failure
YixingZhang007 Dec 2, 2025
f4730f1
code clean up
YixingZhang007 Dec 2, 2025
2f11978
code clean up
YixingZhang007 Dec 2, 2025
fe390ac
code clean up
YixingZhang007 Dec 2, 2025
7232ae9
resolve some of the comment
YixingZhang007 Dec 7, 2025
eea2fd1
resolve comment
YixingZhang007 Dec 7, 2025
fd60624
add the compiler option to device-compiler
YixingZhang007 Dec 10, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
7 changes: 7 additions & 0 deletions clang/include/clang/Driver/Driver.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -626,6 +626,13 @@ class Driver {
/// @name Helper Methods
/// @{

/// Utility function to parse all devices passed via -fsycl-targets.
/// Return 'true' for JIT, AOT Intel CPU/GPUs and NVidia/AMD targets.
/// Otherwise return 'false'.
bool
GetUseNewOffloadDriverForSYCLOffload(Compilation &C,
const llvm::opt::ArgList &Args) const;

/// getSYCLDeviceTriple - Returns the SYCL device triple for the
/// specified subarch
// TODO: Additional Arg input parameter is for diagnostic output information
Expand Down
37 changes: 32 additions & 5 deletions clang/lib/Driver/Driver.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1565,6 +1565,32 @@ static void appendOneArg(InputArgList &Args, const Arg *Opt) {
}
}

// Utility function to parse all devices passed via -fsycl-targets.
// Return 'true' for JIT, AOT Intel CPU/GPUs and NVidia/AMD targets.
// Otherwise return 'false'.
bool Driver::GetUseNewOffloadDriverForSYCLOffload(Compilation &C,
const ArgList &Args) const {
// Check only if enabled with -fsycl
if (!Args.hasFlag(options::OPT_fsycl, options::OPT_fno_sycl, false))
return false;

if (Args.hasFlag(options::OPT_no_offload_new_driver,
options::OPT_offload_new_driver, false))
return false;

if (Args.hasArg(options::OPT_fintelfpga))
return false;

if (const Arg *A = Args.getLastArg(options::OPT_fsycl_targets_EQ)) {
for (const char *Val : A->getValues()) {
llvm::Triple TT(C.getDriver().getSYCLDeviceTriple(Val));
if ((!TT.isSPIROrSPIRV()) || TT.isSPIRAOT())
return false;
}
}
return true;
}

bool Driver::readConfigFile(StringRef FileName,
llvm::cl::ExpansionContext &ExpCtx) {
// Try opening the given file.
Expand Down Expand Up @@ -2195,12 +2221,12 @@ Compilation *Driver::BuildCompilation(ArrayRef<const char *> ArgList) {
// Use new offloading path for OpenMP. This is disabled as the SYCL
// offloading path is not properly setup to use the updated device linking
// scheme.
if ((C->isOffloadingHostKind(Action::OFK_OpenMP) &&
TranslatedArgs->hasFlag(options::OPT_fopenmp_new_driver,
options::OPT_no_offload_new_driver, true)) ||
if (C->isOffloadingHostKind(Action::OFK_OpenMP) ||
TranslatedArgs->hasFlag(options::OPT_offload_new_driver,
options::OPT_no_offload_new_driver, false))
options::OPT_no_offload_new_driver, false) ||
GetUseNewOffloadDriverForSYCLOffload(*C, *TranslatedArgs)) {
setUseNewOffloadingDriver();
}

// Construct the list of abstract actions to perform for this compilation. On
// MachO targets this uses the driver-driver and universal actions.
Expand Down Expand Up @@ -7095,7 +7121,8 @@ void Driver::BuildDefaultActions(Compilation &C, DerivedArgList &Args,
options::OPT_fno_offload_via_llvm, false) ||
Args.hasFlag(options::OPT_offload_new_driver,
options::OPT_no_offload_new_driver,
C.isOffloadingHostKind(Action::OFK_Cuda));
C.isOffloadingHostKind(Action::OFK_Cuda)) ||
GetUseNewOffloadDriverForSYCLOffload(C, Args);

bool HIPNoRDC =
C.isOffloadingHostKind(Action::OFK_HIP) &&
Expand Down
70 changes: 22 additions & 48 deletions clang/lib/Driver/ToolChains/Clang.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5232,7 +5232,9 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
(JA.isHostOffloading(C.getActiveOffloadKinds()) &&
Args.hasFlag(options::OPT_offload_new_driver,
options::OPT_no_offload_new_driver,
C.isOffloadingHostKind(Action::OFK_Cuda)));
C.isOffloadingHostKind(Action::OFK_Cuda))) ||
(JA.isHostOffloading(Action::OFK_SYCL) &&
C.getDriver().GetUseNewOffloadDriverForSYCLOffload(C, Args));

bool IsRDCMode =
Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc, IsSYCL);
Expand Down Expand Up @@ -11245,14 +11247,16 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA,
if (Kind == Action::OFK_OpenMP && !Args.hasArg(OPT_no_offloadlib) &&
(TC->getTriple().isAMDGPU() || TC->getTriple().isNVPTX()))
LinkerArgs.emplace_back("-lompdevice");

// Forward all of these to the appropriate toolchain.
for (StringRef Arg : CompilerArgs)
CmdArgs.push_back(Args.MakeArgString(
"--device-compiler=" + TC->getTripleString() + "=" + Arg));
for (StringRef Arg : LinkerArgs)
CmdArgs.push_back(Args.MakeArgString(
"--device-linker=" + TC->getTripleString() + "=" + Arg));
if (!C.hasOffloadToolChain<Action::OFK_SYCL>()) {
Copy link
Contributor

@YuriPlyakhin YuriPlyakhin Dec 2, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is it possible to avoid this SYCL-specific customization? What happens if we keep this for SYCL as well?
If I understand correctly @mdtoguchi : "The expected behavior is for any device options to be handled at both compile and link time. ", we need to keep this code and pass user-specified options to clang-linker-wrapper same as for any other programming model.

Copy link
Contributor

@YuriPlyakhin YuriPlyakhin Dec 2, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Where does CompilerArgs and LinkerArgs come from? I mean, is it from some clang++ arguments?

Copy link
Contributor Author

@YixingZhang007 YixingZhang007 Dec 8, 2025 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The CompilerArgs comes from the clang++ argument specified after --offload-arch=xxx, such as '-v' and '-O3', as implemented at https://github.com/YixingZhang007/llvm/blob/eea2fd15b33a90c010bf8a8a2629b63c485a816b/clang/lib/Driver/ToolChains/Clang.cpp#L11241. An example can also be found at

llvm/clang/test/Driver/amdgpu-openmp-toolchain.c

Lines 77 to 79 in af40a52

// RUN: %clang -### -target x86_64-pc-linux-gnu -fopenmp --offload-arch=gfx90a \
// RUN: -O3 -nogpulib %s 2>&1 | FileCheck %s --check-prefix=CHECK-OPT
// CHECK-OPT: clang-linker-wrapper{{.*}}"--device-compiler=amdgcn-amd-amdhsa=-O3"
.

If we add CompilerArgs to --device-compiler=, it will later be passed into clang-linker-wrapper and used for options for ocloc, which I don't think we will want that. For SYCL, we only want to pass the argument after -Xsycl-target-backend=spir64_gen and -Xsycl-target-backend=spir64_x86_64 -backend-cpu-opt.

Copy link
Contributor

@YuriPlyakhin YuriPlyakhin Dec 8, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

If we add CompilerArgs to --device-compiler=, it will later be passed into clang-linker-wrapper

Is not it the behavior that is expected?

for (StringRef Arg : CompilerArgs)
CmdArgs.push_back(Args.MakeArgString(
"--device-compiler=" + TC->getTripleString() + "=" + Arg));
for (StringRef Arg : LinkerArgs)
CmdArgs.push_back(Args.MakeArgString(
"--device-linker=" + TC->getTripleString() + "=" + Arg));
}

// Forward the LTO mode relying on the Driver's parsing.
if (C.getDriver().getOffloadLTOMode() == LTOK_Full)
Expand Down Expand Up @@ -11458,57 +11462,27 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA,
CmdArgs.push_back(
Args.MakeArgString("-sycl-allow-device-image-dependencies"));

// Formulate and add any offload-wrapper and AOT specific options. These
// are additional options passed in via -Xsycl-target-linker and
// -Xsycl-target-backend.
// For AOT, pass along backend target args via --device-compiler options
// to the clang-linker-wrapper.
const toolchains::SYCLToolChain &SYCLTC =
static_cast<const toolchains::SYCLToolChain &>(getToolChain());
// Only store compile/link opts in the image descriptor for the SPIR-V
// target. For AOT, pass along the addition options via GPU or CPU
// specific clang-linker-wrapper options.
const ArgList &Args =
C.getArgsForToolChain(nullptr, StringRef(), Action::OFK_SYCL);
for (auto &ToolChainMember :
llvm::make_range(ToolChainRange.first, ToolChainRange.second)) {
const ToolChain *TC = ToolChainMember.second;
bool IsJIT = false;
StringRef WrapperOption;
StringRef WrapperLinkOption;
if (TC->getTriple().isSPIROrSPIRV()) {
if (TC->getTriple().getSubArch() == llvm::Triple::NoSubArch) {
IsJIT = true;
WrapperOption = "--sycl-backend-compile-options=";
}
if (TC->getTriple().getSubArch() == llvm::Triple::SPIRSubArch_gen)
WrapperOption = "--gpu-tool-arg=";
if (TC->getTriple().getSubArch() == llvm::Triple::SPIRSubArch_x86_64)
WrapperOption = "--cpu-tool-arg=";
} else
continue;
ArgStringList BuildArgs;
SmallString<128> BackendOptString;
SmallString<128> LinkOptString;
SYCLTC.TranslateBackendTargetArgs(TC->getTriple(), Args, BuildArgs);
for (const auto &A : BuildArgs)
appendOption(BackendOptString, A);

BuildArgs.clear();
SYCLTC.TranslateLinkerTargetArgs(TC->getTriple(), Args, BuildArgs);
for (const auto &A : BuildArgs) {
if (IsJIT)
appendOption(LinkOptString, A);
else
// For AOT, combine the Backend and Linker strings into one.
SmallString<128> BackendOptString;
if (TC->getTriple().getSubArch() == llvm::Triple::SPIRSubArch_gen || (TC->getTriple().getSubArch() == llvm::Triple::SPIRSubArch_x86_64)) {
for (const auto &A : BuildArgs)
appendOption(BackendOptString, A);
CmdArgs.push_back(Args.MakeArgString(
"--device-compiler=" + TC->getTripleString() + "=" + BackendOptString));
}
if (!BackendOptString.empty())
CmdArgs.push_back(
Args.MakeArgString(Twine(WrapperOption) + BackendOptString));
if (!LinkOptString.empty())
CmdArgs.push_back(
Args.MakeArgString("--sycl-target-link-options=" + LinkOptString));
}

// Add option to enable creating of the .syclbin file.
const ArgList &Args =
C.getArgsForToolChain(nullptr, StringRef(), Action::OFK_SYCL);
if (Arg *A = Args.getLastArg(options::OPT_fsyclbin_EQ))
CmdArgs.push_back(
Args.MakeArgString("--syclbin=" + StringRef{A->getValue()}));
Expand Down
3 changes: 2 additions & 1 deletion clang/test/Driver/clang-linker-wrapper.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -74,7 +74,8 @@
// Check that when --gpu-tool-arg is specified in clang-linker-wrapper
// (happen when AOT device is specified via -Xsycl-target-backend '-device pvc' in clang),
// the target is not passed to sycl-post-link for filtering.
// RUN: clang-linker-wrapper -sycl-embed-ir -sycl-device-libraries=%t1.devicelib.o -sycl-post-link-options="SYCL_POST_LINK_OPTIONS" -llvm-spirv-options="LLVM_SPIRV_OPTIONS" "--host-triple=x86_64-unknown-linux-gnu" "--gpu-tool-arg=-device pvc" "--linker-path=/usr/bin/ld" "--" HOST_LINKER_FLAGS "-dynamic-linker" HOST_DYN_LIB "-o" "a.out" HOST_LIB_PATH HOST_STAT_LIB %t1.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-NO-CMDS-AOT-GEN %s
// RUN: clang-linker-wrapper -sycl-embed-ir -sycl-device-libraries=%t1.devicelib.o -sycl-post-link-options="SYCL_POST_LINK_OPTIONS" -llvm-spirv-options="LLVM_SPIRV_OPTIONS" "--host-triple=x86_64-unknown-linux-gnu" "--device-compiler=spir64_gen-unknown-unknown=-device pvc" "--linker-path=/usr/bin/ld" "--" HOST_LINKER_FLAGS "-dynamic-linker" HOST_DYN_LIB "-o" "a.out" HOST_LIB_PATH HOST_STAT_LIB %t1.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-NO-CMDS-AOT-GEN %s
//
// CHK-NO-CMDS-AOT-GEN: sycl-post-link{{.*}} SYCL_POST_LINK_OPTIONS -o {{[^,]*}}.table {{.*}}.bc

/// Check for list of commands for standalone clang-linker-wrapper run for sycl (AOT for Intel CPU)
Expand Down
17 changes: 2 additions & 15 deletions clang/test/Driver/sycl-offload-new-driver.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -141,27 +141,14 @@
// MULT_TARG_PHASES: 15: backend, {14}, assembler, (host-sycl)
// MULT_TARG_PHASES: 16: assembler, {15}, object, (host-sycl)

/// Test option passing behavior for clang-offload-wrapper options.
// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver \
// RUN: -Xsycl-target-backend -backend-opt -### %s 2>&1 \
Copy link
Contributor

@YuriPlyakhin YuriPlyakhin Dec 2, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Per my understanding, we still need to pass -backend-opt to clang-linker-wrapper through device-compiler option.
so this test should be just updated to use proper clang-linker-wrapper option.
Similar for linker options.

Copy link
Contributor Author

@YixingZhang007 YixingZhang007 Dec 2, 2025 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for pointing this out!

I think for backend option and linker option that are passed at compile time, they are stored in the SYCLImage like --image=file=... compile-opts=... -backend-opt. The option that are stored in the SYCLImage got directly extracted in clang-linker-wrapper (this is done is another PR #19579 and extractSYCLCompileLinkOptions function found at https://github.com/YixingZhang007/llvm/blob/fe390ac353c65ec85ccd2b3bf7883a1efca55a5e/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp#L2218).

Only the -backend-opt that are passed at link time for AOT, which should be through format clang++ ... -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen -backend-gen-opt will added to device-compiler and passed to the clang-linker-wrapper.

Therefore, for this test and the one got deleted below, -backend-opt and -link-opt are passed through the SYCLImage but not directly added as argument for clang-linker-wrapper :)

Copy link
Contributor

@YuriPlyakhin YuriPlyakhin Dec 3, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes, I understand how it works with your implementation. And I challenge that. Despite we extracted and put options to SYCLImage on compile stage, we still want to pass these options to link stage as well.
Example of when this is needed is if we provide both source and object files as an input to clang++.

// RUN: | FileCheck -check-prefix WRAPPER_OPTIONS_BACKEND %s
// WRAPPER_OPTIONS_BACKEND: clang-linker-wrapper{{.*}} "--sycl-backend-compile-options={{.*}}-backend-opt{{.*}}"

// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver \
// RUN: -Xsycl-target-linker -link-opt -### %s 2>&1 \
// RUN: | FileCheck -check-prefix WRAPPER_OPTIONS_LINK %s
// WRAPPER_OPTIONS_LINK: clang-linker-wrapper{{.*}} "--sycl-target-link-options={{.*}}-link-opt{{.*}}"

/// Test option passing behavior for clang-offload-wrapper options for AOT.
/// Test option passing behavior for clang-offload-wrapper options for AOT
// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver \
// RUN: -fsycl-targets=spir64_gen,spir64_x86_64 \
// RUN: -Xsycl-target-backend=spir64_gen -backend-gen-opt \
// RUN: -Xsycl-target-backend=spir64_x86_64 -backend-cpu-opt \
// RUN: -### %s 2>&1 \
// RUN: | FileCheck -check-prefix WRAPPER_OPTIONS_BACKEND_AOT %s
// WRAPPER_OPTIONS_BACKEND_AOT: clang-linker-wrapper{{.*}} "--host-triple=x86_64-unknown-linux-gnu"
// WRAPPER_OPTIONS_BACKEND_AOT-SAME: "--gpu-tool-arg=-backend-gen-opt"
// WRAPPER_OPTIONS_BACKEND_AOT-SAME: "--cpu-tool-arg=-backend-cpu-opt"
// WRAPPER_OPTIONS_BACKEND_AOT: clang-linker-wrapper{{.*}} "--host-triple=x86_64-unknown-linux-gnu" {{.*}} "--device-compiler=spir64_gen-unknown-unknown=-backend-gen-opt" "--device-compiler=spir64_x86_64-unknown-unknown=-backend-cpu-opt"

/// Verify arch settings for nvptx and amdgcn targets
// RUN: %clangxx -fsycl -### -fsycl-targets=amdgcn-amd-amdhsa -fno-sycl-libspirv \
Expand Down
9 changes: 0 additions & 9 deletions clang/test/Driver/sycl-offload.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -325,11 +325,6 @@
// CHK-NO-FSYCL-TARGET-ERROR-NOT: clang{{.*}} error: cannot deduce implicit triple value for '-Xsycl-target-frontend', specify triple using '-Xsycl-target-frontend=<triple>'

/// ###########################################################################

// RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown -Xsycl-target-backend "-DFOO1 -DFOO2" %s 2>&1 \
// RUN: | FileCheck -check-prefix=CHK-TOOLS-OPTS %s
// CHK-TOOLS-OPTS: clang-linker-wrapper{{.*}} "--sycl-backend-compile-options=-DFOO1 -DFOO2"
Copy link
Contributor

@YuriPlyakhin YuriPlyakhin Dec 2, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

similar here, per my understanding, need to pass options using device-compiler


/// Check for implied options (-g -O0)
// RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown -g -O0 -Xsycl-target-backend "-DFOO1 -DFOO2" %s 2>&1 \
// RUN: | FileCheck -check-prefix=CHK-TOOLS-IMPLIED-OPTS %s
Expand All @@ -346,10 +341,6 @@
// RUN: | FileCheck -check-prefix=CHK-TOOLS-IMPLIED-OPTS-O0 %s
// CHK-TOOLS-IMPLIED-OPTS-O0-NOT: llvm-offload-binary{{.*}} {{.*}}compile-opts={{.*}}-cl-opt-disable"

// RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown -Xsycl-target-linker "-DFOO1 -DFOO2" %s 2>&1 \
// RUN: | FileCheck -check-prefix=CHK-TOOLS-OPTS2 %s
// CHK-TOOLS-OPTS2: clang-linker-wrapper{{.*}} "--sycl-target-link-options=-DFOO1 -DFOO2"
Copy link
Contributor

@YuriPlyakhin YuriPlyakhin Dec 2, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

same here, but for device-linker


/// -fsycl-range-rounding settings
///
/// // Check that driver flag is passed to cc1
Expand Down
12 changes: 2 additions & 10 deletions clang/test/Driver/sycl-offload.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -52,21 +52,13 @@
// RUN: | FileCheck -check-prefix DEFAULT_LINK %s
// DEFAULT_LINK: clang-linker-wrapper{{.*}}

/// Passing in the default triple should allow for -Xsycl-target options, both the
/// "=<triple>" and the default spelling
// RUN: %clangxx -### -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver -fsycl-targets=spir64 -Xsycl-target-backend=spir64 -DFOO -Xsycl-target-linker=spir64 -DFOO2 %S/Inputs/SYCL/objlin64.o 2>&1 \
// RUN: | FileCheck -check-prefixes=SYCL_TARGET_OPT %s
// RUN: %clangxx -### -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver -Xsycl-target-backend=spir64 -DFOO -Xsycl-target-linker=spir64 -DFOO2 %S/Inputs/SYCL/objlin64.o 2>&1 \
// RUN: | FileCheck -check-prefixes=SYCL_TARGET_OPT %s
// SYCL_TARGET_OPT: clang-linker-wrapper{{.*}} "--sycl-backend-compile-options=-DFOO" "--sycl-target-link-options=-DFOO2"
Copy link
Contributor

@YuriPlyakhin YuriPlyakhin Dec 2, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

same here...


// RUN: %clangxx -### -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver -fsycl-targets=spir64_x86_64 -Xsycl-target-backend -DFOO %S/Inputs/SYCL/objlin64.o 2>&1 \
// RUN: | FileCheck -check-prefixes=SYCL_TARGET_OPT_AOT,SYCL_TARGET_OPT_CPU %s
// RUN: %clangxx -### -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver -fsycl-targets=spir64_gen -Xsycl-target-backend -DFOO %S/Inputs/SYCL/objlin64.o 2>&1 \
// RUN: | FileCheck -check-prefixes=SYCL_TARGET_OPT_AOT,SYCL_TARGET_OPT_GPU %s
// SYCL_TARGET_OPT_AOT-NOT: error: cannot deduce implicit triple value for '-Xsycl-target-backend'
// SYCL_TARGET_OPT_CPU: clang-linker-wrapper{{.*}} "--cpu-tool-arg=-DFOO"
// SYCL_TARGET_OPT_GPU: clang-linker-wrapper{{.*}} "--gpu-tool-arg=-DFOO"
// SYCL_TARGET_OPT_CPU: clang-linker-wrapper{{.*}} "--device-compiler=spir64_x86_64-unknown-unknown=-DFOO"
// SYCL_TARGET_OPT_GPU: clang-linker-wrapper{{.*}} "--device-compiler=spir64_gen-unknown-unknown=-DFOO"

/// Check -fsycl-targets=spir64 enables addition of -ffine-grained-bitfield-accesses option
// RUN: %clangxx -### -fsycl-device-only --offload-new-driver %s 2>&1 | FileCheck -check-prefixes=CHECK_BITFIELD_OPTION %s
Expand Down
27 changes: 17 additions & 10 deletions clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -716,7 +716,7 @@ getTripleBasedSYCLPostLinkOpts(const ArgList &Args,
/// code and will be parsed to generate options required to be passed into the
/// sycl-post-link tool.
static Expected<std::vector<module_split::SplitModule>>
runSYCLPostLinkTool(ArrayRef<StringRef> InputFiles, const ArgList &Args) {
runSYCLPostLinkTool(ArrayRef<StringRef> InputFiles, const ArgList &Args, bool HasGPUTool) {
Expected<std::string> SYCLPostLinkPath =
findProgram("sycl-post-link", {getMainExecutable("sycl-post-link")});
if (!SYCLPostLinkPath)
Expand All @@ -731,14 +731,13 @@ runSYCLPostLinkTool(ArrayRef<StringRef> InputFiles, const ArgList &Args) {

// Enable the driver to invoke sycl-post-link with the device architecture
// when Intel GPU targets are passed in -fsycl-targets.
// OPT_gpu_tool_arg_EQ is checked to ensure the device architecture is not
// HasGPUTool is checked to ensure the device architecture is not
// passed through -Xsycl-target-backend=spir64_gen "-device <arch>" format
const llvm::Triple Triple(Args.getLastArgValue(OPT_triple_EQ));
StringRef Arch = Args.getLastArgValue(OPT_arch_EQ);
StringRef IsGPUTool = Args.getLastArgValue(OPT_gpu_tool_arg_EQ);

if (Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen && !Arch.empty() &&
IsGPUTool.empty() && Arch != "*")
!HasGPUTool && Arch != "*")
OutputPathWithArch = "intel_gpu_" + Arch.str() + "," + OutputPathWithArch;
else if (Triple.getSubArch() == llvm::Triple::SPIRSubArch_x86_64)
OutputPathWithArch = "spir64_x86_64," + OutputPathWithArch;
Expand Down Expand Up @@ -982,10 +981,6 @@ static void addSYCLBackendOptions(const ArgList &Args,
CmdArgs.push_back(Args.MakeArgString(JoinedOptions));
}
}

StringRef OptTool = (IsCPU) ? Args.getLastArgValue(OPT_cpu_tool_arg_EQ)
: Args.getLastArgValue(OPT_gpu_tool_arg_EQ);
OptTool.split(CmdArgs, " ", /*MaxSplit=*/-1, /*KeepEmpty=*/false);
return;
}

Expand Down Expand Up @@ -2222,6 +2217,15 @@ linkAndWrapDeviceFiles(ArrayRef<SmallVector<OffloadFile>> LinkerInputFiles,
extractSYCLCompileLinkOptions(Input);
if (!CompileLinkOptionsOrErr)
return CompileLinkOptionsOrErr.takeError();

// Append any additional backend compiler options specified at link time.
const llvm::Triple Triple(LinkerArgs.getLastArgValue(OPT_triple_EQ));
for (StringRef Arg : LinkerArgs.getAllArgValues(OPT_device_compiler_args_EQ)) {
auto [ArgTriple, ArgValue] = Arg.split('=');
if (ArgTriple == Triple.getTriple() && !ArgValue.empty()) {
CompileLinkOptionsOrErr->first += Twine(" ", ArgValue).str();
}
}

SmallVector<StringRef> InputFiles;
// Write device inputs to an output file for the linker.
Expand All @@ -2238,16 +2242,19 @@ linkAndWrapDeviceFiles(ArrayRef<SmallVector<OffloadFile>> LinkerInputFiles,
return TmpOutputOrErr.takeError();
SmallVector<StringRef> InputFilesSYCL;
InputFilesSYCL.emplace_back(*TmpOutputOrErr);

SmallVector<StringRef, 16> Args;
StringRef(CompileLinkOptionsOrErr->first).split(Args, ' ');
bool HasGPUTool = std::find(Args.begin(), Args.end(), "-device") != Args.end();
auto SplitModulesOrErr =
UseSYCLPostLinkTool
? sycl::runSYCLPostLinkTool(InputFilesSYCL, LinkerArgs)
? sycl::runSYCLPostLinkTool(InputFilesSYCL, LinkerArgs, HasGPUTool)
: sycl::runSYCLSplitLibrary(InputFilesSYCL, LinkerArgs,
*SYCLModuleSplitMode);
if (!SplitModulesOrErr)
return SplitModulesOrErr.takeError();

auto &SplitModules = *SplitModulesOrErr;
const llvm::Triple Triple(LinkerArgs.getLastArgValue(OPT_triple_EQ));
bool IsJIT = Triple.isSPIROrSPIRV() &&
Triple.getSubArch() == llvm::Triple::NoSubArch;
if ((Triple.isNVPTX() || Triple.isAMDGCN()) &&
Expand Down
10 changes: 0 additions & 10 deletions clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -199,16 +199,6 @@ def sycl_add_default_spec_consts_image : Flag<["--", "-"], "sycl-add-default-spe
def no_sycl_add_default_spec_consts_image : Flag<["--", "-"], "no-sycl-add-default-spec-consts-image">,
Flags<[WrapperOnlyOption, HelpHidden]>;

// Special options to pass backend options required for AOT compilation
def gpu_tool_arg_EQ :
Joined<["--", "-"], "gpu-tool-arg=">,
Flags<[WrapperOnlyOption]>,
HelpText<"Options that are passed to the backend of target device compiler for Intel GPU during AOT compilation">;
def cpu_tool_arg_EQ :
Joined<["--", "-"], "cpu-tool-arg=">,
Flags<[WrapperOnlyOption]>,
HelpText<"Options that are passed to the backend of target device compiler for Intel CPU during AOT compilation">;

def sycl_thin_lto : Flag<["--", "-"], "sycl-thin-lto">,
Flags<[WrapperOnlyOption]>, HelpText<"Link SYCL device code using thinLTO">;

Expand Down
Loading