Skip to content

Commit 3575f77

Browse files
committed
Accept numpy data types for kernel arguments
Scalar kernel arguments are passed by size, so any numpy data type with plain-byte values works, including ml_dtypes types. copy, fill, LocalAccessor, WorkGroupMemory and SpecializationConstant accept any form numpy.dtype does. Adds a bfloat16 pybind11 type caster.
1 parent 2dfa6bb commit 3575f77

21 files changed

Lines changed: 1256 additions & 170 deletions

‎docs/doc_sources/api_reference/dpctl/index.rst‎

Lines changed: 52 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -77,6 +77,58 @@
7777
* - :py:class:`dpctl.global_mem_cache_type`
7878
- An :class:`enum.Enum` of global memory cache types for a device.
7979

80+
.. rubric:: Data types of kernel arguments
81+
82+
.. _dpctl_data_types:
83+
84+
Every :py:mod:`dpctl` function or class taking a data type, such as
85+
:py:meth:`dpctl.SyclQueue.copy`, :py:meth:`dpctl.SyclQueue.fill`,
86+
:py:class:`dpctl.LocalAccessor`, :py:class:`dpctl.WorkGroupMemory` and
87+
:py:class:`dpctl.program.SpecializationConstant`, accepts it in any form
88+
:py:class:`numpy.dtype` accepts, e.g. ``"f4"``, ``"float32"``,
89+
:py:class:`numpy.float32` or a :py:class:`numpy.dtype`. Any data type whose
90+
values are plain bytes is supported, including structured data types and data
91+
types registered with :py:mod:`numpy` by other packages, such as those of
92+
``ml_dtypes``. Extended precision ``numpy.longdouble`` is not supported.
93+
94+
A NumPy scalar kernel argument is passed as its bytes, so the kernel determines
95+
how they are interpreted.
96+
97+
.. code-block:: python
98+
99+
import ml_dtypes
100+
101+
bf16 = ml_dtypes.bfloat16(1.5) # sycl::ext::oneapi::bfloat16
102+
c = ml_dtypes.complex32(1 - 2j) # std::complex<sycl::half>
103+
104+
q.submit(kernel, [x_usm, bf16, c], [n])
105+
106+
Scalars of 1, 2, 4, 8 or 16 bytes are always supported. Scalars of other sizes,
107+
such as a structured scalar matching a C++ ``struct``, require the
108+
``raw_kernel_arg`` DPC++ extension.
109+
110+
Python scalars are converted to a NumPy scalar of a fixed data type, which must
111+
match the type of the kernel parameter:
112+
113+
.. list-table::
114+
:widths: 20 30 30
115+
116+
* - Python type
117+
- NumPy type
118+
- C++ type
119+
* - :py:class:`bool`
120+
- :py:class:`numpy.bool`
121+
- ``bool``
122+
* - :py:class:`int`
123+
- :py:class:`numpy.int64`
124+
- ``std::int64_t``
125+
* - :py:class:`float`
126+
- :py:class:`numpy.float64`
127+
- ``double``
128+
* - :py:class:`complex`
129+
- :py:class:`numpy.complex128`
130+
- ``std::complex<double>``
131+
80132
.. rubric:: Exceptions
81133

82134
.. autosummary::

‎docs/doc_sources/api_reference/dpctl_pybind11.rst‎

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -44,4 +44,26 @@ We can now use it from Python:
4444
# device descriptor
4545
print(dev.name)
4646
47+
Type caster for ``bfloat16``
48+
----------------------------
49+
50+
When the DPC++ compiler used to compile an extension provides the
51+
``sycl_ext_oneapi_bfloat16`` extension, ``dpctl4pybind11.hpp`` defines the macro
52+
``DPCTL_HAS_BFLOAT16`` and a type caster connecting
53+
``sycl::ext::oneapi::bfloat16`` to Python. A function taking such an argument
54+
accepts any Python object convertible to a :py:class:`float`, such as an
55+
``ml_dtypes.bfloat16`` scalar, and rounds its value to the nearest
56+
``bfloat16``. A returned ``bfloat16`` becomes a Python :py:class:`float`.
57+
58+
Since the caster is only defined when the extension is available, code that has
59+
to build with compilers that do not provide it should guard its use:
60+
61+
.. code-block:: c++
62+
63+
#ifdef DPCTL_HAS_BFLOAT16
64+
sycl::ext::oneapi::bfloat16 round_trip(sycl::ext::oneapi::bfloat16 value) {
65+
return value;
66+
}
67+
#endif
68+
4769
.. _pybind11_url: https://pybind11.readthedocs.io/

‎docs/doc_sources/known_words.txt‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -82,3 +82,5 @@ profiler
8282
keepdims
8383
broadcasted
8484
broadcastable
85+
bfloat
86+
dtypes

‎dpctl/_backend.pxd‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -73,6 +73,7 @@ cdef extern from "syclinterface/dpctl_sycl_enum_types.h":
7373
_LOCAL_ACCESSOR "DPCTL_LOCAL_ACCESSOR",
7474
_WORK_GROUP_MEMORY "DPCTL_WORK_GROUP_MEMORY"
7575
_RAW_KERNEL_ARG "DPCTL_RAW_KERNEL_ARG"
76+
_BYTES16 "DPCTL_BYTES16_T"
7677

7778
ctypedef enum _queue_property_type "DPCTLQueuePropertyType":
7879
_DEFAULT_PROPERTY "DPCTL_DEFAULT_PROPERTY"

‎dpctl/_dtype_utils.py‎

