diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def
index ff116dd9bf368..8e3d85e3c7f12 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 b1559e53f410b..a757a01f323a2 100644
--- a/clang/include/clang/Options/Options.td
+++ b/clang/include/clang/Options/Options.td
@@ -1386,6 +1386,7 @@ def no_offload_arch_EQ
 
 def offload_new_driver : Flag<["--"], "offload-new-driver">,
   Visibility<[ClangOption, CC1Option]>, Group<f_Group>,
+  MarshallingInfoFlag<LangOpts<"OffloadingNewDriver">>,
   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<f_Group>,
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index b22cb34f92115..3eb1ababfceba 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -1310,11 +1310,19 @@ 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);
+      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
+      // (so OffloadObjects is not empty).
+      AddGlobalCtor(createSYCLRegisterLibFunc());
+    }
   }
   if (OpenMPRuntime) {
     OpenMPRuntime->createOffloadEntriesAndInfoMetadata();
diff --git a/clang/lib/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h
index 3016da93d9872..333dfdef27b66 100644
--- a/clang/lib/CodeGen/CodeGenModule.h
+++ b/clang/lib/CodeGen/CodeGenModule.h
@@ -2240,6 +2240,17 @@ class CodeGenModule : public CodeGenTypeCache {
   /// if the binary could not be read.
   llvm::Function *embedSYCLDeviceBinary();
 
+  /// Emit a per-module __sycl_registerlib_<unique_identifier> 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
+  /// unique 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 665de1ace5c6d..de135e8d1cd8f 100644
--- a/clang/lib/CodeGen/CodeGenSYCL.cpp
+++ b/clang/lib/CodeGen/CodeGenSYCL.cpp
@@ -132,3 +132,24 @@ llvm::Function *CodeGenModule::embedSYCLDeviceBinary() {
   }
   return RegistrationFunc;
 }
+
+llvm::Function *CodeGenModule::createSYCLRegisterLibFunc() {
+  assert(!LangOpts.SYCLUniquePrefix.empty() &&
+         "Expected non-empty SYCL unique prefix");
+  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 8d3c99c7ebe54..cc4735229c6dd 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -8052,11 +8052,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-no-device-code.cpp b/clang/test/CodeGenSYCL/registerlib-host-no-device-code.cpp
new file mode 100644
index 0000000000000..1dcfe81941c1a
--- /dev/null
+++ b/clang/test/CodeGenSYCL/registerlib-host-no-device-code.cpp
@@ -0,0 +1,6 @@
+// Verify that we don't emit __sycl_registerlib_<unique identifier> 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
+
+// CHECK-NOT: __sycl_
diff --git a/clang/test/CodeGenSYCL/registerlib-host.cpp b/clang/test/CodeGenSYCL/registerlib-host.cpp
new file mode 100644
index 0000000000000..53e9f0456abc4
--- /dev/null
+++ b/clang/test/CodeGenSYCL/registerlib-host.cpp
@@ -0,0 +1,28 @@
+// Check the per-translation-unit __sycl_registerlib_<id> 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: 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 -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_
+
+// 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 0178d0f709d9f..dd0615825e572 100644
--- a/clang/test/Driver/sycl-offload-new-driver.cpp
+++ b/clang/test/Driver/sycl-offload-new-driver.cpp
@@ -27,9 +27,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 -ccc-install-dir %{clang_install_dir} %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
@@ -96,7 +96,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..41790fe7a75cb
--- /dev/null
+++ b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl-registerlib.cpp
@@ -0,0 +1,47 @@
+// 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_<id> 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 <typename t, typename Func>
+__attribute__((sycl_kernel)) void kernel(const Func &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 *) {}
+}
+
+void lib_fn();
+int main() {
+  kernel<class main_kernel>([]() {});
+  lib_fn();
+}
+#else
+void FN() { kernel<class lib_kernel>([]() {}); }
+#endif
diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
index 8260d7f379d17..ae74aae900238 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"
@@ -158,6 +162,20 @@ static bool LinkingSYCLBINFiles = false;
 
 static SmallString<128> OffloadImageDumpDir;
 
+using TripleAndArch = std::pair<std::string, std::string>;
+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<std::string, 4>;
+static DenseMap<TripleAndArch, FuncNameSet, TripleAndArchInfo>
+    SYCLRegisterLibFuncMap;
+
 /// Whether or not to look through symlinks when resolving binaries.
 static bool CanonicalPrefixes = true;
 
@@ -1063,6 +1081,46 @@ runAOTCompileIntelCPU(StringRef InputFile, const ArgList &Args,
   return *TempFileOrErr;
 }
 
+/// Define __sycl_registerlib_<unique_identifier> 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.
+/// \p Args encompasses all arguments required for linking and wrapping device
+/// code.
+static void defineRegisterLibFuncs(Module &M, const ArgList &Args) {
+  // 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. 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(
+          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);
+      }
+    }
+  }
+}
+
 /// 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.
