diff --git a/CHANGELOG.md b/CHANGELOG.md index 6eff2d59ac..f51c7785ed 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -11,6 +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` 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) 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, diff --git a/dpctl/_sycl_queue.pxd b/dpctl/_sycl_queue.pxd index c2102a52e9..ff5e7b8471 100644 --- a/dpctl/_sycl_queue.pxd +++ b/dpctl/_sycl_queue.pxd @@ -107,6 +107,10 @@ 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 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 a22586b71e..4289d25d1c 100644 --- a/dpctl/_sycl_queue.pyx +++ b/dpctl/_sycl_queue.pyx @@ -49,6 +49,8 @@ from ._backend cimport ( # noqa: E211 DPCTLQueue_MemAdvise, DPCTLQueue_Memcpy, DPCTLQueue_MemcpyWithEvents, + DPCTLQueue_Memset, + DPCTLQueue_MemsetWithEvents, DPCTLQueue_Prefetch, DPCTLQueue_SubmitBarrierForEvents, DPCTLQueue_SubmitNDRange, @@ -602,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. """ @@ -1594,6 +1625,103 @@ 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 DPCTLSyclEventRef ERef = NULL + + ERef = _memset_impl(self, mem, val, count, NULL, 0) + if (ERef is NULL): + raise RuntimeError( + "SyclQueue.memset operation encountered an error" + ) + with nogil: + 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 diff --git a/dpctl/tests/test_sycl_queue_memset.py b/dpctl/tests/test_sycl_queue_memset.py new file mode 100644 index 0000000000..f6545fbd47 --- /dev/null +++ b/dpctl/tests/test_sycl_queue_memset.py @@ -0,0 +1,166 @@ +# 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) + + +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]) diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_queue_interface.h b/libsyclinterface/include/syclinterface/dpctl_sycl_queue_interface.h index afd2b3240d..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 @@ -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, 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. + * @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..f9207b6cec 100644 --- a/libsyclinterface/source/dpctl_sycl_queue_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_queue_interface.cpp @@ -899,10 +899,46 @@ 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; + } +}; + +__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 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;