From f34f4ca56e056fe0f1bff3741a8b876533ed15c2 Mon Sep 17 00:00:00 2001 From: abagusetty Date: Mon, 24 Aug 2026 09:47:26 -0500 Subject: [PATCH 1/4] Fix GIL for fft extensions --- CHANGELOG.md | 1 + dpnp/backend/extensions/fft/common.hpp | 8 +- dpnp/backend/extensions/fft/in_place.tpp | 4 + dpnp/backend/extensions/fft/out_of_place.tpp | 4 + dpnp/tests/test_fft_gil.py | 108 +++++++++++++++++++ 5 files changed, 124 insertions(+), 1 deletion(-) create mode 100644 dpnp/tests/test_fft_gil.py diff --git a/CHANGELOG.md b/CHANGELOG.md index 113b0060fdfd..4cb9073b79c7 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -85,6 +85,7 @@ This release is compatible with NumPy 2.5. * Fixed a crash in boolean-mask advanced indexing (`dpnp.ndarray` get/set item) when the selection is empty (e.g. a scalar `False` index that injects a length-0 axis) [#3019](https://github.com/IntelPython/dpnp/pull/3019) * 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) * Fixed `dpnp.repeat` raising an unclear `TypeError` for a nested sequence of `repeats` [#3024](https://github.com/IntelPython/dpnp/pull/3024) +* Released the GIL before the blocking OneMKL DFT calls in the FFT extension [#3040](https://github.com/IntelPython/dpnp/pull/3040) ### Security diff --git a/dpnp/backend/extensions/fft/common.hpp b/dpnp/backend/extensions/fft/common.hpp index b293f14f48a9..06177c0900f7 100644 --- a/dpnp/backend/extensions/fft/common.hpp +++ b/dpnp/backend/extensions/fft/common.hpp @@ -61,7 +61,13 @@ class DescriptorWrapper "device does not support double precision."); } - descr_.commit(q); + { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + + descr_.commit(q); + } queue_ptr_ = std::make_unique(q); } diff --git a/dpnp/backend/extensions/fft/in_place.tpp b/dpnp/backend/extensions/fft/in_place.tpp index fa2ce1e1988b..790960a71fac 100644 --- a/dpnp/backend/extensions/fft/in_place.tpp +++ b/dpnp/backend/extensions/fft/in_place.tpp @@ -90,6 +90,10 @@ std::pair bool is_exception_caught = false; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + if (is_forward) { fft_event = mkl_dft::compute_forward(descr.get_descriptor(), in_out_ptr, depends); diff --git a/dpnp/backend/extensions/fft/out_of_place.tpp b/dpnp/backend/extensions/fft/out_of_place.tpp index 8ceb5f48c28b..bf40dca9b38b 100644 --- a/dpnp/backend/extensions/fft/out_of_place.tpp +++ b/dpnp/backend/extensions/fft/out_of_place.tpp @@ -143,6 +143,10 @@ std::pair bool is_exception_caught = false; try { + // Release GIL to avoid serialization of host task submissions + // to the same queue in OneMKL + py::gil_scoped_release lock{}; + if (is_forward) { using ScaleT_in = typename ScaleType::type_in; using ScaleT_out = typename ScaleType::type_out; diff --git a/dpnp/tests/test_fft_gil.py b/dpnp/tests/test_fft_gil.py new file mode 100644 index 000000000000..acea1d491633 --- /dev/null +++ b/dpnp/tests/test_fft_gil.py @@ -0,0 +1,108 @@ +"""Blocking oneMKL calls in the FFT extension must release the GIL. + +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 pytest + +import dpnp + +from .helper import has_support_aspect64 + +# Smaller sizes stop discriminating: the calls either stay asynchronous or +# block too briefly for a stable measurement. +_BATCH = 512 +_SIZE = 4096 + +_MIN_RATIO = 0.10 + +_BACKLOG = 2 # queued transforms, 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.""" + + 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 has_support_aspect64(), reason="requires fp64 support") +class TestFftReleasesGil: + @pytest.fixture(autouse=True) + def _switch_interval(self): + # Stop CPython handing the GIL over on its own timer. + previous = sys.getswitchinterval() + sys.setswitchinterval(0.001) + yield + sys.setswitchinterval(previous) + + def _assert_releases_gil(self, name, a, fn): + queue = a.sycl_queue + + def rate(ticker, func): + total_ticks = 0 + total_ms = 0.0 + for _ in range(_TRIALS): + for _ in range(_BACKLOG): + dpnp.fft.fft(a) + 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() + + 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." + ) + + @pytest.mark.slow + def test_fft_out_of_place(self): + # a complex input is passed to oneMKL as is, so the transform is + # computed out-of-place + a = dpnp.ones((_BATCH, _SIZE), dtype="c16") + self._assert_releases_gil( + "compute_fft_out_of_place", a, lambda: dpnp.fft.fft(a) + ) + + @pytest.mark.slow + def test_fft_in_place(self): + # a real input is copied to a complex array first, which allows the + # transform to be computed in-place + a = dpnp.ones((_BATCH, _SIZE), dtype="f8") + self._assert_releases_gil( + "compute_fft_in_place", a, lambda: dpnp.fft.fft(a) + ) From f68351148d8f383fe897b55a31baad4ba96ebf86 Mon Sep 17 00:00:00 2001 From: Abhishek Bagusetty Date: Tue, 25 Aug 2026 06:35:32 -0500 Subject: [PATCH 2/4] fix: call get_data() before releasing the GIL in out-of-place FFT usm_ndarray::get_data() goes through the Cython api function UsmNDArray_GetData, which calls into the Python C-API. Calling it after py::gil_scoped_release aborts the interpreter with Fatal Python error: PyThreadState_Get: the function must be called with the GIL held [...] the GIL is released which crashed every test worker that reached an out-of-place transform. Read both pointers before the release, matching in_place.tpp. The GIL test used SyclQueue.wait() as its reference, but by the time it ran the queue was already drained, so the reference measured no ticks at all. Compare against a nogil time.sleep() of the same duration instead. Co-Authored-By: Claude Opus 5 (1M context) --- dpnp/backend/extensions/fft/in_place.tpp | 2 ++ dpnp/backend/extensions/fft/out_of_place.tpp | 14 ++++++--- dpnp/tests/test_fft_gil.py | 31 ++++++++++++-------- 3 files changed, 31 insertions(+), 16 deletions(-) diff --git a/dpnp/backend/extensions/fft/in_place.tpp b/dpnp/backend/extensions/fft/in_place.tpp index 790960a71fac..3c6cf91fc152 100644 --- a/dpnp/backend/extensions/fft/in_place.tpp +++ b/dpnp/backend/extensions/fft/in_place.tpp @@ -83,6 +83,8 @@ std::pair // in-place is only used for c2c FFT at this time, passing true or false is // indifferent using ScaleT = typename ScaleType::type_in; + // get_data() calls into the Python C-API and so must be called while the + // GIL is still held ScaleT *in_out_ptr = in_out.get_data(); sycl::event fft_event = {}; diff --git a/dpnp/backend/extensions/fft/out_of_place.tpp b/dpnp/backend/extensions/fft/out_of_place.tpp index bf40dca9b38b..dcde7d78b10b 100644 --- a/dpnp/backend/extensions/fft/out_of_place.tpp +++ b/dpnp/backend/extensions/fft/out_of_place.tpp @@ -138,6 +138,12 @@ std::pair dpnp::tensor::validation::CheckWritable::throw_if_not_writable(out); dpnp::tensor::validation::AmpleMemory::throw_if_not_ample(out, n_elems); + // the input and output types depend on is_forward, so the untyped + // pointers are cast to the expected type below. get_data() calls into the + // Python C-API and so must be called while the GIL is still held + char *in_data = in.get_data(); + char *out_data = out.get_data(); + sycl::event fft_event = {}; std::stringstream error_msg; bool is_exception_caught = false; @@ -150,16 +156,16 @@ std::pair if (is_forward) { using ScaleT_in = typename ScaleType::type_in; using ScaleT_out = typename ScaleType::type_out; - ScaleT_in *in_ptr = in.get_data(); - ScaleT_out *out_ptr = out.get_data(); + ScaleT_in *in_ptr = reinterpret_cast(in_data); + ScaleT_out *out_ptr = reinterpret_cast(out_data); fft_event = mkl_dft::compute_forward(descr.get_descriptor(), in_ptr, out_ptr, depends); } else { using ScaleT_in = typename ScaleType::type_in; using ScaleT_out = typename ScaleType::type_out; - ScaleT_in *in_ptr = in.get_data(); - ScaleT_out *out_ptr = out.get_data(); + ScaleT_in *in_ptr = reinterpret_cast(in_data); + ScaleT_out *out_ptr = reinterpret_cast(out_data); fft_event = mkl_dft::compute_backward(descr.get_descriptor(), in_ptr, out_ptr, depends); } diff --git a/dpnp/tests/test_fft_gil.py b/dpnp/tests/test_fft_gil.py index acea1d491633..db4e8565085a 100644 --- a/dpnp/tests/test_fft_gil.py +++ b/dpnp/tests/test_fft_gil.py @@ -1,7 +1,8 @@ """Blocking oneMKL calls in the FFT extension must release the GIL. -Progress of a competing thread is compared against ``SyclQueue.wait()``, which -is ``nogil`` and therefore the best rate achievable on the machine. +Progress of a competing Python thread is measured while the FFT call blocks, +and compared against ``time.sleep()`` of the same duration, which is ``nogil`` +and therefore the best rate achievable on the machine. """ import sys @@ -19,7 +20,9 @@ _BATCH = 512 _SIZE = 4096 -_MIN_RATIO = 0.10 +# A call that holds the GIL measures around 0.09 of the reference rate, one +# that releases it around 0.3. The threshold sits between the two. +_MIN_RATIO = 0.18 _BACKLOG = 2 # queued transforms, so the measured call has something to wait on @@ -53,40 +56,44 @@ class TestFftReleasesGil: def _switch_interval(self): # Stop CPython handing the GIL over on its own timer. previous = sys.getswitchinterval() - sys.setswitchinterval(0.001) + sys.setswitchinterval(0.0005) yield sys.setswitchinterval(previous) def _assert_releases_gil(self, name, a, fn): queue = a.sycl_queue - def rate(ticker, func): + def measure(ticker, func): total_ticks = 0 - total_ms = 0.0 + total_s = 0.0 for _ in range(_TRIALS): for _ in range(_BACKLOG): dpnp.fft.fft(a) ticker.ticks = 0 start = time.perf_counter() func() - total_ms += 1000 * (time.perf_counter() - start) + total_s += time.perf_counter() - start total_ticks += ticker.ticks queue.wait() - return total_ticks / max(total_ms, 1e-3) + # ticks per millisecond, and the average duration of a single call + return total_ticks / max(1000 * total_s, 1e-3), total_s / _TRIALS fn() # warm up JIT queue.wait() with _Ticker() as ticker: - measured = rate(ticker, fn) - reference = rate(ticker, queue.wait) + measured, duration = measure(ticker, fn) + # time.sleep() is nogil, so blocking in it for the same amount of + # time gives the best tick rate the machine can produce + reference, _ = measure(ticker, lambda: time.sleep(duration)) 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." + f"{reference:.2f} for a nogil sleep of the same duration (ratio " + f"{ratio:.3f}, need >= {_MIN_RATIO}). The oneMKL call needs " + f"py::gil_scoped_release." ) @pytest.mark.slow From 595c93368c4f8ef22238d926b4473a17f4e7e055 Mon Sep 17 00:00:00 2001 From: Abhishek Bagusetty <59661409+abagusetty@users.noreply.github.com> Date: Tue, 25 Aug 2026 14:06:57 -0500 Subject: [PATCH 3/4] Update dpnp/backend/extensions/fft/out_of_place.tpp Co-authored-by: Anton <100830759+antonwolfy@users.noreply.github.com> --- dpnp/backend/extensions/fft/out_of_place.tpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/dpnp/backend/extensions/fft/out_of_place.tpp b/dpnp/backend/extensions/fft/out_of_place.tpp index dcde7d78b10b..1cda62944126 100644 --- a/dpnp/backend/extensions/fft/out_of_place.tpp +++ b/dpnp/backend/extensions/fft/out_of_place.tpp @@ -140,7 +140,8 @@ std::pair // the input and output types depend on is_forward, so the untyped // pointers are cast to the expected type below. get_data() calls into the - // Python C-API and so must be called while the GIL is still held + // get_data() calls into the Python C-API and so must be called while the + // GIL is still held char *in_data = in.get_data(); char *out_data = out.get_data(); From 43733c153040c0bf112727f6311e03940cd7cbc7 Mon Sep 17 00:00:00 2001 From: Abhishek Bagusetty <59661409+abagusetty@users.noreply.github.com> Date: Tue, 25 Aug 2026 14:23:54 -0500 Subject: [PATCH 4/4] Update dpnp/backend/extensions/fft/out_of_place.tpp Co-authored-by: ndgrigorian <46709016+ndgrigorian@users.noreply.github.com> --- dpnp/backend/extensions/fft/out_of_place.tpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/dpnp/backend/extensions/fft/out_of_place.tpp b/dpnp/backend/extensions/fft/out_of_place.tpp index 1cda62944126..dcde7d78b10b 100644 --- a/dpnp/backend/extensions/fft/out_of_place.tpp +++ b/dpnp/backend/extensions/fft/out_of_place.tpp @@ -140,8 +140,7 @@ std::pair // the input and output types depend on is_forward, so the untyped // pointers are cast to the expected type below. get_data() calls into the - // get_data() calls into the Python C-API and so must be called while the - // GIL is still held + // Python C-API and so must be called while the GIL is still held char *in_data = in.get_data(); char *out_data = out.get_data();