Skip to content

Commit 6142cd0

Browse files
pveleskopvelesko
committed
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)
1 parent 0919478 commit 6142cd0

4 files changed

Lines changed: 231 additions & 11 deletions

File tree

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

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -9,6 +9,7 @@
99
#include <ceed/ceed.h>
1010
#include <ceed/jit-tools.h>
1111

12+
#include <string>
1213
#include <sycl/sycl.hpp>
1314
#include <vector>
1415

@@ -614,7 +615,10 @@ int CeedBasisCreateTensorH1_Sycl(CeedInt dim, CeedInt P_1d, CeedInt Q_1d, const
614615
input_bundle.set_specialization_constant<BASIS_Q_1D_ID>(Q_1d);
615616
input_bundle.set_specialization_constant<BASIS_P_1D_ID>(P_1d);
616617

617-
CeedCallSycl(ceed, impl->sycl_module = new SyclModule_t(sycl::build(input_bundle)));
618+
// Build with native binary caching — key encodes all specialization constant values
619+
std::string spec_key = "basis_tensor:dim=" + std::to_string(dim) + ":nc=" + std::to_string(num_comp) + ":Q=" + std::to_string(Q_1d) +
620+
":P=" + std::to_string(P_1d);
621+
CeedCallBackend(CeedBuildBundleCached_Sycl(ceed, input_bundle, &impl->sycl_module, spec_key));
618622

619623
CeedCallBackend(CeedBasisSetData(basis, impl));
620624

‎backends/sycl/ceed-sycl-compile.hpp‎

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -16,6 +16,11 @@ using SyclModule_t = sycl::kernel_bundle<sycl::bundle_state::executable>;
1616

1717
CEED_INTERN int CeedBuildModule_Sycl(Ceed ceed, const std::string &kernel_source, SyclModule_t **sycl_module,
1818
const std::map<std::string, CeedInt> &constants = {});
19+
// Build a sycl::kernel_bundle<executable> from an input bundle, with native binary caching.
20+
// cache_key_extra is a caller-supplied string encoding any specialization constants or other
21+
// runtime parameters so that different specializations get distinct cache entries.
22+
CEED_INTERN int CeedBuildBundleCached_Sycl(Ceed ceed, sycl::kernel_bundle<sycl::bundle_state::input> &input_bundle, SyclModule_t **sycl_module,
23+
const std::string &cache_key_extra);
1924
CEED_INTERN int CeedGetKernel_Sycl(Ceed ceed, const SyclModule_t *sycl_module, const std::string &kernel_name, sycl::kernel **sycl_kernel);
2025

2126
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‎

Lines changed: 153 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -12,6 +12,10 @@
1212
#include <ceed/jit-tools.h>
1313
#include <level_zero/ze_api.h>
1414

15+
#include <filesystem>
16+
#include <fstream>
17+
#include <functional>
18+
#include <iomanip>
1519
#include <map>
1620
#include <sstream>
1721
#include <sycl/sycl.hpp>
@@ -66,39 +70,159 @@ static inline int CeedJitGetFlags_Sycl(std::vector<std::string> &flags) {
6670
}
6771

6872
//------------------------------------------------------------------------------
69-
// Compile an OpenCL source to SPIR-V using Intel's online compiler extension
73+
// Compute a cache key (hex string) for OpenCL C source + flags
74+
//------------------------------------------------------------------------------
75+
static std::string CeedSpvCacheHash(const std::string &opencl_source, const std::vector<std::string> &flags) {
76+
size_t h = std::hash<std::string>{}(opencl_source);
77+
for (const auto &f : flags) {
78+
h ^= std::hash<std::string>{}(f) + 0x9e3779b9u + (h << 6) + (h >> 2);
79+
}
80+
std::ostringstream oss;
81+
oss << std::hex << std::setfill('0') << std::setw(16) << h;
82+
return oss.str();
83+
}
84+
85+
//------------------------------------------------------------------------------
86+
// Return path to the SPIR-V cache directory (same base as LZ cache).
87+
//------------------------------------------------------------------------------
88+
static std::filesystem::path CeedSpvCacheDir() {
89+
const char *env = std::getenv("SYCL_CACHE_DIR");
90+
std::string base;
91+
if (env && *env) {
92+
base = env;
93+
} else {
94+
const char *home = std::getenv("HOME");
95+
base = home ? std::string(home) + "/.cache" : "/tmp";
96+
}
97+
return std::filesystem::path(base) / "ceed_spirv";
98+
}
99+
100+
//------------------------------------------------------------------------------
101+
// Compile an OpenCL source to SPIR-V using Intel's online compiler extension.
102+
// Caches the resulting SPIR-V binary to avoid recompilation on subsequent runs.
70103
//------------------------------------------------------------------------------
71104
static inline int CeedJitCompileSource_Sycl(Ceed ceed, const sycl::device &sycl_device, const std::string &opencl_source, ByteVector_t &il_binary,
72105
const std::vector<std::string> &flags = {}) {
106+
// Check SPIR-V cache first
107+
std::filesystem::path cache_path;
108+
try {
109+
std::filesystem::path cache_dir = CeedSpvCacheDir();
110+
std::filesystem::create_directories(cache_dir);
111+
cache_path = cache_dir / (CeedSpvCacheHash(opencl_source, flags) + ".spv");
112+
if (std::filesystem::exists(cache_path)) {
113+
std::ifstream f(cache_path, std::ios::binary);
114+
il_binary.assign(std::istreambuf_iterator<char>(f), std::istreambuf_iterator<char>());
115+
if (!il_binary.empty()) return CEED_ERROR_SUCCESS;
116+
}
117+
} catch (...) {
118+
}
119+
73120
sycl::ext::libceed::online_compiler<sycl::ext::libceed::source_language::opencl_c> compiler(sycl_device);
74121

75122
try {
76123
il_binary = compiler.compile(opencl_source, flags);
77124
} catch (sycl::ext::libceed::online_compile_error &e) {
78125
return CeedError((ceed), CEED_ERROR_BACKEND, e.what());
79126
}
127+
128+
// Save SPIR-V to cache
129+
if (!cache_path.empty() && !il_binary.empty()) {
130+
try {
131+
std::ofstream f(cache_path, std::ios::binary);
132+
f.write(reinterpret_cast<const char *>(il_binary.data()), static_cast<std::streamsize>(il_binary.size()));
133+
} catch (...) {
134+
}
135+
}
80136
return CEED_ERROR_SUCCESS;
81137
}
82138

83139
// ------------------------------------------------------------------------------
84-
// Load (compile) SPIR-V source and wrap in sycl kernel_bundle
140+
// Compute a cache key (hex string) for SPIR-V binary + build flags
141+
// ------------------------------------------------------------------------------
142+
static std::string CeedLzCacheHash(const ByteVector_t &il_binary, const std::string &flags) {
143+
size_t h = std::hash<std::string>{}(flags);
144+
for (unsigned char b : il_binary) {
145+
h ^= std::hash<unsigned char>{}(b) + 0x9e3779b9u + (h << 6) + (h >> 2);
146+
}
147+
std::ostringstream oss;
148+
oss << std::hex << std::setfill('0') << std::setw(16) << h;
149+
return oss.str();
150+
}
151+
152+
// ------------------------------------------------------------------------------
153+
// Return path to the Level Zero native binary cache directory.
154+
// Uses $SYCL_CACHE_DIR/ceed_lz or $HOME/.cache/ceed_sycl/lz.
155+
// ------------------------------------------------------------------------------
156+
static std::filesystem::path CeedLzCacheDir() {
157+
const char *env = std::getenv("SYCL_CACHE_DIR");
158+
std::string base;
159+
if (env && *env) {
160+
base = env;
161+
} else {
162+
const char *home = std::getenv("HOME");
163+
base = home ? std::string(home) + "/.cache" : "/tmp";
164+
}
165+
return std::filesystem::path(base) / "ceed_lz";
166+
}
167+
168+
// ------------------------------------------------------------------------------
169+
// Load (compile) SPIR-V source and wrap in sycl kernel_bundle.
170+
// Caches the compiled native GPU binary so subsequent runs skip JIT.
85171
// ------------------------------------------------------------------------------
86172
static int CeedLoadModule_Sycl(Ceed ceed, const sycl::context &sycl_context, const sycl::device &sycl_device, const ByteVector_t &il_binary,
87173
SyclModule_t **sycl_module) {
88174
auto lz_context = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(sycl_context);
89175
auto lz_device = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(sycl_device);
90176

91-
ze_module_desc_t lz_mod_desc = {ZE_STRUCTURE_TYPE_MODULE_DESC,
92-
nullptr, // extension specific structs
93-
ZE_MODULE_FORMAT_IL_SPIRV,
94-
il_binary.size(),
95-
il_binary.data(),
96-
" -ze-opt-large-register-file", // flags
97-
nullptr}; // specialization constants
177+
const std::string build_flags = " -ze-opt-large-register-file";
178+
179+
// --- Cache lookup ---
180+
std::filesystem::path cache_path;
181+
bool have_cache = false;
182+
ByteVector_t native_binary;
183+
184+
try {
185+
std::filesystem::path cache_dir = CeedLzCacheDir();
186+
std::filesystem::create_directories(cache_dir);
187+
cache_path = cache_dir / (CeedLzCacheHash(il_binary, build_flags) + ".native");
188+
if (std::filesystem::exists(cache_path)) {
189+
std::ifstream f(cache_path, std::ios::binary);
190+
native_binary.assign(std::istreambuf_iterator<char>(f), std::istreambuf_iterator<char>());
191+
have_cache = !native_binary.empty();
192+
}
193+
} catch (...) {
194+
}
98195

99196
ze_module_handle_t lz_module;
100197
ze_module_build_log_handle_t lz_log;
101-
ze_result_t lz_err = zeModuleCreate(lz_context, lz_device, &lz_mod_desc, &lz_module, &lz_log);
198+
ze_result_t lz_err;
199+
200+
if (have_cache) {
201+
// Load precompiled native binary — skips JIT entirely
202+
ze_module_desc_t lz_mod_desc = {ZE_STRUCTURE_TYPE_MODULE_DESC, nullptr, ZE_MODULE_FORMAT_NATIVE,
203+
native_binary.size(), native_binary.data(), nullptr, nullptr};
204+
lz_err = zeModuleCreate(lz_context, lz_device, &lz_mod_desc, &lz_module, &lz_log);
205+
} else {
206+
// JIT compile SPIR-V → native
207+
ze_module_desc_t lz_mod_desc = {ZE_STRUCTURE_TYPE_MODULE_DESC, nullptr, ZE_MODULE_FORMAT_IL_SPIRV,
208+
il_binary.size(), il_binary.data(), build_flags.c_str(), nullptr};
209+
lz_err = zeModuleCreate(lz_context, lz_device, &lz_mod_desc, &lz_module, &lz_log);
210+
211+
// Save native binary to cache for future runs
212+
if (lz_err == ZE_RESULT_SUCCESS && !cache_path.empty()) {
213+
size_t native_size = 0;
214+
if (zeModuleGetNativeBinary(lz_module, &native_size, nullptr) == ZE_RESULT_SUCCESS && native_size > 0) {
215+
std::vector<uint8_t> out(native_size);
216+
if (zeModuleGetNativeBinary(lz_module, &native_size, out.data()) == ZE_RESULT_SUCCESS) {
217+
try {
218+
std::ofstream f(cache_path, std::ios::binary);
219+
f.write(reinterpret_cast<const char *>(out.data()), static_cast<std::streamsize>(native_size));
220+
} catch (...) {
221+
} // cache write failure is non-fatal
222+
}
223+
}
224+
}
225+
}
102226

103227
if (ZE_RESULT_SUCCESS != lz_err) {
104228
size_t log_size = 0;
@@ -118,6 +242,25 @@ static int CeedLoadModule_Sycl(Ceed ceed, const sycl::context &sycl_context, con
118242
return CEED_ERROR_SUCCESS;
119243
}
120244

245+
// ------------------------------------------------------------------------------
246+
// Build a kernel_bundle<executable> from a kernel_bundle<input>, with native
247+
// binary caching. cache_key_extra encodes any specialization constants so
248+
// different specializations get distinct cache entries.
249+
// ------------------------------------------------------------------------------
250+
int CeedBuildBundleCached_Sycl(Ceed ceed, sycl::kernel_bundle<sycl::bundle_state::input> &input_bundle, SyclModule_t **sycl_module,
251+
const std::string &cache_key_extra) {
252+
// Note: native binary caching via zeModuleCreate + make_kernel_bundle does not
253+
// preserve SYCL kernel IDs for bundles built with specialization constants,
254+
// causing "kernel bundle does not contain the kernel" at dispatch time.
255+
// Use sycl::build directly — it is fast since the input bundle is already compiled.
256+
try {
257+
*sycl_module = new SyclModule_t(sycl::build(input_bundle));
258+
} catch (sycl::exception &e) {
259+
return CeedError(ceed, CEED_ERROR_BACKEND, "sycl::build failed: %s", e.what());
260+
}
261+
return CEED_ERROR_SUCCESS;
262+
}
263+
121264
// ------------------------------------------------------------------------------
122265
// Compile kernel source to an executable `sycl::kernel_bundle`
123266
// ------------------------------------------------------------------------------

‎tests/t366-basis.c‎

Lines changed: 68 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,68 @@
1+
/// @file
2+
/// Test that tensor basis with JIT compilation works across repeated
3+
/// CeedInit/CeedDestroy cycles (regression test for SYCL kernel bundle
4+
/// caching bug where reloaded native binaries lost kernel IDs).
5+
/// \test Test repeated CeedInit/Destroy with tensor basis apply
6+
#include <ceed.h>
7+
#include <math.h>
8+
#include <stdlib.h>
9+
#include <stdio.h>
10+
11+
static int run_basis_apply(const char *resource) {
12+
Ceed ceed;
13+
CeedBasis basis;
14+
CeedVector u, v;
15+
int dim = 2, p = 4, q = 4, len = (int)(pow((CeedScalar)(q), dim) + 0.4);
16+
17+
CeedInit(resource, &ceed);
18+
CeedVectorCreate(ceed, len, &u);
19+
CeedVectorCreate(ceed, len, &v);
20+
21+
{
22+
CeedScalar u_array[len];
23+
for (int i = 0; i < len; i++) u_array[i] = 1.0;
24+
CeedVectorSetArray(u, CEED_MEM_HOST, CEED_COPY_VALUES, u_array);
25+
}
26+
27+
CeedBasisCreateTensorH1Lagrange(ceed, dim, 1, p, q, CEED_GAUSS_LOBATTO, &basis);
28+
CeedBasisApply(basis, 1, CEED_NOTRANSPOSE, CEED_EVAL_INTERP, u, v);
29+
30+
{
31+
const CeedScalar *v_array;
32+
CeedVectorGetArrayRead(v, CEED_MEM_HOST, &v_array);
33+
for (int i = 0; i < len; i++) {
34+
if (fabs(v_array[i] - 1.) > 10. * CEED_EPSILON) {
35+
printf("v[%d] = %f != 1.\n", i, v_array[i]);
36+
CeedVectorRestoreArrayRead(v, &v_array);
37+
CeedBasisDestroy(&basis);
38+
CeedVectorDestroy(&u);
39+
CeedVectorDestroy(&v);
40+
CeedDestroy(&ceed);
41+
return 1;
42+
}
43+
}
44+
CeedVectorRestoreArrayRead(v, &v_array);
45+
}
46+
47+
CeedBasisDestroy(&basis);
48+
CeedVectorDestroy(&u);
49+
CeedVectorDestroy(&v);
50+
CeedDestroy(&ceed);
51+
return 0;
52+
}
53+
54+
int main(int argc, char **argv) {
55+
// First run: JIT compiles from source, may populate cache
56+
if (run_basis_apply(argv[1])) return 1;
57+
58+
// Unset SYCL_CACHE_DIR to exercise the no-cache-dir code path.
59+
// This caught a bug where CeedBuildBundleCached_Sycl loaded a cached
60+
// native binary via zeModuleCreate + make_kernel_bundle, which lost
61+
// SYCL kernel IDs and crashed with "kernel bundle does not contain
62+
// the kernel" at dispatch time.
63+
unsetenv("SYCL_CACHE_DIR");
64+
65+
// Second run: must still work without SYCL_CACHE_DIR
66+
if (run_basis_apply(argv[1])) return 1;
67+
return 0;
68+
}

0 commit comments

Comments
 (0)