diff --git a/CHANGELOG.md b/CHANGELOG.md index e8ef5cf884..ea649bd635 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.fill` and `dpctl.SyclQueue.fill_async` methods [gh-2365](https://github.com/IntelPython/dpctl/pull/2365) +* Added `DPCTLQueue_Fill8/16/32/64/128WithEvents` C-API functions to support `dpctl.SyclQueue.fill_async` [gh-2365](https://github.com/IntelPython/dpctl/pull/2365) ### 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..0d6a51326e 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -21,7 +21,7 @@ types defined by dpctl's C API. """ -from libc.stdint cimport int64_t, uint32_t, uint64_t +from libc.stdint cimport int64_t, uint8_t, uint16_t, uint32_t, uint64_t from libcpp cimport bool @@ -679,6 +679,66 @@ cdef extern from "syclinterface/dpctl_sycl_queue_interface.h": void *Dest, int Val, size_t Count) + cdef DPCTLSyclEventRef DPCTLQueue_Fill8( + const DPCTLSyclQueueRef Q, + void *Dest, + uint8_t Val, + size_t Count) + cdef DPCTLSyclEventRef DPCTLQueue_Fill8WithEvents( + const DPCTLSyclQueueRef Q, + void *Dest, + uint8_t Val, + size_t Count, + const DPCTLSyclEventRef *depEvents, + size_t depEventsCount) + cdef DPCTLSyclEventRef DPCTLQueue_Fill16( + const DPCTLSyclQueueRef Q, + void *Dest, + uint16_t Val, + size_t Count) + cdef DPCTLSyclEventRef DPCTLQueue_Fill16WithEvents( + const DPCTLSyclQueueRef Q, + void *Dest, + uint16_t Val, + size_t Count, + const DPCTLSyclEventRef *depEvents, + size_t depEventsCount) + cdef DPCTLSyclEventRef DPCTLQueue_Fill32( + const DPCTLSyclQueueRef Q, + void *Dest, + uint32_t Val, + size_t Count) + cdef DPCTLSyclEventRef DPCTLQueue_Fill32WithEvents( + const DPCTLSyclQueueRef Q, + void *Dest, + uint32_t Val, + size_t Count, + const DPCTLSyclEventRef *depEvents, + size_t depEventsCount) + cdef DPCTLSyclEventRef DPCTLQueue_Fill64( + const DPCTLSyclQueueRef Q, + void *Dest, + uint64_t Val, + size_t Count) + cdef DPCTLSyclEventRef DPCTLQueue_Fill64WithEvents( + const DPCTLSyclQueueRef Q, + void *Dest, + uint64_t Val, + size_t Count, + const DPCTLSyclEventRef *depEvents, + size_t depEventsCount) + cdef DPCTLSyclEventRef DPCTLQueue_Fill128( + const DPCTLSyclQueueRef Q, + void *Dest, + uint64_t *Val, + size_t Count) + cdef DPCTLSyclEventRef DPCTLQueue_Fill128WithEvents( + const DPCTLSyclQueueRef Q, + void *Dest, + uint64_t *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..ef0bba2284 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 fill(self, dest, value, size_t count, str dtype=*) + cpdef SyclEvent fill_async( + self, dest, value, size_t count, list dEvents=*, str dtype=* + ) 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..84e00da8fa 100644 --- a/dpctl/_sycl_queue.pyx +++ b/dpctl/_sycl_queue.pyx @@ -40,6 +40,16 @@ from ._backend cimport ( # noqa: E211 DPCTLQueue_CopyDataWithEvents, DPCTLQueue_Create, DPCTLQueue_Delete, + DPCTLQueue_Fill8, + DPCTLQueue_Fill8WithEvents, + DPCTLQueue_Fill16, + DPCTLQueue_Fill16WithEvents, + DPCTLQueue_Fill32, + DPCTLQueue_Fill32WithEvents, + DPCTLQueue_Fill64, + DPCTLQueue_Fill64WithEvents, + DPCTLQueue_Fill128, + DPCTLQueue_Fill128WithEvents, DPCTLQueue_GetBackend, DPCTLQueue_GetContext, DPCTLQueue_GetDevice, @@ -86,10 +96,13 @@ from cpython.buffer cimport ( PyObject_GetBuffer, ) from cpython.ref cimport Py_INCREF, PyObject +from libc.stdint cimport uint8_t, uint16_t, uint32_t, uint64_t from libc.stdlib cimport free, malloc import collections.abc import logging +import struct +import sys cdef extern from "_host_task_util.hpp": @@ -602,6 +615,128 @@ cdef DPCTLSyclEventRef _copy_impl( ) +# struct format per dtype used to pack the fill value into its byte pattern. +# The pattern's length is the element size. +# Complex types are packed as (real, imag). +_fill_dtype_formats = { + "i1": "b", + "u1": "B", + "i2": "h", + "u2": "H", + "i4": "i", + "u4": "I", + "f4": "f", + "i8": "q", + "u8": "Q", + "f8": "d", + "c8": "ff", + "c16": "dd", +} + + +cdef bytes _pack_fill_pattern(object value, str dtype): + """ + Return the native-byte-order byte pattern of ``value`` interpreted as a + single element of ``dtype``. ``sycl::queue::fill`` replicates this pattern + across the destination allocation. + """ + cdef str fmt + if dtype not in _fill_dtype_formats: + raise ValueError( + f"Unrecognized dtype '{dtype}'. Expected one of: " + "i1, u1, i2, u2, i4, u4, i8, u8, f4, f8, c8, c16" + ) + fmt = _fill_dtype_formats[dtype] + try: + if dtype == "c8" or dtype == "c16": + # complex: pack real then imag + cval = complex(value) + return struct.pack("=" + fmt, cval.real, cval.imag) + return struct.pack("=" + fmt, value) + except (struct.error, TypeError, ValueError) as e: + raise ValueError( + f"Value {value!r} cannot be represented as dtype '{dtype}': {e}" + ) + + +cdef DPCTLSyclEventRef _fill_impl( + SyclQueue q, + object dst, + object value, + size_t count, + str dtype, + DPCTLSyclEventRef *dep_events, + size_t dep_events_count +) except *: + cdef void *c_dst_ptr = NULL + cdef DPCTLSyclEventRef ERef = NULL + cdef uint64_t val128[2] + cdef bytes pattern = _pack_fill_pattern(value, dtype) + cdef size_t element_size = len(pattern) + cdef bint with_events = dep_events is not NULL and dep_events_count > 0 + cdef uint8_t v8 + cdef uint16_t v16 + cdef uint32_t v32 + cdef uint64_t v64 + + if isinstance(dst, _Memory): + c_dst_ptr = (<_Memory>dst).get_data_ptr() + else: + raise TypeError( + "Parameter `dest` should have type `dpctl.memory._Memory`" + ) + + if element_size == 1: + v8 = int.from_bytes(pattern, sys.byteorder) + if with_events: + ERef = DPCTLQueue_Fill8WithEvents( + q._queue_ref, c_dst_ptr, v8, count, + dep_events, dep_events_count + ) + else: + ERef = DPCTLQueue_Fill8(q._queue_ref, c_dst_ptr, v8, count) + elif element_size == 2: + v16 = int.from_bytes(pattern, sys.byteorder) + if with_events: + ERef = DPCTLQueue_Fill16WithEvents( + q._queue_ref, c_dst_ptr, v16, count, + dep_events, dep_events_count + ) + else: + ERef = DPCTLQueue_Fill16(q._queue_ref, c_dst_ptr, v16, count) + elif element_size == 4: + v32 = int.from_bytes(pattern, sys.byteorder) + if with_events: + ERef = DPCTLQueue_Fill32WithEvents( + q._queue_ref, c_dst_ptr, v32, count, + dep_events, dep_events_count + ) + else: + ERef = DPCTLQueue_Fill32(q._queue_ref, c_dst_ptr, v32, count) + elif element_size == 8: + v64 = int.from_bytes(pattern, sys.byteorder) + if with_events: + ERef = DPCTLQueue_Fill64WithEvents( + q._queue_ref, c_dst_ptr, v64, count, + dep_events, dep_events_count + ) + else: + ERef = DPCTLQueue_Fill64(q._queue_ref, c_dst_ptr, v64, count) + else: + # 128-bit pattern is passed as two uint64 + val128[0] = int.from_bytes(pattern[:8], sys.byteorder) + val128[1] = int.from_bytes(pattern[8:], sys.byteorder) + if with_events: + ERef = DPCTLQueue_Fill128WithEvents( + q._queue_ref, c_dst_ptr, val128, count, + dep_events, dep_events_count + ) + else: + ERef = DPCTLQueue_Fill128(q._queue_ref, c_dst_ptr, val128, count) + + return ERef + + cdef class _SyclQueue: """ Barebone data owner class used by SyclQueue. """ @@ -1594,6 +1729,135 @@ cdef class SyclQueue(_SyclQueue): return SyclEvent._create(ERef) + cpdef fill(self, dest, value, size_t count, str dtype="u1"): + """Fill ``dest`` with ``count`` copies of ``value`` and wait. + + Internally, this dispatches ``sycl::queue::fill``, which sets each of + the ``count`` elements of ``dest`` to ``value``. The number of bytes + written is ``count`` multiplied by the size of ``dtype``. The default + ``dtype`` of ``"u1"`` (a single byte) makes the default a byte-wise + fill, analogous to ``memset``. + + This is a synchronizing variant corresponding to + :meth:`dpctl.SyclQueue.fill_async`. + + Args: + dest (dpctl.memory._Memory): + Destination USM allocation to fill. + value (int, float, complex): + Value used to fill ``dest``. It is reinterpreted according to + ``dtype``: integral ``dtype`` values expect a Python integer, + ``"f4"`` and ``"f8"`` expect a Python float, and ``"c8"`` and + ``"c16"`` expect a Python complex. + count (int): + Number of elements to fill. + dtype (str, optional): + Data type string of the fill elements. Determines the element + size and how ``value`` is interpreted. + Defaults to ``"u1"`` (one byte per element). + Supported types: i1, u1, i2, u2, i4, u4, i8, u8, f4, f8, c8, + c16. + + Raises: + TypeError: + If ``dest`` is not a ``dpctl.memory._Memory`` USM allocation. + ValueError: + If ``dtype`` is unrecognized, or ``value`` cannot be + represented as ``dtype``. + RuntimeError: + If the fill operation encountered an error. + """ + cdef DPCTLSyclEventRef ERef = NULL + + ERef = _fill_impl(self, dest, value, count, dtype, NULL, 0) + if (ERef is NULL): + raise RuntimeError( + "SyclQueue.fill operation encountered an error" + ) + with nogil: + DPCTLEvent_Wait(ERef) + DPCTLEvent_Delete(ERef) + + cpdef SyclEvent fill_async( + self, dest, value, size_t count, list dEvents=None, str dtype="u1" + ): + """Fill ``dest`` with ``count`` copies of ``value`` asynchronously. + + Internally, this dispatches ``sycl::queue::fill``, which sets each of + the ``count`` elements of ``dest`` to ``value``. The number of bytes + written is ``count`` multiplied by the size of ``dtype``. The default + ``dtype`` of ``"u1"`` (a single byte) makes the default a byte-wise + fill, analogous to ``memset``. + + Args: + dest (dpctl.memory._Memory): + Destination USM allocation to fill. + value (int, float, complex): + Value used to fill ``dest``. It is reinterpreted according to + ``dtype``: integral ``dtype`` values expect a Python integer, + ``"f4"`` and ``"f8"`` expect a Python float, and ``"c8"`` and + ``"c16"`` expect a Python complex. + count (int): + Number of elements to fill. + dEvents (List[dpctl.SyclEvent], optional): + Events that this fill depends on. + dtype (str, optional): + Data type string of the fill elements. Determines the element + size and how ``value`` is interpreted. + Defaults to ``"u1"`` (one byte per element). + Supported types: i1, u1, i2, u2, i4, u4, i8, u8, f4, f8, c8, + c16. + + Returns: + dpctl.SyclEvent: + Event associated with the fill operation. + + Raises: + TypeError: + If ``dest`` is not a ``dpctl.memory._Memory`` USM allocation, + or ``dEvents`` is not a sequence of :class:`dpctl.SyclEvent`. + ValueError: + If ``dtype`` is unrecognized, or ``value`` cannot be + represented as ``dtype``. + RuntimeError: + If the fill operation encountered an error. + """ + cdef DPCTLSyclEventRef ERef = NULL + cdef DPCTLSyclEventRef *depEvents = NULL + cdef size_t nDE = 0 + + if dEvents is None: + ERef = _fill_impl( + self, dest, value, count, dtype, 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 = _fill_impl( + self, dest, value, count, dtype, depEvents, nDE + ) + finally: + free(depEvents) + + if (ERef is NULL): + raise RuntimeError( + "SyclQueue.fill_async 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_fill.py b/dpctl/tests/test_sycl_queue_fill.py new file mode 100644 index 0000000000..da14956486 --- /dev/null +++ b/dpctl/tests/test_sycl_queue_fill.py @@ -0,0 +1,299 @@ +# 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.fill.""" + +import numpy as np +import pytest + +import dpctl +import dpctl.memory + +# Maps a dtype string to the NumPy dtype used to build the expected result. +_dtype_to_np = { + "i1": np.int8, + "u1": np.uint8, + "i2": np.int16, + "u2": np.uint16, + "i4": np.int32, + "u4": np.uint32, + "f4": np.float32, + "i8": np.int64, + "u8": np.uint64, + "f8": np.float64, + "c8": np.complex64, + "c16": np.complex128, +} + +# A representative fill value per dtype, chosen to exercise all bytes and, +# for signed types, the sign bit. +_fill_values = { + "i1": -7, + "u1": 0xAB, + "i2": -1234, + "u2": 0xABCD, + "i4": -123456, + "u4": 0x0BADC0DE, + "f4": 3.5, + "i8": -1234567890123, + "u8": 0x0123456789ABCDEF, + "f8": 2.718281828459045, + "c8": complex(1.5, -2.5), + "c16": complex(3.14159, -2.71828), +} + + +def _read_back(mem, nbytes): + """Return the first ``nbytes`` bytes of a USM allocation as ``bytes``.""" + if isinstance(mem, dpctl.memory.MemoryUSMDevice): + result = bytearray(nbytes) + mem.copy_to_host(result) + return bytes(result) + return memoryview(mem)[:nbytes].tobytes() + + +@pytest.mark.parametrize("dtype", list(_dtype_to_np)) +@pytest.mark.parametrize( + "usm_type", + [ + lambda n, q: dpctl.memory.MemoryUSMShared(n, queue=q), + lambda n, q: dpctl.memory.MemoryUSMHost(n, queue=q), + lambda n, q: dpctl.memory.MemoryUSMDevice(n, queue=q), + ], + ids=["shared", "host", "device"], +) +def test_fill_with_dtype_valid(dtype, usm_type): + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Default constructor for SyclQueue failed") + + np_dt = _dtype_to_np[dtype] + value = _fill_values[dtype] + num_elements = 16 + element_size = np.dtype(np_dt).itemsize + nbytes = num_elements * element_size + + mem = usm_type(nbytes, q) + q.fill(mem, value, num_elements, dtype=dtype) + + expected = np.full(num_elements, value, dtype=np_dt).tobytes() + assert _read_back(mem, nbytes) == expected + + +def test_fill_default_dtype_is_bytewise(): + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Default constructor for SyclQueue failed") + + nbytes = 64 + mem = dpctl.memory.MemoryUSMShared(nbytes, queue=q) + + q.fill(mem, 0, nbytes) + assert memoryview(mem).tobytes() == b"\x00" * nbytes + + q.fill(mem, 0xFF, nbytes) + assert memoryview(mem).tobytes() == b"\xff" * nbytes + + +def test_fill_signed_negative_value(): + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Default constructor for SyclQueue failed") + + num_elements = 8 + mem = dpctl.memory.MemoryUSMShared(num_elements, queue=q) + + q.fill(mem, -1, num_elements, dtype="i1") + assert memoryview(mem).tobytes() == b"\xff" * num_elements + + +def test_fill_async(): + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Default constructor for SyclQueue failed") + + num_elements = 32 + mem = dpctl.memory.MemoryUSMShared(num_elements, queue=q) + + e = q.fill_async(mem, 0xAB, num_elements) + assert isinstance(e, dpctl.SyclEvent) + e.wait() + + assert memoryview(mem).tobytes() == b"\xab" * num_elements + + +def test_fill_async_with_dtype(): + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Default constructor for SyclQueue failed") + + value = 1.5 + num_elements = 10 + nbytes = num_elements * 8 + + mem = dpctl.memory.MemoryUSMShared(nbytes, queue=q) + e = q.fill_async(mem, value, num_elements, dtype="f8") + e.wait() + + expected = np.full(num_elements, value, dtype=np.float64).tobytes() + assert memoryview(mem)[:nbytes].tobytes() == expected + + +def test_fill_async_with_dep_events(): + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Default constructor for SyclQueue failed") + + num_elements = 64 + mem = dpctl.memory.MemoryUSMShared(num_elements, queue=q) + + # The second fill depends on the first, so its value wins deterministically + # even on an out-of-order queue. + e1 = q.fill_async(mem, 0x11, num_elements) + e2 = q.fill_async(mem, 0x22, num_elements, dEvents=[e1]) + e2.wait() + + assert memoryview(mem).tobytes() == b"\x22" * num_elements + + +def test_fill_async_bad_dep_events(): + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Default constructor for SyclQueue failed") + + mem = dpctl.memory.MemoryUSMShared(16, queue=q) + + with pytest.raises(TypeError): + q.fill_async(mem, 0, 16, dEvents=[None]) + + +@pytest.mark.parametrize("dtype", ["c8", "c16"]) +def test_fill_complex_from_real_scalar(dtype): + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Default constructor for SyclQueue failed") + + np_dt = _dtype_to_np[dtype] + num_elements = 8 + nbytes = num_elements * np.dtype(np_dt).itemsize + + mem = dpctl.memory.MemoryUSMShared(nbytes, queue=q) + q.fill(mem, 2.0, num_elements, dtype=dtype) + + expected = np.full(num_elements, complex(2.0, 0.0), dtype=np_dt).tobytes() + assert memoryview(mem)[:nbytes].tobytes() == expected + + +@pytest.mark.parametrize( + "dtype,element_size", + [ + ("i2", 2), + ("i4", 4), + ("f8", 8), + ("u8", 8), + ], +) +def test_fill_count_is_in_elements(dtype, element_size): + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Default constructor for SyclQueue failed") + + num_elements = 8 + nbytes = num_elements * element_size + + mem = dpctl.memory.MemoryUSMShared(nbytes, queue=q) + + # Seed the whole allocation with a known sentinel to verify the untouched + # tail is left intact. + mv = memoryview(mem) + for i in range(nbytes): + mv[i] = 0xAA + + # Filling half the elements writes exactly half the bytes. + q.fill(mem, 1, num_elements // 2, dtype=dtype) + + half_bytes = (num_elements // 2) * element_size + expected_head = np.full( + num_elements // 2, 1, dtype=_dtype_to_np[dtype] + ).tobytes() + assert mv[:half_bytes].tobytes() == expected_head + assert mv[half_bytes:nbytes].tobytes() == b"\xaa" * (nbytes - half_bytes) + + +def test_fill_with_invalid_dtype(): + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Default constructor for SyclQueue failed") + + mem = dpctl.memory.MemoryUSMShared(64, queue=q) + + for bad_dtype in ["i3", "f16", "u16", "x4", "42", "float", ""]: + with pytest.raises(ValueError) as cm: + q.fill(mem, 0, 8, dtype=bad_dtype) + assert ( + "dtype" in str(cm.value).lower() + or "unrecognized" in str(cm.value).lower() + ) + + +@pytest.mark.parametrize( + "dtype,value", + [ + ("u1", 256), + ("u1", -1), + ("i1", 128), + ("u2", 1 << 16), + ("i4", 1.5), + ("u8", 1 << 64), + ], +) +def test_fill_value_out_of_range(dtype, value): + """fill raises ValueError when value cannot be represented as dtype.""" + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Default constructor for SyclQueue failed") + + mem = dpctl.memory.MemoryUSMShared(64, queue=q) + + with pytest.raises(ValueError): + q.fill(mem, value, 8, dtype=dtype) + + +def test_fill_type_error(): + """fill raises TypeError when ``dest`` is not a USM allocation.""" + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Default constructor for SyclQueue failed") + + with pytest.raises(TypeError) as cm: + q.fill(None, 0, 4) + assert "_Memory" in str(cm.value) + + with pytest.raises(TypeError) as cm: + q.fill(bytearray(16), 0, 4) + assert "_Memory" in str(cm.value) diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_queue_interface.h b/libsyclinterface/include/syclinterface/dpctl_sycl_queue_interface.h index afd2b3240d..366abd7b07 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_queue_interface.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_queue_interface.h @@ -499,6 +499,29 @@ DPCTLQueue_Fill8(__dpctl_keep const DPCTLSyclQueueRef QRef, uint8_t Value, size_t Count); +/*! + * @brief C-API wrapper for ``sycl::queue::fill``. + * + * @param QRef An opaque pointer to the ``sycl::queue``. + * @param USMRef An USM pointer to the memory to fill. + * @param Value A uint8_t 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::fill`` function. + * @ingroup QueueInterface + */ +DPCTL_API +__dpctl_give DPCTLSyclEventRef +DPCTLQueue_Fill8WithEvents(__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``. * @@ -517,6 +540,29 @@ DPCTLQueue_Fill16(__dpctl_keep const DPCTLSyclQueueRef QRef, uint16_t Value, size_t Count); +/*! + * @brief C-API wrapper for ``sycl::queue::fill``. + * + * @param QRef An opaque pointer to the ``sycl::queue``. + * @param USMRef An USM pointer to the memory to fill. + * @param Value A uint16_t value to fill. + * @param Count A number of uint16_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::fill`` function. + * @ingroup QueueInterface + */ +DPCTL_API +__dpctl_give DPCTLSyclEventRef +DPCTLQueue_Fill16WithEvents(__dpctl_keep const DPCTLSyclQueueRef QRef, + void *USMRef, + uint16_t Value, + size_t Count, + __dpctl_keep const DPCTLSyclEventRef *DepEvents, + size_t DepEventsCount); + /*! * @brief C-API wrapper for ``sycl::queue::fill``. * @@ -535,6 +581,29 @@ DPCTLQueue_Fill32(__dpctl_keep const DPCTLSyclQueueRef QRef, uint32_t Value, size_t Count); +/*! + * @brief C-API wrapper for ``sycl::queue::fill``. + * + * @param QRef An opaque pointer to the ``sycl::queue``. + * @param USMRef An USM pointer to the memory to fill. + * @param Value A uint32_t value to fill. + * @param Count A number of uint32_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::fill`` function. + * @ingroup QueueInterface + */ +DPCTL_API +__dpctl_give DPCTLSyclEventRef +DPCTLQueue_Fill32WithEvents(__dpctl_keep const DPCTLSyclQueueRef QRef, + void *USMRef, + uint32_t Value, + size_t Count, + __dpctl_keep const DPCTLSyclEventRef *DepEvents, + size_t DepEventsCount); + /*! * @brief C-API wrapper for ``sycl::queue::fill``. * @@ -553,6 +622,29 @@ DPCTLQueue_Fill64(__dpctl_keep const DPCTLSyclQueueRef QRef, uint64_t Value, size_t Count); +/*! + * @brief C-API wrapper for ``sycl::queue::fill``. + * + * @param QRef An opaque pointer to the ``sycl::queue``. + * @param USMRef An USM pointer to the memory to fill. + * @param Value A uint64_t value to fill. + * @param Count A number of uint64_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::fill`` function. + * @ingroup QueueInterface + */ +DPCTL_API +__dpctl_give DPCTLSyclEventRef +DPCTLQueue_Fill64WithEvents(__dpctl_keep const DPCTLSyclQueueRef QRef, + void *USMRef, + uint64_t Value, + size_t Count, + __dpctl_keep const DPCTLSyclEventRef *DepEvents, + size_t DepEventsCount); + /*! * @brief C-API wrapper for ``sycl::queue::fill``. * @@ -572,4 +664,28 @@ DPCTLQueue_Fill128(__dpctl_keep const DPCTLSyclQueueRef QRef, uint64_t *Value, size_t Count); +/*! + * @brief C-API wrapper for ``sycl::queue::fill``. + * + * @param QRef An opaque pointer to the ``sycl::queue``. + * @param USMRef An USM pointer to the memory to fill. + * @param Value A pointer to uint64_t array of 2 elements with value + * to fill. + * @param Count A number of 128-bit 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::fill`` function. + * @ingroup QueueInterface + */ +DPCTL_API +__dpctl_give DPCTLSyclEventRef +DPCTLQueue_Fill128WithEvents(__dpctl_keep const DPCTLSyclQueueRef QRef, + void *USMRef, + uint64_t *Value, + size_t Count, + __dpctl_keep const DPCTLSyclEventRef *DepEvents, + size_t DepEventsCount); + DPCTL_C_EXTERN_C_END diff --git a/libsyclinterface/source/dpctl_sycl_queue_interface.cpp b/libsyclinterface/source/dpctl_sycl_queue_interface.cpp index 575a3c13fa..1b9b5374ff 100644 --- a/libsyclinterface/source/dpctl_sycl_queue_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_queue_interface.cpp @@ -929,6 +929,41 @@ DPCTLQueue_Fill8(__dpctl_keep const DPCTLSyclQueueRef QRef, } } +__dpctl_give DPCTLSyclEventRef +DPCTLQueue_Fill8WithEvents(__dpctl_keep const DPCTLSyclQueueRef QRef, + void *USMRef, + uint8_t Value, + size_t Count, + const DPCTLSyclEventRef *DepEvents, + size_t DepEventsCount) +{ + auto Q = unwrap(QRef); + if (Q && USMRef) { + sycl::event ev; + try { + std::vector dep_events; + if (DepEvents) { + dep_events.reserve(DepEventsCount); + for (size_t i = 0; i < DepEventsCount; ++i) { + event *ei = unwrap(DepEvents[i]); + if (ei) + dep_events.push_back(*ei); + } + } + ev = Q->fill(USMRef, Value, Count, dep_events); + } catch (std::exception const &e) { + error_handler(e, __FILE__, __func__, __LINE__); + return nullptr; + } + return wrap(new event(std::move(ev))); + } + else { + error_handler("QRef or USMRef passed to fill8 were NULL.", __FILE__, + __func__, __LINE__); + return nullptr; + } +} + __dpctl_give DPCTLSyclEventRef DPCTLQueue_Fill16(__dpctl_keep const DPCTLSyclQueueRef QRef, void *USMRef, @@ -953,6 +988,41 @@ DPCTLQueue_Fill16(__dpctl_keep const DPCTLSyclQueueRef QRef, } } +__dpctl_give DPCTLSyclEventRef +DPCTLQueue_Fill16WithEvents(__dpctl_keep const DPCTLSyclQueueRef QRef, + void *USMRef, + uint16_t Value, + size_t Count, + const DPCTLSyclEventRef *DepEvents, + size_t DepEventsCount) +{ + auto Q = unwrap(QRef); + if (Q && USMRef) { + sycl::event ev; + try { + std::vector dep_events; + if (DepEvents) { + dep_events.reserve(DepEventsCount); + for (size_t i = 0; i < DepEventsCount; ++i) { + event *ei = unwrap(DepEvents[i]); + if (ei) + dep_events.push_back(*ei); + } + } + ev = Q->fill(USMRef, Value, Count, dep_events); + } catch (std::exception const &e) { + error_handler(e, __FILE__, __func__, __LINE__); + return nullptr; + } + return wrap(new event(std::move(ev))); + } + else { + error_handler("QRef or USMRef passed to fill16 were NULL.", __FILE__, + __func__, __LINE__); + return nullptr; + } +} + __dpctl_give DPCTLSyclEventRef DPCTLQueue_Fill32(__dpctl_keep const DPCTLSyclQueueRef QRef, void *USMRef, @@ -977,6 +1047,41 @@ DPCTLQueue_Fill32(__dpctl_keep const DPCTLSyclQueueRef QRef, } } +__dpctl_give DPCTLSyclEventRef +DPCTLQueue_Fill32WithEvents(__dpctl_keep const DPCTLSyclQueueRef QRef, + void *USMRef, + uint32_t Value, + size_t Count, + const DPCTLSyclEventRef *DepEvents, + size_t DepEventsCount) +{ + auto Q = unwrap(QRef); + if (Q && USMRef) { + sycl::event ev; + try { + std::vector dep_events; + if (DepEvents) { + dep_events.reserve(DepEventsCount); + for (size_t i = 0; i < DepEventsCount; ++i) { + event *ei = unwrap(DepEvents[i]); + if (ei) + dep_events.push_back(*ei); + } + } + ev = Q->fill(USMRef, Value, Count, dep_events); + } catch (std::exception const &e) { + error_handler(e, __FILE__, __func__, __LINE__); + return nullptr; + } + return wrap(new event(std::move(ev))); + } + else { + error_handler("QRef or USMRef passed to fill32 were NULL.", __FILE__, + __func__, __LINE__); + return nullptr; + } +} + __dpctl_give DPCTLSyclEventRef DPCTLQueue_Fill64(__dpctl_keep const DPCTLSyclQueueRef QRef, void *USMRef, @@ -1001,6 +1106,41 @@ DPCTLQueue_Fill64(__dpctl_keep const DPCTLSyclQueueRef QRef, } } +__dpctl_give DPCTLSyclEventRef +DPCTLQueue_Fill64WithEvents(__dpctl_keep const DPCTLSyclQueueRef QRef, + void *USMRef, + uint64_t Value, + size_t Count, + const DPCTLSyclEventRef *DepEvents, + size_t DepEventsCount) +{ + auto Q = unwrap(QRef); + if (Q && USMRef) { + sycl::event ev; + try { + std::vector dep_events; + if (DepEvents) { + dep_events.reserve(DepEventsCount); + for (size_t i = 0; i < DepEventsCount; ++i) { + event *ei = unwrap(DepEvents[i]); + if (ei) + dep_events.push_back(*ei); + } + } + ev = Q->fill(USMRef, Value, Count, dep_events); + } catch (std::exception const &e) { + error_handler(e, __FILE__, __func__, __LINE__); + return nullptr; + } + return wrap(new event(std::move(ev))); + } + else { + error_handler("QRef or USMRef passed to fill64 were NULL.", __FILE__, + __func__, __LINE__); + return nullptr; + } +} + __dpctl_give DPCTLSyclEventRef DPCTLQueue_Fill128(__dpctl_keep const DPCTLSyclQueueRef QRef, void *USMRef, @@ -1027,3 +1167,41 @@ DPCTLQueue_Fill128(__dpctl_keep const DPCTLSyclQueueRef QRef, return nullptr; } } + +__dpctl_give DPCTLSyclEventRef +DPCTLQueue_Fill128WithEvents(__dpctl_keep const DPCTLSyclQueueRef QRef, + void *USMRef, + uint64_t *Value, + size_t Count, + const DPCTLSyclEventRef *DepEvents, + size_t DepEventsCount) +{ + auto Q = unwrap(QRef); + if (Q && USMRef) { + sycl::event ev; + try { + std::vector dep_events; + if (DepEvents) { + dep_events.reserve(DepEventsCount); + for (size_t i = 0; i < DepEventsCount; ++i) { + event *ei = unwrap(DepEvents[i]); + if (ei) + dep_events.push_back(*ei); + } + } + complexNumber Val; + Val.real = Value[0]; + Val.imag = Value[1]; + ev = Q->fill(USMRef, Val, Count, dep_events); + } catch (std::exception const &e) { + error_handler(e, __FILE__, __func__, __LINE__); + return nullptr; + } + return wrap(new event(std::move(ev))); + } + else { + error_handler("QRef or USMRef passed to fill128 were NULL.", __FILE__, + __func__, __LINE__); + return nullptr; + } +} diff --git a/libsyclinterface/tests/test_sycl_queue_interface.cpp b/libsyclinterface/tests/test_sycl_queue_interface.cpp index b95aae3e76..a4ef26952c 100644 --- a/libsyclinterface/tests/test_sycl_queue_interface.cpp +++ b/libsyclinterface/tests/test_sycl_queue_interface.cpp @@ -559,6 +559,26 @@ TEST(TestDPCTLSyclQueueInterface, CheckFillNullQRef) ASSERT_NO_FATAL_FAILURE(ERef = DPCTLQueue_Fill128(QRef, p, val128, 1)); ASSERT_FALSE(bool(ERef)); + + ASSERT_NO_FATAL_FAILURE( + ERef = DPCTLQueue_Fill8WithEvents(QRef, p, val8, 1, nullptr, 0)); + ASSERT_FALSE(bool(ERef)); + + ASSERT_NO_FATAL_FAILURE( + ERef = DPCTLQueue_Fill16WithEvents(QRef, p, val16, 1, nullptr, 0)); + ASSERT_FALSE(bool(ERef)); + + ASSERT_NO_FATAL_FAILURE( + ERef = DPCTLQueue_Fill32WithEvents(QRef, p, val32, 1, nullptr, 0)); + ASSERT_FALSE(bool(ERef)); + + ASSERT_NO_FATAL_FAILURE( + ERef = DPCTLQueue_Fill64WithEvents(QRef, p, val64, 1, nullptr, 0)); + ASSERT_FALSE(bool(ERef)); + + ASSERT_NO_FATAL_FAILURE( + ERef = DPCTLQueue_Fill128WithEvents(QRef, p, val128, 1, nullptr, 0)); + ASSERT_FALSE(bool(ERef)); } TEST_P(TestDPCTLQueueMemberFunctions, CheckFill8) @@ -751,6 +771,210 @@ TEST_P(TestDPCTLQueueMemberFunctions, CheckFill128) delete[] host_arr; } +TEST_P(TestDPCTLQueueMemberFunctions, CheckFill8WithEvents) +{ + using T = uint8_t; + DPCTLSyclUSMRef p = nullptr; + DPCTLSyclEventRef FillDep_ERef = nullptr; + DPCTLSyclEventRef Fill_ERef = nullptr; + DPCTLSyclEventRef Memcpy_ERef = nullptr; + T val = static_cast(0xB); + size_t nelems = 256; + T *host_arr = new T[nelems]; + size_t nbytes = nelems * sizeof(T); + + ASSERT_FALSE(host_arr == nullptr); + + ASSERT_NO_FATAL_FAILURE(p = DPCTLmalloc_device(nbytes, QRef)); + ASSERT_FALSE(p == nullptr); + + ASSERT_NO_FATAL_FAILURE(FillDep_ERef = + DPCTLQueue_Fill8(QRef, (void *)p, 0, nelems)); + + ASSERT_NO_FATAL_FAILURE( + Fill_ERef = DPCTLQueue_Fill8WithEvents(QRef, (void *)p, val, nelems, + &FillDep_ERef, 1)); + + ASSERT_NO_FATAL_FAILURE(Memcpy_ERef = DPCTLQueue_MemcpyWithEvents( + QRef, host_arr, p, nbytes, &Fill_ERef, 1)); + ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Wait(Memcpy_ERef)); + + ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Delete(FillDep_ERef)); + ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Delete(Fill_ERef)); + ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Delete(Memcpy_ERef)); + + ASSERT_NO_FATAL_FAILURE(DPCTLfree_with_queue(p, QRef)); + + for (size_t i = 0; i < nelems; ++i) { + ASSERT_TRUE(host_arr[i] == val); + } + delete[] host_arr; +} + +TEST_P(TestDPCTLQueueMemberFunctions, CheckFill16WithEvents) +{ + using T = uint16_t; + DPCTLSyclUSMRef p = nullptr; + DPCTLSyclEventRef FillDep_ERef = nullptr; + DPCTLSyclEventRef Fill_ERef = nullptr; + DPCTLSyclEventRef Memcpy_ERef = nullptr; + T val = static_cast(0xAB); + size_t nelems = 256; + T *host_arr = new T[nelems]; + size_t nbytes = nelems * sizeof(T); + + ASSERT_FALSE(host_arr == nullptr); + + ASSERT_NO_FATAL_FAILURE(p = DPCTLmalloc_device(nbytes, QRef)); + ASSERT_FALSE(p == nullptr); + + ASSERT_NO_FATAL_FAILURE(FillDep_ERef = + DPCTLQueue_Fill16(QRef, (void *)p, 0, nelems)); + + ASSERT_NO_FATAL_FAILURE( + Fill_ERef = DPCTLQueue_Fill16WithEvents(QRef, (void *)p, val, nelems, + &FillDep_ERef, 1)); + + ASSERT_NO_FATAL_FAILURE(Memcpy_ERef = DPCTLQueue_MemcpyWithEvents( + QRef, host_arr, p, nbytes, &Fill_ERef, 1)); + ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Wait(Memcpy_ERef)); + + ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Delete(FillDep_ERef)); + ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Delete(Fill_ERef)); + ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Delete(Memcpy_ERef)); + + ASSERT_NO_FATAL_FAILURE(DPCTLfree_with_queue(p, QRef)); + + for (size_t i = 0; i < nelems; ++i) { + ASSERT_TRUE(host_arr[i] == val); + } + delete[] host_arr; +} + +TEST_P(TestDPCTLQueueMemberFunctions, CheckFill32WithEvents) +{ + using T = uint32_t; + DPCTLSyclUSMRef p = nullptr; + DPCTLSyclEventRef FillDep_ERef = nullptr; + DPCTLSyclEventRef Fill_ERef = nullptr; + DPCTLSyclEventRef Memcpy_ERef = nullptr; + T val = static_cast(0xABCD); + size_t nelems = 256; + T *host_arr = new T[nelems]; + size_t nbytes = nelems * sizeof(T); + + ASSERT_FALSE(host_arr == nullptr); + + ASSERT_NO_FATAL_FAILURE(p = DPCTLmalloc_device(nbytes, QRef)); + ASSERT_FALSE(p == nullptr); + + ASSERT_NO_FATAL_FAILURE(FillDep_ERef = + DPCTLQueue_Fill32(QRef, (void *)p, 0, nelems)); + + ASSERT_NO_FATAL_FAILURE( + Fill_ERef = DPCTLQueue_Fill32WithEvents(QRef, (void *)p, val, nelems, + &FillDep_ERef, 1)); + + ASSERT_NO_FATAL_FAILURE(Memcpy_ERef = DPCTLQueue_MemcpyWithEvents( + QRef, host_arr, p, nbytes, &Fill_ERef, 1)); + ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Wait(Memcpy_ERef)); + + ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Delete(FillDep_ERef)); + ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Delete(Fill_ERef)); + ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Delete(Memcpy_ERef)); + + ASSERT_NO_FATAL_FAILURE(DPCTLfree_with_queue(p, QRef)); + + for (size_t i = 0; i < nelems; ++i) { + ASSERT_TRUE(host_arr[i] == val); + } + delete[] host_arr; +} + +TEST_P(TestDPCTLQueueMemberFunctions, CheckFill64WithEvents) +{ + using T = uint64_t; + DPCTLSyclUSMRef p = nullptr; + DPCTLSyclEventRef FillDep_ERef = nullptr; + DPCTLSyclEventRef Fill_ERef = nullptr; + DPCTLSyclEventRef Memcpy_ERef = nullptr; + T val = static_cast(0xABCDEF73); + size_t nelems = 256; + T *host_arr = new T[nelems]; + size_t nbytes = nelems * sizeof(T); + + ASSERT_FALSE(host_arr == nullptr); + + ASSERT_NO_FATAL_FAILURE(p = DPCTLmalloc_device(nbytes, QRef)); + ASSERT_FALSE(p == nullptr); + + ASSERT_NO_FATAL_FAILURE(FillDep_ERef = + DPCTLQueue_Fill64(QRef, (void *)p, 0, nelems)); + + ASSERT_NO_FATAL_FAILURE( + Fill_ERef = DPCTLQueue_Fill64WithEvents(QRef, (void *)p, val, nelems, + &FillDep_ERef, 1)); + + ASSERT_NO_FATAL_FAILURE(Memcpy_ERef = DPCTLQueue_MemcpyWithEvents( + QRef, host_arr, p, nbytes, &Fill_ERef, 1)); + ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Wait(Memcpy_ERef)); + + ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Delete(FillDep_ERef)); + ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Delete(Fill_ERef)); + ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Delete(Memcpy_ERef)); + + ASSERT_NO_FATAL_FAILURE(DPCTLfree_with_queue(p, QRef)); + + for (size_t i = 0; i < nelems; ++i) { + ASSERT_TRUE(host_arr[i] == val); + } + delete[] host_arr; +} + +TEST_P(TestDPCTLQueueMemberFunctions, CheckFill128WithEvents) +{ + using T = value128_t; + DPCTLSyclUSMRef p = nullptr; + DPCTLSyclEventRef FillDep_ERef = nullptr; + DPCTLSyclEventRef Fill_ERef = nullptr; + DPCTLSyclEventRef Memcpy_ERef = nullptr; + T val{static_cast(0xABCDEF73), static_cast(0x3746AF05)}; + T zero{}; + size_t nelems = 256; + T *host_arr = new T[nelems]; + size_t nbytes = nelems * sizeof(T); + + ASSERT_FALSE(host_arr == nullptr); + + ASSERT_NO_FATAL_FAILURE(p = DPCTLmalloc_device(nbytes, QRef)); + ASSERT_FALSE(p == nullptr); + + ASSERT_NO_FATAL_FAILURE( + FillDep_ERef = DPCTLQueue_Fill128( + QRef, (void *)p, reinterpret_cast(&zero), nelems)); + + ASSERT_NO_FATAL_FAILURE(Fill_ERef = DPCTLQueue_Fill128WithEvents( + QRef, (void *)p, + reinterpret_cast(&val), nelems, + &FillDep_ERef, 1)); + + ASSERT_NO_FATAL_FAILURE(Memcpy_ERef = DPCTLQueue_MemcpyWithEvents( + QRef, host_arr, p, nbytes, &Fill_ERef, 1)); + ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Wait(Memcpy_ERef)); + + ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Delete(FillDep_ERef)); + ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Delete(Fill_ERef)); + ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Delete(Memcpy_ERef)); + + ASSERT_NO_FATAL_FAILURE(DPCTLfree_with_queue(p, QRef)); + + for (size_t i = 0; i < nelems; ++i) { + ASSERT_TRUE(host_arr[i].first == val.first); + ASSERT_TRUE(host_arr[i].second == val.second); + } + delete[] host_arr; +} + INSTANTIATE_TEST_SUITE_P( DPCTLQueueMemberFuncTests, TestDPCTLQueueMemberFunctions,