Allow NumPy data types as arguments

CHANGELOG.md

@@ -6,9 +6,19 @@ 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
+* `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
+* `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 and Python `bool`, `int`, `float`, and `complex` scalars [gh-2335](https://github.com/IntelPython/dpctl/pull/2335)
+* Added `DPCTL_BYTES16_T` to `DPCTLKernelArgType` before `DPCTL_UNSUPPORTED_KERNEL_ARG`, which shifts the value of the latter [gh-2335](https://github.com/IntelPython/dpctl/pull/2335)
 
 ### Deprecated
 

docs/doc_sources/api_reference/dpctl/index.rst

@@ -77,6 +77,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.program.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:
+
+.. 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,26 @@ 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, and rounds its value to 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

@@ -82,3 +82,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,116 @@
+#                      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 not a supported data type.
+    """
+    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}'")
+    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(