Allow NumPy data types as arguments

.github/workflows/conda-package-cf.yml

@@ -238,7 +238,7 @@ jobs:
           CHANNELS="-c $GITHUB_WORKSPACE/channel ${{ env.CHANNELS }}"
           export CHANNELS
           # intel-opencl-rt provides the OpenCL CPU device to run the tests on
-          TEST_DEPENDENCIES="pytest cython setuptools intel-opencl-rt"
+          TEST_DEPENDENCIES="pytest cython setuptools intel-opencl-rt ml_dtypes"
           export TEST_DEPENDENCIES
           PACKAGE_VERSION=$(python -c "${VER_SCRIPT1} ${VER_SCRIPT2}")
           PYTHON_VER="${{ matrix.python_spec }}"
@@ -352,7 +352,7 @@ jobs:
           SET "PYTHON_VER=${{ matrix.python_spec }}"
           IF "%PYTHON_VER%"=="" SET "PYTHON_VER=${{ matrix.python }}"
           REM intel-opencl-rt provides the OpenCL CPU device to run the tests on
-          SET "TEST_DEPENDENCIES=pytest cython setuptools intel-opencl-rt"
+          SET "TEST_DEPENDENCIES=pytest cython setuptools intel-opencl-rt ml_dtypes"
           conda install -n ${{ env.TEST_ENV_NAME }} ${{ env.PACKAGE_NAME }}=%PACKAGE_VERSION% %TEST_DEPENDENCIES% python="%PYTHON_VER%" -c ${{ env.workdir }}/channel ${{ env.CHANNELS }}
 
       - name: Report content of test environment

.github/workflows/conda-package.yml

@@ -276,7 +276,7 @@ jobs:
         run: |
           CHANNELS="-c $GITHUB_WORKSPACE/channel ${{ env.CHANNELS }}"
           export CHANNELS
-          TEST_DEPENDENCIES="pytest cython setuptools"
+          TEST_DEPENDENCIES="pytest cython setuptools ml_dtypes"
           export TEST_DEPENDENCIES
           PACKAGE_VERSION=$(python -c "${VER_SCRIPT1} ${VER_SCRIPT2}")
           PYTHON_VER="${{ matrix.python_spec }}"
@@ -452,7 +452,7 @@ jobs:
           )
            SET "PYTHON_VER=${{ matrix.python_spec }}"
            IF "%PYTHON_VER%"=="" SET "PYTHON_VER=${{ matrix.python }}"
-          SET TEST_DEPENDENCIES=pytest"<8" cython setuptools
+          SET TEST_DEPENDENCIES=pytest"<8" cython setuptools ml_dtypes
            conda install -n ${{ env.TEST_ENV_NAME }} ${{ env.PACKAGE_NAME }}=%PACKAGE_VERSION% %TEST_DEPENDENCIES% python="%PYTHON_VER%" -c ${{ env.workdir }}/channel ${{ env.CHANNELS }}
 
       - name: Report content of test environment

.github/workflows/generate-coverage.yaml

@@ -93,7 +93,7 @@ jobs:
       - name: Install dpctl dependencies
         shell: bash -l {0}
         run: |
-          pip install numpy cython setuptools"<80" pytest pytest-cov scikit-build cmake coverage[toml] versioneer[toml]==0.29
+          pip install numpy cython setuptools"<80" pytest pytest-cov ml_dtypes scikit-build cmake coverage[toml] versioneer[toml]==0.29
 
       - name: Build dpctl with coverage
         shell: bash -l {0}

.github/workflows/os-llvm-sycl-build.yml

@@ -155,7 +155,7 @@ jobs:
       - name: Install dpctl dependencies
         shell: bash -l {0}
         run: |
-          pip install numpy cython setuptools"<80" pytest scikit-build cmake ninja versioneer[toml]==0.29
+          pip install numpy cython setuptools"<80" pytest ml_dtypes scikit-build cmake ninja versioneer[toml]==0.29
 
       - name: Create set_allvars.sh
         shell: bash -l {0}

CHANGELOG.md

@@ -6,16 +6,25 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0
 
 ## [dev] - XXX. XX, XXXX
 
