diff --git a/sycl-jit/common/include/Kernel.h b/sycl-jit/common/include/Kernel.h
index bbb9a802995ff..794126a462e20 100644
--- a/sycl-jit/common/include/Kernel.h
+++ b/sycl-jit/common/include/Kernel.h
@@ -17,6 +17,7 @@
 #include <cstdint>
 #include <cstring>
 #include <functional>
+#include <string_view>
 #include <type_traits>
 
 namespace jit_compiler {
@@ -350,11 +351,60 @@ struct SYCLKernelInfo {
       : Name{KernelName}, Args{NumArgs}, Attributes{}, NDR{}, BinaryInfo{} {}
 };
 
+// RTC-related datastructures
+// TODO: Consider moving into separate header.
+
 struct InMemoryFile {
   const char *Path;
   const char *Contents;
 };
 
+using RTCBundleBinaryInfo = SYCLKernelBinaryInfo;
+using FrozenSymbolTable = DynArray<sycl::detail::string>;
+
+// Note: `FrozenPropertyValue` and `FrozenPropertySet` constructors take
+// `std::string_view` arguments instead of `const char *` because they will be
+// created from `llvm::SmallString`s, which don't contain the trailing '\0'
+// byte. Hence obtaining a C-string would cause an additional copy.
+
+struct FrozenPropertyValue {
+  sycl::detail::string Name;
+  bool IsUIntValue;
+  uint32_t UIntValue;
+  DynArray<uint8_t> Bytes;
+
+  FrozenPropertyValue() = default;
+  FrozenPropertyValue(FrozenPropertyValue &&) = default;
+  FrozenPropertyValue &operator=(FrozenPropertyValue &&) = default;
+
+  FrozenPropertyValue(std::string_view Name, uint32_t Value)
+      : Name{Name}, IsUIntValue{true}, UIntValue{Value}, Bytes{0} {}
+  FrozenPropertyValue(std::string_view Name, const uint8_t *Ptr, size_t Size)
+      : Name{Name}, IsUIntValue{false}, Bytes{Size} {
+    std::memcpy(Bytes.begin(), Ptr, Size);
+  }
+};
+
+struct FrozenPropertySet {
+  sycl::detail::string Name;
+  DynArray<FrozenPropertyValue> Values;
+
+  FrozenPropertySet() = default;
+  FrozenPropertySet(FrozenPropertySet &&) = default;
+  FrozenPropertySet &operator=(FrozenPropertySet &&) = default;
+
+  FrozenPropertySet(std::string_view Name, size_t Size)
+      : Name{Name}, Values{Size} {}
+};
+
+using FrozenPropertyRegistry = DynArray<FrozenPropertySet>;
+
+struct RTCBundleInfo {
+  RTCBundleBinaryInfo BinaryInfo;
+  FrozenSymbolTable SymbolTable;
+  FrozenPropertyRegistry Properties;
+};
+
 } // namespace jit_compiler
 
 #endif // SYCL_FUSION_COMMON_KERNEL_H
