diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml
index 164c2d9d..a36e601f 100644
--- a/.github/workflows/ci.yml
+++ b/.github/workflows/ci.yml
@@ -18,6 +18,13 @@ jobs:
     continue-on-error: true
     env:
       GITHUB_TOKEN: ${{ secrets.GITHUB_TOKEN }}
+      # Pin each parallel test worker to a distinct GPU tile instead of
+      # oversubscribing device 0 (see test/runtests.jl).
+      ONEAPI_TEST_SPREAD_GPUS: '1'
+      ONEAPI_LTS: '1'
+      # Synchronize after every command-list submission to work around the
+      # Aurora LTS NEO dropped-tail corruption (see lib/level-zero/cmdlist.jl).
+      ONEAPI_SYNC_EACH_SUBMISSION: '1'
     runs-on: [self-hosted, linux, X64]
     strategy:
       matrix:
@@ -38,5 +45,13 @@ jobs:
       - uses: julia-actions/cache@v3
       - uses: julia-actions/julia-buildpkg@latest
         continue-on-error: true
-      - uses: julia-actions/julia-runtest@latest
+      # Disable AVX512-FP16 host codegen on the Aurora Sapphire Rapids nodes. Under concurrent
+      # oneMKL load the native AVX512-FP16 path silently miscomputes *host* Float16 (e.g. the
+      # GPUArrays `A .* B .+ c` broadcast reference), failing tests even though the GPU result
+      # is correct (single-process clean; MXCSR clean; only the native-FP16 path, not Float32).
+      # `-C native,-avx512fp16` routes Float16 through Float32 and propagates to the Pkg.test
+      # subprocess and its parallel workers via Base.julia_cmd(). `julia-runtest` cannot pass a
+      # cpu-target, so invoke Pkg.test() directly. See repro_bcast_mkl.jl.
+      - name: Run tests (AVX512-FP16 disabled)
         continue-on-error: true
+        run: julia -C "native,-avx512fp16" --color=yes --project=. -e 'import Pkg; Pkg.test()'
diff --git a/Project.toml b/Project.toml
index 1775b329..c6731ca0 100644
--- a/Project.toml
+++ b/Project.toml
@@ -16,18 +16,19 @@ KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c"
 LLVM = "929cbde3-209d-540e-8aea-75f648917ca0"
 Libdl = "8f399da3-3557-5675-b5ff-fb832c97cbdb"
 LinearAlgebra = "37e2e46d-f89d-539d-b4ee-838fcccc9c8e"
-NEO_jll = "700fe977-ac61-5f37-bbc8-c6c4b2b6a9fd"
+NEO_LTS_jll = "a724f90f-ce79-56dd-a1bd-b9de5a61085f"
 Preferences = "21216c6a-2e73-6563-6e65-726566657250"
 Printf = "de0858da-6303-5e67-8744-51eddeeeb8d7"
 Random = "9a3f8284-a2c9-5f02-9a11-845980a1fd5c"
 SPIRVIntrinsics = "71d1d633-e7e8-4a92-83a1-de8814b09ba8"
 SPIRV_LLVM_Backend_jll = "4376b9bf-cff8-51b6-bb48-39421dff0d0c"
+SPIRV_LLVM_Translator_jll = "4a5d46fc-d8cf-5151-a261-86b458210efb"
 SPIRV_Tools_jll = "6ac6d60f-d740-5983-97d7-a4482c0689f4"
 SparseArrays = "2f01184e-e22b-5df5-ae63-d93ebab69eaf"
 SpecialFunctions = "276daf66-3868-5448-9aa4-cd146d93841b"
 StaticArrays = "90137ffa-7385-5640-81b9-e52037218182"
-oneAPI_Level_Zero_Headers_jll = "f4bc562b-d309-54f8-9efb-476e56f0410d"
-oneAPI_Level_Zero_Loader_jll = "13eca655-d68d-5b81-8367-6d99d727ab01"
+oneAPI_Level_Zero_Headers_LTS_jll = "d79c0b2e-896c-561b-aab9-323701ec0314"
+oneAPI_Level_Zero_Loader_LTS_jll = "f6e5cbb4-ba2a-56dc-92a2-9d66f5656ccd"
 oneAPI_Support_jll = "b049733a-a71d-5ed3-8eba-7d323ac00b36"
 
 [compat]
@@ -41,16 +42,22 @@ GPUCompiler = "1.6"
 GPUToolbox = "0.1, 0.2, 0.3, 1, 3"
 KernelAbstractions = "0.9.39"
 LLVM = "6, 7, 8, 9"
-NEO_jll = "=26.18.38308"
+NEO_LTS_jll = "=25.18.33578"
 Preferences = "1"
 SPIRVIntrinsics = "1"
 SPIRV_LLVM_Backend_jll = "22"
+SPIRV_LLVM_Translator_jll = "21"
 SPIRV_Tools_jll = "2025.4.0"
 SpecialFunctions = "1.3, 2"
 StaticArrays = "1"
 julia = "1.10"
-oneAPI_Level_Zero_Loader_jll = "1.29"
+oneAPI_Level_Zero_Loader_LTS_jll = "=1.24"
 oneAPI_Support_jll = "0.9.2"
 
 [extras]
-libigc_jll = "94295238-5935-5bd7-bb0f-b00942e9bdd5"
+libigc_LTS_jll = "9a8258a1-e827-5686-bee9-144461246960"
+
+[sources]
+# Pin to a commit SHA, not the mutable `atomic-float-ops` branch head, so builds are
+# reproducible and a force-push upstream cannot silently change what users resolve.
+SPIRVIntrinsics = {url = "https://github.com/michel2323/OpenCL.jl", rev = "4257d4075162276d9985168923788490e4071e9e", subdir = "lib/intrinsics"}
diff --git a/deps/Project.toml b/deps/Project.toml
index b8446dbe..8833df53 100644
--- a/deps/Project.toml
+++ b/deps/Project.toml
@@ -8,8 +8,8 @@ Ninja_jll = "76642167-d241-5cee-8c94-7a494e8cb7b7"
 Pkg = "44cfe95a-1eb2-52ea-b672-e2afdf69b78f"
 Preferences = "21216c6a-2e73-6563-6e65-726566657250"
 Scratch = "6c6a2e73-6563-6170-7368-637461726353"
-oneAPI_Level_Zero_Headers_jll = "f4bc562b-d309-54f8-9efb-476e56f0410d"
-oneAPI_Support_Headers_jll = "24f86df5-245d-5634-a4cc-32433d9800b3"
+oneAPI_Level_Zero_Headers_LTS_jll = "d79c0b2e-896c-561b-aab9-323701ec0314"
+oneAPI_Support_Headers_LTS_jll = "0e9de0da-c0b6-5d6c-9871-5c996d414ca7"
 
 [compat]
-oneAPI_Support_Headers_jll = "=2025.2.0"
+oneAPI_Support_Headers_LTS_jll = "=2025.3.1"
diff --git a/deps/build_local.jl b/deps/build_local.jl
index 80e67f51..4d613732 100644
--- a/deps/build_local.jl
+++ b/deps/build_local.jl
@@ -20,7 +20,8 @@ if haskey(ENV, "BUILDKITE")
     run(`buildkite-agent annotate 'Using a locally-built support library; A bump of oneAPI_Support_jll is required before releasing this packages.' --style 'warning' --context 'ctx-deps'`)
 end
 
-using Scratch, Preferences, CMake_jll, Ninja_jll, oneAPI_Level_Zero_Headers_jll
+using Scratch, Preferences, CMake_jll, Ninja_jll
+import oneAPI_Level_Zero_Headers_LTS_jll as oneAPI_Level_Zero_Headers_jll
 
 oneAPI = Base.UUID("8f75cd03-7ff8-4ecb-9b8f-daf728133b1b")
 
@@ -62,7 +63,7 @@ if !isfile(joinpath(conda_dir, "condarc-julia.yml"))
     touch(joinpath(conda_dir, "conda-meta", "history"))
 end
 Conda.add_channel("https://software.repos.intel.com/python/conda/", conda_dir)
-Conda.add(["dpcpp_linux-64=2026.0.0", "mkl-devel-dpcpp=2026.0.0"], conda_dir)
+Conda.add(["dpcpp_linux-64=2025.3.1", "mkl-devel-dpcpp=2025.3.1"], conda_dir)
 
 Conda.list(conda_dir)
 
diff --git a/deps/generate_interfaces.jl b/deps/generate_interfaces.jl
index 108518db..8c62b757 100644
--- a/deps/generate_interfaces.jl
+++ b/deps/generate_interfaces.jl
@@ -1,4 +1,4 @@
-using oneAPI_Support_Headers_jll
+import oneAPI_Support_Headers_LTS_jll as oneAPI_Support_Headers_jll
 
 include("generate_helpers.jl")
 
@@ -337,12 +337,34 @@ function generate_headers(library::String, filename::Vector{String}, output::Str
     end
   end
 
+  # Dedup: when two signatures map to the same C function name (because MKL
+  # added an overload), keep the one with more parameters — typically the
+  # newer signature (e.g. set_csr_data gained an `nnz` arg in MKL 2025.3.1).
+  # Without this the generated onemkl.cpp has duplicate function definitions
+  # and won't compile.
+  _fn_name(h) = (pos = findfirst('(', h); strip(split(strip(h[1:pos-1]))[end]))
+  _param_cnt(h) = (pos = findfirst('(', h); ep = findnext(')', h, pos); count(==(','), h[pos+1:ep-1]) + 1)
+  keep_idx = Dict{String,Int}()
+  keep_pc  = Dict{String,Int}()
+  for (i, sig) in enumerate(signatures)
+    (sig[2] in blacklist) && continue
+    fn = _fn_name(sig[1])
+    pc = _param_cnt(sig[1])
+    if !haskey(keep_idx, fn) || pc > keep_pc[fn]
+      keep_idx[fn] = i
+      keep_pc[fn]  = pc
+    end
+  end
+  keep_set = Set(values(keep_idx))
+
   path_oneapi_headers = joinpath(@__DIR__, output)
   oneapi_headers = open(path_oneapi_headers, "w")
 
-  for (header, name_routine, version, type_routine, template) in signatures
+  for (i, (header, name_routine, version, type_routine, template)) in enumerate(signatures)
     # Blacklist
     (name_routine in blacklist) && continue
+    # Dedup
+    (i in keep_set) || continue
 
     # Pass scalars (e.g. alpha/beta inputs) as references instead of values
     for type in ("short", "float", "double", "float _Complex", "double _Complex")
diff --git a/deps/src/onemkl.cpp b/deps/src/onemkl.cpp
index 2b880327..f9479c2f 100644
--- a/deps/src/onemkl.cpp
+++ b/deps/src/onemkl.cpp
@@ -5471,142 +5471,270 @@ extern "C" int64_t onemklZunmqr_batch_scratchpad_size(syclQueue_t device_queue,
 }
 
 // SPARSE
-extern "C" int onemklXsparse_init_matrix_handle(matrix_handle_t *p_spMat) {
-   oneapi::mkl::sparse::init_matrix_handle((oneapi::mkl::sparse::matrix_handle_t*) p_spMat);
+extern "C" int onemklXsparse_init_matrix_handle(matrix_handle_t *p_spmat) {
+   oneapi::mkl::sparse::init_matrix_handle((oneapi::mkl::sparse::matrix_handle_t*) p_spmat);
    return 0;
 }
 
-extern "C" int onemklXsparse_release_matrix_handle(syclQueue_t device_queue, matrix_handle_t *p_spMat) {
+extern "C" int onemklXsparse_release_matrix_handle(syclQueue_t device_queue, matrix_handle_t *p_spmat) {
    try {
-      auto status = oneapi::mkl::sparse::release_matrix_handle(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t*) p_spMat, {});
+      auto status = oneapi::mkl::sparse::release_matrix_handle(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t*) p_spmat, {});
       device_queue->val.wait_and_throw();
    } catch (const sycl::exception& e) { return -1; }
    return 0;
 }
 
-extern "C" int onemklSsparse_set_csr_data(syclQueue_t device_queue, matrix_handle_t spMat, int32_t nrows, int32_t ncols, onemklIndex index, int32_t *row_ptr, int32_t *col_ind, float *values) {
+extern "C" int onemklSsparse_set_csr_data(syclQueue_t device_queue, matrix_handle_t spmat, int64_t nrows, int64_t ncols, int64_t nnz, onemklIndex index, int32_t *row_ptr, int32_t *col_ind, float *values) {
    try {
-      auto status = oneapi::mkl::sparse::set_csr_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spMat, nrows, ncols, convert(index), row_ptr, col_ind, values, {});
+      auto status = oneapi::mkl::sparse::set_csr_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spmat, nrows, ncols, nnz, convert(index), row_ptr, col_ind, values, {});
       device_queue->val.wait_and_throw();
    } catch (const sycl::exception& e) { return -1; }
    return 0;
 }
 
-extern "C" int onemklSsparse_set_csr_data_64(syclQueue_t device_queue, matrix_handle_t spMat, int64_t nrows, int64_t ncols, onemklIndex index, int64_t *row_ptr, int64_t *col_ind, float *values) {
+extern "C" int onemklSsparse_set_csr_data_64(syclQueue_t device_queue, matrix_handle_t spmat, int64_t nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t *row_ptr, int64_t *col_ind, float *values) {
    try {
-      auto status = oneapi::mkl::sparse::set_csr_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spMat, nrows, ncols, convert(index), row_ptr, col_ind, values, {});
+      auto status = oneapi::mkl::sparse::set_csr_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spmat, nrows, ncols, nnz, convert(index), row_ptr, col_ind, values, {});
       device_queue->val.wait_and_throw();
    } catch (const sycl::exception& e) { return -1; }
    return 0;
 }
 
-extern "C" int onemklDsparse_set_csr_data(syclQueue_t device_queue, matrix_handle_t spMat, int32_t nrows, int32_t ncols, onemklIndex index, int32_t *row_ptr, int32_t *col_ind, double *values) {
+extern "C" int onemklDsparse_set_csr_data(syclQueue_t device_queue, matrix_handle_t spmat, int64_t nrows, int64_t ncols, int64_t nnz, onemklIndex index, int32_t *row_ptr, int32_t *col_ind, double *values) {
    try {
-      auto status = oneapi::mkl::sparse::set_csr_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spMat, nrows, ncols, convert(index), row_ptr, col_ind, values, {});
+      auto status = oneapi::mkl::sparse::set_csr_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spmat, nrows, ncols, nnz, convert(index), row_ptr, col_ind, values, {});
       device_queue->val.wait_and_throw();
    } catch (const sycl::exception& e) { return -1; }
    return 0;
 }
 
-extern "C" int onemklDsparse_set_csr_data_64(syclQueue_t device_queue, matrix_handle_t spMat, int64_t nrows, int64_t ncols, onemklIndex index, int64_t *row_ptr, int64_t *col_ind, double *values) {
+extern "C" int onemklDsparse_set_csr_data_64(syclQueue_t device_queue, matrix_handle_t spmat, int64_t nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t *row_ptr, int64_t *col_ind, double *values) {
    try {
-      auto status = oneapi::mkl::sparse::set_csr_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spMat, nrows, ncols, convert(index), row_ptr, col_ind, values, {});
+      auto status = oneapi::mkl::sparse::set_csr_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spmat, nrows, ncols, nnz, convert(index), row_ptr, col_ind, values, {});
       device_queue->val.wait_and_throw();
    } catch (const sycl::exception& e) { return -1; }
    return 0;
 }
 
-extern "C" int onemklCsparse_set_csr_data(syclQueue_t device_queue, matrix_handle_t spMat, int32_t nrows, int32_t ncols, onemklIndex index, int32_t *row_ptr, int32_t *col_ind, float _Complex *values) {
+extern "C" int onemklCsparse_set_csr_data(syclQueue_t device_queue, matrix_handle_t spmat, int64_t nrows, int64_t ncols, int64_t nnz, onemklIndex index, int32_t *row_ptr, int32_t *col_ind, float _Complex *values) {
    try {
-      auto status = oneapi::mkl::sparse::set_csr_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spMat, nrows, ncols, convert(index), row_ptr, col_ind, reinterpret_cast<std::complex<float>*>(values), {});
+      auto status = oneapi::mkl::sparse::set_csr_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spmat, nrows, ncols, nnz, convert(index), row_ptr, col_ind, reinterpret_cast<std::complex<float>*>(values), {});
       device_queue->val.wait_and_throw();
    } catch (const sycl::exception& e) { return -1; }
    return 0;
 }
 
-extern "C" int onemklCsparse_set_csr_data_64(syclQueue_t device_queue, matrix_handle_t spMat, int64_t nrows, int64_t ncols, onemklIndex index, int64_t *row_ptr, int64_t *col_ind, float _Complex *values) {
+extern "C" int onemklCsparse_set_csr_data_64(syclQueue_t device_queue, matrix_handle_t spmat, int64_t nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t *row_ptr, int64_t *col_ind, float _Complex *values) {
    try {
-      auto status = oneapi::mkl::sparse::set_csr_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spMat, nrows, ncols, convert(index), row_ptr, col_ind, reinterpret_cast<std::complex<float>*>(values), {});
+      auto status = oneapi::mkl::sparse::set_csr_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spmat, nrows, ncols, nnz, convert(index), row_ptr, col_ind, reinterpret_cast<std::complex<float>*>(values), {});
       device_queue->val.wait_and_throw();
    } catch (const sycl::exception& e) { return -1; }
    return 0;
 }
 