Lines changed: 116 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,116 @@
1+
# Data Parallel Control (dpctl)
2+
#
3+
# Copyright 2026 Intel Corporation
4+
#
5+
# Licensed under the Apache License, Version 2.0 (the "License");
6+
# you may not use this file except in compliance with the License.
7+
# You may obtain a copy of the License at
8+
#
9+
# http://www.apache.org/licenses/LICENSE-2.0
10+
#
11+
# Unless required by applicable law or agreed to in writing, software
12+
# distributed under the License is distributed on an "AS IS" BASIS,
13+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
14+
# See the License for the specific language governing permissions and
15+
# limitations under the License.
16+
17+
"""Data type handling shared by the dpctl functions taking a data type."""
18+
19+
import struct
20+
21+
import numpy as np
22+
23+
# kinds of NumPy data types whose values are not plain bytes
24+
_NON_POD_KINDS = frozenset("OSUMmT")
25+
26+
# extended precision types, which devices do not support
27+
_EXTENDED_PRECISION_NUMS = frozenset(
28+
np.dtype(t).num
29+
for t in (np.longdouble, np.clongdouble)
30+
if np.finfo(t).nmant > np.finfo(np.float64).nmant
31+
)
32+
33+
# struct formats of the data types whose values are range checked when packed
34+
_STRUCT_FORMATS = {
35+
("b", 1): "?",
36+
("i", 1): "b",
37+
("i", 2): "h",
38+
("i", 4): "i",
39+
("i", 8): "q",
40+
("u", 1): "B",
41+
("u", 2): "H",
42+
("u", 4): "I",
43+
("u", 8): "Q",
44+
("f", 2): "e",
45+
("f", 4): "f",
46+
("f", 8): "d",
47+
("c", 8): "ff",
48+
("c", 16): "dd",
49+
}
50+
51+
52+
def is_supported_dtype(dt):
53+
"""Returns whether values of :class:`numpy.dtype` ``dt`` can be used on a
54+
device."""
55+
return not (
56+
dt.hasobject
57+
or dt.kind in _NON_POD_KINDS
58+
or dt.itemsize == 0
59+
or dt.num in _EXTENDED_PRECISION_NUMS
60+
)
61+
62+
63+
def as_dtype(dtype):
64+
"""Returns ``numpy.dtype(dtype)``.
65+
66+
Raises:
67+
ValueError: if ``dtype`` is not a supported data type.
68+
"""
69+
try:
70+
dt = np.dtype(dtype)
71+
except TypeError:
72+
raise ValueError(f"Unrecognized dtype '{dtype}'") from None
73+
if not is_supported_dtype(dt):
74+
raise ValueError(f"Unsupported dtype '{dt}'")
75+
return dt
76+
77+
78+
def as_scalar(value):
79+
"""Converts a Python ``bool``, ``int``, ``float`` or ``complex`` to a
80+
``numpy.bool``, ``numpy.int64``, ``numpy.float64`` or ``numpy.complex128``
81+
respectively. Other values are returned unchanged.
82+
83+
Raises:
84+
OverflowError: if an ``int`` does not fit in ``numpy.int64``.
85+
"""
86+
if isinstance(value, np.generic):
87+
return value
88+
if isinstance(value, bool):
89+
return np.bool_(value)
90+
if isinstance(value, int):
91+
return np.int64(value)
92+
if isinstance(value, float):
93+
return np.float64(value)
94+
if isinstance(value, complex):
95+
return np.complex128(value)
96+
return value
97+
98+
99+
def pack_value(value, dt):
100+
"""Returns the bytes of ``value`` converted to :class:`numpy.dtype` ``dt``.
101+
102+
Raises:
103+
ValueError: if ``value`` can not be represented in ``dt``.
104+
"""
105+
fmt = _STRUCT_FORMATS.get((dt.kind, dt.itemsize))
106+
try:
107+
if fmt is None:
108+
return np.asarray(value, dtype=dt).tobytes()
109+
if dt.kind == "c":
110+
cval = complex(value)
111+
return struct.pack("=" + fmt, cval.real, cval.imag)
112+
return struct.pack("=" + fmt, value)
113+
except (struct.error, OverflowError, TypeError, ValueError) as e:
114+
raise ValueError(
115+
f"Value {value!r} cannot be represented as dtype '{dt}': {e}"
116+
) from None

‎dpctl/_sycl_queue.pxd‎

Lines changed: 6 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -63,7 +63,8 @@ cdef public api class SyclQueue (_SyclQueue) [
6363
self,
6464
list args,
6565
void **kargs,
66-
_arg_data_type *kargty
66+
_arg_data_type *kargty,
67+
list _converted
6768
)
6869
cdef int _populate_range(self, size_t Range[3], list gS, size_t nGS)
6970

@@ -103,13 +104,13 @@ cdef public api class SyclQueue (_SyclQueue) [
103104
cdef DPCTLSyclQueueRef get_queue_ref(self)
104105
cpdef memcpy(self, dest, src, size_t count)
105106
cpdef SyclEvent memcpy_async(self, dest, src, size_t count, list dEvents=*)
106-
cpdef copy(self, dest, src, size_t count, str dtype=*)
107+
cpdef copy(self, dest, src, size_t count, dtype=*)
107108
cpdef SyclEvent copy_async(
108-
self, dest, src, size_t count, list dEvents=*, str dtype=*
109+
self, dest, src, size_t count, list dEvents=*, dtype=*
109110
)
110-
cpdef fill(self, dest, value, size_t count, str dtype=*)
111+
cpdef fill(self, dest, value, size_t count, dtype=*)
111112
cpdef SyclEvent fill_async(
112-
self, dest, value, size_t count, list dEvents=*, str dtype=*
113+
self, dest, value, size_t count, list dEvents=*, dtype=*
113114
)
114115
cpdef memset(self, mem, int val, size_t count=*)
115116
cpdef SyclEvent memset_async(

0 commit comments

Comments
 (0)