From e8fcb4396d56ac13f0543417363d0f4c4fa1b5ef Mon Sep 17 00:00:00 2001 From: Abhishek Bagusetty Date: Fri, 14 Aug 2026 11:41:42 -0500 Subject: [PATCH 1/6] Release GIL around blocking oneMKL calls in BLAS/LAPACK potrf, getrf, geqrf, gesv, gesvd, heevd, syevd, ungqr and the BLAS dot/gemm family block the calling thread for the whole oneMKL call while holding the GIL. That stalls every other Python thread, including the host tasks dpnp queues to manage object lifetimes, which reacquire the GIL to drop references. Completes the backport of gh-2850, which applied the same fix to orgqr.cpp; gesv_batch, gesvd_batch, heevd_batch and syevd_batch already had it. --- dpnp/backend/extensions/blas/dot.hpp | 8 + dpnp/backend/extensions/blas/dotc.hpp | 8 + dpnp/backend/extensions/blas/dotu.hpp | 8 + dpnp/backend/extensions/blas/gemm.cpp | 5 + dpnp/backend/extensions/blas/gemm_batch.cpp | 5 + dpnp/backend/extensions/lapack/geqrf.cpp | 4 + .../backend/extensions/lapack/geqrf_batch.cpp | 4 + dpnp/backend/extensions/lapack/gesv.cpp | 8 + dpnp/backend/extensions/lapack/gesvd.cpp | 4 + dpnp/backend/extensions/lapack/getrf.cpp | 4 + .../backend/extensions/lapack/getrf_batch.cpp | 4 + dpnp/backend/extensions/lapack/heevd.cpp | 5 + .../backend/extensions/lapack/orgqr_batch.cpp | 4 + dpnp/backend/extensions/lapack/potrf.cpp | 4 + .../backend/extensions/lapack/potrf_batch.cpp | 4 + dpnp/backend/extensions/lapack/syevd.cpp | 5 + dpnp/backend/extensions/lapack/ungqr.cpp | 4 + .../backend/extensions/lapack/ungqr_batch.cpp | 4 + dpnp/tests/test_blas_lapack_gil.py | 176 ++++++++++++++++++ 19 files changed, 268 insertions(+) create mode 100644 dpnp/tests/test_blas_lapack_gil.py diff --git a/dpnp/backend/extensions/blas/dot.hpp b/dpnp/backend/extensions/blas/dot.hpp index 300a5348c53e..f0f90e635f91 100644 --- a/dpnp/backend/extensions/blas/dot.hpp +++ b/dpnp/backend/extensions/blas/dot.hpp @@ -30,12 +30,15 @@ #include +#include + #include "dot_common.hpp" namespace dpnp::extensions::blas { namespace mkl_blas = oneapi::mkl::blas; namespace type_utils = dpnp::tensor::type_utils; +namespace py = pybind11; template static sycl::event dot_impl(sycl::queue &exec_q, @@ -58,6 +61,11 @@ static sycl::event dot_impl(sycl::queue &exec_q, sycl::event dot_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + + dot_event = mkl_blas::column_major::dot(exec_q, n, // size of the input vectors x, // Pointer to vector x. diff --git a/dpnp/backend/extensions/blas/dotc.hpp b/dpnp/backend/extensions/blas/dotc.hpp index 535f79156fb5..273ebfa032cd 100644 --- a/dpnp/backend/extensions/blas/dotc.hpp +++ b/dpnp/backend/extensions/blas/dotc.hpp @@ -30,12 +30,15 @@ #include +#include + #include "dot_common.hpp" namespace dpnp::extensions::blas { namespace mkl_blas = oneapi::mkl::blas; namespace type_utils = dpnp::tensor::type_utils; +namespace py = pybind11; template static sycl::event dotc_impl(sycl::queue &exec_q, @@ -58,6 +61,11 @@ static sycl::event dotc_impl(sycl::queue &exec_q, sycl::event dotc_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + + dotc_event = mkl_blas::column_major::dotc(exec_q, n, // size of the input vectors diff --git a/dpnp/backend/extensions/blas/dotu.hpp b/dpnp/backend/extensions/blas/dotu.hpp index aafe9721339f..485fdd33f938 100644 --- a/dpnp/backend/extensions/blas/dotu.hpp +++ b/dpnp/backend/extensions/blas/dotu.hpp @@ -30,12 +30,15 @@ #include +#include + #include "dot_common.hpp" namespace dpnp::extensions::blas { namespace mkl_blas = oneapi::mkl::blas; namespace type_utils = dpnp::tensor::type_utils; +namespace py = pybind11; template static sycl::event dotu_impl(sycl::queue &exec_q, @@ -58,6 +61,11 @@ static sycl::event dotu_impl(sycl::queue &exec_q, sycl::event dotu_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + + dotu_event = mkl_blas::column_major::dotu(exec_q, n, // size of the input vectors diff --git a/dpnp/backend/extensions/blas/gemm.cpp b/dpnp/backend/extensions/blas/gemm.cpp index bf91f5cfb0bf..d0901097c759 100644 --- a/dpnp/backend/extensions/blas/gemm.cpp +++ b/dpnp/backend/extensions/blas/gemm.cpp @@ -113,6 +113,11 @@ static sycl::event gemm_impl(sycl::queue &exec_q, c, ldc, deps); } }; + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + + gemm_event = gemm_func( exec_q, transA, // Defines the transpose operation for matrix A: diff --git a/dpnp/backend/extensions/blas/gemm_batch.cpp b/dpnp/backend/extensions/blas/gemm_batch.cpp index f1bc9f30b5dd..f3a0c009214e 100644 --- a/dpnp/backend/extensions/blas/gemm_batch.cpp +++ b/dpnp/backend/extensions/blas/gemm_batch.cpp @@ -125,6 +125,11 @@ static sycl::event gemm_batch_impl(sycl::queue &exec_q, strideb, beta, c, ldc, stridec, batch_size, deps); } }; + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + + gemm_batch_event = gemm_batch_func( exec_q, transA, // Defines the transpose operation for matrix A: diff --git a/dpnp/backend/extensions/lapack/geqrf.cpp b/dpnp/backend/extensions/lapack/geqrf.cpp index 083c408ea3d4..03e9c5060fdd 100644 --- a/dpnp/backend/extensions/lapack/geqrf.cpp +++ b/dpnp/backend/extensions/lapack/geqrf.cpp @@ -86,6 +86,10 @@ static sycl::event geqrf_impl(sycl::queue &exec_q, sycl::event geqrf_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + scratchpad = sycl::malloc_device(scratchpad_size, exec_q); geqrf_event = mkl_lapack::geqrf( diff --git a/dpnp/backend/extensions/lapack/geqrf_batch.cpp b/dpnp/backend/extensions/lapack/geqrf_batch.cpp index 6e62d4dde336..739e2a82b543 100644 --- a/dpnp/backend/extensions/lapack/geqrf_batch.cpp +++ b/dpnp/backend/extensions/lapack/geqrf_batch.cpp @@ -94,6 +94,10 @@ static sycl::event geqrf_batch_impl(sycl::queue &exec_q, sycl::event geqrf_batch_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + scratchpad = sycl::malloc_device(scratchpad_size, exec_q); geqrf_batch_event = mkl_lapack::geqrf_batch( diff --git a/dpnp/backend/extensions/lapack/gesv.cpp b/dpnp/backend/extensions/lapack/gesv.cpp index 072efd169c87..3bab2e5b35a3 100644 --- a/dpnp/backend/extensions/lapack/gesv.cpp +++ b/dpnp/backend/extensions/lapack/gesv.cpp @@ -112,6 +112,10 @@ static sycl::event gesv_impl(sycl::queue &exec_q, #if defined(USE_ONEMATH) sycl::event getrf_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + getrf_event = mkl_lapack::getrf( exec_q, n, // The order of the square matrix A (0 ≤ n). @@ -169,6 +173,10 @@ static sycl::event gesv_impl(sycl::queue &exec_q, } #else try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + comp_event = mkl_lapack::gesv( exec_q, n, // The order of the square matrix A diff --git a/dpnp/backend/extensions/lapack/gesvd.cpp b/dpnp/backend/extensions/lapack/gesvd.cpp index ebc6000671b0..02b388977144 100644 --- a/dpnp/backend/extensions/lapack/gesvd.cpp +++ b/dpnp/backend/extensions/lapack/gesvd.cpp @@ -99,6 +99,10 @@ static sycl::event gesvd_impl(sycl::queue &exec_q, sycl::event gesvd_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + gesvd_event = mkl_lapack::gesvd( exec_q, jobu, // Character specifying how to compute the matrix U: diff --git a/dpnp/backend/extensions/lapack/getrf.cpp b/dpnp/backend/extensions/lapack/getrf.cpp index 42021639d63f..c31213b38620 100644 --- a/dpnp/backend/extensions/lapack/getrf.cpp +++ b/dpnp/backend/extensions/lapack/getrf.cpp @@ -87,6 +87,10 @@ static sycl::event getrf_impl(sycl::queue &exec_q, sycl::event getrf_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + scratchpad = sycl::malloc_device(scratchpad_size, exec_q); getrf_event = mkl_lapack::getrf( diff --git a/dpnp/backend/extensions/lapack/getrf_batch.cpp b/dpnp/backend/extensions/lapack/getrf_batch.cpp index b4e30dae57d3..917794557378 100644 --- a/dpnp/backend/extensions/lapack/getrf_batch.cpp +++ b/dpnp/backend/extensions/lapack/getrf_batch.cpp @@ -96,6 +96,10 @@ static sycl::event getrf_batch_impl(sycl::queue &exec_q, sycl::event getrf_batch_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + scratchpad = sycl::malloc_device(scratchpad_size, exec_q); getrf_batch_event = mkl_lapack::getrf_batch( diff --git a/dpnp/backend/extensions/lapack/heevd.cpp b/dpnp/backend/extensions/lapack/heevd.cpp index 002a7ad2ef89..86ba97765884 100644 --- a/dpnp/backend/extensions/lapack/heevd.cpp +++ b/dpnp/backend/extensions/lapack/heevd.cpp @@ -42,6 +42,7 @@ namespace dpnp::extensions::lapack { +namespace py = pybind11; namespace mkl_lapack = oneapi::mkl::lapack; namespace type_utils = dpnp::tensor::type_utils; @@ -73,6 +74,10 @@ static sycl::event heevd_impl(sycl::queue &exec_q, sycl::event heevd_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + heevd_event = mkl_lapack::heevd( exec_q, jobz, // 'jobz == job::vec' means eigenvalues and eigenvectors are diff --git a/dpnp/backend/extensions/lapack/orgqr_batch.cpp b/dpnp/backend/extensions/lapack/orgqr_batch.cpp index 840ef82a09b6..0d996bef2b87 100644 --- a/dpnp/backend/extensions/lapack/orgqr_batch.cpp +++ b/dpnp/backend/extensions/lapack/orgqr_batch.cpp @@ -96,6 +96,10 @@ static sycl::event orgqr_batch_impl(sycl::queue &exec_q, sycl::event orgqr_batch_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + scratchpad = sycl::malloc_device(scratchpad_size, exec_q); orgqr_batch_event = mkl_lapack::orgqr_batch( diff --git a/dpnp/backend/extensions/lapack/potrf.cpp b/dpnp/backend/extensions/lapack/potrf.cpp index d7edc2d244f6..a14ce223458c 100644 --- a/dpnp/backend/extensions/lapack/potrf.cpp +++ b/dpnp/backend/extensions/lapack/potrf.cpp @@ -83,6 +83,10 @@ static sycl::event potrf_impl(sycl::queue &exec_q, sycl::event potrf_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + scratchpad = sycl::malloc_device(scratchpad_size, exec_q); potrf_event = mkl_lapack::potrf( diff --git a/dpnp/backend/extensions/lapack/potrf_batch.cpp b/dpnp/backend/extensions/lapack/potrf_batch.cpp index 58bdad462930..8193b32d4f05 100644 --- a/dpnp/backend/extensions/lapack/potrf_batch.cpp +++ b/dpnp/backend/extensions/lapack/potrf_batch.cpp @@ -90,6 +90,10 @@ static sycl::event potrf_batch_impl(sycl::queue &exec_q, sycl::event potrf_batch_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + scratchpad = sycl::malloc_device(scratchpad_size, exec_q); potrf_batch_event = mkl_lapack::potrf_batch( diff --git a/dpnp/backend/extensions/lapack/syevd.cpp b/dpnp/backend/extensions/lapack/syevd.cpp index 7280e415f271..f2be6214f60a 100644 --- a/dpnp/backend/extensions/lapack/syevd.cpp +++ b/dpnp/backend/extensions/lapack/syevd.cpp @@ -42,6 +42,7 @@ namespace dpnp::extensions::lapack { +namespace py = pybind11; namespace mkl_lapack = oneapi::mkl::lapack; namespace type_utils = dpnp::tensor::type_utils; @@ -73,6 +74,10 @@ static sycl::event syevd_impl(sycl::queue &exec_q, sycl::event syevd_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + syevd_event = mkl_lapack::syevd( exec_q, jobz, // 'jobz == job::vec' means eigenvalues and eigenvectors are diff --git a/dpnp/backend/extensions/lapack/ungqr.cpp b/dpnp/backend/extensions/lapack/ungqr.cpp index 7bea568f048e..42b14059176f 100644 --- a/dpnp/backend/extensions/lapack/ungqr.cpp +++ b/dpnp/backend/extensions/lapack/ungqr.cpp @@ -87,6 +87,10 @@ static sycl::event ungqr_impl(sycl::queue &exec_q, sycl::event ungqr_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + scratchpad = sycl::malloc_device(scratchpad_size, exec_q); ungqr_event = mkl_lapack::ungqr( diff --git a/dpnp/backend/extensions/lapack/ungqr_batch.cpp b/dpnp/backend/extensions/lapack/ungqr_batch.cpp index 6452a30e2a86..1ac7220064be 100644 --- a/dpnp/backend/extensions/lapack/ungqr_batch.cpp +++ b/dpnp/backend/extensions/lapack/ungqr_batch.cpp @@ -96,6 +96,10 @@ static sycl::event ungqr_batch_impl(sycl::queue &exec_q, sycl::event ungqr_batch_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + scratchpad = sycl::malloc_device(scratchpad_size, exec_q); ungqr_batch_event = mkl_lapack::ungqr_batch( diff --git a/dpnp/tests/test_blas_lapack_gil.py b/dpnp/tests/test_blas_lapack_gil.py new file mode 100644 index 000000000000..d017b3422a02 --- /dev/null +++ b/dpnp/tests/test_blas_lapack_gil.py @@ -0,0 +1,176 @@ +"""Blocking oneMKL calls in the BLAS/LAPACK extensions must release the GIL. + +Routines such as ``potrf``, ``getrf``, ``syevd`` and ``gesv`` block the calling +thread for the whole oneMKL call. Holding the GIL across that stalls every +other Python thread, including the host tasks dpnp queues to manage object +lifetimes, which reacquire the GIL. Five sibling files already wrap their call +in ``py::gil_scoped_release`` (gh-2850 did this for ``orgqr``); these tests +cover the rest. + +Progress of a competing thread is compared against ``SyclQueue.wait()``, which +is ``nogil`` and therefore the best rate achievable on the machine. +""" + +import sys +import threading +import time + +import dpctl +import pytest + +import dpnp + +from .helper import has_support_aspect64, is_gpu_device + +# Large enough that every routine here actually blocks. Smaller sizes stop +# discriminating: at n=512 cholesky returns asynchronously, and at n=1024 +# gesv blocks for only ~7 ms and passes even without the fix. +_SIZE = 2048 + +# A releasing call matches the nogil reference within a factor of ~2; a call +# holding the GIL measures ~0.00-0.04 of it. +_MIN_RATIO = 0.10 + +_BACKLOG = 2 # queued matmuls, so the measured call has something to wait on + + +class _Ticker: + """Counts how often a competing Python thread gets scheduled.""" + + def __enter__(self): + self.ticks = 0 + self._stop = False + self._thread = threading.Thread(target=self._run, daemon=True) + self._thread.start() + return self + + def _run(self): + while not self._stop: + self.ticks += 1 + time.sleep(0) + + def __exit__(self, *exc): + self._stop = True + self._thread.join(timeout=5) + return False + + +@pytest.mark.skipif(not is_gpu_device(), reason="requires a GPU device") +@pytest.mark.skipif(not has_support_aspect64(), reason="requires fp64 support") +class TestBlockingCallsReleaseGil: + @pytest.fixture(autouse=True) + def _switch_interval(self): + # Stop CPython handing the GIL over on its own timer, so the + # measurement reflects the extension rather than the interpreter. + previous = sys.getswitchinterval() + sys.setswitchinterval(0.001) + yield + sys.setswitchinterval(previous) + + @pytest.fixture + def mats(self): + spd = dpnp.eye(_SIZE, dtype="f8") * float(_SIZE) # positive definite + return spd, dpnp.eye(_SIZE, dtype="f8") + 0.1 + + def _assert_releases_gil(self, name, spd, fn): + queue = spd.sycl_queue + + def rate(ticker, func): + for _ in range(_BACKLOG): + dpnp.matmul(spd, spd) + ticker.ticks = 0 + start = time.perf_counter() + func() + elapsed_ms = 1000 * (time.perf_counter() - start) + queue.wait() + return ticker.ticks / max(elapsed_ms, 1e-3) + + fn() # warm up JIT + queue.wait() + + with _Ticker() as ticker: + measured = rate(ticker, fn) + reference = rate(ticker, queue.wait) + + assert reference > 0, "reference measurement produced no ticks" + ratio = measured / reference + assert ratio >= _MIN_RATIO, ( + f"{name} holds the GIL while blocking: {measured:.2f} ticks/ms vs " + f"{reference:.2f} for nogil queue.wait() (ratio {ratio:.3f}, need " + f">= {_MIN_RATIO}). The oneMKL call needs py::gil_scoped_release." + ) + + def test_potrf(self, mats): + spd, _ = mats + self._assert_releases_gil("potrf", spd, lambda: dpnp.linalg.cholesky(spd)) + + def test_getrf(self, mats): + spd, gen = mats + self._assert_releases_gil("getrf", spd, lambda: dpnp.linalg.det(gen)) + + def test_syevd(self, mats): + spd, _ = mats + self._assert_releases_gil("syevd", spd, lambda: dpnp.linalg.eigh(spd)) + + # No qr/geqrf case: dpnp.linalg.qr() also calls orgqr, already fixed in + # gh-2850, and that release alone keeps the competing thread running. The + # measurement would pass with or without geqrf's own release. + + def test_gesv(self, mats): + spd, gen = mats + rhs = dpnp.ones(_SIZE, dtype="f8") + self._assert_releases_gil( + "gesv", spd, lambda: dpnp.linalg.solve(gen, rhs) + ) + + +@pytest.mark.skipif(not is_gpu_device(), reason="requires a GPU device") +@pytest.mark.skipif(not has_support_aspect64(), reason="requires fp64 support") +def test_multithreaded_linalg_overlaps(): + """Two threads on two queues must run oneMKL work concurrently. + + Holding the GIL across the blocking call serializes them, so a + multithreaded application gets no scaling from a second queue. + """ + max_ratio = 0.85 + + queues = [ + dpctl.SyclQueue(dpctl.SyclDevice("gpu"), property="in_order") + for _ in range(2) + ] + mats = [ + dpnp.eye(_SIZE, dtype="f8", sycl_queue=q) * float(_SIZE) for q in queues + ] + + def work(idx): + for _ in range(2): + dpnp.linalg.eigh(mats[idx]) + queues[idx].wait() + + for idx in range(2): # warm up JIT + work(idx) + + previous = sys.getswitchinterval() + sys.setswitchinterval(0.001) + try: + start = time.perf_counter() + work(0) + work(1) + serial = time.perf_counter() - start + + start = time.perf_counter() + threads = [threading.Thread(target=work, args=(i,)) for i in range(2)] + for thread in threads: + thread.start() + for thread in threads: + thread.join() + parallel = time.perf_counter() - start + finally: + sys.setswitchinterval(previous) + + ratio = parallel / serial + assert ratio <= max_ratio, ( + f"two threads running eigh did not overlap: {parallel * 1000:.1f} ms " + f"parallel vs {serial * 1000:.1f} ms serial (ratio {ratio:.3f}, need " + f"<= {max_ratio}). The oneMKL call is holding the GIL." + ) From 3dadf6791ee4445214521338ad6aff931668c282 Mon Sep 17 00:00:00 2001 From: Abhishek Bagusetty Date: Fri, 14 Aug 2026 11:41:42 -0500 Subject: [PATCH 2/6] Release GIL around blocking oneMKL calls in BLAS/LAPACK potrf, getrf, geqrf, gesv, gesvd, heevd, syevd, ungqr and the BLAS dot/gemm family block the calling thread for the whole oneMKL call while holding the GIL. That stalls every other Python thread, including the host tasks dpnp queues to manage object lifetimes, which reacquire the GIL to drop references. Completes the backport of gh-2850, which applied the same fix to orgqr.cpp; gesv_batch, gesvd_batch, heevd_batch and syevd_batch already had it. --- CHANGELOG.md | 1 + dpnp/backend/extensions/blas/dot.hpp | 8 + dpnp/backend/extensions/blas/dotc.hpp | 8 + dpnp/backend/extensions/blas/dotu.hpp | 8 + dpnp/backend/extensions/blas/gemm.cpp | 5 + dpnp/backend/extensions/blas/gemm_batch.cpp | 5 + dpnp/backend/extensions/lapack/geqrf.cpp | 4 + .../backend/extensions/lapack/geqrf_batch.cpp | 4 + dpnp/backend/extensions/lapack/gesv.cpp | 8 + dpnp/backend/extensions/lapack/gesvd.cpp | 4 + dpnp/backend/extensions/lapack/getrf.cpp | 4 + .../backend/extensions/lapack/getrf_batch.cpp | 4 + dpnp/backend/extensions/lapack/heevd.cpp | 5 + .../backend/extensions/lapack/orgqr_batch.cpp | 4 + dpnp/backend/extensions/lapack/potrf.cpp | 4 + .../backend/extensions/lapack/potrf_batch.cpp | 4 + dpnp/backend/extensions/lapack/syevd.cpp | 5 + dpnp/backend/extensions/lapack/ungqr.cpp | 4 + .../backend/extensions/lapack/ungqr_batch.cpp | 4 + dpnp/tests/test_blas_lapack_gil.py | 176 ++++++++++++++++++ 20 files changed, 269 insertions(+) create mode 100644 dpnp/tests/test_blas_lapack_gil.py diff --git a/CHANGELOG.md b/CHANGELOG.md index 7dee9a8a90e9..7d194577d1e3 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -79,6 +79,7 @@ This release is compatible with NumPy 2.5. * Fixed missing strides validation in `dpnp.tensor.usm_ndarray` constructor when allocating new memory [#2927](https://github.com/IntelPython/dpnp/pull/2927) * Fixed `dpnp.bincount` raising a `ValueError` on an empty input array instead of returning an empty `intp` array [#3018](https://github.com/IntelPython/dpnp/pull/3018) * Fixed `dpnp.tensor.top_k` aborting for `k=0` by returning empty result arrays without launching a zero-sized kernel [#3022](https://github.com/IntelPython/dpnp/pull/3022) +* Released the GIL before the remaining blocking OneMKL BLAS and LAPACK calls to prevent host tasks contention, completing the work started in [#2850](https://github.com/IntelPython/dpnp/pull/2850) [#3027](https://github.com/IntelPython/dpnp/pull/3027) ### Security diff --git a/dpnp/backend/extensions/blas/dot.hpp b/dpnp/backend/extensions/blas/dot.hpp index 300a5348c53e..f0f90e635f91 100644 --- a/dpnp/backend/extensions/blas/dot.hpp +++ b/dpnp/backend/extensions/blas/dot.hpp @@ -30,12 +30,15 @@ #include +#include + #include "dot_common.hpp" namespace dpnp::extensions::blas { namespace mkl_blas = oneapi::mkl::blas; namespace type_utils = dpnp::tensor::type_utils; +namespace py = pybind11; template static sycl::event dot_impl(sycl::queue &exec_q, @@ -58,6 +61,11 @@ static sycl::event dot_impl(sycl::queue &exec_q, sycl::event dot_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + + dot_event = mkl_blas::column_major::dot(exec_q, n, // size of the input vectors x, // Pointer to vector x. diff --git a/dpnp/backend/extensions/blas/dotc.hpp b/dpnp/backend/extensions/blas/dotc.hpp index 535f79156fb5..273ebfa032cd 100644 --- a/dpnp/backend/extensions/blas/dotc.hpp +++ b/dpnp/backend/extensions/blas/dotc.hpp @@ -30,12 +30,15 @@ #include +#include + #include "dot_common.hpp" namespace dpnp::extensions::blas { namespace mkl_blas = oneapi::mkl::blas; namespace type_utils = dpnp::tensor::type_utils; +namespace py = pybind11; template static sycl::event dotc_impl(sycl::queue &exec_q, @@ -58,6 +61,11 @@ static sycl::event dotc_impl(sycl::queue &exec_q, sycl::event dotc_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + + dotc_event = mkl_blas::column_major::dotc(exec_q, n, // size of the input vectors diff --git a/dpnp/backend/extensions/blas/dotu.hpp b/dpnp/backend/extensions/blas/dotu.hpp index aafe9721339f..485fdd33f938 100644 --- a/dpnp/backend/extensions/blas/dotu.hpp +++ b/dpnp/backend/extensions/blas/dotu.hpp @@ -30,12 +30,15 @@ #include +#include + #include "dot_common.hpp" namespace dpnp::extensions::blas { namespace mkl_blas = oneapi::mkl::blas; namespace type_utils = dpnp::tensor::type_utils; +namespace py = pybind11; template static sycl::event dotu_impl(sycl::queue &exec_q, @@ -58,6 +61,11 @@ static sycl::event dotu_impl(sycl::queue &exec_q, sycl::event dotu_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + + dotu_event = mkl_blas::column_major::dotu(exec_q, n, // size of the input vectors diff --git a/dpnp/backend/extensions/blas/gemm.cpp b/dpnp/backend/extensions/blas/gemm.cpp index bf91f5cfb0bf..d0901097c759 100644 --- a/dpnp/backend/extensions/blas/gemm.cpp +++ b/dpnp/backend/extensions/blas/gemm.cpp @@ -113,6 +113,11 @@ static sycl::event gemm_impl(sycl::queue &exec_q, c, ldc, deps); } }; + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + + gemm_event = gemm_func( exec_q, transA, // Defines the transpose operation for matrix A: diff --git a/dpnp/backend/extensions/blas/gemm_batch.cpp b/dpnp/backend/extensions/blas/gemm_batch.cpp index f1bc9f30b5dd..f3a0c009214e 100644 --- a/dpnp/backend/extensions/blas/gemm_batch.cpp +++ b/dpnp/backend/extensions/blas/gemm_batch.cpp @@ -125,6 +125,11 @@ static sycl::event gemm_batch_impl(sycl::queue &exec_q, strideb, beta, c, ldc, stridec, batch_size, deps); } }; + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + + gemm_batch_event = gemm_batch_func( exec_q, transA, // Defines the transpose operation for matrix A: diff --git a/dpnp/backend/extensions/lapack/geqrf.cpp b/dpnp/backend/extensions/lapack/geqrf.cpp index 083c408ea3d4..03e9c5060fdd 100644 --- a/dpnp/backend/extensions/lapack/geqrf.cpp +++ b/dpnp/backend/extensions/lapack/geqrf.cpp @@ -86,6 +86,10 @@ static sycl::event geqrf_impl(sycl::queue &exec_q, sycl::event geqrf_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + scratchpad = sycl::malloc_device(scratchpad_size, exec_q); geqrf_event = mkl_lapack::geqrf( diff --git a/dpnp/backend/extensions/lapack/geqrf_batch.cpp b/dpnp/backend/extensions/lapack/geqrf_batch.cpp index 6e62d4dde336..739e2a82b543 100644 --- a/dpnp/backend/extensions/lapack/geqrf_batch.cpp +++ b/dpnp/backend/extensions/lapack/geqrf_batch.cpp @@ -94,6 +94,10 @@ static sycl::event geqrf_batch_impl(sycl::queue &exec_q, sycl::event geqrf_batch_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + scratchpad = sycl::malloc_device(scratchpad_size, exec_q); geqrf_batch_event = mkl_lapack::geqrf_batch( diff --git a/dpnp/backend/extensions/lapack/gesv.cpp b/dpnp/backend/extensions/lapack/gesv.cpp index 072efd169c87..3bab2e5b35a3 100644 --- a/dpnp/backend/extensions/lapack/gesv.cpp +++ b/dpnp/backend/extensions/lapack/gesv.cpp @@ -112,6 +112,10 @@ static sycl::event gesv_impl(sycl::queue &exec_q, #if defined(USE_ONEMATH) sycl::event getrf_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + getrf_event = mkl_lapack::getrf( exec_q, n, // The order of the square matrix A (0 ≤ n). @@ -169,6 +173,10 @@ static sycl::event gesv_impl(sycl::queue &exec_q, } #else try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + comp_event = mkl_lapack::gesv( exec_q, n, // The order of the square matrix A diff --git a/dpnp/backend/extensions/lapack/gesvd.cpp b/dpnp/backend/extensions/lapack/gesvd.cpp index ebc6000671b0..02b388977144 100644 --- a/dpnp/backend/extensions/lapack/gesvd.cpp +++ b/dpnp/backend/extensions/lapack/gesvd.cpp @@ -99,6 +99,10 @@ static sycl::event gesvd_impl(sycl::queue &exec_q, sycl::event gesvd_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + gesvd_event = mkl_lapack::gesvd( exec_q, jobu, // Character specifying how to compute the matrix U: diff --git a/dpnp/backend/extensions/lapack/getrf.cpp b/dpnp/backend/extensions/lapack/getrf.cpp index 42021639d63f..c31213b38620 100644 --- a/dpnp/backend/extensions/lapack/getrf.cpp +++ b/dpnp/backend/extensions/lapack/getrf.cpp @@ -87,6 +87,10 @@ static sycl::event getrf_impl(sycl::queue &exec_q, sycl::event getrf_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + scratchpad = sycl::malloc_device(scratchpad_size, exec_q); getrf_event = mkl_lapack::getrf( diff --git a/dpnp/backend/extensions/lapack/getrf_batch.cpp b/dpnp/backend/extensions/lapack/getrf_batch.cpp index b4e30dae57d3..917794557378 100644 --- a/dpnp/backend/extensions/lapack/getrf_batch.cpp +++ b/dpnp/backend/extensions/lapack/getrf_batch.cpp @@ -96,6 +96,10 @@ static sycl::event getrf_batch_impl(sycl::queue &exec_q, sycl::event getrf_batch_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + scratchpad = sycl::malloc_device(scratchpad_size, exec_q); getrf_batch_event = mkl_lapack::getrf_batch( diff --git a/dpnp/backend/extensions/lapack/heevd.cpp b/dpnp/backend/extensions/lapack/heevd.cpp index 002a7ad2ef89..86ba97765884 100644 --- a/dpnp/backend/extensions/lapack/heevd.cpp +++ b/dpnp/backend/extensions/lapack/heevd.cpp @@ -42,6 +42,7 @@ namespace dpnp::extensions::lapack { +namespace py = pybind11; namespace mkl_lapack = oneapi::mkl::lapack; namespace type_utils = dpnp::tensor::type_utils; @@ -73,6 +74,10 @@ static sycl::event heevd_impl(sycl::queue &exec_q, sycl::event heevd_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + heevd_event = mkl_lapack::heevd( exec_q, jobz, // 'jobz == job::vec' means eigenvalues and eigenvectors are diff --git a/dpnp/backend/extensions/lapack/orgqr_batch.cpp b/dpnp/backend/extensions/lapack/orgqr_batch.cpp index 840ef82a09b6..0d996bef2b87 100644 --- a/dpnp/backend/extensions/lapack/orgqr_batch.cpp +++ b/dpnp/backend/extensions/lapack/orgqr_batch.cpp @@ -96,6 +96,10 @@ static sycl::event orgqr_batch_impl(sycl::queue &exec_q, sycl::event orgqr_batch_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + scratchpad = sycl::malloc_device(scratchpad_size, exec_q); orgqr_batch_event = mkl_lapack::orgqr_batch( diff --git a/dpnp/backend/extensions/lapack/potrf.cpp b/dpnp/backend/extensions/lapack/potrf.cpp index d7edc2d244f6..a14ce223458c 100644 --- a/dpnp/backend/extensions/lapack/potrf.cpp +++ b/dpnp/backend/extensions/lapack/potrf.cpp @@ -83,6 +83,10 @@ static sycl::event potrf_impl(sycl::queue &exec_q, sycl::event potrf_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + scratchpad = sycl::malloc_device(scratchpad_size, exec_q); potrf_event = mkl_lapack::potrf( diff --git a/dpnp/backend/extensions/lapack/potrf_batch.cpp b/dpnp/backend/extensions/lapack/potrf_batch.cpp index 58bdad462930..8193b32d4f05 100644 --- a/dpnp/backend/extensions/lapack/potrf_batch.cpp +++ b/dpnp/backend/extensions/lapack/potrf_batch.cpp @@ -90,6 +90,10 @@ static sycl::event potrf_batch_impl(sycl::queue &exec_q, sycl::event potrf_batch_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + scratchpad = sycl::malloc_device(scratchpad_size, exec_q); potrf_batch_event = mkl_lapack::potrf_batch( diff --git a/dpnp/backend/extensions/lapack/syevd.cpp b/dpnp/backend/extensions/lapack/syevd.cpp index 7280e415f271..f2be6214f60a 100644 --- a/dpnp/backend/extensions/lapack/syevd.cpp +++ b/dpnp/backend/extensions/lapack/syevd.cpp @@ -42,6 +42,7 @@ namespace dpnp::extensions::lapack { +namespace py = pybind11; namespace mkl_lapack = oneapi::mkl::lapack; namespace type_utils = dpnp::tensor::type_utils; @@ -73,6 +74,10 @@ static sycl::event syevd_impl(sycl::queue &exec_q, sycl::event syevd_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + syevd_event = mkl_lapack::syevd( exec_q, jobz, // 'jobz == job::vec' means eigenvalues and eigenvectors are diff --git a/dpnp/backend/extensions/lapack/ungqr.cpp b/dpnp/backend/extensions/lapack/ungqr.cpp index 7bea568f048e..42b14059176f 100644 --- a/dpnp/backend/extensions/lapack/ungqr.cpp +++ b/dpnp/backend/extensions/lapack/ungqr.cpp @@ -87,6 +87,10 @@ static sycl::event ungqr_impl(sycl::queue &exec_q, sycl::event ungqr_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + scratchpad = sycl::malloc_device(scratchpad_size, exec_q); ungqr_event = mkl_lapack::ungqr( diff --git a/dpnp/backend/extensions/lapack/ungqr_batch.cpp b/dpnp/backend/extensions/lapack/ungqr_batch.cpp index 6452a30e2a86..1ac7220064be 100644 --- a/dpnp/backend/extensions/lapack/ungqr_batch.cpp +++ b/dpnp/backend/extensions/lapack/ungqr_batch.cpp @@ -96,6 +96,10 @@ static sycl::event ungqr_batch_impl(sycl::queue &exec_q, sycl::event ungqr_batch_event; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + scratchpad = sycl::malloc_device(scratchpad_size, exec_q); ungqr_batch_event = mkl_lapack::ungqr_batch( diff --git a/dpnp/tests/test_blas_lapack_gil.py b/dpnp/tests/test_blas_lapack_gil.py new file mode 100644 index 000000000000..d017b3422a02 --- /dev/null +++ b/dpnp/tests/test_blas_lapack_gil.py @@ -0,0 +1,176 @@ +"""Blocking oneMKL calls in the BLAS/LAPACK extensions must release the GIL. + +Routines such as ``potrf``, ``getrf``, ``syevd`` and ``gesv`` block the calling +thread for the whole oneMKL call. Holding the GIL across that stalls every +other Python thread, including the host tasks dpnp queues to manage object +lifetimes, which reacquire the GIL. Five sibling files already wrap their call +in ``py::gil_scoped_release`` (gh-2850 did this for ``orgqr``); these tests +cover the rest. + +Progress of a competing thread is compared against ``SyclQueue.wait()``, which +is ``nogil`` and therefore the best rate achievable on the machine. +""" + +import sys +import threading +import time + +import dpctl +import pytest + +import dpnp + +from .helper import has_support_aspect64, is_gpu_device + +# Large enough that every routine here actually blocks. Smaller sizes stop +# discriminating: at n=512 cholesky returns asynchronously, and at n=1024 +# gesv blocks for only ~7 ms and passes even without the fix. +_SIZE = 2048 + +# A releasing call matches the nogil reference within a factor of ~2; a call +# holding the GIL measures ~0.00-0.04 of it. +_MIN_RATIO = 0.10 + +_BACKLOG = 2 # queued matmuls, so the measured call has something to wait on + + +class _Ticker: + """Counts how often a competing Python thread gets scheduled.""" + + def __enter__(self): + self.ticks = 0 + self._stop = False + self._thread = threading.Thread(target=self._run, daemon=True) + self._thread.start() + return self + + def _run(self): + while not self._stop: + self.ticks += 1 + time.sleep(0) + + def __exit__(self, *exc): + self._stop = True + self._thread.join(timeout=5) + return False + + +@pytest.mark.skipif(not is_gpu_device(), reason="requires a GPU device") +@pytest.mark.skipif(not has_support_aspect64(), reason="requires fp64 support") +class TestBlockingCallsReleaseGil: + @pytest.fixture(autouse=True) + def _switch_interval(self): + # Stop CPython handing the GIL over on its own timer, so the + # measurement reflects the extension rather than the interpreter. + previous = sys.getswitchinterval() + sys.setswitchinterval(0.001) + yield + sys.setswitchinterval(previous) + + @pytest.fixture + def mats(self): + spd = dpnp.eye(_SIZE, dtype="f8") * float(_SIZE) # positive definite + return spd, dpnp.eye(_SIZE, dtype="f8") + 0.1 + + def _assert_releases_gil(self, name, spd, fn): + queue = spd.sycl_queue + + def rate(ticker, func): + for _ in range(_BACKLOG): + dpnp.matmul(spd, spd) + ticker.ticks = 0 + start = time.perf_counter() + func() + elapsed_ms = 1000 * (time.perf_counter() - start) + queue.wait() + return ticker.ticks / max(elapsed_ms, 1e-3) + + fn() # warm up JIT + queue.wait() + + with _Ticker() as ticker: + measured = rate(ticker, fn) + reference = rate(ticker, queue.wait) + + assert reference > 0, "reference measurement produced no ticks" + ratio = measured / reference + assert ratio >= _MIN_RATIO, ( + f"{name} holds the GIL while blocking: {measured:.2f} ticks/ms vs " + f"{reference:.2f} for nogil queue.wait() (ratio {ratio:.3f}, need " + f">= {_MIN_RATIO}). The oneMKL call needs py::gil_scoped_release." + ) + + def test_potrf(self, mats): + spd, _ = mats + self._assert_releases_gil("potrf", spd, lambda: dpnp.linalg.cholesky(spd)) + + def test_getrf(self, mats): + spd, gen = mats + self._assert_releases_gil("getrf", spd, lambda: dpnp.linalg.det(gen)) + + def test_syevd(self, mats): + spd, _ = mats + self._assert_releases_gil("syevd", spd, lambda: dpnp.linalg.eigh(spd)) + + # No qr/geqrf case: dpnp.linalg.qr() also calls orgqr, already fixed in + # gh-2850, and that release alone keeps the competing thread running. The + # measurement would pass with or without geqrf's own release. + + def test_gesv(self, mats): + spd, gen = mats + rhs = dpnp.ones(_SIZE, dtype="f8") + self._assert_releases_gil( + "gesv", spd, lambda: dpnp.linalg.solve(gen, rhs) + ) + + +@pytest.mark.skipif(not is_gpu_device(), reason="requires a GPU device") +@pytest.mark.skipif(not has_support_aspect64(), reason="requires fp64 support") +def test_multithreaded_linalg_overlaps(): + """Two threads on two queues must run oneMKL work concurrently. + + Holding the GIL across the blocking call serializes them, so a + multithreaded application gets no scaling from a second queue. + """ + max_ratio = 0.85 + + queues = [ + dpctl.SyclQueue(dpctl.SyclDevice("gpu"), property="in_order") + for _ in range(2) + ] + mats = [ + dpnp.eye(_SIZE, dtype="f8", sycl_queue=q) * float(_SIZE) for q in queues + ] + + def work(idx): + for _ in range(2): + dpnp.linalg.eigh(mats[idx]) + queues[idx].wait() + + for idx in range(2): # warm up JIT + work(idx) + + previous = sys.getswitchinterval() + sys.setswitchinterval(0.001) + try: + start = time.perf_counter() + work(0) + work(1) + serial = time.perf_counter() - start + + start = time.perf_counter() + threads = [threading.Thread(target=work, args=(i,)) for i in range(2)] + for thread in threads: + thread.start() + for thread in threads: + thread.join() + parallel = time.perf_counter() - start + finally: + sys.setswitchinterval(previous) + + ratio = parallel / serial + assert ratio <= max_ratio, ( + f"two threads running eigh did not overlap: {parallel * 1000:.1f} ms " + f"parallel vs {serial * 1000:.1f} ms serial (ratio {ratio:.3f}, need " + f"<= {max_ratio}). The oneMKL call is holding the GIL." + ) From dcd255841bb2589a86f77ad8ae38fe52d43630aa Mon Sep 17 00:00:00 2001 From: Abhishek Bagusetty Date: Fri, 14 Aug 2026 13:18:37 -0500 Subject: [PATCH 3/6] Run GIL tests on CPU devices too, so they cover public CI --- dpnp/tests/test_blas_lapack_gil.py | 64 ++++++++++++------------------ 1 file changed, 26 insertions(+), 38 deletions(-) diff --git a/dpnp/tests/test_blas_lapack_gil.py b/dpnp/tests/test_blas_lapack_gil.py index d017b3422a02..cf933941e72b 100644 --- a/dpnp/tests/test_blas_lapack_gil.py +++ b/dpnp/tests/test_blas_lapack_gil.py @@ -1,12 +1,5 @@ """Blocking oneMKL calls in the BLAS/LAPACK extensions must release the GIL. -Routines such as ``potrf``, ``getrf``, ``syevd`` and ``gesv`` block the calling -thread for the whole oneMKL call. Holding the GIL across that stalls every -other Python thread, including the host tasks dpnp queues to manage object -lifetimes, which reacquire the GIL. Five sibling files already wrap their call -in ``py::gil_scoped_release`` (gh-2850 did this for ``orgqr``); these tests -cover the rest. - Progress of a competing thread is compared against ``SyclQueue.wait()``, which is ``nogil`` and therefore the best rate achievable on the machine. """ @@ -22,17 +15,16 @@ from .helper import has_support_aspect64, is_gpu_device -# Large enough that every routine here actually blocks. Smaller sizes stop -# discriminating: at n=512 cholesky returns asynchronously, and at n=1024 -# gesv blocks for only ~7 ms and passes even without the fix. +# Smaller sizes stop discriminating: the calls either stay asynchronous or +# block too briefly for a stable measurement. _SIZE = 2048 -# A releasing call matches the nogil reference within a factor of ~2; a call -# holding the GIL measures ~0.00-0.04 of it. _MIN_RATIO = 0.10 _BACKLOG = 2 # queued matmuls, so the measured call has something to wait on +_TRIALS = 5 # samples averaged per measurement + class _Ticker: """Counts how often a competing Python thread gets scheduled.""" @@ -55,13 +47,11 @@ def __exit__(self, *exc): return False -@pytest.mark.skipif(not is_gpu_device(), reason="requires a GPU device") @pytest.mark.skipif(not has_support_aspect64(), reason="requires fp64 support") class TestBlockingCallsReleaseGil: @pytest.fixture(autouse=True) def _switch_interval(self): - # Stop CPython handing the GIL over on its own timer, so the - # measurement reflects the extension rather than the interpreter. + # Stop CPython handing the GIL over on its own timer. previous = sys.getswitchinterval() sys.setswitchinterval(0.001) yield @@ -69,21 +59,25 @@ def _switch_interval(self): @pytest.fixture def mats(self): - spd = dpnp.eye(_SIZE, dtype="f8") * float(_SIZE) # positive definite + spd = dpnp.eye(_SIZE, dtype="f8") * float(_SIZE) return spd, dpnp.eye(_SIZE, dtype="f8") + 0.1 def _assert_releases_gil(self, name, spd, fn): queue = spd.sycl_queue def rate(ticker, func): - for _ in range(_BACKLOG): - dpnp.matmul(spd, spd) - ticker.ticks = 0 - start = time.perf_counter() - func() - elapsed_ms = 1000 * (time.perf_counter() - start) - queue.wait() - return ticker.ticks / max(elapsed_ms, 1e-3) + total_ticks = 0 + total_ms = 0.0 + for _ in range(_TRIALS): + for _ in range(_BACKLOG): + dpnp.matmul(spd, spd) + ticker.ticks = 0 + start = time.perf_counter() + func() + total_ms += 1000 * (time.perf_counter() - start) + total_ticks += ticker.ticks + queue.wait() + return total_ticks / max(total_ms, 1e-3) fn() # warm up JIT queue.wait() @@ -102,7 +96,9 @@ def rate(ticker, func): def test_potrf(self, mats): spd, _ = mats - self._assert_releases_gil("potrf", spd, lambda: dpnp.linalg.cholesky(spd)) + self._assert_releases_gil( + "potrf", spd, lambda: dpnp.linalg.cholesky(spd) + ) def test_getrf(self, mats): spd, gen = mats @@ -112,10 +108,6 @@ def test_syevd(self, mats): spd, _ = mats self._assert_releases_gil("syevd", spd, lambda: dpnp.linalg.eigh(spd)) - # No qr/geqrf case: dpnp.linalg.qr() also calls orgqr, already fixed in - # gh-2850, and that release alone keeps the competing thread running. The - # measurement would pass with or without geqrf's own release. - def test_gesv(self, mats): spd, gen = mats rhs = dpnp.ones(_SIZE, dtype="f8") @@ -124,20 +116,16 @@ def test_gesv(self, mats): ) +# GPU only: on a CPU device oneMKL already saturates every core, so two threads +# contend for the same hardware and the ratio stays ~1.0 either way. @pytest.mark.skipif(not is_gpu_device(), reason="requires a GPU device") @pytest.mark.skipif(not has_support_aspect64(), reason="requires fp64 support") def test_multithreaded_linalg_overlaps(): - """Two threads on two queues must run oneMKL work concurrently. - - Holding the GIL across the blocking call serializes them, so a - multithreaded application gets no scaling from a second queue. - """ + """Two threads on two queues must run oneMKL work concurrently.""" max_ratio = 0.85 - queues = [ - dpctl.SyclQueue(dpctl.SyclDevice("gpu"), property="in_order") - for _ in range(2) - ] + device = dpctl.select_default_device() + queues = [dpctl.SyclQueue(device, property="in_order") for _ in range(2)] mats = [ dpnp.eye(_SIZE, dtype="f8", sycl_queue=q) * float(_SIZE) for q in queues ] From 63556976367397c052640d6dc789db76181afe5e Mon Sep 17 00:00:00 2001 From: Abhishek Bagusetty Date: Fri, 14 Aug 2026 14:16:45 -0500 Subject: [PATCH 4/6] Drop stray blank line flagged by clang-format --- dpnp/backend/extensions/blas/dot.hpp | 1 - dpnp/backend/extensions/blas/dotc.hpp | 1 - dpnp/backend/extensions/blas/dotu.hpp | 1 - dpnp/backend/extensions/blas/gemm.cpp | 1 - dpnp/backend/extensions/blas/gemm_batch.cpp | 1 - 5 files changed, 5 deletions(-) diff --git a/dpnp/backend/extensions/blas/dot.hpp b/dpnp/backend/extensions/blas/dot.hpp index f0f90e635f91..755dd046d2e9 100644 --- a/dpnp/backend/extensions/blas/dot.hpp +++ b/dpnp/backend/extensions/blas/dot.hpp @@ -65,7 +65,6 @@ static sycl::event dot_impl(sycl::queue &exec_q, // to the same queue in OneMKL py::gil_scoped_release lock{}; - dot_event = mkl_blas::column_major::dot(exec_q, n, // size of the input vectors x, // Pointer to vector x. diff --git a/dpnp/backend/extensions/blas/dotc.hpp b/dpnp/backend/extensions/blas/dotc.hpp index 273ebfa032cd..051c85b7b16f 100644 --- a/dpnp/backend/extensions/blas/dotc.hpp +++ b/dpnp/backend/extensions/blas/dotc.hpp @@ -65,7 +65,6 @@ static sycl::event dotc_impl(sycl::queue &exec_q, // to the same queue in OneMKL py::gil_scoped_release lock{}; - dotc_event = mkl_blas::column_major::dotc(exec_q, n, // size of the input vectors diff --git a/dpnp/backend/extensions/blas/dotu.hpp b/dpnp/backend/extensions/blas/dotu.hpp index 485fdd33f938..f517d89bf211 100644 --- a/dpnp/backend/extensions/blas/dotu.hpp +++ b/dpnp/backend/extensions/blas/dotu.hpp @@ -65,7 +65,6 @@ static sycl::event dotu_impl(sycl::queue &exec_q, // to the same queue in OneMKL py::gil_scoped_release lock{}; - dotu_event = mkl_blas::column_major::dotu(exec_q, n, // size of the input vectors diff --git a/dpnp/backend/extensions/blas/gemm.cpp b/dpnp/backend/extensions/blas/gemm.cpp index d0901097c759..b7b39c3bc660 100644 --- a/dpnp/backend/extensions/blas/gemm.cpp +++ b/dpnp/backend/extensions/blas/gemm.cpp @@ -117,7 +117,6 @@ static sycl::event gemm_impl(sycl::queue &exec_q, // to the same queue in OneMKL py::gil_scoped_release lock{}; - gemm_event = gemm_func( exec_q, transA, // Defines the transpose operation for matrix A: diff --git a/dpnp/backend/extensions/blas/gemm_batch.cpp b/dpnp/backend/extensions/blas/gemm_batch.cpp index f3a0c009214e..d4e09162a6ed 100644 --- a/dpnp/backend/extensions/blas/gemm_batch.cpp +++ b/dpnp/backend/extensions/blas/gemm_batch.cpp @@ -129,7 +129,6 @@ static sycl::event gemm_batch_impl(sycl::queue &exec_q, // to the same queue in OneMKL py::gil_scoped_release lock{}; - gemm_batch_event = gemm_batch_func( exec_q, transA, // Defines the transpose operation for matrix A: From 56fb4d818e7a2c17a39bba68fa727697d1d6e44c Mon Sep 17 00:00:00 2001 From: Abhishek Bagusetty Date: Mon, 17 Aug 2026 11:45:10 -0500 Subject: [PATCH 5/6] Mark tests as slow except for 1 fastest method for functionality testing --- dpnp/tests/test_blas_lapack_gil.py | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/dpnp/tests/test_blas_lapack_gil.py b/dpnp/tests/test_blas_lapack_gil.py index cf933941e72b..5b8860372b12 100644 --- a/dpnp/tests/test_blas_lapack_gil.py +++ b/dpnp/tests/test_blas_lapack_gil.py @@ -100,14 +100,17 @@ def test_potrf(self, mats): "potrf", spd, lambda: dpnp.linalg.cholesky(spd) ) + @pytest.mark.slow def test_getrf(self, mats): spd, gen = mats self._assert_releases_gil("getrf", spd, lambda: dpnp.linalg.det(gen)) + @pytest.mark.slow def test_syevd(self, mats): spd, _ = mats self._assert_releases_gil("syevd", spd, lambda: dpnp.linalg.eigh(spd)) + @pytest.mark.slow def test_gesv(self, mats): spd, gen = mats rhs = dpnp.ones(_SIZE, dtype="f8") @@ -118,6 +121,7 @@ def test_gesv(self, mats): # GPU only: on a CPU device oneMKL already saturates every core, so two threads # contend for the same hardware and the ratio stays ~1.0 either way. +@pytest.mark.slow @pytest.mark.skipif(not is_gpu_device(), reason="requires a GPU device") @pytest.mark.skipif(not has_support_aspect64(), reason="requires fp64 support") def test_multithreaded_linalg_overlaps(): From 8dbabc7f0f7c68f2554bbe8c76eb8349d4abb601 Mon Sep 17 00:00:00 2001 From: Abhishek Bagusetty Date: Mon, 17 Aug 2026 15:56:13 -0500 Subject: [PATCH 6/6] fix the potrf test with slow decorator --- dpnp/tests/test_blas_lapack_gil.py | 1 + 1 file changed, 1 insertion(+) diff --git a/dpnp/tests/test_blas_lapack_gil.py b/dpnp/tests/test_blas_lapack_gil.py index 5b8860372b12..9aec868d0797 100644 --- a/dpnp/tests/test_blas_lapack_gil.py +++ b/dpnp/tests/test_blas_lapack_gil.py @@ -94,6 +94,7 @@ def rate(ticker, func): f">= {_MIN_RATIO}). The oneMKL call needs py::gil_scoped_release." ) + @pytest.mark.slow def test_potrf(self, mats): spd, _ = mats self._assert_releases_gil(