-extern "C" int onemklZsparse_set_csr_data(syclQueue_t device_queue, matrix_handle_t spMat, int32_t nrows, int32_t ncols, onemklIndex index, int32_t *row_ptr, int32_t *col_ind, double _Complex *values) {
+extern "C" int onemklZsparse_set_csr_data(syclQueue_t device_queue, matrix_handle_t spmat, int64_t nrows, int64_t ncols, int64_t nnz, onemklIndex index, int32_t *row_ptr, int32_t *col_ind, double _Complex *values) {
    try {
-      auto status = oneapi::mkl::sparse::set_csr_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spMat, nrows, ncols, convert(index), row_ptr, col_ind, reinterpret_cast<std::complex<double>*>(values), {});
+      auto status = oneapi::mkl::sparse::set_csr_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spmat, nrows, ncols, nnz, convert(index), row_ptr, col_ind, reinterpret_cast<std::complex<double>*>(values), {});
       device_queue->val.wait_and_throw();
    } catch (const sycl::exception& e) { return -1; }
    return 0;
 }
 
-extern "C" int onemklZsparse_set_csr_data_64(syclQueue_t device_queue, matrix_handle_t spMat, int64_t nrows, int64_t ncols, onemklIndex index, int64_t *row_ptr, int64_t *col_ind, double _Complex *values) {
+extern "C" int onemklZsparse_set_csr_data_64(syclQueue_t device_queue, matrix_handle_t spmat, int64_t nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t *row_ptr, int64_t *col_ind, double _Complex *values) {
    try {
-      auto status = oneapi::mkl::sparse::set_csr_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spMat, nrows, ncols, convert(index), row_ptr, col_ind, reinterpret_cast<std::complex<double>*>(values), {});
+      auto status = oneapi::mkl::sparse::set_csr_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spmat, nrows, ncols, nnz, convert(index), row_ptr, col_ind, reinterpret_cast<std::complex<double>*>(values), {});
       device_queue->val.wait_and_throw();
    } catch (const sycl::exception& e) { return -1; }
    return 0;
 }
 
-extern "C" int onemklSsparse_set_coo_data(syclQueue_t device_queue, matrix_handle_t spMat, int32_t nrows, int32_t ncols, int32_t nnz, onemklIndex index, int32_t *row_ind, int32_t *col_ind, float *values) {
+extern "C" int onemklSsparse_set_csc_data(syclQueue_t device_queue, matrix_handle_t spMat, int64_t nrows, int64_t ncols, int64_t nnz, onemklIndex index, int32_t *col_ptr, int32_t *row_ind, float *values) {
    try {
-      auto status = oneapi::mkl::sparse::set_coo_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spMat, nrows, ncols, nnz, convert(index), row_ind, col_ind, values, {});
+      auto status = oneapi::mkl::sparse::set_csc_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spMat, nrows, ncols, nnz, convert(index), col_ptr, row_ind, values, {});
       device_queue->val.wait_and_throw();
    } catch (const sycl::exception& e) { return -1; }
    return 0;
 }
 
-extern "C" int onemklSsparse_set_coo_data_64(syclQueue_t device_queue, matrix_handle_t spMat, int64_t nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t *row_ind, int64_t *col_ind, float *values) {
+extern "C" int onemklSsparse_set_csc_data_64(syclQueue_t device_queue, matrix_handle_t spMat, int64_t nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t *col_ptr, int64_t *row_ind, float *values) {
    try {
-      auto status = oneapi::mkl::sparse::set_coo_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spMat, nrows, ncols, nnz, convert(index), row_ind, col_ind, values, {});
+      auto status = oneapi::mkl::sparse::set_csc_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spMat, nrows, ncols, nnz, convert(index), col_ptr, row_ind, values, {});
       device_queue->val.wait_and_throw();
    } catch (const sycl::exception& e) { return -1; }
    return 0;
 }
 
-extern "C" int onemklDsparse_set_coo_data(syclQueue_t device_queue, matrix_handle_t spMat, int32_t nrows, int32_t ncols, int32_t nnz, onemklIndex index, int32_t *row_ind, int32_t *col_ind, double *values) {
+extern "C" int onemklDsparse_set_csc_data(syclQueue_t device_queue, matrix_handle_t spMat, int64_t nrows, int64_t ncols, int64_t nnz, onemklIndex index, int32_t *col_ptr, int32_t *row_ind, double *values) {
    try {
-      auto status = oneapi::mkl::sparse::set_coo_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spMat, nrows, ncols, nnz, convert(index), row_ind, col_ind, values, {});
+      auto status = oneapi::mkl::sparse::set_csc_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spMat, nrows, ncols, nnz, convert(index), col_ptr, row_ind, values, {});
       device_queue->val.wait_and_throw();
    } catch (const sycl::exception& e) { return -1; }
    return 0;
 }
 
-extern "C" int onemklDsparse_set_coo_data_64(syclQueue_t device_queue, matrix_handle_t spMat, int64_t nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t *row_ind, int64_t *col_ind, double *values) {
+extern "C" int onemklDsparse_set_csc_data_64(syclQueue_t device_queue, matrix_handle_t spMat, int64_t nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t *col_ptr, int64_t *row_ind, double *values) {
    try {
-      auto status = oneapi::mkl::sparse::set_coo_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spMat, nrows, ncols, nnz, convert(index), row_ind, col_ind, values, {});
+      auto status = oneapi::mkl::sparse::set_csc_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spMat, nrows, ncols, nnz, convert(index), col_ptr, row_ind, values, {});
       device_queue->val.wait_and_throw();
    } catch (const sycl::exception& e) { return -1; }
    return 0;
 }
 
-extern "C" int onemklCsparse_set_coo_data(syclQueue_t device_queue, matrix_handle_t spMat, int32_t nrows, int32_t ncols, int32_t nnz, onemklIndex index, int32_t *row_ind, int32_t *col_ind, float _Complex *values) {
+extern "C" int onemklCsparse_set_csc_data(syclQueue_t device_queue, matrix_handle_t spMat, int64_t nrows, int64_t ncols, int64_t nnz, onemklIndex index, int32_t *col_ptr, int32_t *row_ind, float _Complex *values) {
    try {
-      auto status = oneapi::mkl::sparse::set_coo_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spMat, nrows, ncols, nnz, convert(index), row_ind, col_ind, reinterpret_cast<std::complex<float>*>(values), {});
+      auto status = oneapi::mkl::sparse::set_csc_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spMat, nrows, ncols, nnz, convert(index), col_ptr, row_ind, reinterpret_cast<std::complex<float>*>(values), {});
       device_queue->val.wait_and_throw();
    } catch (const sycl::exception& e) { return -1; }
    return 0;
 }
 
-extern "C" int onemklCsparse_set_coo_data_64(syclQueue_t device_queue, matrix_handle_t spMat, int64_t nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t *row_ind, int64_t *col_ind, float _Complex *values) {
+extern "C" int onemklCsparse_set_csc_data_64(syclQueue_t device_queue, matrix_handle_t spMat, int64_t nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t *col_ptr, int64_t *row_ind, float _Complex *values) {
    try {
-      auto status = oneapi::mkl::sparse::set_coo_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spMat, nrows, ncols, nnz, convert(index), row_ind, col_ind, reinterpret_cast<std::complex<float>*>(values), {});
+      auto status = oneapi::mkl::sparse::set_csc_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spMat, nrows, ncols, nnz, convert(index), col_ptr, row_ind, reinterpret_cast<std::complex<float>*>(values), {});
       device_queue->val.wait_and_throw();
    } catch (const sycl::exception& e) { return -1; }
    return 0;
 }
 
-extern "C" int onemklZsparse_set_coo_data(syclQueue_t device_queue, matrix_handle_t spMat, int32_t nrows, int32_t ncols, int32_t nnz, onemklIndex index, int32_t *row_ind, int32_t *col_ind, double _Complex *values) {
+extern "C" int onemklZsparse_set_csc_data(syclQueue_t device_queue, matrix_handle_t spMat, int64_t nrows, int64_t ncols, int64_t nnz, onemklIndex index, int32_t *col_ptr, int32_t *row_ind, double _Complex *values) {
    try {
-      auto status = oneapi::mkl::sparse::set_coo_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spMat, nrows, ncols, nnz, convert(index), row_ind, col_ind, reinterpret_cast<std::complex<double>*>(values), {});
+      auto status = oneapi::mkl::sparse::set_csc_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spMat, nrows, ncols, nnz, convert(index), col_ptr, row_ind, reinterpret_cast<std::complex<double>*>(values), {});
       device_queue->val.wait_and_throw();
    } catch (const sycl::exception& e) { return -1; }
    return 0;
 }
 
