From 30834c0c898e0395f3def1d131cb497e11f3c563 Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 13 Aug 2026 06:58:02 -0700 Subject: [PATCH 01/10] Add SyclQueue.memset method --- dpctl/_sycl_queue.pxd | 1 + dpctl/_sycl_queue.pyx | 43 +++++++++++++++++++++++++++++++++++++++++++ 2 files changed, 44 insertions(+) diff --git a/dpctl/_sycl_queue.pxd b/dpctl/_sycl_queue.pxd index c2102a52e9..66d80178f5 100644 --- a/dpctl/_sycl_queue.pxd +++ b/dpctl/_sycl_queue.pxd @@ -107,6 +107,7 @@ cdef public api class SyclQueue (_SyclQueue) [ cpdef SyclEvent copy_async( self, dest, src, size_t count, list dEvents=*, str dtype=* ) + cpdef memset(self, mem, int val, size_t count=*) cpdef prefetch(self, ptr, size_t count=*) cpdef mem_advise(self, ptr, size_t count, int mem) cpdef SyclEvent submit_barrier(self, dependent_events=*) diff --git a/dpctl/_sycl_queue.pyx b/dpctl/_sycl_queue.pyx index a22586b71e..2c609e788a 100644 --- a/dpctl/_sycl_queue.pyx +++ b/dpctl/_sycl_queue.pyx @@ -49,6 +49,7 @@ from ._backend cimport ( # noqa: E211 DPCTLQueue_MemAdvise, DPCTLQueue_Memcpy, DPCTLQueue_MemcpyWithEvents, + DPCTLQueue_Memset, DPCTLQueue_Prefetch, DPCTLQueue_SubmitBarrierForEvents, DPCTLQueue_SubmitNDRange, @@ -1594,6 +1595,48 @@ cdef class SyclQueue(_SyclQueue): return SyclEvent._create(ERef) + cpdef memset(self, mem, int val, size_t count=0): + """Fill USM allocation ``mem`` with the byte value ``val`` and wait. + + Internally, this dispatches ``sycl::queue::memset``. The operation is + byte-wise: ``count`` bytes are set, each to the same value ``val``. + + Args: + mem: + Destination USM allocation, an instance of + :class:`dpctl.memory._Memory`. + val (int): + Value to fill ``mem`` with. Following ``sycl::queue::memset``, + it is interpreted as an ``unsigned char``, i.e. only the least + significant byte is used. + count (int, optional): + Number of bytes to fill. If ``0`` or greater than the size of + ``mem``, the whole allocation is filled. Default: ``0``. + + Raises: + TypeError: + If ``mem`` is not an instance of :class:`dpctl.memory._Memory`. + """ + cdef void *ptr + cdef DPCTLSyclEventRef ERef = NULL + + if isinstance(mem, _Memory): + ptr = (<_Memory>mem).get_data_ptr() + else: + raise TypeError("Parameter `mem` should have type _Memory") + + if (count <= 0 or count > mem.nbytes): + count = mem.nbytes + + ERef = DPCTLQueue_Memset(self._queue_ref, ptr, val, count) + if (ERef is NULL): + raise RuntimeError( + "SyclQueue.memset operation encountered an error" + ) + with nogil: + DPCTLEvent_Wait(ERef) + DPCTLEvent_Delete(ERef) + cpdef prefetch(self, mem, size_t count=0): cdef void *ptr cdef DPCTLSyclEventRef ERef = NULL From 4f08b2f2416837d4f89247ae3e7efe07af62ea7b Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 13 Aug 2026 06:59:13 -0700 Subject: [PATCH 02/10] Add tests for SyclQueue.memset --- dpctl/tests/test_sycl_queue_memset.py | 106 ++++++++++++++++++++++++++ 1 file changed, 106 insertions(+) create mode 100644 dpctl/tests/test_sycl_queue_memset.py diff --git a/dpctl/tests/test_sycl_queue_memset.py b/dpctl/tests/test_sycl_queue_memset.py new file mode 100644 index 0000000000..c97413a6b9 --- /dev/null +++ b/dpctl/tests/test_sycl_queue_memset.py @@ -0,0 +1,106 @@ +# Data Parallel Control (dpctl) +# +# Copyright 2026 Intel Corporation +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +"""Defines unit test cases for the SyclQueue.memset.""" + +import pytest + +import dpctl +import dpctl.memory + + +def _create_memory(q, nbytes=1024): + return dpctl.memory.MemoryUSMShared(nbytes, queue=q) + + +def test_memset_fills_whole_allocation(): + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Default constructor for SyclQueue failed") + nbytes = 256 + mobj = _create_memory(q, nbytes) + + q.memset(mobj, 0xAB) + + assert bytes(memoryview(mobj)) == b"\xab" * nbytes + + +def test_memset_zero_count_fills_whole_allocation(): + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Default constructor for SyclQueue failed") + nbytes = 64 + mobj = _create_memory(q, nbytes) + + q.memset(mobj, 0x01, 0) + + assert bytes(memoryview(mobj)) == b"\x01" * nbytes + + +def test_memset_partial_count(): + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Default constructor for SyclQueue failed") + nbytes = 16 + mobj = _create_memory(q, nbytes) + + # zero-out first, then fill only the leading 4 bytes + q.memset(mobj, 0x00) + q.memset(mobj, 0x7F, 4) + + assert bytes(memoryview(mobj)) == b"\x7f" * 4 + b"\x00" * (nbytes - 4) + + +def test_memset_count_clamped_to_allocation(): + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Default constructor for SyclQueue failed") + nbytes = 8 + mobj = _create_memory(q, nbytes) + + # requesting more bytes than allocated must not overrun; it is clamped + q.memset(mobj, 0x02, 4 * nbytes) + + assert bytes(memoryview(mobj)) == b"\x02" * nbytes + + +def test_memset_zero_value(): + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Default constructor for SyclQueue failed") + nbytes = 32 + mobj = _create_memory(q, nbytes) + + q.memset(mobj, 0xFF) + q.memset(mobj, 0) + + assert bytes(memoryview(mobj)) == b"\x00" * nbytes + + +def test_memset_type_error(): + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Default constructor for SyclQueue failed") + + with pytest.raises(TypeError) as cm: + q.memset(None, 1) + assert "_Memory" in str(cm.value) From a591bcf31aaf2a3d90c45ac1f2e9bfa535afcbcf Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 13 Aug 2026 07:04:10 -0700 Subject: [PATCH 03/10] Add gh-2361 to changelog --- CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 6eff2d59ac..a23ce6c315 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -11,6 +11,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 * Added a number of `sycl::device` info queries to `dpctl.SyclDevice` [gh-2324](https://github.com/IntelPython/dpctl/pull/2324) * Added `sycl::info::context` queries `sycl_platform`, `atomic_memory_order_capabilities`, `atomic_fence_order_capabilities`, `atomic_memory_scope_capabilities`, and `atomic_fence_scope_capabilities` to `dpctl.SyclContext` [gh-2354](https://github.com/IntelPython/dpctl/pull/2354) * Added `create_kernel_bundle_from_sycl_source`, `is_sycl_source_compilation_available`, and `dpctl.SyclDevice.can_compile` for supporting the creation of `dpctl.SyclKernelBundle`s from SYCL source strings via DPC++ extension, as well as corresponding C-API functions to support it [gh-2206](https://github.com/IntelPython/dpctl/pull/2206) +* Added `dpctl.SyclQueue.memset` method [gh-2361](https://github.com/IntelPython/dpctl/pull/2361) ### Changed * Bump minimum NumPy version to 1.26 [gh-2192](https://github.com/IntelPython/dpctl/pull/2192) From af2da99124a405c6082065a5293fe8d7a6529b4d Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 13 Aug 2026 08:17:08 -0700 Subject: [PATCH 04/10] Add DPCTLQueue_MemsetWithEvents C-API function --- .../dpctl_sycl_queue_interface.h | 23 ++++++++++++ .../source/dpctl_sycl_queue_interface.cpp | 36 +++++++++++++++++++ 2 files changed, 59 insertions(+) diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_queue_interface.h b/libsyclinterface/include/syclinterface/dpctl_sycl_queue_interface.h index afd2b3240d..6d35984abf 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_queue_interface.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_queue_interface.h @@ -481,6 +481,29 @@ DPCTLQueue_Memset(__dpctl_keep const DPCTLSyclQueueRef QRef, uint8_t Value, size_t Count); +/*! + * @brief C-API wrapper for ``sycl::queue::memset``. + * + * @param QRef An opaque pointer to the ``sycl::queue``. + * @param USMRef An USM pointer to the memory to fill. + * @param Value A value to fill. + * @param Count A number of uint8_t elements to fill. + * @param DepEvents A pointer to array of DPCTLSyclEventRef opaque + * pointers to dependent events. + * @param DepEventsCount A number of dependent events. + * @return An opaque pointer to the ``sycl::event`` returned by the + * ``sycl::queue::memset`` function. + * @ingroup QueueInterface + */ +DPCTL_API +__dpctl_give DPCTLSyclEventRef +DPCTLQueue_MemsetWithEvents(__dpctl_keep const DPCTLSyclQueueRef QRef, + void *USMRef, + uint8_t Value, + size_t Count, + __dpctl_keep const DPCTLSyclEventRef *DepEvents, + size_t DepEventsCount); + /*! * @brief C-API wrapper for ``sycl::queue::fill``. * diff --git a/libsyclinterface/source/dpctl_sycl_queue_interface.cpp b/libsyclinterface/source/dpctl_sycl_queue_interface.cpp index 575a3c13fa..c2f6701999 100644 --- a/libsyclinterface/source/dpctl_sycl_queue_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_queue_interface.cpp @@ -905,6 +905,42 @@ DPCTLQueue_Memset(__dpctl_keep const DPCTLSyclQueueRef QRef, } }; +__dpctl_give DPCTLSyclEventRef +DPCTLQueue_MemsetWithEvents(__dpctl_keep const DPCTLSyclQueueRef QRef, + void *USMRef, + uint8_t Value, + size_t Count, + const DPCTLSyclEventRef *DepEvents, + size_t DepEventsCount) +{ + event ev; + auto Q = unwrap(QRef); + if (Q && USMRef) { + try { + ev = Q->submit([&](handler &cgh) { + if (DepEvents) + for (size_t i = 0; i < DepEventsCount; ++i) { + event *ei = unwrap(DepEvents[i]); + if (ei) + cgh.depends_on(*ei); + } + + cgh.memset(USMRef, static_cast(Value), Count); + }); + } catch (const std::exception &ex) { + error_handler(ex, __FILE__, __func__, __LINE__); + return nullptr; + } + } + else { + error_handler("QRef or USMRef passed to memset were NULL.", __FILE__, + __func__, __LINE__); + return nullptr; + } + + return wrap(new event(ev)); +}; + __dpctl_give DPCTLSyclEventRef DPCTLQueue_Fill8(__dpctl_keep const DPCTLSyclQueueRef QRef, void *USMRef, From 4b197d5847640d66305f0cc936dd66aaaaf665d9 Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 13 Aug 2026 08:18:51 -0700 Subject: [PATCH 05/10] Add tests for DPCTLQueue_MemsetWithEvents --- .../tests/test_sycl_queue_interface.cpp | 44 +++++++++++++++++++ 1 file changed, 44 insertions(+) diff --git a/libsyclinterface/tests/test_sycl_queue_interface.cpp b/libsyclinterface/tests/test_sycl_queue_interface.cpp index b95aae3e76..e140008ff3 100644 --- a/libsyclinterface/tests/test_sycl_queue_interface.cpp +++ b/libsyclinterface/tests/test_sycl_queue_interface.cpp @@ -467,6 +467,10 @@ TEST(TestDPCTLSyclQueueInterface, CheckMemsetNullQRef) ASSERT_NO_FATAL_FAILURE(ERef = DPCTLQueue_Memset(QRef, p, val8, 1)); ASSERT_FALSE(bool(ERef)); + + ASSERT_NO_FATAL_FAILURE( + ERef = DPCTLQueue_MemsetWithEvents(QRef, p, val8, 1, NULL, 0)); + ASSERT_FALSE(bool(ERef)); } TEST_P(TestDPCTLQueueMemberFunctions, CheckMemset) @@ -534,6 +538,46 @@ TEST_P(TestDPCTLQueueMemberFunctions, CheckMemset2) delete[] host_arr; } +TEST_P(TestDPCTLQueueMemberFunctions, CheckMemsetWithEvents) +{ + DPCTLSyclUSMRef p = nullptr; + DPCTLSyclEventRef Memset_ERef = nullptr; + DPCTLSyclEventRef MemsetWithEvents_ERef = nullptr; + DPCTLSyclEventRef Memcpy_ERef = nullptr; + uint8_t val1 = 42; + uint8_t val2 = 73; + size_t nbytes = 256; + uint8_t *host_arr = new uint8_t[nbytes]; + + ASSERT_FALSE(host_arr == nullptr); + + ASSERT_NO_FATAL_FAILURE(p = DPCTLmalloc_device(nbytes, QRef)); + ASSERT_FALSE(p == nullptr); + + ASSERT_NO_FATAL_FAILURE( + Memset_ERef = DPCTLQueue_Memset(QRef, (void *)p, val1, nbytes)); + + ASSERT_NO_FATAL_FAILURE( + MemsetWithEvents_ERef = DPCTLQueue_MemsetWithEvents( + QRef, (void *)p, val2, nbytes, &Memset_ERef, 1)); + + ASSERT_NO_FATAL_FAILURE( + Memcpy_ERef = DPCTLQueue_MemcpyWithEvents(QRef, host_arr, p, nbytes, + &MemsetWithEvents_ERef, 1)); + ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Wait(Memcpy_ERef)); + + ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Delete(Memset_ERef)); + ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Delete(MemsetWithEvents_ERef)); + ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Delete(Memcpy_ERef)); + + ASSERT_NO_FATAL_FAILURE(DPCTLfree_with_queue(p, QRef)); + + for (size_t i = 0; i < nbytes; ++i) { + ASSERT_TRUE(host_arr[i] == val2); + } + delete[] host_arr; +} + TEST(TestDPCTLSyclQueueInterface, CheckFillNullQRef) { DPCTLSyclQueueRef QRef = nullptr; From 9d63c23e822d3da0d33f22d5abe3ec77e655ccf7 Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 13 Aug 2026 08:21:30 -0700 Subject: [PATCH 06/10] Add DPCTLQueue_MemsetWithEvents declaration to _backend.pxd --- dpctl/_backend.pxd | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index 21a301d94d..29da03ece7 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -679,6 +679,13 @@ cdef extern from "syclinterface/dpctl_sycl_queue_interface.h": void *Dest, int Val, size_t Count) + cdef DPCTLSyclEventRef DPCTLQueue_MemsetWithEvents( + const DPCTLSyclQueueRef Q, + void *Dest, + int Val, + size_t Count, + const DPCTLSyclEventRef *depEvents, + size_t depEventsCount) cdef DPCTLSyclEventRef DPCTLQueue_Prefetch( const DPCTLSyclQueueRef Q, const void *Src, From fd541c008a91c65e29d2445ee5ba24b265c090a0 Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 13 Aug 2026 08:22:56 -0700 Subject: [PATCH 07/10] Add SyclQueue.memset_async method --- dpctl/_sycl_queue.pxd | 3 ++ dpctl/_sycl_queue.pyx | 105 ++++++++++++++++++++++++++++++++++++++---- 2 files changed, 98 insertions(+), 10 deletions(-) diff --git a/dpctl/_sycl_queue.pxd b/dpctl/_sycl_queue.pxd index 66d80178f5..ff5e7b8471 100644 --- a/dpctl/_sycl_queue.pxd +++ b/dpctl/_sycl_queue.pxd @@ -108,6 +108,9 @@ cdef public api class SyclQueue (_SyclQueue) [ self, dest, src, size_t count, list dEvents=*, str dtype=* ) cpdef memset(self, mem, int val, size_t count=*) + cpdef SyclEvent memset_async( + self, mem, int val, size_t count=*, list dEvents=* + ) cpdef prefetch(self, ptr, size_t count=*) cpdef mem_advise(self, ptr, size_t count, int mem) cpdef SyclEvent submit_barrier(self, dependent_events=*) diff --git a/dpctl/_sycl_queue.pyx b/dpctl/_sycl_queue.pyx index 2c609e788a..4289d25d1c 100644 --- a/dpctl/_sycl_queue.pyx +++ b/dpctl/_sycl_queue.pyx @@ -50,6 +50,7 @@ from ._backend cimport ( # noqa: E211 DPCTLQueue_Memcpy, DPCTLQueue_MemcpyWithEvents, DPCTLQueue_Memset, + DPCTLQueue_MemsetWithEvents, DPCTLQueue_Prefetch, DPCTLQueue_SubmitBarrierForEvents, DPCTLQueue_SubmitNDRange, @@ -603,6 +604,35 @@ cdef DPCTLSyclEventRef _copy_impl( ) +cdef DPCTLSyclEventRef _memset_impl( + SyclQueue q, + object mem, + int val, + size_t count, + DPCTLSyclEventRef *dep_events, + size_t dep_events_count, +) except *: + cdef void *ptr = NULL + cdef DPCTLSyclEventRef ERef = NULL + + if isinstance(mem, _Memory): + ptr = (<_Memory>mem).get_data_ptr() + else: + raise TypeError("Parameter `mem` should have type _Memory") + + if count <= 0 or count > mem.nbytes: + count = mem.nbytes + + if dep_events_count == 0 or dep_events is NULL: + ERef = DPCTLQueue_Memset(q._queue_ref, ptr, val, count) + else: + ERef = DPCTLQueue_MemsetWithEvents( + q._queue_ref, ptr, val, count, dep_events, dep_events_count + ) + + return ERef + + cdef class _SyclQueue: """ Barebone data owner class used by SyclQueue. """ @@ -1617,18 +1647,9 @@ cdef class SyclQueue(_SyclQueue): TypeError: If ``mem`` is not an instance of :class:`dpctl.memory._Memory`. """ - cdef void *ptr cdef DPCTLSyclEventRef ERef = NULL - if isinstance(mem, _Memory): - ptr = (<_Memory>mem).get_data_ptr() - else: - raise TypeError("Parameter `mem` should have type _Memory") - - if (count <= 0 or count > mem.nbytes): - count = mem.nbytes - - ERef = DPCTLQueue_Memset(self._queue_ref, ptr, val, count) + ERef = _memset_impl(self, mem, val, count, NULL, 0) if (ERef is NULL): raise RuntimeError( "SyclQueue.memset operation encountered an error" @@ -1637,6 +1658,70 @@ cdef class SyclQueue(_SyclQueue): DPCTLEvent_Wait(ERef) DPCTLEvent_Delete(ERef) + cpdef SyclEvent memset_async( + self, mem, int val, size_t count=0, list dEvents=None + ): + """Fill USM allocation ``mem`` with the byte value ``val`` + asynchronously. + + Internally, this dispatches ``sycl::queue::memset``. The operation is + byte-wise: ``count`` bytes are set, each to the same value ``val``. + + Args: + mem: + Destination USM allocation, an instance of + :class:`dpctl.memory._Memory`. + val (int): + Value to fill ``mem`` with. Following ``sycl::queue::memset``, + it is interpreted as an ``unsigned char``, i.e. only the least + significant byte is used. + count (int, optional): + Number of bytes to fill. If ``0`` or greater than the size of + ``mem``, the whole allocation is filled. Default: ``0``. + dEvents (List[dpctl.SyclEvent], optional): + Events that this operation depends on. + + Returns: + dpctl.SyclEvent: + Event associated with the memset operation. + + Raises: + TypeError: + If ``mem`` is not an instance of :class:`dpctl.memory._Memory`, + or ``dEvents`` is not a sequence of :class:`dpctl.SyclEvent`. + """ + cdef DPCTLSyclEventRef ERef = NULL + cdef DPCTLSyclEventRef *depEvents = NULL + cdef size_t nDE = 0 + + if dEvents is None: + ERef = _memset_impl(self, mem, val, count, NULL, 0) + else: + nDE = len(dEvents) + depEvents = ( + malloc(nDE*sizeof(DPCTLSyclEventRef)) + ) + if depEvents is NULL: + raise MemoryError() + try: + for idx, de in enumerate(dEvents): + if isinstance(de, SyclEvent): + depEvents[idx] = (de).get_event_ref() + else: + raise TypeError( + "A sequence of dpctl.SyclEvent is expected" + ) + ERef = _memset_impl(self, mem, val, count, depEvents, nDE) + finally: + free(depEvents) + + if (ERef is NULL): + raise RuntimeError( + "SyclQueue.memset operation encountered an error" + ) + + return SyclEvent._create(ERef) + cpdef prefetch(self, mem, size_t count=0): cdef void *ptr cdef DPCTLSyclEventRef ERef = NULL From 9a3fbe17c200de42f378881b0ad01b04314837fb Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 13 Aug 2026 08:33:14 -0700 Subject: [PATCH 08/10] Fix typos and improve DPCTLQueue_Memset doc --- .../include/syclinterface/dpctl_sycl_queue_interface.h | 6 +++--- libsyclinterface/source/dpctl_sycl_queue_interface.cpp | 2 +- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_queue_interface.h b/libsyclinterface/include/syclinterface/dpctl_sycl_queue_interface.h index 6d35984abf..ee0b37888c 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_queue_interface.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_queue_interface.h @@ -468,10 +468,10 @@ __dpctl_give DPCTLSyclEventRef DPCTLQueue_SubmitBarrierForEvents( * * @param QRef An opaque pointer to the ``sycl::queue``. * @param USMRef An USM pointer to the memory to fill. - * @param Value A value to fill. + * @param Value A value to fill, interpreted as an unsigned char. * @param Count A number of uint8_t elements to fill. * @return An opaque pointer to the ``sycl::event`` returned by the - * ``sycl::queue::fill`` function. + * ``sycl::queue::memset`` function. * @ingroup QueueInterface */ DPCTL_API @@ -486,7 +486,7 @@ DPCTLQueue_Memset(__dpctl_keep const DPCTLSyclQueueRef QRef, * * @param QRef An opaque pointer to the ``sycl::queue``. * @param USMRef An USM pointer to the memory to fill. - * @param Value A value to fill. + * @param Value A value to fill, interpreted as an unsigned char. * @param Count A number of uint8_t elements to fill. * @param DepEvents A pointer to array of DPCTLSyclEventRef opaque * pointers to dependent events. diff --git a/libsyclinterface/source/dpctl_sycl_queue_interface.cpp b/libsyclinterface/source/dpctl_sycl_queue_interface.cpp index c2f6701999..f9207b6cec 100644 --- a/libsyclinterface/source/dpctl_sycl_queue_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_queue_interface.cpp @@ -899,7 +899,7 @@ DPCTLQueue_Memset(__dpctl_keep const DPCTLSyclQueueRef QRef, return wrap(new event(std::move(ev))); } else { - error_handler("QRef or USMRef passed to fill8 were NULL.", __FILE__, + error_handler("QRef or USMRef passed to memset were NULL.", __FILE__, __func__, __LINE__); return nullptr; } From c905349fc3fc4b94029796848e39788c8286a93e Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 13 Aug 2026 08:33:59 -0700 Subject: [PATCH 09/10] Add tests for SyclQueue.memset_async --- dpctl/tests/test_sycl_queue_memset.py | 60 +++++++++++++++++++++++++++ 1 file changed, 60 insertions(+) diff --git a/dpctl/tests/test_sycl_queue_memset.py b/dpctl/tests/test_sycl_queue_memset.py index c97413a6b9..f6545fbd47 100644 --- a/dpctl/tests/test_sycl_queue_memset.py +++ b/dpctl/tests/test_sycl_queue_memset.py @@ -104,3 +104,63 @@ def test_memset_type_error(): with pytest.raises(TypeError) as cm: q.memset(None, 1) assert "_Memory" in str(cm.value) + + +def test_memset_async(): + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Default constructor for SyclQueue failed") + nbytes = 64 + mobj = _create_memory(q, nbytes) + + e = q.memset_async(mobj, 0xAB) + assert isinstance(e, dpctl.SyclEvent) + e.wait() + + assert bytes(memoryview(mobj)) == b"\xab" * nbytes + + +def test_memset_async_with_dependent_events(): + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Default constructor for SyclQueue failed") + nbytes = 64 + mobj = _create_memory(q, nbytes) + + e1 = q.memset_async(mobj, 0x01) + e2 = q.memset_async(mobj, 0x02, nbytes, [e1]) + e2.wait() + + assert bytes(memoryview(mobj)) == b"\x02" * nbytes + + +def test_memset_async_partial_count(): + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Default constructor for SyclQueue failed") + nbytes = 16 + mobj = _create_memory(q, nbytes) + + q.memset(mobj, 0x00) + e = q.memset_async(mobj, 0x7F, 4) + e.wait() + + assert bytes(memoryview(mobj)) == b"\x7f" * 4 + b"\x00" * (nbytes - 4) + + +def test_memset_async_type_error(): + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Default constructor for SyclQueue failed") + mobj = _create_memory(q) + + with pytest.raises(TypeError) as cm: + q.memset_async(None, 1) + assert "_Memory" in str(cm.value) + + with pytest.raises(TypeError): + q.memset_async(mobj, 1, 0, [None]) From 8e1cfde7e3731e357e0bfa3de28d48f31035dcaf Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 13 Aug 2026 08:37:10 -0700 Subject: [PATCH 10/10] Update gh-2361 changelog entry --- CHANGELOG.md | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index a23ce6c315..f51c7785ed 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -11,7 +11,8 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 * Added a number of `sycl::device` info queries to `dpctl.SyclDevice` [gh-2324](https://github.com/IntelPython/dpctl/pull/2324) * Added `sycl::info::context` queries `sycl_platform`, `atomic_memory_order_capabilities`, `atomic_fence_order_capabilities`, `atomic_memory_scope_capabilities`, and `atomic_fence_scope_capabilities` to `dpctl.SyclContext` [gh-2354](https://github.com/IntelPython/dpctl/pull/2354) * Added `create_kernel_bundle_from_sycl_source`, `is_sycl_source_compilation_available`, and `dpctl.SyclDevice.can_compile` for supporting the creation of `dpctl.SyclKernelBundle`s from SYCL source strings via DPC++ extension, as well as corresponding C-API functions to support it [gh-2206](https://github.com/IntelPython/dpctl/pull/2206) -* Added `dpctl.SyclQueue.memset` method [gh-2361](https://github.com/IntelPython/dpctl/pull/2361) +* Added `dpctl.SyclQueue.memset` and `dpctl.SyclQueue.memset_async` methods [gh-2361](https://github.com/IntelPython/dpctl/pull/2361) +* Added `DPCTLQueue_MemsetWithEvents` C-API function to support `dpctl.SyclQueue.memset_async` [gh-2361](https://github.com/IntelPython/dpctl/pull/2361) ### Changed * Bump minimum NumPy version to 1.26 [gh-2192](https://github.com/IntelPython/dpctl/pull/2192)