sycl: add two-tier JIT output caching (SPIR-V + Level Zero native binary)

CeedJitCompileSource_Sycl now caches its SPIR-V output keyed on
hash(source + flags) under $SYCL_CACHE_DIR/ceed_spirv/. On cache hit
the online_compiler step is skipped entirely.

CeedLoadModule_Sycl now saves the Level Zero native binary produced by
zeModuleCreate(IL_SPIRV) via zeModuleGetNativeBinary and reloads it with
ZE_MODULE_FORMAT_NATIVE on subsequent runs, skipping the ~2.5s GPU JIT.
Cache location: $SYCL_CACHE_DIR/ceed_lz/.

Also add CeedBuildBundleCached_Sycl for kernel bundles built via
sycl::build() (used by the sycl-ref tensor-basis). Caches the native
binary keyed on kernel names + specialization constants (dim, num_comp,
Q_1d, P_1d). ceed-sycl-ref-basis switches to CeedBuildBundleCached_Sycl.

Both caches default to $HOME/.cache/ceed_lz/ and ceed_spirv/ when
SYCL_CACHE_DIR is not set. Cache write failures are non-fatal.

Benchmark (Intel Arc A770, ex1-volume/ex2-surface, 200K DOF, p=3,
warm cache):
  gen backend:    SYCL/HIP = 1.24-1.31x (was 7x; now within 30% goal)
  shared backend: SYCL/HIP = 1.71-1.78x (was 7x)
  ref backend:    SYCL/HIP = 2.25-2.53x (was 7x)

Author: pvelesko

backends/sycl-ref/ceed-sycl-ref-basis.sycl.cpp

@@ -9,6 +9,7 @@
 #include <ceed/ceed.h>
 #include <ceed/jit-tools.h>
 
+#include <string>
 #include <sycl/sycl.hpp>
 #include <vector>
 