-extern "C" int onemklZsparse_set_coo_data_64(syclQueue_t device_queue, matrix_handle_t spMat, int64_t nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t *row_ind, int64_t *col_ind, double _Complex *values) {
+extern "C" int onemklZsparse_set_csc_data_64(syclQueue_t device_queue, matrix_handle_t spMat, int64_t nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t *col_ptr, int64_t *row_ind, double _Complex *values) {
    try {
-      auto status = oneapi::mkl::sparse::set_coo_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spMat, nrows, ncols, nnz, convert(index), row_ind, col_ind, reinterpret_cast<std::complex<double>*>(values), {});
+      auto status = oneapi::mkl::sparse::set_csc_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spMat, nrows, ncols, nnz, convert(index), col_ptr, row_ind, reinterpret_cast<std::complex<double>*>(values), {});
+      device_queue->val.wait_and_throw();
+   } catch (const sycl::exception& e) { return -1; }
+   return 0;
+}
+
+extern "C" int onemklSsparse_set_coo_data(syclQueue_t device_queue, matrix_handle_t spmat, int32_t nrows, int32_t ncols, int32_t nnz, onemklIndex index, int32_t *row_ind, int32_t *col_ind, float *values) {
+   try {
+      auto status = oneapi::mkl::sparse::set_coo_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spmat, nrows, ncols, nnz, convert(index), row_ind, col_ind, values, {});
+      device_queue->val.wait_and_throw();
+   } catch (const sycl::exception& e) { return -1; }
+   return 0;
+}
+
+extern "C" int onemklSsparse_set_coo_data_64(syclQueue_t device_queue, matrix_handle_t spmat, int64_t nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t *row_ind, int64_t *col_ind, float *values) {
+   try {
+      auto status = oneapi::mkl::sparse::set_coo_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spmat, nrows, ncols, nnz, convert(index), row_ind, col_ind, values, {});
+      device_queue->val.wait_and_throw();
+   } catch (const sycl::exception& e) { return -1; }
+   return 0;
+}
+
+extern "C" int onemklDsparse_set_coo_data(syclQueue_t device_queue, matrix_handle_t spmat, int32_t nrows, int32_t ncols, int32_t nnz, onemklIndex index, int32_t *row_ind, int32_t *col_ind, double *values) {
+   try {
+      auto status = oneapi::mkl::sparse::set_coo_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spmat, nrows, ncols, nnz, convert(index), row_ind, col_ind, values, {});
+      device_queue->val.wait_and_throw();
+   } catch (const sycl::exception& e) { return -1; }
+   return 0;
+}
+
+extern "C" int onemklDsparse_set_coo_data_64(syclQueue_t device_queue, matrix_handle_t spmat, int64_t nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t *row_ind, int64_t *col_ind, double *values) {
+   try {
+      auto status = oneapi::mkl::sparse::set_coo_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spmat, nrows, ncols, nnz, convert(index), row_ind, col_ind, values, {});
+      device_queue->val.wait_and_throw();
+   } catch (const sycl::exception& e) { return -1; }
+   return 0;
+}
+
+extern "C" int onemklCsparse_set_coo_data(syclQueue_t device_queue, matrix_handle_t spmat, int32_t nrows, int32_t ncols, int32_t nnz, onemklIndex index, int32_t *row_ind, int32_t *col_ind, float _Complex *values) {
+   try {
+      auto status = oneapi::mkl::sparse::set_coo_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spmat, nrows, ncols, nnz, convert(index), row_ind, col_ind, reinterpret_cast<std::complex<float>*>(values), {});
+      device_queue->val.wait_and_throw();
+   } catch (const sycl::exception& e) { return -1; }
+   return 0;
+}
+
+extern "C" int onemklCsparse_set_coo_data_64(syclQueue_t device_queue, matrix_handle_t spmat, int64_t nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t *row_ind, int64_t *col_ind, float _Complex *values) {
+   try {
+      auto status = oneapi::mkl::sparse::set_coo_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spmat, nrows, ncols, nnz, convert(index), row_ind, col_ind, reinterpret_cast<std::complex<float>*>(values), {});
+      device_queue->val.wait_and_throw();
+   } catch (const sycl::exception& e) { return -1; }
+   return 0;
+}
+
+extern "C" int onemklZsparse_set_coo_data(syclQueue_t device_queue, matrix_handle_t spmat, int32_t nrows, int32_t ncols, int32_t nnz, onemklIndex index, int32_t *row_ind, int32_t *col_ind, double _Complex *values) {
+   try {
+      auto status = oneapi::mkl::sparse::set_coo_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spmat, nrows, ncols, nnz, convert(index), row_ind, col_ind, reinterpret_cast<std::complex<double>*>(values), {});
+      device_queue->val.wait_and_throw();
+   } catch (const sycl::exception& e) { return -1; }
+   return 0;
+}
+
+extern "C" int onemklZsparse_set_coo_data_64(syclQueue_t device_queue, matrix_handle_t spmat, int64_t nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t *row_ind, int64_t *col_ind, double _Complex *values) {
+   try {
+      auto status = oneapi::mkl::sparse::set_coo_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spmat, nrows, ncols, nnz, convert(index), row_ind, col_ind, reinterpret_cast<std::complex<double>*>(values), {});
+      device_queue->val.wait_and_throw();
+   } catch (const sycl::exception& e) { return -1; }
+   return 0;
+}
+
+extern "C" int onemklSsparse_set_bsr_data(syclQueue_t device_queue, matrix_handle_t spmat, int64_t blk_nrows, int64_t blk_ncols, int64_t blk_nnz, int64_t row_blk_size, int64_t col_blk_size, onemklLayout blk_layout, onemklIndex index, int32_t *bsr_row_ptr, int32_t *bsr_col_ind, float *bsr_values) {
+   try {
+      auto status = oneapi::mkl::sparse::set_bsr_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spmat, blk_nrows, blk_ncols, blk_nnz, row_blk_size, col_blk_size, convert(blk_layout), convert(index), bsr_row_ptr, bsr_col_ind, bsr_values, {});
+      device_queue->val.wait_and_throw();
+   } catch (const sycl::exception& e) { return -1; }
+   return 0;
+}
+
+extern "C" int onemklSsparse_set_bsr_data_64(syclQueue_t device_queue, matrix_handle_t spmat, int64_t blk_nrows, int64_t blk_ncols, int64_t blk_nnz, int64_t row_blk_size, int64_t col_blk_size, onemklLayout blk_layout, onemklIndex index, int64_t *bsr_row_ptr, int64_t *bsr_col_ind, float *bsr_values) {
+   try {
+      auto status = oneapi::mkl::sparse::set_bsr_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spmat, blk_nrows, blk_ncols, blk_nnz, row_blk_size, col_blk_size, convert(blk_layout), convert(index), bsr_row_ptr, bsr_col_ind, bsr_values, {});
+      device_queue->val.wait_and_throw();
+   } catch (const sycl::exception& e) { return -1; }
+   return 0;
+}
+
+extern "C" int onemklDsparse_set_bsr_data(syclQueue_t device_queue, matrix_handle_t spmat, int64_t blk_nrows, int64_t blk_ncols, int64_t blk_nnz, int64_t row_blk_size, int64_t col_blk_size, onemklLayout blk_layout, onemklIndex index, int32_t *bsr_row_ptr, int32_t *bsr_col_ind, double *bsr_values) {
+   try {
+      auto status = oneapi::mkl::sparse::set_bsr_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spmat, blk_nrows, blk_ncols, blk_nnz, row_blk_size, col_blk_size, convert(blk_layout), convert(index), bsr_row_ptr, bsr_col_ind, bsr_values, {});
+      device_queue->val.wait_and_throw();
+   } catch (const sycl::exception& e) { return -1; }
+   return 0;
+}
+
+extern "C" int onemklDsparse_set_bsr_data_64(syclQueue_t device_queue, matrix_handle_t spmat, int64_t blk_nrows, int64_t blk_ncols, int64_t blk_nnz, int64_t row_blk_size, int64_t col_blk_size, onemklLayout blk_layout, onemklIndex index, int64_t *bsr_row_ptr, int64_t *bsr_col_ind, double *bsr_values) {
+   try {
+      auto status = oneapi::mkl::sparse::set_bsr_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spmat, blk_nrows, blk_ncols, blk_nnz, row_blk_size, col_blk_size, convert(blk_layout), convert(index), bsr_row_ptr, bsr_col_ind, bsr_values, {});
+      device_queue->val.wait_and_throw();
+   } catch (const sycl::exception& e) { return -1; }
+   return 0;
+}
+
+extern "C" int onemklCsparse_set_bsr_data(syclQueue_t device_queue, matrix_handle_t spmat, int64_t blk_nrows, int64_t blk_ncols, int64_t blk_nnz, int64_t row_blk_size, int64_t col_blk_size, onemklLayout blk_layout, onemklIndex index, int32_t *bsr_row_ptr, int32_t *bsr_col_ind, float _Complex *bsr_values) {
+   try {
+      auto status = oneapi::mkl::sparse::set_bsr_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spmat, blk_nrows, blk_ncols, blk_nnz, row_blk_size, col_blk_size, convert(blk_layout), convert(index), bsr_row_ptr, bsr_col_ind, reinterpret_cast<std::complex<float>*>(bsr_values), {});
+      device_queue->val.wait_and_throw();
+   } catch (const sycl::exception& e) { return -1; }
+   return 0;
+}
+
+extern "C" int onemklCsparse_set_bsr_data_64(syclQueue_t device_queue, matrix_handle_t spmat, int64_t blk_nrows, int64_t blk_ncols, int64_t blk_nnz, int64_t row_blk_size, int64_t col_blk_size, onemklLayout blk_layout, onemklIndex index, int64_t *bsr_row_ptr, int64_t *bsr_col_ind, float _Complex *bsr_values) {
+   try {
+      auto status = oneapi::mkl::sparse::set_bsr_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spmat, blk_nrows, blk_ncols, blk_nnz, row_blk_size, col_blk_size, convert(blk_layout), convert(index), bsr_row_ptr, bsr_col_ind, reinterpret_cast<std::complex<float>*>(bsr_values), {});
+      device_queue->val.wait_and_throw();
+   } catch (const sycl::exception& e) { return -1; }
+   return 0;
+}
+
+extern "C" int onemklZsparse_set_bsr_data(syclQueue_t device_queue, matrix_handle_t spmat, int64_t blk_nrows, int64_t blk_ncols, int64_t blk_nnz, int64_t row_blk_size, int64_t col_blk_size, onemklLayout blk_layout, onemklIndex index, int32_t *bsr_row_ptr, int32_t *bsr_col_ind, double _Complex *bsr_values) {
+   try {
+      auto status = oneapi::mkl::sparse::set_bsr_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spmat, blk_nrows, blk_ncols, blk_nnz, row_blk_size, col_blk_size, convert(blk_layout), convert(index), bsr_row_ptr, bsr_col_ind, reinterpret_cast<std::complex<double>*>(bsr_values), {});
+      device_queue->val.wait_and_throw();
+   } catch (const sycl::exception& e) { return -1; }
+   return 0;
+}
+
+extern "C" int onemklZsparse_set_bsr_data_64(syclQueue_t device_queue, matrix_handle_t spmat, int64_t blk_nrows, int64_t blk_ncols, int64_t blk_nnz, int64_t row_blk_size, int64_t col_blk_size, onemklLayout blk_layout, onemklIndex index, int64_t *bsr_row_ptr, int64_t *bsr_col_ind, double _Complex *bsr_values) {
+   try {
+      auto status = oneapi::mkl::sparse::set_bsr_data(device_queue->val, (oneapi::mkl::sparse::matrix_handle_t) spmat, blk_nrows, blk_ncols, blk_nnz, row_blk_size, col_blk_size, convert(blk_layout), convert(index), bsr_row_ptr, bsr_col_ind, reinterpret_cast<std::complex<double>*>(bsr_values), {});
       device_queue->val.wait_and_throw();
    } catch (const sycl::exception& e) { return -1; }
    return 0;
diff --git a/deps/src/onemkl.h b/deps/src/onemkl.h
index 1f850bf2..4d1caa08 100644
--- a/deps/src/onemkl.h
+++ b/deps/src/onemkl.h
@@ -140,9 +140,9 @@ typedef struct omatconvert_descr *omatconvert_descr_t;
 struct omatadd_descr;
 typedef struct omatadd_descr *omatadd_descr_t;
 
-const int64_t ONEMKL_VERSION_MAJOR = 2026;
-const int64_t ONEMKL_VERSION_MINOR = 0;
-const int64_t ONEMKL_VERSION_PATCH = 0;
+const int64_t ONEMKL_VERSION_MAJOR = 2025;
+const int64_t ONEMKL_VERSION_MINOR = 3;
+const int64_t ONEMKL_VERSION_PATCH = 1;
 void onemkl_version(int64_t *major, int64_t *minor, int64_t *patch);
 
 int onemklHgemm_batch(syclQueue_t device_queue, onemklTranspose transa,
@@ -2737,74 +2737,154 @@ int64_t onemklZunmqr_batch_scratchpad_size(syclQueue_t device_queue, onemklSide*
                                            group_count, int64_t* group_sizes);
 
 // SPARSE
-int onemklXsparse_init_matrix_handle(matrix_handle_t *p_spMat);
+int onemklXsparse_init_matrix_handle(matrix_handle_t *p_spmat);
 
-int onemklXsparse_release_matrix_handle(syclQueue_t device_queue, matrix_handle_t *p_spMat);
+int onemklXsparse_release_matrix_handle(syclQueue_t device_queue, matrix_handle_t *p_spmat);
 
-int onemklSsparse_set_csr_data(syclQueue_t device_queue, matrix_handle_t spMat, int32_t nrows,
-                               int32_t ncols, onemklIndex index, int32_t *row_ptr, int32_t
-                               *col_ind, float *values);
+int onemklSsparse_set_csr_data(syclQueue_t device_queue, matrix_handle_t spmat, int64_t nrows,
+                               int64_t ncols, int64_t nnz, onemklIndex index, int32_t *row_ptr,
+                               int32_t *col_ind, float *values);
 
-int onemklSsparse_set_csr_data_64(syclQueue_t device_queue, matrix_handle_t spMat, int64_t
-                                  nrows, int64_t ncols, onemklIndex index, int64_t *row_ptr,
-                                  int64_t *col_ind, float *values);
+int onemklSsparse_set_csr_data_64(syclQueue_t device_queue, matrix_handle_t spmat, int64_t
+                                  nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t
+                                  *row_ptr, int64_t *col_ind, float *values);
 
-int onemklDsparse_set_csr_data(syclQueue_t device_queue, matrix_handle_t spMat, int32_t nrows,
-                               int32_t ncols, onemklIndex index, int32_t *row_ptr, int32_t
-                               *col_ind, double *values);
+int onemklDsparse_set_csr_data(syclQueue_t device_queue, matrix_handle_t spmat, int64_t nrows,
+                               int64_t ncols, int64_t nnz, onemklIndex index, int32_t *row_ptr,
+                               int32_t *col_ind, double *values);
 
-int onemklDsparse_set_csr_data_64(syclQueue_t device_queue, matrix_handle_t spMat, int64_t
-                                  nrows, int64_t ncols, onemklIndex index, int64_t *row_ptr,
-                                  int64_t *col_ind, double *values);
+int onemklDsparse_set_csr_data_64(syclQueue_t device_queue, matrix_handle_t spmat, int64_t
+                                  nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t
+                                  *row_ptr, int64_t *col_ind, double *values);
 
-int onemklCsparse_set_csr_data(syclQueue_t device_queue, matrix_handle_t spMat, int32_t nrows,
-                               int32_t ncols, onemklIndex index, int32_t *row_ptr, int32_t
-                               *col_ind, float _Complex *values);
+int onemklCsparse_set_csr_data(syclQueue_t device_queue, matrix_handle_t spmat, int64_t nrows,
+                               int64_t ncols, int64_t nnz, onemklIndex index, int32_t *row_ptr,
+                               int32_t *col_ind, float _Complex *values);
 
-int onemklCsparse_set_csr_data_64(syclQueue_t device_queue, matrix_handle_t spMat, int64_t
-                                  nrows, int64_t ncols, onemklIndex index, int64_t *row_ptr,
-                                  int64_t *col_ind, float _Complex *values);
+int onemklCsparse_set_csr_data_64(syclQueue_t device_queue, matrix_handle_t spmat, int64_t
+                                  nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t
+                                  *row_ptr, int64_t *col_ind, float _Complex *values);
 
-int onemklZsparse_set_csr_data(syclQueue_t device_queue, matrix_handle_t spMat, int32_t nrows,
-                               int32_t ncols, onemklIndex index, int32_t *row_ptr, int32_t
-                               *col_ind, double _Complex *values);
+int onemklZsparse_set_csr_data(syclQueue_t device_queue, matrix_handle_t spmat, int64_t nrows,
+                               int64_t ncols, int64_t nnz, onemklIndex index, int32_t *row_ptr,
+                               int32_t *col_ind, double _Complex *values);
+
+int onemklZsparse_set_csr_data_64(syclQueue_t device_queue, matrix_handle_t spmat, int64_t
+                                  nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t
+                                  *row_ptr, int64_t *col_ind, double _Complex *values);
 
-int onemklZsparse_set_csr_data_64(syclQueue_t device_queue, matrix_handle_t spMat, int64_t
-                                  nrows, int64_t ncols, onemklIndex index, int64_t *row_ptr,
-                                  int64_t *col_ind, double _Complex *values);
+int onemklSsparse_set_csc_data(syclQueue_t device_queue, matrix_handle_t spMat, int64_t nrows,
+                               int64_t ncols, int64_t nnz, onemklIndex index, int32_t *col_ptr,
+                               int32_t *row_ind, float *values);
+
+int onemklSsparse_set_csc_data_64(syclQueue_t device_queue, matrix_handle_t spMat, int64_t
+                                  nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t
+                                  *col_ptr, int64_t *row_ind, float *values);
 
-int onemklSsparse_set_coo_data(syclQueue_t device_queue, matrix_handle_t spMat, int32_t nrows,
+int onemklDsparse_set_csc_data(syclQueue_t device_queue, matrix_handle_t spMat, int64_t nrows,
+                               int64_t ncols, int64_t nnz, onemklIndex index, int32_t *col_ptr,
+                               int32_t *row_ind, double *values);
+
+int onemklDsparse_set_csc_data_64(syclQueue_t device_queue, matrix_handle_t spMat, int64_t
+                                  nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t
+                                  *col_ptr, int64_t *row_ind, double *values);
+
+int onemklCsparse_set_csc_data(syclQueue_t device_queue, matrix_handle_t spMat, int64_t nrows,
+                               int64_t ncols, int64_t nnz, onemklIndex index, int32_t *col_ptr,
+                               int32_t *row_ind, float _Complex *values);
+
+int onemklCsparse_set_csc_data_64(syclQueue_t device_queue, matrix_handle_t spMat, int64_t
+                                  nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t
+                                  *col_ptr, int64_t *row_ind, float _Complex *values);
+
+int onemklZsparse_set_csc_data(syclQueue_t device_queue, matrix_handle_t spMat, int64_t nrows,
+                               int64_t ncols, int64_t nnz, onemklIndex index, int32_t *col_ptr,
+                               int32_t *row_ind, double _Complex *values);
+
+int onemklZsparse_set_csc_data_64(syclQueue_t device_queue, matrix_handle_t spMat, int64_t
+                                  nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t
+                                  *col_ptr, int64_t *row_ind, double _Complex *values);
+
+int onemklSsparse_set_coo_data(syclQueue_t device_queue, matrix_handle_t spmat, int32_t nrows,
                                int32_t ncols, int32_t nnz, onemklIndex index, int32_t *row_ind,
                                int32_t *col_ind, float *values);
 
-int onemklSsparse_set_coo_data_64(syclQueue_t device_queue, matrix_handle_t spMat, int64_t
+int onemklSsparse_set_coo_data_64(syclQueue_t device_queue, matrix_handle_t spmat, int64_t
                                   nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t
                                   *row_ind, int64_t *col_ind, float *values);
 
-int onemklDsparse_set_coo_data(syclQueue_t device_queue, matrix_handle_t spMat, int32_t nrows,
+int onemklDsparse_set_coo_data(syclQueue_t device_queue, matrix_handle_t spmat, int32_t nrows,
                                int32_t ncols, int32_t nnz, onemklIndex index, int32_t *row_ind,
                                int32_t *col_ind, double *values);
 
-int onemklDsparse_set_coo_data_64(syclQueue_t device_queue, matrix_handle_t spMat, int64_t
+int onemklDsparse_set_coo_data_64(syclQueue_t device_queue, matrix_handle_t spmat, int64_t
                                   nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t
                                   *row_ind, int64_t *col_ind, double *values);
 
-int onemklCsparse_set_coo_data(syclQueue_t device_queue, matrix_handle_t spMat, int32_t nrows,
+int onemklCsparse_set_coo_data(syclQueue_t device_queue, matrix_handle_t spmat, int32_t nrows,
                                int32_t ncols, int32_t nnz, onemklIndex index, int32_t *row_ind,
                                int32_t *col_ind, float _Complex *values);
 
-int onemklCsparse_set_coo_data_64(syclQueue_t device_queue, matrix_handle_t spMat, int64_t
+int onemklCsparse_set_coo_data_64(syclQueue_t device_queue, matrix_handle_t spmat, int64_t
                                   nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t
                                   *row_ind, int64_t *col_ind, float _Complex *values);
 
-int onemklZsparse_set_coo_data(syclQueue_t device_queue, matrix_handle_t spMat, int32_t nrows,
+int onemklZsparse_set_coo_data(syclQueue_t device_queue, matrix_handle_t spmat, int32_t nrows,
                                int32_t ncols, int32_t nnz, onemklIndex index, int32_t *row_ind,
                                int32_t *col_ind, double _Complex *values);
 
-int onemklZsparse_set_coo_data_64(syclQueue_t device_queue, matrix_handle_t spMat, int64_t
+int onemklZsparse_set_coo_data_64(syclQueue_t device_queue, matrix_handle_t spmat, int64_t
                                   nrows, int64_t ncols, int64_t nnz, onemklIndex index, int64_t
                                   *row_ind, int64_t *col_ind, double _Complex *values);
 
+int onemklSsparse_set_bsr_data(syclQueue_t device_queue, matrix_handle_t spmat, int64_t
+                               blk_nrows, int64_t blk_ncols, int64_t blk_nnz, int64_t
+                               row_blk_size, int64_t col_blk_size, onemklLayout blk_layout,
+                               onemklIndex index, int32_t *bsr_row_ptr, int32_t *bsr_col_ind,
+                               float *bsr_values);
+
+int onemklSsparse_set_bsr_data_64(syclQueue_t device_queue, matrix_handle_t spmat, int64_t
+                                  blk_nrows, int64_t blk_ncols, int64_t blk_nnz, int64_t
+                                  row_blk_size, int64_t col_blk_size, onemklLayout blk_layout,
+                                  onemklIndex index, int64_t *bsr_row_ptr, int64_t
+                                  *bsr_col_ind, float *bsr_values);
+
+int onemklDsparse_set_bsr_data(syclQueue_t device_queue, matrix_handle_t spmat, int64_t
+                               blk_nrows, int64_t blk_ncols, int64_t blk_nnz, int64_t
+                               row_blk_size, int64_t col_blk_size, onemklLayout blk_layout,
+                               onemklIndex index, int32_t *bsr_row_ptr, int32_t *bsr_col_ind,
+                               double *bsr_values);
+
+int onemklDsparse_set_bsr_data_64(syclQueue_t device_queue, matrix_handle_t spmat, int64_t
+                                  blk_nrows, int64_t blk_ncols, int64_t blk_nnz, int64_t
+                                  row_blk_size, int64_t col_blk_size, onemklLayout blk_layout,
+                                  onemklIndex index, int64_t *bsr_row_ptr, int64_t
+                                  *bsr_col_ind, double *bsr_values);
+
+int onemklCsparse_set_bsr_data(syclQueue_t device_queue, matrix_handle_t spmat, int64_t
+                               blk_nrows, int64_t blk_ncols, int64_t blk_nnz, int64_t
+                               row_blk_size, int64_t col_blk_size, onemklLayout blk_layout,
+                               onemklIndex index, int32_t *bsr_row_ptr, int32_t *bsr_col_ind,
+                               float _Complex *bsr_values);
+
+int onemklCsparse_set_bsr_data_64(syclQueue_t device_queue, matrix_handle_t spmat, int64_t
+                                  blk_nrows, int64_t blk_ncols, int64_t blk_nnz, int64_t
+                                  row_blk_size, int64_t col_blk_size, onemklLayout blk_layout,
+                                  onemklIndex index, int64_t *bsr_row_ptr, int64_t
+                                  *bsr_col_ind, float _Complex *bsr_values);
+
+int onemklZsparse_set_bsr_data(syclQueue_t device_queue, matrix_handle_t spmat, int64_t
+                               blk_nrows, int64_t blk_ncols, int64_t blk_nnz, int64_t
+                               row_blk_size, int64_t col_blk_size, onemklLayout blk_layout,
+                               onemklIndex index, int32_t *bsr_row_ptr, int32_t *bsr_col_ind,
+                               double _Complex *bsr_values);
+
+int onemklZsparse_set_bsr_data_64(syclQueue_t device_queue, matrix_handle_t spmat, int64_t
+                                  blk_nrows, int64_t blk_ncols, int64_t blk_nnz, int64_t
+                                  row_blk_size, int64_t col_blk_size, onemklLayout blk_layout,
+                                  onemklIndex index, int64_t *bsr_row_ptr, int64_t
+                                  *bsr_col_ind, double _Complex *bsr_values);
+
 int onemklXsparse_init_matmat_descr(matmat_descr_t *p_desc);
 
 int onemklXsparse_release_matmat_descr(matmat_descr_t *p_desc);
diff --git a/deps/src/onemkl_dft.cpp b/deps/src/onemkl_dft.cpp
index 9f39c127..8c10ffb7 100644
--- a/deps/src/onemkl_dft.cpp
+++ b/deps/src/onemkl_dft.cpp
@@ -116,10 +116,8 @@ static inline config_param to_param(onemklDftConfigParam p) {
         case ONEMKL_DFT_PARAM_NUMBER_OF_TRANSFORMS: return config_param::NUMBER_OF_TRANSFORMS;
         case ONEMKL_DFT_PARAM_COMPLEX_STORAGE: return config_param::COMPLEX_STORAGE;
         case ONEMKL_DFT_PARAM_PLACEMENT: return config_param::PLACEMENT;
-        // oneMKL >= 2026.0 dropped the deprecated INPUT_STRIDES/OUTPUT_STRIDES;
-        // map the legacy parameters onto their FWD_STRIDES/BWD_STRIDES successors.
-        case ONEMKL_DFT_PARAM_INPUT_STRIDES: return config_param::FWD_STRIDES;
-        case ONEMKL_DFT_PARAM_OUTPUT_STRIDES: return config_param::BWD_STRIDES;
+        case ONEMKL_DFT_PARAM_INPUT_STRIDES: return config_param::INPUT_STRIDES;
+        case ONEMKL_DFT_PARAM_OUTPUT_STRIDES: return config_param::OUTPUT_STRIDES;
         case ONEMKL_DFT_PARAM_FWD_DISTANCE: return config_param::FWD_DISTANCE;
         case ONEMKL_DFT_PARAM_BWD_DISTANCE: return config_param::BWD_DISTANCE;
         case ONEMKL_DFT_PARAM_WORKSPACE: return config_param::WORKSPACE;
@@ -212,28 +210,7 @@ int onemklDftGetValueInt64(onemklDftDescriptor_t desc, onemklDftConfigParam para
 
 int onemklDftGetValueDouble(onemklDftDescriptor_t desc, onemklDftConfigParam param, double *value) {
     if (!desc || !value) return -2; if (!desc->ptr) return -3;
-    try {
-        config_param p = to_param(param);
-        // oneMKL >= 2026.0 requires a pointer to the descriptor's real scalar type
-        // (float for single precision, double for double); use a matching temporary
-        // and widen to double for the C interface.
-        if (desc->prec == precision::SINGLE) {
-            float tmp;
-            if (desc->dom == domain::REAL)
-                static_cast< descriptor<precision::SINGLE, domain::REAL>* >(desc->ptr)->get_value(p, &tmp);
-            else
-                static_cast< descriptor<precision::SINGLE, domain::COMPLEX>* >(desc->ptr)->get_value(p, &tmp);
-            *value = static_cast<double>(tmp);
-        } else {
-            double tmp;
-            if (desc->dom == domain::REAL)
-                static_cast< descriptor<precision::DOUBLE, domain::REAL>* >(desc->ptr)->get_value(p, &tmp);
-            else
-                static_cast< descriptor<precision::DOUBLE, domain::COMPLEX>* >(desc->ptr)->get_value(p, &tmp);
-            *value = tmp;
-        }
-        return 0;
-    } catch (...) { return -1; }
+    try { ONEMKL_DFT_DISPATCH_CFG(desc->ptr, d->get_value(to_param(param), value)); return 0; } catch (...) { return -1; }
 }
 
 int onemklDftGetValueInt64Array(onemklDftDescriptor_t desc, onemklDftConfigParam param, int64_t *values, int64_t *n) {
@@ -466,8 +443,8 @@ int onemklDftQueryParamIndices(int64_t *out, int64_t n) {
             config_param::NUMBER_OF_TRANSFORMS,
             config_param::COMPLEX_STORAGE,
             config_param::PLACEMENT,
-            config_param::FWD_STRIDES,  // was INPUT_STRIDES (removed in oneMKL 2026.0)
-            config_param::BWD_STRIDES,  // was OUTPUT_STRIDES (removed in oneMKL 2026.0)
+            config_param::INPUT_STRIDES,
+            config_param::OUTPUT_STRIDES,
             config_param::FWD_DISTANCE,
             config_param::BWD_DISTANCE,
             config_param::WORKSPACE,
diff --git a/examples/bfloat16.jl b/examples/bfloat16.jl
index 3fd88da9..581bd0f3 100644
--- a/examples/bfloat16.jl
+++ b/examples/bfloat16.jl
@@ -14,6 +14,15 @@ if !bfloat16_supported
     exit()
 end
 
+# _device_supports_bfloat16() reflects hardware conversion support, but the Aurora LTS
+# SPIR-V stack (Khronos translator + NEO/IGC) cannot translate the `bfloat` LLVM type at
+# all: even a load/store forces SPV_KHR_bfloat16, which the LTS runtime rejects. So bf16
+# kernels can't be compiled there regardless of device support.
+if oneAPI.oneL0.LTS[]
+    @info "BFloat16 kernels are unsupported on the Aurora LTS SPIR-V stack, skipping."
+    exit()
+end
+
 # Conversions: Core.BFloat16 in Julia 1.12 may not have Float32 constructors yet
 float32_to_bf16(x::Float32) = reinterpret(Core.BFloat16, (reinterpret(UInt32, x) >> 16) % UInt16)
 bf16_to_float32(x::Core.BFloat16) = reinterpret(Float32, UInt32(reinterpret(UInt16, x)) << 16)
diff --git a/lib/level-zero/cmdlist.jl b/lib/level-zero/cmdlist.jl
index 133d2551..1df4bb90 100644
--- a/lib/level-zero/cmdlist.jl
+++ b/lib/level-zero/cmdlist.jl
@@ -47,8 +47,58 @@ function ZeCommandList(f::Base.Callable, args...; kwargs...)
     return list
 end
 
-execute!(queue::ZeCommandQueue, lists::Vector{ZeCommandList}, fence=nothing) =
-    zeCommandQueueExecuteCommandLists(queue, length(lists), lists, something(fence, C_NULL))
+# Opt-in workaround for the Aurora LTS NEO stack (set ONEAPI_SYNC_EACH_SUBMISSION=1).
+# Under heavy multi-process oversubscription of a single tile, a whole-queue
+# `zeCommandQueueSynchronize` does not reliably retire the tail of an earlier,
+# separately-submitted command list — producing silent "dropped tail" corruption (the
+# last work-item of a kernel / last element of a copy is missing). See
+# ISSUE_dropped_tail.md. Synchronizing after *every* submission eliminates it, at a large
+# throughput cost (~3x), so it is off by default and only enabled when correctness under
+# oversubscription matters more than speed.
+const SYNC_EACH_SUBMISSION = Ref{Bool}(false)
+
+"""
+    sync_each_submission() -> Bool
+
+Whether [`execute!`](@ref) follows every command-list submission with a full
+`zeCommandQueueSynchronize` (the Aurora LTS "dropped tail" workaround). See
+[`sync_each_submission!`](@ref).
+"""
+sync_each_submission() = SYNC_EACH_SUBMISSION[]
+
+"""
+    sync_each_submission!(enable::Bool) -> Bool
+
+Enable or disable synchronizing after every submission, returning the previous setting.
+Initialized from the `ONEAPI_SYNC_EACH_SUBMISSION` environment variable.
+"""
+function sync_each_submission!(enable::Bool)
+    old = SYNC_EACH_SUBMISSION[]
+    SYNC_EACH_SUBMISSION[] = enable
+    return old
+end
+
+"""
+    sync_each_submission(f, enable::Bool)
+
+Run `f()` with the workaround temporarily set to `enable`, restoring the previous setting
+afterwards. Use this for submit-then-signal patterns that would otherwise deadlock, where a
+synchronize is forced before the event gating the submitted work is signaled.
+"""
+function sync_each_submission(f::Base.Callable, enable::Bool)
+    old = sync_each_submission!(enable)
+    try
+        f()
+    finally
+        sync_each_submission!(old)
+    end
+end
+
+function execute!(queue::ZeCommandQueue, lists::Vector{ZeCommandList}, fence=nothing)
+    r = zeCommandQueueExecuteCommandLists(queue, length(lists), lists, something(fence, C_NULL))
+    sync_each_submission() && synchronize(queue)
+    return r
+end
 
 """
     execute!(queue::ZeCommandQueue, ...) do list
diff --git a/lib/level-zero/cmdqueue.jl b/lib/level-zero/cmdqueue.jl
index 3004a145..75ddb63b 100644
--- a/lib/level-zero/cmdqueue.jl
+++ b/lib/level-zero/cmdqueue.jl
@@ -2,6 +2,12 @@
 
 export ZeCommandQueue, synchronize
 
+# Bound on how long the LTS drain-before-destroy finalizer waits for in-flight work before
+# giving up and leaking the queue (see the finalizer below). 10 s comfortably covers any
+# kernel still legitimately running at finalization while keeping a never-signaled event
+# from hanging GC or process exit forever.
+const FINALIZER_SYNC_TIMEOUT_NS = UInt64(10_000_000_000)
+
 mutable struct ZeCommandQueue
     handle::ze_command_queue_handle_t
 
@@ -20,7 +26,30 @@ mutable struct ZeCommandQueue
         zeCommandQueueCreate(ctx, dev, desc_ref, handle_ref)
         obj = new(handle_ref[], ctx, dev, ordinal)
         finalizer(obj) do obj
+            if LTS[]
+                # the queue may still have work in flight (nothing requires a task to
+                # synchronize before dying), and zeCommandQueueDestroy does not wait for
+                # it: on the LTS NEO stack the still-running work then faults as soon as
+                # a referenced allocation is freed, getting the context banned. drain the
+                # queue first; unchecked, as sync on a banned context returns an error.
+                #
+                # Bounded wait, not typemax(UInt64): event-gated work whose event is never
+                # signaled (a task that submits work with a wait event and dies before
+                # signaling it) would make an infinite wait hang the finalizer forever, and
+                # with it GC and process exit. On timeout, leak the queue deliberately —
+                # destroying it now would trigger the very fault+ban this drain prevents.
+                if unchecked_zeCommandQueueSynchronize(obj, FINALIZER_SYNC_TIMEOUT_NS) == RESULT_NOT_READY
+                    @warn "Leaking a command queue still busy after $(FINALIZER_SYNC_TIMEOUT_NS ÷ 1_000_000_000)s to avoid blocking finalization (event-gated work whose event was never signaled?)" maxlog = 1
+                    return
+                end
+            end
             zeCommandQueueDestroy(obj)
+            if LTS[]
+                # mark the queue as destroyed: it can still be weakly reachable (e.g. from
+                # the queue registry used by `synchronize_all_queues`), and synchronizing a
+                # destroyed handle crashes in the driver.
+                obj.handle = ze_command_queue_handle_t(C_NULL)
+            end
         end
         obj
     end
diff --git a/lib/level-zero/oneL0.jl b/lib/level-zero/oneL0.jl
index 900f80de..69636412 100644
--- a/lib/level-zero/oneL0.jl
+++ b/lib/level-zero/oneL0.jl
@@ -11,8 +11,10 @@ using Libdl
 if Sys.iswindows()
     const libze_loader = "ze_loader"
 else
-    using NEO_jll
-    using oneAPI_Level_Zero_Loader_jll
+    using NEO_LTS_jll
+    using oneAPI_Level_Zero_Loader_LTS_jll
+    const NEO_jll = NEO_LTS_jll
+    const oneAPI_Level_Zero_Loader_jll = oneAPI_Level_Zero_Loader_LTS_jll
 end
 
 include("utils.jl")
@@ -137,10 +139,41 @@ const functional = Ref{Bool}(false)
 const validation_layer = Ref{Bool}()
 const parameter_validation = Ref{Bool}()
 
+# Master switch for the Intel LTS-stack workarounds (driver/IGC quirks on the Aurora
+# LTS NEO 25.18 stack): the SPIR-V translator codegen path, strided-reduction
+# materialization, and command-queue drain-before-free. All such code paths are gated
+# on `LTS[]`, so with it disabled the package behaves like the upstream rolling stack.
+# Default OFF (rolling stack), matching upstream: an LTS deployment opts in with
+# `ONEAPI_LTS=1`. On Aurora that variable is set in the environment, and the GitHub
+# Actions self-hosted (Aurora) runner sets it in ci.yml; the rolling-stack buildkite
+# runner gets the default and thus actually exercises the non-LTS (:llvm back-end) path.
+const LTS = Ref{Bool}(false)
+
+# Parse a boolean-valued environment variable, accepting the same spellings for every
+# oneAPI flag. Returns `default` when the variable is unset or empty. Warns (and returns
+# `default`) on an unrecognized value, so a deployment that writes `ONEAPI_SYNC_EACH_SUBMISSION=on`
+# — valid for the sibling `ONEAPI_LTS` — no longer silently reads as off, and a typo like
+# `ONEAPI_LTS=treu` no longer silently disables the LTS path.
+function parse_env_bool(name::AbstractString, default::Bool)
+    haskey(ENV, name) || return default
+    val = lowercase(strip(ENV[name]))
+    isempty(val) && return default
+    val in ("1", "true", "yes", "on") && return true
+    val in ("0", "false", "no", "off") && return false
+    @warn "Ignoring unrecognized boolean value for $name; using default" value = ENV[name] default
+    return default
+end
+
 function __init__()
     precompiling = ccall(:jl_generating_output, Cint, ()) != 0
     precompiling && return
 
+    # Resolve the LTS master switch up front, before the driver-availability early
+    # returns below: it gates codegen and behavior and must be set even on hosts
+    # without a functional GPU. Default off (rolling stack); an LTS deployment such as
+    # Aurora opts in with ONEAPI_LTS=1.
+    LTS[] = parse_env_bool("ONEAPI_LTS", false)
+
     if Sys.iswindows()
         if Libdl.dlopen(libze_loader; throw_error=false) === nothing
             @error "The oneAPI Level Zero loader was not found. Please ensure the Intel GPU drivers are installed."
@@ -199,6 +232,7 @@ function __init__()
 
     validation_layer[] = parse(Bool, get(ENV, "ZE_ENABLE_VALIDATION_LAYER", "false"))
     parameter_validation[] = parse(Bool, get(ENV, "ZE_ENABLE_PARAMETER_VALIDATION", "false"))
+    sync_each_submission!(parse_env_bool("ONEAPI_SYNC_EACH_SUBMISSION", false))
 end
 
 end
diff --git a/lib/mkl/fft.jl b/lib/mkl/fft.jl
index 4429801b..0ccdca46 100644
--- a/lib/mkl/fft.jl
+++ b/lib/mkl/fft.jl
@@ -93,23 +93,22 @@ function _create_descriptor(sz::NTuple{N,Int}, T::Type, complex::Bool) where {N}
     prec = T<:Float64 || T<:ComplexF64 ? ONEMKL_DFT_PRECISION_DOUBLE : ONEMKL_DFT_PRECISION_SINGLE
     dom = complex ? ONEMKL_DFT_DOMAIN_COMPLEX : ONEMKL_DFT_DOMAIN_REAL
     desc_ref = Ref{onemklDftDescriptor_t}()
-    # Create descriptor for the full array dimensions. `lengths` must stay rooted
-    # across the ccall: oneMKL copies the dimensions out of `pointer(lengths)`, and
-    # without GC.@preserve the array can be collected first, leaving the descriptor
-    # with garbage dimensions (commit then fails with FFT_INVALID_DESCRIPTOR).
+    # Create descriptor for the full array dimensions
     lengths = collect(Int64, sz)
-    st = GC.@preserve lengths (length(lengths) == 1 ?
-        onemklDftCreate1D(desc_ref, prec, dom, lengths[1]) :
-        onemklDftCreateND(desc_ref, prec, dom, length(lengths), pointer(lengths)))
+    # NB: pass the arrays themselves to the ccall wrappers (NOT `pointer(...)`) so they
+    # are rooted for the duration of the call; passing a raw Ptr lets the GC collect the
+    # vector mid-call, which made MKL read garbage lengths/strides and fail `commit` with
+    # invalid_descriptor_exception or a SIGFPE, depending on heap reuse.
+    st = length(lengths) == 1 ? onemklDftCreate1D(desc_ref, prec, dom, lengths[1]) : onemklDftCreateND(desc_ref, prec, dom, length(lengths), lengths)
     st == 0 || error("onemkl DFT create failed (status $st)")
     desc = desc_ref[]
     # Do not program descriptor scaling; we'll perform inverse normalization manually.
     # Set placement explicitly based on plan type later
-    # Use the task-local cached SYCL queue wrapping the global Level Zero queue, like the
-    # other oneMKL wrappers do. Creating fresh syclContext/syclQueue objects per plan is
-    # unsound: once they become garbage their finalizers (syclQueueDestroy etc.) tear down
-    # SYCL runtime state for the still-in-use underlying queue, corrupting later DFT
-    # commits and crashing at process exit.
+    # Use the task-local cached SYCL queue wrapping the global Level Zero queue, like
+    # the other oneMKL wrappers do. Creating fresh syclContext/syclQueue objects per
+    # plan is unsound: once they become garbage their finalizers (syclQueueDestroy etc.)
+    # tear down SYCL runtime state for the still-in-use underlying queue, corrupting
+    # later DFT commits (SIGFPE) and crashing at process exit.
     ze_ctx = oneAPI.context(); ze_dev = oneAPI.device()
     q = oneAPI.sycl_queue(oneAPI.global_queue(ze_ctx, ze_dev))
     return desc, q
@@ -132,10 +131,8 @@ function plan_fft(X::oneAPI.oneArray{T,N}, region) where {T<:Union{ComplexF32,Co
             strides[i+1] = prod
             prod *= size(X,i)
         end
-        GC.@preserve strides begin
-            onemklDftSetValueInt64Array(desc, ONEMKL_DFT_PARAM_FWD_STRIDES, pointer(strides), length(strides))
-            onemklDftSetValueInt64Array(desc, ONEMKL_DFT_PARAM_BWD_STRIDES, pointer(strides), length(strides))
-        end
+        onemklDftSetValueInt64Array(desc, ONEMKL_DFT_PARAM_FWD_STRIDES, strides, length(strides))
+        onemklDftSetValueInt64Array(desc, ONEMKL_DFT_PARAM_BWD_STRIDES, strides, length(strides))
     end
     stc = onemklDftCommit(desc, q); stc == 0 || error("commit failed ($stc)")
     return cMKLFFTPlan{T,MKLFFT_FORWARD,false,N,R,Nothing}(desc,q,size(X),size(X),false,reg,nothing,nothing)
@@ -153,10 +150,8 @@ function plan_bfft(X::oneAPI.oneArray{T,N}, region) where {T<:Union{ComplexF32,C
         @inbounds for i in 1:N
             strides[i+1]=prod; prod*=size(X,i)
         end
-        GC.@preserve strides begin
-            onemklDftSetValueInt64Array(desc, ONEMKL_DFT_PARAM_FWD_STRIDES, pointer(strides), length(strides))
-            onemklDftSetValueInt64Array(desc, ONEMKL_DFT_PARAM_BWD_STRIDES, pointer(strides), length(strides))
-        end
+        onemklDftSetValueInt64Array(desc, ONEMKL_DFT_PARAM_FWD_STRIDES, strides, length(strides))
+        onemklDftSetValueInt64Array(desc, ONEMKL_DFT_PARAM_BWD_STRIDES, strides, length(strides))
     end
     stc = onemklDftCommit(desc, q); stc == 0 || error("commit failed ($stc)")
     return cMKLFFTPlan{T,MKLFFT_INVERSE,false,N,R,Nothing}(desc,q,size(X),size(X),false,reg,nothing,nothing)
@@ -176,10 +171,8 @@ function plan_fft!(X::oneAPI.oneArray{T,N}, region) where {T<:Union{ComplexF32,C
         @inbounds for i in 1:N
             strides[i+1]=prod; prod*=size(X,i)
         end
-        GC.@preserve strides begin
-            onemklDftSetValueInt64Array(desc, ONEMKL_DFT_PARAM_FWD_STRIDES, pointer(strides), length(strides))
-            onemklDftSetValueInt64Array(desc, ONEMKL_DFT_PARAM_BWD_STRIDES, pointer(strides), length(strides))
-        end
+        onemklDftSetValueInt64Array(desc, ONEMKL_DFT_PARAM_FWD_STRIDES, strides, length(strides))
+        onemklDftSetValueInt64Array(desc, ONEMKL_DFT_PARAM_BWD_STRIDES, strides, length(strides))
     end
     stc = onemklDftCommit(desc, q); stc == 0 || error("commit failed ($stc)")
     cMKLFFTPlan{T,MKLFFT_FORWARD,true,N,R,Nothing}(desc,q,size(X),size(X),false,reg,nothing,nothing)
@@ -197,10 +190,8 @@ function plan_bfft!(X::oneAPI.oneArray{T,N}, region) where {T<:Union{ComplexF32,
         @inbounds for i in 1:N
             strides[i+1]=prod; prod*=size(X,i)
         end
-        GC.@preserve strides begin
-            onemklDftSetValueInt64Array(desc, ONEMKL_DFT_PARAM_FWD_STRIDES, pointer(strides), length(strides))
-            onemklDftSetValueInt64Array(desc, ONEMKL_DFT_PARAM_BWD_STRIDES, pointer(strides), length(strides))
-        end
+        onemklDftSetValueInt64Array(desc, ONEMKL_DFT_PARAM_FWD_STRIDES, strides, length(strides))
+        onemklDftSetValueInt64Array(desc, ONEMKL_DFT_PARAM_BWD_STRIDES, strides, length(strides))
     end
     stc = onemklDftCommit(desc, q); stc == 0 || error("commit failed ($stc)")
     cMKLFFTPlan{T,MKLFFT_INVERSE,true,N,R,Nothing}(desc,q,size(X),size(X),false,reg,nothing,nothing)
diff --git a/lib/mkl/wrappers_sparse.jl b/lib/mkl/wrappers_sparse.jl
index 8e58956b..3c112027 100644
--- a/lib/mkl/wrappers_sparse.jl
+++ b/lib/mkl/wrappers_sparse.jl
@@ -60,7 +60,8 @@ for (fname, elty, intty) in ((:onemklSsparse_set_csr_data   , :Float32   , :Int3
             queue = global_queue(context(nzVal), device(nzVal))
             # Don't update handle if matrix is empty
             if m != 0 && n != 0
-                $fname(sycl_queue(queue), handle_ptr[], m, n, 'O', rowPtr, colVal, nzVal)
+                Support._check_sparse_abi()
+                $fname(sycl_queue(queue), handle_ptr[], m, n, nnzA, 'O', rowPtr, colVal, nzVal)
                 dA = oneSparseMatrixCSR{$elty, $intty}(handle_ptr[], rowPtr, colVal, nzVal, (m, n), nnzA)
                 finalizer(sparse_release_matrix_handle, dA)
             else
@@ -81,7 +82,8 @@ for (fname, elty, intty) in ((:onemklSsparse_set_csr_data   , :Float32   , :Int3
             nnzA = length(nzVal)
             # Don't update handle if matrix is empty
             if m != 0 && n != 0
-                $fname(sycl_queue(queue), handle_ptr[], n, m, 'O', colPtr, rowVal, nzVal)  # CSC of A is CSR of Aᵀ
+                Support._check_sparse_abi()
+                $fname(sycl_queue(queue), handle_ptr[], n, m, nnzA, 'O', colPtr, rowVal, nzVal)  # CSC of A is CSR of Aᵀ
                 dA = oneSparseMatrixCSC{$elty, $intty}(handle_ptr[], colPtr, rowVal, nzVal, (m, n), nnzA)
                 finalizer(sparse_release_matrix_handle, dA)
             else
@@ -144,6 +146,7 @@ for (fname, elty, intty) in ((:onemklSsparse_set_coo_data   , :Float32   , :Int3
             nnzA = length(val)
             queue = global_queue(context(nzVal), device(nzVal))
             if m != 0 && n != 0
+                Support._check_sparse_abi()
                 $fname(sycl_queue(queue), handle_ptr[], m, n, nnzA, 'O', rowInd, colInd, nzVal)
                 dA = oneSparseMatrixCOO{$elty, $intty}(handle_ptr[], rowInd, colInd, nzVal, (m, n), nnzA)
                 finalizer(sparse_release_matrix_handle, dA)
diff --git a/lib/support/Support.jl b/lib/support/Support.jl
index 418eac4e..5fe9fc35 100644
--- a/lib/support/Support.jl
+++ b/lib/support/Support.jl
@@ -10,6 +10,8 @@ using ..oneL0:
 
 using oneAPI_Support_jll
 
+import Libdl
+
 include("liboneapi_support.jl")
 
 # export everything
@@ -19,14 +21,42 @@ for n in names(@__MODULE__; all=true)
     end
 end
 
+# The sparse wrappers were regenerated with a wider ABI: the onemklX sparse_set_{csr,csc,coo}
+# setters gained a 64-bit `nnz` argument and 64-bit dims, and new set_csc_data / set_bsr_data
+# entry points were added. These only match a liboneapi_support built from the updated
+# onemkl.cpp/onemkl.h. The oneAPI_Support_jll 0.9.2 binary shipped in the registry predates
+# them, so calling a sparse setter against it shifts every pointer one slot and MKL
+# dereferences garbage -> segfault or silently corrupted sparse data. `deps/build_local.jl`
+# rebuilds a matching binary; a plain `Pkg.add` / CI `julia-buildpkg` does not. Probe once at
+# load for a symbol that only the rebuilt binary exports and turn the crash into a clear error.
+const _sparse_abi_ok = Ref(false)
+
+function _check_sparse_abi()
+    _sparse_abi_ok[] && return
+    error(
+        "The loaded oneAPI_Support_jll binary predates the current sparse wrappers and has an " *
+        "incompatible ABI; calling sparse routines against it would corrupt memory or crash. " *
+        "Rebuild the support library from source with `julia --project deps/build_local.jl` " *
+        "(or install a oneAPI_Support_jll built from the updated onemkl.cpp)."
+    )
+end
+
 function __init__()
     precompiling = ccall(:jl_generating_output, Cint, ()) != 0
     precompiling && return
-    
+
     if !oneAPI_Support_jll.is_available()
         @error """oneAPI support wrapper not available for your platform."""
         return
     end
+
+    # Probe the loaded support binary for a symbol that only the regenerated (wide-ABI)
+    # sparse wrappers export; `_check_sparse_abi` gates the affected calls on the result.
+    handle = Libdl.dlopen(oneAPI_Support_jll.liboneapi_support; throw_error = false)
+    if handle !== nothing
+        _sparse_abi_ok[] =
+            Libdl.dlsym(handle, :onemklSsparse_set_csc_data; throw_error = false) !== nothing
+    end
 end
 
 end
diff --git a/lib/support/liboneapi_support.jl b/lib/support/liboneapi_support.jl
index 91146c99..e1b7327f 100644
--- a/lib/support/liboneapi_support.jl
+++ b/lib/support/liboneapi_support.jl
@@ -6428,107 +6428,195 @@ function onemklZunmqr_batch_scratchpad_size(device_queue, side, trans, m, n, k,
                                                                 group_sizes::Ptr{Int64})::Int64
 end
 
-function onemklXsparse_init_matrix_handle(p_spMat)
-    @ccall liboneapi_support.onemklXsparse_init_matrix_handle(p_spMat::Ptr{matrix_handle_t})::Cint
+function onemklXsparse_init_matrix_handle(p_spmat)
+    @ccall liboneapi_support.onemklXsparse_init_matrix_handle(p_spmat::Ptr{matrix_handle_t})::Cint
 end
 
-function onemklXsparse_release_matrix_handle(device_queue, p_spMat)
+function onemklXsparse_release_matrix_handle(device_queue, p_spmat)
     @ccall liboneapi_support.onemklXsparse_release_matrix_handle(device_queue::syclQueue_t,
-                                                                 p_spMat::Ptr{matrix_handle_t})::Cint
+                                                                 p_spmat::Ptr{matrix_handle_t})::Cint
 end
 
-function onemklSsparse_set_csr_data(device_queue, spMat, nrows, ncols, index, row_ptr,
+function onemklSsparse_set_csr_data(device_queue, spmat, nrows, ncols, nnz, index, row_ptr,
                                     col_ind, values)
     @ccall liboneapi_support.onemklSsparse_set_csr_data(device_queue::syclQueue_t,
-                                                        spMat::matrix_handle_t,
-                                                        nrows::Int32, ncols::Int32,
-                                                        index::onemklIndex,
+                                                        spmat::matrix_handle_t,
+                                                        nrows::Int64, ncols::Int64,
+                                                        nnz::Int64, index::onemklIndex,
                                                         row_ptr::ZePtr{Int32},
                                                         col_ind::ZePtr{Int32},
                                                         values::ZePtr{Cfloat})::Cint
 end
 
-function onemklSsparse_set_csr_data_64(device_queue, spMat, nrows, ncols, index, row_ptr,
-                                       col_ind, values)
+function onemklSsparse_set_csr_data_64(device_queue, spmat, nrows, ncols, nnz, index,
+                                       row_ptr, col_ind, values)
     @ccall liboneapi_support.onemklSsparse_set_csr_data_64(device_queue::syclQueue_t,
-                                                           spMat::matrix_handle_t,
+                                                           spmat::matrix_handle_t,
                                                            nrows::Int64, ncols::Int64,
-                                                           index::onemklIndex,
+                                                           nnz::Int64, index::onemklIndex,
                                                            row_ptr::ZePtr{Int64},
                                                            col_ind::ZePtr{Int64},
                                                            values::ZePtr{Cfloat})::Cint
 end
 
-function onemklDsparse_set_csr_data(device_queue, spMat, nrows, ncols, index, row_ptr,
+function onemklDsparse_set_csr_data(device_queue, spmat, nrows, ncols, nnz, index, row_ptr,
                                     col_ind, values)
     @ccall liboneapi_support.onemklDsparse_set_csr_data(device_queue::syclQueue_t,
-                                                        spMat::matrix_handle_t,
-                                                        nrows::Int32, ncols::Int32,
-                                                        index::onemklIndex,
+                                                        spmat::matrix_handle_t,
+                                                        nrows::Int64, ncols::Int64,
+                                                        nnz::Int64, index::onemklIndex,
                                                         row_ptr::ZePtr{Int32},
                                                         col_ind::ZePtr{Int32},
                                                         values::ZePtr{Cdouble})::Cint
 end
 
-function onemklDsparse_set_csr_data_64(device_queue, spMat, nrows, ncols, index, row_ptr,
-                                       col_ind, values)
+function onemklDsparse_set_csr_data_64(device_queue, spmat, nrows, ncols, nnz, index,
+                                       row_ptr, col_ind, values)
     @ccall liboneapi_support.onemklDsparse_set_csr_data_64(device_queue::syclQueue_t,
-                                                           spMat::matrix_handle_t,
+                                                           spmat::matrix_handle_t,
                                                            nrows::Int64, ncols::Int64,
-                                                           index::onemklIndex,
+                                                           nnz::Int64, index::onemklIndex,
                                                            row_ptr::ZePtr{Int64},
                                                            col_ind::ZePtr{Int64},
                                                            values::ZePtr{Cdouble})::Cint
 end
 
-function onemklCsparse_set_csr_data(device_queue, spMat, nrows, ncols, index, row_ptr,
+function onemklCsparse_set_csr_data(device_queue, spmat, nrows, ncols, nnz, index, row_ptr,
                                     col_ind, values)
     @ccall liboneapi_support.onemklCsparse_set_csr_data(device_queue::syclQueue_t,
-                                                        spMat::matrix_handle_t,
-                                                        nrows::Int32, ncols::Int32,
-                                                        index::onemklIndex,
+                                                        spmat::matrix_handle_t,
+                                                        nrows::Int64, ncols::Int64,
+                                                        nnz::Int64, index::onemklIndex,
                                                         row_ptr::ZePtr{Int32},
                                                         col_ind::ZePtr{Int32},
                                                         values::ZePtr{ComplexF32})::Cint
 end
 
-function onemklCsparse_set_csr_data_64(device_queue, spMat, nrows, ncols, index, row_ptr,
-                                       col_ind, values)
+function onemklCsparse_set_csr_data_64(device_queue, spmat, nrows, ncols, nnz, index,
+                                       row_ptr, col_ind, values)
     @ccall liboneapi_support.onemklCsparse_set_csr_data_64(device_queue::syclQueue_t,
-                                                           spMat::matrix_handle_t,
+                                                           spmat::matrix_handle_t,
                                                            nrows::Int64, ncols::Int64,
-                                                           index::onemklIndex,
+                                                           nnz::Int64, index::onemklIndex,
                                                            row_ptr::ZePtr{Int64},
                                                            col_ind::ZePtr{Int64},
                                                            values::ZePtr{ComplexF32})::Cint
 end
 
-function onemklZsparse_set_csr_data(device_queue, spMat, nrows, ncols, index, row_ptr,
+function onemklZsparse_set_csr_data(device_queue, spmat, nrows, ncols, nnz, index, row_ptr,
                                     col_ind, values)
     @ccall liboneapi_support.onemklZsparse_set_csr_data(device_queue::syclQueue_t,
-                                                        spMat::matrix_handle_t,
-                                                        nrows::Int32, ncols::Int32,
-                                                        index::onemklIndex,
+                                                        spmat::matrix_handle_t,
+                                                        nrows::Int64, ncols::Int64,
+                                                        nnz::Int64, index::onemklIndex,
                                                         row_ptr::ZePtr{Int32},
                                                         col_ind::ZePtr{Int32},
                                                         values::ZePtr{ComplexF64})::Cint
 end
 
-function onemklZsparse_set_csr_data_64(device_queue, spMat, nrows, ncols, index, row_ptr,
-                                       col_ind, values)
+function onemklZsparse_set_csr_data_64(device_queue, spmat, nrows, ncols, nnz, index,
+                                       row_ptr, col_ind, values)
     @ccall liboneapi_support.onemklZsparse_set_csr_data_64(device_queue::syclQueue_t,
-                                                           spMat::matrix_handle_t,
+                                                           spmat::matrix_handle_t,
                                                            nrows::Int64, ncols::Int64,
-                                                           index::onemklIndex,
+                                                           nnz::Int64, index::onemklIndex,
                                                            row_ptr::ZePtr{Int64},
                                                            col_ind::ZePtr{Int64},
                                                            values::ZePtr{ComplexF64})::Cint
 end
 
-function onemklSsparse_set_coo_data(device_queue, spMat, nrows, ncols, nnz, index, row_ind,
+function onemklSsparse_set_csc_data(device_queue, spMat, nrows, ncols, nnz, index, col_ptr,
+                                    row_ind, values)
+    @ccall liboneapi_support.onemklSsparse_set_csc_data(device_queue::syclQueue_t,
+                                                        spMat::matrix_handle_t,
+                                                        nrows::Int64, ncols::Int64,
+                                                        nnz::Int64, index::onemklIndex,
+                                                        col_ptr::Ptr{Int32},
+                                                        row_ind::Ptr{Int32},
+                                                        values::Ptr{Cfloat})::Cint
+end
+
+function onemklSsparse_set_csc_data_64(device_queue, spMat, nrows, ncols, nnz, index,
+                                       col_ptr, row_ind, values)
+    @ccall liboneapi_support.onemklSsparse_set_csc_data_64(device_queue::syclQueue_t,
+                                                           spMat::matrix_handle_t,
+                                                           nrows::Int64, ncols::Int64,
+                                                           nnz::Int64, index::onemklIndex,
+                                                           col_ptr::Ptr{Int64},
+                                                           row_ind::Ptr{Int64},
+                                                           values::Ptr{Cfloat})::Cint
+end
+
+function onemklDsparse_set_csc_data(device_queue, spMat, nrows, ncols, nnz, index, col_ptr,
+                                    row_ind, values)
+    @ccall liboneapi_support.onemklDsparse_set_csc_data(device_queue::syclQueue_t,
+                                                        spMat::matrix_handle_t,
+                                                        nrows::Int64, ncols::Int64,
+                                                        nnz::Int64, index::onemklIndex,
+                                                        col_ptr::Ptr{Int32},
+                                                        row_ind::Ptr{Int32},
+                                                        values::Ptr{Cdouble})::Cint
+end
+
+function onemklDsparse_set_csc_data_64(device_queue, spMat, nrows, ncols, nnz, index,
+                                       col_ptr, row_ind, values)
+    @ccall liboneapi_support.onemklDsparse_set_csc_data_64(device_queue::syclQueue_t,
+                                                           spMat::matrix_handle_t,
+                                                           nrows::Int64, ncols::Int64,
+                                                           nnz::Int64, index::onemklIndex,
+                                                           col_ptr::Ptr{Int64},
+                                                           row_ind::Ptr{Int64},
+                                                           values::Ptr{Cdouble})::Cint
+end
+
+function onemklCsparse_set_csc_data(device_queue, spMat, nrows, ncols, nnz, index, col_ptr,
+                                    row_ind, values)
+    @ccall liboneapi_support.onemklCsparse_set_csc_data(device_queue::syclQueue_t,
+                                                        spMat::matrix_handle_t,
+                                                        nrows::Int64, ncols::Int64,
+                                                        nnz::Int64, index::onemklIndex,
+                                                        col_ptr::Ptr{Int32},
+                                                        row_ind::Ptr{Int32},
+                                                        values::Ptr{ComplexF32})::Cint
+end
+
+function onemklCsparse_set_csc_data_64(device_queue, spMat, nrows, ncols, nnz, index,
+                                       col_ptr, row_ind, values)
+    @ccall liboneapi_support.onemklCsparse_set_csc_data_64(device_queue::syclQueue_t,
+                                                           spMat::matrix_handle_t,
+                                                           nrows::Int64, ncols::Int64,
+                                                           nnz::Int64, index::onemklIndex,
+                                                           col_ptr::Ptr{Int64},
+                                                           row_ind::Ptr{Int64},
+                                                           values::Ptr{ComplexF32})::Cint
+end
+
+function onemklZsparse_set_csc_data(device_queue, spMat, nrows, ncols, nnz, index, col_ptr,
+                                    row_ind, values)
+    @ccall liboneapi_support.onemklZsparse_set_csc_data(device_queue::syclQueue_t,
+                                                        spMat::matrix_handle_t,
+                                                        nrows::Int64, ncols::Int64,
+                                                        nnz::Int64, index::onemklIndex,
+                                                        col_ptr::Ptr{Int32},
+                                                        row_ind::Ptr{Int32},
+                                                        values::Ptr{ComplexF32})::Cint
+end
+
+function onemklZsparse_set_csc_data_64(device_queue, spMat, nrows, ncols, nnz, index,
+                                       col_ptr, row_ind, values)
+    @ccall liboneapi_support.onemklZsparse_set_csc_data_64(device_queue::syclQueue_t,
+                                                           spMat::matrix_handle_t,
+                                                           nrows::Int64, ncols::Int64,
+                                                           nnz::Int64, index::onemklIndex,
+                                                           col_ptr::Ptr{Int64},
+                                                           row_ind::Ptr{Int64},
+                                                           values::Ptr{ComplexF32})::Cint
+end
+
+function onemklSsparse_set_coo_data(device_queue, spmat, nrows, ncols, nnz, index, row_ind,
                                     col_ind, values)
     @ccall liboneapi_support.onemklSsparse_set_coo_data(device_queue::syclQueue_t,
-                                                        spMat::matrix_handle_t,
+                                                        spmat::matrix_handle_t,
                                                         nrows::Int32, ncols::Int32,
                                                         nnz::Int32, index::onemklIndex,
                                                         row_ind::ZePtr{Int32},
@@ -6536,10 +6624,10 @@ function onemklSsparse_set_coo_data(device_queue, spMat, nrows, ncols, nnz, inde
                                                         values::ZePtr{Cfloat})::Cint
 end
 
-function onemklSsparse_set_coo_data_64(device_queue, spMat, nrows, ncols, nnz, index,
+function onemklSsparse_set_coo_data_64(device_queue, spmat, nrows, ncols, nnz, index,
                                        row_ind, col_ind, values)
     @ccall liboneapi_support.onemklSsparse_set_coo_data_64(device_queue::syclQueue_t,
-                                                           spMat::matrix_handle_t,
+                                                           spmat::matrix_handle_t,
                                                            nrows::Int64, ncols::Int64,
                                                            nnz::Int64, index::onemklIndex,
                                                            row_ind::ZePtr{Int64},
@@ -6547,10 +6635,10 @@ function onemklSsparse_set_coo_data_64(device_queue, spMat, nrows, ncols, nnz, i
                                                            values::ZePtr{Cfloat})::Cint
 end
 
-function onemklDsparse_set_coo_data(device_queue, spMat, nrows, ncols, nnz, index, row_ind,
+function onemklDsparse_set_coo_data(device_queue, spmat, nrows, ncols, nnz, index, row_ind,
                                     col_ind, values)
     @ccall liboneapi_support.onemklDsparse_set_coo_data(device_queue::syclQueue_t,
-                                                        spMat::matrix_handle_t,
+                                                        spmat::matrix_handle_t,
                                                         nrows::Int32, ncols::Int32,
                                                         nnz::Int32, index::onemklIndex,
                                                         row_ind::ZePtr{Int32},
@@ -6558,10 +6646,10 @@ function onemklDsparse_set_coo_data(device_queue, spMat, nrows, ncols, nnz, inde
                                                         values::ZePtr{Cdouble})::Cint
 end
 
-function onemklDsparse_set_coo_data_64(device_queue, spMat, nrows, ncols, nnz, index,
+function onemklDsparse_set_coo_data_64(device_queue, spmat, nrows, ncols, nnz, index,
                                        row_ind, col_ind, values)
     @ccall liboneapi_support.onemklDsparse_set_coo_data_64(device_queue::syclQueue_t,
-                                                           spMat::matrix_handle_t,
+                                                           spmat::matrix_handle_t,
                                                            nrows::Int64, ncols::Int64,
                                                            nnz::Int64, index::onemklIndex,
                                                            row_ind::ZePtr{Int64},
@@ -6569,10 +6657,10 @@ function onemklDsparse_set_coo_data_64(device_queue, spMat, nrows, ncols, nnz, i
                                                            values::ZePtr{Cdouble})::Cint
 end
 
-function onemklCsparse_set_coo_data(device_queue, spMat, nrows, ncols, nnz, index, row_ind,
+function onemklCsparse_set_coo_data(device_queue, spmat, nrows, ncols, nnz, index, row_ind,
                                     col_ind, values)
     @ccall liboneapi_support.onemklCsparse_set_coo_data(device_queue::syclQueue_t,
-                                                        spMat::matrix_handle_t,
+                                                        spmat::matrix_handle_t,
                                                         nrows::Int32, ncols::Int32,
                                                         nnz::Int32, index::onemklIndex,
                                                         row_ind::ZePtr{Int32},
@@ -6580,10 +6668,10 @@ function onemklCsparse_set_coo_data(device_queue, spMat, nrows, ncols, nnz, inde
                                                         values::ZePtr{ComplexF32})::Cint
 end
 
-function onemklCsparse_set_coo_data_64(device_queue, spMat, nrows, ncols, nnz, index,
+function onemklCsparse_set_coo_data_64(device_queue, spmat, nrows, ncols, nnz, index,
                                        row_ind, col_ind, values)
     @ccall liboneapi_support.onemklCsparse_set_coo_data_64(device_queue::syclQueue_t,
-                                                           spMat::matrix_handle_t,
+                                                           spmat::matrix_handle_t,
                                                            nrows::Int64, ncols::Int64,
                                                            nnz::Int64, index::onemklIndex,
                                                            row_ind::ZePtr{Int64},
@@ -6591,10 +6679,10 @@ function onemklCsparse_set_coo_data_64(device_queue, spMat, nrows, ncols, nnz, i
                                                            values::ZePtr{ComplexF32})::Cint
 end
 
-function onemklZsparse_set_coo_data(device_queue, spMat, nrows, ncols, nnz, index, row_ind,
+function onemklZsparse_set_coo_data(device_queue, spmat, nrows, ncols, nnz, index, row_ind,
                                     col_ind, values)
     @ccall liboneapi_support.onemklZsparse_set_coo_data(device_queue::syclQueue_t,
-                                                        spMat::matrix_handle_t,
+                                                        spmat::matrix_handle_t,
                                                         nrows::Int32, ncols::Int32,
                                                         nnz::Int32, index::onemklIndex,
                                                         row_ind::ZePtr{Int32},
@@ -6602,10 +6690,10 @@ function onemklZsparse_set_coo_data(device_queue, spMat, nrows, ncols, nnz, inde
                                                         values::ZePtr{ComplexF64})::Cint
 end
 
-function onemklZsparse_set_coo_data_64(device_queue, spMat, nrows, ncols, nnz, index,
+function onemklZsparse_set_coo_data_64(device_queue, spmat, nrows, ncols, nnz, index,
                                        row_ind, col_ind, values)
     @ccall liboneapi_support.onemklZsparse_set_coo_data_64(device_queue::syclQueue_t,
-                                                           spMat::matrix_handle_t,
+                                                           spmat::matrix_handle_t,
                                                            nrows::Int64, ncols::Int64,
                                                            nnz::Int64, index::onemklIndex,
                                                            row_ind::ZePtr{Int64},
@@ -6613,6 +6701,130 @@ function onemklZsparse_set_coo_data_64(device_queue, spMat, nrows, ncols, nnz, i
                                                            values::ZePtr{ComplexF64})::Cint
 end
 
+function onemklSsparse_set_bsr_data(device_queue, spmat, blk_nrows, blk_ncols, blk_nnz,
+                                    row_blk_size, col_blk_size, blk_layout, index,
+                                    bsr_row_ptr, bsr_col_ind, bsr_values)
+    @ccall liboneapi_support.onemklSsparse_set_bsr_data(device_queue::syclQueue_t,
+                                                        spmat::matrix_handle_t,
+                                                        blk_nrows::Int64, blk_ncols::Int64,
+                                                        blk_nnz::Int64, row_blk_size::Int64,
+                                                        col_blk_size::Int64,
+                                                        blk_layout::onemklLayout,
+                                                        index::onemklIndex,
+                                                        bsr_row_ptr::Ptr{Int32},
+                                                        bsr_col_ind::Ptr{Int32},
+                                                        bsr_values::Ptr{Cfloat})::Cint
+end
+
+function onemklSsparse_set_bsr_data_64(device_queue, spmat, blk_nrows, blk_ncols, blk_nnz,
+                                       row_blk_size, col_blk_size, blk_layout, index,
+                                       bsr_row_ptr, bsr_col_ind, bsr_values)
+    @ccall liboneapi_support.onemklSsparse_set_bsr_data_64(device_queue::syclQueue_t,
+                                                           spmat::matrix_handle_t,
+                                                           blk_nrows::Int64,
+                                                           blk_ncols::Int64, blk_nnz::Int64,
+                                                           row_blk_size::Int64,
+                                                           col_blk_size::Int64,
+                                                           blk_layout::onemklLayout,
+                                                           index::onemklIndex,
+                                                           bsr_row_ptr::Ptr{Int64},
+                                                           bsr_col_ind::Ptr{Int64},
+                                                           bsr_values::Ptr{Cfloat})::Cint
+end
+
+function onemklDsparse_set_bsr_data(device_queue, spmat, blk_nrows, blk_ncols, blk_nnz,
+                                    row_blk_size, col_blk_size, blk_layout, index,
+                                    bsr_row_ptr, bsr_col_ind, bsr_values)
+    @ccall liboneapi_support.onemklDsparse_set_bsr_data(device_queue::syclQueue_t,
+                                                        spmat::matrix_handle_t,
+                                                        blk_nrows::Int64, blk_ncols::Int64,
+                                                        blk_nnz::Int64, row_blk_size::Int64,
+                                                        col_blk_size::Int64,
+                                                        blk_layout::onemklLayout,
+                                                        index::onemklIndex,
+                                                        bsr_row_ptr::Ptr{Int32},
+                                                        bsr_col_ind::Ptr{Int32},
+                                                        bsr_values::Ptr{Cdouble})::Cint
+end
+
+function onemklDsparse_set_bsr_data_64(device_queue, spmat, blk_nrows, blk_ncols, blk_nnz,
+                                       row_blk_size, col_blk_size, blk_layout, index,
+                                       bsr_row_ptr, bsr_col_ind, bsr_values)
+    @ccall liboneapi_support.onemklDsparse_set_bsr_data_64(device_queue::syclQueue_t,
+                                                           spmat::matrix_handle_t,
+                                                           blk_nrows::Int64,
+                                                           blk_ncols::Int64, blk_nnz::Int64,
+                                                           row_blk_size::Int64,
+                                                           col_blk_size::Int64,
+                                                           blk_layout::onemklLayout,
+                                                           index::onemklIndex,
+                                                           bsr_row_ptr::Ptr{Int64},
+                                                           bsr_col_ind::Ptr{Int64},
+                                                           bsr_values::Ptr{Cdouble})::Cint
+end
+
+function onemklCsparse_set_bsr_data(device_queue, spmat, blk_nrows, blk_ncols, blk_nnz,
+                                    row_blk_size, col_blk_size, blk_layout, index,
+                                    bsr_row_ptr, bsr_col_ind, bsr_values)
+    @ccall liboneapi_support.onemklCsparse_set_bsr_data(device_queue::syclQueue_t,
+                                                        spmat::matrix_handle_t,
+                                                        blk_nrows::Int64, blk_ncols::Int64,
+                                                        blk_nnz::Int64, row_blk_size::Int64,
+                                                        col_blk_size::Int64,
+                                                        blk_layout::onemklLayout,
+                                                        index::onemklIndex,
+                                                        bsr_row_ptr::Ptr{Int32},
+                                                        bsr_col_ind::Ptr{Int32},
+                                                        bsr_values::Ptr{ComplexF32})::Cint
+end
+
+function onemklCsparse_set_bsr_data_64(device_queue, spmat, blk_nrows, blk_ncols, blk_nnz,
+                                       row_blk_size, col_blk_size, blk_layout, index,
+                                       bsr_row_ptr, bsr_col_ind, bsr_values)
+    @ccall liboneapi_support.onemklCsparse_set_bsr_data_64(device_queue::syclQueue_t,
+                                                           spmat::matrix_handle_t,
+                                                           blk_nrows::Int64,
+                                                           blk_ncols::Int64, blk_nnz::Int64,
+                                                           row_blk_size::Int64,
+                                                           col_blk_size::Int64,
+                                                           blk_layout::onemklLayout,
+                                                           index::onemklIndex,
+                                                           bsr_row_ptr::Ptr{Int64},
+                                                           bsr_col_ind::Ptr{Int64},
+                                                           bsr_values::Ptr{ComplexF32})::Cint
+end
+
+function onemklZsparse_set_bsr_data(device_queue, spmat, blk_nrows, blk_ncols, blk_nnz,
+                                    row_blk_size, col_blk_size, blk_layout, index,
+                                    bsr_row_ptr, bsr_col_ind, bsr_values)
+    @ccall liboneapi_support.onemklZsparse_set_bsr_data(device_queue::syclQueue_t,
+                                                        spmat::matrix_handle_t,
+                                                        blk_nrows::Int64, blk_ncols::Int64,
+                                                        blk_nnz::Int64, row_blk_size::Int64,
+                                                        col_blk_size::Int64,
+                                                        blk_layout::onemklLayout,
+                                                        index::onemklIndex,
+                                                        bsr_row_ptr::Ptr{Int32},
+                                                        bsr_col_ind::Ptr{Int32},
+                                                        bsr_values::Ptr{ComplexF32})::Cint
+end
+
+function onemklZsparse_set_bsr_data_64(device_queue, spmat, blk_nrows, blk_ncols, blk_nnz,
+                                       row_blk_size, col_blk_size, blk_layout, index,
+                                       bsr_row_ptr, bsr_col_ind, bsr_values)
+    @ccall liboneapi_support.onemklZsparse_set_bsr_data_64(device_queue::syclQueue_t,
+                                                           spmat::matrix_handle_t,
+                                                           blk_nrows::Int64,
+                                                           blk_ncols::Int64, blk_nnz::Int64,
+                                                           row_blk_size::Int64,
+                                                           col_blk_size::Int64,
+                                                           blk_layout::onemklLayout,
+                                                           index::onemklIndex,
+                                                           bsr_row_ptr::Ptr{Int64},
+                                                           bsr_col_ind::Ptr{Int64},
+                                                           bsr_values::Ptr{ComplexF32})::Cint
+end
+
 function onemklXsparse_init_matmat_descr(p_desc)
     @ccall liboneapi_support.onemklXsparse_init_matmat_descr(p_desc::Ptr{matmat_descr_t})::Cint
 end
diff --git a/res/Project.toml b/res/Project.toml
index f7b4a7fe..20a2ab05 100644
--- a/res/Project.toml
+++ b/res/Project.toml
@@ -1,4 +1,4 @@
 [deps]
 Clang = "40e3b903-d033-50b4-a0cc-940c62c95e31"
 JuliaFormatter = "98e50ef6-434e-11e9-1051-2b60c6c9e899"
-oneAPI_Level_Zero_Headers_jll = "f4bc562b-d309-54f8-9efb-476e56f0410d"
+oneAPI_Level_Zero_Headers_LTS_jll = "d79c0b2e-896c-561b-aab9-323701ec0314"
diff --git a/res/support.toml b/res/support.toml
index 230ac406..ac641b90 100644
--- a/res/support.toml
+++ b/res/support.toml
@@ -370,15 +370,17 @@ use_ccall_macro = true
 6 = "ZePtr{Float32}"
 8 = "Ref{Float32}"
 
+# Argument positions account for the `nnz` parameter (arg 5) and the
+# `index::onemklIndex` enum (arg 6); the device pointers are args 7-9.
 [api.onemklXsparse_set_csr_data.argtypes]
-6 = "ZePtr{Int32}"
 7 = "ZePtr{Int32}"
-8 = "ZePtr{T}"
+8 = "ZePtr{Int32}"
+9 = "ZePtr{T}"
 
 [api.onemklXsparse_set_csr_data_64.argtypes]
-6 = "ZePtr{Int64}"
 7 = "ZePtr{Int64}"
-8 = "ZePtr{T}"
+8 = "ZePtr{Int64}"
+9 = "ZePtr{T}"
 
 [api.onemklXsparse_set_coo_data.argtypes]
 7 = "ZePtr{Int32}"
diff --git a/res/wrap.jl b/res/wrap.jl
index 1d48315e..a86428d2 100644
--- a/res/wrap.jl
+++ b/res/wrap.jl
@@ -108,7 +108,7 @@ end
 # Main application
 #
 
-using oneAPI_Level_Zero_Headers_jll
+import oneAPI_Level_Zero_Headers_LTS_jll as oneAPI_Level_Zero_Headers_jll
 
 function main()
     wrap("ze", oneAPI_Level_Zero_Headers_jll.ze_api)
diff --git a/src/compiler/compilation.jl b/src/compiler/compilation.jl
index 142b75cb..3b572c9d 100644
--- a/src/compiler/compilation.jl
+++ b/src/compiler/compilation.jl
@@ -51,9 +51,12 @@ function GPUCompiler.finish_ir!(job::oneAPICompilerJob, mod::LLVM.Module,
     flatten_nested_insertvalue!(mod)
 
     # When the device supports BFloat16 but the SPIR-V runtime doesn't accept
-    # SPV_KHR_bfloat16, lower all bfloat types to i16 so the translator can
-    # handle the module without the extension.
-    if @static(isdefined(Core, :BFloat16) && isdefined(LLVM, :BFloatType)) &&
+    # SPV_KHR_bfloat16, lower all bfloat types to i16 so the back-end/translator can
+    # handle the module without the extension. Rolling stack only: bf16 is forced off on
+    # the LTS stack (`supports_bfloat16 = false`, see `_compiler_config`), so bf16 values
+    # never reach this pass there and lowering would just be a no-op scan.
+    if !oneL0.LTS[] &&
+            @static(isdefined(Core, :BFloat16) && isdefined(LLVM, :BFloatType)) &&
             _device_supports_bfloat16() && !_driver_supports_bfloat16_spirv()
         lower_bfloat_to_i16!(mod)
     end
@@ -306,18 +309,37 @@ end
 @noinline function _compiler_config(dev; kernel=true, name=nothing, always_inline=false, kwargs...)
     supports_fp16 = oneL0.module_properties(device()).fp16flags & oneL0.ZE_DEVICE_MODULE_FLAG_FP16 == oneL0.ZE_DEVICE_MODULE_FLAG_FP16
     supports_fp64 = oneL0.module_properties(device()).fp64flags & oneL0.ZE_DEVICE_MODULE_FLAG_FP64 == oneL0.ZE_DEVICE_MODULE_FLAG_FP64
-    # Allow BFloat16 in IR if the device supports it (even if the SPIR-V runtime doesn't
-    # advertise the extension). We lower bfloat→i16 in finish_ir! when needed.
-    supports_bfloat16 = _device_supports_bfloat16()
-
-    extensions = String[]
-    # Only add the SPIR-V extension if the runtime actually supports it
-    if _driver_supports_bfloat16_spirv()
-        push!(extensions, "SPV_KHR_bfloat16")
+
+    # SPIR-V codegen path. The Aurora LTS NEO/IGC runtime only accepts SPIR-V from the
+    # Khronos translator and needs these extensions declared explicitly; the rolling stack
+    # uses the LLVM SPIR-V back-end (which handles the extensions itself). GPUCompiler picks
+    # the tool from the target's `backend` field and loads the JLL lazily, so both can be
+    # listed as deps and the choice is made here at compile time.
+    if oneL0.LTS[]
+        backend = :khronos
+        # TODO: emit printf format strings in constant memory
+        extensions = String[
+            "SPV_EXT_relaxed_printf_string_address_space",
+            "SPV_EXT_shader_atomic_float_add"
+        ]
+        # The LTS SPIR-V stack cannot codegen native bfloat in generic kernels (clamp! ->
+        # InvalidIRError, and SPV_KHR_bfloat16 segfaults NEO), so keep bf16 off entirely on
+        # this path; the test suite gates bf16 eltypes on `!oneL0.LTS[]` to match.
+        supports_bfloat16 = false
+    else
+        backend = :llvm
+        extensions = String[]
+        # Allow BFloat16 in IR if the device supports it (even if the SPIR-V runtime doesn't
+        # advertise the extension); finish_ir! lowers bfloat->i16 when the extension is absent.
+        supports_bfloat16 = _device_supports_bfloat16()
+        # Only declare the SPIR-V extension when the runtime actually accepts it.
+        if _driver_supports_bfloat16_spirv()
+            push!(extensions, "SPV_KHR_bfloat16")
+        end
     end
 
     # create GPUCompiler objects
-    target = SPIRVCompilerTarget(; extensions, supports_fp16, supports_fp64, supports_bfloat16, kwargs...)
+    target = SPIRVCompilerTarget(; backend, extensions, supports_fp16, supports_fp64, supports_bfloat16, kwargs...)
     params = oneAPICompilerParams()
     CompilerConfig(target, params; kernel, name, always_inline)
 end
diff --git a/src/context.jl b/src/context.jl
index 88035e74..17219d7a 100644
--- a/src/context.jl
+++ b/src/context.jl
@@ -224,10 +224,101 @@ function global_queue(ctx::ZeContext, dev::ZeDevice)
     # NOTE: dev purposefully does not default to context() or device() to stress that
     #       objects should track ownership, and not rely on implicit global state.
     get!(task_local_storage(), (:ZeCommandQueue, ctx, dev)) do
-        ZeCommandQueue(ctx, dev; flags = oneL0.ZE_COMMAND_QUEUE_FLAG_IN_ORDER)
+        queue = ZeCommandQueue(ctx, dev; flags = oneL0.ZE_COMMAND_QUEUE_FLAG_IN_ORDER)
+        register_queue!(ctx, dev, queue)
     end
 end
 
+# Register `queue` as a queue targeting (ctx, dev) so `synchronize_all_queues`/`release`
+# can find and drain it before freeing buffers whose in-flight work it may still reference.
+# EVERY queue that becomes a task's active queue must go through here — not just the one
+# `global_queue` creates but also the replacement `KA.priority!` installs — or the
+# unregistered queue's in-flight work can outlive a freed buffer (a use-after-free that
+# faults and bans the context on the LTS NEO stack). Only the LTS stack maintains the
+# registry; on the rolling stack this is a no-op. Returns `queue`.
+function register_queue!(ctx::ZeContext, dev::ZeDevice, queue::ZeCommandQueue)
+    oneL0.LTS[] || return queue
+    # disable finalizers while mutating the registry: a GC-driven finalizer on this
+    # task could call back into `synchronize_all_queues` (the lock is reentrant) and
+    # observe/mutate the registry mid-update.
+    GC.enable_finalizers(false)
+    try
+        @lock queue_registry_lock begin
+            push!(get!(Vector{Tuple{WeakRef,ZeCommandQueue}}, queue_registry, (ctx, dev)),
+                  (WeakRef(current_task()), queue))
+        end
+    finally
+        GC.enable_finalizers(true)
+    end
+    return queue
+end
+
+# Registry of all queues created through `global_queue`, across tasks. Buffers can be
+# freed from any task (GC finalizers), so `release` needs to be able to find the queues
+# that may still have work in flight referencing the buffer; queues themselves are
+# cached task-locally and would otherwise be unreachable from the finalizing task.
+#
+# Entries reference the queue *strongly*: the GC clears WeakRefs to a dead queue in the
+# same cycle that queues its finalizer, i.e., before the finalizer runs, so a WeakRef
+# would hide the queue from `release` exactly when its in-flight work still references
+# buffers about to be freed. The owning task is tracked weakly instead: queues are
+# task-local, so once their task is dead no new work can reach them, and the entry can
+# be dropped (allowing the queue to be finalized) after a final synchronize.
+const queue_registry_lock = ReentrantLock()
+const queue_registry = Dict{Tuple{ZeContext,ZeDevice},Vector{Tuple{WeakRef,ZeCommandQueue}}}()
+
+# synchronize all known queues that target the given context (and device, if specified),
+# i.e., all queues whose in-flight work could possibly reference an allocation that is
+# about to be freed.
+function synchronize_all_queues(ctx::ZeContext, dev::Union{ZeDevice,Nothing})
+    # only the LTS stack populates the queue registry (see `global_queue`); on the
+    # rolling stack this is a no-op and `release` frees directly.
+    oneL0.LTS[] || return
+    queues = ZeCommandQueue[]
+    stale = Tuple{WeakRef,ZeCommandQueue}[]
+    GC.enable_finalizers(false)
+    try
+        @lock queue_registry_lock begin
+            for ((qctx, qdev), entries) in queue_registry
+                qctx == ctx || continue
+                (dev === nothing || qdev == dev) || continue
+                for entry in entries
+                    (task, queue) = entry
+                    queue.handle == C_NULL && continue  # finalized, handle destroyed
+                    push!(queues, queue)
+                    # entries whose task was already dead at this point cannot
+                    # receive new work, so they are safe to retire after the sync
+                    if task.value === nothing || istaskdone(task.value::Task)
+                        push!(stale, entry)
+                    end
+                end
+            end
+        end
+        # synchronize outside the lock: this can block for as long as a kernel runs,
+        # and finalizers running concurrently also need to take the lock. Keep
+        # finalizers disabled so none of the collected queues can be destroyed
+        # between collection and synchronization.
+        for queue in queues
+            oneL0.synchronize(queue)
+        end
+        # retire drained queues of dead tasks, allowing them to be finalized (the
+        # finalizer synchronizes once more before destroying the queue, in case
+        # the queue is dropped through other means).
+        if !isempty(stale)
+            @lock queue_registry_lock begin
+                for ((qctx, qdev), entries) in queue_registry
+                    qctx == ctx || continue
+                    (dev === nothing || qdev == dev) || continue
+                    filter!(entry -> !any(s -> s === entry, stale), entries)
+                end
+            end
+        end
+    finally
+        GC.enable_finalizers(true)
+    end
+    return
+end
+
 """
     synchronize()
 
diff --git a/src/mapreduce.jl b/src/mapreduce.jl
index 645db2cd..f37e7867 100644
--- a/src/mapreduce.jl
+++ b/src/mapreduce.jl
@@ -109,14 +109,68 @@ function partial_mapreduce_device(f, op, neutral, maxitems, Rreduce, Rother, R,
     return
 end
 
+# Coalesced reduction for when the contiguous leading dimension is NOT reduced (the reduced
+# axes are strided). One work-item per output slice (Rother element), grid-strided; each
+# serially reduces over Rreduce. Consecutive work-items map to consecutive output slices,
+# which are consecutive in memory, so global reads are coalesced across lanes — the access
+# pattern a `dims=1` reduction already uses. On the Aurora LTS stack the workgroup-per-slice
+# kernel above reads a *strided* reduced dimension non-coalesced, which silently corrupts
+# large reductions (e.g. `sum(A; dims=2)`); this path avoids that pattern entirely.
+function coalesced_mapreduce_device(f, op, neutral, Rreduce, Rother, R, As...)
+    iother = (get_group_id() - 1) * get_local_size() + get_local_id()
+    gstride = get_num_groups() * get_local_size()
+    @inbounds while iother <= length(Rother)
+        Iother = Rother[iother]
+        Iout = CartesianIndex(Tuple(Iother)..., 1)
+        neut = neutral === nothing ? R[Iout] : neutral
+        val = op(neut, neut)
+        for ireduce in 1:length(Rreduce)
+            Ireduce = Rreduce[ireduce]
+            J = max(Iother, Ireduce)
+            val = op(val, f(_map_getindex(As, J)...))
+        end
+        R[Iout] = val
+        iother += gstride
+    end
+    return
+end
+
 ## COV_EXCL_STOP
 
+# Aurora LTS workaround: the NEO/IGC LTS stack miscompiles *strided* (non-coalesced) global
+# reads inside the reduction kernel, silently corrupting results whenever an input is read
+# along a non-contiguous axis (e.g. `a == transpose(b)`, `sum(transpose(x))`, `ishermitian`).
+# Elementwise copies are NOT affected. `_dense_reduce_input` returns false for any input that
+# reads non-contiguous memory (a transposed/permuted/strided view, or a broadcast containing
+# one), so such inputs get materialized to a dense `oneArray` before reducing. See
+# ISSUE_mapreduce_corruption.md and the `naive_transpose` (`a == transpose(b)`) failure.
+@inline _dense_reduce_input(::oneArray) = true
+@inline _dense_reduce_input(x::Base.ReshapedArray) = _dense_reduce_input(parent(x))
+@inline _dense_reduce_input(::AbstractArray) = false   # Transpose/Adjoint/PermutedDims/SubArray/…
+@inline _dense_reduce_input(::Any) = true              # scalars/Refs/tuples carried in a broadcast
+@inline _dense_reduce_input(bc::Broadcast.Broadcasted) = all(_dense_reduce_input, bc.args)
+
 function GPUArrays.mapreducedim!(f::F, op::OP, R::oneWrappedArray{T},
                                  A::Union{AbstractArray,Broadcast.Broadcasted};
                                  init=nothing) where {F, OP, T}
     Base.check_reducedims(R, A)
     length(A) == 0 && return R # isempty(::Broadcasted) iterates
 
+    # Aurora LTS workaround (see `_dense_reduce_input` above): materialize strided inputs to a
+    # dense array first so every global read in the reduction kernel is coalesced.
+    if oneL0.LTS[] && !_dense_reduce_input(A)
+        Acontig = Broadcast.materialize(Broadcast.broadcasted(f, A))
+        # Only recurse if materialization actually produced a dense `oneArray`. A strided
+        # *device* view (Transpose/PermutedDims of a oneArray, or a broadcast over one)
+        # materializes to a oneArray, but a plain *host* AbstractArray materializes to
+        # another host array that is still not `_dense_reduce_input` — recursing on it would
+        # loop forever (StackOverflowError). Fall through instead so the normal path raises
+        # the usual kernel-conversion error for the unsupported host input.
+        if Acontig isa oneArray
+            return GPUArrays.mapreducedim!(identity, op, R, Acontig; init=init)
+        end
+    end
+
     # add singleton dimensions to the output container, if needed
     if ndims(R) < ndims(A)
         dims = Base.fill_to_length(size(R), 1, Val(ndims(A)))
@@ -137,6 +191,40 @@ function GPUArrays.mapreducedim!(f::F, op::OP, R::oneWrappedArray{T},
     # but allows us to write a generalized kernel supporting partial reductions.
     R′ = reshape(R, (size(R)..., 1))
 
+    # Aurora LTS workaround: the workgroup-per-slice kernel below reads a *strided* reduced
+    # dimension non-coalesced, which silently corrupts reductions on this stack (regardless of
+    # output count — it depends on the reduction length, not the number of slices). Whenever
+    # the contiguous leading dimension is NOT reduced (`size(Rreduce, 1) == 1`), use the
+    # coalesced one-work-item-per-slice kernel, whose lanes read consecutive memory. Few-slice
+    # reductions get less parallelism but stay correct; the common many-slice case is also fast.
+    #
+    # `size(Rreduce, 1) == 1` (i.e. dim 1 kept) is the correct predicate, not just a special
+    # case. In `partial_mapreduce_device` adjacent work-item lanes step through adjacent
+    # `Rreduce` entries, whose *fastest* axis is array dim 1. The miscompile bites only when
+    # adjacent lanes land a *full stride* apart on every lane — i.e. when the innermost reduced
+    # axis is itself strided, which is exactly the dim-1-kept case routed to the coalesced
+    # kernel above. When dim 1 is among the reduced dims, adjacent lanes instead read
+    # consecutive memory (stride 1) within each length-n1 run, with only an occasional jump at
+    # a run boundary — not the every-lane full-stride pattern — so the reads stay well-formed
+    # even when a strided dim (e.g. dim 3) is *also* reduced. Verified with exact Int32
+    # reductions (immune to the accumulation-order rounding that makes >2^24-length Float32
+    # sums differ between GPU and CPU): `dims=(1,3)`, `(1,4)` and full `(1,2,3)` reductions at
+    # large strided sizes match the CPU exactly, so these need no extra materialization (see
+    # PR_REVIEW.md #7 and the "strided mixed reductions" regression test).
+    if oneL0.LTS[] && size(Rreduce, 1) == 1
+        # cap the group size at what this kernel actually supports on this device (queried
+        # from the driver), rather than a hardcoded 256 that can exceed the kernel's max
+        # work-group size on some devices and fail the launch.
+        cargs = (f, op, init, Rreduce, Rother, R′, A)
+        ckernel = zefunction(coalesced_mapreduce_device,
+                             Tuple{Core.Typeof.(kernel_convert.(cargs))...})
+        items = clamp(length(Rother), 1, launch_configuration(ckernel))
+        groups = min(cld(length(Rother), items), 1024)
+        @oneapi items=items groups=groups coalesced_mapreduce_device(
+            f, op, init, Rreduce, Rother, R′, A)
+        return R
+    end
+
     # how many items do we want?
     #
     # items in a group work together to reduce values across the reduction dimensions;
diff --git a/src/oneAPI.jl b/src/oneAPI.jl
index f36f8171..3098025c 100644
--- a/src/oneAPI.jl
+++ b/src/oneAPI.jl
@@ -17,7 +17,13 @@ using LLVM
 using LLVM.Interop
 using Core: LLVMPtr
 
-using SPIRV_LLVM_Backend_jll, SPIRV_Tools_jll
+import Libdl
+
+# Load both SPIR-V codegen back-ends: GPUCompiler resolves the tool from the target's
+# `backend` field (`:khronos` -> translator, `:llvm` -> LLVM back-end) via a LazyModule that
+# looks the JLL up in `Base.loaded_modules`, so both must be loaded for either path to work.
+# The LTS stack uses the translator; the rolling stack (ONEAPI_LTS=0) uses the LLVM back-end.
+using SPIRV_LLVM_Translator_jll, SPIRV_LLVM_Backend_jll, SPIRV_Tools_jll
 using oneAPI_Support_jll
 
 export oneL0
@@ -139,6 +145,26 @@ function __init__()
             if oneL0.NEO_jll.is_available()
                 # ensure that the OpenCL loader finds the ICD files from our artifacts
                 ENV["OCL_ICD_FILENAMES"] = oneL0.NEO_jll.libigdrcl
+
+                # libsycl (loaded later for oneMKL) carries its own bundled Level Zero loader
+                # that rediscovers the NEO driver by dlopen'ing it by soname rather than
+                # reusing our JLL-loaded module. dlopen NEO here by full path so it is already
+                # resident in this process before libsycl loads: a later
+                # dlopen("libze_intel_gpu.so.1") then resolves to this module by soname without
+                # needing a search path. Required when no system NEO is installed.
+                Libdl.dlopen(oneL0.NEO_jll.libze_intel_gpu; throw_error = false)
+
+                # Extend LD_LIBRARY_PATH with the NEO directory as well. NOTE: this does NOT
+                # affect the running process's own dlopen search — glibc captures
+                # LD_LIBRARY_PATH once at startup — so it serves only *child* processes (e.g.
+                # the multi-worker test suite) that inherit the environment and do their own
+                # path-based driver discovery. In-process discovery is handled by the dlopen
+                # above (and by ZE_ENABLE_ALT_DRIVERS, set in oneL0's __init__).
+                neo_libdir = dirname(oneL0.NEO_jll.libze_intel_gpu)
+                ld = get(ENV, "LD_LIBRARY_PATH", "")
+                if !occursin(neo_libdir, ld)
+                    ENV["LD_LIBRARY_PATH"] = isempty(ld) ? neo_libdir : "$neo_libdir:$ld"
+                end
             end
         end
 
@@ -150,7 +176,7 @@ function __init__()
 end
 
 function set_debug!(debug::Bool)
-    for jll in [oneL0.NEO_jll, oneL0.NEO_jll.libigc_jll]
+    for jll in [oneL0.NEO_jll, oneL0.NEO_jll.libigc_LTS_jll]
         Preferences.set_preferences!(jll, "debug" => string(debug); force=true)
     end
     @info "oneAPI debug mode $(debug ? "enabled" : "disabled"); please re-start Julia."
diff --git a/src/oneAPIKernels.jl b/src/oneAPIKernels.jl
index bc6f3218..0a4b70c8 100644
--- a/src/oneAPIKernels.jl
+++ b/src/oneAPIKernels.jl
@@ -264,6 +264,12 @@ function KA.priority!(::oneAPIBackend, prio::Symbol)
         priority = priority_enum
     )
 
+    # Register the replacement queue so `synchronize_all_queues`/`release` can drain it
+    # before freeing a buffer whose in-flight work it references; otherwise all work after
+    # `priority!` runs on an unregistered queue and a freed buffer can be reused while its
+    # kernel is still running (use-after-free → banned context on the LTS NEO stack).
+    oneAPI.register_queue!(ctx, dev, new_queue)
+
     task_local_storage((:ZeCommandQueue, ctx, dev), new_queue)
 
     return nothing
diff --git a/src/pool.jl b/src/pool.jl
index ed3b75ba..4430f351 100644
--- a/src/pool.jl
+++ b/src/pool.jl
@@ -83,6 +83,18 @@ function release(buf::oneL0.AbstractBuffer)
     #    evict(ctx, dev, buf)
     #end
 
+    # NEO (at least the 25.18 LTS release) does not honor the BLOCKING_FREE/DEFER_FREE
+    # policies of zeMemFreeExt: it advertises ZE_extension_memory_free_policies but
+    # unmaps the allocation immediately, even with work in flight that references it.
+    # That turns a GC-driven free of a dead array whose last kernel/copy hasn't retired
+    # into a GPU pagefault, which gets the kernel context banned and makes every later
+    # submission fail with ZE_RESULT_ERROR_UNKNOWN. Synchronize the queues that could
+    # reference this buffer before freeing. (No-op on the rolling stack, which honors
+    # BLOCKING_FREE.)
+    if oneL0.LTS[]
+        synchronize_all_queues(oneL0.context(buf), oneL0.device(buf))
+    end
+
     free(buf; policy=oneL0.ZE_DRIVER_MEMORY_FREE_POLICY_EXT_FLAG_BLOCKING_FREE)
 
     # TODO: queue-ordered free from non-finalizer tasks once we have
diff --git a/src/utils.jl b/src/utils.jl
index f8076a79..3c14edcb 100644
--- a/src/utils.jl
+++ b/src/utils.jl
@@ -2,8 +2,8 @@
 function versioninfo(io::IO=stdout)
     if Sys.islinux()
         println(io, "Binary dependencies:")
-        for jll in [oneL0.NEO_jll, oneL0.NEO_jll.libigc_jll, oneL0.NEO_jll.gmmlib_jll,
-                    SPIRV_LLVM_Backend_jll, SPIRV_Tools_jll, oneAPI_Support_jll]
+        for jll in [oneL0.NEO_jll, oneL0.NEO_jll.libigc_LTS_jll, oneL0.NEO_jll.gmmlib_jll,
+                    SPIRV_LLVM_Backend_jll, SPIRV_LLVM_Translator_jll, SPIRV_Tools_jll, oneAPI_Support_jll]
             name = string(jll)
             print(io, "- $(name[1:end-4]): $(Base.pkgversion(jll))")
             if jll.host_platform !== nothing
diff --git a/test/Project.toml b/test/Project.toml
index 4dd3fee0..8f585b8d 100644
--- a/test/Project.toml
+++ b/test/Project.toml
@@ -1,6 +1,7 @@
 [deps]
 AbstractFFTs = "621f4979-c628-5d54-868e-fcf4e3e8185c"
 Adapt = "79e6a3ab-5dfb-504d-930d-738a2a938a0e"
+BFloat16s = "ab4f0b2a-ad5b-11e8-123f-65d77653426b"
 Dates = "ade2ca70-3891-5945-98fb-dc099432e06a"
 Distributed = "8ba89e20-285c-5b6f-9357-94700520ee1b"
 FFTW = "7a1cc6ca-52ef-59f5-83cd-3a7055c09341"
@@ -9,18 +10,17 @@ InteractiveUtils = "b77e0a4c-d291-57a0-90e8-8db25a27a240"
 JLD2 = "033835bb-8acc-5ee8-8aae-3f567f8a3819"
 KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c"
 LinearAlgebra = "37e2e46d-f89d-539d-b4ee-838fcccc9c8e"
-NEO_jll = "700fe977-ac61-5f37-bbc8-c6c4b2b6a9fd"
+NEO_LTS_jll = "a724f90f-ce79-56dd-a1bd-b9de5a61085f"
 ParallelTestRunner = "d3525ed8-44d0-4b2c-a655-542cee43accc"
 Printf = "de0858da-6303-5e67-8744-51eddeeeb8d7"
 REPL = "3fa0cd96-eef1-5676-8a61-b3b8758bbffb"
 Random = "9a3f8284-a2c9-5f02-9a11-845980a1fd5c"
-SPIRV_LLVM_Backend_jll = "4376b9bf-cff8-51b6-bb48-39421dff0d0c"
 SparseArrays = "2f01184e-e22b-5df5-ae63-d93ebab69eaf"
 SpecialFunctions = "276daf66-3868-5448-9aa4-cd146d93841b"
 StaticArrays = "90137ffa-7385-5640-81b9-e52037218182"
 Statistics = "10745b16-79ce-11e8-11f9-7d13ad32a3b2"
 Test = "8dfed614-e22c-5e08-85e1-65c5234f0b40"
-libigc_jll = "94295238-5935-5bd7-bb0f-b00942e9bdd5"
+libigc_LTS_jll = "9a8258a1-e827-5686-bee9-144461246960"
 oneAPI = "8f75cd03-7ff8-4ecb-9b8f-daf728133b1b"
 oneAPI_Support_jll = "b049733a-a71d-5ed3-8eba-7d323ac00b36"
 
diff --git a/test/array.jl b/test/array.jl
index aadc8928..e4be18bd 100644
--- a/test/array.jl
+++ b/test/array.jl
@@ -104,3 +104,24 @@ end
   resize!(b, 1)
   @test length(b) == 1
 end
+
+@testset "strided mixed reductions" begin
+  # The Aurora LTS IGC miscompiles a reduction kernel's global reads when the *innermost*
+  # reduced axis is strided (dim 1 kept, e.g. `dims=2`); mapreducedim! routes those to a
+  # coalesced kernel. Reductions that also reduce dim 1 (e.g. `dims=(1,3)`) keep a contiguous
+  # innermost axis and stay correct on the workgroup-per-slice kernel — including with a small
+  # leading dim, where the contiguous run is short. Use Int32 (exact, associative) so the
+  # check is immune to Float32 accumulation-order rounding. Regression for PR_REVIEW.md #7.
+  for (sz, dts) in (
+      ((2, 512, 64),    ((1, 3), (2,), (3,), (2, 3), (1, 2, 3), 1)),
+      ((3, 256, 48),    ((1, 3), (2,), (1, 2, 3))),
+      ((7, 300, 20),    ((1, 3), (2,), (3,))),
+      ((2, 128, 8, 32), ((1, 3), (1, 4), (2, 4), (1, 2, 4))),
+    )
+    A = rand(Int32(1):Int32(4), sz...)
+    dA = oneArray(A)
+    for dt in dts
+      @test Array(sum(dA; dims = dt)) == sum(A; dims = dt)
+    end
+  end
+end
diff --git a/test/level-zero.jl b/test/level-zero.jl
index ed7b2838..bbf8f465 100644
--- a/test/level-zero.jl
+++ b/test/level-zero.jl
@@ -210,14 +210,21 @@ pool = ZeEventPool(ctx, 2)
 signal_event = pool[1]
 wait_event = pool[2]
 
-execute!(queue) do list
-    append_launch!(list, kernel, 1, signal_event, wait_event)
-end
-@test !Base.isdone(signal_event)
+# This is a submit-then-signal pattern: the kernel is gated on `wait_event`, which is
+# only signaled *after* submission. The ONEAPI_SYNC_EACH_SUBMISSION=1 workaround (Aurora
+# LTS) makes `execute!` block in zeCommandQueueSynchronize right after submitting, which
+# would deadlock here since the kernel cannot retire before `wait_event` is signaled. No
+# production code path submits event-gated work, so disable the workaround just here.
+oneL0.sync_each_submission(false) do
+    execute!(queue) do list
+        append_launch!(list, kernel, 1, signal_event, wait_event)
+    end
+    @test !Base.isdone(signal_event)
 
-signal(wait_event)
-synchronize(queue)
-@test Base.isdone(signal_event)
+    signal(wait_event)
+    synchronize(queue)
+    @test Base.isdone(signal_event)
+end
 
 end
 
diff --git a/test/onemkl.jl b/test/onemkl.jl
index 98774404..4a4c7fcd 100644
--- a/test/onemkl.jl
+++ b/test/onemkl.jl
@@ -14,7 +14,7 @@ k = 13
 
 @testset "Version" begin
     version_onemkl = oneMKL.version()
-    @test version_onemkl ≥ v"2026.0.0"
+    @test version_onemkl ≥ v"2025.2.0"
 end
 
 ############################################################################################
diff --git a/test/runtests.jl b/test/runtests.jl
index b0d9f555..bf656c45 100644
--- a/test/runtests.jl
+++ b/test/runtests.jl
@@ -64,6 +64,9 @@ end
 init_worker_code = quote
     $device_claim_code
     using oneAPI, Adapt
+    # BFloat16s.jl supplies the host-side numeric conversions (e.g. BFloat16(::Real))
+    # that Core.BFloat16 lacks on its own; the testsuite's CPU reference path needs them.
+    using BFloat16s
 
     import GPUArrays
     include($gpuarrays_testsuite)
@@ -83,7 +86,12 @@ init_worker_code = quote
         append!(eltypes, [Float64, ComplexF64])
     end
     @static if isdefined(Core, :BFloat16)
-        const bfloat16_supported = oneAPI._device_supports_bfloat16()
+        # _device_supports_bfloat16() is a hardware check (bf16 conversions ext / PVC device),
+        # but the Aurora LTS SPIR-V stack (Khronos translator + NEO/IGC) cannot codegen native
+        # `bfloat` in generic kernels: any kernel that keeps a bfloat value (e.g. clamp!) fails
+        # with InvalidIRError, and declaring SPV_KHR_bfloat16 crashes the LTS runtime. So only
+        # exercise BFloat16 as a generic element type off the LTS stack.
+        const bfloat16_supported = oneAPI._device_supports_bfloat16() && !oneL0.LTS[]
         if bfloat16_supported
             push!(eltypes, Core.BFloat16)
         end