From 34014acb20dd5258b39715ae89a672b223fbd2d2 Mon Sep 17 00:00:00 2001 From: Nick Sarnie Date: Thu, 1 Oct 2026 11:59:01 -0700 Subject: [PATCH 01/19] [SYCL] Fix linking of fat static archives Signed-off-by: Nick Sarnie Co-Authored-By: Claude Sonnet 5.5 --- clang/include/clang/Basic/LangOptions.def | 1 + clang/include/clang/Options/Options.td | 1 + clang/lib/CodeGen/CodeGenModule.cpp | 14 +- clang/lib/CodeGen/CodeGenModule.h | 11 ++ clang/lib/CodeGen/CodeGenSYCL.cpp | 19 +++ clang/lib/Driver/ToolChains/Clang.cpp | 15 +- clang/test/CodeGenSYCL/registerlib-host.cpp | 26 ++++ clang/test/Driver/sycl-offload-new-driver.cpp | 6 +- .../linker-wrapper-sycl-registerlib.cpp | 45 ++++++ .../ClangLinkerWrapper.cpp | 129 ++++++++++++++++++ .../Inputs/static-lib-registerlib-bmg.cpp | 21 +++ .../Inputs/static-lib-registerlib-dg2.cpp | 21 +++ .../Inputs/static-lib-registerlib-jit.cpp | 21 +++ .../Inputs/static-lib-registerlib.hpp | 18 +++ ...-multiple-objects-multiple-targets-win.cpp | 32 +++++ ...-lib-multiple-objects-multiple-targets.cpp | 100 ++++++++++++++ 16 files changed, 468 insertions(+), 12 deletions(-) create mode 100644 clang/test/CodeGenSYCL/registerlib-host.cpp create mode 100644 clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl-registerlib.cpp create mode 100644 sycl/test-e2e/NewOffloadDriver/Inputs/static-lib-registerlib-bmg.cpp create mode 100644 sycl/test-e2e/NewOffloadDriver/Inputs/static-lib-registerlib-dg2.cpp create mode 100644 sycl/test-e2e/NewOffloadDriver/Inputs/static-lib-registerlib-jit.cpp create mode 100644 sycl/test-e2e/NewOffloadDriver/Inputs/static-lib-registerlib.hpp create mode 100644 sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets-win.cpp create mode 100644 sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets.cpp diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def index 5294a1832ee25..8f69b51a03e26 100644 --- a/clang/include/clang/Basic/LangOptions.def +++ b/clang/include/clang/Basic/LangOptions.def @@ -264,6 +264,7 @@ LANGOPT(GPUAllowDeviceInit, 1, 0, NotCompatible, "allowing device side global in LANGOPT(GPUMaxThreadsPerBlock, 32, 1024, NotCompatible, "default max threads per block for kernel launch bounds for HIP") LANGOPT(GPUDeferDiag, 1, 0, NotCompatible, "defer host/device related diagnostic messages for CUDA/HIP") LANGOPT(GPUExcludeWrongSideOverloads, 1, 0, NotCompatible, "always exclude wrong side overloads in overloading resolution for CUDA/HIP") +LANGOPT(OffloadingNewDriver, 1, 0, NotCompatible, "use the new driver for generating offloading code.") LANGOPT(OffloadViaLLVM, 1, 0, NotCompatible, "target LLVM/Offload as portable offloading runtime.") LANGOPT(CUDANVCCABI, 1, 0, NotCompatible, "emit the NVCC-compatible host registration ABI for CUDA") LANGOPT(SPIRVDefaultAddrSpaceIsGeneric, 1, 0, NotCompatible, "use generic address space as default address space for SPIR-V targets") diff --git a/clang/include/clang/Options/Options.td b/clang/include/clang/Options/Options.td index cc183f7258661..235a2918e2233 100644 --- a/clang/include/clang/Options/Options.td +++ b/clang/include/clang/Options/Options.td @@ -1385,6 +1385,7 @@ def no_offload_arch_EQ def offload_new_driver : Flag<["--"], "offload-new-driver">, Visibility<[ClangOption, CC1Option]>, Group, + MarshallingInfoFlag>, HelpText<"Deprecated; the new driver is always used for offloading compilation.">; def no_offload_new_driver : Flag<["--"], "no-offload-new-driver">, Visibility<[ClangOption, CC1Option]>, Group, diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 6275281061750..8393c5c51e025 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1285,11 +1285,15 @@ void CodeGenModule::Release() { if (llvm::Function *CudaCtorFunction = CUDARuntime->finalizeModule()) AddGlobalCtor(CudaCtorFunction); } - if (LangOpts.SYCLIsHost && !CodeGenOpts.OffloadBinaryToEmbedFile.empty()) { - if (llvm::Function *SYCLCtorFunction = embedSYCLDeviceBinary()) - // A static initializer may launch a kernel, so the device binary has to - // be registered before any of them run, hence a priority. - AddGlobalCtor(SYCLCtorFunction, /*Priority=*/101); + if (LangOpts.SYCLIsHost) { + if (!CodeGenOpts.OffloadBinaryToEmbedFile.empty()) { + if (llvm::Function *SYCLCtorFunction = embedSYCLDeviceBinary()) + // A static initializer may launch a kernel, so the device binary has to + // be registered before any of them run, hence a priority. + AddGlobalCtor(SYCLCtorFunction, /*Priority=*/101); + } else if (LangOpts.OffloadingNewDriver) { + AddGlobalCtor(createSYCLRegisterLibFunc()); + } } if (OpenMPRuntime) { OpenMPRuntime->createOffloadEntriesAndInfoMetadata(); diff --git a/clang/lib/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h index f29a6baadcc99..a48779081ed3c 100644 --- a/clang/lib/CodeGen/CodeGenModule.h +++ b/clang/lib/CodeGen/CodeGenModule.h @@ -2232,6 +2232,17 @@ class CodeGenModule : public CodeGenTypeCache { /// if the binary could not be read. llvm::Function *embedSYCLDeviceBinary(); + /// Emit a per-module __sycl_registerlib_ reference that + /// ties the host and device sides of a translation unit together. This is + /// used to force linking of static libraries: we emit a global constructor + /// that references (calls) the symbol, creating an undefined reference that + /// forces the corresponding device static library member (which defines the + /// matching symbol) to be pulled in at link time. The device-side definition + /// is emitted by clang-linker-wrapper, so this is new offload model only. The + /// hash of the module identifier is what ties the two sides together. + /// \return the global constructor that calls registerlib function. + llvm::Function *createSYCLRegisterLibFunc(); + /// Determine whether the definition must be emitted; if this returns \c /// false, the definition can be emitted lazily if it's used. bool MustBeEmitted(const ValueDecl *D); diff --git a/clang/lib/CodeGen/CodeGenSYCL.cpp b/clang/lib/CodeGen/CodeGenSYCL.cpp index 1cce5ee27064d..af043aab31f38 100644 --- a/clang/lib/CodeGen/CodeGenSYCL.cpp +++ b/clang/lib/CodeGen/CodeGenSYCL.cpp @@ -123,3 +123,22 @@ llvm::Function *CodeGenModule::embedSYCLDeviceBinary() { } return RegistrationFunc; } + +llvm::Function *CodeGenModule::createSYCLRegisterLibFunc() { + std::string RegisterFuncName = + "__sycl_registerlib_" + LangOpts.SYCLUniquePrefix; + llvm::FunctionType *RegisterFTy = + llvm::FunctionType::get(VoidTy, /*isVarArg=*/false); + // Declare the symbol and emit a global constructor that calls it. + llvm::FunctionCallee RegisterFunc = + getModule().getOrInsertFunction(RegisterFuncName, RegisterFTy); + llvm::Function *SYCLRegisterLibCtor = CreateGlobalInitOrCleanUpFunction( + RegisterFTy, "__sycl_registerlib_ctor", + getTypes().arrangeNullaryFunction(), SourceLocation()); + llvm::BasicBlock *Entry = + llvm::BasicBlock::Create(getLLVMContext(), "entry", SYCLRegisterLibCtor); + llvm::IRBuilder<> Builder(Entry); + Builder.CreateCall(RegisterFunc); + Builder.CreateRetVoid(); + return SYCLRegisterLibCtor; +} diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 1b72069f7b4d2..33c9b04ba133b 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -8038,11 +8038,18 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, Args.addOptOutFlag(CmdArgs, options::OPT_fopenmp_extensions, options::OPT_fno_openmp_extensions); } - // Forward '-foffload-via-llvm' to code generation to target the LLVM/Offload - // runtime. + // Forward the offload runtime change to code generation, liboffload implies + // new driver. Otherwise, check if we should forward the new driver to change + // offloading code generation. if (Args.hasFlag(options::OPT_foffload_via_llvm, - options::OPT_fno_offload_via_llvm, false)) - CmdArgs.push_back("-foffload-via-llvm"); + options::OPT_fno_offload_via_llvm, false)) { + CmdArgs.append({"--offload-new-driver", "-foffload-via-llvm"}); + } else if (Args.hasFlag(options::OPT_offload_new_driver, + options::OPT_no_offload_new_driver, + (C.getActiveOffloadKinds() != Action::OFK_None && + C.getActiveOffloadKinds() != Action::OFK_SYCL))) { + CmdArgs.push_back("--offload-new-driver"); + } const XRayArgs &XRay = TC.getXRayArgs(Args); XRay.addArgs(TC, Args, CmdArgs, InputType); diff --git a/clang/test/CodeGenSYCL/registerlib-host.cpp b/clang/test/CodeGenSYCL/registerlib-host.cpp new file mode 100644 index 0000000000000..07cb84139aef4 --- /dev/null +++ b/clang/test/CodeGenSYCL/registerlib-host.cpp @@ -0,0 +1,26 @@ +// Check the per-translation-unit __sycl_registerlib_ reference that SYCL +// host code generation emits under the new offload driver. The id is the +// translation unit's -fsycl-unique-prefix, which the driver generates. + +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsycl-is-host --offload-new-driver -fsycl-unique-prefix=uid42 -emit-llvm %s -o - | FileCheck %s + +// Nothing is emitted without the new offload driver. +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsycl-is-host -fsycl-unique-prefix=uid42 -emit-llvm %s -o - | FileCheck %s --check-prefix=OLD + +// Nor when the module embeds its device binary (-fno-sycl-rdc), as it then +// registers the binary itself. +// RUN: echo -n 'FAKE_SYCL_DEVICE_IMAGE' > %t.bin +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsycl-is-host --offload-new-driver -foffload-include-binary %t.bin -fsycl-unique-prefix=uid42 -emit-llvm %s -o - | FileCheck %s --check-prefix=EMBED --implicit-check-not=__sycl_registerlib_ + +// CHECK: @llvm.global_ctors = appending global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 65535, ptr @__sycl_registerlib_ctor, ptr null }] +// CHECK: declare void @__sycl_registerlib_uid42() +// CHECK: define internal void @__sycl_registerlib_ctor() +// CHECK-NEXT: entry: +// CHECK-NEXT: call void @__sycl_registerlib_uid42() +// CHECK-NEXT: ret void + +// OLD-NOT: __sycl_registerlib_ + +// EMBED: @.sycl_offloading.binary + +void foo() {} diff --git a/clang/test/Driver/sycl-offload-new-driver.cpp b/clang/test/Driver/sycl-offload-new-driver.cpp index 93fc6f65d7a63..e520ca028e436 100644 --- a/clang/test/Driver/sycl-offload-new-driver.cpp +++ b/clang/test/Driver/sycl-offload-new-driver.cpp @@ -24,9 +24,9 @@ /// Check the toolflow for SYCL compilation using new offload model // RUN: %clangxx -### --target=x86_64-unknown-linux-gnu -fsycl -fsycl-targets=spir64 --offload-new-driver --sysroot=%S/Inputs/SYCL %s 2>&1 | FileCheck -check-prefix=CHK-FLOW %s -// CHK-FLOW: clang{{.*}} "-cc1" "-triple" "spir64-unknown-unknown" "-aux-triple" "x86_64-unknown-linux-gnu" "-fsycl-is-device" {{.*}} "-fsycl-int-header=[[HEADER:.*]].h" "-fsycl-int-footer=[[FOOTER:.*]].h" {{.*}} "-o" "[[CC1DEVOUT:.*]]" "-x" "c++" "[[INPUT:.*]]" +// CHK-FLOW: clang{{.*}} "-cc1" "-triple" "spir64-unknown-unknown" "-aux-triple" "x86_64-unknown-linux-gnu" "-fsycl-is-device" {{.*}} "-fsycl-int-header=[[HEADER:.*]].h" "-fsycl-int-footer=[[FOOTER:.*]].h" {{.*}} "--offload-new-driver" {{.*}} "-o" "[[CC1DEVOUT:.*]]" "-x" "c++" "[[INPUT:.*]]" // CHK-FLOW-NEXT: llvm-offload-binary{{.*}} "-o" "[[PACKOUT:.*]]" "--image=file=[[CC1DEVOUT]],triple=spir64-unknown-unknown,arch=generic,kind=sycl{{.*}}" -// CHK-FLOW-NEXT: clang{{.*}} "-cc1" "-triple" "x86_64-unknown-linux-gnu"{{.*}} "-fsycl-is-host"{{.*}} "-include-internal-header" "[[HEADER]].h" "-dependency-filter" "[[HEADER]].h" {{.*}} "-include-internal-footer" "[[FOOTER]].h" "-dependency-filter" "[[FOOTER]].h"{{.*}} "-fembed-offload-object=[[PACKOUT]]" {{.*}} "-o" "[[CC1FINALOUT:.*]]" "-x" "c++" "[[INPUT]]" +// CHK-FLOW-NEXT: clang{{.*}} "-cc1" "-triple" "x86_64-unknown-linux-gnu"{{.*}} "-fsycl-is-host"{{.*}} "-include-internal-header" "[[HEADER]].h" "-dependency-filter" "[[HEADER]].h" {{.*}} "-include-internal-footer" "[[FOOTER]].h" "-dependency-filter" "[[FOOTER]].h"{{.*}} "--offload-new-driver" {{.*}} "-fembed-offload-object=[[PACKOUT]]" {{.*}} "-o" "[[CC1FINALOUT:.*]]" "-x" "c++" "[[INPUT]]" // CHK-FLOW-NEXT: clang-linker-wrapper{{.*}} "--host-triple=x86_64-unknown-linux-gnu"{{.*}} "--linker-path={{.*}}/ld" {{.*}} "[[CC1FINALOUT]]" /// Verify phases used to generate SPIR-V instead of LLVM-IR @@ -93,7 +93,7 @@ // RUN: | FileCheck -check-prefix=CHK_ARCH \ // RUN: -DTRIPLE=nvptx64-nvidia-cuda -DARCH=sm_75 %s // CHK_ARCH: clang{{.*}} "-triple" "[[TRIPLE]]" -// CHK_ARCH-SAME: "-fsycl-is-device" {{.*}} "-o" "[[CC1DEVOUT:.+\.bc]]" +// CHK_ARCH-SAME: "-fsycl-is-device" {{.*}} "--offload-new-driver"{{.*}} "-o" "[[CC1DEVOUT:.+\.bc]]" // CHK_ARCH-NEXT: llvm-offload-binary{{.*}} "--image=file=[[CC1DEVOUT]],triple=[[TRIPLE]],arch=[[ARCH]]{{.*}},kind=sycl{{.*}}" // Verify offload-packager option values diff --git a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl-registerlib.cpp b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl-registerlib.cpp new file mode 100644 index 0000000000000..a2ce083b4aa74 --- /dev/null +++ b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl-registerlib.cpp @@ -0,0 +1,45 @@ +// REQUIRES: system-linux, x86-registered-target, spirv-to-ir-wrapper, sycl-post-link + +// Check that the SYCL wrapper object for a target defines the +// __sycl_registerlib_ symbols referenced by the host objects whose device +// code is for that target, including archive members found through -l. + +// RUN: %clang_cc1 -fsycl-is-device -disable-llvm-passes -triple=spir64-unknown-unknown %s -emit-llvm-bc -o %t.bc +// RUN: llvm-offload-binary -o %t.fat --image=file=%t.bc,kind=sycl,triple=spir64-unknown-unknown +// RUN: llvm-offload-binary -o %t.gen.fat --image=file=%t.bc,kind=sycl,triple=spir64_gen-unknown-unknown,arch=pvc +// RUN: %clang_cc1 -DMAIN %s -triple=x86_64-unknown-linux-gnu -fsycl-is-host --offload-new-driver -fsycl-unique-prefix=main -emit-obj -fembed-offload-object=%t.fat -o %t.main.o +// RUN: %clang_cc1 -DFN=lib_fn %s -triple=x86_64-unknown-linux-gnu -fsycl-is-host --offload-new-driver -fsycl-unique-prefix=lib -emit-obj -fembed-offload-object=%t.fat -o %t.lib.o +// RUN: %clang_cc1 -DFN=gen_fn %s -triple=x86_64-unknown-linux-gnu -fsycl-is-host --offload-new-driver -fsycl-unique-prefix=gen -emit-obj -fembed-offload-object=%t.gen.fat -o %t.gen.o +// RUN: rm -rf %t.dir && mkdir -p %t.dir +// RUN: llvm-ar rc %t.dir/libsyclreg.a %t.lib.o %t.gen.o +// RUN: touch %t.devicelib.bc +// RUN: clang-linker-wrapper --print-wrapped-module --host-triple=x86_64-unknown-linux-gnu \ +// RUN: --bitcode-library=spir64-unknown-unknown=%t.devicelib.bc \ +// RUN: -sycl-post-link-options=-split=auto -sycl-post-link-options=-symbols \ +// RUN: -sycl-post-link-options=-properties --linker-path=/usr/bin/ld \ +// RUN: %t.main.o -L%t.dir -lsyclreg -o %t.out 2>&1 \ +// RUN: | FileCheck %s --implicit-check-not=__sycl_registerlib_gen + +// No spir64_gen device code is linked, so nothing defines that reference. +// CHECK: define weak void @__sycl_registerlib_main() +// CHECK: define weak void @__sycl_registerlib_lib() + +template +__attribute__((sycl_kernel)) void kernel(const Func &func) { + func(); +} + +#ifdef MAIN +extern "C" { +void __sycl_register_lib(void *) {} +void __sycl_unregister_lib(void *) {} +} + +void lib_fn(); +int main() { + kernel([]() {}); + lib_fn(); +} +#else +void FN() { kernel([]() {}); } +#endif diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index a01947931e548..c83b341cb9806 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -18,8 +18,10 @@ #include "clang/Basic/TargetID.h" #include "clang/Basic/Version.h" #include "llvm/ADT/BitVector.h" +#include "llvm/ADT/DenseMap.h" #include "llvm/ADT/MapVector.h" #include "llvm/ADT/STLFunctionalExtras.h" +#include "llvm/ADT/SmallSet.h" #include "llvm/ADT/StringSet.h" #include "llvm/BinaryFormat/Magic.h" #include "llvm/Bitcode/BitcodeWriter.h" @@ -30,10 +32,12 @@ #include "llvm/Frontend/Offloading/SYCLOffloadWrapper.h" #include "llvm/Frontend/Offloading/Utility.h" #include "llvm/IR/DiagnosticPrinter.h" +#include "llvm/IR/IRBuilder.h" #include "llvm/IR/Module.h" #include "llvm/IRReader/IRReader.h" #include "llvm/LTO/LTO.h" #include "llvm/MC/TargetRegistry.h" +#include "llvm/Object/Archive.h" #include "llvm/Object/Binary.h" #include "llvm/Object/IRObjectFile.h" #include "llvm/Object/ObjectFile.h" @@ -163,6 +167,10 @@ static bool LinkingSYCLBINFiles = false; static SmallString<128> OffloadImageDumpDir; +using TripleAndArch = std::pair; +using FuncNameSet = SmallSet; +static DenseMap SYCLRegisterLibFuncMap; + /// Whether or not to look through symlinks when resolving binaries. static bool CanonicalPrefixes = true; @@ -184,6 +192,14 @@ template <> struct DenseMapInfo { return LHS == RHS; } }; +template <> struct DenseMapInfo { + static unsigned getHashValue(const std::string &S) { + return hash_value(StringRef(S)); + } + static bool isEqual(const std::string &L, const std::string &R) { + return L == R; + } +}; } // namespace llvm namespace { @@ -1112,6 +1128,38 @@ runAOTCompileIntelCPU(StringRef InputFile, const ArgList &Args, return *TempFileOrErr; } +/// Define __sycl_registerlib_ functions defined by the +/// host compile to force all device code images to be linked in, which +/// may not happen for static archives with -fsycl-link. +/// \p Module The module. +/// \p Args encompasses all arguments required for linking and wrapping device +/// code. +static void defineRegisterLibFuncs(Module &M, const ArgList &Args) { + const llvm::Triple Triple(Args.getLastArgValue(OPT_triple_EQ)); + const std::string Arch(Args.getLastArgValue(OPT_arch_EQ)); + // In the case we have a fat object input with multiple device images of + // differing triples, and we are not doing -fsycl-link (so multiple linked + // device images are allowed), the host object defined a single registerlib + // symbol, so we need to make sure to only define the function once, not once + // per device image from that fat object. We need to keep track of the triple + // through, because in the -fsycl-link case, we may not use one of those + // device images if it doesn't match the -fsycl-link triple. + static StringSet TotalDefinedFuncs; + for (auto Name : SYCLRegisterLibFuncMap[std::make_pair(Triple.str(), Arch)]) { + if (TotalDefinedFuncs.find(Name) != TotalDefinedFuncs.end()) + continue; + llvm::FunctionType *FTy = llvm::FunctionType::get( + llvm::Type::getVoidTy(M.getContext()), /*isVarArg=*/false); + auto *Fn = llvm::Function::Create(FTy, llvm::GlobalValue::WeakAnyLinkage, + Name, &M); + llvm::BasicBlock *Entry = + llvm::BasicBlock::Create(M.getContext(), "entry", Fn); + llvm::IRBuilder<> Builder(Entry); + Builder.CreateRetVoid(); + TotalDefinedFuncs.insert(Name); + } +} + /// Run AOT compilation for Intel GPU /// Calls ocloc tool to generate device code for Intel GPU backend. /// \p InputFile is the input SPIR-V file. @@ -1294,6 +1342,8 @@ wrapSYCLBinariesFromFile(ArrayRef SplitModules, M.setTargetTriple(Triple( Args.getLastArgValue(OPT_host_triple_EQ, sys::getDefaultTargetTriple()))); + defineRegisterLibFuncs(M, Args); + if (Error E = offloading::wrapSYCLBinaries(M, Images, offloading::SYCLWrappingOptions())) return E; @@ -3303,6 +3353,83 @@ Error emitExtractCommands( return Error::success(); } +/// See comments in collectSYCLRegisterLibSymbols. +static Error collectSYCLRegisterLibSymbolsFromObject(MemoryBufferRef Buffer) { + static constexpr StringRef SYCLRegisterLibPrefix = "__sycl_registerlib_"; + Expected> ObjOrErr = + ObjectFile::createObjectFile(Buffer); + if (!ObjOrErr) { + // Not something we can read symbols from; ignore it. + consumeError(ObjOrErr.takeError()); + return Error::success(); + } + SmallSet FuncNames; + for (const SymbolRef &Sym : (*ObjOrErr)->symbols()) { + Expected FlagsOrErr = Sym.getFlags(); + if (!FlagsOrErr) + return FlagsOrErr.takeError(); + Expected NameOrErr = Sym.getName(); + if (!NameOrErr) + return NameOrErr.takeError(); + if (!NameOrErr->starts_with(SYCLRegisterLibPrefix)) + continue; + if (*FlagsOrErr & SymbolRef::SF_Undefined) + FuncNames.insert(NameOrErr->str()); + } + if (!FuncNames.empty()) { + SmallVector Binaries; + if (Error Err = extractOffloadBinaries(Buffer, Binaries)) + return std::move(Err); + for (const auto &Binary : Binaries) { + std::string Triple = Binary.getBinary()->getTriple().str(); + std::string Arch = Binary.getBinary()->getArch().str(); + if (Arch == "generic") + Arch = ""; + SYCLRegisterLibFuncMap[{Triple, Arch}].insert(FuncNames.begin(), + FuncNames.end()); + } + } + return Error::success(); +} + +/// Scan the host input in \p Buffer for +/// undefined symbols named __sycl_registerlib_ and record them in the +/// global SYCLRegisterLibFuncMap. Each such reference is satisfied by defining +/// the symbol in the generated SYCL device-image wrapper object. The references +/// may live in bare object inputs or in members of a static archive. +static Error collectSYCLRegisterLibSymbols(MemoryBufferRef Buffer) { + switch (identify_magic(Buffer.getBuffer())) { + case file_magic::elf_relocatable: + case file_magic::elf_shared_object: + case file_magic::elf_executable: + case file_magic::coff_object: + case file_magic::macho_object: + if (Error Err = collectSYCLRegisterLibSymbolsFromObject(Buffer)) + return Err; + break; + case file_magic::archive: { + Expected> ArOrErr = + object::Archive::create(Buffer); + if (!ArOrErr) + return ArOrErr.takeError(); + Error Err = Error::success(); + for (const object::Archive::Child &Child : (*ArOrErr)->children(Err)) { + Expected ChildBufOrErr = Child.getMemoryBufferRef(); + if (!ChildBufOrErr) + return ChildBufOrErr.takeError(); + if (Error E = collectSYCLRegisterLibSymbolsFromObject(*ChildBufOrErr)) + return E; + } + if (Err) + return Err; + break; + } + default: + break; + } + return Error::success(); +} + /// Search the input files and libraries for embedded device offloading code /// and add it to the list of files to be linked. Files coming from static /// libraries are only added to the input if they are used by an existing @@ -3379,6 +3506,8 @@ getDeviceInput(const ArgList &Args) { MemoryBufferRef Buffer = **BufferOrErr; if (identify_magic(Buffer.getBuffer()) == file_magic::elf_shared_object) continue; + if (Error Err = collectSYCLRegisterLibSymbols(Buffer)) + return std::move(Err); SmallVector Binaries; size_t OldSize = Binaries.size(); if (Error Err = extractOffloadBinaries(Buffer, Binaries)) { diff --git a/sycl/test-e2e/NewOffloadDriver/Inputs/static-lib-registerlib-bmg.cpp b/sycl/test-e2e/NewOffloadDriver/Inputs/static-lib-registerlib-bmg.cpp new file mode 100644 index 0000000000000..9a29e4608cbce --- /dev/null +++ b/sycl/test-e2e/NewOffloadDriver/Inputs/static-lib-registerlib-bmg.cpp @@ -0,0 +1,21 @@ +//===----------------------------------------------------------------------===// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +//===----------------------------------------------------------------------===// + +#include +#include + +#include "static-lib-registerlib.hpp" + +void run_bmg(std::size_t *out, std::size_t N) { + sycl::queue q; + std::size_t *p = sycl::malloc_shared(N, q); + q.parallel_for(sycl::range<1>{N}, [=](sycl::id<1> idx) { + p[idx] = idx[0] + BMGOffset; + }).wait(); + for (std::size_t i = 0; i < N; ++i) + out[i] = p[i]; + sycl::free(p, q); +} diff --git a/sycl/test-e2e/NewOffloadDriver/Inputs/static-lib-registerlib-dg2.cpp b/sycl/test-e2e/NewOffloadDriver/Inputs/static-lib-registerlib-dg2.cpp new file mode 100644 index 0000000000000..b9e85957e37b5 --- /dev/null +++ b/sycl/test-e2e/NewOffloadDriver/Inputs/static-lib-registerlib-dg2.cpp @@ -0,0 +1,21 @@ +//===----------------------------------------------------------------------===// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +//===----------------------------------------------------------------------===// + +#include +#include + +#include "static-lib-registerlib.hpp" + +void run_dg2(std::size_t *out, std::size_t N) { + sycl::queue q; + std::size_t *p = sycl::malloc_shared(N, q); + q.parallel_for(sycl::range<1>{N}, [=](sycl::id<1> idx) { + p[idx] = idx[0] + DG2Offset; + }).wait(); + for (std::size_t i = 0; i < N; ++i) + out[i] = p[i]; + sycl::free(p, q); +} diff --git a/sycl/test-e2e/NewOffloadDriver/Inputs/static-lib-registerlib-jit.cpp b/sycl/test-e2e/NewOffloadDriver/Inputs/static-lib-registerlib-jit.cpp new file mode 100644 index 0000000000000..d3d075df927e6 --- /dev/null +++ b/sycl/test-e2e/NewOffloadDriver/Inputs/static-lib-registerlib-jit.cpp @@ -0,0 +1,21 @@ +//===----------------------------------------------------------------------===// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +//===----------------------------------------------------------------------===// + +#include +#include + +#include "static-lib-registerlib.hpp" + +void run_jit(std::size_t *out, std::size_t N) { + sycl::queue q; + std::size_t *p = sycl::malloc_shared(N, q); + q.parallel_for(sycl::range<1>{N}, [=](sycl::id<1> idx) { + p[idx] = idx[0] + JITOffset; + }).wait(); + for (std::size_t i = 0; i < N; ++i) + out[i] = p[i]; + sycl::free(p, q); +} diff --git a/sycl/test-e2e/NewOffloadDriver/Inputs/static-lib-registerlib.hpp b/sycl/test-e2e/NewOffloadDriver/Inputs/static-lib-registerlib.hpp new file mode 100644 index 0000000000000..35e1bcbc89001 --- /dev/null +++ b/sycl/test-e2e/NewOffloadDriver/Inputs/static-lib-registerlib.hpp @@ -0,0 +1,18 @@ +//===----------------------------------------------------------------------===// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +//===----------------------------------------------------------------------===// +// Shared declarations for static-lib-multiple-objects-multiple-targets.cpp. + +#pragma once + +#include + +void run_jit(std::size_t *out, std::size_t N); +void run_dg2(std::size_t *out, std::size_t N); +void run_bmg(std::size_t *out, std::size_t N); + +inline constexpr std::size_t JITOffset = 0; +inline constexpr std::size_t DG2Offset = 100; +inline constexpr std::size_t BMGOffset = 1000; diff --git a/sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets-win.cpp b/sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets-win.cpp new file mode 100644 index 0000000000000..189e649453fcc --- /dev/null +++ b/sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets-win.cpp @@ -0,0 +1,32 @@ +//==- static-lib-multiple-objects-multiple-targets-win.cpp ----------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// Windows counterpart of static-lib-multiple-objects-multiple-targets.cpp: same +// scenario, built in MSVC mode (clang-cl, /MD) with cl for the final host link. +// The kernels and main are reused via the #include below. + +// REQUIRES: ocloc, target-spir, windows +// REQUIRES: gpu-intel-dg2 || arch-intel_gpu_bmg_g21 + +// DEFINE: %{clcxx} = %clangxx --driver-mode=cl -fsycl --offload-new-driver /MD -Wno-error=unused-command-line-argument + +// RUN: %{clcxx} -fsycl-targets=spir64 -c %S/Inputs/static-lib-registerlib-jit.cpp -o %t.jit.obj +// RUN: %{clcxx} -fsycl-link -fsycl-targets=spir64 %t.jit.obj -o %t.jit.pre.obj +// RUN: %{clcxx} -fsycl-targets=intel_gpu_dg2 -c %S/Inputs/static-lib-registerlib-dg2.cpp -o %t.dg2.obj +// RUN: %{clcxx} -fsycl-link -fsycl-targets=intel_gpu_dg2 %t.dg2.obj -o %t.dg2.pre.obj +// RUN: %{clcxx} -fsycl-targets=intel_gpu_bmg_g21 -c %S/Inputs/static-lib-registerlib-bmg.cpp -o %t.bmg.obj +// RUN: %{clcxx} -fsycl-link -fsycl-targets=intel_gpu_bmg_g21 %t.bmg.obj -o %t.bmg.pre.obj + +// RUN: rm -f %t.a +// RUN: llvm-ar crv %t.a %t.jit.obj %t.jit.pre.obj %t.dg2.obj %t.dg2.pre.obj %t.bmg.obj %t.bmg.pre.obj + +// RUN: %clangxx --driver-mode=cl /MD /std:c++17 %sycl_include -c %s -o %t.main.obj +// RUN: cl -nologo /MD %t.main.obj -Fe%t.exe -link %t.a /defaultlib:%sycl_static_libs_dir/sycl.lib +// RUN: %{run} %t.exe + +// Reuse the kernels and main() from the Linux test. +#include "static-lib-multiple-objects-multiple-targets.cpp" diff --git a/sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets.cpp b/sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets.cpp new file mode 100644 index 0000000000000..c4dc9d063ecbc --- /dev/null +++ b/sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets.cpp @@ -0,0 +1,100 @@ +//==- static-lib-multiple-objects-multiple-targets.cpp --------------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// Pull SYCL device code out of a static library, without --whole-archive, using +// the new offloading model and gcc for the final host link. The archive holds +// three translation units, each pre-linked with -fsycl-link for a single +// distinct target (spir64 / intel_gpu_dg2 / intel_gpu_bmg_g21), so each +// contributes its own __sycl_registerlib_. +// +// The Windows counterpart is +// static-lib-multiple-objects-multiple-targets-win.cpp. + +// REQUIRES: ocloc, target-spir, linux +// REQUIRES: gpu-intel-dg2 || arch-intel_gpu_bmg_g21 + +// DEFINE: %{ond} = --offload-new-driver -Wno-error=unused-command-line-argument + +// RUN: %clangxx -fsycl %{ond} -fsycl-targets=spir64 -fPIC -c %S/Inputs/static-lib-registerlib-jit.cpp -o %t.jit.o +// RUN: %clangxx -fsycl %{ond} -fsycl-link -fsycl-targets=spir64 -fPIC %t.jit.o -o %t.jit.pre.o +// RUN: %clangxx -fsycl %{ond} -fsycl-targets=intel_gpu_dg2 -fPIC -c %S/Inputs/static-lib-registerlib-dg2.cpp -o %t.dg2.o +// RUN: %clangxx -fsycl %{ond} -fsycl-link -fsycl-targets=intel_gpu_dg2 -fPIC %t.dg2.o -o %t.dg2.pre.o +// RUN: %clangxx -fsycl %{ond} -fsycl-targets=intel_gpu_bmg_g21 -fPIC -c %S/Inputs/static-lib-registerlib-bmg.cpp -o %t.bmg.o +// RUN: %clangxx -fsycl %{ond} -fsycl-link -fsycl-targets=intel_gpu_bmg_g21 -fPIC %t.bmg.o -o %t.bmg.pre.o + +// RUN: rm -f %t.a +// RUN: llvm-ar crv %t.a %t.jit.o %t.jit.pre.o %t.dg2.o %t.dg2.pre.o %t.bmg.o %t.bmg.pre.o + +// RUN: %clangxx -Wno-error=unused-command-line-argument %sycl_include %cxx_std_optionc++17 -fPIC -c %s -o %t.main.o +// RUN: g++ %t.main.o %t.a -L%sycl_libs_dir -lsycl -Wl,-rpath,%sycl_libs_dir -o %t.exe +// RUN: %{run} %t.exe + +// The same library, found through -l and linked by the SYCL driver. With a +// plain host main nothing pulls the members' device code into the link, so +// their pre-linked objects have to be pulled in by the +// __sycl_registerlib_ references instead. +// RUN: mkdir -p %t.dir && cp %t.a %t.dir/libregisterlib.a +// RUN: %clangxx -fsycl %{ond} %t.main.o -L%t.dir -lregisterlib -o %t.l.exe +// RUN: %{run} %t.l.exe + +// With a SYCL main the spir64 member's device code is linked into the wrapper +// object, while the AOT members' still comes from their pre-linked objects. +// RUN: %clangxx -fsycl %{ond} %cxx_std_optionc++17 -fPIC -c %s -o %t.main.sycl.o +// RUN: %clangxx -fsycl %{ond} %t.main.sycl.o -L%t.dir -lregisterlib -o %t.l.sycl.exe +// RUN: %{run} %t.l.sycl.exe + +#include +#include + +#include +#include + +#include "Inputs/static-lib-registerlib.hpp" + +namespace syclex = sycl::ext::oneapi::experimental; + +static bool check(const char *Name, std::size_t Offset, + void (*Run)(std::size_t *, std::size_t)) { + constexpr std::size_t N = 8; + std::size_t out[N] = {}; + Run(out, N); + for (std::size_t I = 0; I < N; ++I) { + if (out[I] != I + Offset) { + std::cout << "fail " << Name << ": out[" << I << "] == " << out[I] + << ", expected " << (I + Offset) << "\n"; + return false; + } + } + std::cout << Name << " ran\n"; + return true; +} + +int main() { + // The JIT (spir64) kernel runs on any SPIR device; the AOT kernels only have + // an image for their target, so only invoke the one matching this device. + bool Ok = check("jit", JITOffset, run_jit); + auto Arch = + sycl::queue{}.get_device().get_info(); + using arch = syclex::architecture; + switch (Arch) { + case arch::intel_gpu_dg2_g10: + case arch::intel_gpu_dg2_g11: + case arch::intel_gpu_dg2_g12: + Ok &= check("dg2", DG2Offset, run_dg2); + break; + case arch::intel_gpu_bmg_g21: + Ok &= check("bmg", BMGOffset, run_bmg); + break; + default: + std::cout << "no matching AOT kernel for this device\n"; + break; + } + + if (Ok) + std::cout << "pass\n"; + return Ok ? 0 : 1; +} From 72b85a59b3efa427cec401e9930a27478d6adb88 Mon Sep 17 00:00:00 2001 From: Nick Sarnie Date: Thu, 8 Oct 2026 13:14:04 -0700 Subject: [PATCH 02/19] Address feedback and fix -fsycl-host-only Signed-off-by: Nick Sarnie --- clang/lib/CodeGen/CodeGenModule.cpp | 5 +++- clang/test/CodeGenSYCL/registerlib-host.cpp | 8 ++++--- .../ClangLinkerWrapper.cpp | 24 +++++++++---------- 3 files changed, 21 insertions(+), 16 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 8393c5c51e025..acaec5ec97815 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1291,7 +1291,10 @@ void CodeGenModule::Release() { // A static initializer may launch a kernel, so the device binary has to // be registered before any of them run, hence a priority. AddGlobalCtor(SYCLCtorFunction, /*Priority=*/101); - } else if (LangOpts.OffloadingNewDriver) { + } else if (LangOpts.OffloadingNewDriver && + !CodeGenOpts.OffloadObjects.empty()) { + // Only add registration functions if there has been a device compiler (so + // OffloadBinaryToEmbedFile or OffloadObjects is not empty). AddGlobalCtor(createSYCLRegisterLibFunc()); } } diff --git a/clang/test/CodeGenSYCL/registerlib-host.cpp b/clang/test/CodeGenSYCL/registerlib-host.cpp index 07cb84139aef4..aca1042dc60bf 100644 --- a/clang/test/CodeGenSYCL/registerlib-host.cpp +++ b/clang/test/CodeGenSYCL/registerlib-host.cpp @@ -2,15 +2,17 @@ // host code generation emits under the new offload driver. The id is the // translation unit's -fsycl-unique-prefix, which the driver generates. -// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsycl-is-host --offload-new-driver -fsycl-unique-prefix=uid42 -emit-llvm %s -o - | FileCheck %s +// RUN: touch %t.o + +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsycl-is-host --offload-new-driver -fsycl-unique-prefix=uid42 -emit-llvm -fembed-offload-object=%t.o %s -o - | FileCheck %s // Nothing is emitted without the new offload driver. -// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsycl-is-host -fsycl-unique-prefix=uid42 -emit-llvm %s -o - | FileCheck %s --check-prefix=OLD +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsycl-is-host -fsycl-unique-prefix=uid42 -emit-llvm -fembed-offload-object=%t.o %s -o - | FileCheck %s --check-prefix=OLD // Nor when the module embeds its device binary (-fno-sycl-rdc), as it then // registers the binary itself. // RUN: echo -n 'FAKE_SYCL_DEVICE_IMAGE' > %t.bin -// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsycl-is-host --offload-new-driver -foffload-include-binary %t.bin -fsycl-unique-prefix=uid42 -emit-llvm %s -o - | FileCheck %s --check-prefix=EMBED --implicit-check-not=__sycl_registerlib_ +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsycl-is-host --offload-new-driver -foffload-include-binary %t.bin -fsycl-unique-prefix=uid42 -emit-llvm -fembed-offload-object=%t.o %s -o - | FileCheck %s --check-prefix=EMBED --implicit-check-not=__sycl_registerlib_ // CHECK: @llvm.global_ctors = appending global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 65535, ptr @__sycl_registerlib_ctor, ptr null }] // CHECK: declare void @__sycl_registerlib_uid42() diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index c83b341cb9806..bce27eead4ccc 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -1137,17 +1137,13 @@ runAOTCompileIntelCPU(StringRef InputFile, const ArgList &Args, static void defineRegisterLibFuncs(Module &M, const ArgList &Args) { const llvm::Triple Triple(Args.getLastArgValue(OPT_triple_EQ)); const std::string Arch(Args.getLastArgValue(OPT_arch_EQ)); - // In the case we have a fat object input with multiple device images of - // differing triples, and we are not doing -fsycl-link (so multiple linked - // device images are allowed), the host object defined a single registerlib - // symbol, so we need to make sure to only define the function once, not once - // per device image from that fat object. We need to keep track of the triple - // through, because in the -fsycl-link case, we may not use one of those - // device images if it doesn't match the -fsycl-link triple. - static StringSet TotalDefinedFuncs; - for (auto Name : SYCLRegisterLibFuncMap[std::make_pair(Triple.str(), Arch)]) { - if (TotalDefinedFuncs.find(Name) != TotalDefinedFuncs.end()) - continue; + // We need to keep track of the triple and arch because we may not use all + // device images passed to the tool and we don't want to force unused objects + // to be linked in. + auto It = SYCLRegisterLibFuncMap.find(std::make_pair(Triple.str(), Arch)); + if (It == SYCLRegisterLibFuncMap.end()) + return; + for (std::string Name : It->second) { llvm::FunctionType *FTy = llvm::FunctionType::get( llvm::Type::getVoidTy(M.getContext()), /*isVarArg=*/false); auto *Fn = llvm::Function::Create(FTy, llvm::GlobalValue::WeakAnyLinkage, @@ -1156,7 +1152,11 @@ static void defineRegisterLibFuncs(Module &M, const ArgList &Args) { llvm::BasicBlock::Create(M.getContext(), "entry", Fn); llvm::IRBuilder<> Builder(Entry); Builder.CreateRetVoid(); - TotalDefinedFuncs.insert(Name); + if (M.getTargetTriple().supportsCOMDAT()) { + llvm::Comdat *C = M.getOrInsertComdat(Name); + C->setSelectionKind(llvm::Comdat::Any); + Fn->setComdat(C); + } } } From c7101ffbdcd24c1823c634271b4d5f5c973de2e3 Mon Sep 17 00:00:00 2001 From: Nick Sarnie Date: Thu, 8 Oct 2026 13:33:47 -0700 Subject: [PATCH 03/19] fix typo Signed-off-by: Nick Sarnie --- clang/lib/CodeGen/CodeGenModule.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index acaec5ec97815..4d76d0fadd363 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1293,8 +1293,8 @@ void CodeGenModule::Release() { AddGlobalCtor(SYCLCtorFunction, /*Priority=*/101); } else if (LangOpts.OffloadingNewDriver && !CodeGenOpts.OffloadObjects.empty()) { - // Only add registration functions if there has been a device compiler (so - // OffloadBinaryToEmbedFile or OffloadObjects is not empty). + // Only add registration functions if there has been a device compile (so + // OffloadObjects is not empty). AddGlobalCtor(createSYCLRegisterLibFunc()); } } From 629ecd3ccc3b578ebaa418602fc5d039f3073dbb Mon Sep 17 00:00:00 2001 From: Nick Sarnie Date: Fri, 9 Oct 2026 09:01:21 -0700 Subject: [PATCH 04/19] fix lto and spir64_gen no arch Signed-off-by: Nick Sarnie --- .../ClangLinkerWrapper.cpp | 73 +++++++++++++------ 1 file changed, 49 insertions(+), 24 deletions(-) diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index c4e76986317f6..16e791ed40389 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -1137,26 +1137,48 @@ runAOTCompileIntelCPU(StringRef InputFile, const ArgList &Args, static void defineRegisterLibFuncs(Module &M, const ArgList &Args) { const llvm::Triple Triple(Args.getLastArgValue(OPT_triple_EQ)); const std::string Arch(Args.getLastArgValue(OPT_arch_EQ)); + bool IsGPUAOT = Triple.isSPIROrSPIRV() && + Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen; // We need to keep track of the triple and arch because we may not use all // device images passed to the tool and we don't want to force unused objects // to be linked in. auto It = SYCLRegisterLibFuncMap.find(std::make_pair(Triple.str(), Arch)); if (It == SYCLRegisterLibFuncMap.end()) return; - for (std::string Name : It->second) { - llvm::FunctionType *FTy = llvm::FunctionType::get( - llvm::Type::getVoidTy(M.getContext()), /*isVarArg=*/false); - auto *Fn = llvm::Function::Create(FTy, llvm::GlobalValue::WeakAnyLinkage, - Name, &M); - llvm::BasicBlock *Entry = - llvm::BasicBlock::Create(M.getContext(), "entry", Fn); - llvm::IRBuilder<> Builder(Entry); - Builder.CreateRetVoid(); - if (M.getTargetTriple().supportsCOMDAT()) { - llvm::Comdat *C = M.getOrInsertComdat(Name); - C->setSelectionKind(llvm::Comdat::Any); - Fn->setComdat(C); + auto DefineFuncs = [&](auto &&Range) { + for (const std::string &Name : Range) { + if (M.getFunction(Name)) + continue; + llvm::FunctionType *FTy = llvm::FunctionType::get( + llvm::Type::getVoidTy(M.getContext()), /*isVarArg=*/false); + auto *Fn = llvm::Function::Create(FTy, llvm::GlobalValue::WeakAnyLinkage, + Name, &M); + llvm::BasicBlock *Entry = + llvm::BasicBlock::Create(M.getContext(), "entry", Fn); + llvm::IRBuilder<> Builder(Entry); + Builder.CreateRetVoid(); + if (M.getTargetTriple().supportsCOMDAT()) { + llvm::Comdat *C = M.getOrInsertComdat(Name); + C->setSelectionKind(llvm::Comdat::Any); + Fn->setComdat(C); + } } + }; + // There is a special case when the current target being linked in spir64_gen + // with non-empty arch. We need to consider the cases where there are other + // device objects with spir64_gen triple with empty (generic) arch because + // this is considered a compatible target with spir64_gen non-empty arch and + // will be linked together with it. In this case, we need to also define + // symbols from the empty-arch device objects. + auto GenericArchIt = + SYCLRegisterLibFuncMap.find(std::make_pair(Triple.str(), std::string())); + if (!IsGPUAOT || GenericArchIt == SYCLRegisterLibFuncMap.end() || + GenericArchIt->second.empty()) { + DefineFuncs(It->second); + } else { + // Just iterate over both sets, duplcates are handled in DefineFuncs. + DefineFuncs( + llvm::concat(It->second, GenericArchIt->second)); } } @@ -3364,26 +3386,31 @@ Error emitExtractCommands( /// See comments in collectSYCLRegisterLibSymbols. static Error collectSYCLRegisterLibSymbolsFromObject(MemoryBufferRef Buffer) { + LLVMContext Context; static constexpr StringRef SYCLRegisterLibPrefix = "__sycl_registerlib_"; - Expected> ObjOrErr = - ObjectFile::createObjectFile(Buffer); + llvm::file_magic Magic = identify_magic(Buffer.getBuffer()); + Expected> ObjOrErr = + SymbolicFile::createSymbolicFile(Buffer, Magic, &Context, + /*InitContent=*/Magic != + llvm::file_magic::bitcode); if (!ObjOrErr) { // Not something we can read symbols from; ignore it. consumeError(ObjOrErr.takeError()); return Error::success(); } SmallSet FuncNames; - for (const SymbolRef &Sym : (*ObjOrErr)->symbols()) { + for (const BasicSymbolRef &Sym : (*ObjOrErr)->symbols()) { Expected FlagsOrErr = Sym.getFlags(); if (!FlagsOrErr) return FlagsOrErr.takeError(); - Expected NameOrErr = Sym.getName(); - if (!NameOrErr) - return NameOrErr.takeError(); - if (!NameOrErr->starts_with(SYCLRegisterLibPrefix)) + std::string Name; + llvm::raw_string_ostream OS(Name); + if (Error Err = Sym.printName(OS)) + return Err; + if (!StringRef(Name).starts_with(SYCLRegisterLibPrefix)) continue; if (*FlagsOrErr & SymbolRef::SF_Undefined) - FuncNames.insert(NameOrErr->str()); + FuncNames.insert(Name); } if (!FuncNames.empty()) { SmallVector Binaries; @@ -3409,10 +3436,8 @@ static Error collectSYCLRegisterLibSymbolsFromObject(MemoryBufferRef Buffer) { static Error collectSYCLRegisterLibSymbols(MemoryBufferRef Buffer) { switch (identify_magic(Buffer.getBuffer())) { case file_magic::elf_relocatable: - case file_magic::elf_shared_object: - case file_magic::elf_executable: case file_magic::coff_object: - case file_magic::macho_object: + case file_magic::bitcode: if (Error Err = collectSYCLRegisterLibSymbolsFromObject(Buffer)) return Err; break; From 92b57f779a02f22877461bed50b21fcb3eba0310 Mon Sep 17 00:00:00 2001 From: Nick Sarnie Date: Mon, 12 Oct 2026 08:44:49 -0700 Subject: [PATCH 05/19] make clang link test case clearer Signed-off-by: Nick Sarnie --- .../static-lib-multiple-objects-multiple-targets.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets.cpp b/sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets.cpp index c4dc9d063ecbc..95a0cb18d1cee 100644 --- a/sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets.cpp +++ b/sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets.cpp @@ -33,12 +33,12 @@ // RUN: g++ %t.main.o %t.a -L%sycl_libs_dir -lsycl -Wl,-rpath,%sycl_libs_dir -o %t.exe // RUN: %{run} %t.exe -// The same library, found through -l and linked by the SYCL driver. With a +// The same library, found through -l and linked by clang (non-SYCL mode). With a // plain host main nothing pulls the members' device code into the link, so // their pre-linked objects have to be pulled in by the // __sycl_registerlib_ references instead. // RUN: mkdir -p %t.dir && cp %t.a %t.dir/libregisterlib.a -// RUN: %clangxx -fsycl %{ond} %t.main.o -L%t.dir -lregisterlib -o %t.l.exe +// RUN: %clangxx %{ond} %t.main.o -L%sycl_libs_dir -lsycl -Wl,-rpath,%sycl_libs_dir -L%t.dir -lregisterlib -o %t.l.exe // RUN: %{run} %t.l.exe // With a SYCL main the spir64 member's device code is linked into the wrapper From 99608f0a26093fe2d1d78dff414d9973f4d26190 Mon Sep 17 00:00:00 2001 From: Nick Sarnie Date: Mon, 12 Oct 2026 09:09:55 -0700 Subject: [PATCH 06/19] add lto test Signed-off-by: Nick Sarnie --- sycl/test-e2e/E2EExpr.py | 1 + ...-multiple-objects-multiple-targets-lto.cpp | 33 +++++++++++++++++++ ...-lib-multiple-objects-multiple-targets.cpp | 4 +-- sycl/test-e2e/lit.cfg.py | 1 + 4 files changed, 37 insertions(+), 2 deletions(-) create mode 100644 sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets-lto.cpp diff --git a/sycl/test-e2e/E2EExpr.py b/sycl/test-e2e/E2EExpr.py index 19348ebe0384f..3e5e1a8c5e312 100644 --- a/sycl/test-e2e/E2EExpr.py +++ b/sycl/test-e2e/E2EExpr.py @@ -58,6 +58,7 @@ class E2EExpr(BooleanExpression): "sg-8", "sg-16", "sg-32", + "lld", } def __init__(self, string, variables, build_only_mode, final_unknown_value): diff --git a/sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets-lto.cpp b/sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets-lto.cpp new file mode 100644 index 0000000000000..006cf6f516af5 --- /dev/null +++ b/sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets-lto.cpp @@ -0,0 +1,33 @@ +//==- static-lib-multiple-objects-multiple-targets-lto.cpp--==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===------------------------------------------------------===// +// Same case as static-lib-multiple-objects-multiple-targets.cpp but using LTO +// objects. + +// REQUIRES: ocloc, target-spir, linux, lld +// REQUIRES: gpu-intel-dg2 || arch-intel_gpu_bmg_g21 + +// DEFINE: %{ond} = --offload-new-driver -Wno-error=unused-command-line-argument + +// RUN: %clangxx -fsycl %{ond} -flto -fsycl-targets=spir64 -fPIC -c %S/Inputs/static-lib-registerlib-jit.cpp -o %t.jit.o +// RUN: %clangxx -fsycl %{ond} -flto -fsycl-link -fsycl-targets=spir64 -fPIC %t.jit.o -o %t.jit.pre.o +// RUN: %clangxx -fsycl %{ond} -flto -fsycl-targets=intel_gpu_dg2 -fPIC -c %S/Inputs/static-lib-registerlib-dg2.cpp -o %t.dg2.o +// RUN: %clangxx -fsycl %{ond} -flto -fsycl-link -fsycl-targets=intel_gpu_dg2 -fPIC %t.dg2.o -o %t.dg2.pre.o +// RUN: %clangxx -fsycl %{ond} -flto -fsycl-targets=intel_gpu_bmg_g21 -fPIC -c %S/Inputs/static-lib-registerlib-bmg.cpp -o %t.bmg.o +// RUN: %clangxx -fsycl %{ond} -flto -fsycl-link -fsycl-targets=intel_gpu_bmg_g21 -fPIC %t.bmg.o -o %t.bmg.pre.o + +// RUN: rm -f %t.a +// RUN: llvm-ar crv %t.a %t.jit.o %t.jit.pre.o %t.dg2.o %t.dg2.pre.o %t.bmg.o %t.bmg.pre.o + +// RUN: %clangxx %{ond} -flto -Wno-error=unused-command-line-argument %sycl_include %cxx_std_optionc++17 -fPIC -c %s -o %t.main.o + +// Note this is a non-SYCL compile. +// RUN: %clangxx %{ond} -flto -fuse-ld=lld %t.main.o %t.a -L%sycl_libs_dir -lsycl -Wl,-rpath,%sycl_libs_dir -o %t.exe + +// RUN: %{run} %t.exe + +#include "static-lib-multiple-objects-multiple-targets.cpp" diff --git a/sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets.cpp b/sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets.cpp index 95a0cb18d1cee..6c2cea5f6a1c7 100644 --- a/sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets.cpp +++ b/sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets.cpp @@ -33,8 +33,8 @@ // RUN: g++ %t.main.o %t.a -L%sycl_libs_dir -lsycl -Wl,-rpath,%sycl_libs_dir -o %t.exe // RUN: %{run} %t.exe -// The same library, found through -l and linked by clang (non-SYCL mode). With a -// plain host main nothing pulls the members' device code into the link, so +// The same library, found through -l and linked by clang (non-SYCL mode). With +// a plain host main nothing pulls the members' device code into the link, so // their pre-linked objects have to be pulled in by the // __sycl_registerlib_ references instead. // RUN: mkdir -p %t.dir && cp %t.a %t.dir/libregisterlib.a diff --git a/sycl/test-e2e/lit.cfg.py b/sycl/test-e2e/lit.cfg.py index 99977019f962f..fadf0919078f5 100644 --- a/sycl/test-e2e/lit.cfg.py +++ b/sycl/test-e2e/lit.cfg.py @@ -914,6 +914,7 @@ def remove_level_zero_suffix(devices): ToolSubst("llvm-spirv", unresolved="ignore"), ToolSubst("opencl-aot", unresolved="ignore"), ToolSubst("ocloc", unresolved="ignore"), + ToolSubst("lld", command=FindTool("ld.lld"), unresolved="ignore"), ] tools = [ From 116f73267b5c6c07182675c9a3884e68014151a4 Mon Sep 17 00:00:00 2001 From: Nick Sarnie Date: Mon, 12 Oct 2026 12:01:26 -0700 Subject: [PATCH 07/19] reuse Binaries vector Signed-off-by: Nick Sarnie --- .../ClangLinkerWrapper.cpp | 24 ++++++++++++------- 1 file changed, 15 insertions(+), 9 deletions(-) diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 16e791ed40389..2a08fb12fbe29 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -3385,7 +3385,9 @@ Error emitExtractCommands( } /// See comments in collectSYCLRegisterLibSymbols. -static Error collectSYCLRegisterLibSymbolsFromObject(MemoryBufferRef Buffer) { +static Error +collectSYCLRegisterLibSymbolsFromObject(MemoryBufferRef Buffer, + ArrayRef Binaries) { LLVMContext Context; static constexpr StringRef SYCLRegisterLibPrefix = "__sycl_registerlib_"; llvm::file_magic Magic = identify_magic(Buffer.getBuffer()); @@ -3413,10 +3415,11 @@ static Error collectSYCLRegisterLibSymbolsFromObject(MemoryBufferRef Buffer) { FuncNames.insert(Name); } if (!FuncNames.empty()) { - SmallVector Binaries; - if (Error Err = extractOffloadBinaries(Buffer, Binaries)) - return std::move(Err); for (const auto &Binary : Binaries) { + // Only use the binaries embedded in this object. + if (Binary.getBinary()->getMemoryBufferRef().getBufferIdentifier() != + Buffer.getBufferIdentifier()) + continue; std::string Triple = Binary.getBinary()->getTriple().str(); std::string Arch = Binary.getBinary()->getArch().str(); if (Arch == "generic") @@ -3433,12 +3436,14 @@ static Error collectSYCLRegisterLibSymbolsFromObject(MemoryBufferRef Buffer) { /// global SYCLRegisterLibFuncMap. Each such reference is satisfied by defining /// the symbol in the generated SYCL device-image wrapper object. The references /// may live in bare object inputs or in members of a static archive. -static Error collectSYCLRegisterLibSymbols(MemoryBufferRef Buffer) { +/// \p Binaries are the offload binaries already extracted from \p Buffer. +static Error collectSYCLRegisterLibSymbols(MemoryBufferRef Buffer, + ArrayRef Binaries) { switch (identify_magic(Buffer.getBuffer())) { case file_magic::elf_relocatable: case file_magic::coff_object: case file_magic::bitcode: - if (Error Err = collectSYCLRegisterLibSymbolsFromObject(Buffer)) + if (Error Err = collectSYCLRegisterLibSymbolsFromObject(Buffer, Binaries)) return Err; break; case file_magic::archive: { @@ -3451,7 +3456,8 @@ static Error collectSYCLRegisterLibSymbols(MemoryBufferRef Buffer) { Expected ChildBufOrErr = Child.getMemoryBufferRef(); if (!ChildBufOrErr) return ChildBufOrErr.takeError(); - if (Error E = collectSYCLRegisterLibSymbolsFromObject(*ChildBufOrErr)) + if (Error E = + collectSYCLRegisterLibSymbolsFromObject(*ChildBufOrErr, Binaries)) return E; } if (Err) @@ -3540,8 +3546,6 @@ getDeviceInput(const ArgList &Args) { MemoryBufferRef Buffer = **BufferOrErr; if (identify_magic(Buffer.getBuffer()) == file_magic::elf_shared_object) continue; - if (Error Err = collectSYCLRegisterLibSymbols(Buffer)) - return std::move(Err); SmallVector Binaries; size_t OldSize = Binaries.size(); if (Error Err = extractOffloadBinaries(Buffer, Binaries)) { @@ -3550,6 +3554,8 @@ getDeviceInput(const ArgList &Args) { // fall through to the SYCL bundled objects extraction path. consumeError(std::move(Err)); } + if (Error Err = collectSYCLRegisterLibSymbols(Buffer, Binaries)) + return std::move(Err); if (Binaries.size() == OldSize) { if (Error Err = sycl::extractBundledObjects(*Filename, Args, Binaries)) return std::move(Err); From e8cf02cde0c46a48382fe3f228ce23eaa565c410 Mon Sep 17 00:00:00 2001 From: Nick Sarnie Date: Mon, 12 Oct 2026 13:00:11 -0700 Subject: [PATCH 08/19] avoid odr violation and cleanup comments Signed-off-by: Nick Sarnie --- clang/lib/CodeGen/CodeGenModule.h | 4 +-- .../ClangLinkerWrapper.cpp | 33 ++++++++++--------- ...-lib-multiple-objects-multiple-targets.cpp | 4 +-- 3 files changed, 22 insertions(+), 19 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h index a48779081ed3c..884774566adc6 100644 --- a/clang/lib/CodeGen/CodeGenModule.h +++ b/clang/lib/CodeGen/CodeGenModule.h @@ -2232,14 +2232,14 @@ class CodeGenModule : public CodeGenTypeCache { /// if the binary could not be read. llvm::Function *embedSYCLDeviceBinary(); - /// Emit a per-module __sycl_registerlib_ reference that + /// Emit a per-module __sycl_registerlib_ reference that /// ties the host and device sides of a translation unit together. This is /// used to force linking of static libraries: we emit a global constructor /// that references (calls) the symbol, creating an undefined reference that /// forces the corresponding device static library member (which defines the /// matching symbol) to be pulled in at link time. The device-side definition /// is emitted by clang-linker-wrapper, so this is new offload model only. The - /// hash of the module identifier is what ties the two sides together. + /// unique identifier is what ties the two sides together. /// \return the global constructor that calls registerlib function. llvm::Function *createSYCLRegisterLibFunc(); diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 2a08fb12fbe29..6b2205902687a 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -168,8 +168,18 @@ static bool LinkingSYCLBINFiles = false; static SmallString<128> OffloadImageDumpDir; using TripleAndArch = std::pair; +struct TripleAndArchInfo { + static unsigned getHashValue(const TripleAndArch &K) { + return hash_combine(K.first, K.second); + } + static bool isEqual(const TripleAndArch &L, const TripleAndArch &R) { + return L == R; + } +}; + using FuncNameSet = SmallSet; -static DenseMap SYCLRegisterLibFuncMap; +static DenseMap + SYCLRegisterLibFuncMap; /// Whether or not to look through symlinks when resolving binaries. static bool CanonicalPrefixes = true; @@ -192,14 +202,6 @@ template <> struct DenseMapInfo { return LHS == RHS; } }; -template <> struct DenseMapInfo { - static unsigned getHashValue(const std::string &S) { - return hash_value(StringRef(S)); - } - static bool isEqual(const std::string &L, const std::string &R) { - return L == R; - } -}; } // namespace llvm namespace { @@ -1128,10 +1130,10 @@ runAOTCompileIntelCPU(StringRef InputFile, const ArgList &Args, return *TempFileOrErr; } -/// Define __sycl_registerlib_ functions defined by the +/// Define __sycl_registerlib_ functions referenced by the /// host compile to force all device code images to be linked in, which /// may not happen for static archives with -fsycl-link. -/// \p Module The module. +/// \p M The module. /// \p Args encompasses all arguments required for linking and wrapping device /// code. static void defineRegisterLibFuncs(Module &M, const ArgList &Args) { @@ -3432,10 +3434,11 @@ collectSYCLRegisterLibSymbolsFromObject(MemoryBufferRef Buffer, } /// Scan the host input in \p Buffer for -/// undefined symbols named __sycl_registerlib_ and record them in the -/// global SYCLRegisterLibFuncMap. Each such reference is satisfied by defining -/// the symbol in the generated SYCL device-image wrapper object. The references -/// may live in bare object inputs or in members of a static archive. +/// undefined symbols named __sycl_registerlib_ and record +/// them in the global SYCLRegisterLibFuncMap. Each such reference is satisfied +/// by defining the symbol in the generated SYCL device-image wrapper object. +/// The references may live in bare object inputs or in members of a static +/// archive. /// \p Binaries are the offload binaries already extracted from \p Buffer. static Error collectSYCLRegisterLibSymbols(MemoryBufferRef Buffer, ArrayRef Binaries) { diff --git a/sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets.cpp b/sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets.cpp index 6c2cea5f6a1c7..87c8cdfdc4674 100644 --- a/sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets.cpp +++ b/sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets.cpp @@ -9,7 +9,7 @@ // the new offloading model and gcc for the final host link. The archive holds // three translation units, each pre-linked with -fsycl-link for a single // distinct target (spir64 / intel_gpu_dg2 / intel_gpu_bmg_g21), so each -// contributes its own __sycl_registerlib_. +// contributes its own __sycl_registerlib_. // // The Windows counterpart is // static-lib-multiple-objects-multiple-targets-win.cpp. @@ -36,7 +36,7 @@ // The same library, found through -l and linked by clang (non-SYCL mode). With // a plain host main nothing pulls the members' device code into the link, so // their pre-linked objects have to be pulled in by the -// __sycl_registerlib_ references instead. +// __sycl_registerlib_ references instead. // RUN: mkdir -p %t.dir && cp %t.a %t.dir/libregisterlib.a // RUN: %clangxx %{ond} %t.main.o -L%sycl_libs_dir -lsycl -Wl,-rpath,%sycl_libs_dir -L%t.dir -lregisterlib -o %t.l.exe // RUN: %{run} %t.l.exe From ca731f74b450d3c75204e9252b79ad5ced8ef8d7 Mon Sep 17 00:00:00 2001 From: Nick Sarnie Date: Mon, 12 Oct 2026 13:13:48 -0700 Subject: [PATCH 09/19] add spir64_gen no device test Signed-off-by: Nick Sarnie --- .../static-lib-registerlib-gen-generic.cpp | 22 ++++++++ .../Inputs/static-lib-registerlib.hpp | 4 +- .../static-lib-gen-generic-arch.cpp | 56 +++++++++++++++++++ 3 files changed, 81 insertions(+), 1 deletion(-) create mode 100644 sycl/test-e2e/NewOffloadDriver/Inputs/static-lib-registerlib-gen-generic.cpp create mode 100644 sycl/test-e2e/NewOffloadDriver/static-lib-gen-generic-arch.cpp diff --git a/sycl/test-e2e/NewOffloadDriver/Inputs/static-lib-registerlib-gen-generic.cpp b/sycl/test-e2e/NewOffloadDriver/Inputs/static-lib-registerlib-gen-generic.cpp new file mode 100644 index 0000000000000..d37f0c7d9d5b1 --- /dev/null +++ b/sycl/test-e2e/NewOffloadDriver/Inputs/static-lib-registerlib-gen-generic.cpp @@ -0,0 +1,22 @@ +//===----------------------------------------------------------------------===// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +//===----------------------------------------------------------------------===// + +#include +#include + +#include "static-lib-registerlib.hpp" + +void run_gen_generic(std::size_t *out, std::size_t N) { + sycl::queue q; + std::size_t *p = sycl::malloc_shared(N, q); + q.parallel_for( + sycl::range<1>{N}, + [=](sycl::id<1> idx) { p[idx] = idx[0] + GenGenericOffset; }) + .wait(); + for (std::size_t i = 0; i < N; ++i) + out[i] = p[i]; + sycl::free(p, q); +} diff --git a/sycl/test-e2e/NewOffloadDriver/Inputs/static-lib-registerlib.hpp b/sycl/test-e2e/NewOffloadDriver/Inputs/static-lib-registerlib.hpp index 35e1bcbc89001..4c49935729947 100644 --- a/sycl/test-e2e/NewOffloadDriver/Inputs/static-lib-registerlib.hpp +++ b/sycl/test-e2e/NewOffloadDriver/Inputs/static-lib-registerlib.hpp @@ -3,7 +3,7 @@ // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception //===----------------------------------------------------------------------===// -// Shared declarations for static-lib-multiple-objects-multiple-targets.cpp. +// Shared declarations for the static-lib-* registerlib tests. #pragma once @@ -12,7 +12,9 @@ void run_jit(std::size_t *out, std::size_t N); void run_dg2(std::size_t *out, std::size_t N); void run_bmg(std::size_t *out, std::size_t N); +void run_gen_generic(std::size_t *out, std::size_t N); inline constexpr std::size_t JITOffset = 0; inline constexpr std::size_t DG2Offset = 100; inline constexpr std::size_t BMGOffset = 1000; +inline constexpr std::size_t GenGenericOffset = 10000; diff --git a/sycl/test-e2e/NewOffloadDriver/static-lib-gen-generic-arch.cpp b/sycl/test-e2e/NewOffloadDriver/static-lib-gen-generic-arch.cpp new file mode 100644 index 0000000000000..de4ec8f962ff6 --- /dev/null +++ b/sycl/test-e2e/NewOffloadDriver/static-lib-gen-generic-arch.cpp @@ -0,0 +1,56 @@ +//==- static-lib-gen-generic-arch.cpp -------------------------------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// Link an object built for bmg against a static library whose member +// was built for spir64_gen with no device specified, using the new offloading +// model. The bmg wrapper object must also define the generic code's +// __sycl_registerlib_ in addition to its own. + +// REQUIRES: ocloc, target-spir, linux, arch-intel_gpu_bmg_g21 + +// DEFINE: %{ond} = --offload-new-driver -Wno-error=unused-command-line-argument + +// Notice no device is specified in the below compile command. +// RUN: %clangxx -fsycl %{ond} -fsycl-targets=spir64_gen -c %S/Inputs/static-lib-registerlib-gen-generic.cpp -o %t.gen.o +// RUN: rm -f %t.a +// RUN: llvm-ar crv %t.a %t.gen.o + +// RUN: %clangxx -fsycl %{ond} -fsycl-targets=intel_gpu_bmg_g21 -c %s -o %t.main.o +// RUN: %clangxx -fsycl %{ond} -fsycl-targets=intel_gpu_bmg_g21 %t.main.o %t.a -o %t.exe +// RUN: %{run} %t.exe + +#include +#include + +#include +#include + +#include "Inputs/static-lib-registerlib.hpp" + +int main() { + constexpr std::size_t N = 8; + sycl::queue q; + std::size_t *p = sycl::malloc_shared(N, q); + q.parallel_for(sycl::range<1>{N}, [=](sycl::id<1> idx) { + p[idx] = idx[0]; + }).wait(); + std::size_t out[N] = {}; + run_gen_generic(out, N); + + bool Ok = true; + for (std::size_t I = 0; I < N; ++I) { + if (p[I] != I || out[I] != I + GenGenericOffset) { + std::cout << "fail: p[" << I << "] == " << p[I] << ", out[" << I + << "] == " << out[I] << "\n"; + Ok = false; + } + } + sycl::free(p, q); + if (Ok) + std::cout << "pass\n"; + return Ok ? 0 : 1; +} From e8e09e4256c682ba7811a9b592b57882762a0852 Mon Sep 17 00:00:00 2001 From: Nick Sarnie Date: Wed, 14 Oct 2026 08:47:17 -0700 Subject: [PATCH 10/19] dedup Signed-off-by: Nick Sarnie --- .../ClangLinkerWrapper.cpp | 68 +++++++++++-------- 1 file changed, 40 insertions(+), 28 deletions(-) diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 6b2205902687a..797dc4bd97925 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -3386,10 +3386,12 @@ Error emitExtractCommands( return Error::success(); } -/// See comments in collectSYCLRegisterLibSymbols. -static Error -collectSYCLRegisterLibSymbolsFromObject(MemoryBufferRef Buffer, - ArrayRef Binaries) { +/// Extract the offload binaries embedded in the single host object \p Buffer +/// into \p Binaries. If any were found, also record the object's undefined +/// __sycl_registerlib_ references in SYCLRegisterLibFuncMap +/// for the triple and arch of each of those binaries. +static Error extractOffloadBinariesAndRegisterLibSymbolsFromObject( + MemoryBufferRef Buffer, llvm::SmallVectorImpl &Binaries) { LLVMContext Context; static constexpr StringRef SYCLRegisterLibPrefix = "__sycl_registerlib_"; llvm::file_magic Magic = identify_magic(Buffer.getBuffer()); @@ -3397,11 +3399,22 @@ collectSYCLRegisterLibSymbolsFromObject(MemoryBufferRef Buffer, SymbolicFile::createSymbolicFile(Buffer, Magic, &Context, /*InitContent=*/Magic != llvm::file_magic::bitcode); + size_t OldNumBinaries = Binaries.size(); + if (Error Err = extractOffloadBinaries(Buffer, Binaries)) { + // The SYCL pipeline embeds raw bitcode in .llvm.offloading sections + // which is not in OffloadBinary format. Consume the parse error and + // fall through to the SYCL bundled objects extraction path. + consumeError(std::move(Err)); + } if (!ObjOrErr) { // Not something we can read symbols from; ignore it. consumeError(ObjOrErr.takeError()); return Error::success(); } + // An object without device code has no device image to define its + // references. + if (OldNumBinaries == Binaries.size()) + return Error::success(); SmallSet FuncNames; for (const BasicSymbolRef &Sym : (*ObjOrErr)->symbols()) { Expected FlagsOrErr = Sym.getFlags(); @@ -3417,11 +3430,8 @@ collectSYCLRegisterLibSymbolsFromObject(MemoryBufferRef Buffer, FuncNames.insert(Name); } if (!FuncNames.empty()) { - for (const auto &Binary : Binaries) { - // Only use the binaries embedded in this object. - if (Binary.getBinary()->getMemoryBufferRef().getBufferIdentifier() != - Buffer.getBufferIdentifier()) - continue; + for (size_t Idx = OldNumBinaries; Idx < Binaries.size(); Idx++) { + const OffloadFile &Binary = Binaries[Idx]; std::string Triple = Binary.getBinary()->getTriple().str(); std::string Arch = Binary.getBinary()->getArch().str(); if (Arch == "generic") @@ -3433,20 +3443,20 @@ collectSYCLRegisterLibSymbolsFromObject(MemoryBufferRef Buffer, return Error::success(); } -/// Scan the host input in \p Buffer for -/// undefined symbols named __sycl_registerlib_ and record -/// them in the global SYCLRegisterLibFuncMap. Each such reference is satisfied -/// by defining the symbol in the generated SYCL device-image wrapper object. -/// The references may live in bare object inputs or in members of a static -/// archive. -/// \p Binaries are the offload binaries already extracted from \p Buffer. -static Error collectSYCLRegisterLibSymbols(MemoryBufferRef Buffer, - ArrayRef Binaries) { +/// Extract the offload binaries embedded in the host input \p Buffer into +/// \p Binaries, like extractOffloadBinaries. For each host object, or each +/// member of a static archive, that carries device code, also record its +/// undefined __sycl_registerlib_ references in the global +/// SYCLRegisterLibFuncMap. Each such reference is satisfied by defining the +/// symbol in the generated SYCL device-image wrapper object. +static Error extractOffloadBinariesAndRegisterLibSymbols( + MemoryBufferRef Buffer, llvm::SmallVectorImpl &Binaries) { switch (identify_magic(Buffer.getBuffer())) { case file_magic::elf_relocatable: case file_magic::coff_object: case file_magic::bitcode: - if (Error Err = collectSYCLRegisterLibSymbolsFromObject(Buffer, Binaries)) + if (Error Err = extractOffloadBinariesAndRegisterLibSymbolsFromObject( + Buffer, Binaries)) return Err; break; case file_magic::archive: { @@ -3459,8 +3469,8 @@ static Error collectSYCLRegisterLibSymbols(MemoryBufferRef Buffer, Expected ChildBufOrErr = Child.getMemoryBufferRef(); if (!ChildBufOrErr) return ChildBufOrErr.takeError(); - if (Error E = - collectSYCLRegisterLibSymbolsFromObject(*ChildBufOrErr, Binaries)) + if (Error E = extractOffloadBinariesAndRegisterLibSymbolsFromObject( + *ChildBufOrErr, Binaries)) return E; } if (Err) @@ -3468,6 +3478,13 @@ static Error collectSYCLRegisterLibSymbols(MemoryBufferRef Buffer, break; } default: + // No host symbols to scan, only extract the device code. + if (Error Err = extractOffloadBinaries(Buffer, Binaries)) { + // The SYCL pipeline embeds raw bitcode in .llvm.offloading sections + // which is not in OffloadBinary format. Consume the parse error and + // fall through to the SYCL bundled objects extraction path. + consumeError(std::move(Err)); + } break; } return Error::success(); @@ -3551,13 +3568,8 @@ getDeviceInput(const ArgList &Args) { continue; SmallVector Binaries; size_t OldSize = Binaries.size(); - if (Error Err = extractOffloadBinaries(Buffer, Binaries)) { - // The SYCL pipeline embeds raw bitcode in .llvm.offloading sections - // which is not in OffloadBinary format. Consume the parse error and - // fall through to the SYCL bundled objects extraction path. - consumeError(std::move(Err)); - } - if (Error Err = collectSYCLRegisterLibSymbols(Buffer, Binaries)) + if (Error Err = + extractOffloadBinariesAndRegisterLibSymbols(Buffer, Binaries)) return std::move(Err); if (Binaries.size() == OldSize) { if (Error Err = sycl::extractBundledObjects(*Filename, Args, Binaries)) From 69e17a7dea3976841db83396aa501654b217a998 Mon Sep 17 00:00:00 2001 From: Nick Sarnie Date: Wed, 14 Oct 2026 10:40:27 -0700 Subject: [PATCH 11/19] fix generic case and add test Signed-off-by: Nick Sarnie Co-Authored-By: Claude Opus 5.5 --- .../ClangLinkerWrapper.cpp | 40 ++++++------------- .../static-lib-gen-generic-arch.cpp | 18 +++++++++ 2 files changed, 30 insertions(+), 28 deletions(-) diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 797dc4bd97925..e3f3ebee8fd51 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -1137,18 +1137,20 @@ runAOTCompileIntelCPU(StringRef InputFile, const ArgList &Args, /// \p Args encompasses all arguments required for linking and wrapping device /// code. static void defineRegisterLibFuncs(Module &M, const ArgList &Args) { - const llvm::Triple Triple(Args.getLastArgValue(OPT_triple_EQ)); - const std::string Arch(Args.getLastArgValue(OPT_arch_EQ)); - bool IsGPUAOT = Triple.isSPIROrSPIRV() && - Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen; + // getLinkerArgs passes an empty arch for device images with no arch + // specified, but we need it to be "generic" for the areTargetsCompatible + // call. + StringRef Arch = Args.getLastArgValue(OPT_arch_EQ); + OffloadFile::TargetID Target(Args.getLastArgValue(OPT_triple_EQ), + Arch.empty() ? "generic" : Arch); // We need to keep track of the triple and arch because we may not use all // device images passed to the tool and we don't want to force unused objects - // to be linked in. - auto It = SYCLRegisterLibFuncMap.find(std::make_pair(Triple.str(), Arch)); - if (It == SYCLRegisterLibFuncMap.end()) - return; - auto DefineFuncs = [&](auto &&Range) { - for (const std::string &Name : Range) { + // to be linked in. Define the symbols for every device image that can be + // linked into this target, using the same check as getDeviceInput. + for (const auto &[Key, Names] : SYCLRegisterLibFuncMap) { + if (!object::areTargetsCompatible(Key, Target)) + continue; + for (const std::string &Name : Names) { if (M.getFunction(Name)) continue; llvm::FunctionType *FTy = llvm::FunctionType::get( @@ -1165,22 +1167,6 @@ static void defineRegisterLibFuncs(Module &M, const ArgList &Args) { Fn->setComdat(C); } } - }; - // There is a special case when the current target being linked in spir64_gen - // with non-empty arch. We need to consider the cases where there are other - // device objects with spir64_gen triple with empty (generic) arch because - // this is considered a compatible target with spir64_gen non-empty arch and - // will be linked together with it. In this case, we need to also define - // symbols from the empty-arch device objects. - auto GenericArchIt = - SYCLRegisterLibFuncMap.find(std::make_pair(Triple.str(), std::string())); - if (!IsGPUAOT || GenericArchIt == SYCLRegisterLibFuncMap.end() || - GenericArchIt->second.empty()) { - DefineFuncs(It->second); - } else { - // Just iterate over both sets, duplcates are handled in DefineFuncs. - DefineFuncs( - llvm::concat(It->second, GenericArchIt->second)); } } @@ -3434,8 +3420,6 @@ static Error extractOffloadBinariesAndRegisterLibSymbolsFromObject( const OffloadFile &Binary = Binaries[Idx]; std::string Triple = Binary.getBinary()->getTriple().str(); std::string Arch = Binary.getBinary()->getArch().str(); - if (Arch == "generic") - Arch = ""; SYCLRegisterLibFuncMap[{Triple, Arch}].insert(FuncNames.begin(), FuncNames.end()); } diff --git a/sycl/test-e2e/NewOffloadDriver/static-lib-gen-generic-arch.cpp b/sycl/test-e2e/NewOffloadDriver/static-lib-gen-generic-arch.cpp index de4ec8f962ff6..09bb0f166f742 100644 --- a/sycl/test-e2e/NewOffloadDriver/static-lib-gen-generic-arch.cpp +++ b/sycl/test-e2e/NewOffloadDriver/static-lib-gen-generic-arch.cpp @@ -23,6 +23,24 @@ // RUN: %clangxx -fsycl %{ond} -fsycl-targets=intel_gpu_bmg_g21 %t.main.o %t.a -o %t.exe // RUN: %{run} %t.exe +// Also test the opposite case. + +// RUN: %clangxx -fsycl %{ond} -fsycl-targets=intel_gpu_bmg_g21 -c %S/Inputs/static-lib-registerlib-gen-generic.cpp -o %t.opp.gen.o +// RUN: rm -f %t.opp.a +// RUN: llvm-ar crv %t.opp.a %t.opp.gen.o + +// Notice no device is specified in the below compile command. +// RUN: %clangxx -fsycl %{ond} -fsycl-targets=spir64_gen -c %s -o %t.opp.main.o + +// The device is only given at link time, so the object file's device code +// still has no arch and the bmg archive member is linked together with it. +// The device must be given with -Xsycl-target-backend: the device code is +// linked for the arch recorded in the object file, which is generic, and +// -fsycl-targets=intel_gpu_bmg_g21 at link time would not pass a device to +// ocloc for it. +// RUN: %clangxx -fsycl %{ond} -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device bmg-g21" %t.opp.main.o %t.opp.a -o %t.opp.exe +// RUN: %{run} %t.opp.exe + #include #include From 3250e52e14b88a56475ec8173ac44f1b544a2f44 Mon Sep 17 00:00:00 2001 From: Nick Sarnie Date: Wed, 14 Oct 2026 12:32:17 -0700 Subject: [PATCH 12/19] return early and remove magic arg Signed-off-by: Nick Sarnie --- .../ClangLinkerWrapper.cpp | 18 +++++++----------- 1 file changed, 7 insertions(+), 11 deletions(-) diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index e3f3ebee8fd51..15a6188656b05 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -3378,13 +3378,6 @@ Error emitExtractCommands( /// for the triple and arch of each of those binaries. static Error extractOffloadBinariesAndRegisterLibSymbolsFromObject( MemoryBufferRef Buffer, llvm::SmallVectorImpl &Binaries) { - LLVMContext Context; - static constexpr StringRef SYCLRegisterLibPrefix = "__sycl_registerlib_"; - llvm::file_magic Magic = identify_magic(Buffer.getBuffer()); - Expected> ObjOrErr = - SymbolicFile::createSymbolicFile(Buffer, Magic, &Context, - /*InitContent=*/Magic != - llvm::file_magic::bitcode); size_t OldNumBinaries = Binaries.size(); if (Error Err = extractOffloadBinaries(Buffer, Binaries)) { // The SYCL pipeline embeds raw bitcode in .llvm.offloading sections @@ -3392,15 +3385,18 @@ static Error extractOffloadBinariesAndRegisterLibSymbolsFromObject( // fall through to the SYCL bundled objects extraction path. consumeError(std::move(Err)); } + // No need to process the symbols if the object contains no device code. + if (OldNumBinaries == Binaries.size()) + return Error::success(); + LLVMContext Context; + static constexpr StringRef SYCLRegisterLibPrefix = "__sycl_registerlib_"; + Expected> ObjOrErr = + SymbolicFile::createSymbolicFile(Buffer, file_magic::unknown, &Context); if (!ObjOrErr) { // Not something we can read symbols from; ignore it. consumeError(ObjOrErr.takeError()); return Error::success(); } - // An object without device code has no device image to define its - // references. - if (OldNumBinaries == Binaries.size()) - return Error::success(); SmallSet FuncNames; for (const BasicSymbolRef &Sym : (*ObjOrErr)->symbols()) { Expected FlagsOrErr = Sym.getFlags(); From 9b3b23f0eca9e9601ce778ff026e5b5e5ddd786b Mon Sep 17 00:00:00 2001 From: Nick Sarnie Date: Wed, 14 Oct 2026 13:04:20 -0700 Subject: [PATCH 13/19] add lit test for no device code Signed-off-by: Nick Sarnie --- clang/test/CodeGenSYCL/registerlib-host-no-device-code.cpp | 6 ++++++ 1 file changed, 6 insertions(+) create mode 100644 clang/test/CodeGenSYCL/registerlib-host-no-device-code.cpp diff --git a/clang/test/CodeGenSYCL/registerlib-host-no-device-code.cpp b/clang/test/CodeGenSYCL/registerlib-host-no-device-code.cpp new file mode 100644 index 0000000000000..c88be79405685 --- /dev/null +++ b/clang/test/CodeGenSYCL/registerlib-host-no-device-code.cpp @@ -0,0 +1,6 @@ +// Verify that we don't emit __sycl_registerlib_ nor __sycl_registerlib_ctor if -fembed-offload-object is not passed, as +// there is no device code in that case. + +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsycl-is-host --offload-new-driver -fsycl-unique-prefix=uid42 -emit-llvm %s -o - | FileCheck %s + +// CHECK-NOT: __sycl_ From 343d4513d4cb9c8de400709ebc2c093d2b3dc56a Mon Sep 17 00:00:00 2001 From: Nick Sarnie Date: Wed, 14 Oct 2026 13:14:35 -0700 Subject: [PATCH 14/19] address nits Signed-off-by: Nick Sarnie --- clang/lib/CodeGen/CodeGenModule.cpp | 4 ++-- .../clang-linker-wrapper/ClangLinkerWrapper.cpp | 16 ++++++++-------- ...tic-lib-multiple-objects-multiple-targets.cpp | 2 +- 3 files changed, 11 insertions(+), 11 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 4d76d0fadd363..b413458b13cbe 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1293,8 +1293,8 @@ void CodeGenModule::Release() { AddGlobalCtor(SYCLCtorFunction, /*Priority=*/101); } else if (LangOpts.OffloadingNewDriver && !CodeGenOpts.OffloadObjects.empty()) { - // Only add registration functions if there has been a device compile (so - // OffloadObjects is not empty). + // Only emit the registerlib reference if there has been a device compile + // (so OffloadObjects is not empty). AddGlobalCtor(createSYCLRegisterLibFunc()); } } diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 15a6188656b05..90ce564abd5c5 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -3411,14 +3411,14 @@ static Error extractOffloadBinariesAndRegisterLibSymbolsFromObject( if (*FlagsOrErr & SymbolRef::SF_Undefined) FuncNames.insert(Name); } - if (!FuncNames.empty()) { - for (size_t Idx = OldNumBinaries; Idx < Binaries.size(); Idx++) { - const OffloadFile &Binary = Binaries[Idx]; - std::string Triple = Binary.getBinary()->getTriple().str(); - std::string Arch = Binary.getBinary()->getArch().str(); - SYCLRegisterLibFuncMap[{Triple, Arch}].insert(FuncNames.begin(), - FuncNames.end()); - } + if (FuncNames.empty()) + return Error::success(); + for (size_t Idx = OldNumBinaries; Idx < Binaries.size(); Idx++) { + const OffloadFile &Binary = Binaries[Idx]; + std::string Triple = Binary.getBinary()->getTriple().str(); + std::string Arch = Binary.getBinary()->getArch().str(); + SYCLRegisterLibFuncMap[{Triple, Arch}].insert(FuncNames.begin(), + FuncNames.end()); } return Error::success(); } diff --git a/sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets.cpp b/sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets.cpp index 87c8cdfdc4674..1b6aef43b17ba 100644 --- a/sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets.cpp +++ b/sycl/test-e2e/NewOffloadDriver/static-lib-multiple-objects-multiple-targets.cpp @@ -9,7 +9,7 @@ // the new offloading model and gcc for the final host link. The archive holds // three translation units, each pre-linked with -fsycl-link for a single // distinct target (spir64 / intel_gpu_dg2 / intel_gpu_bmg_g21), so each -// contributes its own __sycl_registerlib_. +// contributes its own __sycl_registerlib_. // // The Windows counterpart is // static-lib-multiple-objects-multiple-targets-win.cpp. From 040ede7dc4e9732b92bf69ae251280571b4eb3c7 Mon Sep 17 00:00:00 2001 From: Nick Sarnie Date: Wed, 14 Oct 2026 13:21:46 -0700 Subject: [PATCH 15/19] fix comment Signed-off-by: Nick Sarnie --- clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp | 8 +++----- 1 file changed, 3 insertions(+), 5 deletions(-) diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 90ce564abd5c5..5a0f40d5718d7 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -3424,11 +3424,10 @@ static Error extractOffloadBinariesAndRegisterLibSymbolsFromObject( } /// Extract the offload binaries embedded in the host input \p Buffer into -/// \p Binaries, like extractOffloadBinaries. For each host object, or each -/// member of a static archive, that carries device code, also record its +/// \p Binaries. For each host object, or each +/// member of a static archive, that contains device code, also record its /// undefined __sycl_registerlib_ references in the global -/// SYCLRegisterLibFuncMap. Each such reference is satisfied by defining the -/// symbol in the generated SYCL device-image wrapper object. +/// SYCLRegisterLibFuncMap. static Error extractOffloadBinariesAndRegisterLibSymbols( MemoryBufferRef Buffer, llvm::SmallVectorImpl &Binaries) { switch (identify_magic(Buffer.getBuffer())) { @@ -3458,7 +3457,6 @@ static Error extractOffloadBinariesAndRegisterLibSymbols( break; } default: - // No host symbols to scan, only extract the device code. if (Error Err = extractOffloadBinaries(Buffer, Binaries)) { // The SYCL pipeline embeds raw bitcode in .llvm.offloading sections // which is not in OffloadBinary format. Consume the parse error and From 6f020362710e7f1341ff316f8c13bf7bc742eca5 Mon Sep 17 00:00:00 2001 From: Nick Sarnie Date: Tue, 20 Oct 2026 10:23:59 -0700 Subject: [PATCH 16/19] fix nits Signed-off-by: Nick Sarnie --- clang/test/CodeGenSYCL/registerlib-host-no-device-code.cpp | 2 +- clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/test/CodeGenSYCL/registerlib-host-no-device-code.cpp b/clang/test/CodeGenSYCL/registerlib-host-no-device-code.cpp index c88be79405685..1dcfe81941c1a 100644 --- a/clang/test/CodeGenSYCL/registerlib-host-no-device-code.cpp +++ b/clang/test/CodeGenSYCL/registerlib-host-no-device-code.cpp @@ -1,4 +1,4 @@ -// Verify that we don't emit __sycl_registerlib_ nor __sycl_registerlib_ctor if -fembed-offload-object is not passed, as +// Verify that we don't emit __sycl_registerlib_ or __sycl_registerlib_ctor if -fembed-offload-object is not passed, as // there is no device code in that case. // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsycl-is-host --offload-new-driver -fsycl-unique-prefix=uid42 -emit-llvm %s -o - | FileCheck %s diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 5a0f40d5718d7..468cd9ce278bc 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -1130,7 +1130,7 @@ runAOTCompileIntelCPU(StringRef InputFile, const ArgList &Args, return *TempFileOrErr; } -/// Define __sycl_registerlib_ functions referenced by the +/// Define __sycl_registerlib_ functions referenced by the /// host compile to force all device code images to be linked in, which /// may not happen for static archives with -fsycl-link. /// \p M The module. From f31d8a9e10bbd49943bc00049b3ae3b3d0ef196e Mon Sep 17 00:00:00 2001 From: Nick Sarnie Date: Wed, 21 Oct 2026 13:36:08 -0700 Subject: [PATCH 17/19] remove spir64_gen no device test Signed-off-by: Nick Sarnie --- .../static-lib-gen-generic-arch.cpp | 74 ------------------- 1 file changed, 74 deletions(-) delete mode 100644 sycl/test-e2e/NewOffloadDriver/static-lib-gen-generic-arch.cpp diff --git a/sycl/test-e2e/NewOffloadDriver/static-lib-gen-generic-arch.cpp b/sycl/test-e2e/NewOffloadDriver/static-lib-gen-generic-arch.cpp deleted file mode 100644 index 09bb0f166f742..0000000000000 --- a/sycl/test-e2e/NewOffloadDriver/static-lib-gen-generic-arch.cpp +++ /dev/null @@ -1,74 +0,0 @@ -//==- static-lib-gen-generic-arch.cpp -------------------------------------==// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// -// Link an object built for bmg against a static library whose member -// was built for spir64_gen with no device specified, using the new offloading -// model. The bmg wrapper object must also define the generic code's -// __sycl_registerlib_ in addition to its own. - -// REQUIRES: ocloc, target-spir, linux, arch-intel_gpu_bmg_g21 - -// DEFINE: %{ond} = --offload-new-driver -Wno-error=unused-command-line-argument - -// Notice no device is specified in the below compile command. -// RUN: %clangxx -fsycl %{ond} -fsycl-targets=spir64_gen -c %S/Inputs/static-lib-registerlib-gen-generic.cpp -o %t.gen.o -// RUN: rm -f %t.a -// RUN: llvm-ar crv %t.a %t.gen.o - -// RUN: %clangxx -fsycl %{ond} -fsycl-targets=intel_gpu_bmg_g21 -c %s -o %t.main.o -// RUN: %clangxx -fsycl %{ond} -fsycl-targets=intel_gpu_bmg_g21 %t.main.o %t.a -o %t.exe -// RUN: %{run} %t.exe - -// Also test the opposite case. - -// RUN: %clangxx -fsycl %{ond} -fsycl-targets=intel_gpu_bmg_g21 -c %S/Inputs/static-lib-registerlib-gen-generic.cpp -o %t.opp.gen.o -// RUN: rm -f %t.opp.a -// RUN: llvm-ar crv %t.opp.a %t.opp.gen.o - -// Notice no device is specified in the below compile command. -// RUN: %clangxx -fsycl %{ond} -fsycl-targets=spir64_gen -c %s -o %t.opp.main.o - -// The device is only given at link time, so the object file's device code -// still has no arch and the bmg archive member is linked together with it. -// The device must be given with -Xsycl-target-backend: the device code is -// linked for the arch recorded in the object file, which is generic, and -// -fsycl-targets=intel_gpu_bmg_g21 at link time would not pass a device to -// ocloc for it. -// RUN: %clangxx -fsycl %{ond} -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device bmg-g21" %t.opp.main.o %t.opp.a -o %t.opp.exe -// RUN: %{run} %t.opp.exe - -#include -#include - -#include -#include - -#include "Inputs/static-lib-registerlib.hpp" - -int main() { - constexpr std::size_t N = 8; - sycl::queue q; - std::size_t *p = sycl::malloc_shared(N, q); - q.parallel_for(sycl::range<1>{N}, [=](sycl::id<1> idx) { - p[idx] = idx[0]; - }).wait(); - std::size_t out[N] = {}; - run_gen_generic(out, N); - - bool Ok = true; - for (std::size_t I = 0; I < N; ++I) { - if (p[I] != I || out[I] != I + GenGenericOffset) { - std::cout << "fail: p[" << I << "] == " << p[I] << ", out[" << I - << "] == " << out[I] << "\n"; - Ok = false; - } - } - sycl::free(p, q); - if (Ok) - std::cout << "pass\n"; - return Ok ? 0 : 1; -} From 5513605055905c940e6c0af7e69acfd028f438cb Mon Sep 17 00:00:00 2001 From: Nick Sarnie Date: Fri, 23 Oct 2026 07:21:50 -0700 Subject: [PATCH 18/19] address tom's feedback and fix test bug Signed-off-by: Nick Sarnie --- clang/lib/CodeGen/CodeGenModule.cpp | 1 + clang/test/CodeGenSYCL/registerlib-host.cpp | 2 +- .../clang-linker-wrapper/linker-wrapper-sycl-registerlib.cpp | 2 ++ 3 files changed, 4 insertions(+), 1 deletion(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 98b04bfb6b770..3eb1ababfceba 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1316,6 +1316,7 @@ void CodeGenModule::Release() { // A static initializer may launch a kernel, so the device binary has to // be registered before any of them run, hence a priority. AddGlobalCtor(SYCLCtorFunction, /*Priority=*/101); + assert(CodeGenOpts.OffloadObjects.empty() && "Unexpected offload object"); } else if (LangOpts.OffloadingNewDriver && !CodeGenOpts.OffloadObjects.empty()) { // Only emit the registerlib reference if there has been a device compile diff --git a/clang/test/CodeGenSYCL/registerlib-host.cpp b/clang/test/CodeGenSYCL/registerlib-host.cpp index aca1042dc60bf..53e9f0456abc4 100644 --- a/clang/test/CodeGenSYCL/registerlib-host.cpp +++ b/clang/test/CodeGenSYCL/registerlib-host.cpp @@ -12,7 +12,7 @@ // Nor when the module embeds its device binary (-fno-sycl-rdc), as it then // registers the binary itself. // RUN: echo -n 'FAKE_SYCL_DEVICE_IMAGE' > %t.bin -// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsycl-is-host --offload-new-driver -foffload-include-binary %t.bin -fsycl-unique-prefix=uid42 -emit-llvm -fembed-offload-object=%t.o %s -o - | FileCheck %s --check-prefix=EMBED --implicit-check-not=__sycl_registerlib_ +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsycl-is-host --offload-new-driver -foffload-include-binary %t.bin -fsycl-unique-prefix=uid42 -emit-llvm %s -o - | FileCheck %s --check-prefix=EMBED --implicit-check-not=__sycl_registerlib_ // CHECK: @llvm.global_ctors = appending global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 65535, ptr @__sycl_registerlib_ctor, ptr null }] // CHECK: declare void @__sycl_registerlib_uid42() diff --git a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl-registerlib.cpp b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl-registerlib.cpp index a2ce083b4aa74..41790fe7a75cb 100644 --- a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl-registerlib.cpp +++ b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl-registerlib.cpp @@ -31,6 +31,8 @@ __attribute__((sycl_kernel)) void kernel(const Func &func) { #ifdef MAIN extern "C" { +// Define stub versions of the SYCL runtime image registration functions that are called by +// the wrapped host code so we don't need to link against the real SYCL runtime. void __sycl_register_lib(void *) {} void __sycl_unregister_lib(void *) {} } From d176840cdcdd7e0e4933da4e88abbdc29236cc80 Mon Sep 17 00:00:00 2001 From: Nick Sarnie Date: Mon, 26 Oct 2026 07:05:25 -0700 Subject: [PATCH 19/19] add assert for empty unique ident Signed-off-by: Nick Sarnie --- clang/lib/CodeGen/CodeGenSYCL.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/clang/lib/CodeGen/CodeGenSYCL.cpp b/clang/lib/CodeGen/CodeGenSYCL.cpp index 23ac73a64c47f..3aecb4d85cd5c 100644 --- a/clang/lib/CodeGen/CodeGenSYCL.cpp +++ b/clang/lib/CodeGen/CodeGenSYCL.cpp @@ -135,6 +135,8 @@ llvm::Function *CodeGenModule::embedSYCLDeviceBinary() { } llvm::Function *CodeGenModule::createSYCLRegisterLibFunc() { + assert(!LangOpts.SYCLUniquePrefix.empty() && + "Expected non-empty SYCL unique prefix"); std::string RegisterFuncName = "__sycl_registerlib_" + LangOpts.SYCLUniquePrefix; llvm::FunctionType *RegisterFTy =