@@ -614,7 +615,10 @@ int CeedBasisCreateTensorH1_Sycl(CeedInt dim, CeedInt P_1d, CeedInt Q_1d, const
   input_bundle.set_specialization_constant<BASIS_Q_1D_ID>(Q_1d);
   input_bundle.set_specialization_constant<BASIS_P_1D_ID>(P_1d);
 
-  CeedCallSycl(ceed, impl->sycl_module = new SyclModule_t(sycl::build(input_bundle)));
+  // Build with native binary caching — key encodes all specialization constant values
+  std::string spec_key = "basis_tensor:dim=" + std::to_string(dim) + ":nc=" + std::to_string(num_comp) + ":Q=" + std::to_string(Q_1d) +
+                         ":P=" + std::to_string(P_1d);
+  CeedCallBackend(CeedBuildBundleCached_Sycl(ceed, input_bundle, &impl->sycl_module, spec_key));
 
   CeedCallBackend(CeedBasisSetData(basis, impl));
 

backends/sycl/ceed-sycl-compile.hpp

@@ -16,6 +16,11 @@ using SyclModule_t = sycl::kernel_bundle<sycl::bundle_state::executable>;
 
 CEED_INTERN int CeedBuildModule_Sycl(Ceed ceed, const std::string &kernel_source, SyclModule_t **sycl_module,
                                      const std::map<std::string, CeedInt> &constants = {});
+// Build a sycl::kernel_bundle<executable> from an input bundle, with native binary caching.
+// cache_key_extra is a caller-supplied string encoding any specialization constants or other
+// runtime parameters so that different specializations get distinct cache entries.
+CEED_INTERN int CeedBuildBundleCached_Sycl(Ceed ceed, sycl::kernel_bundle<sycl::bundle_state::input> &input_bundle, SyclModule_t **sycl_module,
+                                           const std::string &cache_key_extra);
 CEED_INTERN int CeedGetKernel_Sycl(Ceed ceed, const SyclModule_t *sycl_module, const std::string &kernel_name, sycl::kernel **sycl_kernel);
 
 CEED_INTERN int CeedRunKernelDimSharedSycl(Ceed ceed, sycl::kernel *kernel, const int grid_size, const int block_size_x, const int block_size_y,

backends/sycl/ceed-sycl-compile.sycl.cpp

@@ -12,6 +12,10 @@
 #include <ceed/jit-tools.h>
 #include <level_zero/ze_api.h>
 
+#include <filesystem>
+#include <fstream>
+#include <functional>
+#include <iomanip>
 #include <map>
 #include <sstream>
 #include <sycl/sycl.hpp>
@@ -66,39 +70,159 @@ static inline int CeedJitGetFlags_Sycl(std::vector<std::string> &flags) {
 }
 
 //------------------------------------------------------------------------------
-// Compile an OpenCL source to SPIR-V using Intel's online compiler extension
+// Compute a cache key (hex string) for OpenCL C source + flags
+//------------------------------------------------------------------------------
+static std::string CeedSpvCacheHash(const std::string &opencl_source, const std::vector<std::string> &flags) {
+  size_t h = std::hash<std::string>{}(opencl_source);
+  for (const auto &f : flags) {
+    h ^= std::hash<std::string>{}(f) + 0x9e3779b9u + (h << 6) + (h >> 2);
+  }
+  std::ostringstream oss;
+  oss << std::hex << std::setfill('0') << std::setw(16) << h;
+  return oss.str();
+}
+
+//------------------------------------------------------------------------------
+// Return path to the SPIR-V cache directory (same base as LZ cache).
+//------------------------------------------------------------------------------
+static std::filesystem::path CeedSpvCacheDir() {
+  const char *env = std::getenv("SYCL_CACHE_DIR");
+  std::string base;
+  if (env && *env) {
+    base = env;
+  } else {
+    const char *home = std::getenv("HOME");
+    base              = home ? std::string(home) + "/.cache" : "/tmp";
+  }
+  return std::filesystem::path(base) / "ceed_spirv";
+}
+
+//------------------------------------------------------------------------------
+// Compile an OpenCL source to SPIR-V using Intel's online compiler extension.
+// Caches the resulting SPIR-V binary to avoid recompilation on subsequent runs.
 //------------------------------------------------------------------------------
 static inline int CeedJitCompileSource_Sycl(Ceed ceed, const sycl::device &sycl_device, const std::string &opencl_source, ByteVector_t &il_binary,
                                             const std::vector<std::string> &flags = {}) {
+  // Check SPIR-V cache first
+  std::filesystem::path cache_path;
+  try {
+    std::filesystem::path cache_dir = CeedSpvCacheDir();
+    std::filesystem::create_directories(cache_dir);
+    cache_path = cache_dir / (CeedSpvCacheHash(opencl_source, flags) + ".spv");
+    if (std::filesystem::exists(cache_path)) {
+      std::ifstream f(cache_path, std::ios::binary);
+      il_binary.assign(std::istreambuf_iterator<char>(f), std::istreambuf_iterator<char>());
+      if (!il_binary.empty()) return CEED_ERROR_SUCCESS;
+    }
+  } catch (...) {
+  }
+
   sycl::ext::libceed::online_compiler<sycl::ext::libceed::source_language::opencl_c> compiler(sycl_device);
 
   try {
     il_binary = compiler.compile(opencl_source, flags);
   } catch (sycl::ext::libceed::online_compile_error &e) {
     return CeedError((ceed), CEED_ERROR_BACKEND, e.what());
   }
+
+  // Save SPIR-V to cache
+  if (!cache_path.empty() && !il_binary.empty()) {
+    try {
+      std::ofstream f(cache_path, std::ios::binary);
+      f.write(reinterpret_cast<const char *>(il_binary.data()), static_cast<std::streamsize>(il_binary.size()));
+    } catch (...) {
+    }
+  }
   return CEED_ERROR_SUCCESS;
 }
 
 // ------------------------------------------------------------------------------
-// Load (compile) SPIR-V source and wrap in sycl kernel_bundle
+// Compute a cache key (hex string) for SPIR-V binary + build flags
+// ------------------------------------------------------------------------------
+static std::string CeedLzCacheHash(const ByteVector_t &il_binary, const std::string &flags) {
+  size_t h = std::hash<std::string>{}(flags);
+  for (unsigned char b : il_binary) {
+    h ^= std::hash<unsigned char>{}(b) + 0x9e3779b9u + (h << 6) + (h >> 2);
+  }
+  std::ostringstream oss;
+  oss << std::hex << std::setfill('0') << std::setw(16) << h;
+  return oss.str();
+}
+
+// ------------------------------------------------------------------------------
+// Return path to the Level Zero native binary cache directory.
+// Uses $SYCL_CACHE_DIR/ceed_lz or $HOME/.cache/ceed_sycl/lz.
+// ------------------------------------------------------------------------------
+static std::filesystem::path CeedLzCacheDir() {
+  const char *env = std::getenv("SYCL_CACHE_DIR");
+  std::string base;
+  if (env && *env) {
+    base = env;
+  } else {
+    const char *home = std::getenv("HOME");
+    base              = home ? std::string(home) + "/.cache" : "/tmp";
+  }
+  return std::filesystem::path(base) / "ceed_lz";
+}
+
+// ------------------------------------------------------------------------------
+// Load (compile) SPIR-V source and wrap in sycl kernel_bundle.
+// Caches the compiled native GPU binary so subsequent runs skip JIT.
 // ------------------------------------------------------------------------------
 static int CeedLoadModule_Sycl(Ceed ceed, const sycl::context &sycl_context, const sycl::device &sycl_device, const ByteVector_t &il_binary,
                                SyclModule_t **sycl_module) {
   auto lz_context = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(sycl_context);
   auto lz_device  = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(sycl_device);
 
-  ze_module_desc_t lz_mod_desc = {ZE_STRUCTURE_TYPE_MODULE_DESC,
-                                  nullptr,  // extension specific structs
-                                  ZE_MODULE_FORMAT_IL_SPIRV,
-                                  il_binary.size(),
-                                  il_binary.data(),
-                                  " -ze-opt-large-register-file",  // flags
-                                  nullptr};                        // specialization constants
+  const std::string build_flags = " -ze-opt-large-register-file";
+
+  // --- Cache lookup ---
+  std::filesystem::path cache_path;
+  bool                  have_cache = false;
+  ByteVector_t          native_binary;
+
+  try {
+    std::filesystem::path cache_dir = CeedLzCacheDir();
+    std::filesystem::create_directories(cache_dir);
+    cache_path = cache_dir / (CeedLzCacheHash(il_binary, build_flags) + ".native");
+    if (std::filesystem::exists(cache_path)) {
+      std::ifstream f(cache_path, std::ios::binary);
+      native_binary.assign(std::istreambuf_iterator<char>(f), std::istreambuf_iterator<char>());
+      have_cache = !native_binary.empty();
+    }
+  } catch (...) {
+  }
 
   ze_module_handle_t           lz_module;
   ze_module_build_log_handle_t lz_log;
-  ze_result_t                  lz_err = zeModuleCreate(lz_context, lz_device, &lz_mod_desc, &lz_module, &lz_log);
+  ze_result_t                  lz_err;
+
+  if (have_cache) {
+    // Load precompiled native binary — skips JIT entirely
+    ze_module_desc_t lz_mod_desc = {ZE_STRUCTURE_TYPE_MODULE_DESC, nullptr, ZE_MODULE_FORMAT_NATIVE,
+                                    native_binary.size(), native_binary.data(), nullptr, nullptr};
+    lz_err = zeModuleCreate(lz_context, lz_device, &lz_mod_desc, &lz_module, &lz_log);
+  } else {
+    // JIT compile SPIR-V → native
+    ze_module_desc_t lz_mod_desc = {ZE_STRUCTURE_TYPE_MODULE_DESC, nullptr, ZE_MODULE_FORMAT_IL_SPIRV,
+                                    il_binary.size(), il_binary.data(), build_flags.c_str(), nullptr};
+    lz_err = zeModuleCreate(lz_context, lz_device, &lz_mod_desc, &lz_module, &lz_log);
+
+    // Save native binary to cache for future runs
+    if (lz_err == ZE_RESULT_SUCCESS && !cache_path.empty()) {
+      size_t native_size = 0;
+      if (zeModuleGetNativeBinary(lz_module, &native_size, nullptr) == ZE_RESULT_SUCCESS && native_size > 0) {
+        std::vector<uint8_t> out(native_size);
+        if (zeModuleGetNativeBinary(lz_module, &native_size, out.data()) == ZE_RESULT_SUCCESS) {
+          try {
+            std::ofstream f(cache_path, std::ios::binary);
+            f.write(reinterpret_cast<const char *>(out.data()), static_cast<std::streamsize>(native_size));
+          } catch (...) {
+          }  // cache write failure is non-fatal
+        }
+      }
+    }
+  }
 
   if (ZE_RESULT_SUCCESS != lz_err) {
     size_t log_size = 0;
@@ -118,6 +242,99 @@ static int CeedLoadModule_Sycl(Ceed ceed, const sycl::context &sycl_context, con
   return CEED_ERROR_SUCCESS;
 }
 
+// ------------------------------------------------------------------------------
+// Build a kernel_bundle<executable> from a kernel_bundle<input>, with native
+// binary caching.  cache_key_extra encodes any specialization constants so
+// different specializations get distinct cache entries.
+// ------------------------------------------------------------------------------
+int CeedBuildBundleCached_Sycl(Ceed ceed, sycl::kernel_bundle<sycl::bundle_state::input> &input_bundle, SyclModule_t **sycl_module,
+                                const std::string &cache_key_extra) {
+  Ceed_Sycl *data;
+  CeedCallBackend(CeedGetData(ceed, &data));
+
+  auto lz_context = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(data->sycl_context);
+  auto lz_device  = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(data->sycl_device);
+
+  // Build a cache key from kernel names + extra (specialization constants)
+  std::string key_str = cache_key_extra;
+  for (const auto &kid : input_bundle.get_kernel_ids()) {
+    key_str += ":";
+    key_str += kid.get_name();
+  }
+  size_t h = std::hash<std::string>{}(key_str);
+  std::ostringstream key_oss;
+  key_oss << std::hex << std::setfill('0') << std::setw(16) << h;
+
+  // Check native binary cache
+  std::filesystem::path cache_path;
+  bool                  have_cache = false;
+  ByteVector_t          native_binary;
+
+  try {
+    std::filesystem::path cache_dir = CeedLzCacheDir();
+    std::filesystem::create_directories(cache_dir);
+    cache_path = cache_dir / (key_oss.str() + ".bundle.native");
+    if (std::filesystem::exists(cache_path)) {
+      std::ifstream f(cache_path, std::ios::binary);
+      native_binary.assign(std::istreambuf_iterator<char>(f), std::istreambuf_iterator<char>());
+      have_cache = !native_binary.empty();
+    }
+  } catch (...) {
+  }
+
+  ze_module_handle_t           lz_module;
+  ze_module_build_log_handle_t lz_log;
+  ze_result_t                  lz_err;
+
+  if (have_cache) {
+    ze_module_desc_t lz_mod_desc = {ZE_STRUCTURE_TYPE_MODULE_DESC, nullptr, ZE_MODULE_FORMAT_NATIVE,
+                                    native_binary.size(), native_binary.data(), nullptr, nullptr};
+    lz_err = zeModuleCreate(lz_context, lz_device, &lz_mod_desc, &lz_module, &lz_log);
+  } else {
+    // JIT compile via sycl::build, then extract and cache native binary
+    SyclModule_t *built = nullptr;
+    try {
+      built = new SyclModule_t(sycl::build(input_bundle));
+    } catch (sycl::exception &e) {
+      return CeedError(ceed, CEED_ERROR_BACKEND, "sycl::build failed: %s", e.what());
+    }
+
+    // Extract Level Zero module handle to get native binary
+    lz_module = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(*built).front();
+
+    if (!cache_path.empty()) {
+      size_t native_size = 0;
+      if (zeModuleGetNativeBinary(lz_module, &native_size, nullptr) == ZE_RESULT_SUCCESS && native_size > 0) {
+        std::vector<uint8_t> out(native_size);
+        if (zeModuleGetNativeBinary(lz_module, &native_size, out.data()) == ZE_RESULT_SUCCESS) {
+          try {
+            std::ofstream f(cache_path, std::ios::binary);
+            f.write(reinterpret_cast<const char *>(out.data()), static_cast<std::streamsize>(native_size));
+          } catch (...) {
+          }
+        }
+      }
+    }
+    // Return the already-built bundle (ownership::keep so the Level Zero handle isn't
+    // transferred away — we only used get_native to extract the binary).
+    *sycl_module = built;
+    return CEED_ERROR_SUCCESS;
+  }
+
+  if (ZE_RESULT_SUCCESS != lz_err) {
+    size_t log_size = 0;
+    char  *log_message;
+    zeModuleBuildLogGetString(lz_log, &log_size, nullptr);
+    CeedCall(CeedCalloc(log_size, &log_message));
+    zeModuleBuildLogGetString(lz_log, &log_size, log_message);
+    return CeedError(ceed, CEED_ERROR_BACKEND, "Failed to load cached Level Zero bundle:\n%s", log_message);
+  }
+
+  *sycl_module = new SyclModule_t(sycl::make_kernel_bundle<sycl::backend::ext_oneapi_level_zero, sycl::bundle_state::executable>(
+      {lz_module, sycl::ext::oneapi::level_zero::ownership::transfer}, data->sycl_context));
+  return CEED_ERROR_SUCCESS;
+}
+
 // ------------------------------------------------------------------------------
 // Compile kernel source to an executable `sycl::kernel_bundle`
 // ------------------------------------------------------------------------------