diff --git a/clang/lib/CodeGen/CMakeLists.txt b/clang/lib/CodeGen/CMakeLists.txt
index f2d256043571..f7252f299119 100644
--- a/clang/lib/CodeGen/CMakeLists.txt
+++ b/clang/lib/CodeGen/CMakeLists.txt
@@ -2,6 +2,7 @@ set(LLVM_LINK_COMPONENTS
   ABI
   AggressiveInstCombine
   Analysis
+  BinaryFormat
   BitReader
   BitWriter
   CodeGen
diff --git a/clang/lib/CodeGen/CodeGenSYCL.cpp b/clang/lib/CodeGen/CodeGenSYCL.cpp
index 793ad85b8c1f..665de1ace5c6 100644
--- a/clang/lib/CodeGen/CodeGenSYCL.cpp
+++ b/clang/lib/CodeGen/CodeGenSYCL.cpp
@@ -14,7 +14,7 @@
 #include "CodeGenFunction.h"
 #include "CodeGenModule.h"
 #include "clang/Basic/DiagnosticFrontend.h"
-#include "llvm/Bitcode/BitcodeReader.h"
+#include "llvm/BinaryFormat/Magic.h"
 #include "llvm/Frontend/Offloading/OffloadWrapper.h"
 #include "llvm/Support/MemoryBuffer.h"
 #include "llvm/Support/VirtualFileSystem.h"
@@ -112,14 +112,13 @@ llvm::Function *CodeGenModule::embedSYCLDeviceBinary() {
     return nullptr;
   }
   std::unique_ptr<llvm::MemoryBuffer> Buffer = std::move(BufferOrErr.get());
-  // A wrapper module in bitcode form already carries the registration
-  // descriptor the SYCL runtime expects. It is linked into the host module by
-  // loadLinkModules, so there is nothing to embed here. In SYCL CUDA
-  // compatibility mode the file is the CUDA fat binary, which CUDA registers.
-  if (getLangOpts().CUDA ||
-      llvm::isBitcode(
-          reinterpret_cast<const unsigned char *>(Buffer->getBufferStart()),
-          reinterpret_cast<const unsigned char *>(Buffer->getBufferEnd())))
+  // The SYCL runtime registers device images through a descriptor that also
+  // carries their properties, kernel entries and options. loadLinkModules
+  // wraps an offload binary container into such a descriptor, so there is
+  // nothing to embed here. In SYCL CUDA compatibility mode the file is the
+  // CUDA fat binary, which CUDA registers.
+  if (getLangOpts().CUDA || llvm::identify_magic(Buffer->getBuffer()) ==
+                                llvm::file_magic::offload_binary)
     return nullptr;
   llvm::Function *RegistrationFunc = nullptr;
   if (llvm::Error Err = llvm::offloading::wrapSYCLBinaries(
diff --git a/clang/lib/CodeGen/ModuleLinker.cpp b/clang/lib/CodeGen/ModuleLinker.cpp
index 0cfd0cc32b9f..f7f91d9319db 100644
--- a/clang/lib/CodeGen/ModuleLinker.cpp
+++ b/clang/lib/CodeGen/ModuleLinker.cpp
@@ -11,7 +11,9 @@
 #include "clang/Basic/CodeGenOptions.h"
 #include "clang/Basic/DiagnosticFrontend.h"
 #include "clang/Frontend/CompilerInstance.h"
+#include "llvm/BinaryFormat/Magic.h"
 #include "llvm/Bitcode/BitcodeReader.h"
+#include "llvm/Frontend/Offloading/SYCLOffloadWrapper.h"
 #include "llvm/IR/Module.h"
 
 using namespace clang;
@@ -47,39 +49,45 @@ bool clang::loadLinkModules(CompilerInstance &CI, llvm::LLVMContext &Ctx,
                            F.Internalize, F.LinkFlags});
   }
 
-  // For SYCL no-RDC, link the per-TU wrapper module produced by
-  // clang-linker-wrapper --emit-fatbin-only into the host module so the SYCL
-  // runtime finds the device image at program startup. A raw device binary is
-  // embedded by CodeGenModule::embedSYCLDeviceBinary instead.
-  if (CI.getLangOpts().SYCLIsHost &&
-      !CI.getCodeGenOpts().OffloadBinaryToEmbedFile.empty()) {
-    auto BCBuf = CI.getFileManager().getBufferForFile(
-        CI.getCodeGenOpts().OffloadBinaryToEmbedFile);
-    if (!BCBuf) {
+  // For SYCL no-RDC, wrap the device images of the fat binary produced by
+  // clang-linker-wrapper --emit-fatbin-only into a module that registers them
+  // with the SYCL runtime, and link it into the host module. Other binaries
+  // are embedded by CodeGenModule::embedSYCLDeviceBinary instead.
+  StringRef FatbinFile = CI.getCodeGenOpts().OffloadBinaryToEmbedFile;
+  if (CI.getLangOpts().SYCLIsHost && !CI.getLangOpts().CUDA &&
+      !FatbinFile.empty()) {
+    auto FatbinBuf = CI.getFileManager().getBufferForFile(FatbinFile);
+    if (!FatbinBuf) {
       CI.getDiagnostics().Report(diag::err_cannot_open_file)
-          << CI.getCodeGenOpts().OffloadBinaryToEmbedFile
-          << BCBuf.getError().message();
+          << FatbinFile << FatbinBuf.getError().message();
       LinkModules.clear();
       return true;
     }
-    if (!llvm::isBitcode(
-            reinterpret_cast<const unsigned char *>((*BCBuf)->getBufferStart()),
-            reinterpret_cast<const unsigned char *>((*BCBuf)->getBufferEnd())))
+    if (llvm::identify_magic((*FatbinBuf)->getBuffer()) !=
+        llvm::file_magic::offload_binary)
       return false;
-    llvm::Expected<std::unique_ptr<llvm::Module>> MOrErr =
-        llvm::parseBitcodeFile((*BCBuf)->getMemBufferRef(), Ctx);
-    if (!MOrErr) {
-      llvm::handleAllErrors(MOrErr.takeError(), [&](llvm::ErrorInfoBase &EIB) {
+
+    auto ReportError = [&](llvm::Error Err) {
+      llvm::handleAllErrors(std::move(Err), [&](llvm::ErrorInfoBase &EIB) {
         CI.getDiagnostics().Report(diag::err_fe_linking_module)
-            << CI.getCodeGenOpts().OffloadBinaryToEmbedFile << EIB.message();
+            << FatbinFile << EIB.message();
       });
       LinkModules.clear();
       return true;
-    }
-    // The wrapper .bc from clang-linker-wrapper has no data layout set.
-    // Stamp the host layout so linkModules doesn't warn about a mismatch.
-    (*MOrErr)->setDataLayout(CI.getTarget().getDataLayoutString());
-    LinkModules.push_back({std::move(*MOrErr),
+    };
+    auto ImagesOrErr =
+        llvm::offloading::readSYCLImages((*FatbinBuf)->getMemBufferRef());
+    if (!ImagesOrErr)
+      return ReportError(ImagesOrErr.takeError());
+
+    auto WrapperM =
+        std::make_unique<llvm::Module>("offload.wrapper.object", Ctx);
+    WrapperM->setTargetTriple(CI.getTarget().getTriple());
+    WrapperM->setDataLayout(CI.getTarget().getDataLayoutString());
+    if (llvm::Error Err =
+            llvm::offloading::wrapSYCLBinaries(*WrapperM, *ImagesOrErr))
+      return ReportError(std::move(Err));
+    LinkModules.push_back({std::move(WrapperM),
                            /*PropagateAttrs=*/false,
                            /*Internalize=*/false,
                            /*LinkFlags=*/0});
diff --git a/clang/test/CodeGen/sycl-no-rdc-embed-binary.cpp b/clang/test/CodeGen/sycl-no-rdc-embed-binary.cpp
index f43301be7f0c..79b1b12ed749 100644
--- a/clang/test/CodeGen/sycl-no-rdc-embed-binary.cpp
+++ b/clang/test/CodeGen/sycl-no-rdc-embed-binary.cpp
@@ -1,5 +1,5 @@
-// Tests for -foffload-include-binary with a SYCL wrapper module in bitcode form
-// (SYCL -fno-sycl-rdc compile-step device embedding via loadLinkModules).
+// Tests for -foffload-include-binary with a SYCL fat binary in offload binary
+// form (SYCL -fno-sycl-rdc compile-step device embedding via loadLinkModules).
 
 // Missing file: err_cannot_open_file
 // RUN: not %clang_cc1 -triple x86_64-unknown-linux-gnu -fsycl-is-host \
@@ -7,19 +7,42 @@
 // RUN:   | FileCheck --check-prefix=CHECK-NO-FILE %s
 // CHECK-NO-FILE: fatal error: cannot open file 'no-such-file.bc'
 
-// Wrapper module bitcode: linked into the host module as is, so its
-// registration constructor is kept and no raw device binary is embedded.
+// Offload binary: each image is wrapped into the descriptor the SYCL runtime
+// registers, together with its target, options and kernel entries. The raw
+// container itself is not embedded.
 // RUN: split-file %s %t
-// RUN: llvm-as %t/wrapper.ll -o %t/wrapper.bc
+// RUN: echo -n 'FAKE_SYCL_DEVICE_IMAGE' > %t/image.bin
+// RUN: llvm-offload-binary -o %t/fatbin.syclfb \
+// RUN:   --image=file=%t/image.bin,kind=sycl,triple=spir64-unknown-unknown,arch=,sycl-target=spir64,sycl-image-format=0,symbols=_ZTS1K,compile-opts=-g,link-opts=-cl-opt
 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsycl-is-host \
-// RUN:   -foffload-include-binary %t/wrapper.bc -emit-llvm %t/host.cpp -o - \
+// RUN:   -foffload-include-binary %t/fatbin.syclfb -emit-llvm %t/host.cpp -o - \
 // RUN:   | FileCheck --check-prefix=CHECK-LINK %s \
 // RUN:     --implicit-check-not='.sycl_fatbin' \
 // RUN:     --implicit-check-not='.sycl_offloading.binary'
-// CHECK-LINK: @llvm.global_ctors = {{.*}} ptr @sycl.descriptor_reg
+// CHECK-LINK-DAG: @.sycl_offloading.target.0 = internal unnamed_addr constant [7 x i8] c"spir64\00"
+// CHECK-LINK-DAG: @.sycl_offloading.opts.compile.0 = internal unnamed_addr constant [3 x i8] c"-g\00"
+// CHECK-LINK-DAG: @.sycl_offloading.opts.link.0 = internal unnamed_addr constant [8 x i8] c"-cl-opt\00"
+// CHECK-LINK-DAG: @.sycl_offloading.0.data = internal unnamed_addr constant [22 x i8] c"FAKE_SYCL_DEVICE_IMAGE"
+// CHECK-LINK-DAG: @__sycl_offload_entry_name = internal unnamed_addr constant [7 x i8] c"_ZTS1K\00"
+// CHECK-LINK-DAG: @llvm.global_ctors = {{.*}} ptr @sycl.descriptor_reg
 // CHECK-LINK: define internal void @sycl.descriptor_reg()
 // CHECK-LINK: call void @__sycl_register_lib(ptr @.sycl_offloading.descriptor)
 
+// Malformed runtime metadata is diagnosed.
+// RUN: llvm-offload-binary -o %t/no-format.syclfb \
+// RUN:   --image=file=%t/image.bin,kind=sycl,triple=spir64-unknown-unknown
+// RUN: not %clang_cc1 -triple x86_64-unknown-linux-gnu -fsycl-is-host \
+// RUN:   -foffload-include-binary %t/no-format.syclfb -emit-llvm %t/host.cpp \
+// RUN:   -o /dev/null 2>&1 | FileCheck --check-prefix=CHECK-NO-FORMAT %s
+// CHECK-NO-FORMAT: error: cannot link module '{{.*}}no-format.syclfb': invalid SYCL image format ''
+
+// RUN: llvm-offload-binary -o %t/openmp.syclfb \
+// RUN:   --image=file=%t/image.bin,kind=openmp,triple=spir64-unknown-unknown
+// RUN: not %clang_cc1 -triple x86_64-unknown-linux-gnu -fsycl-is-host \
+// RUN:   -foffload-include-binary %t/openmp.syclfb -emit-llvm %t/host.cpp \
+// RUN:   -o /dev/null 2>&1 | FileCheck --check-prefix=CHECK-NOT-SYCL %s
+// CHECK-NOT-SYCL: error: cannot link module '{{.*}}openmp.syclfb': offload binary entry 0 is not a SYCL device image
+
 // SYCL CUDA compatibility mode: the host compilation is also a CUDA host
 // compilation and the file is the CUDA fat binary. Only CUDA registers it; it
 // is neither linked nor embedded as a SYCL device image.
@@ -33,17 +56,6 @@
 // CHECK-SYCL-CUDA: section ".nv_fatbin"
 // CHECK-SYCL-CUDA: call ptr @__cudaRegisterFatBinary(
 
-//--- wrapper.ll
-@.sycl_offloading.descriptor = internal constant [1 x i8] zeroinitializer
-@llvm.global_ctors = appending global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 1, ptr @sycl.descriptor_reg, ptr null }]
-
-define internal void @sycl.descriptor_reg() {
-  call void @__sycl_register_lib(ptr @.sycl_offloading.descriptor)
-  ret void
-}
-
-declare void @__sycl_register_lib(ptr)
-
 //--- host.cpp
 void f() {}
 
diff --git a/clang/test/Driver/sycl-nordc-fatbin.cpp b/clang/test/Driver/sycl-nordc-fatbin.cpp
index 48f42a9a497b..9c07c9d76127 100644
--- a/clang/test/Driver/sycl-nordc-fatbin.cpp
+++ b/clang/test/Driver/sycl-nordc-fatbin.cpp
@@ -1,42 +1,38 @@
 // REQUIRES: spirv-registered-target, x86-registered-target
 
-// TODO: With the intel/llvm SYCL driver the .sycl_fatbin section holds LLVM
-// bitcode rather than an offload binary wrapping a finalized SPIR-V module.
-// Enable once the -fno-sycl-rdc flow matches upstream.
-// XFAIL: *
+// REQUIRES: spirv-to-ir-wrapper, sycl-post-link
 
-/// Verify which section carries the SYCL device binary in the host object.
-// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl -fno-gpu-rdc \
+/// Verify that no-RDC emits image address/size records in .tgtimg, while RDC
+/// embeds device code for linking in .llvm.offloading.
+// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver --no-offloadlib -fno-gpu-rdc \
 // RUN:   -c %s -o %t.nordc.o
 // RUN: llvm-readelf -S %t.nordc.o \
 // RUN:   | FileCheck -check-prefix=NORDC %s --implicit-check-not='.llvm.offloading'
 
-// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl -fgpu-rdc \
+/// -fsycl-link also forwards --sycl-device-link to the per-TU no-RDC job.
+/// Fatbin-only completion must skip the device-link copy branch.
+// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver --no-offloadlib -fno-gpu-rdc \
+// RUN:   -fsycl-link -c %s -o %t.nordc.devlink.o
+// RUN: llvm-readelf -S %t.nordc.devlink.o \
+// RUN:   | FileCheck -check-prefix=NORDC %s --implicit-check-not='.llvm.offloading'
+
+// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver --no-offloadlib -fgpu-rdc \
 // RUN:   -c %s -o %t.rdc.o
 // RUN: llvm-readelf -S %t.rdc.o \
-// RUN:   | FileCheck -check-prefix=RDC %s --implicit-check-not='.sycl_fatbin'
+// RUN:   | FileCheck -check-prefix=RDC %s --implicit-check-not='.tgtimg'
 
 /// The choice does not depend on the object format.
-// RUN: %clangxx --target=x86_64-pc-windows-msvc -fsycl -fno-gpu-rdc \
+// RUN: %clangxx --target=x86_64-pc-windows-msvc -fsycl --offload-new-driver --no-offloadlib -fno-gpu-rdc \
 // RUN:   -c %s -o %t.nordc.obj
 // RUN: llvm-readobj --sections %t.nordc.obj \
 // RUN:   | FileCheck -check-prefix=NORDC %s --implicit-check-not='.llvm.offloading'
 
-// RUN: %clangxx --target=x86_64-pc-windows-msvc -fsycl -fgpu-rdc \
+// RUN: %clangxx --target=x86_64-pc-windows-msvc -fsycl --offload-new-driver --no-offloadlib -fgpu-rdc \
 // RUN:   -c %s -o %t.rdc.obj
 // RUN: llvm-readobj --sections %t.rdc.obj \
-// RUN:   | FileCheck -check-prefix=RDC %s --implicit-check-not='.sycl_fatbin'
+// RUN:   | FileCheck -check-prefix=RDC %s --implicit-check-not='.tgtimg'
 
-// NORDC: .sycl_fatbin
+// NORDC: .tgtimg
 // RDC: .llvm.offloading
 
-/// The section holds an offload binary (magic 0x10FF10AD) whose image is a
-/// finalized SPIR-V module (magic 0x07230203), both shown little endian by the
-/// hex dump.
-// RUN: llvm-readelf --hex-dump=.sycl_fatbin %t.nordc.o \
-// RUN:   | FileCheck -check-prefix=BINARY %s
-
-// BINARY: 10ff10ad
-// BINARY: 03022307
-
 void f() {}
diff --git a/clang/test/Driver/sycl-nordc-wrapper-object.cpp b/clang/test/Driver/sycl-nordc-wrapper-object.cpp
index a98899baa5f5..92458837df29 100644
--- a/clang/test/Driver/sycl-nordc-wrapper-object.cpp
+++ b/clang/test/Driver/sycl-nordc-wrapper-object.cpp
@@ -1,39 +1,15 @@
-// REQUIRES: spirv-registered-target, x86-registered-target
-// REQUIRES: spirv-to-ir-wrapper, sycl-post-link
+// REQUIRES: sycl-post-link
 
-/// intel/llvm counterpart of LLORG's sycl-nordc-fatbin.cpp. In intel/llvm the
-/// non-RDC host object carries the wrapper module produced at compile time:
-/// the device image in ".tgtimg" plus the __sycl_register_lib registration,
-/// and no ".llvm.offloading", so the final link does not device-link it again.
-/// The RDC host object is the reverse.
+/// Verify that the no-RDC driver flow registers and unregisters the descriptor
+/// expected by the Intel SYCL runtime.
 // RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver \
-// RUN:   --no-offloadlib -fno-gpu-rdc -c %s -o %t.nordc.o
-// RUN: llvm-readelf -S %t.nordc.o \
-// RUN:   | FileCheck -check-prefix=NORDC %s --implicit-check-not='.llvm.offloading'
-// RUN: llvm-nm %t.nordc.o | FileCheck -check-prefix=REG %s
-
-// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver \
-// RUN:   --no-offloadlib -fgpu-rdc -c %s -o %t.rdc.o
-// RUN: llvm-readelf -S %t.rdc.o \
-// RUN:   | FileCheck -check-prefix=RDC %s --implicit-check-not='.tgtimg'
-// RUN: llvm-nm %t.rdc.o | FileCheck -check-prefix=NOREG %s
-
-/// The choice does not depend on the object format.
-// RUN: %clangxx --target=x86_64-pc-windows-msvc -fsycl --offload-new-driver \
-// RUN:   --no-offloadlib -fno-gpu-rdc -c %s -o %t.nordc.obj
-// RUN: llvm-readobj --sections %t.nordc.obj \
-// RUN:   | FileCheck -check-prefix=NORDC %s --implicit-check-not='.llvm.offloading'
-// RUN: llvm-nm %t.nordc.obj | FileCheck -check-prefix=REG %s
-
+// RUN:   --no-offloadlib -fno-gpu-rdc -S -emit-llvm %s -o - \
+// RUN:   | FileCheck %s
 // RUN: %clangxx --target=x86_64-pc-windows-msvc -fsycl --offload-new-driver \
-// RUN:   --no-offloadlib -fgpu-rdc -c %s -o %t.rdc.obj
-// RUN: llvm-readobj --sections %t.rdc.obj \
-// RUN:   | FileCheck -check-prefix=RDC %s --implicit-check-not='.tgtimg'
-// RUN: llvm-nm %t.rdc.obj | FileCheck -check-prefix=NOREG %s
+// RUN:   --no-offloadlib -fno-gpu-rdc -S -emit-llvm %s -o - \
+// RUN:   | FileCheck %s
 
-// NORDC: .tgtimg
-// RDC: .llvm.offloading
-// REG: U __sycl_register_lib
-// NOREG-NOT: __sycl_register_lib
+// CHECK: call void @__sycl_register_lib(ptr @.sycl_offloading.descriptor)
+// CHECK: call void @__sycl_unregister_lib(ptr @.sycl_offloading.descriptor)
 
 void f() {}
diff --git a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl-fatbin.cpp b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl-fatbin.cpp
new file mode 100644
index 000000000000..aa49a3e728fc
--- /dev/null
+++ b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl-fatbin.cpp
@@ -0,0 +1,53 @@
+// REQUIRES: system-linux
+
+// Fatbin-only output must not generate or compile a host wrapper.
+// RUN: touch %t.o
+// RUN: llvm-offload-binary -o %t.pack --image=file=%t.o,triple=spir64-unknown-unknown,arch=generic,kind=sycl
+// RUN: clang-linker-wrapper --dry-run --emit-fatbin-only --linker-path=/usr/bin/ld %t.pack -o %t.fatbin 2>&1 | FileCheck %s
+// CHECK: sycl-post-link
+// CHECK: llvm-spirv
+// CHECK-NOT: offload-wrapper
+// CHECK-NOT: "{{.*}}clang"{{.*}} -c
+// CHECK-NOT: "/usr/bin/ld"
+// RUN: llvm-objdump --offloading %t.fatbin | FileCheck %s --check-prefix=CONTAINER
+// CONTAINER: OFFLOADING IMAGE [0]:
+// CONTAINER: kind elf
+// CONTAINER: triple spir64-unknown-unknown
+// CONTAINER: producer sycl
+
+// Fatbin-only completion takes precedence over the device-link copy branch.
+// RUN: clang-linker-wrapper --dry-run --emit-fatbin-only --sycl-device-link --linker-path=/usr/bin/ld %t.pack -o %t.devlink.fatbin 2>&1 | FileCheck %s
+// RUN: llvm-objdump --offloading %t.devlink.fatbin | FileCheck %s --check-prefix=CONTAINER
+
+// Keep all independent no-RDC modules in a single container.
+// RUN: llvm-offload-binary -o %t.multi.pack --image=file=%t.o,triple=spir64-unknown-unknown,arch=generic,kind=sycl --image=file=%t.o,triple=spir64-unknown-unknown,arch=generic,kind=sycl
+// RUN: clang-linker-wrapper --dry-run --emit-fatbin-only --no-sycl-rdc --linker-path=/usr/bin/ld %t.multi.pack -o %t.multi.fatbin 2>&1 | FileCheck %s
+// RUN: llvm-objdump --offloading %t.multi.fatbin | FileCheck %s --check-prefix=MULTI
+// MULTI: OFFLOADING IMAGE [0]:
+// MULTI: producer sycl
+// MULTI: OFFLOADING IMAGE [1]:
+// MULTI: producer sycl
+// MULTI-NOT: OFFLOADING IMAGE
+
+// Reject modes that add host objects through the pipeline callback rather
+// than allowing them to overwrite the device-image container.
+// RUN: llvm-offload-binary -o %t.native.pack --image=file=%t.o,triple=native_cpu,arch=generic,kind=sycl
+// RUN: not clang-linker-wrapper --dry-run --emit-fatbin-only --linker-path=/usr/bin/ld %t.native.pack -o %t.native.fatbin 2>&1 | FileCheck %s --check-prefix=NATIVE
+// RUN: test ! -e %t.native.fatbin
+// NATIVE-NOT: sycl-post-link
+// NATIVE: error: --emit-fatbin-only is not supported for NativeCPU
+// NATIVE-NOT: sycl-post-link
+
+// RUN: llvm-offload-binary -o %t.nvptx.pack --image=file=%t.o,triple=nvptx64-nvidia-cuda,arch=sm_50,kind=sycl
+// RUN: not clang-linker-wrapper --dry-run --emit-fatbin-only --sycl-embed-ir --linker-path=/usr/bin/ld %t.nvptx.pack -o %t.nvptx.fatbin 2>&1 | FileCheck %s --check-prefix=EMBED
+// RUN: test ! -e %t.nvptx.fatbin
+// RUN: llvm-offload-binary -o %t.amdgcn.pack --image=file=%t.o,triple=amdgcn-amd-amdhsa,arch=gfx908,kind=sycl
+// RUN: not clang-linker-wrapper --dry-run --emit-fatbin-only --sycl-embed-ir --linker-path=/usr/bin/ld %t.amdgcn.pack -o %t.amdgcn.fatbin 2>&1 | FileCheck %s --check-prefix=EMBED
+// RUN: test ! -e %t.amdgcn.fatbin
+// EMBED-NOT: sycl-post-link
+// EMBED: error: --emit-fatbin-only is not supported with --sycl-embed-ir for NVPTX or AMDGCN
+// EMBED-NOT: sycl-post-link
+
+// Embedding IR does not add host objects for SPIR targets.
+// RUN: clang-linker-wrapper --dry-run --emit-fatbin-only --sycl-embed-ir --linker-path=/usr/bin/ld %t.pack -o %t.spir.fatbin 2>&1 | FileCheck %s
+// RUN: llvm-objdump --offloading %t.spir.fatbin | FileCheck %s --check-prefix=CONTAINER
diff --git a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl-image.cpp b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl-image.cpp
index 88eaed808227..d0adc12e10d0 100644
--- a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl-image.cpp
+++ b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl-image.cpp
@@ -15,21 +15,21 @@
 // RUN:   -sycl-post-link-options=-properties --linker-path=/usr/bin/ld \
 // RUN:   %t.o -o %t.out 2>&1 | FileCheck %s
 
-// With --emit-fatbin-only (-fno-sycl-rdc compile step) the same wrapper module
-// is written to the output as bitcode for the host compilation to link in.
+// With --emit-fatbin-only, the output is an offload container holding the
+// finalized SPIR-V with its runtime metadata. The host compilation wraps it
+// into the same descriptor-based registration.
 // RUN: clang-linker-wrapper --emit-fatbin-only --no-sycl-rdc \
 // RUN:   --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.o -o %t.syclfb
-// RUN: llvm-dis %t.syclfb -o - | FileCheck %s --check-prefix=NORDC
-// NORDC-DAG: @SYCL_PropSetName.1 = internal unnamed_addr constant [22 x i8] c"SYCL/kernel param opt\00"
-// NORDC-DAG: @.sycl_offloading.descriptor = internal constant %__sycl.tgt_bin_desc { i16 1, i16 1, ptr @.sycl_offloading.device_images, ptr null, ptr null }
-// NORDC-DAG: @llvm.global_ctors = {{.*}} { i32 1, ptr @sycl.descriptor_reg, ptr null }]
-//     NORDC: define internal void @sycl.descriptor_reg() section ".text.startup" {
-// NORDC-NEXT: entry:
-// NORDC-NEXT:   call void @__sycl_register_lib(ptr @.sycl_offloading.descriptor)
+// RUN: llvm-objdump --offloading %t.syclfb | FileCheck %s --check-prefix=NORDC
+// NORDC: OFFLOADING IMAGE [0]:
+// NORDC: triple spir64-unknown-unknown
+// NORDC: producer sycl
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsycl-is-host \
+// RUN:   -foffload-include-binary %t.syclfb -emit-llvm %s -o - | FileCheck %s
 
 // With thin LTO, sycl-post-link emits a two-column [Code|Properties] table.
 // RUN: clang-linker-wrapper --print-wrapped-module --host-triple=x86_64-unknown-linux-gnu \
diff --git a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp
index eff9831b8ffb..136589201be5 100644
--- a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp
+++ b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp
@@ -342,32 +342,40 @@
 // CHK-DEVLINK-CMDS-NEXT: "{{.*cp|copy}}" [[CLANGOUT]] /dev/null
 // CHK-DEVLINK-CMDS-NOT: "{{.*}}/ld"
 
-// Verify that with --emit-fatbin-only (-fno-sycl-rdc compile step) the wrapper
-// module is copied to the output as bitcode, without compiling it or calling
-// the host linker.
+// Verify that --emit-fatbin-only writes a raw fatbin without generating
+// wrapper bitcode, compiling it, or host linking.
 // RUN: clang-linker-wrapper --emit-fatbin-only --no-sycl-rdc --bitcode-library=spir64-unknown-unknown=%t.devicelib.bc -sycl-post-link-options=SYCL_POST_LINK_OPTIONS -llvm-spirv-options=LLVM_SPIRV_OPTIONS --host-triple=x86_64-unknown-linux-gnu --linker-path=/usr/bin/ld -o %t.syclfb %t.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-FATBIN-ONLY %s
+// CHK-FATBIN-ONLY-NOT: offload-wrapper:
+// CHK-FATBIN-ONLY-NOT: "{{.*}}clang"{{.*}} -c
+// CHK-FATBIN-ONLY-NOT: "{{.*cp|copy}}"
 // CHK-FATBIN-ONLY: llvm-spirv{{.*}} LLVM_SPIRV_OPTIONS -o {{.*}}
-// CHK-FATBIN-ONLY-NEXT: offload-wrapper: output: [[WRAPPEROUT:.*]].bc, input: {{.*}}
-// CHK-FATBIN-ONLY-NEXT: "{{.*cp|copy}}" [[WRAPPEROUT]].bc {{.*}}.syclfb
+// CHK-FATBIN-ONLY-NOT: offload-wrapper:
+// CHK-FATBIN-ONLY-NOT: "{{.*}}clang"{{.*}} -c
+// CHK-FATBIN-ONLY-NOT: "{{.*cp|copy}}"
 // CHK-FATBIN-ONLY-NOT: "{{.*}}/ld"
 
-// With several GPU architectures, all their images are wrapped into a single
-// wrapper module.
+// With several GPU architectures, all their images are packaged into a single
+// fatbin without wrapping or host linking.
 // RUN: clang-linker-wrapper --emit-fatbin-only --no-sycl-rdc --host-triple=x86_64-unknown-linux-gnu --linker-path=/usr/bin/ld -o %t.multi.syclfb %t_aot_gpu_multi.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-FATBIN-ONLY-MULTI %s
-// CHK-FATBIN-ONLY-MULTI-DAG: ocloc{{.*}} -device pvc{{.*}} -output [[PVC:[^ ]+]]
-// CHK-FATBIN-ONLY-MULTI-DAG: ocloc{{.*}} -device dg2{{.*}} -output [[DG2:[^ ]+]]
-// CHK-FATBIN-ONLY-MULTI: offload-wrapper: output: [[WRAPPEROUT:.*]].bc,{{.*}} input: [[DG2]],{{.*}} input: [[PVC]],
 // CHK-FATBIN-ONLY-MULTI-NOT: offload-wrapper:
-// CHK-FATBIN-ONLY-MULTI: "{{.*cp|copy}}" [[WRAPPEROUT]].bc {{.*}}.multi.syclfb
+// CHK-FATBIN-ONLY-MULTI-DAG: ocloc{{.*}} -device pvc
+// CHK-FATBIN-ONLY-MULTI-DAG: ocloc{{.*}} -device dg2
+// CHK-FATBIN-ONLY-MULTI-NOT: offload-wrapper:
+// CHK-FATBIN-ONLY-MULTI-NOT: "{{.*cp|copy}}"
 // CHK-FATBIN-ONLY-MULTI-NOT: "{{.*}}/ld"
-// RUN: not clang-linker-wrapper --emit-fatbin-only --no-sycl-rdc --host-triple=x86_64-unknown-linux-gnu --linker-path=/usr/bin/ld -o %t.mixed.syclfb %t_aot_gpu.o %t_aot_cpu.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-FATBIN-ONLY-MIXED %s
-// CHK-FATBIN-ONLY-MIXED: error: --no-sycl-rdc does not support multiple SYCL target triples ('{{.*}}' and '{{.*}}')
+// A fatbin can contain images for different SYCL target triples.
+// RUN: clang-linker-wrapper --emit-fatbin-only --no-sycl-rdc --host-triple=x86_64-unknown-linux-gnu --linker-path=/usr/bin/ld -o %t.mixed.syclfb %t_aot_gpu.o %t_aot_cpu.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-FATBIN-ONLY-MIXED %s
+// CHK-FATBIN-ONLY-MIXED-NOT: offload-wrapper:
+// CHK-FATBIN-ONLY-MIXED-DAG: ocloc
+// CHK-FATBIN-ONLY-MIXED-DAG: opencl-aot
+// CHK-FATBIN-ONLY-MIXED-NOT: offload-wrapper:
+// CHK-FATBIN-ONLY-MIXED-NOT: "{{.*}}/ld"
 // A SYCL input without a target triple is diagnosed.
 // RUN: %clang -cc1 -triple spir64-unknown-unknown -emit-llvm-bc -x c /dev/null -o %t_notriple.bc
 // RUN: llvm-offload-binary -o %t_notriple.fat --image=file=%t_notriple.bc,kind=sycl,triple=
 // RUN: %clang -cc1 -triple x86_64-unknown-linux-gnu -emit-obj -x c /dev/null -o %t_notriple.o -fembed-offload-object=%t_notriple.fat
 // RUN: not clang-linker-wrapper --emit-fatbin-only --no-sycl-rdc --host-triple=x86_64-unknown-linux-gnu --linker-path=/usr/bin/ld -o %t.notriple.syclfb %t_notriple.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-FATBIN-ONLY-NOTRIPLE %s
-// CHK-FATBIN-ONLY-NOTRIPLE: error: --no-sycl-rdc requires a target triple for every SYCL input
+// CHK-FATBIN-ONLY-NOTRIPLE: error: can't compile a SYCL device code. target is unknown
 
 // Verify list of commands when syclbin is used
 // RUN: clang-linker-wrapper --bitcode-library=spir64-unknown-unknown=%t.devicelib.bc -sycl-post-link-options=SYCL_POST_LINK_OPTIONS --host-triple=x86_64-unknown-linux-gnu --linker-path=/usr/bin/ld -o /dev/null %t.o --dry-run -syclbin=executable 2>&1 | FileCheck --check-prefix CHK-SYCLBIN-CMDS %s
diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
index dda61aef2c0a..6e514687d141 100644
--- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
+++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
@@ -1212,6 +1212,54 @@ static Error compressImages(SmallVectorImpl<offloading::SYCLImage> &Images,
   return Error::success();
 }
 
+/// Reads the post-link device images in \p SplitModules and prepares them for
+/// SYCL runtime registration: dumps them if requested, assigns the runtime
+/// target and compresses them if requested.
+static Expected<SmallVector<offloading::SYCLImage>>
+createSYCLImages(ArrayRef<module_split::SplitModule> SplitModules,
+                 const ArgList &Args, bool IsEmbeddedIR,
+                 const llvm::Triple &Target) {
+  SmallVector<llvm::offloading::SYCLImage> Images;
+  // SYCL runtime currently works for spir64 target triple and not for
+  // spir64-unknown-unknown/spirv64-unknown-unknown/spirv64.
+  // TODO: Fix SYCL runtime to accept other triples
+  std::string EmbeddedIRTarget("llvm_");
+  EmbeddedIRTarget.append(Target.getArchName());
+  StringRef RegularTarget(Target.getArchName());
+  if (RegularTarget == "spirv64")
+    RegularTarget = "spir64";
+
+  for (const module_split::SplitModule &SI : SplitModules) {
+    if (!OffloadImageDumpDir.empty() && !DryRun) {
+      StringRef CopyFrom = SI.ModuleFilePath;
+      SmallString<128> CopyTo = OffloadImageDumpDir;
+      StringRef Filename = sys::path::filename(CopyFrom);
+      CopyTo.append(Filename);
+      std::error_code EC = sys::fs::copy_file(CopyFrom, CopyTo);
+      if (EC)
+        return createStringError(EC, formatv("failed to copy file. From: "
+                                             "{0} to: {1}, error_code: {2}",
+                                             CopyFrom, CopyTo, EC.value()));
+    }
+
+    auto MBOrDesc = MemoryBuffer::getFile(SI.ModuleFilePath);
+    if (std::error_code EC = MBOrDesc.getError()) {
+      if (!DryRun)
+        return createFileError(SI.ModuleFilePath, EC);
+      MBOrDesc = MemoryBuffer::getMemBuffer("", SI.ModuleFilePath);
+    }
+
+    StringRef ImageTarget =
+        IsEmbeddedIR ? StringRef(EmbeddedIRTarget) : StringRef(RegularTarget);
+    Images.emplace_back(std::move(*MBOrDesc), SI.Properties, SI.Symbols,
+                        ImageTarget, SI.CompileOptions, SI.LinkOptions);
+  }
+
+  if (Error E = compressImages(Images, Args))
+    return std::move(E);
+  return std::move(Images);
+}
+
 /// Wraps the device images of \p SplitModules, built for the device target
 /// triple \p Target, in one LLVM IR module as constant data. If
 /// \p IsEmbeddedIR is true, the images are LLVM IR to be compiled at runtime.
@@ -1247,41 +1295,11 @@ wrapSYCLBinariesFromFile(ArrayRef<module_split::SplitModule> SplitModules,
       return OutputFilePath;
   }
 
-  SmallVector<llvm::offloading::SYCLImage> Images;
-  // SYCL runtime currently works for spir64 target triple and not for
-  // spir64-unknown-unknown/spirv64-unknown-unknown/spirv64.
-  // TODO: Fix SYCL runtime to accept other triples
-  std::string EmbeddedIRTarget("llvm_");
-  EmbeddedIRTarget.append(Target.getArchName());
-  StringRef RegularTarget(Target.getArchName());
-  if (RegularTarget == "spirv64")
-    RegularTarget = "spir64";
-
-  for (const module_split::SplitModule &SI : SplitModules) {
-    if (!OffloadImageDumpDir.empty()) {
-      StringRef CopyFrom = SI.ModuleFilePath;
-      SmallString<128> CopyTo = OffloadImageDumpDir;
-      StringRef Filename = sys::path::filename(CopyFrom);
-      CopyTo.append(Filename);
-      std::error_code EC = sys::fs::copy_file(CopyFrom, CopyTo);
-      if (EC)
-        return createStringError(EC, formatv("failed to copy file. From: "
-                                             "{0} to: {1}, error_code: {2}",
-                                             CopyFrom, CopyTo, EC.value()));
-    }
-
-    auto MBOrDesc = MemoryBuffer::getFile(SI.ModuleFilePath);
-    if (!MBOrDesc)
-      return createFileError(SI.ModuleFilePath, MBOrDesc.getError());
-
-    StringRef ImageTarget =
-        IsEmbeddedIR ? StringRef(EmbeddedIRTarget) : StringRef(RegularTarget);
-    Images.emplace_back(std::move(*MBOrDesc), SI.Properties, SI.Symbols,
-                        ImageTarget, SI.CompileOptions, SI.LinkOptions);
-  }
-
-  if (Error E = compressImages(Images, Args))
-    return std::move(E);
+  Expected<SmallVector<offloading::SYCLImage>> ImagesOrErr =
+      createSYCLImages(SplitModules, Args, IsEmbeddedIR, Target);
+  if (!ImagesOrErr)
+    return ImagesOrErr.takeError();
+  SmallVector<offloading::SYCLImage> &Images = *ImagesOrErr;
 
   LLVMContext C;
   Module M("offload.wrapper.object", C);
@@ -2935,16 +2953,19 @@ extractSYCLCompileLinkOptions(ArrayRef<OffloadFile> OffloadFiles) {
 }
 
 /// Transforms all the extracted offloading input files into an image that can
-/// be registered by the runtime. If NeedsWrapping is false, writes bundled
+/// be registered by the runtime. If EmitFatbinOnly is true, writes bundled
 /// output directly without wrapping or host linking.
 Expected<SmallVector<StringRef>>
 linkAndWrapDeviceFiles(ArrayRef<SmallVector<OffloadFile>> LinkerInputFiles,
                        const InputArgList &Args, char **Argv, int Argc,
-                       bool NeedsWrapping) {
+                       bool EmitFatbinOnly) {
   llvm::TimeTraceScope TimeScope("Handle all device input");
 
   std::mutex ImageMtx;
   MapVector<OffloadKind, SmallVector<OffloadingImage, 0>> Images;
+  // Owns the strings of the SYCL image entries; guarded by ImageMtx.
+  BumpPtrAllocator SYCLImageStringAlloc;
+  StringSaver SYCLImageStringSaver(SYCLImageStringAlloc);
   // Create a binary image of each offloading image and embed it into a new
   // object file.
   SmallVector<StringRef> WrappedOutput;
@@ -2954,36 +2975,6 @@ linkAndWrapDeviceFiles(ArrayRef<SmallVector<OffloadFile>> LinkerInputFiles,
   std::mutex SYCLBINModulesMtx;
   SmallVector<SYCLBIN::SYCLBINModuleDesc> SYCLBINModules;
 
-  // The -fno-sycl-rdc compile step wraps the modules of all architectures into
-  // a single wrapper module. One vector of split modules per input group, so
-  // the wrap order doesn't depend on thread scheduling. Each group only writes
-  // its own slot, so no lock is needed.
-  using SplitModulesVec = std::vector<module_split::SplitModule>;
-  std::vector<SplitModulesVec> NoRDCModulesPerGroup(LinkerInputFiles.size());
-
-  // The single wrapper module has one target triple, so check before any
-  // device compilation that all SYCL input groups share it.
-  std::optional<std::string> NoRDCTriple;
-  if (!NeedsWrapping && !OutputSYCLBIN) {
-    for (const SmallVector<OffloadFile> &Input : LinkerInputFiles) {
-      const auto *It = llvm::find_if(Input, [](const OffloadFile &F) {
-        return F.getBinary()->getOffloadKind() == OFK_SYCL;
-      });
-      if (It == Input.end())
-        continue;
-      StringRef Triple = It->getBinary()->getTriple();
-      if (Triple.empty())
-        return createStringError(
-            "--no-sycl-rdc requires a target triple for every SYCL input");
-      if (!NoRDCTriple)
-        NoRDCTriple = Triple.str();
-      else if (*NoRDCTriple != Triple)
-        return createStringError(
-            "--no-sycl-rdc does not support multiple SYCL target triples ('" +
-            *NoRDCTriple + "' and '" + Triple + "')");
-    }
-  }
-
   // Initialize the images with any overriding inputs.
   if (Args.hasArg(OPT_override_image))
     if (Error Err = handleOverrideImages(Args, Images))
@@ -3040,6 +3031,17 @@ linkAndWrapDeviceFiles(ArrayRef<SmallVector<OffloadFile>> LinkerInputFiles,
       // check in runSYCLOffloadingPipeline().
       const llvm::Triple TargetTriple(
           LinkerArgs.getLastArgValue(OPT_triple_EQ));
+      if (EmitFatbinOnly && !OutputSYCLBIN) {
+        // These modes add host objects through the pipeline callback, which
+        // cannot be represented in the device-image container.
+        if (TargetTriple.isNativeCPU())
+          return createStringError(
+              "--emit-fatbin-only is not supported for NativeCPU");
+        if ((TargetTriple.isNVPTX() || TargetTriple.isAMDGCN()) &&
+            LinkerArgs.hasArg(OPT_sycl_embed_ir))
+          return createStringError("--emit-fatbin-only is not supported with "
+                                   "--sycl-embed-ir for NVPTX or AMDGCN");
+      }
       const bool IsAOT = TargetTriple.isSPIRAOT();
       std::vector<std::string> AOTDeviceArgs;
       auto AddDeviceOptions = [&](opt::OptSpecifier OptID, bool IsLink,
@@ -3088,11 +3090,31 @@ linkAndWrapDeviceFiles(ArrayRef<SmallVector<OffloadFile>> LinkerInputFiles,
         MD.SplitModules = std::move(Modules);
         std::scoped_lock<std::mutex> Guard(SYCLBINModulesMtx);
         SYCLBINModules.emplace_back(std::move(MD));
-      } else if (!NeedsWrapping) {
-        // -fno-sycl-rdc compile step: the modules are wrapped once all inputs
-        // are processed.
-        size_t GroupIdx = std::distance(LinkerInputFiles.begin(), &Input);
-        NoRDCModulesPerGroup[GroupIdx] = std::move(Modules);
+      } else if (EmitFatbinOnly) {
+        // Package the finalized images with their runtime metadata; the host
+        // compilation wraps them for registration with the SYCL runtime.
+        Expected<SmallVector<offloading::SYCLImage>> SYCLImagesOrErr =
+            sycl::createSYCLImages(Modules, LinkerArgs,
+                                   /*IsEmbeddedIR=*/false, TargetTriple);
+        if (!SYCLImagesOrErr)
+          return SYCLImagesOrErr.takeError();
+
+        std::scoped_lock<decltype(ImageMtx)> Guard(ImageMtx);
+        for (offloading::SYCLImage &SYCLImage : *SYCLImagesOrErr) {
+          OffloadingImage TheImage{};
+          TheImage.TheImageKind =
+              Args.hasArg(OPT_embed_bitcode) ? IMG_Bitcode : IMG_Object;
+          TheImage.TheOffloadKind = OFK_SYCL;
+          TheImage.StringData["triple"] =
+              Args.MakeArgString(LinkerArgs.getLastArgValue(OPT_triple_EQ));
+          TheImage.StringData["arch"] =
+              Args.MakeArgString(LinkerArgs.getLastArgValue(OPT_arch_EQ));
+          offloading::addSYCLImageStrings(SYCLImage, TheImage.StringData,
+                                          SYCLImageStringSaver);
+          TheImage.Image = std::move(SYCLImage.Image);
+
+          Images[OFK_SYCL].emplace_back(std::move(TheImage));
+        }
       } else {
         // TODO(NOM7): Remove this call and use community flow for bundle/wrap
         Expected<StringRef> OutputFile =
@@ -3165,20 +3187,6 @@ linkAndWrapDeviceFiles(ArrayRef<SmallVector<OffloadFile>> LinkerInputFiles,
   if (Err)
     return std::move(Err);
 
-  // -fno-sycl-rdc compile step: emit the wrapper module as bitcode; the host
-  // compilation links it in via -foffload-include-binary.
-  if (NoRDCTriple) {
-    SplitModulesVec Modules;
-    for (SplitModulesVec &GroupModules : NoRDCModulesPerGroup)
-      std::move(GroupModules.begin(), GroupModules.end(),
-                std::back_inserter(Modules));
-    Expected<StringRef> OutputFile = sycl::wrapSYCLBinariesFromFile(
-        Modules, Args, /*IsEmbeddedIR=*/false, llvm::Triple(*NoRDCTriple));
-    if (!OutputFile)
-      return OutputFile.takeError();
-    WrappedOutput.push_back(*OutputFile);
-  }
-
   if (OutputSYCLBIN) {
     auto OutputOrErr = sycl::packageSYCLBIN(SYCLBINState, SYCLBINModules);
     if (!OutputOrErr)
@@ -3187,7 +3195,7 @@ linkAndWrapDeviceFiles(ArrayRef<SmallVector<OffloadFile>> LinkerInputFiles,
   }
 
   for (auto &[Kind, Input] : Images) {
-    if (Kind == OFK_SYCL)
+    if (Kind == OFK_SYCL && !EmitFatbinOnly)
       continue;
     // We sort the entries before bundling so they appear in a deterministic
     // order in the final binary.
@@ -3202,11 +3210,23 @@ linkAndWrapDeviceFiles(ArrayRef<SmallVector<OffloadFile>> LinkerInputFiles,
         return ArchA > ArchB;
       return A.TheOffloadKind < B.TheOffloadKind;
     });
+    if (Kind == OFK_SYCL) {
+      // Package all post-link images into a single offload container.
+      SmallString<0> Buffer = OffloadBinary::write(Input);
+      auto FOBOrErr = FileOutputBuffer::create(ExecutableName, Buffer.size());
+      if (!FOBOrErr)
+        return FOBOrErr.takeError();
+      llvm::copy(Buffer, (*FOBOrErr)->getBufferStart());
+      if (Error E = (*FOBOrErr)->commit())
+        return std::move(E);
+      continue;
+    }
+
     auto BundledImagesOrErr = bundleLinkedOutput(Input, Args, Kind);
     if (!BundledImagesOrErr)
       return BundledImagesOrErr.takeError();
 
-    if (!NeedsWrapping) {
+    if (EmitFatbinOnly) {
       if (BundledImagesOrErr->size() != 1)
         return createStringError(
             "Expected a single bundled image for direct fat binary output");
@@ -3703,34 +3723,27 @@ int main(int Argc, char **Argv) {
     // Link and process the device images. The function may emit a direct fat
     // binary if --emit-fatbin-only is specified.
     auto FilesOrErr = linkAndWrapDeviceFiles(*DeviceInputFiles, Args, Argv,
-                                             Argc, !EmitFatbinOnly);
+                                             Argc, EmitFatbinOnly);
     if (!FilesOrErr)
       reportError(FilesOrErr.takeError());
 
     if (OutputSYCLBIN) {
       if (Error Err = sycl::mergeSYCLBIN(*FilesOrErr, Args))
         reportError(std::move(Err));
-    } else if (Args.hasArg(OPT_sycl_device_link)) {
-      // Skip host linker if --sycl-device-link option is set.
-      // Just copy the output of device linking and wrapping action.
-      if (FilesOrErr->size() != 1)
-        reportError(
-            createStringError("Expect single output from the device linker."));
-      if (Error Err = sycl::copyFileToFinalExecutable((*FilesOrErr)[0], Args))
-        reportError(std::move(Err));
-    } else {
-      // Run the host linking job with the rendered arguments.
-      if (!EmitFatbinOnly) {
-        if (Error Err = runLinker(*FilesOrErr, Args))
-          reportError(std::move(Err));
-      } else if (!FilesOrErr->empty()) {
-        // Other offload kinds write the fat binary directly; the SYCL wrapper
-        // module still has to be copied to the requested output.
+    } else if (!EmitFatbinOnly) {
+      // Fatbin-only output has already been written by linkAndWrapDeviceFiles.
+      if (Args.hasArg(OPT_sycl_device_link)) {
+        // Skip host linker if --sycl-device-link option is set.
+        // Just copy the output of device linking and wrapping action.
         if (FilesOrErr->size() != 1)
           reportError(createStringError(
-              "Expect single output from the SYCL device linker."));
+              "Expect single output from the device linker."));
         if (Error Err = sycl::copyFileToFinalExecutable((*FilesOrErr)[0], Args))
           reportError(std::move(Err));
+      } else {
+        // Run the host linking job with the rendered arguments.
+        if (Error Err = runLinker(*FilesOrErr, Args))
+          reportError(std::move(Err));
       }
     }
   }
diff --git a/llvm/include/llvm/Frontend/Offloading/SYCLOffloadWrapper.h b/llvm/include/llvm/Frontend/Offloading/SYCLOffloadWrapper.h
index 4e1765a56a00..ce260454694d 100644
--- a/llvm/include/llvm/Frontend/Offloading/SYCLOffloadWrapper.h
+++ b/llvm/include/llvm/Frontend/Offloading/SYCLOffloadWrapper.h
@@ -9,11 +9,14 @@
 #ifndef LLVM_CLANG_TOOLS_CLANG_LINKER_WRAPPER_SYCL_OFFLOAD_WRAPPER_H
 #define LLVM_CLANG_TOOLS_CLANG_LINKER_WRAPPER_SYCL_OFFLOAD_WRAPPER_H
 
+#include "llvm/ADT/MapVector.h"
 #include "llvm/ADT/SmallVector.h"
 #include "llvm/ADT/StringRef.h"
 #include "llvm/IR/Module.h"
+#include "llvm/Support/Error.h"
 #include "llvm/Support/MemoryBuffer.h"
 #include "llvm/Support/PropertySetIO.h"
+#include "llvm/Support/StringSaver.h"
 
 #include <string>
 #include <utility>
@@ -39,6 +42,9 @@ enum class SYCLBinaryImageFormat {
 /// Returns the ABI value of \p Format shared with the SYCL runtime.
 int8_t binaryImageFormatToInt8(SYCLBinaryImageFormat Format);
 
+/// Returns the SYCL image format for an ABI value, or an error if invalid.
+Expected<SYCLBinaryImageFormat> int8ToBinaryImageFormat(int8_t Format);
+
 /// Misc property with the pre-compression format; mirrored in SYCL RT.
 constexpr char OriginalFormatPropName[] = "originalFormat";
 
@@ -88,6 +94,21 @@ llvm::Error
 wrapSYCLBinaries(llvm::Module &M, const llvm::SmallVector<SYCLImage> &Images,
                  SYCLWrappingOptions Options = SYCLWrappingOptions());
 
+/// Records the fields of \p Image other than the image bytes in the string
+/// table \p StringData of an offload binary entry, so that a fat binary can be
+/// wrapped later by readSYCLImages() and wrapSYCLBinaries(). \p Saver owns the
+/// serialized strings, which must outlive the offload binary serialization.
+void addSYCLImageStrings(
+    const SYCLImage &Image,
+    llvm::MapVector<llvm::StringRef, llvm::StringRef> &StringData,
+    llvm::StringSaver &Saver);
+
+/// Reads SYCL images from the offload binary container in \p Buffer, whose
+/// entries carry the strings recorded by addSYCLImageStrings(). The returned
+/// images reference the data of \p Buffer, which must outlive them.
+llvm::Expected<llvm::SmallVector<SYCLImage>>
+readSYCLImages(llvm::MemoryBufferRef Buffer);
+
 } // namespace offloading
 } // namespace llvm
 
diff --git a/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp b/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp
index 9457b278e034..19c0f7fcd136 100644
--- a/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp
+++ b/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp
@@ -28,12 +28,14 @@
 #include "llvm/IR/LLVMContext.h"
 #include "llvm/IR/Module.h"
 #include "llvm/IR/Type.h"
+#include "llvm/Object/OffloadBinary.h"
 #include "llvm/SYCLLowerIR/SYCLUtils.h"
 #include "llvm/SYCLLowerIR/UtilsSYCLNativeCPU.h"
 #include "llvm/Support/Error.h"
 #include "llvm/Support/ErrorHandling.h"
 #include "llvm/Support/LineIterator.h"
 #include "llvm/Support/PropertySetIO.h"
+#include "llvm/Support/raw_ostream.h"
 #include "llvm/TargetParser/Triple.h"
 #include "llvm/Transforms/Utils/ModuleUtils.h"
 #include <memory>
@@ -63,6 +65,25 @@ int8_t llvm::offloading::binaryImageFormatToInt8(SYCLBinaryImageFormat Format) {
   }
 }
 
+Expected<SYCLBinaryImageFormat>
+llvm::offloading::int8ToBinaryImageFormat(int8_t Format) {
+  switch (Format) {
+  case 0:
+    return SYCLBinaryImageFormat::BIF_None;
+  case 1:
+    return SYCLBinaryImageFormat::BIF_Native;
+  case 2:
+    return SYCLBinaryImageFormat::BIF_SPIRV;
+  case 3:
+    return SYCLBinaryImageFormat::BIF_LLVMBC;
+  case 4:
+    return SYCLBinaryImageFormat::BIF_Compressed;
+  default:
+    return createStringError("invalid SYCL image format '" +
+                             Twine(static_cast<int>(Format)) + "'");
+  }
+}
+
 namespace {
 
 /// Wrapper helper class that creates all LLVM IRs wrapping given images.
@@ -787,3 +808,68 @@ Error llvm::offloading::wrapSYCLBinaries(llvm::Module &M,
   }
   return Error::success();
 }
+
+namespace {
+// Offload binary string keys that carry the SYCLImage fields.
+constexpr char TargetKey[] = "sycl-target";
+constexpr char FormatKey[] = "sycl-image-format";
+constexpr char EntriesKey[] = "symbols";
+constexpr char PropertiesKey[] = "properties";
+constexpr char CompileOptionsKey[] = "compile-opts";
+constexpr char LinkOptionsKey[] = "link-opts";
+} // anonymous namespace
+
+void llvm::offloading::addSYCLImageStrings(
+    const SYCLImage &Image, MapVector<StringRef, StringRef> &StringData,
+    StringSaver &Saver) {
+  std::string Properties;
+  raw_string_ostream PropertiesOS(Properties);
+  Image.PropertyRegistry.write(PropertiesOS);
+  StringData[TargetKey] = Saver.save(Image.Target);
+  StringData[FormatKey] = Saver.save(
+      Twine(static_cast<int>(binaryImageFormatToInt8(Image.Format))));
+  StringData[EntriesKey] = Saver.save(Image.Entries);
+  StringData[PropertiesKey] = Saver.save(Properties);
+  StringData[CompileOptionsKey] = Saver.save(Image.CompileOptions);
+  StringData[LinkOptionsKey] = Saver.save(Image.LinkOptions);
+}
+
+Expected<SmallVector<SYCLImage>>
+llvm::offloading::readSYCLImages(MemoryBufferRef Buffer) {
+  auto BinariesOrErr = object::OffloadBinary::create(Buffer);
+  if (!BinariesOrErr)
+    return BinariesOrErr.takeError();
+
+  SmallVector<SYCLImage> Images;
+  for (const std::unique_ptr<object::OffloadBinary> &Binary : *BinariesOrErr) {
+    if (Binary->getOffloadKind() != object::OFK_SYCL)
+      return createStringError("offload binary entry " +
+                               Twine(Binary->getIndex()) +
+                               " is not a SYCL device image");
+
+    std::unique_ptr<MemoryBuffer> PropertiesBuffer =
+        MemoryBuffer::getMemBufferCopy(Binary->getString(PropertiesKey));
+    auto PropertiesOrErr = PropertySetRegistry::read(PropertiesBuffer.get());
+    if (!PropertiesOrErr)
+      return PropertiesOrErr.takeError();
+
+    SYCLImage Image(MemoryBuffer::getMemBuffer(
+                        Binary->getImage(), Buffer.getBufferIdentifier(),
+                        /*RequiresNullTerminator=*/false),
+                    **PropertiesOrErr, Binary->getString(EntriesKey),
+                    Binary->getString(TargetKey),
+                    Binary->getString(CompileOptionsKey),
+                    Binary->getString(LinkOptionsKey));
+
+    int8_t Format = 0;
+    if (Binary->getString(FormatKey).getAsInteger(10, Format))
+      return createStringError("invalid SYCL image format '" +
+                               Binary->getString(FormatKey) + "'");
+    auto FormatOrErr = int8ToBinaryImageFormat(Format);
+    if (!FormatOrErr)
+      return FormatOrErr.takeError();
+    Image.Format = *FormatOrErr;
+    Images.emplace_back(std::move(Image));
+  }
+  return std::move(Images);
+}