+**NOTE**: This release contains breaking changes:
+
+* The value of the `DPCTL_UNSUPPORTED_KERNEL_ARG` enumerator of `DPCTLKernelArgType` increased by one, as `DPCTL_BYTES16_T` was added before it [gh-2335](https://github.com/IntelPython/dpctl/pull/2335)
+
 ### Added
 * Added `dpctl.SyclQueue.keep_args_alive` method, which keeps Python objects alive until the work on that queue using them completes [gh-2359](https://github.com/IntelPython/dpctl/pull/2359)
 * Added `add_event` and `add_cleanup_event` methods and `cleanup_events` and `num_cleanup_events` properties to the order manager, for tracking the events that gate the release of objects used by offloaded tasks [gh-2359](https://github.com/IntelPython/dpctl/pull/2359)
 * Added `SyclQueue_DeferRelease` C-API function, which defers a release callback until work submitted to a queue completes, and through which `dpctl::utils::keep_args_alive` keeps its arguments alive [gh-2359](https://github.com/IntelPython/dpctl/pull/2359)
 * Added `dpctl::utils::defer_release` to `dpctl4pybind11.hpp`, which calls a release callback once the given events have completed [gh-2359](https://github.com/IntelPython/dpctl/pull/2359)
+* `dpctl.SyclQueue.submit` and `dpctl.SyclQueue.submit_async` accept NumPy scalars of any data type whose values are plain bytes, including structured data types and those of `ml_dtypes` such as `bfloat16` and `complex32`, as well as Python `bool`, `int`, `float`, and `complex` scalars, which are passed as `numpy.bool`, `numpy.int64`, `numpy.float64`, and `numpy.complex128` respectively [gh-2335](https://github.com/IntelPython/dpctl/pull/2335)
+* Added `DPCTL_BYTES16_T` enumerator to `DPCTLKernelArgType` for 16-byte scalar and local accessor kernel arguments [gh-2335](https://github.com/IntelPython/dpctl/pull/2335)
+* Added a `sycl::ext::oneapi::bfloat16` type caster to `dpctl4pybind11.hpp`, defined when the compiler provides the `sycl_ext_oneapi_bfloat16` extension [gh-2335](https://github.com/IntelPython/dpctl/pull/2335)
 
 ### Changed
 * Object lifetime during offload is now managed by a list of releases deferred until the events that gate them have completed, rather than by a `host_task`. `dpctl` keeps one thread for looking over the list and running the releases that have come due, started by the first release deferred and again in the child of a `fork`, which takes the GIL to do so and lets a release call into the interpreter and free USM allocations [gh-2359](https://github.com/IntelPython/dpctl/pull/2359)
 * `import_dpctl()` now returns 0 on success and -1 with a Python exception set on failure, and `dpctl4pybind11.hpp` raises that exception rather than failing on the first C-API call [gh-2359](https://github.com/IntelPython/dpctl/pull/2359)
 * The event returned by `dpctl.SyclQueue.keep_args_alive`, `dpctl::utils::keep_args_alive` and `dpctl.SyclQueue._submit_keep_args_alive` is now for an empty kernel that gates the deferred release rather than for a `host_task` that performs it, so its completion means that the objects are no longer in use rather than that they have been released [gh-2359](https://github.com/IntelPython/dpctl/pull/2359)
+* `dpctl.SyclQueue.copy`, `dpctl.SyclQueue.copy_async`, `dpctl.SyclQueue.fill`, `dpctl.SyclQueue.fill_async`, `dpctl.LocalAccessor`, `dpctl.WorkGroupMemory`, and `dpctl.compiler.SpecializationConstant` accept a data type in any form `numpy.dtype` accepts, for any data type whose values are plain bytes, rather than a fixed set of type strings [gh-2335](https://github.com/IntelPython/dpctl/pull/2335)
+* `dpctl.compiler.SpecializationConstant` accepts NumPy scalars [gh-2335](https://github.com/IntelPython/dpctl/pull/2335)
 
 ### Deprecated
 * Deprecated `dpctl.SyclQueue._submit_keep_args_alive` in favor of `dpctl.SyclQueue.keep_args_alive`, which it now calls and which behaves identically [gh-2359](https://github.com/IntelPython/dpctl/pull/2359)

conda-recipe/meta.yaml

@@ -64,6 +64,7 @@ test:
         - cython
         - setuptools
         - pytest
+        - ml_dtypes
 
 about:
     home: https://github.com/IntelPython/dpctl.git

docs/doc_sources/api_reference/dpctl/index.rst

@@ -75,6 +75,58 @@
     * - :py:class:`dpctl.global_mem_cache_type`
       - An :class:`enum.Enum` of global memory cache types for a device.
 
+.. rubric:: Data types of kernel arguments
+
+.. _dpctl_data_types:
+
+Every :py:mod:`dpctl` function or class taking a data type, such as
+:py:meth:`dpctl.SyclQueue.copy`, :py:meth:`dpctl.SyclQueue.fill`,
+:py:class:`dpctl.LocalAccessor`, :py:class:`dpctl.WorkGroupMemory` and
+:py:class:`dpctl.compiler.SpecializationConstant`, accepts it in any form
+:py:class:`numpy.dtype` accepts, e.g. ``"f4"``, ``"float32"``,
+:py:class:`numpy.float32` or a :py:class:`numpy.dtype`. Any data type whose
+values are plain bytes is supported, including structured data types and data
+types registered with :py:mod:`numpy` by other packages, such as those of
+``ml_dtypes``. Extended precision ``numpy.longdouble`` is not supported.
+
+A NumPy scalar kernel argument is passed as its bytes, so the kernel determines
+how they are interpreted.
+
+.. code-block:: python
+
+    import ml_dtypes
+
+    bf16 = ml_dtypes.bfloat16(1.5)   # sycl::ext::oneapi::bfloat16
+    c = ml_dtypes.complex32(1 - 2j)  # std::complex<sycl::half>
+
+    q.submit(kernel, [x_usm, bf16, c], [n])
+
+Scalars of 1, 2, 4, 8 or 16 bytes are always supported. Scalars of other sizes,
+such as a structured scalar matching a C++ ``struct``, require the
+``raw_kernel_arg`` DPC++ extension.
+
+Python scalars are converted to a NumPy scalar of a fixed data type, which must
+match the type of the kernel parameter, as a mismatch is not detected:
+
+.. list-table::
+    :widths: 20 30 30
+
+    * - Python type
+      - NumPy type
+      - C++ type
+    * - :py:class:`bool`
+      - :py:class:`numpy.bool`
+      - ``bool``
+    * - :py:class:`int`
+      - :py:class:`numpy.int64`
+      - ``std::int64_t``
+    * - :py:class:`float`
+      - :py:class:`numpy.float64`
+      - ``double``
+    * - :py:class:`complex`
+      - :py:class:`numpy.complex128`
+      - ``std::complex<double>``
+
 .. rubric:: Exceptions
 
 .. autosummary::

docs/doc_sources/api_reference/dpctl_pybind11.rst

@@ -44,4 +44,27 @@ We can now use it from Python:
     # device descriptor
     print(dev.name)
 
+Type caster for ``bfloat16``
+----------------------------
+
+When the DPC++ compiler used to compile an extension provides the
+``sycl_ext_oneapi_bfloat16`` extension, ``dpctl4pybind11.hpp`` defines the macro
+``DPCTL_HAS_BFLOAT16`` and a type caster connecting
+``sycl::ext::oneapi::bfloat16`` to Python. A function taking such an argument
+accepts any Python object convertible to a :py:class:`float`, such as an
+``ml_dtypes.bfloat16`` scalar. Its value is rounded to C++ ``float`` and then
+to ``bfloat16``, so it is not always the nearest ``bfloat16``. A returned
+``bfloat16`` becomes a Python :py:class:`float`.
+
+Since the caster is only defined when the extension is available, code that has
+to build with compilers that do not provide it should guard its use:
+
+.. code-block:: c++
+
+    #ifdef DPCTL_HAS_BFLOAT16
+    sycl::ext::oneapi::bfloat16 round_trip(sycl::ext::oneapi::bfloat16 value) {
+        return value;
+    }
+    #endif
+
 .. _pybind11_url: https://pybind11.readthedocs.io/

docs/doc_sources/known_words.txt

@@ -81,3 +81,5 @@ profiler
 keepdims
 broadcasted
 broadcastable
+bfloat
+dtypes

dpctl/_backend.pxd

@@ -73,6 +73,7 @@ cdef extern from "syclinterface/dpctl_sycl_enum_types.h":
         _LOCAL_ACCESSOR     "DPCTL_LOCAL_ACCESSOR",
         _WORK_GROUP_MEMORY  "DPCTL_WORK_GROUP_MEMORY"
         _RAW_KERNEL_ARG     "DPCTL_RAW_KERNEL_ARG"
+        _BYTES16            "DPCTL_BYTES16_T"
 
     ctypedef enum _queue_property_type "DPCTLQueuePropertyType":
         _DEFAULT_PROPERTY   "DPCTL_DEFAULT_PROPERTY"

dpctl/_dtype_utils.py

@@ -0,0 +1,121 @@
+#                      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.
+
+"""Data type handling shared by the dpctl functions taking a data type."""
+
+import struct
+
+import numpy as np
+
+# kinds of NumPy data types whose values are not plain bytes
+_NON_POD_KINDS = frozenset("OSUMmT")
+
+# extended precision types, which devices do not support
+_EXTENDED_PRECISION_NUMS = frozenset(
+    np.dtype(t).num
+    for t in (np.longdouble, np.clongdouble)
+    if np.finfo(t).nmant > np.finfo(np.float64).nmant
+)
+
+# struct formats of the data types whose values are range checked when packed
+_STRUCT_FORMATS = {
+    ("b", 1): "?",
+    ("i", 1): "b",
+    ("i", 2): "h",
+    ("i", 4): "i",
+    ("i", 8): "q",
+    ("u", 1): "B",
+    ("u", 2): "H",
+    ("u", 4): "I",
+    ("u", 8): "Q",
+    ("f", 2): "e",
+    ("f", 4): "f",
+    ("f", 8): "d",
+    ("c", 8): "ff",
+    ("c", 16): "dd",
+}
+
+
+def is_supported_dtype(dt):
+    """Returns whether values of :class:`numpy.dtype` ``dt`` can be used on a
+    device."""
+    return not (
+        dt.hasobject
+        or dt.kind in _NON_POD_KINDS
+        or dt.itemsize == 0
+        or dt.num in _EXTENDED_PRECISION_NUMS
+    )
+
+
+def as_dtype(dtype):
+    """Returns ``numpy.dtype(dtype)``.
+
+    Raises:
+        ValueError: if ``dtype`` is ``None``, not a supported data type, or not
+            in native byte order.
+    """
+    if dtype is None:
+        raise ValueError("dtype must not be None")
+    try:
+        dt = np.dtype(dtype)
+    except TypeError:
+        raise ValueError(f"Unrecognized dtype '{dtype}'") from None
+    if not is_supported_dtype(dt):
+        raise ValueError(f"Unsupported dtype '{dt}'")
+    if not dt.isnative:
+        raise ValueError(f"dtype '{dt}' is not in native byte order")
+    return dt
+
+
+def as_scalar(value):
+    """Converts a Python ``bool``, ``int``, ``float`` or ``complex`` to a
+    ``numpy.bool``, ``numpy.int64``, ``numpy.float64`` or ``numpy.complex128``
+    respectively. Other values are returned unchanged.
+
+    Raises:
+        OverflowError: if an ``int`` does not fit in ``numpy.int64``.
+    """
+    if isinstance(value, np.generic):
+        return value
+    if isinstance(value, bool):
+        return np.bool_(value)
+    if isinstance(value, int):
+        return np.int64(value)
+    if isinstance(value, float):
+        return np.float64(value)
+    if isinstance(value, complex):
+        return np.complex128(value)
+    return value
+
+
+def pack_value(value, dt):
+    """Returns the bytes of ``value`` converted to :class:`numpy.dtype` ``dt``.
+
+    Raises:
+        ValueError: if ``value`` can not be represented in ``dt``.
+    """
+    fmt = _STRUCT_FORMATS.get((dt.kind, dt.itemsize))
+    try:
+        if fmt is None:
+            return np.asarray(value, dtype=dt).tobytes()
+        if dt.kind == "c":
+            cval = complex(value)
+            return struct.pack("=" + fmt, cval.real, cval.imag)
+        return struct.pack("=" + fmt, value)
+    except (struct.error, OverflowError, TypeError, ValueError) as e:
+        raise ValueError(
+            f"Value {value!r} cannot be represented as dtype '{dt}': {e}"
+        ) from None

dpctl/_sycl_queue.pxd

@@ -63,7 +63,8 @@ cdef public api class SyclQueue (_SyclQueue) [
         self,
         list args,
         void **kargs,
-        _arg_data_type *kargty
+        _arg_data_type *kargty,
+        list _converted
     )
     cdef int _populate_range(self, size_t Range[3], list gS, size_t nGS)
 
@@ -103,13 +104,13 @@ cdef public api class SyclQueue (_SyclQueue) [
     cdef DPCTLSyclQueueRef get_queue_ref(self)
     cpdef memcpy(self, dest, src, size_t count)
     cpdef SyclEvent memcpy_async(self, dest, src, size_t count, list dEvents=*)
-    cpdef copy(self, dest, src, size_t count, str dtype=*)
+    cpdef copy(self, dest, src, size_t count, dtype=*)
     cpdef SyclEvent copy_async(
-        self, dest, src, size_t count, list dEvents=*, str dtype=*
+        self, dest, src, size_t count, list dEvents=*, dtype=*
     )
-    cpdef fill(self, dest, value, size_t count, str dtype=*)
+    cpdef fill(self, dest, value, size_t count, dtype=*)
     cpdef SyclEvent fill_async(
-        self, dest, value, size_t count, list dEvents=*, str dtype=*
+        self, dest, value, size_t count, list dEvents=*, dtype=*
     )
     cpdef memset(self, mem, int val, size_t count=*)
     cpdef SyclEvent memset_async(