From 5440241829b8bea99348b922f568be240035a908 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Tue, 23 Jun 2026 15:45:58 +0000 Subject: [PATCH 01/26] Build against and run on the Intel GPU LTS stack Aurora ships the Intel "LTS" GPU software stack rather than the rolling release that upstream oneAPI.jl targets. Pin the whole toolchain to it: - JLLs: NEO_LTS_jll =25.18.33578, oneAPI_Level_Zero_Loader_LTS_jll =1.24, oneAPI_Level_Zero_Headers_LTS_jll, libigc_LTS_jll, and the oneMKL 2025.3.1 support library. lib/level-zero/oneL0.jl aliases the *_LTS_jll packages back to their canonical names so the rest of the code is unchanged. - Keep the SPIR-V Translator (SPIRV_LLVM_Translator_jll 21, SPIRVIntrinsics 0.5). The LTS NEO/IGC runtime does not accept the output of upstream's LLVM SPIR-V back-end (#491), so revert the back-end switch in src/compiler/compilation.jl, src/oneAPI.jl and src/utils.jl and restore the explicit SPIR-V extension list. - Regenerate the oneMKL C bindings against 2025.3.1 (deps/src/onemkl.*, lib/support/liboneapi_support.jl, deps/generate_interfaces.jl) and fix the sparse CSR argument types (ZePtr/onemklIndex) in the regenerated bindings and res/support.toml. Stay on oneMKL 2025.3.1, declining the upstream oneAPI 2026.0.0 migration (#582). - Prepend NEO's libze_intel_gpu directory to LD_LIBRARY_PATH in __init__ so libsycl's bundled ze_lib finds the driver by path when no system NEO is installed. Co-Authored-By: Claude Opus 4.8 (1M context) --- Project.toml | 22 +-- deps/Project.toml | 6 +- deps/build_local.jl | 5 +- deps/generate_interfaces.jl | 26 ++- deps/src/onemkl.cpp | 200 +++++++++++++++---- deps/src/onemkl.h | 154 +++++++++++---- deps/src/onemkl_dft.cpp | 33 +--- lib/level-zero/oneL0.jl | 6 +- lib/mkl/wrappers_sparse.jl | 4 +- lib/support/liboneapi_support.jl | 316 ++++++++++++++++++++++++++----- res/Project.toml | 2 +- res/support.toml | 10 +- res/wrap.jl | 2 +- src/compiler/compilation.jl | 137 +------------- src/oneAPI.jl | 13 +- src/utils.jl | 4 +- test/Project.toml | 5 +- test/onemkl.jl | 2 +- 18 files changed, 628 insertions(+), 319 deletions(-) diff --git a/Project.toml b/Project.toml index 1775b329..c997a054 100644 --- a/Project.toml +++ b/Project.toml @@ -1,7 +1,7 @@ name = "oneAPI" uuid = "8f75cd03-7ff8-4ecb-9b8f-daf728133b1b" authors = ["Tim Besard ", "Alexis Montoison", "Michel Schanen "] -version = "2.6.2" +version = "2.6.1" [deps] AbstractFFTs = "621f4979-c628-5d54-868e-fcf4e3e8185c" @@ -16,18 +16,18 @@ 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] @@ -38,19 +38,19 @@ CEnum = "0.4, 0.5" ExprTools = "0.1" GPUArrays = "11.2.1" GPUCompiler = "1.6" -GPUToolbox = "0.1, 0.2, 0.3, 1, 3" +GPUToolbox = "0.1, 0.2, 0.3, 1" 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" +SPIRVIntrinsics = "0.5" +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" 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*>(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*>(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*>(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*>(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*>(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*>(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*>(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*>(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*>(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*>(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*>(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*>(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*>(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*>(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*>(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*>(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*>(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*>(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*>(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*>(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*>(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*>(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*>(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*>(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* >(desc->ptr)->get_value(p, &tmp); - else - static_cast< descriptor* >(desc->ptr)->get_value(p, &tmp); - *value = static_cast(tmp); - } else { - double tmp; - if (desc->dom == domain::REAL) - static_cast< descriptor* >(desc->ptr)->get_value(p, &tmp); - else - static_cast< descriptor* >(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/lib/level-zero/oneL0.jl b/lib/level-zero/oneL0.jl index 900f80de..14c43cc5 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") diff --git a/lib/mkl/wrappers_sparse.jl b/lib/mkl/wrappers_sparse.jl index 8e58956b..92cd0d43 100644 --- a/lib/mkl/wrappers_sparse.jl +++ b/lib/mkl/wrappers_sparse.jl @@ -60,7 +60,7 @@ 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) + $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 +81,7 @@ 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ᵀ + $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 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..4ca04b0d 100644 --- a/src/compiler/compilation.jl +++ b/src/compiler/compilation.jl @@ -50,14 +50,6 @@ function GPUCompiler.finish_ir!(job::oneAPICompilerJob, mod::LLVM.Module, # indices (e.g., "1 0") corrupts adjacent struct fields. 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)) && - _device_supports_bfloat16() && !_driver_supports_bfloat16_spirv() - lower_bfloat_to_i16!(mod) - end - return entry end @@ -166,106 +158,6 @@ function flatten_insert!(inst::LLVM.Instruction) end -# Lower bfloat types to i16 in the LLVM IR. -# This is needed when the device supports BFloat16 but the SPIR-V runtime/translator -# doesn't support SPV_KHR_bfloat16. Since sizeof(bfloat)==sizeof(i16)==2, the memory -# layout is identical. -# -# TODO: Julia 1.12's Core.BFloat16 is a bare primitive (no Float32 conversion, no -# arithmetic), so fptrunc/fpext instructions never appear in practice. If Julia adds -# BFloat16 conversion methods in the future, this pass should be extended to handle -# fptrunc float→bfloat and fpext bfloat→float, either via inline RNE bit manipulation -# or calls to __devicelib_ConvertFToBF16INTEL / __devicelib_ConvertBF16ToFINTEL. -function lower_bfloat_to_i16!(mod::LLVM.Module) - T_bf16 = LLVM.BFloatType() - T_i16 = LLVM.Int16Type() - - # Phase 1: Eliminate all bitcasts between i16 and bfloat (same bit width). - eliminate_bf16_bitcasts!(mod, T_bf16, T_i16) - - # Phase 2: Replace remaining bfloat GEPs, loads, and stores with i16 equivalents. - for f in functions(mod) - isempty(blocks(f)) && continue - for bb in blocks(f) - to_replace = LLVM.Instruction[] - for inst in instructions(bb) - opcode = LLVM.API.LLVMGetInstructionOpcode(inst) - if opcode == LLVM.API.LLVMGetElementPtr - src_ty = LLVMType(LLVM.API.LLVMGetGEPSourceElementType(inst)) - src_ty == T_bf16 && push!(to_replace, inst) - elseif opcode == LLVM.API.LLVMLoad - value_type(inst) == T_bf16 && push!(to_replace, inst) - elseif opcode == LLVM.API.LLVMStore - value_type(LLVM.operands(inst)[1]) == T_bf16 && push!(to_replace, inst) - end - end - - for inst in to_replace - opcode = LLVM.API.LLVMGetInstructionOpcode(inst) - builder = LLVM.IRBuilder() - LLVM.position!(builder, inst) - - if opcode == LLVM.API.LLVMGetElementPtr - ptr = LLVM.operands(inst)[1] - indices = LLVM.Value[LLVM.operands(inst)[i] for i in 2:length(LLVM.operands(inst))] - new_gep = if LLVM.API.LLVMIsInBounds(inst) != 0 - LLVM.inbounds_gep!(builder, T_i16, ptr, indices) - else - LLVM.gep!(builder, T_i16, ptr, indices) - end - LLVM.replace_uses!(inst, new_gep) - elseif opcode == LLVM.API.LLVMLoad - ptr = LLVM.operands(inst)[1] - new_load = LLVM.load!(builder, T_i16, ptr) - LLVM.replace_uses!(inst, new_load) - elseif opcode == LLVM.API.LLVMStore - val = LLVM.operands(inst)[1] - ptr = LLVM.operands(inst)[2] - LLVM.store!(builder, val, ptr) - end - - LLVM.API.LLVMInstructionEraseFromParent(inst) - LLVM.dispose(builder) - end - end - end - - return true -end - -# Iteratively eliminate bitcasts between i16 and bfloat (same bit representation). -function eliminate_bf16_bitcasts!(mod::LLVM.Module, T_bf16::LLVMType, T_i16::LLVMType) - changed = true - while changed - changed = false - for f in functions(mod) - isempty(blocks(f)) && continue - for bb in blocks(f) - to_delete = LLVM.Instruction[] - for inst in instructions(bb) - if LLVM.API.LLVMGetInstructionOpcode(inst) == LLVM.API.LLVMBitCast - src = LLVM.operands(inst)[1] - src_ty = value_type(src) - dst_ty = value_type(inst) - if (src_ty == T_i16 && dst_ty == T_bf16) || - (src_ty == T_bf16 && dst_ty == T_i16) || - (src_ty == dst_ty) - LLVM.replace_uses!(inst, src) - push!(to_delete, inst) - changed = true - end - end - end - for inst in to_delete - LLVM.API.LLVMInstructionEraseFromParent(inst) - end - end - end - end - return -end - - ## compiler implementation (cache, configure, compile, and link) # cache of compilation caches, per device @@ -291,33 +183,18 @@ function compiler_config(dev; kwargs...) end return config end -# Whether the driver's SPIR-V runtime accepts the SPV_KHR_bfloat16 extension. -function _driver_supports_bfloat16_spirv() - return @static if isdefined(Core, :BFloat16) - haskey( - oneL0.extension_properties(driver()), - oneL0.ZE_BFLOAT16_CONVERSIONS_EXT_NAME - ) - else - false - end -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") - end + + # TODO: emit printf format strings in constant memory + extensions = String[ + "SPV_EXT_relaxed_printf_string_address_space", + "SPV_EXT_shader_atomic_float_add" + ] # create GPUCompiler objects - target = SPIRVCompilerTarget(; extensions, supports_fp16, supports_fp64, supports_bfloat16, kwargs...) + target = SPIRVCompilerTarget(; extensions, supports_fp16, supports_fp64, kwargs...) params = oneAPICompilerParams() CompilerConfig(target, params; kernel, name, always_inline) end diff --git a/src/oneAPI.jl b/src/oneAPI.jl index f36f8171..45ef5f40 100644 --- a/src/oneAPI.jl +++ b/src/oneAPI.jl @@ -17,7 +17,7 @@ using LLVM using LLVM.Interop using Core: LLVMPtr -using SPIRV_LLVM_Backend_jll, SPIRV_Tools_jll +using SPIRV_LLVM_Translator_jll, SPIRV_Tools_jll using oneAPI_Support_jll export oneL0 @@ -139,6 +139,15 @@ 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 + + # ensure that libsycl's bundled ze_lib finds NEO's libze_intel_gpu via + # path-based driver discovery (it does not reuse the JLL-loaded module). + # Required when no system NEO is installed. + 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 +159,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/utils.jl b/src/utils.jl index f8076a79..e7d232ec 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_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..fa5557e9 100644 --- a/test/Project.toml +++ b/test/Project.toml @@ -9,18 +9,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/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 ############################################################################################ From 9199940cf6a9bb80b98a348b664396232822eb78 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Tue, 23 Jun 2026 15:45:59 +0000 Subject: [PATCH 02/26] Work around strided mapreducedim! miscompiles on the LTS IGC The LTS IGC silently miscompiles strided memory accesses inside the reduction kernels, producing wrong results with no error. Add two complementary guards in src/mapreduce.jl: - A coalesced reduction kernel for the case where the contiguous leading dimension is reduced (size(Rreduce, 1) == 1, e.g. sum(A; dims=2)), avoiding the strided per-thread loads. - Materialize strided / non-dense inputs (Transpose, Adjoint, PermutedDimsArray, SubArray, ...) to a dense array before reducing, via _dense_reduce_input; this fixes e.g. `a == transpose(b)`. Co-Authored-By: Claude Opus 4.8 (1M context) --- src/mapreduce.jl | 60 ++++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 60 insertions(+) diff --git a/src/mapreduce.jl b/src/mapreduce.jl index 645db2cd..f44905fc 100644 --- a/src/mapreduce.jl +++ b/src/mapreduce.jl @@ -109,14 +109,60 @@ 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 !_dense_reduce_input(A) + Acontig = Broadcast.materialize(Broadcast.broadcasted(f, A)) + return GPUArrays.mapreducedim!(identity, op, R, Acontig; init=init) + 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 +183,20 @@ 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. + if size(Rreduce, 1) == 1 + items = clamp(length(Rother), 1, 256) + 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; From 6aa93e094b073e65aed7507947becdeebc6ee996 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Tue, 23 Jun 2026 15:46:00 +0000 Subject: [PATCH 03/26] Harden command-queue lifetime to prevent banned-context faults on free On the LTS NEO stack (25.18) freeing a buffer does not drain queues that still have work in flight referencing it: the in-flight kernel then faults and the context is banned, surfacing later as a ZE_RESULT_ERROR_UNKNOWN at an unrelated op. global_queue is task-local, so a test file's task can also die with work still queued. - Register every global queue in a per-(context,device) registry that holds the queue strongly, keyed by a weak reference to the owning task (src/context.jl). A WeakRef to the queue would be cleared in the same GC cycle that queues its finalizer, hiding it from release exactly when its in-flight work still references buffers being freed. - Before any BLOCKING_FREE, synchronize all queues that could reference the buffer (src/pool.jl, synchronize_all_queues). - Synchronize outside the registry lock with finalizers disabled, skip already-finalized queues, and retire queues of dead tasks. - In the queue finalizer, drain (unchecked, since a banned context returns an error) then destroy, and null the handle so a concurrent synchronize_all_queues skips it (lib/level-zero/cmdqueue.jl). Co-Authored-By: Claude Opus 4.8 (1M context) --- lib/level-zero/cmdqueue.jl | 10 +++++ src/context.jl | 78 +++++++++++++++++++++++++++++++++++++- src/pool.jl | 9 +++++ 3 files changed, 96 insertions(+), 1 deletion(-) diff --git a/lib/level-zero/cmdqueue.jl b/lib/level-zero/cmdqueue.jl index 3004a145..bdefeb44 100644 --- a/lib/level-zero/cmdqueue.jl +++ b/lib/level-zero/cmdqueue.jl @@ -20,7 +20,17 @@ mutable struct ZeCommandQueue zeCommandQueueCreate(ctx, dev, desc_ref, handle_ref) obj = new(handle_ref[], ctx, dev, ordinal) finalizer(obj) do obj + # 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. + unchecked_zeCommandQueueSynchronize(obj, typemax(UInt64)) zeCommandQueueDestroy(obj) + # 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 obj end diff --git a/src/context.jl b/src/context.jl index 88035e74..9150b367 100644 --- a/src/context.jl +++ b/src/context.jl @@ -224,8 +224,84 @@ 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) + # 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 + queue + end +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}) + 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 """ diff --git a/src/pool.jl b/src/pool.jl index ed3b75ba..ad7e2bb0 100644 --- a/src/pool.jl +++ b/src/pool.jl @@ -83,6 +83,15 @@ 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. + synchronize_all_queues(oneL0.context(buf), oneL0.device(buf)) + free(buf; policy=oneL0.ZE_DRIVER_MEMORY_FREE_POLICY_EXT_FLAG_BLOCKING_FREE) # TODO: queue-ordered free from non-finalizer tasks once we have From 2adadbc4fe74b4460138bdba393203a5d42600aa Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Tue, 23 Jun 2026 15:46:00 +0000 Subject: [PATCH 04/26] Root oneMKL FFT descriptor arrays across the C call The DFT wrappers passed pointer(lengths) / pointer(strides) to the descriptor create/set calls. A raw Ptr does not keep the backing vector alive, so the GC could collect it mid-call and oneMKL would read garbage dimensions/strides, failing commit with FFT_INVALID_DESCRIPTOR or a SIGFPE depending on heap reuse. Pass the arrays themselves so ccall roots them for the duration of the call (lib/mkl/fft.jl). Co-Authored-By: Claude Opus 4.8 (1M context) --- lib/mkl/fft.jl | 47 +++++++++++++++++++---------------------------- 1 file changed, 19 insertions(+), 28 deletions(-) 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) From 5108b77abd33ab1c3efed4fa6dc124ef3a6cfbbf Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Tue, 23 Jun 2026 15:46:00 +0000 Subject: [PATCH 05/26] Add an opt-in per-submission synchronize for the LTS dropped-tail bug The Aurora LTS NEO stack intermittently drops the tail of a command list, silently corrupting results. Add an off-by-default workaround (ONEAPI_SYNC_EACH_SUBMISSION=1) that synchronizes the queue after every command-list submission (lib/level-zero/cmdlist.jl, parsed in lib/level-zero/oneL0.jl __init__). Co-Authored-By: Claude Opus 4.8 (1M context) --- lib/level-zero/cmdlist.jl | 17 +++++++++++++++-- lib/level-zero/oneL0.jl | 1 + 2 files changed, 16 insertions(+), 2 deletions(-) diff --git a/lib/level-zero/cmdlist.jl b/lib/level-zero/cmdlist.jl index 133d2551..24d1d3a5 100644 --- a/lib/level-zero/cmdlist.jl +++ b/lib/level-zero/cmdlist.jl @@ -47,8 +47,21 @@ 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) + +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/oneL0.jl b/lib/level-zero/oneL0.jl index 14c43cc5..e8624ec7 100644 --- a/lib/level-zero/oneL0.jl +++ b/lib/level-zero/oneL0.jl @@ -201,6 +201,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[] = lowercase(get(ENV, "ONEAPI_SYNC_EACH_SUBMISSION", "")) in ("1", "true", "yes") end end From 4fd200faaa65e8215f412e57d53aabd21b3143c7 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Tue, 23 Jun 2026 15:46:00 +0000 Subject: [PATCH 06/26] Configure the Aurora LTS self-hosted CI runner Set up the GitHub Actions job for the Aurora LTS self-runner (.github/workflows/ci.yml): - Enable per-worker GPU spreading (ONEAPI_TEST_SPREAD_GPUS=1) and the per-submission synchronize workaround (ONEAPI_SYNC_EACH_SUBMISSION=1). - Run the test step through `julia -C native,-avx512fp16` to avoid Sapphire Rapids AVX512-FP16 host miscompilation of Float16 work under concurrent oneMKL load (the GPU results are correct; the corruption is on the host CPU). Co-Authored-By: Claude Opus 4.8 (1M context) --- .github/workflows/ci.yml | 16 +++++++++++++++- 1 file changed, 15 insertions(+), 1 deletion(-) diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 164c2d9d..c8932682 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -18,6 +18,12 @@ 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' + # 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 +44,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()' From 2dbe13fa10e8df5bc837eaffb2113de031023936 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Tue, 23 Jun 2026 16:35:58 +0000 Subject: [PATCH 07/26] Gate the LTS workarounds behind a runtime ONEAPI_LTS flag (layers 1-2) Make the Intel LTS-stack support opt-in so this branch can eventually merge into main without affecting the rolling stack: with the flag off, every LTS code path is bypassed and behavior matches upstream. Introduce oneL0.LTS[], resolved at the top of oneL0.__init__ from ONEAPI_LTS (default on for this branch; set ONEAPI_LTS=0 for the rolling-stack paths). Resolved before the driver-availability early returns so it is set even on a host without a functional GPU. Gate on it: - Layer 1 (behavior): the strided-reduction materialization and the coalesced-kernel dispatch in mapreducedim!; the command-queue registry (global_queue / synchronize_all_queues), the pre-free synchronize in release, and the queue-finalizer drain + handle-null in cmdqueue.jl. - Layer 2 (codegen): _compiler_config now selects the SPIRVCompilerTarget backend (:khronos translator vs :llvm back-end) and the SPIR-V extension list from the flag. GPUCompiler loads the tool lazily, so both SPIR-V JLLs can coexist as deps and the choice is made at compile time. Deferred to layer 3: driver-JLL selection (NEO/loader/headers/igc) via Preferences and the oneMKL support-library ABI. The oneMKL FFT GC-rooting is left as-is (already fixed upstream via GC.@preserve). Verified: precompiles and loads; ONEAPI_LTS unset/1 -> on, ONEAPI_LTS=0 -> off. Co-Authored-By: Claude Opus 4.8 (1M context) --- lib/level-zero/cmdqueue.jl | 24 ++++++++++++++---------- lib/level-zero/oneL0.jl | 16 ++++++++++++++++ src/compiler/compilation.jl | 23 +++++++++++++++++------ src/context.jl | 25 +++++++++++++++---------- src/mapreduce.jl | 4 ++-- src/pool.jl | 7 +++++-- 6 files changed, 69 insertions(+), 30 deletions(-) diff --git a/lib/level-zero/cmdqueue.jl b/lib/level-zero/cmdqueue.jl index bdefeb44..bb3374e5 100644 --- a/lib/level-zero/cmdqueue.jl +++ b/lib/level-zero/cmdqueue.jl @@ -20,17 +20,21 @@ mutable struct ZeCommandQueue zeCommandQueueCreate(ctx, dev, desc_ref, handle_ref) obj = new(handle_ref[], ctx, dev, ordinal) finalizer(obj) do obj - # 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. - unchecked_zeCommandQueueSynchronize(obj, typemax(UInt64)) + 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. + unchecked_zeCommandQueueSynchronize(obj, typemax(UInt64)) + end zeCommandQueueDestroy(obj) - # 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) + 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 e8624ec7..87a5e830 100644 --- a/lib/level-zero/oneL0.jl +++ b/lib/level-zero/oneL0.jl @@ -139,10 +139,26 @@ 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 exactly like the upstream rolling +# stack. Default on for this branch (which pins the LTS toolchain); set `ONEAPI_LTS=0` +# to exercise the rolling-stack paths. When merged upstream the default flips to false +# and an LTS deployment opts in with `ONEAPI_LTS=1`. +const LTS = Ref{Bool}(true) + 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 on for this (LTS) branch; flip the "true" + # default to "false" when merged onto the rolling stack. ONEAPI_LTS=0 forces the + # rolling-stack code paths. + LTS[] = lowercase(get(ENV, "ONEAPI_LTS", "true")) in ("1", "true", "yes", "on") + 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." diff --git a/src/compiler/compilation.jl b/src/compiler/compilation.jl index 4ca04b0d..adc17d62 100644 --- a/src/compiler/compilation.jl +++ b/src/compiler/compilation.jl @@ -187,14 +187,25 @@ end 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 - # TODO: emit printf format strings in constant memory - extensions = String[ - "SPV_EXT_relaxed_printf_string_address_space", - "SPV_EXT_shader_atomic_float_add" - ] + # 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" + ] + else + backend = :llvm + extensions = String[] + end # create GPUCompiler objects - target = SPIRVCompilerTarget(; extensions, supports_fp16, supports_fp64, kwargs...) + target = SPIRVCompilerTarget(; backend, extensions, supports_fp16, supports_fp64, kwargs...) params = oneAPICompilerParams() CompilerConfig(target, params; kernel, name, always_inline) end diff --git a/src/context.jl b/src/context.jl index 9150b367..bc5ac70b 100644 --- a/src/context.jl +++ b/src/context.jl @@ -225,17 +225,19 @@ function global_queue(ctx::ZeContext, dev::ZeDevice) # objects should track ownership, and not rely on implicit global state. get!(task_local_storage(), (:ZeCommandQueue, ctx, dev)) do queue = ZeCommandQueue(ctx, dev; flags = oneL0.ZE_COMMAND_QUEUE_FLAG_IN_ORDER) - # 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)) + if oneL0.LTS[] + # 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 - finally - GC.enable_finalizers(true) end queue end @@ -259,6 +261,9 @@ const queue_registry = Dict{Tuple{ZeContext,ZeDevice},Vector{Tuple{WeakRef,ZeCom # 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) diff --git a/src/mapreduce.jl b/src/mapreduce.jl index f44905fc..d03d72fd 100644 --- a/src/mapreduce.jl +++ b/src/mapreduce.jl @@ -158,7 +158,7 @@ function GPUArrays.mapreducedim!(f::F, op::OP, R::oneWrappedArray{T}, # 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 !_dense_reduce_input(A) + if oneL0.LTS[] && !_dense_reduce_input(A) Acontig = Broadcast.materialize(Broadcast.broadcasted(f, A)) return GPUArrays.mapreducedim!(identity, op, R, Acontig; init=init) end @@ -189,7 +189,7 @@ function GPUArrays.mapreducedim!(f::F, op::OP, R::oneWrappedArray{T}, # 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. - if size(Rreduce, 1) == 1 + if oneL0.LTS[] && size(Rreduce, 1) == 1 items = clamp(length(Rother), 1, 256) groups = min(cld(length(Rother), items), 1024) @oneapi items=items groups=groups coalesced_mapreduce_device( diff --git a/src/pool.jl b/src/pool.jl index ad7e2bb0..4430f351 100644 --- a/src/pool.jl +++ b/src/pool.jl @@ -89,8 +89,11 @@ function release(buf::oneL0.AbstractBuffer) # 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. - synchronize_all_queues(oneL0.context(buf), oneL0.device(buf)) + # 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) From d0d3b096be1f17b7f875e86c5c7778b5365e0d4f Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 1 Jul 2026 14:48:27 +0000 Subject: [PATCH 08/26] Compat: SPIRVIntrsinics v1 --- .github/workflows/ci.yml | 1 + Project.toml | 2 +- 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index c8932682..a36e601f 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -21,6 +21,7 @@ jobs: # 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' diff --git a/Project.toml b/Project.toml index c997a054..9ca0c3e2 100644 --- a/Project.toml +++ b/Project.toml @@ -43,7 +43,7 @@ KernelAbstractions = "0.9.39" LLVM = "6, 7, 8, 9" NEO_LTS_jll = "=25.18.33578" Preferences = "1" -SPIRVIntrinsics = "0.5" +SPIRVIntrinsics = "1" SPIRV_LLVM_Translator_jll = "21" SPIRV_Tools_jll = "2025.4.0" SpecialFunctions = "1.3, 2" From 25d5236d2b2a289f445dd2640c12196609f2ba4c Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Thu, 2 Jul 2026 15:29:33 +0000 Subject: [PATCH 09/26] Use SPIRVIntrinsics fork with fixes --- Project.toml | 3 +++ 1 file changed, 3 insertions(+) diff --git a/Project.toml b/Project.toml index 9ca0c3e2..16851806 100644 --- a/Project.toml +++ b/Project.toml @@ -54,3 +54,6 @@ oneAPI_Support_jll = "0.9.2" [extras] libigc_LTS_jll = "9a8258a1-e827-5686-bee9-144461246960" + +[sources] +SPIRVIntrinsics = {url = "https://github.com/michel2323/OpenCL.jl", rev = "atomic-float-ops", subdir = "lib/intrinsics"} From f158ea6d31a7bfb2c66894f5327a866af5bd55f8 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Mon, 6 Jul 2026 15:05:17 +0000 Subject: [PATCH 10/26] test: fix BFloat16 math test on the Aurora LTS stack The gpuarrays/math/intrinsics BFloat16 subtest failed two ways: 1. Host-side: Core.BFloat16 exists on Julia 1.12 but has no Real->bf16 constructors (those live in BFloat16s.jl), so ET(-1) threw a MethodError before any GPU work. Add BFloat16s as a test dep and load it in the worker so the CPU reference path works. 2. Codegen: _device_supports_bfloat16() is a hardware check, but the Aurora LTS SPIR-V stack (Khronos translator + NEO/IGC) cannot codegen native bfloat in generic kernels -- clamp! fails with InvalidIRError, and declaring SPV_KHR_bfloat16 crashes the LTS runtime. Gate the bf16 eltypes push on !oneL0.LTS[] so BFloat16 is only exercised as a generic element type off the LTS stack (storage/conversions/oneMKL bf16 are unaffected). --- test/Project.toml | 1 + test/runtests.jl | 10 +++++++++- 2 files changed, 10 insertions(+), 1 deletion(-) diff --git a/test/Project.toml b/test/Project.toml index fa5557e9..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" 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 From 299b6a411e2d93b7193d2ae1fb153f8934f99532 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Tue, 7 Jul 2026 13:48:02 +0000 Subject: [PATCH 11/26] examples: skip bfloat16.jl on the Aurora LTS stack The examples testset runs bfloat16.jl as a subprocess, and it unconditionally compiles the scale_bf16 kernel. The example uses Core.BFloat16 as storage only (all math via Float32 reinterpret round-trips), but even a bfloat load/store forces the bfloat LLVM type into the module, which the LTS Khronos translator cannot emit without SPV_KHR_bfloat16 (rejected by the LTS runtime). So it fails with "Failed to translate LLVM code to SPIR-V" regardless of device support. Add an oneL0.LTS[] skip-guard alongside the existing old-Julia and unsupported-device guards, so the example exits cleanly on the LTS stack. Claude-Session: https://claude.ai/code/session_01QM51M7h61UVk7uKU5o5sCZ --- examples/bfloat16.jl | 9 +++++++++ 1 file changed, 9 insertions(+) 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) From 0ddd27a2fe1deed087690ed26b4b5a4e56ec721d Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Tue, 7 Jul 2026 14:38:01 +0000 Subject: [PATCH 12/26] level-zero: fix sync-each-submission deadlock in kernel-execution test Wrap the SYNC_EACH_SUBMISSION Ref in sync_each_submission() getter, sync_each_submission!(enable) setter, and a scoped sync_each_submission(f, enable) do-block form. Use the scoped form to disable the workaround around the submit-then-signal block in test/level-zero.jl. That block launches a kernel gated on a wait_event that is only signaled after submission; under ONEAPI_SYNC_EACH_SUBMISSION=1, execute! synchronizes on submission and deadlocks waiting for a kernel that cannot retire yet. No production path submits event-gated work, so the workaround stays safe everywhere it matters. --- lib/level-zero/cmdlist.jl | 41 +++++++++++++++++++++++++++++++++++++-- lib/level-zero/oneL0.jl | 2 +- test/level-zero.jl | 21 +++++++++++++------- 3 files changed, 54 insertions(+), 10 deletions(-) diff --git a/lib/level-zero/cmdlist.jl b/lib/level-zero/cmdlist.jl index 24d1d3a5..1df4bb90 100644 --- a/lib/level-zero/cmdlist.jl +++ b/lib/level-zero/cmdlist.jl @@ -55,11 +55,48 @@ end # 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) +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) + sync_each_submission() && synchronize(queue) return r end diff --git a/lib/level-zero/oneL0.jl b/lib/level-zero/oneL0.jl index 87a5e830..7f549f29 100644 --- a/lib/level-zero/oneL0.jl +++ b/lib/level-zero/oneL0.jl @@ -217,7 +217,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[] = lowercase(get(ENV, "ONEAPI_SYNC_EACH_SUBMISSION", "")) in ("1", "true", "yes") + sync_each_submission!(lowercase(get(ENV, "ONEAPI_SYNC_EACH_SUBMISSION", "")) in ("1", "true", "yes")) 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 From 698c2438b1d922006b392f20979960ef7dfbee0a Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 14:51:01 +0000 Subject: [PATCH 13/26] Project.toml: restore version 2.6.2 and GPUToolbox 3 compat Rebase artifacts regressed the package version 2.6.2 -> 2.6.1 (which would make the registry reject a tag and Pkg treat the branch as a downgrade) and dropped GPUToolbox 3 from the compat entry (GPUToolbox 3.0.0 is registered, so environments holding v3 could no longer resolve this package). Restore both to match main. Addresses PR_REVIEW.md finding #9. --- Project.toml | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/Project.toml b/Project.toml index 16851806..8b05a619 100644 --- a/Project.toml +++ b/Project.toml @@ -1,7 +1,7 @@ name = "oneAPI" uuid = "8f75cd03-7ff8-4ecb-9b8f-daf728133b1b" authors = ["Tim Besard ", "Alexis Montoison", "Michel Schanen "] -version = "2.6.1" +version = "2.6.2" [deps] AbstractFFTs = "621f4979-c628-5d54-868e-fcf4e3e8185c" @@ -38,7 +38,7 @@ CEnum = "0.4, 0.5" ExprTools = "0.1" GPUArrays = "11.2.1" GPUCompiler = "1.6" -GPUToolbox = "0.1, 0.2, 0.3, 1" +GPUToolbox = "0.1, 0.2, 0.3, 1, 3" KernelAbstractions = "0.9.39" LLVM = "6, 7, 8, 9" NEO_LTS_jll = "=25.18.33578" From e2d3a5343c721e63fe2b5039c12f87616dd35b53 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 14:52:12 +0000 Subject: [PATCH 14/26] level-zero: share one env-bool parser across ONEAPI_* flags ONEAPI_SYNC_EACH_SUBMISSION accepted only ("1","true","yes") while the sibling ONEAPI_LTS also accepted "on". A deployment setting ONEAPI_SYNC_EACH_SUBMISSION=on therefore ran with the dropped-tail workaround silently disabled. Route both flags through a shared parse_env_bool helper that accepts on/off/yes/no/1/0/true/false and warns on an unrecognized value instead of silently defaulting to off. Addresses PR_REVIEW.md finding #8. --- lib/level-zero/oneL0.jl | 19 +++++++++++++++++-- 1 file changed, 17 insertions(+), 2 deletions(-) diff --git a/lib/level-zero/oneL0.jl b/lib/level-zero/oneL0.jl index 7f549f29..a9469658 100644 --- a/lib/level-zero/oneL0.jl +++ b/lib/level-zero/oneL0.jl @@ -148,6 +148,21 @@ const parameter_validation = Ref{Bool}() # and an LTS deployment opts in with `ONEAPI_LTS=1`. const LTS = Ref{Bool}(true) +# 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 @@ -157,7 +172,7 @@ function __init__() # without a functional GPU. Default on for this (LTS) branch; flip the "true" # default to "false" when merged onto the rolling stack. ONEAPI_LTS=0 forces the # rolling-stack code paths. - LTS[] = lowercase(get(ENV, "ONEAPI_LTS", "true")) in ("1", "true", "yes", "on") + LTS[] = parse_env_bool("ONEAPI_LTS", true) if Sys.iswindows() if Libdl.dlopen(libze_loader; throw_error=false) === nothing @@ -217,7 +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!(lowercase(get(ENV, "ONEAPI_SYNC_EACH_SUBMISSION", "")) in ("1", "true", "yes")) + sync_each_submission!(parse_env_bool("ONEAPI_SYNC_EACH_SUBMISSION", false)) end end From 9b11a48a03af4c8daeacba1b016e9d8f2c6df712 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 14:53:24 +0000 Subject: [PATCH 15/26] Register KA.priority! replacement queue in the queue registry The queue registry that synchronize_all_queues/release rely on to drain in-flight work before freeing a buffer was only populated by global_queue. KA.priority! replaced the task-local queue with a directly constructed, unregistered ZeCommandQueue, so after priority!(backend, :high) all work ran on a queue release() never synchronizes: freeing a dead oneArray whose last kernel had not retired reused the buffer under the running kernel, the exact use-after-free (GPU pagefault -> banned context) the registry exists to prevent on the LTS NEO stack. Extract the registry insertion into a shared register_queue! helper and call it from both global_queue and priority!. Addresses PR_REVIEW.md finding #4. --- src/context.jl | 38 ++++++++++++++++++++++++-------------- src/oneAPIKernels.jl | 6 ++++++ 2 files changed, 30 insertions(+), 14 deletions(-) diff --git a/src/context.jl b/src/context.jl index bc5ac70b..17219d7a 100644 --- a/src/context.jl +++ b/src/context.jl @@ -225,22 +225,32 @@ function global_queue(ctx::ZeContext, dev::ZeDevice) # objects should track ownership, and not rely on implicit global state. get!(task_local_storage(), (:ZeCommandQueue, ctx, dev)) do queue = ZeCommandQueue(ctx, dev; flags = oneL0.ZE_COMMAND_QUEUE_FLAG_IN_ORDER) - if oneL0.LTS[] - # 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 + 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 - queue + finally + GC.enable_finalizers(true) end + return queue end # Registry of all queues created through `global_queue`, across tasks. Buffers can be 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 From bd3b19ae295810680e19dc41e5bc343146eacc94 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 14:54:14 +0000 Subject: [PATCH 16/26] mapreduce: guard LTS strided-materialization fallback against host inputs _dense_reduce_input(::AbstractArray) is false for a plain CPU array, so the LTS materialization fallback materialized it (to another CPU array) and re-called GPUArrays.mapreducedim!, which failed the same predicate and recursed forever -> StackOverflowError, where main failed fast with a kernel-conversion error. Only recurse when materialization actually produced a dense oneArray; otherwise fall through to the normal path so the usual conversion error surfaces for an unsupported host input. Addresses PR_REVIEW.md finding #5. --- src/mapreduce.jl | 10 +++++++++- 1 file changed, 9 insertions(+), 1 deletion(-) diff --git a/src/mapreduce.jl b/src/mapreduce.jl index d03d72fd..73fe2a79 100644 --- a/src/mapreduce.jl +++ b/src/mapreduce.jl @@ -160,7 +160,15 @@ function GPUArrays.mapreducedim!(f::F, op::OP, R::oneWrappedArray{T}, # 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)) - return GPUArrays.mapreducedim!(identity, op, R, Acontig; init=init) + # 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 From d3cf2da1f6a3bc3d8750e7bf8bbdaf40e3de42d7 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 14:56:44 +0000 Subject: [PATCH 17/26] level-zero: bound the LTS queue-finalizer drain instead of waiting forever MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The LTS finalizer drained the queue with unchecked_zeCommandQueueSynchronize( obj, typemax(UInt64)) before destroy, to avoid the destroy-with-work-in-flight fault that bans the context on the LTS NEO stack. But a task that submits event-gated work and dies before signaling the event leaves the queue permanently unfinishable, so that infinite wait hangs the finalizer — and with it GC and process exit — forever. Wait with a bounded 10 s timeout instead; on timeout, warn and leak the queue deliberately rather than destroy it (which would trigger the very fault+ban the drain exists to prevent). Addresses PR_REVIEW.md finding #6. --- lib/level-zero/cmdqueue.jl | 17 ++++++++++++++++- 1 file changed, 16 insertions(+), 1 deletion(-) diff --git a/lib/level-zero/cmdqueue.jl b/lib/level-zero/cmdqueue.jl index bb3374e5..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 @@ -26,7 +32,16 @@ mutable struct ZeCommandQueue # 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. - unchecked_zeCommandQueueSynchronize(obj, typemax(UInt64)) + # + # 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[] From f1b6d450c8cefbd05716f34a111527224229849c Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 15:04:42 +0000 Subject: [PATCH 18/26] mapreduce: document that mixed dims=(1,3) reductions need no workaround MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Finding #7 conjectured that a reduction including dim 1 plus a strided dim (e.g. sum(A; dims=(1,3)) with a small leading dim) would still hit the LTS IGC strided-load miscompile, since size(Rreduce,1) != 1 routes it to partial_mapreduce_device. Empirically it does not: consecutive work-item lanes step through consecutive Rreduce entries whose first axis is dim 1, so when dim 1 is reduced the innermost reduced axis is contiguous and those reads coalesce — correct even when a strided dim is also reduced. A sweep of dims=(1,3)-style reductions with small leading dims (n1 in 2..7) matched the CPU exactly. The size(Rreduce,1)==1 guard is therefore the correct predicate; record the reasoning and verification rather than add needless materialization. Addresses PR_REVIEW.md finding #7. --- src/mapreduce.jl | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/src/mapreduce.jl b/src/mapreduce.jl index 73fe2a79..06cc5447 100644 --- a/src/mapreduce.jl +++ b/src/mapreduce.jl @@ -197,6 +197,15 @@ function GPUArrays.mapreducedim!(f::F, op::OP, R::oneWrappedArray{T}, # 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` consecutive work-item lanes step through consecutive + # `Rreduce` entries, whose *first* axis is dim 1. When dim 1 is among the reduced dims the + # innermost reduced axis is contiguous, so those lane reads coalesce and the result is + # correct even if a strided dim (e.g. dim 3) is *also* reduced — the miscompile only bites + # when the innermost reduced axis is itself strided (dim 1 kept). Empirically verified: a + # sweep of `dims=(1,3)`-style mixed reductions with small leading dims (n1 ∈ 2..7) matched + # the CPU exactly, so no extra materialization is needed for those (see PR_REVIEW.md #7). if oneL0.LTS[] && size(Rreduce, 1) == 1 items = clamp(length(Rother), 1, 256) groups = min(cld(length(Rother), items), 1024) From 13dd61cc67a06a21110c4928ca806716a973c2fa Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 15:06:39 +0000 Subject: [PATCH 19/26] Restore SPIRV_LLVM_Backend_jll so the rolling stack can compile kernels The branch removed SPIRV_LLVM_Backend_jll from deps and replaced using SPIRV_LLVM_Backend_jll with using SPIRV_LLVM_Translator_jll, but the non-LTS path (ONEAPI_LTS=0) still selects backend=:llvm. GPUCompiler resolves that back-end through a LazyModule that checks Base.loaded_modules, so with the JLL neither a dep nor loaded, the first @oneapi launch on the rolling stack failed in mcgen and no kernel could compile at all. Re-add SPIRV_LLVM_Backend_jll to deps/compat and load it alongside the translator: GPUCompiler picks the tool from the target's backend field, so both coexist (LTS -> :khronos translator, rolling -> :llvm back-end). Verified that oneAPI still precompiles, both JLLs load, and kernels compile/run on the LTS stack (Backend JLL loaded but unused there). Addresses PR_REVIEW.md finding #2. --- Project.toml | 2 ++ src/oneAPI.jl | 6 +++++- 2 files changed, 7 insertions(+), 1 deletion(-) diff --git a/Project.toml b/Project.toml index 8b05a619..1878538f 100644 --- a/Project.toml +++ b/Project.toml @@ -21,6 +21,7 @@ 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" @@ -44,6 +45,7 @@ LLVM = "6, 7, 8, 9" 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" diff --git a/src/oneAPI.jl b/src/oneAPI.jl index 45ef5f40..0d83872f 100644 --- a/src/oneAPI.jl +++ b/src/oneAPI.jl @@ -17,7 +17,11 @@ using LLVM using LLVM.Interop using Core: LLVMPtr -using SPIRV_LLVM_Translator_jll, SPIRV_Tools_jll +# 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 From 30e5300af8dbe5f83356a215655455b714082bc8 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 15:10:34 +0000 Subject: [PATCH 20/26] Restore BFloat16 kernel support on the rolling stack The refactor dropped supports_bfloat16 from the SPIRVCompilerTarget call on both branches of the oneL0.LTS[] check, so GPUCompiler defaulted it to false and validate_ir rejected any bfloat value with InvalidIRError even on a bf16-capable rolling-stack device (e.g. PVC) where main compiled and ran it. Restore the main behavior on the non-LTS branch: supports_bfloat16 = _device_supports_bfloat16(), declare SPV_KHR_bfloat16 when the runtime advertises it, and otherwise lower bfloat->i16 in finish_ir! (restored _driver_supports_bfloat16_spirv / lower_bfloat_to_i16! / eliminate_bf16_bitcasts!). Keep bf16 forced off on the LTS branch, where the SPIR-V stack cannot codegen native bfloat (the finish_ir! lowering is gated !oneL0.LTS[] so it stays inert there). Verified on the LTS stack: supports_bfloat16=false, the restored helpers run without error, kernels still compile and run. Addresses PR_REVIEW.md finding #3. --- src/compiler/compilation.jl | 136 +++++++++++++++++++++++++++++++++++- 1 file changed, 135 insertions(+), 1 deletion(-) diff --git a/src/compiler/compilation.jl b/src/compiler/compilation.jl index adc17d62..3b572c9d 100644 --- a/src/compiler/compilation.jl +++ b/src/compiler/compilation.jl @@ -50,6 +50,17 @@ function GPUCompiler.finish_ir!(job::oneAPICompilerJob, mod::LLVM.Module, # indices (e.g., "1 0") corrupts adjacent struct fields. 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 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 + return entry end @@ -158,6 +169,106 @@ function flatten_insert!(inst::LLVM.Instruction) end +# Lower bfloat types to i16 in the LLVM IR. +# This is needed when the device supports BFloat16 but the SPIR-V runtime/translator +# doesn't support SPV_KHR_bfloat16. Since sizeof(bfloat)==sizeof(i16)==2, the memory +# layout is identical. +# +# TODO: Julia 1.12's Core.BFloat16 is a bare primitive (no Float32 conversion, no +# arithmetic), so fptrunc/fpext instructions never appear in practice. If Julia adds +# BFloat16 conversion methods in the future, this pass should be extended to handle +# fptrunc float→bfloat and fpext bfloat→float, either via inline RNE bit manipulation +# or calls to __devicelib_ConvertFToBF16INTEL / __devicelib_ConvertBF16ToFINTEL. +function lower_bfloat_to_i16!(mod::LLVM.Module) + T_bf16 = LLVM.BFloatType() + T_i16 = LLVM.Int16Type() + + # Phase 1: Eliminate all bitcasts between i16 and bfloat (same bit width). + eliminate_bf16_bitcasts!(mod, T_bf16, T_i16) + + # Phase 2: Replace remaining bfloat GEPs, loads, and stores with i16 equivalents. + for f in functions(mod) + isempty(blocks(f)) && continue + for bb in blocks(f) + to_replace = LLVM.Instruction[] + for inst in instructions(bb) + opcode = LLVM.API.LLVMGetInstructionOpcode(inst) + if opcode == LLVM.API.LLVMGetElementPtr + src_ty = LLVMType(LLVM.API.LLVMGetGEPSourceElementType(inst)) + src_ty == T_bf16 && push!(to_replace, inst) + elseif opcode == LLVM.API.LLVMLoad + value_type(inst) == T_bf16 && push!(to_replace, inst) + elseif opcode == LLVM.API.LLVMStore + value_type(LLVM.operands(inst)[1]) == T_bf16 && push!(to_replace, inst) + end + end + + for inst in to_replace + opcode = LLVM.API.LLVMGetInstructionOpcode(inst) + builder = LLVM.IRBuilder() + LLVM.position!(builder, inst) + + if opcode == LLVM.API.LLVMGetElementPtr + ptr = LLVM.operands(inst)[1] + indices = LLVM.Value[LLVM.operands(inst)[i] for i in 2:length(LLVM.operands(inst))] + new_gep = if LLVM.API.LLVMIsInBounds(inst) != 0 + LLVM.inbounds_gep!(builder, T_i16, ptr, indices) + else + LLVM.gep!(builder, T_i16, ptr, indices) + end + LLVM.replace_uses!(inst, new_gep) + elseif opcode == LLVM.API.LLVMLoad + ptr = LLVM.operands(inst)[1] + new_load = LLVM.load!(builder, T_i16, ptr) + LLVM.replace_uses!(inst, new_load) + elseif opcode == LLVM.API.LLVMStore + val = LLVM.operands(inst)[1] + ptr = LLVM.operands(inst)[2] + LLVM.store!(builder, val, ptr) + end + + LLVM.API.LLVMInstructionEraseFromParent(inst) + LLVM.dispose(builder) + end + end + end + + return true +end + +# Iteratively eliminate bitcasts between i16 and bfloat (same bit representation). +function eliminate_bf16_bitcasts!(mod::LLVM.Module, T_bf16::LLVMType, T_i16::LLVMType) + changed = true + while changed + changed = false + for f in functions(mod) + isempty(blocks(f)) && continue + for bb in blocks(f) + to_delete = LLVM.Instruction[] + for inst in instructions(bb) + if LLVM.API.LLVMGetInstructionOpcode(inst) == LLVM.API.LLVMBitCast + src = LLVM.operands(inst)[1] + src_ty = value_type(src) + dst_ty = value_type(inst) + if (src_ty == T_i16 && dst_ty == T_bf16) || + (src_ty == T_bf16 && dst_ty == T_i16) || + (src_ty == dst_ty) + LLVM.replace_uses!(inst, src) + push!(to_delete, inst) + changed = true + end + end + end + for inst in to_delete + LLVM.API.LLVMInstructionEraseFromParent(inst) + end + end + end + end + return +end + + ## compiler implementation (cache, configure, compile, and link) # cache of compilation caches, per device @@ -183,6 +294,18 @@ function compiler_config(dev; kwargs...) end return config end +# Whether the driver's SPIR-V runtime accepts the SPV_KHR_bfloat16 extension. +function _driver_supports_bfloat16_spirv() + return @static if isdefined(Core, :BFloat16) + haskey( + oneL0.extension_properties(driver()), + oneL0.ZE_BFLOAT16_CONVERSIONS_EXT_NAME + ) + else + false + end +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 @@ -199,13 +322,24 @@ end "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(; backend, extensions, supports_fp16, supports_fp64, kwargs...) + target = SPIRVCompilerTarget(; backend, extensions, supports_fp16, supports_fp64, supports_bfloat16, kwargs...) params = oneAPICompilerParams() CompilerConfig(target, params; kernel, name, always_inline) end From 4374f681887893244a713a9c173de4a26d8bfb5d Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 15:16:05 +0000 Subject: [PATCH 21/26] mkl: guard sparse setters against a stale oneAPI_Support_jll ABI MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The regenerated sparse wrappers widened the ABI: 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; the registry's oneAPI_Support_jll 0.9.2 binary predates them, so a stock install (or CI, which does not run deps/build_local.jl) calls a sparse setter against the old 8-arg C function — nnz lands in an index slot, every pointer shifts, and MKL dereferences garbage (segfault / silent corruption). Probe the loaded support binary once at load for onemklSsparse_set_csc_data (a symbol only the rebuilt binary exports) and, when it is absent, make the sparse CSR/CSC/COO constructors raise an actionable error pointing at deps/build_local.jl instead of crashing. Verified: probe returns true against a locally rebuilt binary, construction works, and the guard throws the expected message when the flag is cleared. Addresses PR_REVIEW.md finding #1. --- lib/mkl/wrappers_sparse.jl | 3 +++ lib/support/Support.jl | 32 +++++++++++++++++++++++++++++++- 2 files changed, 34 insertions(+), 1 deletion(-) diff --git a/lib/mkl/wrappers_sparse.jl b/lib/mkl/wrappers_sparse.jl index 92cd0d43..3c112027 100644 --- a/lib/mkl/wrappers_sparse.jl +++ b/lib/mkl/wrappers_sparse.jl @@ -60,6 +60,7 @@ 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 + 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) @@ -81,6 +82,7 @@ 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 + 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) @@ -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 From a784431f7a2fb291f0c37b7e270cee68fe49cc2a Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 15:17:44 +0000 Subject: [PATCH 22/26] oneAPI: make in-process NEO discovery real, fix LD_LIBRARY_PATH comment MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Setting ENV["LD_LIBRARY_PATH"] inside __init__ cannot affect the running process's dlopen search paths: glibc captures LD_LIBRARY_PATH once at process startup, so the stated purpose (letting libsycl's bundled ze_lib find NEO's libze_intel_gpu in-process) was never served for the current process — only for child processes that inherit the environment. dlopen the NEO driver by full path in __init__ so it is resident before libsycl loads; a later dlopen("libze_intel_gpu.so.1") by libsycl's bundled loader then resolves to it by soname without a path search. Keep extending LD_LIBRARY_PATH but correct the comment to say it only covers spawned worker processes. Verified oneMKL (which loads libsycl and needs NEO discovery) still works. Addresses PR_REVIEW.md finding #10. --- src/oneAPI.jl | 19 ++++++++++++++++--- 1 file changed, 16 insertions(+), 3 deletions(-) diff --git a/src/oneAPI.jl b/src/oneAPI.jl index 0d83872f..3098025c 100644 --- a/src/oneAPI.jl +++ b/src/oneAPI.jl @@ -17,6 +17,8 @@ using LLVM using LLVM.Interop using Core: LLVMPtr +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. @@ -144,9 +146,20 @@ function __init__() # ensure that the OpenCL loader finds the ICD files from our artifacts ENV["OCL_ICD_FILENAMES"] = oneL0.NEO_jll.libigdrcl - # ensure that libsycl's bundled ze_lib finds NEO's libze_intel_gpu via - # path-based driver discovery (it does not reuse the JLL-loaded module). - # Required when no system NEO is installed. + # 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) From e2ddd99302a963e2937d5197771acf0bbac7c6da Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 15:19:11 +0000 Subject: [PATCH 23/26] Project.toml: pin SPIRVIntrinsics [sources] to a commit SHA rev = "atomic-float-ops" tracked a mutable branch head, so builds were not reproducible and a force-push to the fork would silently change what users resolve. Pin to commit 4257d40, whose lib/intrinsics subtree matches the currently resolved Manifest exactly (git-tree-sha1 87509a31), so this is a no-op today but freezes the input going forward. Addresses PR_REVIEW.md additional observation (mutable [sources] branch pin). --- Project.toml | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/Project.toml b/Project.toml index 1878538f..c6731ca0 100644 --- a/Project.toml +++ b/Project.toml @@ -58,4 +58,6 @@ oneAPI_Support_jll = "0.9.2" libigc_LTS_jll = "9a8258a1-e827-5686-bee9-144461246960" [sources] -SPIRVIntrinsics = {url = "https://github.com/michel2323/OpenCL.jl", rev = "atomic-float-ops", subdir = "lib/intrinsics"} +# 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"} From df6c1a41ca6cae247a0d8675106c7d56db00613f Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 15:20:57 +0000 Subject: [PATCH 24/26] mapreduce: size the coalesced LTS kernel via launch_configuration The coalesced-reduction launch used a hardcoded items = clamp(length(Rother), 1, 256), which can exceed the kernel's max work-group size on some devices and fail the launch. Query launch_configuration(kernel) for the actual limit, like the main partial_mapreduce_device path already does. Verified the coalesced path (dim-1-kept reductions) stays correct across a range of sizes, including large Rother. Addresses PR_REVIEW.md additional observation (hardcoded coalesced items). --- src/mapreduce.jl | 8 +++++++- 1 file changed, 7 insertions(+), 1 deletion(-) diff --git a/src/mapreduce.jl b/src/mapreduce.jl index 06cc5447..be36c082 100644 --- a/src/mapreduce.jl +++ b/src/mapreduce.jl @@ -207,7 +207,13 @@ function GPUArrays.mapreducedim!(f::F, op::OP, R::oneWrappedArray{T}, # sweep of `dims=(1,3)`-style mixed reductions with small leading dims (n1 ∈ 2..7) matched # the CPU exactly, so no extra materialization is needed for those (see PR_REVIEW.md #7). if oneL0.LTS[] && size(Rreduce, 1) == 1 - items = clamp(length(Rother), 1, 256) + # 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) From ac130edee272d99024094f39558b8f8f57ec986c Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 15:33:42 +0000 Subject: [PATCH 25/26] mapreduce: add exact-arithmetic regression test for strided mixed reductions Follow-up to the #7 resolution, addressing the adversarial-review caveats: the prior comment overstated coalescing for a small leading dim, and the empirical sweep left nothing committed and used Float32 (which can pass or fail for the wrong reason on large sums). Tighten the comment to the precise mechanism (the miscompile needs adjacent lanes a *full* stride apart on every lane; reducing dim 1 keeps a contiguous innermost axis, so lanes read stride-1 within each length-n1 run) and add an Int32 regression test covering dims=(1,3)/(1,4)/(1,2,3) plus the coalesced dims=2/3 cases. Int32 is exact and associative, so it is immune to the Float32 accumulation-order rounding that made a stress sweep of >2^24-length sums look corrupted when it was not; verified all cases match the CPU exactly, including small leading dims and large strided sizes. Refines PR_REVIEW.md finding #7. --- src/mapreduce.jl | 19 ++++++++++++------- test/array.jl | 21 +++++++++++++++++++++ 2 files changed, 33 insertions(+), 7 deletions(-) diff --git a/src/mapreduce.jl b/src/mapreduce.jl index be36c082..f37e7867 100644 --- a/src/mapreduce.jl +++ b/src/mapreduce.jl @@ -199,13 +199,18 @@ function GPUArrays.mapreducedim!(f::F, op::OP, R::oneWrappedArray{T}, # 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` consecutive work-item lanes step through consecutive - # `Rreduce` entries, whose *first* axis is dim 1. When dim 1 is among the reduced dims the - # innermost reduced axis is contiguous, so those lane reads coalesce and the result is - # correct even if a strided dim (e.g. dim 3) is *also* reduced — the miscompile only bites - # when the innermost reduced axis is itself strided (dim 1 kept). Empirically verified: a - # sweep of `dims=(1,3)`-style mixed reductions with small leading dims (n1 ∈ 2..7) matched - # the CPU exactly, so no extra materialization is needed for those (see PR_REVIEW.md #7). + # 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 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 From b32eeb20941ef48afd0d73fea63d618fda1c3f4b Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 16:39:32 +0000 Subject: [PATCH 26/26] Set default ONEAPI_LTS=0 --- lib/level-zero/oneL0.jl | 18 +++++++++--------- src/utils.jl | 2 +- 2 files changed, 10 insertions(+), 10 deletions(-) diff --git a/lib/level-zero/oneL0.jl b/lib/level-zero/oneL0.jl index a9469658..69636412 100644 --- a/lib/level-zero/oneL0.jl +++ b/lib/level-zero/oneL0.jl @@ -142,11 +142,12 @@ 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 exactly like the upstream rolling -# stack. Default on for this branch (which pins the LTS toolchain); set `ONEAPI_LTS=0` -# to exercise the rolling-stack paths. When merged upstream the default flips to false -# and an LTS deployment opts in with `ONEAPI_LTS=1`. -const LTS = Ref{Bool}(true) +# 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 @@ -169,10 +170,9 @@ function __init__() # 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 on for this (LTS) branch; flip the "true" - # default to "false" when merged onto the rolling stack. ONEAPI_LTS=0 forces the - # rolling-stack code paths. - LTS[] = parse_env_bool("ONEAPI_LTS", true) + # 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 diff --git a/src/utils.jl b/src/utils.jl index e7d232ec..3c14edcb 100644 --- a/src/utils.jl +++ b/src/utils.jl @@ -3,7 +3,7 @@ function versioninfo(io::IO=stdout) if Sys.islinux() println(io, "Binary dependencies:") for jll in [oneL0.NEO_jll, oneL0.NEO_jll.libigc_LTS_jll, oneL0.NEO_jll.gmmlib_jll, - SPIRV_LLVM_Translator_jll, SPIRV_Tools_jll, oneAPI_Support_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