diff --git a/dpctl/__init__.py b/dpctl/__init__.py index 5f57cdf90e..65048a8535 100644 --- a/dpctl/__init__.py +++ b/dpctl/__init__.py @@ -29,6 +29,24 @@ from . import _init_helper from ._device_selection import select_device_with_aspects +from ._dtypes import ( + DpctlScalar, + bool, + complex64, + complex128, + dtype, + float16, + float32, + float64, + int8, + int16, + int32, + int64, + uint8, + uint16, + uint32, + uint64, +) from ._sycl_context import SyclContext, SyclContextCreationError from ._sycl_device import ( SyclDevice, @@ -122,6 +140,24 @@ __all__ += [ "get_include", ] +__all__ += [ + "dtype", + "DpctlScalar", + "bool", + "int8", + "uint8", + "int16", + "uint16", + "int32", + "uint32", + "int64", + "uint64", + "float16", + "float32", + "float64", + "complex64", + "complex128", +] # add submodules __all__ += [ "memory", diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index 452ff8f255..ddaaeb6504 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -72,6 +72,9 @@ 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" + _FLOAT16 "DPCTL_FLOAT16_T" + _COMPLEX64 "DPCTL_COMPLEX64_T" + _COMPLEX128 "DPCTL_COMPLEX128_T" ctypedef enum _queue_property_type "DPCTLQueuePropertyType": _DEFAULT_PROPERTY "DPCTL_DEFAULT_PROPERTY" diff --git a/dpctl/_dtypes.py b/dpctl/_dtypes.py new file mode 100644 index 0000000000..7c613d8a08 --- /dev/null +++ b/dpctl/_dtypes.py @@ -0,0 +1,346 @@ +# 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. + +"""Scalar types and dtype descriptors for dpctl kernel arguments.""" + +import builtins + +import numpy as np + +__all__ = [ + "dtype", + "DpctlScalar", + "bool", + "int8", + "uint8", + "int16", + "uint16", + "int32", + "uint32", + "int64", + "uint64", + "float16", + "float32", + "float64", + "complex64", + "complex128", +] + +_DPCTL_INT8_T = 0 +_DPCTL_UINT8_T = 1 +_DPCTL_INT16_T = 2 +_DPCTL_UINT16_T = 3 +_DPCTL_INT32_T = 4 +_DPCTL_UINT32_T = 5 +_DPCTL_INT64_T = 6 +_DPCTL_UINT64_T = 7 +_DPCTL_FLOAT32_T = 8 +_DPCTL_FLOAT64_T = 9 +_DPCTL_FLOAT16_T = 14 +_DPCTL_COMPLEX64_T = 15 +_DPCTL_COMPLEX128_T = 16 + + +class DpctlScalar: + """Base class for dpctl scalar kernel argument types.""" + + _cast = None + _arg_type_id = None + + __slots__ = ("_value",) + + def __init__(self, value): + self._value = self._cast(value) + + @property + def value(self): + """The underlying numpy scalar value.""" + return self._value + + @property + def dtype(self): + """The numpy dtype of this scalar.""" + return self._value.dtype + + def __repr__(self): + return f"dpctl.{type(self).__name__}({self._value})" + + +class bool(DpctlScalar): + """Boolean scalar.""" + + _cast = np.int8 + _arg_type_id = _DPCTL_INT8_T + + def __init__(self, value): + super().__init__(builtins.bool(value)) + + +class int8(DpctlScalar): + """Signed 8-bit integer scalar.""" + + _cast = np.int8 + _arg_type_id = _DPCTL_INT8_T + + +class uint8(DpctlScalar): + """Unsigned 8-bit integer scalar.""" + + _cast = np.uint8 + _arg_type_id = _DPCTL_UINT8_T + + +class int16(DpctlScalar): + """Signed 16-bit integer scalar.""" + + _cast = np.int16 + _arg_type_id = _DPCTL_INT16_T + + +class uint16(DpctlScalar): + """Unsigned 16-bit integer scalar.""" + + _cast = np.uint16 + _arg_type_id = _DPCTL_UINT16_T + + +class int32(DpctlScalar): + """Signed 32-bit integer scalar.""" + + _cast = np.int32 + _arg_type_id = _DPCTL_INT32_T + + +class uint32(DpctlScalar): + """Unsigned 32-bit integer scalar.""" + + _cast = np.uint32 + _arg_type_id = _DPCTL_UINT32_T + + +class int64(DpctlScalar): + """Signed 64-bit integer scalar.""" + + _cast = np.int64 + _arg_type_id = _DPCTL_INT64_T + + +class uint64(DpctlScalar): + """Unsigned 64-bit integer scalar.""" + + _cast = np.uint64 + _arg_type_id = _DPCTL_UINT64_T + + +class float16(DpctlScalar): + """IEEE 754 half-precision (16-bit) floating-point scalar.""" + + _cast = np.float16 + _arg_type_id = _DPCTL_FLOAT16_T + + +class float32(DpctlScalar): + """IEEE 754 single-precision (32-bit) floating-point scalar.""" + + _cast = np.float32 + _arg_type_id = _DPCTL_FLOAT32_T + + +class float64(DpctlScalar): + """IEEE 754 double-precision (64-bit) floating-point scalar.""" + + _cast = np.float64 + _arg_type_id = _DPCTL_FLOAT64_T + + +class complex64(DpctlScalar): + """Complex type with two float32 components (64 bits total).""" + + _cast = np.complex64 + _arg_type_id = _DPCTL_COMPLEX64_T + + +class complex128(DpctlScalar): + """Complex type with two float64 components (128 bits total).""" + + _cast = np.complex128 + _arg_type_id = _DPCTL_COMPLEX128_T + + +# dtype metadata +_scalar_type_info = { + bool: (1, "i", "i1"), + int8: (1, "i", "i1"), + uint8: (1, "u", "u1"), + int16: (2, "i", "i2"), + uint16: (2, "u", "u2"), + int32: (4, "i", "i4"), + uint32: (4, "u", "u4"), + int64: (8, "i", "i8"), + uint64: (8, "u", "u8"), + float16: (2, "f", "f2"), + float32: (4, "f", "f4"), + float64: (8, "f", "f8"), + complex64: (8, "c", "c8"), + complex128: (16, "c", "c16"), +} + +_name_to_scalar_type = { + "bool": bool, + "int8": int8, + "uint8": uint8, + "int16": int16, + "uint16": uint16, + "int32": int32, + "uint32": uint32, + "int64": int64, + "uint64": uint64, + "float16": float16, + "float32": float32, + "float64": float64, + "complex64": complex64, + "complex128": complex128, +} + +_numpy_dtype_num_to_scalar_type = { + np.dtype(np.bool_).num: bool, + np.dtype(np.int8).num: int8, + np.dtype(np.uint8).num: uint8, + np.dtype(np.int16).num: int16, + np.dtype(np.uint16).num: uint16, + np.dtype(np.int32).num: int32, + np.dtype(np.uint32).num: uint32, + np.dtype(np.int64).num: int64, + np.dtype(np.uint64).num: uint64, + np.dtype(np.longlong).num: int64, + np.dtype(np.ulonglong).num: uint64, + np.dtype(np.float16).num: float16, + np.dtype(np.float32).num: float32, + np.dtype(np.float64).num: float64, + np.dtype(np.complex64).num: complex64, + np.dtype(np.complex128).num: complex128, +} + + +def _resolve_numpy_arg(arg): + try: + np_dt = np.dtype(arg) + except TypeError: + raise TypeError(f"Cannot create a dpctl.dtype from {arg!r}") + if np_dt.byteorder in ("<", ">"): + raise TypeError("Only native byte order is supported.") + if np_dt.num in _numpy_dtype_num_to_scalar_type: + return _numpy_dtype_num_to_scalar_type[np_dt.num] + raise TypeError(f"Cannot create a dpctl.dtype from {arg!r}") + + +class dtype: + """ + Dpctl dtype descriptor class. + + Analogous to :class:`numpy.dtype`, a :class:`dpctl.dtype` identifies a + data type and carries static metadata. + + A ``dpctl.dtype`` can be constructed from: + + - A dpctl scalar class + - A string name + - A NumPy dtype or type + - Another :class:`dpctl.dtype` + + Non-native byte orders are rejected. + """ + + __slots__ = ( + "_name", + "_itemsize", + "_kind", + "_str", + "_arg_type_id", + "_scalar_type", + ) + + def __init__(self, arg): + if isinstance(arg, dtype): + self._name = arg._name + self._itemsize = arg._itemsize + self._kind = arg._kind + self._str = arg._str + self._arg_type_id = arg._arg_type_id + self._scalar_type = arg._scalar_type + return + + if isinstance(arg, type) and issubclass(arg, DpctlScalar): + scalar_cls = arg + elif isinstance(arg, str) and arg in _name_to_scalar_type: + scalar_cls = _name_to_scalar_type[arg] + elif isinstance(arg, str): + scalar_cls = _resolve_numpy_arg(arg) + else: + scalar_cls = _resolve_numpy_arg(arg) + + self._scalar_type = scalar_cls + self._name = scalar_cls.__name__ + self._arg_type_id = scalar_cls._arg_type_id + info = _scalar_type_info[scalar_cls] + self._itemsize = info[0] + self._kind = info[1] + self._str = info[2] + + @property + def name(self): + """The short name of this dtype""" + return self._name + + @property + def itemsize(self): + """Element size in bytes.""" + return self._itemsize + + @property + def kind(self): + """Single-character kind code""" + return self._kind + + @property + def str(self): + """Short dtype string""" + return self._str + + @property + def scalar_type(self): + """The dpctl scalar class for this dtype""" + return self._scalar_type + + @property + def arg_type_id(self): + """The internal DPCTLKernelArgType enum value""" + return self._arg_type_id + + def __repr__(self): + return f"dpctl.dtype('{self._name}')" + + def __eq__(self, other): + if isinstance(other, dtype): + return self._arg_type_id == other._arg_type_id + return NotImplemented + + def __hash__(self): + return hash(self._arg_type_id) + + def __call__(self, value): + """Construct a scalar of this dtype""" + return self._scalar_type(value) diff --git a/dpctl/_sycl_queue.pxd b/dpctl/_sycl_queue.pxd index 9ca1d7c703..fddb718f58 100644 --- a/dpctl/_sycl_queue.pxd +++ b/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) diff --git a/dpctl/_sycl_queue.pyx b/dpctl/_sycl_queue.pyx index a22586b71e..c19d699aba 100644 --- a/dpctl/_sycl_queue.pyx +++ b/dpctl/_sycl_queue.pyx @@ -73,6 +73,9 @@ from .memory._memory cimport _Memory import ctypes import numbers +import numpy as np + +from ._dtypes import DpctlScalar from .enum_types import backend_type from cpython cimport pycapsule @@ -109,6 +112,25 @@ __all__ = [ _logger = logging.getLogger(__name__) +_numpy_dtype_num_to_arg_type = { + np.dtype(np.bool_).num: _arg_data_type._INT8_T, + np.dtype(np.int8).num: _arg_data_type._INT8_T, + np.dtype(np.uint8).num: _arg_data_type._UINT8_T, + np.dtype(np.int16).num: _arg_data_type._INT16_T, + np.dtype(np.uint16).num: _arg_data_type._UINT16_T, + np.dtype(np.int32).num: _arg_data_type._INT32_T, + np.dtype(np.uint32).num: _arg_data_type._UINT32_T, + np.dtype(np.int64).num: _arg_data_type._INT64_T, + np.dtype(np.uint64).num: _arg_data_type._UINT64_T, + np.dtype(np.longlong).num: _arg_data_type._INT64_T, + np.dtype(np.ulonglong).num: _arg_data_type._UINT64_T, + np.dtype(np.float16).num: _arg_data_type._FLOAT16, + np.dtype(np.float32).num: _arg_data_type._FLOAT, + np.dtype(np.float64).num: _arg_data_type._DOUBLE, + np.dtype(np.complex64).num: _arg_data_type._COMPLEX64, + np.dtype(np.complex128).num: _arg_data_type._COMPLEX128, +} + cdef class kernel_arg_type_attribute: cdef str parent_name @@ -152,9 +174,11 @@ cdef class LocalAccessor: `'u1'`, `'u2'`, `'u4'`, `'u8'` unsigned integral types uint8_t, uint16_t, uint32_t, uint64_t - `'f4'`, `'f8'`, - single- and double-precision floating-point types float and - double + `'f2'`, `'f4'`, `'f8'`: + half-, single-, and double-precision floating-point types + sycl::half, float, and double + `'c8'`, `'c16'`: + complex types with float or double components shape (tuple, list): Size of LocalAccessor dimensions. Dimension of the LocalAccessor is determined by the length of the tuple. Must be of length 1, 2, or 3, @@ -215,10 +239,16 @@ cdef class LocalAccessor: self.lacc.dpctl_type_id = _arg_data_type._INT64_T elif dtype == "u8": self.lacc.dpctl_type_id = _arg_data_type._UINT64_T + elif dtype == "f2": + self.lacc.dpctl_type_id = _arg_data_type._FLOAT16 elif dtype == "f4": self.lacc.dpctl_type_id = _arg_data_type._FLOAT elif dtype == "f8": self.lacc.dpctl_type_id = _arg_data_type._DOUBLE + elif dtype == "c8": + self.lacc.dpctl_type_id = _arg_data_type._COMPLEX64 + elif dtype == "c16": + self.lacc.dpctl_type_id = _arg_data_type._COMPLEX128 else: raise ValueError(f"Unrecognized type value: '{dtype}'") @@ -379,6 +409,33 @@ cdef class _kernel_arg_type: _arg_data_type._RAW_KERNEL_ARG ) + @property + def dpctl_float16(self): + cdef str p_name = "dpctl_float16" + return kernel_arg_type_attribute( + self._name, + p_name, + _arg_data_type._FLOAT16 + ) + + @property + def dpctl_complex64(self): + cdef str p_name = "dpctl_complex64" + return kernel_arg_type_attribute( + self._name, + p_name, + _arg_data_type._COMPLEX64 + ) + + @property + def dpctl_complex128(self): + cdef str p_name = "dpctl_complex128" + return kernel_arg_type_attribute( + self._name, + p_name, + _arg_data_type._COMPLEX128 + ) + kernel_arg_type = _kernel_arg_type() @@ -476,20 +533,23 @@ ctypedef DPCTLSyclEventRef (*queue_copy_with_events_fn)( cdef size_t _get_dtype_size(str dtype) except *: """ Parse numpy-style dtype string and return element size in bytes. - Supports: i1, u1, i2, u2, i4, u4, i8, u8, f4, f8 + Supports: i1, u1, i2, u2, i4, u4, i8, u8, f2, f4, f8, c8, c16 """ if dtype == "i1" or dtype == "u1": return 1 - elif dtype == "i2" or dtype == "u2": + elif dtype == "i2" or dtype == "u2" or dtype == "f2": return 2 elif dtype == "i4" or dtype == "u4" or dtype == "f4": return 4 - elif dtype == "i8" or dtype == "u8" or dtype == "f8": + elif dtype == "i8" or dtype == "u8" or dtype == "f8" or dtype == "c8": return 8 + elif dtype == "c16": + return 16 else: raise ValueError( f"Unrecognized dtype '{dtype}'. " - "Expected one of: i1, u1, i2, u2, i4, u4, i8, u8, f4, f8" + "Expected one of: i1, u1, i2, u2, i4, u4, i8, u8, " + "f2, f4, f8, c8, c16" ) @@ -1008,11 +1068,72 @@ cdef class SyclQueue(_SyclQueue): self, list args, void **kargs, - _arg_data_type *kargty + _arg_data_type *kargty, + list _converted ): cdef int ret = 0 + cdef Py_buffer _buf + cdef object np_scalar + cdef int dtype_num for idx, arg in enumerate(args): - if isinstance(arg, ctypes.c_char): + if isinstance(arg, DpctlScalar): + np_scalar = arg._value + if PyObject_GetBuffer(np_scalar, &_buf, PyBUF_SIMPLE) != 0: + ret = -1 + break + kargs[idx] = _buf.buf + kargty[idx] = <_arg_data_type>arg._arg_type_id + PyBuffer_Release(&_buf) + elif isinstance(arg, np.generic): + dtype_num = arg.dtype.num + if dtype_num not in _numpy_dtype_num_to_arg_type: + ret = -1 + break + if PyObject_GetBuffer(arg, &_buf, PyBUF_SIMPLE) != 0: + ret = -1 + break + kargs[idx] = _buf.buf + kargty[idx] = <_arg_data_type>( + _numpy_dtype_num_to_arg_type[dtype_num] + ) + PyBuffer_Release(&_buf) + elif isinstance(arg, bool): + np_scalar = np.int8(arg) + _converted.append(np_scalar) + if PyObject_GetBuffer(np_scalar, &_buf, PyBUF_SIMPLE) != 0: + ret = -1 + break + kargs[idx] = _buf.buf + kargty[idx] = _arg_data_type._INT8_T + PyBuffer_Release(&_buf) + elif isinstance(arg, int): + np_scalar = np.int64(arg) + _converted.append(np_scalar) + if PyObject_GetBuffer(np_scalar, &_buf, PyBUF_SIMPLE) != 0: + ret = -1 + break + kargs[idx] = _buf.buf + kargty[idx] = _arg_data_type._INT64_T + PyBuffer_Release(&_buf) + elif isinstance(arg, float): + np_scalar = np.float64(arg) + _converted.append(np_scalar) + if PyObject_GetBuffer(np_scalar, &_buf, PyBUF_SIMPLE) != 0: + ret = -1 + break + kargs[idx] = _buf.buf + kargty[idx] = _arg_data_type._DOUBLE + PyBuffer_Release(&_buf) + elif isinstance(arg, complex): + np_scalar = np.complex128(arg) + _converted.append(np_scalar) + if PyObject_GetBuffer(np_scalar, &_buf, PyBUF_SIMPLE) != 0: + ret = -1 + break + kargs[idx] = _buf.buf + kargty[idx] = _arg_data_type._COMPLEX128 + PyBuffer_Release(&_buf) + elif isinstance(arg, ctypes.c_char): kargs[idx] = (ctypes.addressof(arg)) kargty[idx] = _arg_data_type._INT8_T elif isinstance(arg, ctypes.c_uint8): @@ -1043,7 +1164,7 @@ cdef class SyclQueue(_SyclQueue): kargs[idx] = (ctypes.addressof(arg)) kargty[idx] = _arg_data_type._DOUBLE elif isinstance(arg, _Memory): - kargs[idx]= (arg._pointer) + kargs[idx] = (arg._pointer) kargty[idx] = _arg_data_type._VOID_PTR elif isinstance(arg, WorkGroupMemory): kargs[idx] = (arg._ref) @@ -1290,6 +1411,7 @@ cdef class SyclQueue(_SyclQueue): cdef size_t nGS = len(gS) cdef size_t nLS = len(lS) if lS is not None else 0 cdef size_t nDE = len(dEvents) if dEvents is not None else 0 + cdef list _converted = [] # Allocate the arrays to be sent to DPCTLQueue_Submit kargs = malloc(len(args) * sizeof(void*)) @@ -1323,7 +1445,7 @@ cdef class SyclQueue(_SyclQueue): ) # populate the args and argstype arrays - ret = self._populate_args(args, kargs, kargty) + ret = self._populate_args(args, kargs, kargty, _converted) if ret == -1: free(kargs) free(kargty) @@ -1900,7 +2022,7 @@ cdef class WorkGroupMemory: ) dtype = (args[0]) count = (args[1]) - if not dtype[0] in ["i", "u", "f"]: + if not dtype[0] in ["i", "u", "f", "c"]: raise TypeError(f"Unrecognized type value: '{dtype}'") try: bit_width = int(dtype[1:]) diff --git a/dpctl/tests/test_dtypes.py b/dpctl/tests/test_dtypes.py new file mode 100644 index 0000000000..2154364e45 --- /dev/null +++ b/dpctl/tests/test_dtypes.py @@ -0,0 +1,290 @@ +# 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 dpctl scalar types and dtype class""" + +import numpy as np +import pytest + +import dpctl +from dpctl._dtypes import DpctlScalar, dtype + + +@pytest.mark.parametrize( + "dpctl_type,value,expected_dtype", + [ + (dpctl.bool, True, np.dtype(np.int8)), + (dpctl.bool, False, np.dtype(np.int8)), + (dpctl.int8, -3, np.dtype(np.int8)), + (dpctl.uint8, 200, np.dtype(np.uint8)), + (dpctl.int16, -1000, np.dtype(np.int16)), + (dpctl.uint16, 60000, np.dtype(np.uint16)), + (dpctl.int32, -100000, np.dtype(np.int32)), + (dpctl.uint32, 3000000000, np.dtype(np.uint32)), + (dpctl.int64, -(2**40), np.dtype(np.int64)), + (dpctl.uint64, 2**40, np.dtype(np.uint64)), + (dpctl.float16, 3.14, np.dtype(np.float16)), + (dpctl.float32, 2.718, np.dtype(np.float32)), + (dpctl.float64, 1.41421356, np.dtype(np.float64)), + (dpctl.complex64, 1 + 2j, np.dtype(np.complex64)), + (dpctl.complex128, 3 + 4j, np.dtype(np.complex128)), + ], +) +def test_scalar_creation(dpctl_type, value, expected_dtype): + scalar = dpctl_type(value) + assert isinstance(scalar, DpctlScalar) + assert scalar.dtype == expected_dtype + + +@pytest.mark.parametrize( + "dpctl_type,value", + [ + (dpctl.int32, 42), + (dpctl.float64, 3.14), + (dpctl.complex128, 1 + 2j), + ], +) +def test_scalar_value(dpctl_type, value): + scalar = dpctl_type(value) + assert scalar.value == dpctl_type._cast(value) + + +def test_bool_conversion(): + t = dpctl.bool(1) + f = dpctl.bool(0) + assert t.value == np.int8(1) + assert f.value == np.int8(0) + truthy = dpctl.bool([1, 2, 3]) + assert truthy.value == np.int8(1) + + +def test_scalar_repr(): + s = dpctl.int32(5) + assert repr(s) == "dpctl.int32(5)" + s = dpctl.float16(1.5) + assert repr(s) == "dpctl.float16(1.5)" + + +@pytest.mark.parametrize( + "dpctl_type", + [ + dpctl.bool, + dpctl.int8, + dpctl.uint8, + dpctl.int16, + dpctl.uint16, + dpctl.int32, + dpctl.uint32, + dpctl.int64, + dpctl.uint64, + dpctl.float16, + dpctl.float32, + dpctl.float64, + dpctl.complex64, + dpctl.complex128, + ], +) +def test_scalar_has_arg_type_id(dpctl_type): + assert dpctl_type._arg_type_id is not None + assert isinstance(dpctl_type._arg_type_id, int) + + +@pytest.mark.parametrize( + "dpctl_type", + [ + dpctl.bool, + dpctl.int8, + dpctl.uint8, + dpctl.int16, + dpctl.uint16, + dpctl.int32, + dpctl.uint32, + dpctl.int64, + dpctl.uint64, + dpctl.float16, + dpctl.float32, + dpctl.float64, + dpctl.complex64, + dpctl.complex128, + ], +) +def test_scalar_numpy_buffer_protocol(dpctl_type): + """Verify that the underlying numpy scalar supports the buffer protocol.""" + scalar = dpctl_type(1) + mv = memoryview(scalar.value) + assert mv.nbytes == scalar.value.dtype.itemsize + + +@pytest.mark.parametrize( + "scalar_cls,expected_name,expected_itemsize,expected_kind", + [ + (dpctl.bool, "bool", 1, "i"), + (dpctl.int8, "int8", 1, "i"), + (dpctl.uint8, "uint8", 1, "u"), + (dpctl.int16, "int16", 2, "i"), + (dpctl.uint16, "uint16", 2, "u"), + (dpctl.int32, "int32", 4, "i"), + (dpctl.uint32, "uint32", 4, "u"), + (dpctl.int64, "int64", 8, "i"), + (dpctl.uint64, "uint64", 8, "u"), + (dpctl.float16, "float16", 2, "f"), + (dpctl.float32, "float32", 4, "f"), + (dpctl.float64, "float64", 8, "f"), + (dpctl.complex64, "complex64", 8, "c"), + (dpctl.complex128, "complex128", 16, "c"), + ], +) +def test_dtype_from_scalar_class( + scalar_cls, expected_name, expected_itemsize, expected_kind +): + dt = dtype(scalar_cls) + assert dt.name == expected_name + assert dt.itemsize == expected_itemsize + assert dt.kind == expected_kind + assert dt.scalar_type is scalar_cls + + +@pytest.mark.parametrize( + "name", + [ + "bool", + "int8", + "uint8", + "int16", + "uint16", + "int32", + "uint32", + "int64", + "uint64", + "float16", + "float32", + "float64", + "complex64", + "complex128", + ], +) +def test_dtype_from_string(name): + dt = dtype(name) + assert dt.name == name + assert dt.itemsize > 0 + + +@pytest.mark.parametrize( + "np_type,expected_name", + [ + (np.int8, "int8"), + (np.uint8, "uint8"), + (np.int16, "int16"), + (np.int32, "int32"), + (np.float16, "float16"), + (np.float32, "float32"), + (np.float64, "float64"), + (np.complex64, "complex64"), + (np.complex128, "complex128"), + ], +) +def test_dtype_from_numpy(np_type, expected_name): + dt = dtype(np_type) + assert dt.name == expected_name + + dt2 = dtype(np.dtype(np_type)) + assert dt == dt2 + + +def test_dtype_from_dtype(): + dt1 = dtype(dpctl.float32) + dt2 = dtype(dt1) + assert dt1 == dt2 + assert dt1.name == dt2.name + + +def test_dtype_equality_and_hash(): + dt1 = dtype(dpctl.int32) + dt2 = dtype("int32") + dt3 = dtype(np.int32) + assert dt1 == dt2 + assert dt2 == dt3 + assert hash(dt1) == hash(dt2) == hash(dt3) + + dt4 = dtype(dpctl.float32) + assert dt1 != dt4 + + +def test_dtype_repr(): + dt = dtype(dpctl.float32) + assert repr(dt) == "dpctl.dtype('float32')" + + +def test_dtype_callable(): + dt = dtype("float32") + scalar = dt(3.14) + assert isinstance(scalar, DpctlScalar) + assert isinstance(scalar, dpctl.float32) + + +def test_dtype_invalid(): + with pytest.raises(TypeError): + dtype("not_a_type") + + with pytest.raises(TypeError): + dtype(object) + + +def test_dtype_str_property(): + assert dtype(dpctl.int32).str == "i4" + assert dtype(dpctl.float16).str == "f2" + assert dtype(dpctl.complex128).str == "c16" + + +def _non_native_prefix(): + """Return the byte-order prefix that is non-native on this system.""" + import sys + + return ">" if sys.byteorder == "little" else "<" + + +@pytest.mark.parametrize( + "suffix", + ["f4", "i4", "u8", "i2"], +) +def test_dtype_rejects_non_native_endianness_string(suffix): + arg = _non_native_prefix() + suffix + with pytest.raises(TypeError): + dtype(arg) + + +@pytest.mark.parametrize( + "suffix", + ["f4", "i2"], +) +def test_dtype_rejects_non_native_endianness_numpy(suffix): + np_dt = np.dtype(_non_native_prefix() + suffix) + with pytest.raises(TypeError): + dtype(np_dt) + + +@pytest.mark.parametrize( + "arg", + [ + "=f4", + "|i1", + "=i4", + "=c16", + ], +) +def test_dtype_accepts_native_endianness(arg): + dt = dtype(arg) + assert dt.itemsize > 0 diff --git a/dpctl/tests/test_sycl_kernel_submit.py b/dpctl/tests/test_sycl_kernel_submit.py index e2d0971111..2357c0d1b1 100644 --- a/dpctl/tests/test_sycl_kernel_submit.py +++ b/dpctl/tests/test_sycl_kernel_submit.py @@ -273,6 +273,391 @@ def test_submit_async(): assert np.array_equal(x[:, :n], x_ref[:, :n]) +@pytest.mark.parametrize( + "ctype_str,dtype,scalar_ctor", + [ + ("int", np.dtype("i4"), lambda v: np.int32(v)), + ("unsigned int", np.dtype("u4"), lambda v: np.uint32(v)), + ("long", np.dtype(np.longlong), lambda v: np.longlong(v)), + ("float", np.dtype("f4"), lambda v: np.float32(v)), + ("double", np.dtype("f8"), lambda v: np.float64(v)), + ], +) +def test_kernel_submit_numpy_scalar_args(ctype_str, dtype, scalar_ctor): + """Test kernel submission with numpy scalars as kernel arguments""" + try: + q = dpctl.SyclQueue("opencl") + except dpctl.SyclQueueCreationError: + pytest.skip("OpenCL queue could not be created") + if dtype == np.dtype("f8") and q.sycl_device.has_aspect_fp64 is False: + pytest.skip( + "Device does not support double precision floating point type" + ) + oclSrc = ( + "kernel void axpy(" + " global " + ctype_str + " *a, global " + ctype_str + " *b," + " global " + ctype_str + " *c, " + ctype_str + " d) {" + " size_t index = get_global_id(0);" + " c[index] = d * a[index] + b[index];" + "}" + ) + kb = dpctl_prog.create_kernel_bundle_from_source(q, oclSrc) + axpyKernel = kb.get_sycl_kernel("axpy") + + n_elems = 1024 + a = np.arange(n_elems, dtype=dtype) + b = np.arange(n_elems, stop=0, step=-1, dtype=dtype) + c = np.zeros(n_elems, dtype=dtype) + + a_usm = dpm.MemoryUSMDevice(a.nbytes, queue=q) + b_usm = dpm.MemoryUSMDevice(b.nbytes, queue=q) + c_usm = dpm.MemoryUSMDevice(c.nbytes, queue=q) + + q.memcpy(dest=a_usm, src=a, count=a.nbytes) + q.memcpy(dest=b_usm, src=b, count=b.nbytes) + + d = 2 + args = [a_usm, b_usm, c_usm, scalar_ctor(d)] + q.submit(axpyKernel, args, [n_elems]).wait() + + q.memcpy(c, c_usm, c.nbytes) + ref_c = a * np.array(d, dtype=dtype) + b + assert np.allclose(c, ref_c) + + +@pytest.mark.parametrize( + "ctype_str,dtype,dpctl_ctor", + [ + ("int", np.dtype("i4"), dpctl.int32), + ("unsigned int", np.dtype("u4"), dpctl.uint32), + ("float", np.dtype("f4"), dpctl.float32), + ("double", np.dtype("f8"), dpctl.float64), + ], +) +def test_kernel_submit_dpctl_scalar_args(ctype_str, dtype, dpctl_ctor): + """Test kernel submission with dpctl scalar types as kernel arguments""" + try: + q = dpctl.SyclQueue("opencl") + except dpctl.SyclQueueCreationError: + pytest.skip("OpenCL queue could not be created") + if dtype == np.dtype("f8") and q.sycl_device.has_aspect_fp64 is False: + pytest.skip( + "Device does not support double precision floating point type" + ) + oclSrc = ( + "kernel void axpy(" + " global " + ctype_str + " *a, global " + ctype_str + " *b," + " global " + ctype_str + " *c, " + ctype_str + " d) {" + " size_t index = get_global_id(0);" + " c[index] = d * a[index] + b[index];" + "}" + ) + kb = dpctl_prog.create_kernel_bundle_from_source(q, oclSrc) + axpyKernel = kb.get_sycl_kernel("axpy") + + n_elems = 1024 + a = np.arange(n_elems, dtype=dtype) + b = np.arange(n_elems, stop=0, step=-1, dtype=dtype) + c = np.zeros(n_elems, dtype=dtype) + + a_usm = dpm.MemoryUSMDevice(a.nbytes, queue=q) + b_usm = dpm.MemoryUSMDevice(b.nbytes, queue=q) + c_usm = dpm.MemoryUSMDevice(c.nbytes, queue=q) + + q.memcpy(dest=a_usm, src=a, count=a.nbytes) + q.memcpy(dest=b_usm, src=b, count=b.nbytes) + + d = 2 + args = [a_usm, b_usm, c_usm, dpctl_ctor(d)] + q.submit(axpyKernel, args, [n_elems]).wait() + + q.memcpy(c, c_usm, c.nbytes) + ref_c = a * np.array(d, dtype=dtype) + b + assert np.allclose(c, ref_c) + + +@pytest.mark.parametrize( + "ctype_str,dtype,python_value", + [ + ("long", np.dtype(np.longlong), 2), + ("double", np.dtype("f8"), 2.0), + ], +) +def test_kernel_submit_python_builtin_args(ctype_str, dtype, python_value): + """Test kernel submission with Python int/float as kernel arguments""" + try: + q = dpctl.SyclQueue("opencl") + except dpctl.SyclQueueCreationError: + pytest.skip("OpenCL queue could not be created") + if dtype == np.dtype("f8") and q.sycl_device.has_aspect_fp64 is False: + pytest.skip( + "Device does not support double precision floating point type" + ) + oclSrc = ( + "kernel void axpy(" + " global " + ctype_str + " *a, global " + ctype_str + " *b," + " global " + ctype_str + " *c, " + ctype_str + " d) {" + " size_t index = get_global_id(0);" + " c[index] = d * a[index] + b[index];" + "}" + ) + kb = dpctl_prog.create_kernel_bundle_from_source(q, oclSrc) + axpyKernel = kb.get_sycl_kernel("axpy") + + n_elems = 1024 + a = np.arange(n_elems, dtype=dtype) + b = np.arange(n_elems, stop=0, step=-1, dtype=dtype) + c = np.zeros(n_elems, dtype=dtype) + + a_usm = dpm.MemoryUSMDevice(a.nbytes, queue=q) + b_usm = dpm.MemoryUSMDevice(b.nbytes, queue=q) + c_usm = dpm.MemoryUSMDevice(c.nbytes, queue=q) + + q.memcpy(dest=a_usm, src=a, count=a.nbytes) + q.memcpy(dest=b_usm, src=b, count=b.nbytes) + + args = [a_usm, b_usm, c_usm, python_value] + q.submit(axpyKernel, args, [n_elems]).wait() + + q.memcpy(c, c_usm, c.nbytes) + ref_c = a * np.array(python_value, dtype=dtype) + b + assert np.allclose(c, ref_c) + + +def test_kernel_submit_float16_arg(): + """Test kernel submission with a half-precision scalar argument""" + try: + q = dpctl.SyclQueue("opencl") + except dpctl.SyclQueueCreationError: + pytest.skip("OpenCL queue could not be created") + if not q.sycl_device.has_aspect_fp16: + pytest.skip("Device does not support half precision floating point") + + oclSrc = ( + "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n" + "kernel void scale(" + " global half *a, global half *b, half d) {" + " size_t index = get_global_id(0);" + " b[index] = d * a[index];" + "}" + ) + kb = dpctl_prog.create_kernel_bundle_from_source(q, oclSrc) + scaleKernel = kb.get_sycl_kernel("scale") + + n_elems = 1024 + a = np.ones(n_elems, dtype=np.float16) * np.float16(2.0) + b = np.zeros(n_elems, dtype=np.float16) + + a_usm = dpm.MemoryUSMDevice(a.nbytes, queue=q) + b_usm = dpm.MemoryUSMDevice(b.nbytes, queue=q) + + q.memcpy(dest=a_usm, src=a, count=a.nbytes) + + d = dpctl.float16(3.0) + args = [a_usm, b_usm, d] + q.submit(scaleKernel, args, [n_elems]).wait() + + q.memcpy(b, b_usm, b.nbytes) + expected = a * np.float16(3.0) + assert np.allclose(b, expected) + + +def test_kernel_submit_float16_numpy_scalar(): + """Test kernel submission with a numpy float16 scalar argument""" + try: + q = dpctl.SyclQueue("opencl") + except dpctl.SyclQueueCreationError: + pytest.skip("OpenCL queue could not be created") + if not q.sycl_device.has_aspect_fp16: + pytest.skip("Device does not support half precision floating point") + + oclSrc = ( + "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n" + "kernel void scale(" + " global half *a, global half *b, half d) {" + " size_t index = get_global_id(0);" + " b[index] = d * a[index];" + "}" + ) + kb = dpctl_prog.create_kernel_bundle_from_source(q, oclSrc) + scaleKernel = kb.get_sycl_kernel("scale") + + n_elems = 1024 + a = np.ones(n_elems, dtype=np.float16) * np.float16(4.0) + b = np.zeros(n_elems, dtype=np.float16) + + a_usm = dpm.MemoryUSMDevice(a.nbytes, queue=q) + b_usm = dpm.MemoryUSMDevice(b.nbytes, queue=q) + + q.memcpy(dest=a_usm, src=a, count=a.nbytes) + + args = [a_usm, b_usm, np.float16(2.5)] + q.submit(scaleKernel, args, [n_elems]).wait() + + q.memcpy(b, b_usm, b.nbytes) + expected = a * np.float16(2.5) + assert np.allclose(b, expected) + + +def test_kernel_submit_complex64_arg(): + """Test kernel submission with a complex64 scalar argument""" + try: + q = dpctl.SyclQueue("opencl") + except dpctl.SyclQueueCreationError: + pytest.skip("OpenCL queue could not be created") + + # OpenCL float2 has the same layout as complex + oclSrc = ( + "kernel void scale_complex(" + " global float2 *a, global float2 *b, float2 d) {" + " size_t index = get_global_id(0);" + " float2 ai = a[index];" + " b[index] = (float2)(ai.x * d.x - ai.y * d.y," + " ai.x * d.y + ai.y * d.x);" + "}" + ) + kb = dpctl_prog.create_kernel_bundle_from_source(q, oclSrc) + kernel = kb.get_sycl_kernel("scale_complex") + + n_elems = 512 + a = np.ones(n_elems, dtype=np.complex64) * np.complex64(1 + 1j) + b = np.zeros(n_elems, dtype=np.complex64) + + a_usm = dpm.MemoryUSMDevice(a.nbytes, queue=q) + b_usm = dpm.MemoryUSMDevice(b.nbytes, queue=q) + + q.memcpy(dest=a_usm, src=a, count=a.nbytes) + + d = dpctl.complex64(2 + 3j) + args = [a_usm, b_usm, d] + q.submit(kernel, args, [n_elems]).wait() + + q.memcpy(b, b_usm, b.nbytes) + expected = a * np.complex64(2 + 3j) + assert np.allclose(b, expected) + + +def test_kernel_submit_complex64_numpy_scalar(): + """Test kernel submission with a numpy complex64 scalar argument""" + try: + q = dpctl.SyclQueue("opencl") + except dpctl.SyclQueueCreationError: + pytest.skip("OpenCL queue could not be created") + + oclSrc = ( + "kernel void scale_complex(" + " global float2 *a, global float2 *b, float2 d) {" + " size_t index = get_global_id(0);" + " float2 ai = a[index];" + " b[index] = (float2)(ai.x * d.x - ai.y * d.y," + " ai.x * d.y + ai.y * d.x);" + "}" + ) + kb = dpctl_prog.create_kernel_bundle_from_source(q, oclSrc) + kernel = kb.get_sycl_kernel("scale_complex") + + n_elems = 512 + a = np.ones(n_elems, dtype=np.complex64) * np.complex64(2 + 0j) + b = np.zeros(n_elems, dtype=np.complex64) + + a_usm = dpm.MemoryUSMDevice(a.nbytes, queue=q) + b_usm = dpm.MemoryUSMDevice(b.nbytes, queue=q) + + q.memcpy(dest=a_usm, src=a, count=a.nbytes) + + args = [a_usm, b_usm, np.complex64(0.5 - 1j)] + q.submit(kernel, args, [n_elems]).wait() + + q.memcpy(b, b_usm, b.nbytes) + expected = a * np.complex64(0.5 - 1j) + assert np.allclose(b, expected) + + +def test_kernel_submit_complex128_arg(): + """Test kernel submission with a complex128 scalar argument""" + try: + q = dpctl.SyclQueue("opencl") + except dpctl.SyclQueueCreationError: + pytest.skip("OpenCL queue could not be created") + if not q.sycl_device.has_aspect_fp64: + pytest.skip( + "Device does not support double precision floating point type" + ) + + # OpenCL double2 has the same layout as complex + oclSrc = ( + "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n" + "kernel void scale_complex(" + " global double2 *a, global double2 *b, double2 d) {" + " size_t index = get_global_id(0);" + " double2 ai = a[index];" + " b[index] = (double2)(ai.x * d.x - ai.y * d.y," + " ai.x * d.y + ai.y * d.x);" + "}" + ) + kb = dpctl_prog.create_kernel_bundle_from_source(q, oclSrc) + kernel = kb.get_sycl_kernel("scale_complex") + + n_elems = 512 + a = np.ones(n_elems, dtype=np.complex128) * np.complex128(1 + 2j) + b = np.zeros(n_elems, dtype=np.complex128) + + a_usm = dpm.MemoryUSMDevice(a.nbytes, queue=q) + b_usm = dpm.MemoryUSMDevice(b.nbytes, queue=q) + + q.memcpy(dest=a_usm, src=a, count=a.nbytes) + + d = dpctl.complex128(3 - 1j) + args = [a_usm, b_usm, d] + q.submit(kernel, args, [n_elems]).wait() + + q.memcpy(b, b_usm, b.nbytes) + expected = a * np.complex128(3 - 1j) + assert np.allclose(b, expected) + + +def test_kernel_submit_complex128_python_builtin(): + """Test kernel submission with a Python complex as kernel argument""" + try: + q = dpctl.SyclQueue("opencl") + except dpctl.SyclQueueCreationError: + pytest.skip("OpenCL queue could not be created") + if not q.sycl_device.has_aspect_fp64: + pytest.skip( + "Device does not support double precision floating point type" + ) + + oclSrc = ( + "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n" + "kernel void scale_complex(" + " global double2 *a, global double2 *b, double2 d) {" + " size_t index = get_global_id(0);" + " double2 ai = a[index];" + " b[index] = (double2)(ai.x * d.x - ai.y * d.y," + " ai.x * d.y + ai.y * d.x);" + "}" + ) + kb = dpctl_prog.create_kernel_bundle_from_source(q, oclSrc) + kernel = kb.get_sycl_kernel("scale_complex") + + n_elems = 512 + a = np.ones(n_elems, dtype=np.complex128) * np.complex128(2 + 1j) + b = np.zeros(n_elems, dtype=np.complex128) + + a_usm = dpm.MemoryUSMDevice(a.nbytes, queue=q) + b_usm = dpm.MemoryUSMDevice(b.nbytes, queue=q) + + q.memcpy(dest=a_usm, src=a, count=a.nbytes) + + # Pass a raw Python complex directly + args = [a_usm, b_usm, (1 + 1j)] + q.submit(kernel, args, [n_elems]).wait() + + q.memcpy(b, b_usm, b.nbytes) + expected = a * np.complex128(1 + 1j) + assert np.allclose(b, expected) + + def _check_kernel_arg_type_instance(kati): assert isinstance(kati.name, str) assert isinstance(kati.value, int) @@ -303,6 +688,9 @@ def test_kernel_arg_type(): _check_kernel_arg_type_instance(kernel_arg_type.dpctl_local_accessor) _check_kernel_arg_type_instance(kernel_arg_type.dpctl_work_group_memory) _check_kernel_arg_type_instance(kernel_arg_type.dpctl_raw_kernel_arg) + _check_kernel_arg_type_instance(kernel_arg_type.dpctl_float16) + _check_kernel_arg_type_instance(kernel_arg_type.dpctl_complex64) + _check_kernel_arg_type_instance(kernel_arg_type.dpctl_complex128) def get_spirv_abspath(fn): diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_enum_types.h b/libsyclinterface/include/syclinterface/dpctl_sycl_enum_types.h index e6ee311980..a718b83417 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_enum_types.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_enum_types.h @@ -102,6 +102,9 @@ typedef enum DPCTL_LOCAL_ACCESSOR, DPCTL_WORK_GROUP_MEMORY, DPCTL_RAW_KERNEL_ARG, + DPCTL_FLOAT16_T, + DPCTL_COMPLEX64_T, + DPCTL_COMPLEX128_T, DPCTL_UNSUPPORTED_KERNEL_ARG } DPCTLKernelArgType; diff --git a/libsyclinterface/source/dpctl_sycl_queue_interface.cpp b/libsyclinterface/source/dpctl_sycl_queue_interface.cpp index 575a3c13fa..05d7a5f39c 100644 --- a/libsyclinterface/source/dpctl_sycl_queue_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_queue_interface.cpp @@ -35,6 +35,7 @@ #include #include +#include #include #include #include @@ -113,6 +114,24 @@ using namespace sycl; CGH.set_arg(IDX, la); \ return true; \ } \ + case DPCTL_FLOAT16_T: \ + { \ + auto la = local_accessor(R, CGH); \ + CGH.set_arg(IDX, la); \ + return true; \ + } \ + case DPCTL_COMPLEX64_T: \ + { \ + auto la = local_accessor, NDIM>(R, CGH); \ + CGH.set_arg(IDX, la); \ + return true; \ + } \ + case DPCTL_COMPLEX128_T: \ + { \ + auto la = local_accessor, NDIM>(R, CGH); \ + CGH.set_arg(IDX, la); \ + return true; \ + } \ default: \ error_handler("Kernel argument could not be created.", __FILE__, \ __func__, __LINE__, error_level::error); \ @@ -216,6 +235,15 @@ bool set_kernel_arg(handler &cgh, case DPCTL_FLOAT64_T: cgh.set_arg(idx, *(double *)Arg); break; + case DPCTL_FLOAT16_T: + cgh.set_arg(idx, *(sycl::half *)Arg); + break; + case DPCTL_COMPLEX64_T: + cgh.set_arg(idx, *(std::complex *)Arg); + break; + case DPCTL_COMPLEX128_T: + cgh.set_arg(idx, *(std::complex *)Arg); + break; case DPCTL_VOID_PTR: cgh.set_arg(idx, Arg); break;