Add dpctl.SyclQueue.memset() method

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)

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,

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=*)

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 = <void*>(<_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(<SyclQueue>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(<SyclQueue>self, mem, val, count, NULL, 0)
+        else:
+            nDE = len(dEvents)
+            depEvents = (
+                <DPCTLSyclEventRef*>malloc(nDE*sizeof(DPCTLSyclEventRef))
+            )
+            if depEvents is NULL:
+                raise MemoryError()
+            try:
+                for idx, de in enumerate(dEvents):
+                    if isinstance(de, SyclEvent):
+                        depEvents[idx] = (<SyclEvent>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

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])

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``.
  *