From 1427e51493e823c687063b8621fa173dc56bdb61 Mon Sep 17 00:00:00 2001 From: "yixing.zhang" Date: Mon, 26 Jan 2026 16:14:45 +0100 Subject: [PATCH 01/14] initial draft --- clang/lib/Driver/ToolChains/Clang.cpp | 31 ++++++++++++++++++--------- clang/lib/Driver/ToolChains/SYCL.cpp | 10 +++++++-- 2 files changed, 29 insertions(+), 12 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 4fe5f0fd09e15..1f66ce693eb86 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -11614,27 +11614,38 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, if (!TC->getTriple().isSPIROrSPIRV()) continue; ArgStringList BuildArgs; - SmallString<128> BackendOptString; + SmallVector, 4> BackendOptStrings; SmallString<128> LinkOptString; - SYCLTC.TranslateBackendTargetArgs(TC->getTriple(), Args, BuildArgs); - for (const auto &A : BuildArgs) - appendOption(BackendOptString, A); - BuildArgs.clear(); + // Process each -Xsycl-target-backend individually + for (const Arg *A : Args.filtered(options::OPT_Xsycl_backend_EQ)) { + StringRef Device = SYCL::gen::resolveGenDevice(A->getValue(0)); + SmallString<128> BackendString; + + SYCLTC.TranslateBackendTargetArgs(TC->getTriple(), Args, BuildArgs, Device); + for (const auto &BA : BuildArgs) { + appendOption(BackendString, BA); + } + + BuildArgs.clear(); + BackendOptStrings.push_back(std::move(BackendString)); + } + SYCLTC.TranslateLinkerTargetArgs(TC->getTriple(), Args, BuildArgs); for (const auto &A : BuildArgs) { if (TC->getTriple().getSubArch() == llvm::Triple::NoSubArch) appendOption(LinkOptString, A); else - // For AOT, combine the Backend and Linker strings into one. - appendOption(BackendOptString, A); + // For AOT, append linker args to every backend string + for (auto &BOS : BackendOptStrings) + appendOption(BOS, A); } - if (!BackendOptString.empty()) { - CmdArgs.push_back(Args.MakeArgString( + for (auto &BOS : BackendOptStrings) { + CmdArgs.push_back(Args.MakeArgString( "--device-compiler=" + Action::GetOffloadKindName(Action::OFK_SYCL) + ":" + - TC->getTripleString() + "=" + BackendOptString)); + TC->getTripleString() + "=" + BOS)); } if (!LinkOptString.empty()) { CmdArgs.push_back(Args.MakeArgString( diff --git a/clang/lib/Driver/ToolChains/SYCL.cpp b/clang/lib/Driver/ToolChains/SYCL.cpp index a4db63c4fb52f..0bb9e7b2a7766 100644 --- a/clang/lib/Driver/ToolChains/SYCL.cpp +++ b/clang/lib/Driver/ToolChains/SYCL.cpp @@ -1558,8 +1558,12 @@ void SYCLToolChain::TranslateTargetOpt(const llvm::Triple &Triple, StringRef GenDevice = SYCL::gen::resolveGenDevice(A->getValue()); bool IsGenTriple = Triple.isSPIR() && Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen; + llvm::errs() << "[DEBUG] Args: " << A->getAsString(Args) + << " OptTargetTriple: " << OptTargetTriple.str() + << ", Triple: " << Triple.str() << ", Device: " << Device + << ", GenDevice: " << GenDevice << "\n"; if (IsGenTriple) { - if (Device != GenDevice && !Device.empty()) + if (Device != GenDevice) continue; if (OptTargetTriple != Triple && GenDevice.empty()) // Triples do not match, but only skip when we know we are not @@ -1678,7 +1682,9 @@ void SYCLToolChain::AddImpliedTargetArgs(const llvm::Triple &Triple, } // Check for any -device settings. std::string DevArg; - if (IsJIT || Device == "pvc" || hasPVCDevice(TargArgs, DevArg)) { + // if (IsJIT || Device == "pvc" || hasPVCDevice(TargArgs, DevArg)) { + llvm::errs() << "[DEBUG] AddImpliedTargetArgs Device " << Device << "\n"; + if (IsJIT || Device == "pvc") { // The -device option passed in by the user may not be 'pvc'. Use the // value provided by the user if it was specified. StringRef DeviceName = "pvc"; From 187dca0932086e09a4c6ed316d57e83ab5fecf02 Mon Sep 17 00:00:00 2001 From: "yixing.zhang" Date: Mon, 26 Jan 2026 18:47:28 +0100 Subject: [PATCH 02/14] now the passing of compile-opt passed for llvm-offload-binary is correct --- clang/lib/Driver/ToolChains/Clang.cpp | 6 ++++-- clang/lib/Driver/ToolChains/SYCL.cpp | 25 ++++++++++++++----------- 2 files changed, 18 insertions(+), 13 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 1f66ce693eb86..867bee416107d 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -10450,12 +10450,14 @@ void OffloadPackager::ConstructJob(Compilation &C, const JobAction &JA, const ToolChain *HostTC = C.getSingleOffloadToolChain(); const toolchains::SYCLToolChain &SYCLTC = static_cast(*TC); + + SYCLTC.AddImpliedTargetArgs(TC->getTriple(), Args, BuildArgs, JA, *HostTC, Arch); - SYCLTC.TranslateBackendTargetArgs(TC->getTriple(), Args, BuildArgs); + SYCLTC.TranslateBackendTargetArgs(TC->getTriple(), Args, BuildArgs, OffloadAction->getOffloadingArch()); createArgString("compile-opts="); BuildArgs.clear(); - SYCLTC.TranslateLinkerTargetArgs(TC->getTriple(), Args, BuildArgs); + SYCLTC.TranslateLinkerTargetArgs(TC->getTriple(), Args, BuildArgs, OffloadAction->getOffloadingArch()); createArgString("link-opts="); } diff --git a/clang/lib/Driver/ToolChains/SYCL.cpp b/clang/lib/Driver/ToolChains/SYCL.cpp index 0bb9e7b2a7766..62cd3abd4b27d 100644 --- a/clang/lib/Driver/ToolChains/SYCL.cpp +++ b/clang/lib/Driver/ToolChains/SYCL.cpp @@ -1558,20 +1558,25 @@ void SYCLToolChain::TranslateTargetOpt(const llvm::Triple &Triple, StringRef GenDevice = SYCL::gen::resolveGenDevice(A->getValue()); bool IsGenTriple = Triple.isSPIR() && Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen; - llvm::errs() << "[DEBUG] Args: " << A->getAsString(Args) - << " OptTargetTriple: " << OptTargetTriple.str() - << ", Triple: " << Triple.str() << ", Device: " << Device - << ", GenDevice: " << GenDevice << "\n"; + // llvm::errs() << "[DEBUG] Args: " << A->getAsString(Args) + // << " OptTargetTriple: " << OptTargetTriple.str() + // << ", Triple: " << Triple.str() << ", Device: " << Device + // << ", GenDevice: " << GenDevice << "\n"; if (IsGenTriple) { - if (Device != GenDevice) - continue; + if (Device != GenDevice) { + if(!GenDevice.empty()) { + continue; + } + SmallString<64> DeviceOpt("-device "); + DeviceOpt += Device; + if (A->getAsString(Args).find(DeviceOpt.str()) == std::string::npos) + continue; + } + // if (Device != GenDevice && !Device.empty()) if (OptTargetTriple != Triple && GenDevice.empty()) // Triples do not match, but only skip when we know we are not // comparing against intel_gpu_* continue; - if (OptTargetTriple == Triple && !Device.empty()) - // Triples match, but we are expecting a specific device to be set. - continue; } else if (OptTargetTriple != Triple) continue; } else if (!OptNoTriple) @@ -1682,8 +1687,6 @@ void SYCLToolChain::AddImpliedTargetArgs(const llvm::Triple &Triple, } // Check for any -device settings. std::string DevArg; - // if (IsJIT || Device == "pvc" || hasPVCDevice(TargArgs, DevArg)) { - llvm::errs() << "[DEBUG] AddImpliedTargetArgs Device " << Device << "\n"; if (IsJIT || Device == "pvc") { // The -device option passed in by the user may not be 'pvc'. Use the // value provided by the user if it was specified. From 774859c2c2e76f9cef7410dd01116ca46e5df42e Mon Sep 17 00:00:00 2001 From: "yixing.zhang" Date: Mon, 26 Jan 2026 22:50:05 +0100 Subject: [PATCH 03/14] pass the arch name to --device-compiler --- clang/lib/Driver/ToolChains/Clang.cpp | 32 ++++++++++++++------- clang/lib/Driver/ToolChains/SYCL.cpp | 41 +++++++++++++++++---------- clang/lib/Driver/ToolChains/SYCL.h | 1 + 3 files changed, 48 insertions(+), 26 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 867bee416107d..6f7eb0e6dfed3 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -11616,38 +11616,48 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, if (!TC->getTriple().isSPIROrSPIRV()) continue; ArgStringList BuildArgs; - SmallVector, 4> BackendOptStrings; + llvm::SmallMapVector, 4> BackendOptMap; SmallString<128> LinkOptString; // Process each -Xsycl-target-backend individually for (const Arg *A : Args.filtered(options::OPT_Xsycl_backend_EQ)) { - StringRef Device = SYCL::gen::resolveGenDevice(A->getValue(0)); - SmallString<128> BackendString; + StringRef Device; + StringRef TargetBackend = A->getValue(); + if((TargetBackend == "spir64_gen")) { + Device = SYCL::gen::extractDeviceFromArgSimple(A->getValue(1)); + llvm::errs() << "[DEBUG] Resolved device: " << Device << "\n"; + } else { + Device = SYCL::gen::resolveGenDevice(TargetBackend); + } + SmallString<128> &BackendString = BackendOptMap[Device]; SYCLTC.TranslateBackendTargetArgs(TC->getTriple(), Args, BuildArgs, Device); for (const auto &BA : BuildArgs) { appendOption(BackendString, BA); } BuildArgs.clear(); - BackendOptStrings.push_back(std::move(BackendString)); } SYCLTC.TranslateLinkerTargetArgs(TC->getTriple(), Args, BuildArgs); for (const auto &A : BuildArgs) { if (TC->getTriple().getSubArch() == llvm::Triple::NoSubArch) appendOption(LinkOptString, A); - else - // For AOT, append linker args to every backend string - for (auto &BOS : BackendOptStrings) - appendOption(BOS, A); + else { + // For AOT, append linker args to every backend string + for (auto &KV : BackendOptMap) + appendOption(KV.second, A); + } } - for (auto &BOS : BackendOptStrings) { - CmdArgs.push_back(Args.MakeArgString( + for (auto &KV : BackendOptMap) { + llvm::StringRef Device = KV.first; + llvm::SmallString<128> &Opts = KV.second; + + CmdArgs.push_back(Args.MakeArgString( "--device-compiler=" + Action::GetOffloadKindName(Action::OFK_SYCL) + ":" + - TC->getTripleString() + "=" + BOS)); + TC->getTripleString() + ":" + Device + ":""=" + Opts)); } if (!LinkOptString.empty()) { CmdArgs.push_back(Args.MakeArgString( diff --git a/clang/lib/Driver/ToolChains/SYCL.cpp b/clang/lib/Driver/ToolChains/SYCL.cpp index 62cd3abd4b27d..0141a190eca65 100644 --- a/clang/lib/Driver/ToolChains/SYCL.cpp +++ b/clang/lib/Driver/ToolChains/SYCL.cpp @@ -1083,6 +1083,18 @@ void SYCL::gen::BackendCompiler::ConstructJob(Compilation &C, C.addCommand(std::move(Cmd)); } +StringRef SYCL::gen::extractDeviceFromArgSimple(llvm::StringRef Arg) { + llvm::SmallVector Tokens; + Arg.split(Tokens, ' '); + + for (size_t I = 0; I + 1 < Tokens.size(); ++I) { + if (Tokens[I] == "-device") + return Tokens[I + 1]; + } + + return ""; +} + StringRef SYCL::gen::resolveGenDevice(StringRef DeviceName) { StringRef Device; Device = @@ -1555,24 +1567,23 @@ void SYCLToolChain::TranslateTargetOpt(const llvm::Triple &Triple, const llvm::Triple OptTargetTriple = getDriver().getSYCLDeviceTriple(A->getValue(), A); // Passing device args: -X= -opt=val. - StringRef GenDevice = SYCL::gen::resolveGenDevice(A->getValue()); + StringRef GenDevice; + StringRef TargetBackend = A->getValue(); + if((TargetBackend == "spir64_gen")) { + GenDevice = SYCL::gen::extractDeviceFromArgSimple(A->getValue(1)); + } else { + GenDevice = SYCL::gen::resolveGenDevice(TargetBackend); + } + bool IsGenTriple = Triple.isSPIR() && Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen; - // llvm::errs() << "[DEBUG] Args: " << A->getAsString(Args) - // << " OptTargetTriple: " << OptTargetTriple.str() - // << ", Triple: " << Triple.str() << ", Device: " << Device - // << ", GenDevice: " << GenDevice << "\n"; + llvm::errs() << "[DEBUG] Args: " << A->getAsString(Args) + << " OptTargetTriple: " << OptTargetTriple.str() + << ", Triple: " << Triple.str() << ", Device: " << Device + << ", GenDevice: " << GenDevice << "\n"; if (IsGenTriple) { - if (Device != GenDevice) { - if(!GenDevice.empty()) { - continue; - } - SmallString<64> DeviceOpt("-device "); - DeviceOpt += Device; - if (A->getAsString(Args).find(DeviceOpt.str()) == std::string::npos) - continue; - } - // if (Device != GenDevice && !Device.empty()) + if (Device != GenDevice) + continue; if (OptTargetTriple != Triple && GenDevice.empty()) // Triples do not match, but only skip when we know we are not // comparing against intel_gpu_* diff --git a/clang/lib/Driver/ToolChains/SYCL.h b/clang/lib/Driver/ToolChains/SYCL.h index c1403ae1d5036..74720ea1eb8d3 100644 --- a/clang/lib/Driver/ToolChains/SYCL.h +++ b/clang/lib/Driver/ToolChains/SYCL.h @@ -84,6 +84,7 @@ class LLVM_LIBRARY_VISIBILITY BackendCompiler : public Tool { const char *LinkingOutput) const override; }; +StringRef extractDeviceFromArgSimple(StringRef Arg); StringRef resolveGenDevice(StringRef DeviceName); SmallString<64> getGenDeviceMacro(StringRef DeviceName); StringRef getGenGRFFlag(StringRef GRFMode); From 6df1d44d4b77cfd59ec39f5f333b4384307eda23 Mon Sep 17 00:00:00 2001 From: "yixing.zhang" Date: Sun, 1 Feb 2026 19:07:07 +0100 Subject: [PATCH 04/14] remove the extra arch field, passing -device arch instead --- clang/lib/Driver/ToolChains/Clang.cpp | 20 +++++++++----------- 1 file changed, 9 insertions(+), 11 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 6f7eb0e6dfed3..9e9419d95794e 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -11616,26 +11616,27 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, if (!TC->getTriple().isSPIROrSPIRV()) continue; ArgStringList BuildArgs; - llvm::SmallMapVector, 4> BackendOptMap; + std::vector> BackendOptVec; SmallString<128> LinkOptString; // Process each -Xsycl-target-backend individually for (const Arg *A : Args.filtered(options::OPT_Xsycl_backend_EQ)) { StringRef Device; StringRef TargetBackend = A->getValue(); + SmallString<128> BackendArgs; if((TargetBackend == "spir64_gen")) { Device = SYCL::gen::extractDeviceFromArgSimple(A->getValue(1)); - llvm::errs() << "[DEBUG] Resolved device: " << Device << "\n"; } else { Device = SYCL::gen::resolveGenDevice(TargetBackend); + appendOption(BackendArgs, "-device " + Device.str()); } - SmallString<128> &BackendString = BackendOptMap[Device]; SYCLTC.TranslateBackendTargetArgs(TC->getTriple(), Args, BuildArgs, Device); for (const auto &BA : BuildArgs) { - appendOption(BackendString, BA); + appendOption(BackendArgs, BA); } + BackendOptVec.push_back(std::move(BackendArgs)); BuildArgs.clear(); } @@ -11645,19 +11646,16 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, appendOption(LinkOptString, A); else { // For AOT, append linker args to every backend string - for (auto &KV : BackendOptMap) - appendOption(KV.second, A); + for (auto &BackendArgs : BackendOptVec) + appendOption(BackendArgs, A); } } - for (auto &KV : BackendOptMap) { - llvm::StringRef Device = KV.first; - llvm::SmallString<128> &Opts = KV.second; - + for (auto &BackendArgs : BackendOptVec) { CmdArgs.push_back(Args.MakeArgString( "--device-compiler=" + Action::GetOffloadKindName(Action::OFK_SYCL) + ":" + - TC->getTripleString() + ":" + Device + ":""=" + Opts)); + TC->getTripleString() + "=" + BackendArgs)); } if (!LinkOptString.empty()) { CmdArgs.push_back(Args.MakeArgString( From 082b3f82f7e9d14090818d85aeca5894ad0eb090 Mon Sep 17 00:00:00 2001 From: "yixing.zhang" Date: Sun, 1 Feb 2026 22:50:38 +0100 Subject: [PATCH 05/14] code cleanup --- clang/lib/Driver/ToolChains/Clang.cpp | 2 -- clang/lib/Driver/ToolChains/SYCL.cpp | 4 ---- 2 files changed, 6 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 9e9419d95794e..ce3bd6af2385a 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -10450,8 +10450,6 @@ void OffloadPackager::ConstructJob(Compilation &C, const JobAction &JA, const ToolChain *HostTC = C.getSingleOffloadToolChain(); const toolchains::SYCLToolChain &SYCLTC = static_cast(*TC); - - SYCLTC.AddImpliedTargetArgs(TC->getTriple(), Args, BuildArgs, JA, *HostTC, Arch); SYCLTC.TranslateBackendTargetArgs(TC->getTriple(), Args, BuildArgs, OffloadAction->getOffloadingArch()); diff --git a/clang/lib/Driver/ToolChains/SYCL.cpp b/clang/lib/Driver/ToolChains/SYCL.cpp index 0141a190eca65..fe76cade606f4 100644 --- a/clang/lib/Driver/ToolChains/SYCL.cpp +++ b/clang/lib/Driver/ToolChains/SYCL.cpp @@ -1577,10 +1577,6 @@ void SYCLToolChain::TranslateTargetOpt(const llvm::Triple &Triple, bool IsGenTriple = Triple.isSPIR() && Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen; - llvm::errs() << "[DEBUG] Args: " << A->getAsString(Args) - << " OptTargetTriple: " << OptTargetTriple.str() - << ", Triple: " << Triple.str() << ", Device: " << Device - << ", GenDevice: " << GenDevice << "\n"; if (IsGenTriple) { if (Device != GenDevice) continue; From 27b765ce5474fe50506cf30de3a239abf505f7f1 Mon Sep 17 00:00:00 2001 From: "yixing.zhang" Date: Mon, 2 Feb 2026 04:58:31 +0100 Subject: [PATCH 06/14] correct options are now passing to ocloc for each target --- clang/lib/Driver/ToolChains/SYCL.cpp | 2 +- .../ClangLinkerWrapper.cpp | 26 ++++++++++++++----- 2 files changed, 21 insertions(+), 7 deletions(-) diff --git a/clang/lib/Driver/ToolChains/SYCL.cpp b/clang/lib/Driver/ToolChains/SYCL.cpp index fe76cade606f4..a610932fd4932 100644 --- a/clang/lib/Driver/ToolChains/SYCL.cpp +++ b/clang/lib/Driver/ToolChains/SYCL.cpp @@ -1578,7 +1578,7 @@ void SYCLToolChain::TranslateTargetOpt(const llvm::Triple &Triple, bool IsGenTriple = Triple.isSPIR() && Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen; if (IsGenTriple) { - if (Device != GenDevice) + if (Device != GenDevice && !Device.empty()) continue; if (OptTargetTriple != Triple && GenDevice.empty()) // Triples do not match, but only skip when we know we are not diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 52be1a00788c8..d24459aae84e9 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -754,10 +754,19 @@ runSYCLPostLinkTool(ArrayRef InputFiles, const ArgList &Args, const llvm::Triple Triple(Args.getLastArgValue(OPT_triple_EQ)); StringRef Arch = Args.getLastArgValue(OPT_arch_EQ); - if (Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen && !Arch.empty() && - !IsDevicePassedWithSyclTargetBackend && Arch != "*") - OutputPathWithArch = "intel_gpu_" + Arch.str() + "," + OutputPathWithArch; - else if (Triple.getSubArch() == llvm::Triple::SPIRSubArch_x86_64) + if (Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen && !Arch.empty() && Arch != "*") { + SmallVector ArchList; + Arch.split(ArchList, ','); + + std::string ArchPaths; + for (StringRef SingleArch : ArchList) { + if (!ArchPaths.empty()) + ArchPaths += ","; + ArchPaths += "intel_gpu_" + SingleArch.str(); + } + if (!ArchPaths.empty()) + OutputPathWithArch = ArchPaths + "," + OutputPathWithArch; + } else if (Triple.getSubArch() == llvm::Triple::SPIRSubArch_x86_64) OutputPathWithArch = "spir64_x86_64," + OutputPathWithArch; SmallVector CmdArgs; @@ -2188,8 +2197,13 @@ static void appendSYCLDeviceOptionsAtLinkTime(const DerivedArgList &LinkerArgs, std::string &CompileOptions, std::string &LinkOptions) { const StringRef TripleStr = LinkerArgs.getLastArgValue(OPT_triple_EQ); - auto processDeviceArgs = [&](unsigned OptID, std::string &Options) { + auto processDeviceArgs = [&](unsigned OptID, std::string &Options, StringRef TargetArch = StringRef()) { for (StringRef Arg : LinkerArgs.getAllArgValues(OptID)) { + if (!TargetArch.empty()) { + std::string DeviceArchPattern = "-device " + TargetArch.str(); + if (Arg.find(DeviceArchPattern) == StringRef::npos) + continue; + } size_t ColonPos = Arg.find(':'); if (ColonPos != StringRef::npos) { StringRef Kind = Arg.substr(0, ColonPos); @@ -2211,7 +2225,7 @@ static void appendSYCLDeviceOptionsAtLinkTime(const DerivedArgList &LinkerArgs, } }; - processDeviceArgs(OPT_device_compiler_args_EQ, CompileOptions); + processDeviceArgs(OPT_device_compiler_args_EQ, CompileOptions, LinkerArgs.getLastArgValue(OPT_arch_EQ)); processDeviceArgs(OPT_device_linker_args_EQ, LinkOptions); } From f8bcb2c9fd52bebdd2c41660e6ff31d2ef10adc4 Mon Sep 17 00:00:00 2001 From: "yixing.zhang" Date: Wed, 4 Feb 2026 16:56:49 +0100 Subject: [PATCH 07/14] update sycl-post-link to support multiple architecture --- llvm/tools/sycl-post-link/sycl-post-link.cpp | 54 +++++++++++++++----- 1 file changed, 40 insertions(+), 14 deletions(-) diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index c49c552c1d2db..345a43b7d6714 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -97,7 +97,7 @@ cl::opt DeviceLibDir{ cl::value_desc("dirname"), cl::cat(PostLinkCat)}; struct TargetFilenamePair { - std::string Target; + std::vector Targets; std::string Filename; }; @@ -105,10 +105,16 @@ struct TargetFilenamePairParser : public cl::basic_parser { using cl::basic_parser::basic_parser; bool parse(cl::Option &O, StringRef ArgName, StringRef &ArgValue, TargetFilenamePair &Val) const { - auto [Target, Filename] = ArgValue.split(","); - if (Filename == "") - std::swap(Target, Filename); - Val = {Target.str(), Filename.str()}; + SmallVector Parts; + ArgValue.split(Parts, ","); + if (Parts.size() == 1) { + Val.Targets = {}; + Val.Filename = Parts[0].str(); + return false; + } + Val.Filename = Parts.back().str(); + for (size_t i = 0; i + 1 < Parts.size(); ++i) + Val.Targets.push_back(Parts[i].str()); return false; } }; @@ -285,7 +291,10 @@ static void writeToFile(const StringRef Filename, const StringRef Content) { void saveModuleIR(Module &M, const StringRef Filename) { std::error_code EC; raw_fd_ostream Out{Filename, EC, sys::fs::OF_None}; + llvm::errs() << "[DEBUG] trying to open file: " << Filename << "\n"; checkError(EC, "error opening the file '" + Filename + "'"); + llvm::errs() << "[DEBUG] successfully open file: " << Filename << "\n"; + ModulePassManager MPM; ModuleAnalysisManager MAM; @@ -300,7 +309,7 @@ void saveModuleIR(Module &M, const StringRef Filename) { void saveModuleProperties(const module_split::ModuleDesc &MD, const GlobalBinImageProps &GlobProps, - const StringRef Filename, StringRef Target = "") { + const StringRef Filename, const std::vector &Targets) { PropSetRegTy PropSet; @@ -323,9 +332,9 @@ void saveModuleProperties(const module_split::ModuleDesc &MD, PropSet.remove(PropSetRegTy::SYCL_DEVICE_REQUIREMENTS, PropSetRegTy::PROPERTY_REQD_WORK_GROUP_SIZE); - if (!Target.empty()) - PropSet.add(PropSetRegTy::SYCL_DEVICE_REQUIREMENTS, "compile_target", - Target); + for (const auto &T : Targets) + PropSet.add(PropSetRegTy::SYCL_DEVICE_REQUIREMENTS, + "compile_target", T); std::error_code EC; raw_fd_ostream SCOut(Filename, EC); @@ -337,6 +346,7 @@ void saveModuleProperties(const module_split::ModuleDesc &MD, void saveModuleSymbolTable(const module_split::ModuleDesc &MD, const StringRef Filename) { std::string SymT = computeModuleSymbolTable(MD.getModule(), MD.entries()); + llvm::errs() << "[DEBUG] Saving module symbol table to file: " << Filename << "\n"; writeToFile(Filename, SymT); } @@ -384,21 +394,36 @@ void saveModule( } for (const auto &[Table, OutputFile] : zip_equal(OutTables, OutputFiles)) { - if (!isTargetCompatibleWithModule(OutputFile.Target, MD)) + bool Compatible = false; + if (OutputFile.Targets.empty()) { + Compatible = true; + } else { + for (const auto &T : OutputFile.Targets) { + if (isTargetCompatibleWithModule(T, MD)) { + Compatible = true; + break; + } + } + } + + if (!Compatible) continue; auto CopyTriple = BaseTriple; if (DoPropGen) { GlobalBinImageProps Props = {EmitKernelParamInfo, EmitProgramMetadata, EmitKernelNames, EmitExportedSymbols, EmitImportedSymbols, DeviceGlobals}; - StringRef Target = OutputFile.Target; std::string NewSuff = Suffix.str(); - if (!Target.empty()) - NewSuff = (Twine("_") + Target).str(); + if (!OutputFile.Targets.empty()) { + std::string Joined; + for (auto &T : OutputFile.Targets) + Joined += "_" + T; + NewSuff = Joined; + } CopyTriple.Prop = (OutputPrefix + NewSuff + "_" + Twine(I) + ".prop").str(); - saveModuleProperties(MD, Props, CopyTriple.Prop, Target); + saveModuleProperties(MD, Props, CopyTriple.Prop, OutputFile.Targets); } addTableRow(*Table, CopyTriple); } @@ -457,6 +482,7 @@ bool isTargetCompatibleWithModule(const std::string &Target, // (e.g. -o out.table compared to -o intel_gpu_pvc,out-pvc.table) // Target will be empty and we will not want to perform any filtering, so // we return true here. + llvm::errs() << "[DEBUG] target: '" << Target << "'\n"; if (Target.empty()) return true; From a9eb00097c19677f7de638aec405ed6960494b89 Mon Sep 17 00:00:00 2001 From: "yixing.zhang" Date: Wed, 4 Feb 2026 17:43:34 +0100 Subject: [PATCH 08/14] fix bug in sycl-post-link --- llvm/tools/sycl-post-link/sycl-post-link.cpp | 73 +++++++++----------- 1 file changed, 32 insertions(+), 41 deletions(-) diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 345a43b7d6714..8475a9a8ddca1 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -350,7 +350,7 @@ void saveModuleSymbolTable(const module_split::ModuleDesc &MD, writeToFile(Filename, SymT); } -bool isTargetCompatibleWithModule(const std::string &Target, +bool isTargetCompatibleWithModule(const std::vector &Targets, module_split::ModuleDesc &IrMD); void addTableRow(util::SimpleTable &Table, @@ -394,19 +394,7 @@ void saveModule( } for (const auto &[Table, OutputFile] : zip_equal(OutTables, OutputFiles)) { - bool Compatible = false; - if (OutputFile.Targets.empty()) { - Compatible = true; - } else { - for (const auto &T : OutputFile.Targets) { - if (isTargetCompatibleWithModule(T, MD)) { - Compatible = true; - break; - } - } - } - - if (!Compatible) + if (!isTargetCompatibleWithModule(OutputFile.Targets, MD)) continue; auto CopyTriple = BaseTriple; if (DoPropGen) { @@ -476,45 +464,48 @@ void addTableRow(util::SimpleTable &Table, // information comes from the device config file (DeviceConfigFile.td). // For example, the intel_gpu_tgllp target does not support fp64 - therefore, // a module using fp64 would *not* be compatible with intel_gpu_tgllp. -bool isTargetCompatibleWithModule(const std::string &Target, +bool isTargetCompatibleWithModule(const std::vector &Targets, module_split::ModuleDesc &IrMD) { // When the user does not specify a target, // (e.g. -o out.table compared to -o intel_gpu_pvc,out-pvc.table) // Target will be empty and we will not want to perform any filtering, so // we return true here. - llvm::errs() << "[DEBUG] target: '" << Target << "'\n"; - if (Target.empty()) + if (Targets.empty()) return true; + + for (const std::string &Target : Targets) { + llvm::errs() << "[DEBUG] Checking compatibility of target '" << Target + << "' with module '" << IrMD.Name << "'\n"; + // TODO: If a target not found in the device config file is passed, + // to sycl-post-link, then we should probably throw an error. However, + // since not all the information for all the targets is filled out + // right now, we return true, having the affect that unrecognized + // targets have no filtering applied to them. + if (!is_contained(DeviceConfigFile::TargetTable, Target)) + continue; - // TODO: If a target not found in the device config file is passed, - // to sycl-post-link, then we should probably throw an error. However, - // since not all the information for all the targets is filled out - // right now, we return true, having the affect that unrecognized - // targets have no filtering applied to them. - if (!is_contained(DeviceConfigFile::TargetTable, Target)) - return true; + const DeviceConfigFile::TargetInfo &TargetInfo = + DeviceConfigFile::TargetTable[Target]; + const SYCLDeviceRequirements &ModuleReqs = + IrMD.getOrComputeDeviceRequirements(); - const DeviceConfigFile::TargetInfo &TargetInfo = - DeviceConfigFile::TargetTable[Target]; - const SYCLDeviceRequirements &ModuleReqs = - IrMD.getOrComputeDeviceRequirements(); + // Check to see if all the requirements of the input module + // are compatbile with the target. + for (const auto &Aspect : ModuleReqs.Aspects) { + if (!is_contained(TargetInfo.aspects, Aspect.Name)) + return false; + } - // Check to see if all the requirements of the input module - // are compatbile with the target. - for (const auto &Aspect : ModuleReqs.Aspects) { - if (!is_contained(TargetInfo.aspects, Aspect.Name)) + // Check if module sub group size is compatible with the target. + // For ESIMD, the reqd_sub_group_size will be 1; this is not + // a supported by any backend (e.g. no backend can support a kernel + // with sycl::reqd_sub_group_size(1)), but for ESIMD, this is + // a special case. + if (!IrMD.isESIMD() && ModuleReqs.SubGroupSize.has_value() && + !is_contained(TargetInfo.subGroupSizes, *ModuleReqs.SubGroupSize)) return false; } - // Check if module sub group size is compatible with the target. - // For ESIMD, the reqd_sub_group_size will be 1; this is not - // a supported by any backend (e.g. no backend can support a kernel - // with sycl::reqd_sub_group_size(1)), but for ESIMD, this is - // a special case. - if (!IrMD.isESIMD() && ModuleReqs.SubGroupSize.has_value() && - !is_contained(TargetInfo.subGroupSizes, *ModuleReqs.SubGroupSize)) - return false; - return true; } From 30179dbe0499e7ca621ccd94c0d67be6bb0bc167 Mon Sep 17 00:00:00 2001 From: "yixing.zhang" Date: Wed, 4 Feb 2026 21:19:09 +0100 Subject: [PATCH 09/14] code clean up --- clang/lib/Driver/ToolChains/Clang.cpp | 20 ++++++------- clang/lib/Driver/ToolChains/SYCL.cpp | 25 +++++++---------- clang/lib/Driver/ToolChains/SYCL.h | 2 +- .../ClangLinkerWrapper.cpp | 15 ++++------ llvm/tools/sycl-post-link/sycl-post-link.cpp | 28 +++++++------------ 5 files changed, 34 insertions(+), 56 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index ce3bd6af2385a..da01299e6cfa6 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -10455,7 +10455,7 @@ void OffloadPackager::ConstructJob(Compilation &C, const JobAction &JA, SYCLTC.TranslateBackendTargetArgs(TC->getTriple(), Args, BuildArgs, OffloadAction->getOffloadingArch()); createArgString("compile-opts="); BuildArgs.clear(); - SYCLTC.TranslateLinkerTargetArgs(TC->getTriple(), Args, BuildArgs, OffloadAction->getOffloadingArch()); + SYCLTC.TranslateLinkerTargetArgs(TC->getTriple(), Args, BuildArgs); createArgString("link-opts="); } @@ -11619,21 +11619,17 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, // Process each -Xsycl-target-backend individually for (const Arg *A : Args.filtered(options::OPT_Xsycl_backend_EQ)) { - StringRef Device; - StringRef TargetBackend = A->getValue(); SmallString<128> BackendArgs; - if((TargetBackend == "spir64_gen")) { - Device = SYCL::gen::extractDeviceFromArgSimple(A->getValue(1)); - } else { - Device = SYCL::gen::resolveGenDevice(TargetBackend); + StringRef Device = SYCL::gen::resolveGenDevice(A->getValue()); + if(!Device.empty()) appendOption(BackendArgs, "-device " + Device.str()); - } + else + Device = SYCL::gen::extractDeviceFromArg(A->getValue(1)); SYCLTC.TranslateBackendTargetArgs(TC->getTriple(), Args, BuildArgs, Device); for (const auto &BA : BuildArgs) { appendOption(BackendArgs, BA); } - BackendOptVec.push_back(std::move(BackendArgs)); BuildArgs.clear(); } @@ -11643,13 +11639,13 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, if (TC->getTriple().getSubArch() == llvm::Triple::NoSubArch) appendOption(LinkOptString, A); else { - // For AOT, append linker args to every backend string - for (auto &BackendArgs : BackendOptVec) + // For AOT, combine the Backend and Linker strings into one. + for (SmallString<128> &BackendArgs : BackendOptVec) appendOption(BackendArgs, A); } } - for (auto &BackendArgs : BackendOptVec) { + for (SmallString<128> &BackendArgs : BackendOptVec) { CmdArgs.push_back(Args.MakeArgString( "--device-compiler=" + Action::GetOffloadKindName(Action::OFK_SYCL) + ":" + diff --git a/clang/lib/Driver/ToolChains/SYCL.cpp b/clang/lib/Driver/ToolChains/SYCL.cpp index a610932fd4932..2556ad1abab15 100644 --- a/clang/lib/Driver/ToolChains/SYCL.cpp +++ b/clang/lib/Driver/ToolChains/SYCL.cpp @@ -1083,15 +1083,13 @@ void SYCL::gen::BackendCompiler::ConstructJob(Compilation &C, C.addCommand(std::move(Cmd)); } -StringRef SYCL::gen::extractDeviceFromArgSimple(llvm::StringRef Arg) { - llvm::SmallVector Tokens; - Arg.split(Tokens, ' '); - - for (size_t I = 0; I + 1 < Tokens.size(); ++I) { - if (Tokens[I] == "-device") - return Tokens[I + 1]; +StringRef SYCL::gen::extractDeviceFromArg(llvm::StringRef Arg) { + llvm::SmallVector Arglist; + Arg.split(Arglist, ' '); + for (size_t i = 0; i + 1 < Arglist.size(); ++i) { + if (Arglist[i] == "-device") + return Arglist[i + 1]; } - return ""; } @@ -1566,14 +1564,11 @@ void SYCLToolChain::TranslateTargetOpt(const llvm::Triple &Triple, if (A->getOption().matches(Opt_EQ)) { const llvm::Triple OptTargetTriple = getDriver().getSYCLDeviceTriple(A->getValue(), A); + // Passing device args: -X= -opt=val. - StringRef GenDevice; - StringRef TargetBackend = A->getValue(); - if((TargetBackend == "spir64_gen")) { - GenDevice = SYCL::gen::extractDeviceFromArgSimple(A->getValue(1)); - } else { - GenDevice = SYCL::gen::resolveGenDevice(TargetBackend); - } + StringRef GenDevice = SYCL::gen::resolveGenDevice(A->getValue()); + if(GenDevice.empty()) + GenDevice = SYCL::gen::extractDeviceFromArg(A->getValue(1)); bool IsGenTriple = Triple.isSPIR() && Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen; diff --git a/clang/lib/Driver/ToolChains/SYCL.h b/clang/lib/Driver/ToolChains/SYCL.h index 74720ea1eb8d3..66e5f31f23725 100644 --- a/clang/lib/Driver/ToolChains/SYCL.h +++ b/clang/lib/Driver/ToolChains/SYCL.h @@ -84,7 +84,7 @@ class LLVM_LIBRARY_VISIBILITY BackendCompiler : public Tool { const char *LinkingOutput) const override; }; -StringRef extractDeviceFromArgSimple(StringRef Arg); +StringRef extractDeviceFromArg(StringRef Arg); StringRef resolveGenDevice(StringRef DeviceName); SmallString<64> getGenDeviceMacro(StringRef DeviceName); StringRef getGenGRFFlag(StringRef GRFMode); diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index d24459aae84e9..1332ee2c56213 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -755,17 +755,12 @@ runSYCLPostLinkTool(ArrayRef InputFiles, const ArgList &Args, StringRef Arch = Args.getLastArgValue(OPT_arch_EQ); if (Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen && !Arch.empty() && Arch != "*") { - SmallVector ArchList; + SmallVector ArchList; Arch.split(ArchList, ','); - - std::string ArchPaths; - for (StringRef SingleArch : ArchList) { - if (!ArchPaths.empty()) - ArchPaths += ","; - ArchPaths += "intel_gpu_" + SingleArch.str(); - } - if (!ArchPaths.empty()) - OutputPathWithArch = ArchPaths + "," + OutputPathWithArch; + std::string ArchString; + for (StringRef SingleArch : ArchList) + ArchString += "intel_gpu_" + SingleArch.str() + ","; + OutputPathWithArch = ArchString + OutputPathWithArch; } else if (Triple.getSubArch() == llvm::Triple::SPIRSubArch_x86_64) OutputPathWithArch = "spir64_x86_64," + OutputPathWithArch; diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 8475a9a8ddca1..585afbb76c710 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -105,16 +105,11 @@ struct TargetFilenamePairParser : public cl::basic_parser { using cl::basic_parser::basic_parser; bool parse(cl::Option &O, StringRef ArgName, StringRef &ArgValue, TargetFilenamePair &Val) const { - SmallVector Parts; - ArgValue.split(Parts, ","); - if (Parts.size() == 1) { - Val.Targets = {}; - Val.Filename = Parts[0].str(); - return false; - } - Val.Filename = Parts.back().str(); - for (size_t i = 0; i + 1 < Parts.size(); ++i) - Val.Targets.push_back(Parts[i].str()); + SmallVector ArgList; + ArgValue.split(ArgList, ","); + Val.Filename = ArgList.back().str(); + for (size_t i = 0; i + 1 < ArgList.size(); ++i) + Val.Targets.push_back(ArgList[i].str()); return false; } }; @@ -291,10 +286,7 @@ static void writeToFile(const StringRef Filename, const StringRef Content) { void saveModuleIR(Module &M, const StringRef Filename) { std::error_code EC; raw_fd_ostream Out{Filename, EC, sys::fs::OF_None}; - llvm::errs() << "[DEBUG] trying to open file: " << Filename << "\n"; checkError(EC, "error opening the file '" + Filename + "'"); - llvm::errs() << "[DEBUG] successfully open file: " << Filename << "\n"; - ModulePassManager MPM; ModuleAnalysisManager MAM; @@ -346,7 +338,6 @@ void saveModuleProperties(const module_split::ModuleDesc &MD, void saveModuleSymbolTable(const module_split::ModuleDesc &MD, const StringRef Filename) { std::string SymT = computeModuleSymbolTable(MD.getModule(), MD.entries()); - llvm::errs() << "[DEBUG] Saving module symbol table to file: " << Filename << "\n"; writeToFile(Filename, SymT); } @@ -401,17 +392,18 @@ void saveModule( GlobalBinImageProps Props = {EmitKernelParamInfo, EmitProgramMetadata, EmitKernelNames, EmitExportedSymbols, EmitImportedSymbols, DeviceGlobals}; + std::vector Targets = OutputFile.Targets; std::string NewSuff = Suffix.str(); - if (!OutputFile.Targets.empty()) { + if (!Targets.empty()) { std::string Joined; - for (auto &T : OutputFile.Targets) + for (auto &T : Targets) Joined += "_" + T; NewSuff = Joined; } CopyTriple.Prop = (OutputPrefix + NewSuff + "_" + Twine(I) + ".prop").str(); - saveModuleProperties(MD, Props, CopyTriple.Prop, OutputFile.Targets); + saveModuleProperties(MD, Props, CopyTriple.Prop, Targets); } addTableRow(*Table, CopyTriple); } @@ -473,7 +465,7 @@ bool isTargetCompatibleWithModule(const std::vector &Targets, if (Targets.empty()) return true; - for (const std::string &Target : Targets) { + for (const std::string Target : Targets) { llvm::errs() << "[DEBUG] Checking compatibility of target '" << Target << "' with module '" << IrMD.Name << "'\n"; // TODO: If a target not found in the device config file is passed, From d12c9af7ccd91a7b6fc6854c42b031baa9515ca3 Mon Sep 17 00:00:00 2001 From: "yixing.zhang" Date: Wed, 4 Feb 2026 22:14:49 +0100 Subject: [PATCH 10/14] add comment --- clang/lib/Driver/ToolChains/Clang.cpp | 13 +++++++++---- clang/lib/Driver/ToolChains/SYCL.cpp | 4 ++-- .../clang-linker-wrapper/ClangLinkerWrapper.cpp | 6 ++++-- 3 files changed, 15 insertions(+), 8 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index da01299e6cfa6..c213bdaf70159 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -11617,14 +11617,19 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, std::vector> BackendOptVec; SmallString<128> LinkOptString; - // Process each -Xsycl-target-backend individually + // Construct backend options for each target passed via -Xsycl-target-backend + // in the form: "-device " for (const Arg *A : Args.filtered(options::OPT_Xsycl_backend_EQ)) { SmallString<128> BackendArgs; StringRef Device = SYCL::gen::resolveGenDevice(A->getValue()); - if(!Device.empty()) - appendOption(BackendArgs, "-device " + Device.str()); - else + if(Device.empty()) + // If target is spir64_gen, the device name needs to be extracted + // from the arguments. Device = SYCL::gen::extractDeviceFromArg(A->getValue(1)); + else + // If target is intel_gpu_*, "-device " + // is appended to BackendArgs. + appendOption(BackendArgs, "-device " + Device.str()); SYCLTC.TranslateBackendTargetArgs(TC->getTriple(), Args, BuildArgs, Device); for (const auto &BA : BuildArgs) { diff --git a/clang/lib/Driver/ToolChains/SYCL.cpp b/clang/lib/Driver/ToolChains/SYCL.cpp index 2556ad1abab15..6659e78113c36 100644 --- a/clang/lib/Driver/ToolChains/SYCL.cpp +++ b/clang/lib/Driver/ToolChains/SYCL.cpp @@ -1083,6 +1083,8 @@ void SYCL::gen::BackendCompiler::ConstructJob(Compilation &C, C.addCommand(std::move(Cmd)); } +// Extracts the device specified after "-device" from the backend +// argument string provided via -Xsycl-target-backend. StringRef SYCL::gen::extractDeviceFromArg(llvm::StringRef Arg) { llvm::SmallVector Arglist; Arg.split(Arglist, ' '); @@ -1564,12 +1566,10 @@ void SYCLToolChain::TranslateTargetOpt(const llvm::Triple &Triple, if (A->getOption().matches(Opt_EQ)) { const llvm::Triple OptTargetTriple = getDriver().getSYCLDeviceTriple(A->getValue(), A); - // Passing device args: -X= -opt=val. StringRef GenDevice = SYCL::gen::resolveGenDevice(A->getValue()); if(GenDevice.empty()) GenDevice = SYCL::gen::extractDeviceFromArg(A->getValue(1)); - bool IsGenTriple = Triple.isSPIR() && Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen; if (IsGenTriple) { diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 1332ee2c56213..e95c075f6247a 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -753,7 +753,9 @@ runSYCLPostLinkTool(ArrayRef InputFiles, const ArgList &Args, // when Intel GPU targets are passed in -fsycl-targets. const llvm::Triple Triple(Args.getLastArgValue(OPT_triple_EQ)); StringRef Arch = Args.getLastArgValue(OPT_arch_EQ); - + + // Prefix the output path with the target architecture(s), + // e.g. intel_gpu_dg2,intel_gpu_pvc for arch = "dg2,pvc". if (Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen && !Arch.empty() && Arch != "*") { SmallVector ArchList; Arch.split(ArchList, ','); @@ -2184,7 +2186,7 @@ extractSYCLCompileLinkOptions(ArrayRef OffloadFiles) { } // Append SYCL device compiler and linker options specified at link time, -// filtering by target triple and offload kind. +// filtering by target triple, offload kind, and device architecture. // TODO: Consider how to refactor this function to merge it with getLinkerArgs() // and determine if it's possible to use OPT_compiler_arg_EQ and // OPT_linker_arg_EQ to handle device compiler/linker options From 7b41c45adb0c969f4670c7e19f41d6388091ce2d Mon Sep 17 00:00:00 2001 From: "yixing.zhang" Date: Wed, 4 Feb 2026 23:27:42 +0100 Subject: [PATCH 11/14] add the test for the change --- clang/test/Driver/sycl-offload-new-driver.cpp | 13 +++++++++++-- 1 file changed, 11 insertions(+), 2 deletions(-) diff --git a/clang/test/Driver/sycl-offload-new-driver.cpp b/clang/test/Driver/sycl-offload-new-driver.cpp index 0e8b394b1aa3b..f0b681922eceb 100644 --- a/clang/test/Driver/sycl-offload-new-driver.cpp +++ b/clang/test/Driver/sycl-offload-new-driver.cpp @@ -145,11 +145,11 @@ // RUN: -Xsycl-target-backend -backend-opt -### %s 2>&1 \ // RUN: | FileCheck -check-prefix WRAPPER_OPTIONS_BACKEND %s // WRAPPER_OPTIONS_BACKEND: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-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{{.*}} "--device-linker=sycl:spir64-unknown-unknown=-link-opt" +-// WRAPPER_OPTIONS_LINK: clang-linker-wrapper{{.*}} "--device-linker=sycl:spir64-unknown-unknown=-link-opt" /// Test option passing behavior for clang-offload-wrapper options for AOT. // RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver \ @@ -162,6 +162,15 @@ // WRAPPER_OPTIONS_BACKEND_AOT-SAME: "--device-compiler=sycl:spir64_gen-unknown-unknown=-backend-gen-opt" // WRAPPER_OPTIONS_BACKEND_AOT-SAME: "--device-compiler=sycl:spir64_x86_64-unknown-unknown=-backend-cpu-opt" +// Check that AOT backend compiler options passed via +// -Xsycl-target-backend=,