@@ -1267,6 +1325,8 @@ wrapSYCLBinariesFromFile(ArrayRef<module_split::SplitModule> 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;
@@ -3363,6 +3423,102 @@ Error emitExtractCommands(
   return Error::success();
 }
 
+/// 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_<unique_identifier> references in SYCLRegisterLibFuncMap
+/// for the triple and arch of each of those binaries.
+static Error extractOffloadBinariesAndRegisterLibSymbolsFromObject(
+    MemoryBufferRef Buffer, llvm::SmallVectorImpl<OffloadFile> &Binaries) {
+  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));
+  }
+  // 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<std::unique_ptr<SymbolicFile>> 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();
+  }
+  SmallSet<std::string, 4> FuncNames;
+  for (const BasicSymbolRef &Sym : (*ObjOrErr)->symbols()) {
+    Expected<uint32_t> FlagsOrErr = Sym.getFlags();
+    if (!FlagsOrErr)
+      return FlagsOrErr.takeError();
+    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(Name);
+  }
+  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();
+}
+
+/// Extract the offload binaries embedded in the host input \p Buffer into
+/// \p Binaries. For each host object, or each
+/// member of a static archive, that contains device code, also record its
+/// undefined __sycl_registerlib_<unique_identifier> references in the global
+/// SYCLRegisterLibFuncMap.
+static Error extractOffloadBinariesAndRegisterLibSymbols(
+    MemoryBufferRef Buffer, llvm::SmallVectorImpl<OffloadFile> &Binaries) {
+  switch (identify_magic(Buffer.getBuffer())) {
+  case file_magic::elf_relocatable:
+  case file_magic::coff_object:
+  case file_magic::bitcode:
+    if (Error Err = extractOffloadBinariesAndRegisterLibSymbolsFromObject(
+            Buffer, Binaries))
+      return Err;
+    break;
+  case file_magic::archive: {
+    Expected<std::unique_ptr<object::Archive>> ArOrErr =
+        object::Archive::create(Buffer);
+    if (!ArOrErr)
+      return ArOrErr.takeError();
+    Error Err = Error::success();
+    for (const object::Archive::Child &Child : (*ArOrErr)->children(Err)) {
+      Expected<MemoryBufferRef> ChildBufOrErr = Child.getMemoryBufferRef();
+      if (!ChildBufOrErr)
+        return ChildBufOrErr.takeError();
+      if (Error E = extractOffloadBinariesAndRegisterLibSymbolsFromObject(
+              *ChildBufOrErr, Binaries))
+        return E;
+    }
+    if (Err)
+      return Err;
+    break;
+  }
+  default:
+    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();
+}
+
 /// 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
@@ -3441,12 +3597,9 @@ getDeviceInput(const ArgList &Args) {
       continue;
     SmallVector<OffloadFile> 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 =
+            extractOffloadBinariesAndRegisterLibSymbols(Buffer, Binaries))
+      return std::move(Err);
     if (Binaries.size() == OldSize) {
       if (Error Err = sycl::extractBundledObjects(*Filename, Args, Binaries))
         return std::move(Err);
diff --git a/sycl/test-e2e/E2EExpr.py b/sycl/test-e2e/E2EExpr.py
index 6f2224c04630a..13e184089f169 100644
--- a/sycl/test-e2e/E2EExpr.py
+++ b/sycl/test-e2e/E2EExpr.py
@@ -59,6 +59,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/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 <sycl/detail/core.hpp>
+#include <sycl/usm.hpp>
+
+#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<std::size_t>(N, q);
+  q.parallel_for<class KernelBMG>(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 <sycl/detail/core.hpp>
+#include <sycl/usm.hpp>
+
+#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<std::size_t>(N, q);
+  q.parallel_for<class KernelDG2>(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-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 <sycl/detail/core.hpp>
+#include <sycl/usm.hpp>
+
+#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<std::size_t>(N, q);
+  q.parallel_for<class KernelGenGeneric>(
+       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-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 <sycl/detail/core.hpp>
+#include <sycl/usm.hpp>
+
+#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<std::size_t>(N, q);
+  q.parallel_for<class KernelJIT>(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..4c49935729947
--- /dev/null
+++ b/sycl/test-e2e/NewOffloadDriver/Inputs/static-lib-registerlib.hpp
@@ -0,0 +1,20 @@
+//===----------------------------------------------------------------------===//
+// 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 the static-lib-* registerlib tests.
+
+#pragma once
+
+#include <cstddef>
+
+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-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-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..1b6aef43b17ba
--- /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_<unique_identifier>.
+//
+// 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 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_<unique_identifier> 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
+
+// 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 <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/experimental/device_architecture.hpp>
+
+#include <cstddef>
+#include <iostream>
+
+#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<syclex::info::device::architecture>();
+  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;
+}
diff --git a/sycl/test-e2e/lit.cfg.py b/sycl/test-e2e/lit.cfg.py
index 55aaa7d3c644c..b00469338d03b 100644
--- a/sycl/test-e2e/lit.cfg.py
+++ b/sycl/test-e2e/lit.cfg.py
@@ -908,6 +908,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 = [