From ef0827e891e367e7d573091b31ea455385e6f58f Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Tue, 21 Jul 2026 09:57:26 -0500 Subject: [PATCH 01/29] Add ONEAPI_MEMORY_LIMIT to budget device memory across processes MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The proactive-GC heuristic in the memory pool computes its pressure thresholds against the device's total memory, but _allocated_bytes only tracks the current process, and GC in one process cannot free buffers held by another. When several processes share one device — the parallel test suite being the prime example — each process happily grows to 80% of the card before feeling any pressure, so their combined footprint exhausts physical memory and allocations start failing with OutOfGPUMemoryError (hundreds of such errors on an 8GB Arc A750 with --jobs=2). Introduce ONEAPI_MEMORY_LIMIT (a byte count, or a percentage of device memory like "50%") as a per-process soft budget: _maybe_gc now scales its 40%/80% thresholds by the limit instead of the full card. This is deliberately soft — oversized single allocations still go through, with the existing retry_reclaim ladder as the reactive backstop. --- docs/src/memory.md | 18 ++++++++++++++++++ src/pool.jl | 29 ++++++++++++++++++++++++++++- 2 files changed, 46 insertions(+), 1 deletion(-) diff --git a/docs/src/memory.md b/docs/src/memory.md index f97146e8..0e4ea1e9 100644 --- a/docs/src/memory.md +++ b/docs/src/memory.md @@ -49,3 +49,21 @@ oneAPI.unsafe_free!(a) **Warning**: Only use `unsafe_free!` if you are sure the array is no longer used, including by any pending GPU operations. +## Memory Limit + +When multiple processes share one GPU, each process only tracks its own allocations, +and garbage collection in one process cannot free memory held by another. To keep the +processes from collectively exhausting the device, give each a budget with the +`ONEAPI_MEMORY_LIMIT` environment variable, either as a byte count or as a percentage +of device memory: + +``` +ONEAPI_MEMORY_LIMIT=50% # this process budgets half the device +ONEAPI_MEMORY_LIMIT=4000000000 # ... or an absolute byte count +``` + +This is a soft limit: it scales the thresholds at which oneAPI.jl proactively runs the +garbage collector to free stale GPU buffers, before falling back to a full collection +when an allocation actually fails. Allocations larger than the limit are still allowed +to succeed. The test suite sets this automatically for its parallel workers. + diff --git a/src/pool.jl b/src/pool.jl index ed3b75ba..fcc6a3a4 100644 --- a/src/pool.jl +++ b/src/pool.jl @@ -13,10 +13,37 @@ function _get_total_mem(dev) return _total_mem_cache[] end +# Per-process memory budget for the proactive-GC thresholds below. `_allocated_bytes` +# only tracks this process, and GC in this process cannot free buffers held by another, +# so when several processes share one device (e.g. parallel test workers) each must +# budget against its share of the card, not the whole card — otherwise their combined +# usage exceeds physical memory before any of them feels pressure. ONEAPI_MEMORY_LIMIT +# accepts a byte count ("4000000000") or a percentage of device memory ("50%"); unset +# means the full device. +const _memory_limit_cache = Threads.Atomic{Int64}(-1) + +function _memory_limit(dev) + cached = _memory_limit_cache[] + cached >= 0 && return cached + total = _get_total_mem(dev) + str = get(ENV, "ONEAPI_MEMORY_LIMIT", "") + limit = if isempty(str) + total + elseif endswith(str, "%") + pct = parse(Int, chop(str)) + 0 < pct <= 100 || error("ONEAPI_MEMORY_LIMIT percentage must be in (0, 100]") + total * pct ÷ 100 + else + parse(Int64, str) + end + Threads.atomic_cas!(_memory_limit_cache, Int64(-1), Int64(limit)) + return _memory_limit_cache[] +end + function _maybe_gc(dev, bytes) allocated = _allocated_bytes[] allocated <= 0 && return - total_mem = _get_total_mem(dev) + total_mem = _memory_limit(dev) return if allocated + bytes > total_mem * 0.8 # Flush deferred resource releases (e.g., MKL sparse handles) from previous GC # cycles first — these are safe to release now because they were deferred earlier. From cb9ee324689fcde3c9b6d74313c1bda2d30a79b5 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Tue, 21 Jul 2026 09:57:26 -0500 Subject: [PATCH 02/29] Harden the test suite for memory-constrained GPUs MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Three changes, motivated by a full-suite run on an 8GB Arc A750 with --jobs=2 failing 520 tests (366 of them OutOfGPUMemoryError, cascading across ~20 test files): - give each test worker an equal ONEAPI_MEMORY_LIMIT share of device memory, so the pool's proactive GC starts collecting before the workers collectively exhaust the card; - pass recycle_on_failure/retries to ParallelTestRunner (once it supports them, >= 2.7): a failing test file can leave its worker — or the driver — in a state where every subsequent allocation fails, so workers are recycled on failure and failed files get one retry, sequentially, on a fresh worker with an otherwise-idle device; - load BFloat16s.jl on the workers (Core.BFloat16 alone lacks the host-side conversions and rand sampling the testsuite's CPU reference path needs; this alone accounted for ~150 errors), and only include BFloat16 in the generic test element types when the host-side support is actually complete — probed at runtime, so the eltype enables itself once BFloat16s.jl implements the missing div/rem family. Device-side BFloat16 (the bfloat16.jl example) is exercised regardless. --- test/Project.toml | 1 + test/runtests.jl | 40 ++++++++++++++++++++++++++++++++++++++-- 2 files changed, 39 insertions(+), 2 deletions(-) diff --git a/test/Project.toml b/test/Project.toml index 4dd3fee0..788581e1 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..37bf165d 100644 --- a/test/runtests.jl +++ b/test/runtests.jl @@ -35,6 +35,13 @@ args = parse_args(ARGS) # contention/oversubscription bugs. const spread_gpus = lowercase(get(ENV, "ONEAPI_TEST_SPREAD_GPUS", "")) in ("1", "true", "yes") worker_env = Vector{Pair{String, String}}() + +# Workers share one device but cannot see each other's allocations, so give each an equal +# share of device memory: the pool's proactive GC then kicks in before the workers +# collectively exhaust the card (important on 8GB consumer GPUs). +njobs = something(args.jobs, Some(ParallelTestRunner.default_njobs()))[] +push!(worker_env, "ONEAPI_MEMORY_LIMIT" => "$(max(1, 100 ÷ njobs))%") + device_claim_code = :() if spread_gpus ndev = length(oneAPI.devices()) @@ -64,6 +71,10 @@ end init_worker_code = quote $device_claim_code using oneAPI, Adapt + # BFloat16s.jl supplies the host-side numeric support (conversions, rand sampling) + # that Core.BFloat16 lacks on its own; the testsuite's CPU reference path needs it + # when BFloat16 is part of the tested element types. + using BFloat16s import GPUArrays include($gpuarrays_testsuite) @@ -83,7 +94,22 @@ init_worker_code = quote append!(eltypes, [Float64, ComplexF64]) end @static if isdefined(Core, :BFloat16) - const bfloat16_supported = oneAPI._device_supports_bfloat16() + # Exercising BFloat16 as a generic element type needs the testsuite's CPU + # reference path to work, which requires more host-side BFloat16 support than + # BFloat16s.jl currently implements (e.g. the div/rem family throws "rem not + # defined for BFloat16"). Probe the operations the testsuite relies on, so the + # element type enables itself automatically once BFloat16s.jl catches up. + # Device-side BFloat16 (oneMKL gemm/gemv, the bfloat16.jl example) is tested + # regardless. + bfloat16_host_support() = try + x, y = BFloat16(3), BFloat16(2) + mod(x, y); fld(x, y); rand(BFloat16) + true + catch + false + end + const bfloat16_supported = oneAPI._device_supports_bfloat16() && + bfloat16_host_support() if bfloat16_supported push!(eltypes, Core.BFloat16) end @@ -151,4 +177,14 @@ init_code = quote ..@grab_output, ..@on_device, ..sink end -runtests(oneAPI, args; testsuite, init_code, init_worker_code, env = worker_env) +# On memory-pressured GPUs (e.g. 8GB cards) a failing test can leave the worker — or even +# the driver — in a state where every subsequent allocation fails, so recycle workers on +# failure and give failed files one exclusive retry on an otherwise-idle device. +failure_handling = if pkgversion(ParallelTestRunner) >= v"2.7.0" + (; recycle_on_failure = true, retries = 1) +else + (;) +end + +runtests(oneAPI, args; testsuite, init_code, init_worker_code, env = worker_env, + failure_handling...) From fe9e95c240fcf5041343e11f8ba1c3c2f9179d2f Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Tue, 21 Jul 2026 10:57:58 -0500 Subject: [PATCH 03/29] Make floating-point atomics work on the LLVM back-end path Two pieces, without which oneAPI.atomic_add!/atomic_sub! on Float32 fail on discrete GPUs (4 errors in device/intrinsics on an Arc A750): - use the SPIRVIntrinsics atomic-float-ops revision, whose atomic intrinsics lower properly instead of falling back to an unsupported runtime call ("unsupported call to an unknown function gpu_malloc"); - declare SPV_EXT_shader_atomic_float_add (and the relaxed-printf extension) for the LLVM SPIR-V back-end: extensions must be declared to permit the corresponding instructions during translation, and without them the atomic float instructions fail to translate. --- Project.toml | 5 +++++ src/compiler/compilation.jl | 9 ++++++++- 2 files changed, 13 insertions(+), 1 deletion(-) diff --git a/Project.toml b/Project.toml index 29523d7f..04d16b7b 100644 --- a/Project.toml +++ b/Project.toml @@ -56,3 +56,8 @@ oneAPI_Support_jll = "0.9.2" [extras] libigc_jll = "94295238-5935-5bd7-bb0f-b00942e9bdd5" + +[sources] +# Pin to a commit SHA, not the mutable `atomic-float-ops` branch head, so builds are +# reproducible and a force-push upstream cannot silently change what users resolve. +SPIRVIntrinsics = {url = "https://github.com/michel2323/OpenCL.jl", rev = "4257d4075162276d9985168923788490e4071e9e", subdir = "lib/intrinsics"} diff --git a/src/compiler/compilation.jl b/src/compiler/compilation.jl index 5a398a65..928c1de9 100644 --- a/src/compiler/compilation.jl +++ b/src/compiler/compilation.jl @@ -328,7 +328,14 @@ end # advertise the extension). We lower bfloat→i16 in finish_ir! when needed. supports_bfloat16 = _device_supports_bfloat16(dev) - extensions = String[] + # SPIR-V extensions the LLVM back-end may emit. Declaring them permits the + # corresponding instructions during translation: without + # SPV_EXT_shader_atomic_float_add, floating-point atomic operations fail to + # translate ("The atomic float instruction requires ... SPV_EXT_shader_atomic_float_add"). + extensions = String[ + "SPV_EXT_relaxed_printf_string_address_space", + "SPV_EXT_shader_atomic_float_add", + ] # Only add the SPIR-V extension if the runtime actually supports it if _driver_supports_bfloat16_spirv(dev) push!(extensions, "SPV_KHR_bfloat16") From a38644f1f46e29595f2941c72da86e6d877f7b75 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Tue, 21 Jul 2026 10:57:58 -0500 Subject: [PATCH 04/29] onemkl tests: accept oneMKL 2025.2+ --- test/onemkl.jl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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 4438561fa2ae96a7dce4f7dc387671531d42e6b3 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Tue, 23 Jun 2026 15:45:58 +0000 Subject: [PATCH 05/29] 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 | 18 +- 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 | 145 +------------- src/compiler/precompile.jl | 9 +- src/oneAPI.jl | 13 +- src/utils.jl | 4 +- test/Project.toml | 5 +- 18 files changed, 630 insertions(+), 328 deletions(-) diff --git a/Project.toml b/Project.toml index 04d16b7b..9256ddfd 100644 --- a/Project.toml +++ b/Project.toml @@ -16,19 +16,19 @@ KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c" LLVM = "929cbde3-209d-540e-8aea-75f648917ca0" Libdl = "8f399da3-3557-5675-b5ff-fb832c97cbdb" LinearAlgebra = "37e2e46d-f89d-539d-b4ee-838fcccc9c8e" -NEO_jll = "700fe977-ac61-5f37-bbc8-c6c4b2b6a9fd" +NEO_LTS_jll = "a724f90f-ce79-56dd-a1bd-b9de5a61085f" PrecompileTools = "aea7be01-6a6a-4083-8856-8a6e6704d82a" 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] @@ -42,20 +42,20 @@ GPUCompiler = "2" GPUToolbox = "0.1, 0.2, 0.3, 1, 3" KernelAbstractions = "0.9.39" LLVM = "6, 7, 8, 9" -NEO_jll = "=26.18.38308" +NEO_LTS_jll = "=25.18.33578" PrecompileTools = "1" 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" [sources] # Pin to a commit SHA, not the mutable `atomic-float-ops` branch head, so builds are 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 928c1de9..8d6d02f1 100644 --- a/src/compiler/compilation.jl +++ b/src/compiler/compilation.jl @@ -74,18 +74,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. Both conditions are read from the - # (device-independent) compiler config so this stays valid without a live - # device: `_compiler_config` sets `supports_bfloat16` from the device and adds - # the `SPV_KHR_bfloat16` extension iff the driver's SPIR-V runtime accepts it. - target = job.config.target - if @static(isdefined(Core, :BFloat16) && isdefined(LLVM, :BFloatType)) && - target.supports_bfloat16 && !occursin("SPV_KHR_bfloat16", target.extensions) - lower_bfloat_to_i16!(mod) - end - return entry end @@ -194,106 +182,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 (configure, compile, and link) # cache of compiler configurations, per device (but additionally configurable via kwargs) @@ -308,42 +196,23 @@ 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(dev=device()) - return @static if isdefined(Core, :BFloat16) - haskey( - oneL0.extension_properties(dev.driver), - oneL0.ZE_BFLOAT16_CONVERSIONS_EXT_NAME - ) - else - false - end -end - @noinline function _compiler_config(dev; kernel=true, name=nothing, always_inline=false, kwargs...) properties = oneL0.module_properties(dev) supports_fp16 = properties.fp16flags & oneL0.ZE_DEVICE_MODULE_FLAG_FP16 == oneL0.ZE_DEVICE_MODULE_FLAG_FP16 supports_fp64 = properties.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(dev) - - # SPIR-V extensions the LLVM back-end may emit. Declaring them permits the - # corresponding instructions during translation: without - # SPV_EXT_shader_atomic_float_add, floating-point atomic operations fail to - # translate ("The atomic float instruction requires ... SPV_EXT_shader_atomic_float_add"). + + # The LTS NEO/IGC runtime only accepts SPIR-V from the Khronos translator and needs + # these extensions declared explicitly. + 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", + "SPV_EXT_shader_atomic_float_add" ] - # Only add the SPIR-V extension if the runtime actually supports it - if _driver_supports_bfloat16_spirv(dev) - push!(extensions, "SPV_KHR_bfloat16") - end extensions_str = join(map(ext -> "+$ext", extensions), ",") # create GPUCompiler objects - target = SPIRVCompilerTarget(; extensions=extensions_str, supports_fp16, supports_fp64, supports_bfloat16, kwargs...) + target = SPIRVCompilerTarget(; backend, extensions=extensions_str, supports_fp16, supports_fp64, kwargs...) params = oneAPICompilerParams() CompilerConfig(target, params; kernel, name, always_inline) end diff --git a/src/compiler/precompile.jl b/src/compiler/precompile.jl index 4a8493c6..be68ae4e 100644 --- a/src/compiler/precompile.jl +++ b/src/compiler/precompile.jl @@ -3,9 +3,9 @@ using PrecompileTools: @compile_workload # Warm up the GPUCompiler -> SPIR-V pipeline during precompilation so the first real # `zefunction` call is cheap. This doesn't need a GPU: SPIR-V codegen is device-independent # (the target only carries version/extension/capability knobs), and the workload never -# launches anything. It is gated on the SPIR-V LLVM back-end being available so the package +# launches anything. It is gated on the SPIR-V translator being available so the package # still precompiles on platforms/toolchains without it. -if SPIRV_LLVM_Backend_jll.is_available() +if SPIRV_LLVM_Translator_jll.is_available() @compile_workload begin let function _precompile_kernel(a) @@ -16,8 +16,9 @@ if SPIRV_LLVM_Backend_jll.is_available() # Build a device-independent compiler config. `_compiler_config` normally derives # these knobs from the device; here we use conservative, portable defaults (the # workload only exercises the pipeline, it does not target a specific device). - target = SPIRVCompilerTarget(; extensions="", supports_fp16=true, - supports_fp64=true, supports_bfloat16=false) + target = SPIRVCompilerTarget(; backend=:khronos, extensions="", + supports_fp16=true, supports_fp64=true, + supports_bfloat16=false) params = oneAPICompilerParams() config = CompilerConfig(target, params; kernel=true, name=nothing, always_inline=false) diff --git a/src/oneAPI.jl b/src/oneAPI.jl index 2abcb64e..15c2426d 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 @@ -142,6 +142,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 @@ -153,7 +162,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 788581e1..8f585b8d 100644 --- a/test/Project.toml +++ b/test/Project.toml @@ -10,18 +10,17 @@ InteractiveUtils = "b77e0a4c-d291-57a0-90e8-8db25a27a240" JLD2 = "033835bb-8acc-5ee8-8aae-3f567f8a3819" KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c" LinearAlgebra = "37e2e46d-f89d-539d-b4ee-838fcccc9c8e" -NEO_jll = "700fe977-ac61-5f37-bbc8-c6c4b2b6a9fd" +NEO_LTS_jll = "a724f90f-ce79-56dd-a1bd-b9de5a61085f" ParallelTestRunner = "d3525ed8-44d0-4b2c-a655-542cee43accc" Printf = "de0858da-6303-5e67-8744-51eddeeeb8d7" REPL = "3fa0cd96-eef1-5676-8a61-b3b8758bbffb" Random = "9a3f8284-a2c9-5f02-9a11-845980a1fd5c" -SPIRV_LLVM_Backend_jll = "4376b9bf-cff8-51b6-bb48-39421dff0d0c" SparseArrays = "2f01184e-e22b-5df5-ae63-d93ebab69eaf" SpecialFunctions = "276daf66-3868-5448-9aa4-cd146d93841b" StaticArrays = "90137ffa-7385-5640-81b9-e52037218182" Statistics = "10745b16-79ce-11e8-11f9-7d13ad32a3b2" Test = "8dfed614-e22c-5e08-85e1-65c5234f0b40" -libigc_jll = "94295238-5935-5bd7-bb0f-b00942e9bdd5" +libigc_LTS_jll = "9a8258a1-e827-5686-bee9-144461246960" oneAPI = "8f75cd03-7ff8-4ecb-9b8f-daf728133b1b" oneAPI_Support_jll = "b049733a-a71d-5ed3-8eba-7d323ac00b36" From 916f1e533d8cbd2d3733f867a508c5dbee05d7c9 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Tue, 23 Jun 2026 15:45:59 +0000 Subject: [PATCH 06/29] 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 9521e79befa0db34d7427ecae116a473da461a9a Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Tue, 23 Jun 2026 15:46:00 +0000 Subject: [PATCH 07/29] 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 fcc6a3a4..488df2c3 100644 --- a/src/pool.jl +++ b/src/pool.jl @@ -110,6 +110,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 b23d6b22f67852cd418a77991d864317273ec7cb Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Tue, 23 Jun 2026 15:46:00 +0000 Subject: [PATCH 08/29] 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 ff5fe3dfe6ce3c4cc0c7d32188419ac997122ca3 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Tue, 23 Jun 2026 15:46:00 +0000 Subject: [PATCH 09/29] 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 318847e86a66b4856fb2bcc8be6da3bc4638026f Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Tue, 23 Jun 2026 15:46:00 +0000 Subject: [PATCH 10/29] 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 75349da16dc0006ede6b643184bc36b8cbe7df39 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Tue, 23 Jun 2026 16:35:58 +0000 Subject: [PATCH 11/29] 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 | 24 ++++++++++++++++-------- src/context.jl | 25 +++++++++++++++---------- src/mapreduce.jl | 4 ++-- src/pool.jl | 7 +++++-- 6 files changed, 68 insertions(+), 32 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 8d6d02f1..6b527501 100644 --- a/src/compiler/compilation.jl +++ b/src/compiler/compilation.jl @@ -201,14 +201,22 @@ end supports_fp16 = properties.fp16flags & oneL0.ZE_DEVICE_MODULE_FLAG_FP16 == oneL0.ZE_DEVICE_MODULE_FLAG_FP16 supports_fp64 = properties.fp64flags & oneL0.ZE_DEVICE_MODULE_FLAG_FP64 == oneL0.ZE_DEVICE_MODULE_FLAG_FP64 - # The LTS NEO/IGC runtime only accepts SPIR-V from the Khronos translator and needs - # these extensions declared explicitly. - 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" - ] + # 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 extensions_str = join(map(ext -> "+$ext", extensions), ",") # create GPUCompiler objects 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 488df2c3..d67f58c8 100644 --- a/src/pool.jl +++ b/src/pool.jl @@ -116,8 +116,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 7f53f32af9f0399e6921cc6d3103721f97196e8a Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 1 Jul 2026 14:48:27 +0000 Subject: [PATCH 12/29] 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 9256ddfd..4844968c 100644 --- a/Project.toml +++ b/Project.toml @@ -45,7 +45,7 @@ LLVM = "6, 7, 8, 9" NEO_LTS_jll = "=25.18.33578" PrecompileTools = "1" Preferences = "1" -SPIRVIntrinsics = "0.5" +SPIRVIntrinsics = "1" SPIRV_LLVM_Translator_jll = "21" SPIRV_Tools_jll = "2025.4.0" SpecialFunctions = "1.3, 2" From 874789cb0c299fef759d464517c286a05768e47d Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Mon, 6 Jul 2026 15:05:17 +0000 Subject: [PATCH 13/29] 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/runtests.jl | 19 +++++++++++-------- 1 file changed, 11 insertions(+), 8 deletions(-) diff --git a/test/runtests.jl b/test/runtests.jl index 37bf165d..911b976d 100644 --- a/test/runtests.jl +++ b/test/runtests.jl @@ -94,13 +94,16 @@ init_worker_code = quote append!(eltypes, [Float64, ComplexF64]) end @static if isdefined(Core, :BFloat16) - # Exercising BFloat16 as a generic element type needs the testsuite's CPU - # reference path to work, which requires more host-side BFloat16 support than - # BFloat16s.jl currently implements (e.g. the div/rem family throws "rem not - # defined for BFloat16"). Probe the operations the testsuite relies on, so the - # element type enables itself automatically once BFloat16s.jl catches up. - # Device-side BFloat16 (oneMKL gemm/gemv, the bfloat16.jl example) is tested - # regardless. + # Exercising BFloat16 as a generic element type needs (1) hardware support, + # (2) a non-LTS stack — 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 — and (3) a working CPU reference + # path, which requires more host-side BFloat16 support than BFloat16s.jl + # currently implements (e.g. the div/rem family throws "rem not defined for + # BFloat16"). Probe the operations the testsuite relies on, so the element + # type enables itself automatically once BFloat16s.jl catches up. Device-side + # BFloat16 (oneMKL gemm/gemv, the bfloat16.jl example) is tested regardless. bfloat16_host_support() = try x, y = BFloat16(3), BFloat16(2) mod(x, y); fld(x, y); rand(BFloat16) @@ -109,7 +112,7 @@ init_worker_code = quote false end const bfloat16_supported = oneAPI._device_supports_bfloat16() && - bfloat16_host_support() + !oneL0.LTS[] && bfloat16_host_support() if bfloat16_supported push!(eltypes, Core.BFloat16) end From cc6a38b87acdc211d19937dd0ca4dd5648cf8b57 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Tue, 7 Jul 2026 13:48:02 +0000 Subject: [PATCH 14/29] 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 ced6b43ef603a1f41c8dd05b18da2ca1c8bca384 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Tue, 7 Jul 2026 14:38:01 +0000 Subject: [PATCH 15/29] 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 0ea6efba7661a0a3bc8465b3c427ae97c3dbcbc5 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 14:52:12 +0000 Subject: [PATCH 16/29] 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 aff9f1d4b10073e264a222548c012cf1a715399f Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 14:53:24 +0000 Subject: [PATCH 17/29] 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 b6aa24f75b0dc6f84225159e236e06177198fa62 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 14:54:14 +0000 Subject: [PATCH 18/29] 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 5224bd32c5e0047ebbd0e95e03002d0f20b24174 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 14:56:44 +0000 Subject: [PATCH 19/29] 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 a145078710d1cfdcd6599a69e0bc3123c43e6ff8 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 15:04:42 +0000 Subject: [PATCH 20/29] 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 33d6f777d9a656693788b4731a3314ecbda77c56 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 15:06:39 +0000 Subject: [PATCH 21/29] 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/compiler/precompile.jl | 19 +++++++++++++------ src/oneAPI.jl | 6 +++++- 3 files changed, 20 insertions(+), 7 deletions(-) diff --git a/Project.toml b/Project.toml index 4844968c..73d5a8b7 100644 --- a/Project.toml +++ b/Project.toml @@ -22,6 +22,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" @@ -46,6 +47,7 @@ NEO_LTS_jll = "=25.18.33578" PrecompileTools = "1" 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/compiler/precompile.jl b/src/compiler/precompile.jl index be68ae4e..b20a2f7b 100644 --- a/src/compiler/precompile.jl +++ b/src/compiler/precompile.jl @@ -3,9 +3,16 @@ using PrecompileTools: @compile_workload # Warm up the GPUCompiler -> SPIR-V pipeline during precompilation so the first real # `zefunction` call is cheap. This doesn't need a GPU: SPIR-V codegen is device-independent # (the target only carries version/extension/capability knobs), and the workload never -# launches anything. It is gated on the SPIR-V translator being available so the package -# still precompiles on platforms/toolchains without it. -if SPIRV_LLVM_Translator_jll.is_available() +# launches anything. It is gated on the SPIR-V tool for the selected codegen path being +# available so the package still precompiles on platforms/toolchains without it. +# +# The workload mirrors the runtime backend choice (see `_compiler_config`): the LTS stack +# compiles through the Khronos translator, the rolling stack through the LLVM SPIR-V +# back-end. `oneL0.__init__` has not run at precompile time, so resolve ONEAPI_LTS +# directly from the environment with the same parser. +const _precompile_lts = oneL0.parse_env_bool("ONEAPI_LTS", false) +if _precompile_lts ? SPIRV_LLVM_Translator_jll.is_available() : + SPIRV_LLVM_Backend_jll.is_available() @compile_workload begin let function _precompile_kernel(a) @@ -16,9 +23,9 @@ if SPIRV_LLVM_Translator_jll.is_available() # Build a device-independent compiler config. `_compiler_config` normally derives # these knobs from the device; here we use conservative, portable defaults (the # workload only exercises the pipeline, it does not target a specific device). - target = SPIRVCompilerTarget(; backend=:khronos, extensions="", - supports_fp16=true, supports_fp64=true, - supports_bfloat16=false) + target = SPIRVCompilerTarget(; backend=(_precompile_lts ? :khronos : :llvm), + extensions="", supports_fp16=true, + supports_fp64=true, supports_bfloat16=false) params = oneAPICompilerParams() config = CompilerConfig(target, params; kernel=true, name=nothing, always_inline=false) diff --git a/src/oneAPI.jl b/src/oneAPI.jl index 15c2426d..0b2272e8 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 11d800ae8182a5e14e099183e436400652257c37 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 15:10:34 +0000 Subject: [PATCH 22/29] 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 | 140 +++++++++++++++++++++++++++++++++++- 1 file changed, 139 insertions(+), 1 deletion(-) diff --git a/src/compiler/compilation.jl b/src/compiler/compilation.jl index 6b527501..1fdbb742 100644 --- a/src/compiler/compilation.jl +++ b/src/compiler/compilation.jl @@ -74,6 +74,21 @@ 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. Both conditions are read from the + # (device-independent) compiler config so this stays valid without a live + # device: `_compiler_config` sets `supports_bfloat16` from the device and adds + # the `SPV_KHR_bfloat16` extension iff the driver's SPIR-V runtime accepts it. + # Rolling stack only: bf16 is forced off on the LTS stack + # (`supports_bfloat16 = false`, see `_compiler_config`), so this pass never + # runs there. + target = job.config.target + if @static(isdefined(Core, :BFloat16) && isdefined(LLVM, :BFloatType)) && + target.supports_bfloat16 && !occursin("SPV_KHR_bfloat16", target.extensions) + lower_bfloat_to_i16!(mod) + end + return entry end @@ -182,6 +197,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 (configure, compile, and link) # cache of compiler configurations, per device (but additionally configurable via kwargs) @@ -196,6 +311,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(dev=device()) + return @static if isdefined(Core, :BFloat16) + haskey( + oneL0.extension_properties(dev.driver), + oneL0.ZE_BFLOAT16_CONVERSIONS_EXT_NAME + ) + else + false + end +end + @noinline function _compiler_config(dev; kernel=true, name=nothing, always_inline=false, kwargs...) properties = oneL0.module_properties(dev) supports_fp16 = properties.fp16flags & oneL0.ZE_DEVICE_MODULE_FLAG_FP16 == oneL0.ZE_DEVICE_MODULE_FLAG_FP16 @@ -213,14 +340,25 @@ 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(dev) + # Only declare the SPIR-V extension when the runtime actually accepts it. + if _driver_supports_bfloat16_spirv(dev) + push!(extensions, "SPV_KHR_bfloat16") + end end extensions_str = join(map(ext -> "+$ext", extensions), ",") # create GPUCompiler objects - target = SPIRVCompilerTarget(; backend, extensions=extensions_str, supports_fp16, supports_fp64, kwargs...) + target = SPIRVCompilerTarget(; backend, extensions=extensions_str, supports_fp16, supports_fp64, supports_bfloat16, kwargs...) params = oneAPICompilerParams() CompilerConfig(target, params; kernel, name, always_inline) end From 722549652687881165c45ef6f7feb155527d0fd7 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 15:16:05 +0000 Subject: [PATCH 23/29] 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 6fde199483099dcb58c3e5724552869e2b5cef8a Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 15:17:44 +0000 Subject: [PATCH 24/29] 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 0b2272e8..7786e1a8 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. @@ -147,9 +149,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 26ed8e1a3791479449803a5067e902994aba80f9 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 15:20:57 +0000 Subject: [PATCH 25/29] 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 bd8c70f266e788a549e19d18dbd73e9eed840763 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 15:33:42 +0000 Subject: [PATCH 26/29] 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 534fdd1383d892cdeb21b12e92bb5f1d3308055f Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 8 Jul 2026 16:39:32 +0000 Subject: [PATCH 27/29] 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 From f298063ed5c163a8f293f08e873a10d770e1fd51 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 22 Jul 2026 11:18:57 -0500 Subject: [PATCH 28/29] Declare the atomic-float extensions on the rolling-stack path too The LTS gating commit rebuilt the extension list per code path and left the rolling (LLVM back-end) path empty, regressing what landed on main: without SPV_EXT_shader_atomic_float_add declared, floating-point atomic operations fail to translate. Declare the same two extensions as the LTS path. --- src/compiler/compilation.jl | 9 ++++++++- 1 file changed, 8 insertions(+), 1 deletion(-) diff --git a/src/compiler/compilation.jl b/src/compiler/compilation.jl index 1fdbb742..6ea8c827 100644 --- a/src/compiler/compilation.jl +++ b/src/compiler/compilation.jl @@ -346,7 +346,14 @@ end supports_bfloat16 = false else backend = :llvm - extensions = String[] + # SPIR-V extensions the LLVM back-end may emit. Declaring them permits the + # corresponding instructions during translation: without + # SPV_EXT_shader_atomic_float_add, floating-point atomic operations fail to + # translate ("The atomic float instruction requires ... SPV_EXT_shader_atomic_float_add"). + extensions = String[ + "SPV_EXT_relaxed_printf_string_address_space", + "SPV_EXT_shader_atomic_float_add", + ] # 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(dev) From ddaa4184a90e8390f025ac24e4caad8012f65de6 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Wed, 22 Jul 2026 11:18:57 -0500 Subject: [PATCH 29/29] Default to the rolling driver stack; LTS deployments override externally Shipping NEO_LTS_jll unconditionally breaks every non-LTS system: the LTS NEO 25.18/IGC 2.11 driver cannot link kernels emitted by the LLVM back-end on consumer GPUs (undefined __spirv_IAddCarry in UniformScaling kernels) and segfaults on modules declaring SPV_KHR_bfloat16. Making the JLL choice dynamic is not possible either: NEO_LTS_jll pins gmmlib =22.7.2 while the rolling NEO pins =22.10.0, so the two drivers cannot coexist in one environment. Restore the rolling driver stack as the package dependency, matching main. All ONEAPI_LTS=1 runtime gating (Khronos translator codegen, queue-drain workarounds, sync-each-submission) is unaffected. An LTS deployment such as the Aurora runner provides its driver outside Pkg: either the system driver stack (setting ZE_ENABLE_ALT_DRIVERS to the LTS libze_intel_gpu, which oneL0.__init__ already respects), or the res/local.jl preference mechanism. The deps/ and res/ developer tooling still targets the LTS headers for regenerating wrappers and building the support library. --- Project.toml | 12 ++++++------ lib/level-zero/oneL0.jl | 6 ++---- src/oneAPI.jl | 2 +- src/utils.jl | 2 +- test/Project.toml | 4 ++-- 5 files changed, 12 insertions(+), 14 deletions(-) diff --git a/Project.toml b/Project.toml index 73d5a8b7..951498bf 100644 --- a/Project.toml +++ b/Project.toml @@ -16,7 +16,7 @@ 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_LTS_jll = "a724f90f-ce79-56dd-a1bd-b9de5a61085f" +NEO_jll = "700fe977-ac61-5f37-bbc8-c6c4b2b6a9fd" PrecompileTools = "aea7be01-6a6a-4083-8856-8a6e6704d82a" Preferences = "21216c6a-2e73-6563-6e65-726566657250" Printf = "de0858da-6303-5e67-8744-51eddeeeb8d7" @@ -28,8 +28,8 @@ 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_LTS_jll = "d79c0b2e-896c-561b-aab9-323701ec0314" -oneAPI_Level_Zero_Loader_LTS_jll = "f6e5cbb4-ba2a-56dc-92a2-9d66f5656ccd" +oneAPI_Level_Zero_Headers_jll = "f4bc562b-d309-54f8-9efb-476e56f0410d" +oneAPI_Level_Zero_Loader_jll = "13eca655-d68d-5b81-8367-6d99d727ab01" oneAPI_Support_jll = "b049733a-a71d-5ed3-8eba-7d323ac00b36" [compat] @@ -43,7 +43,7 @@ GPUCompiler = "2" GPUToolbox = "0.1, 0.2, 0.3, 1, 3" KernelAbstractions = "0.9.39" LLVM = "6, 7, 8, 9" -NEO_LTS_jll = "=25.18.33578" +NEO_jll = "=26.18.38308" PrecompileTools = "1" Preferences = "1" SPIRVIntrinsics = "1" @@ -53,11 +53,11 @@ SPIRV_Tools_jll = "2025.4.0" SpecialFunctions = "1.3, 2" StaticArrays = "1" julia = "1.10" -oneAPI_Level_Zero_Loader_LTS_jll = "=1.24" +oneAPI_Level_Zero_Loader_jll = "1.29" oneAPI_Support_jll = "0.9.2" [extras] -libigc_LTS_jll = "9a8258a1-e827-5686-bee9-144461246960" +libigc_jll = "94295238-5935-5bd7-bb0f-b00942e9bdd5" [sources] # Pin to a commit SHA, not the mutable `atomic-float-ops` branch head, so builds are diff --git a/lib/level-zero/oneL0.jl b/lib/level-zero/oneL0.jl index 69636412..e736f631 100644 --- a/lib/level-zero/oneL0.jl +++ b/lib/level-zero/oneL0.jl @@ -11,10 +11,8 @@ using Libdl if Sys.iswindows() const libze_loader = "ze_loader" else - 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 + using NEO_jll + using oneAPI_Level_Zero_Loader_jll end include("utils.jl") diff --git a/src/oneAPI.jl b/src/oneAPI.jl index 7786e1a8..7baebfa4 100644 --- a/src/oneAPI.jl +++ b/src/oneAPI.jl @@ -179,7 +179,7 @@ function __init__() end function set_debug!(debug::Bool) - for jll in [oneL0.NEO_jll, oneL0.NEO_jll.libigc_LTS_jll] + for jll in [oneL0.NEO_jll, oneL0.NEO_jll.libigc_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 3c14edcb..8dc385eb 100644 --- a/src/utils.jl +++ b/src/utils.jl @@ -2,7 +2,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, + for jll in [oneL0.NEO_jll, oneL0.NEO_jll.libigc_jll, oneL0.NEO_jll.gmmlib_jll, SPIRV_LLVM_Backend_jll, SPIRV_LLVM_Translator_jll, SPIRV_Tools_jll, oneAPI_Support_jll] name = string(jll) print(io, "- $(name[1:end-4]): $(Base.pkgversion(jll))") diff --git a/test/Project.toml b/test/Project.toml index 8f585b8d..426e2db7 100644 --- a/test/Project.toml +++ b/test/Project.toml @@ -10,7 +10,7 @@ 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_LTS_jll = "a724f90f-ce79-56dd-a1bd-b9de5a61085f" +NEO_jll = "700fe977-ac61-5f37-bbc8-c6c4b2b6a9fd" ParallelTestRunner = "d3525ed8-44d0-4b2c-a655-542cee43accc" Printf = "de0858da-6303-5e67-8744-51eddeeeb8d7" REPL = "3fa0cd96-eef1-5676-8a61-b3b8758bbffb" @@ -20,7 +20,7 @@ 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_LTS_jll = "9a8258a1-e827-5686-bee9-144461246960" +libigc_jll = "94295238-5935-5bd7-bb0f-b00942e9bdd5" oneAPI = "8f75cd03-7ff8-4ecb-9b8f-daf728133b1b" oneAPI_Support_jll = "b049733a-a71d-5ed3-8eba-7d323ac00b36"