diff --git a/sycl-jit/jit-compiler/CMakeLists.txt b/sycl-jit/jit-compiler/CMakeLists.txt
index 98385e8ee4f2a..82d2356691c3f 100644
--- a/sycl-jit/jit-compiler/CMakeLists.txt
+++ b/sycl-jit/jit-compiler/CMakeLists.txt
@@ -31,6 +31,7 @@ add_llvm_library(sycl-jit
    Target
    TargetParser
    MC
+   SYCLLowerIR
    ${LLVM_TARGETS_TO_BUILD}
 
    LINK_LIBS
diff --git a/sycl-jit/jit-compiler/include/KernelFusion.h b/sycl-jit/jit-compiler/include/KernelFusion.h
index d3575f33189aa..76636beb937d9 100644
--- a/sycl-jit/jit-compiler/include/KernelFusion.h
+++ b/sycl-jit/jit-compiler/include/KernelFusion.h
@@ -56,6 +56,32 @@ class JITResult {
   sycl::detail::string ErrorMessage;
 };
 
+class RTCResult {
+public:
+  explicit RTCResult(const char *ErrorMessage)
+      : Failed{true}, BundleInfo{}, ErrorMessage{ErrorMessage} {}
+
+  explicit RTCResult(RTCBundleInfo &&BundleInfo)
+      : Failed{false}, BundleInfo{std::move(BundleInfo)}, ErrorMessage{} {}
+
+  bool failed() const { return Failed; }
+
+  const char *getErrorMessage() const {
+    assert(failed() && "No error message present");
+    return ErrorMessage.c_str();
+  }
+
+  const RTCBundleInfo &getBundleInfo() const {
+    assert(!failed() && "No bundle info");
+    return BundleInfo;
+  }
+
+private:
+  bool Failed;
+  RTCBundleInfo BundleInfo;
+  sycl::detail::string ErrorMessage;
+};
+
 extern "C" {
 
 #ifdef __clang__
@@ -77,7 +103,7 @@ KF_EXPORT_SYMBOL JITResult materializeSpecConstants(
     const char *KernelName, jit_compiler::SYCLKernelBinaryInfo &BinInfo,
     View<unsigned char> SpecConstBlob);
 
-KF_EXPORT_SYMBOL JITResult compileSYCL(InMemoryFile SourceFile,
+KF_EXPORT_SYMBOL RTCResult compileSYCL(InMemoryFile SourceFile,
                                        View<InMemoryFile> IncludeFiles,
                                        View<const char *> UserArgs);
 
diff --git a/sycl-jit/jit-compiler/lib/KernelFusion.cpp b/sycl-jit/jit-compiler/lib/KernelFusion.cpp
index 9f8bb09ec9f26..fae9a3c29dcf4 100644
--- a/sycl-jit/jit-compiler/lib/KernelFusion.cpp
+++ b/sycl-jit/jit-compiler/lib/KernelFusion.cpp
@@ -25,8 +25,8 @@ using namespace jit_compiler;
 using FusedFunction = helper::FusionHelper::FusedFunction;
 using FusedFunctionList = std::vector<FusedFunction>;
 
-static JITResult errorToFusionResult(llvm::Error &&Err,
-                                     const std::string &Msg) {
+template <typename ResultType>
+static ResultType errorTo(llvm::Error &&Err, const std::string &Msg) {
   std::stringstream ErrMsg;
   ErrMsg << Msg << "\nDetailed information:\n";
   llvm::handleAllErrors(std::move(Err),
@@ -35,7 +35,7 @@ static JITResult errorToFusionResult(llvm::Error &&Err,
                           // compiled without exception support.
                           ErrMsg << "\t" << StrErr.getMessage() << "\n";
                         });
-  return JITResult{ErrMsg.str().c_str()};
+  return ResultType{ErrMsg.str().c_str()};
 }
 
 static std::vector<jit_compiler::NDRange>
@@ -95,7 +95,7 @@ extern "C" KF_EXPORT_SYMBOL JITResult materializeSpecConstants(
       translation::KernelTranslator::loadKernels(*JITCtx.getLLVMContext(),
                                                  ModuleInfo.kernels());
   if (auto Error = ModOrError.takeError()) {
-    return errorToFusionResult(std::move(Error), "Failed to load kernels");
+    return errorTo<JITResult>(std::move(Error), "Failed to load kernels");
   }
   std::unique_ptr<llvm::Module> NewMod = std::move(*ModOrError);
   if (!fusion::FusionPipeline::runMaterializerPasses(
@@ -107,8 +107,8 @@ extern "C" KF_EXPORT_SYMBOL JITResult materializeSpecConstants(
   SYCLKernelInfo &MaterializerKernelInfo = *ModuleInfo.getKernelFor(KernelName);
   if (auto Error = translation::KernelTranslator::translateKernel(
           MaterializerKernelInfo, *NewMod, JITCtx, TargetFormat)) {
-    return errorToFusionResult(std::move(Error),
-                               "Translation to output format failed");
+    return errorTo<JITResult>(std::move(Error),
+                              "Translation to output format failed");
   }
 
   return JITResult{MaterializerKernelInfo};
@@ -133,7 +133,7 @@ fuseKernels(View<SYCLKernelInfo> KernelInformation, const char *FusedKernelName,
   llvm::Expected<jit_compiler::FusedNDRange> FusedNDR =
       jit_compiler::FusedNDRange::get(NDRanges);
   if (llvm::Error Err = FusedNDR.takeError()) {
-    return errorToFusionResult(std::move(Err), "Illegal ND-range combination");
+    return errorTo<JITResult>(std::move(Err), "Illegal ND-range combination");
   }
 
   if (!isTargetFormatSupported(TargetFormat)) {
@@ -180,7 +180,7 @@ fuseKernels(View<SYCLKernelInfo> KernelInformation, const char *FusedKernelName,
       translation::KernelTranslator::loadKernels(*JITCtx.getLLVMContext(),
                                                  ModuleInfo.kernels());
   if (auto Error = ModOrError.takeError()) {
-    return errorToFusionResult(std::move(Error), "SPIR-V translation failed");
+    return errorTo<JITResult>(std::move(Error), "SPIR-V translation failed");
   }
   std::unique_ptr<llvm::Module> LLVMMod = std::move(*ModOrError);
 
@@ -197,8 +197,8 @@ fuseKernels(View<SYCLKernelInfo> KernelInformation, const char *FusedKernelName,
   llvm::Expected<std::unique_ptr<llvm::Module>> NewModOrError =
       helper::FusionHelper::addFusedKernel(LLVMMod.get(), FusedKernelList);
   if (auto Error = NewModOrError.takeError()) {
-    return errorToFusionResult(std::move(Error),
-                               "Insertion of fused kernel stub failed");
+    return errorTo<JITResult>(std::move(Error),
+                              "Insertion of fused kernel stub failed");
   }
   std::unique_ptr<llvm::Module> NewMod = std::move(*NewModOrError);
 
@@ -221,8 +221,8 @@ fuseKernels(View<SYCLKernelInfo> KernelInformation, const char *FusedKernelName,
 
   if (auto Error = translation::KernelTranslator::translateKernel(
           FusedKernelInfo, *NewMod, JITCtx, TargetFormat)) {
-    return errorToFusionResult(std::move(Error),
-                               "Translation to output format failed");
+    return errorTo<JITResult>(std::move(Error),
+                              "Translation to output format failed");
   }
 
   FusedKernelInfo.NDR = FusedNDR->getNDR();
@@ -234,20 +234,20 @@ fuseKernels(View<SYCLKernelInfo> KernelInformation, const char *FusedKernelName,
   return JITResult{FusedKernelInfo};
 }
 
-extern "C" KF_EXPORT_SYMBOL JITResult
+extern "C" KF_EXPORT_SYMBOL RTCResult
 compileSYCL(InMemoryFile SourceFile, View<InMemoryFile> IncludeFiles,
             View<const char *> UserArgs) {
   auto UserArgListOrErr = parseUserArgs(UserArgs);
   if (!UserArgListOrErr) {
-    return errorToFusionResult(UserArgListOrErr.takeError(),
-                               "Parsing of user arguments failed");
+    return errorTo<RTCResult>(UserArgListOrErr.takeError(),
+                              "Parsing of user arguments failed");
   }
   llvm::opt::InputArgList UserArgList = std::move(*UserArgListOrErr);
 
   auto ModuleOrErr = compileDeviceCode(SourceFile, IncludeFiles, UserArgList);
   if (!ModuleOrErr) {
-    return errorToFusionResult(ModuleOrErr.takeError(),
-                               "Device compilation failed");
+    return errorTo<RTCResult>(ModuleOrErr.takeError(),
+                              "Device compilation failed");
   }
 
   std::unique_ptr<llvm::LLVMContext> Context;
@@ -255,16 +255,26 @@ compileSYCL(InMemoryFile SourceFile, View<InMemoryFile> IncludeFiles,
   Context.reset(&Module->getContext());
 
   if (auto Error = linkDeviceLibraries(*Module, UserArgList)) {
-    return errorToFusionResult(std::move(Error), "Device linking failed");
+    return errorTo<RTCResult>(std::move(Error), "Device linking failed");
   }
 
-  SYCLKernelInfo Kernel;
-  if (auto Error = translation::KernelTranslator::translateKernel(
-          Kernel, *Module, JITContext::getInstance(), BinaryFormat::SPIRV)) {
-    return errorToFusionResult(std::move(Error), "SPIR-V translation failed");
+  auto BundleInfoOrError = performPostLink(*Module, UserArgList);
+  if (!BundleInfoOrError) {
+    return errorTo<RTCResult>(BundleInfoOrError.takeError(),
+                              "Post-link phase failed");
+  }
+  auto BundleInfo = std::move(*BundleInfoOrError);
+
+  auto BinaryInfoOrError =
+      translation::KernelTranslator::translateBundleToSPIRV(
+          *Module, JITContext::getInstance());
+  if (!BinaryInfoOrError) {
+    return errorTo<RTCResult>(BinaryInfoOrError.takeError(),
+                              "SPIR-V translation failed");
   }
+  BundleInfo.BinaryInfo = std::move(*BinaryInfoOrError);
 
-  return JITResult{Kernel};
+  return RTCResult{std::move(BundleInfo)};
 }
 
 extern "C" KF_EXPORT_SYMBOL void resetJITConfiguration() {
diff --git a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp
index 1bdfe7d63b641..a89635b79ed46 100644
--- a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp
+++ b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp
@@ -18,10 +18,14 @@
 #include <clang/Tooling/CompilationDatabase.h>
 #include <clang/Tooling/Tooling.h>
 
+#include <llvm/IR/PassInstrumentation.h>
+#include <llvm/IR/PassManager.h>
 #include <llvm/IRReader/IRReader.h>
 #include <llvm/Linker/Linker.h>
-
-#include <array>
+#include <llvm/SYCLLowerIR/ComputeModuleRuntimeInfo.h>
+#include <llvm/SYCLLowerIR/ModuleSplitter.h>
+#include <llvm/SYCLLowerIR/SYCLJointMatrixTransform.h>
+#include <llvm/Support/PropertySetIO.h>
 
 using namespace clang;
 using namespace clang::tooling;
@@ -29,6 +33,10 @@ using namespace clang::driver;
 using namespace clang::driver::options;
 using namespace llvm;
 using namespace llvm::opt;
+using namespace llvm::sycl;
+using namespace llvm::module_split;
+using namespace llvm::util;
+using namespace jit_compiler;
 
 #ifdef _GNU_SOURCE
 #include <dlfcn.h>
@@ -356,6 +364,95 @@ Error jit_compiler::linkDeviceLibraries(llvm::Module &Module,
   return Error::success();
 }
 
+template <class PassClass> static bool runModulePass(llvm::Module &M) {
+  ModulePassManager MPM;
+  ModuleAnalysisManager MAM;
+  // Register required analysis
+  MAM.registerPass([&] { return PassInstrumentationAnalysis(); });
+  MPM.addPass(PassClass{});
+  PreservedAnalyses Res = MPM.run(M, MAM);
+  return !Res.areAllPreserved();
+}
+
+Expected<RTCBundleInfo> jit_compiler::performPostLink(
+    llvm::Module &Module, [[maybe_unused]] const InputArgList &UserArgList) {
+  // This is a simplified version of `processInputModule` in
+  // `llvm/tools/sycl-post-link.cpp`. Assertions/TODOs point to functionality
+  // left out of the algorithm for now.
+
+  assert(!Module.getGlobalVariable("llvm.used") &&
+         !Module.getGlobalVariable("llvm.compiler.used"));
+  // Otherwise: Port over the `removeSYCLKernelsConstRefArray` and
+  // `removeDeviceGlobalFromCompilerUsed` methods.
+
+  assert(!isModuleUsingAsan(Module));
+  // Otherwise: Need to instrument each image scope device globals if the module
+  // has been instrumented by sanitizer pass.
+
+  // Transform Joint Matrix builtin calls to align them with SPIR-V friendly
+  // LLVM IR specification.
+  runModulePass<SYCLJointMatrixTransformPass>(Module);
+
+  // TODO: Implement actual device code splitting. We're just using the splitter
+  //       to obtain additional information about the module for now.
+  // TODO: EmitOnlyKernelsAsEntryPoints is controlled by
+  //       `shouldEmitOnlyKernelsAsEntryPoints` in
+  //       `clang/lib/Driver/ToolChains/Clang.cpp`.
+  std::unique_ptr<ModuleSplitterBase> Splitter = getDeviceCodeSplitter(
+      ModuleDesc{std::unique_ptr<llvm::Module>{&Module}}, SPLIT_NONE,
+      /*IROutputOnly=*/false,
+      /*EmitOnlyKernelsAsEntryPoints=*/true);
+  bool SplitOccurred = Splitter->remainingSplits() > 1;
+  assert(!SplitOccurred);
+
+  // TODO: Call `verifyNoCrossModuleDeviceGlobalUsage` if device globals shall
+  //       be processed.
+
+  assert(Splitter->hasMoreSplits());
+  ModuleDesc MDesc = Splitter->nextSplit();
+  assert(&Module == &MDesc.getModule());
+  MDesc.saveSplitInformationAsMetadata();
+
+  RTCBundleInfo BundleInfo;
+  BundleInfo.SymbolTable = FrozenSymbolTable{MDesc.entries().size()};
+  transform(MDesc.entries(), BundleInfo.SymbolTable.begin(),
+            [](Function *F) { return F->getName(); });
+
+  // TODO: Determine what is requested.
+  GlobalBinImageProps PropReq{
+      /*EmitKernelParamInfo=*/true, /*EmitProgramMetadata=*/true,
+      /*EmitExportedSymbols=*/true, /*EmitImportedSymbols=*/true,
+      /*DeviceGlobals=*/false};
+  PropertySetRegistry Properties =
+      computeModuleProperties(MDesc.getModule(), MDesc.entries(), PropReq);
+  // TODO: Manually add `compile_target` property as in
+  //       `saveModuleProperties`?
+  const auto &PropertySets = Properties.getPropSets();
+
+  BundleInfo.Properties = FrozenPropertyRegistry{PropertySets.size()};
+  for (auto &&[KV, FrozenPropSet] : zip(PropertySets, BundleInfo.Properties)) {
+    const auto &PropertySetName = KV.first;
+    const auto &PropertySet = KV.second;
+    FrozenPropSet =
+        FrozenPropertySet{PropertySetName.str(), PropertySet.size()};
+    for (auto &&[KV2, FrozenProp] : zip(PropertySet, FrozenPropSet.Values)) {
+      const auto &PropertyName = KV2.first;
+      const auto &PropertyValue = KV2.second;
+      FrozenProp = PropertyValue.getType() == PropertyValue::Type::UINT32
+                       ? FrozenPropertyValue{PropertyName.str(),
+                                             PropertyValue.asUint32()}
+                       : FrozenPropertyValue{
+                             PropertyName.str(), PropertyValue.asRawByteArray(),
+                             PropertyValue.getRawByteArraySize()};
+    }
+  };
+
+  // Regain ownership of the module.
+  MDesc.releaseModulePtr().release();
+
+  return std::move(BundleInfo);
+}
+
 Expected<InputArgList>
 jit_compiler::parseUserArgs(View<const char *> UserArgs) {
   unsigned MissingArgIndex, MissingArgCount;
@@ -410,5 +507,17 @@ jit_compiler::parseUserArgs(View<const char *> UserArgs) {
     }
   }
 
-  return Expected<InputArgList>{std::move(AL)};
+  if (auto DCSMode = AL.getLastArgValue(OPT_fsycl_device_code_split_EQ, "none");
+      DCSMode != "none" && DCSMode != "auto") {
+    return createStringError("Device code splitting is not yet supported");
+  }
+
+  if (AL.hasArg(OPT_fsycl_device_code_split_esimd,
+                OPT_fno_sycl_device_code_split_esimd)) {
+    // TODO: There are more ESIMD-related options.
+    return createStringError(
+        "Runtime compilation of ESIMD kernels is not yet supported");
+  }
+
+  return std::move(AL);
 }
diff --git a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.h b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.h
index ec890a8213827..8aa47939e3b1d 100644
--- a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.h
+++ b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.h
@@ -27,6 +27,10 @@ compileDeviceCode(InMemoryFile SourceFile, View<InMemoryFile> IncludeFiles,
 llvm::Error linkDeviceLibraries(llvm::Module &Module,
                                 const llvm::opt::InputArgList &UserArgList);
 
+llvm::Expected<RTCBundleInfo>
+performPostLink(llvm::Module &Module,
+                const llvm::opt::InputArgList &UserArgList);
+
 llvm::Expected<llvm::opt::InputArgList>
 parseUserArgs(View<const char *> UserArgs);
 
diff --git a/sycl-jit/jit-compiler/lib/translation/KernelTranslation.cpp b/sycl-jit/jit-compiler/lib/translation/KernelTranslation.cpp
index 40eec5d241ea0..61ca038dea4cf 100644
--- a/sycl-jit/jit-compiler/lib/translation/KernelTranslation.cpp
+++ b/sycl-jit/jit-compiler/lib/translation/KernelTranslation.cpp
@@ -222,6 +222,20 @@ llvm::Error KernelTranslator::translateKernel(SYCLKernelInfo &Kernel,
   return Error::success();
 }
 
+llvm::Expected<RTCBundleBinaryInfo>
+KernelTranslator::translateBundleToSPIRV(llvm::Module &Mod,
+                                         JITContext &JITCtx) {
+  llvm::Expected<KernelBinary *> BinaryOrError = translateToSPIRV(Mod, JITCtx);
+  if (auto Error = BinaryOrError.takeError()) {
+    return Error;
+  }
+  KernelBinary *Binary = *BinaryOrError;
+  RTCBundleBinaryInfo BBI{BinaryFormat::SPIRV,
+                          Mod.getDataLayout().getPointerSizeInBits(),
+                          Binary->address(), Binary->size()};
+  return BBI;
+}
+
 llvm::Expected<KernelBinary *>
 KernelTranslator::translateToSPIRV(llvm::Module &Mod, JITContext &JITCtx) {
   return SPIRVLLVMTranslator::translateLLVMtoSPIRV(Mod, JITCtx);
diff --git a/sycl-jit/jit-compiler/lib/translation/KernelTranslation.h b/sycl-jit/jit-compiler/lib/translation/KernelTranslation.h
index 809c8fab2e42f..e71aa2b0d19f3 100644
--- a/sycl-jit/jit-compiler/lib/translation/KernelTranslation.h
+++ b/sycl-jit/jit-compiler/lib/translation/KernelTranslation.h
@@ -27,6 +27,9 @@ class KernelTranslator {
   static llvm::Error translateKernel(SYCLKernelInfo &Kernel, llvm::Module &Mod,
                                      JITContext &JITCtx, BinaryFormat Format);
 
+  static llvm::Expected<RTCBundleBinaryInfo>
+  translateBundleToSPIRV(llvm::Module &Mod, JITContext &JITCtx);
+
 private:
   ///
   /// Pair of address and size to represent a binary blob.
diff --git a/sycl/source/detail/jit_compiler.cpp b/sycl/source/detail/jit_compiler.cpp
index 19f1915943f05..8e466b97d75cf 100644
--- a/sycl/source/detail/jit_compiler.cpp
+++ b/sycl/source/detail/jit_compiler.cpp
@@ -1119,6 +1119,44 @@ sycl_device_binaries jit_compiler::createPIDeviceBinary(
   return JITDeviceBinaries.back().getPIDeviceStruct();
 }
 
+sycl_device_binaries jit_compiler::createDeviceBinaryImage(
+    const ::jit_compiler::RTCBundleInfo &BundleInfo) {
+  DeviceBinaryContainer Binary;
+  for (const auto &Symbol : BundleInfo.SymbolTable) {
+    // Create an offload entry for each kernel.
+    // It seems to be OK to set zero for most of the information here, at least
+    // that is the case for compiled SPIR-V binaries.
+    OffloadEntryContainer Entry{Symbol.c_str(), /*Addr=*/nullptr, /*Size=*/0,
+                                /*Flags=*/0, /*Reserved=*/0};
+    Binary.addOffloadEntry(std::move(Entry));
+  }
+
+  for (const auto &FPS : BundleInfo.Properties) {
+    PropertySetContainer PropSet{FPS.Name.c_str()};
+    for (const auto &FPV : FPS.Values) {
+      if (FPV.IsUIntValue) {
+        PropSet.addProperty(PropertyContainer{FPV.Name.c_str(), FPV.UIntValue});
+      } else {
+        PropSet.addProperty(PropertyContainer{
+            FPV.Name.c_str(), FPV.Bytes.begin(), FPV.Bytes.size(),
+            sycl_property_type::SYCL_PROPERTY_TYPE_BYTE_ARRAY});
+      }
+    }
+    Binary.addProperty(std::move(PropSet));
+  }
+
+  DeviceBinariesCollection Collection;
+  Collection.addDeviceBinary(std::move(Binary),
+                             BundleInfo.BinaryInfo.BinaryStart,
+                             BundleInfo.BinaryInfo.BinarySize,
+                             (BundleInfo.BinaryInfo.AddressBits == 64)
+                                 ? __SYCL_DEVICE_BINARY_TARGET_SPIRV64
+                                 : __SYCL_DEVICE_BINARY_TARGET_SPIRV32,
+                             SYCL_DEVICE_BINARY_TYPE_SPIRV);
+  JITDeviceBinaries.push_back(std::move(Collection));
+  return JITDeviceBinaries.back().getPIDeviceStruct();
+}
+
 std::vector<uint8_t> jit_compiler::encodeArgUsageMask(
     const ::jit_compiler::ArgUsageMask &Mask) const {
   // This must match the decoding logic in program_manager.cpp.
@@ -1167,7 +1205,7 @@ std::vector<uint8_t> jit_compiler::encodeReqdWorkGroupSize(
   return Encoded;
 }
 
-std::vector<uint8_t> jit_compiler::compileSYCL(
+sycl_device_binaries jit_compiler::compileSYCL(
     const std::string &Id, const std::string &SYCLSource,
     const std::vector<std::pair<std::string, std::string>> &IncludePairs,
     const std::vector<std::string> &UserArgs, std::string *LogPtr,
@@ -1207,10 +1245,7 @@ std::vector<uint8_t> jit_compiler::compileSYCL(
   // TODO: We currently don't have a meaningful build log.
   (void)LogPtr;
 
-  const auto &BI = Result.getKernelInfo().BinaryInfo;
-  assert(BI.Format == ::jit_compiler::BinaryFormat::SPIRV);
-  std::vector<uint8_t> SPV(BI.BinaryStart, BI.BinaryStart + BI.BinarySize);
-  return SPV;
+  return createDeviceBinaryImage(Result.getBundleInfo());
 }
 
 } // namespace detail
diff --git a/sycl/source/detail/jit_compiler.hpp b/sycl/source/detail/jit_compiler.hpp
index 1908defa42e77..b673e4d37b8fa 100644
--- a/sycl/source/detail/jit_compiler.hpp
+++ b/sycl/source/detail/jit_compiler.hpp
@@ -23,6 +23,7 @@ enum class BinaryFormat : uint32_t;
 class JITContext;
 struct SYCLKernelInfo;
 struct SYCLKernelAttribute;
+struct RTCBundleInfo;
 template <typename T> class DynArray;
 using ArgUsageMask = DynArray<uint8_t>;
 using JITEnvVar = DynArray<char>;
@@ -44,7 +45,7 @@ class jit_compiler {
                            const std::string &KernelName,
                            const std::vector<unsigned char> &SpecConstBlob);
 
-  std::vector<uint8_t> compileSYCL(
+  sycl_device_binaries compileSYCL(
       const std::string &Id, const std::string &SYCLSource,
       const std::vector<std::pair<std::string, std::string>> &IncludePairs,
       const std::vector<std::string> &UserArgs, std::string *LogPtr,
@@ -69,6 +70,9 @@ class jit_compiler {
   createPIDeviceBinary(const ::jit_compiler::SYCLKernelInfo &FusedKernelInfo,
                        ::jit_compiler::BinaryFormat Format);
 
+  sycl_device_binaries
+  createDeviceBinaryImage(const ::jit_compiler::RTCBundleInfo &BundleInfo);
+
   std::vector<uint8_t>
   encodeArgUsageMask(const ::jit_compiler::ArgUsageMask &Mask) const;
 
diff --git a/sycl/source/detail/jit_device_binaries.cpp b/sycl/source/detail/jit_device_binaries.cpp
index b5f3b67c08ecd..f90be2c27ec3a 100644
--- a/sycl/source/detail/jit_device_binaries.cpp
+++ b/sycl/source/detail/jit_device_binaries.cpp
@@ -27,7 +27,7 @@ _sycl_offload_entry_struct OffloadEntryContainer::getPIOffloadEntry() {
                                     EntryFlags, EntryReserved};
 }
 
-PropertyContainer::PropertyContainer(const std::string &Name, void *Data,
+PropertyContainer::PropertyContainer(const std::string &Name, const void *Data,
                                      size_t Size, uint32_t Type)
     : PropName{new char[Name.length() + 1]}, Value{new unsigned char[Size]},
       ValueSize{Size}, PropType{Type} {
diff --git a/sycl/source/detail/jit_device_binaries.hpp b/sycl/source/detail/jit_device_binaries.hpp
index 49445e07920e7..bca83839f39e8 100644
--- a/sycl/source/detail/jit_device_binaries.hpp
+++ b/sycl/source/detail/jit_device_binaries.hpp
@@ -50,7 +50,7 @@ class OffloadEntryContainer {
 class PropertyContainer {
 
 public:
-  PropertyContainer(const std::string &Name, void *Data, size_t Size,
+  PropertyContainer(const std::string &Name, const void *Data, size_t Size,
                     uint32_t Type);
   // Set a UR_PROPERTY_TYPE_UINT32 property
   PropertyContainer(const std::string &Name, uint32_t Data);
diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp
index 7ce5971711291..58e605e85c458 100644
--- a/sycl/source/detail/kernel_bundle_impl.hpp
+++ b/sycl/source/detail/kernel_bundle_impl.hpp
@@ -499,10 +499,13 @@ class kernel_bundle_impl {
                                                RegisteredKernelNames);
         }
         if (Language == syclex::source_language::sycl_jit) {
-          const auto &SourceStr = std::get<std::string>(this->Source);
-          return syclex::detail::SYCL_JIT_to_SPIRV(SourceStr, IncludePairs,
-                                                   BuildOptions, LogPtr,
-                                                   RegisteredKernelNames);
+          auto *Binaries = syclex::detail::SYCL_JIT_to_SPIRV(
+              *SourceStrPtr, IncludePairs, BuildOptions, LogPtr,
+              RegisteredKernelNames);
+          assert(Binaries->NumDeviceBinaries == 1 &&
+                 "Device code splitting is not yet supported");
+          return std::vector<uint8_t>(Binaries->DeviceBinaries->BinaryStart,
+                                      Binaries->DeviceBinaries->BinaryEnd);
         }
         throw sycl::exception(
             make_error_code(errc::invalid),
diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp
index 89a8a548a6d8a..6362bf355cfc5 100644
--- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp
+++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp
@@ -342,7 +342,7 @@ bool SYCL_JIT_Compilation_Available() {
 #endif
 }
 
-spirv_vec_t SYCL_JIT_to_SPIRV(
+sycl_device_binaries SYCL_JIT_to_SPIRV(
     [[maybe_unused]] const std::string &SYCLSource,
     [[maybe_unused]] include_pairs_t IncludePairs,
     [[maybe_unused]] const std::vector<std::string> &UserArgs,
diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp
index 948f199c1b3cc..eee0f8a78f602 100644
--- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp
+++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp
@@ -12,6 +12,8 @@
 #include <sycl/detail/export.hpp> // __SYCL_EXPORT
 #include <sycl/device.hpp>
 
+#include <detail/compiler.hpp> // sycl_device_binaries
+
 #include <numeric> // std::accumulate
 #include <string>
 #include <vector>
@@ -33,7 +35,7 @@ bool SYCL_Compilation_Available();
 
 std::string userArgsAsString(const std::vector<std::string> &UserArguments);
 
-spirv_vec_t
+sycl_device_binaries
 SYCL_JIT_to_SPIRV(const std::string &Source, include_pairs_t IncludePairs,
                   const std::vector<std::string> &UserArgs, std::string *LogPtr,
                   const std::vector<std::string> &RegisteredKernelNames);
diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp
index 01f25f813b826..8977caf7d82d4 100644
--- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp
+++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp
@@ -179,6 +179,8 @@ int test_unsupported_options() {
   CheckUnsupported({"-Xsycl-target-frontend", "-fsanitize=address"});
   CheckUnsupported({"-Xsycl-target-frontend=spir64", "-fsanitize=address"});
   CheckUnsupported({"-Xarch_device", "-fsanitize=address"});
+  CheckUnsupported({"-fsycl-device-code-split=kernel"});
+  CheckUnsupported({"-fsycl-device-code-split-esimd"});
 
   return 0;
 }