diff --git a/CHANGELOG.md b/CHANGELOG.md index ab377b2b69cb..8e77700afcc4 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -12,8 +12,8 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 ### Changed * Changed `dpnp.meshgrid` and `dpnp.tensor.meshgrid` to return a tuple instead of a list, aligning with NumPy 2.5+ behavior and 2025.12 version of the Python array API standard [#2854](https://github.com/IntelPython/dpnp/pull/2854) - * Updated `searchsorted` implementations to align with the 2025.12 array API spec [gh-2902](https://github.com/IntelPython/dpnp/pull/2902) +* Replaced `.pxi` includes in `dpnp.tensor` with modular `.pxd`/`.pyx` Cython imports [#2913](https://github.com/IntelPython/dpnp/pull/2913) ### Deprecated diff --git a/dpnp/backend/include/dpnp4pybind11.hpp b/dpnp/backend/include/dpnp4pybind11.hpp index 6394987a5f82..986392409be7 100644 --- a/dpnp/backend/include/dpnp4pybind11.hpp +++ b/dpnp/backend/include/dpnp4pybind11.hpp @@ -32,9 +32,10 @@ #include "dpctl4pybind11.hpp" // Include generated Cython headers for usm_ndarray -// (struct definition and constants only) #include "dpnp/tensor/_usmarray.h" #include "dpnp/tensor/_usmarray_api.h" +// Include usm_ndarray constants (flags, type numbers) +#include "../../tensor/include/usm_ndarray_constants.h" #include #include @@ -191,7 +192,7 @@ class dpnp_capi this->UsmNDArray_MakeSimpleFromPtr_ = UsmNDArray_MakeSimpleFromPtr; this->UsmNDArray_MakeFromPtr_ = UsmNDArray_MakeFromPtr; - // constants + // constants from usm_ndarray_constants.h this->USM_ARRAY_C_CONTIGUOUS_ = USM_ARRAY_C_CONTIGUOUS; this->USM_ARRAY_F_CONTIGUOUS_ = USM_ARRAY_F_CONTIGUOUS; this->USM_ARRAY_WRITABLE_ = USM_ARRAY_WRITABLE; diff --git a/dpnp/tensor/_slicing.pxd b/dpnp/tensor/_slicing.pxd new file mode 100644 index 000000000000..e556ff57f4a4 --- /dev/null +++ b/dpnp/tensor/_slicing.pxd @@ -0,0 +1,42 @@ +# ***************************************************************************** +# Copyright (c) 2026, Intel Corporation +# All rights reserved. +# +# Redistribution and use in source and binary forms, with or without +# modification, are permitted provided that the following conditions are met: +# - Redistributions of source code must retain the above copyright notice, +# this list of conditions and the following disclaimer. +# - Redistributions in binary form must reproduce the above copyright notice, +# this list of conditions and the following disclaimer in the documentation +# and/or other materials provided with the distribution. +# - Neither the name of the copyright holder nor the names of its contributors +# may be used to endorse or promote products derived from this software +# without specific prior written permission. +# +# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE +# ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE +# LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR +# CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF +# SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS +# INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN +# CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) +# ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF +# THE POSSIBILITY OF SUCH DAMAGE. +# ***************************************************************************** + +# distutils: language = c++ +# cython: language_level=3 + +cdef bint _is_buffer(object o) + +cdef Py_ssize_t _slice_len( + Py_ssize_t sl_start, + Py_ssize_t sl_stop, + Py_ssize_t sl_step +) + +cdef bint _is_integral(object x) except * + +cdef bint _is_boolean(object x) except * diff --git a/dpnp/tensor/_slicing.pxi b/dpnp/tensor/_slicing.pyx similarity index 99% rename from dpnp/tensor/_slicing.pxi rename to dpnp/tensor/_slicing.pyx index 2f22894c4b19..2e785951bc6d 100644 --- a/dpnp/tensor/_slicing.pxi +++ b/dpnp/tensor/_slicing.pyx @@ -26,11 +26,16 @@ # THE POSSIBILITY OF SUCH DAMAGE. # ***************************************************************************** +# distutils: language = c++ +# cython: language_level=3 + import numbers from operator import index from cpython.buffer cimport PyObject_CheckBuffer from numpy import ndarray +from ._usmarray cimport usm_ndarray + cdef bint _is_buffer(object o): return PyObject_CheckBuffer(o) diff --git a/dpnp/tensor/_stride_utils.pxd b/dpnp/tensor/_stride_utils.pxd new file mode 100644 index 000000000000..b4f5eac4cd1c --- /dev/null +++ b/dpnp/tensor/_stride_utils.pxd @@ -0,0 +1,59 @@ +# ***************************************************************************** +# Copyright (c) 2026, Intel Corporation +# All rights reserved. +# +# Redistribution and use in source and binary forms, with or without +# modification, are permitted provided that the following conditions are met: +# - Redistributions of source code must retain the above copyright notice, +# this list of conditions and the following disclaimer. +# - Redistributions in binary form must reproduce the above copyright notice, +# this list of conditions and the following disclaimer in the documentation +# and/or other materials provided with the distribution. +# - Neither the name of the copyright holder nor the names of its contributors +# may be used to endorse or promote products derived from this software +# without specific prior written permission. +# +# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE +# ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE +# LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR +# CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF +# SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS +# INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN +# CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) +# ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF +# THE POSSIBILITY OF SUCH DAMAGE. +# ***************************************************************************** + +# distutils: language = c++ +# cython: language_level=3 + +cdef extern from "usm_ndarray_constants.h": + int USM_ARRAY_C_CONTIGUOUS + int USM_ARRAY_F_CONTIGUOUS + int USM_ARRAY_WRITABLE + +cdef int ERROR_MALLOC +cdef int ERROR_INTERNAL +cdef int ERROR_INCORRECT_ORDER +cdef int ERROR_UNEXPECTED_STRIDES + +cdef Py_ssize_t shape_to_elem_count(int nd, Py_ssize_t *shape_arr) + +cdef int _from_input_shape_strides( + int nd, object shape, object strides, int itemsize, char order, + Py_ssize_t **shape_ptr, Py_ssize_t **strides_ptr, + Py_ssize_t *nelems, Py_ssize_t *min_disp, Py_ssize_t *max_disp, + int *contig +) + +cdef object _make_int_tuple(int nd, const Py_ssize_t *ary) + +cdef object _make_reversed_int_tuple(int nd, const Py_ssize_t *ary) + +cdef object _c_contig_strides(int nd, Py_ssize_t *shape) + +cdef object _f_contig_strides(int nd, Py_ssize_t *shape) + +cdef object _swap_last_two(tuple t) diff --git a/dpnp/tensor/_stride_utils.pxi b/dpnp/tensor/_stride_utils.pyx similarity index 98% rename from dpnp/tensor/_stride_utils.pxi rename to dpnp/tensor/_stride_utils.pyx index 3caf8dd8fd1f..7d770e30e704 100644 --- a/dpnp/tensor/_stride_utils.pxi +++ b/dpnp/tensor/_stride_utils.pyx @@ -29,7 +29,7 @@ # distutils: language = c++ # cython: language_level=3 -from cpython.mem cimport PyMem_Malloc +from cpython.mem cimport PyMem_Free, PyMem_Malloc from cpython.ref cimport Py_INCREF from cpython.tuple cimport PyTuple_New, PyTuple_SetItem @@ -39,10 +39,6 @@ cdef int ERROR_INTERNAL = -1 cdef int ERROR_INCORRECT_ORDER = 2 cdef int ERROR_UNEXPECTED_STRIDES = 3 -cdef int USM_ARRAY_C_CONTIGUOUS = 1 -cdef int USM_ARRAY_F_CONTIGUOUS = 2 -cdef int USM_ARRAY_WRITABLE = 4 - cdef Py_ssize_t shape_to_elem_count(int nd, Py_ssize_t *shape_arr): """ diff --git a/dpnp/tensor/_types.pxd b/dpnp/tensor/_types.pxd new file mode 100644 index 000000000000..5230688d4ea6 --- /dev/null +++ b/dpnp/tensor/_types.pxd @@ -0,0 +1,59 @@ +# ***************************************************************************** +# Copyright (c) 2026, Intel Corporation +# All rights reserved. +# +# Redistribution and use in source and binary forms, with or without +# modification, are permitted provided that the following conditions are met: +# - Redistributions of source code must retain the above copyright notice, +# this list of conditions and the following disclaimer. +# - Redistributions in binary form must reproduce the above copyright notice, +# this list of conditions and the following disclaimer in the documentation +# and/or other materials provided with the distribution. +# - Neither the name of the copyright holder nor the names of its contributors +# may be used to endorse or promote products derived from this software +# without specific prior written permission. +# +# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE +# ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE +# LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR +# CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF +# SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS +# INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN +# CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) +# ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF +# THE POSSIBILITY OF SUCH DAMAGE. +# ***************************************************************************** + +# distutils: language = c++ +# cython: language_level=3 + +cdef extern from "usm_ndarray_constants.h": + int UAR_BOOL + int UAR_BYTE + int UAR_UBYTE + int UAR_SHORT + int UAR_USHORT + int UAR_INT + int UAR_UINT + int UAR_LONG + int UAR_ULONG + int UAR_LONGLONG + int UAR_ULONGLONG + int UAR_FLOAT + int UAR_DOUBLE + int UAR_CFLOAT + int UAR_CDOUBLE + int UAR_TYPE_SENTINEL + int UAR_HALF + +cdef int type_bytesize(int typenum) + +cdef str _make_typestr(int typenum) + +cdef int typenum_from_format(str s) + +cdef int descr_to_typenum(object dtype) + +cdef int dtype_to_typenum(dtype) diff --git a/dpnp/tensor/_types.pxi b/dpnp/tensor/_types.pyx similarity index 86% rename from dpnp/tensor/_types.pxi rename to dpnp/tensor/_types.pyx index 090750658f4b..34cea76cfd0d 100644 --- a/dpnp/tensor/_types.pxi +++ b/dpnp/tensor/_types.pyx @@ -26,25 +26,11 @@ # THE POSSIBILITY OF SUCH DAMAGE. # ***************************************************************************** -# these typenum values are aligned to values in NumPy -cdef: - int UAR_BOOL = 0 # pragma: no cover - int UAR_BYTE = 1 # pragma: no cover - int UAR_UBYTE = 2 # pragma: no cover - int UAR_SHORT = 3 # pragma: no cover - int UAR_USHORT = 4 # pragma: no cover - int UAR_INT = 5 # pragma: no cover - int UAR_UINT = 6 # pragma: no cover - int UAR_LONG = 7 # pragma: no cover - int UAR_ULONG = 8 # pragma: no cover - int UAR_LONGLONG = 9 # pragma: no cover - int UAR_ULONGLONG = 10 # pragma: no cover - int UAR_FLOAT = 11 # pragma: no cover - int UAR_DOUBLE = 12 # pragma: no cover - int UAR_CFLOAT = 14 # pragma: no cover - int UAR_CDOUBLE = 15 # pragma: no cover - int UAR_TYPE_SENTINEL = 17 # pragma: no cover - int UAR_HALF = 23 # pragma: no cover +# distutils: language = c++ +# cython: language_level=3 + +import numpy as np + cdef int type_bytesize(int typenum): """ diff --git a/dpnp/tensor/_usmarray.pxd b/dpnp/tensor/_usmarray.pxd index ccb8f4c796b7..516a31cba711 100644 --- a/dpnp/tensor/_usmarray.pxd +++ b/dpnp/tensor/_usmarray.pxd @@ -32,27 +32,28 @@ cimport dpctl -cdef public api int USM_ARRAY_C_CONTIGUOUS -cdef public api int USM_ARRAY_F_CONTIGUOUS -cdef public api int USM_ARRAY_WRITABLE +cdef extern from "usm_ndarray_constants.h": + int USM_ARRAY_C_CONTIGUOUS + int USM_ARRAY_F_CONTIGUOUS + int USM_ARRAY_WRITABLE -cdef public api int UAR_BOOL -cdef public api int UAR_BYTE -cdef public api int UAR_UBYTE -cdef public api int UAR_SHORT -cdef public api int UAR_USHORT -cdef public api int UAR_INT -cdef public api int UAR_UINT -cdef public api int UAR_LONG -cdef public api int UAR_ULONG -cdef public api int UAR_LONGLONG -cdef public api int UAR_ULONGLONG -cdef public api int UAR_FLOAT -cdef public api int UAR_DOUBLE -cdef public api int UAR_CFLOAT -cdef public api int UAR_CDOUBLE -cdef public api int UAR_TYPE_SENTINEL -cdef public api int UAR_HALF + int UAR_BOOL + int UAR_BYTE + int UAR_UBYTE + int UAR_SHORT + int UAR_USHORT + int UAR_INT + int UAR_UINT + int UAR_LONG + int UAR_ULONG + int UAR_LONGLONG + int UAR_ULONGLONG + int UAR_FLOAT + int UAR_DOUBLE + int UAR_CFLOAT + int UAR_CDOUBLE + int UAR_TYPE_SENTINEL + int UAR_HALF cdef api class usm_ndarray [object PyUSMArrayObject, type PyUSMArrayType]: diff --git a/dpnp/tensor/_usmarray.pyx b/dpnp/tensor/_usmarray.pyx index 7d90ffeb05a1..5c3c9430a848 100644 --- a/dpnp/tensor/_usmarray.pyx +++ b/dpnp/tensor/_usmarray.pyx @@ -46,7 +46,6 @@ from ._print import usm_ndarray_repr, usm_ndarray_str cimport dpctl as c_dpctl cimport dpctl.memory as c_dpmem from cpython.mem cimport PyMem_Free -from cpython.tuple cimport PyTuple_New, PyTuple_SetItem from . cimport _dlpack as c_dlpack @@ -56,9 +55,35 @@ from . import _flags from ._dlpack import get_build_dlpack_version from ._tensor_impl import default_device_fp_type -include "_stride_utils.pxi" -include "_types.pxi" -include "_slicing.pxi" +from ._slicing cimport ( + _is_boolean, + _is_buffer, + _is_integral, + _slice_len, +) + +from ._slicing import _basic_slice_meta + +from ._stride_utils cimport ( + ERROR_INCORRECT_ORDER, + ERROR_INTERNAL, + ERROR_MALLOC, + ERROR_UNEXPECTED_STRIDES, + _c_contig_strides, + _f_contig_strides, + _from_input_shape_strides, + _make_int_tuple, + _make_reversed_int_tuple, + _swap_last_two, + shape_to_elem_count, +) +from ._types cimport ( + _make_typestr, + descr_to_typenum, + dtype_to_typenum, + type_bytesize, + typenum_from_format, +) class DLDeviceType(IntEnum): diff --git a/dpnp/tensor/include/usm_ndarray_constants.h b/dpnp/tensor/include/usm_ndarray_constants.h new file mode 100644 index 000000000000..dbe1f9da8b82 --- /dev/null +++ b/dpnp/tensor/include/usm_ndarray_constants.h @@ -0,0 +1,62 @@ +/* ***************************************************************************** + * Copyright (c) 2026, Intel Corporation + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * - Redistributions of source code must retain the above copyright notice, + * this list of conditions and the following disclaimer. + * - Redistributions in binary form must reproduce the above copyright notice, + * this list of conditions and the following disclaimer in the documentation + * and/or other materials provided with the distribution. + * - Neither the name of the copyright holder nor the names of its contributors + * may be used to endorse or promote products derived from this software + * without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" + * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE + * ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE + * LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR + * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF + * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS + * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN + * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) + * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF + * THE POSSIBILITY OF SUCH DAMAGE. + * ****************************************************************************/ + +#ifndef DPNP_TENSOR_USM_NDARRAY_CONSTANTS_H +#define DPNP_TENSOR_USM_NDARRAY_CONSTANTS_H + +/* Array contiguity flags */ +enum +{ + USM_ARRAY_C_CONTIGUOUS = 1, + USM_ARRAY_F_CONTIGUOUS = 2, + USM_ARRAY_WRITABLE = 4 +}; + +/* These typenum values are aligned to values in NumPy */ +enum +{ + UAR_BOOL = 0, + UAR_BYTE = 1, + UAR_UBYTE = 2, + UAR_SHORT = 3, + UAR_USHORT = 4, + UAR_INT = 5, + UAR_UINT = 6, + UAR_LONG = 7, + UAR_ULONG = 8, + UAR_LONGLONG = 9, + UAR_ULONGLONG = 10, + UAR_FLOAT = 11, + UAR_DOUBLE = 12, + UAR_CFLOAT = 14, + UAR_CDOUBLE = 15, + UAR_TYPE_SENTINEL = 17, + UAR_HALF = 23 +}; + +#endif /* DPNP_TENSOR_USM_NDARRAY_CONSTANTS_H */ diff --git a/dpnp/tensor/libtensor/include/kernels/elementwise_functions/expm1.hpp b/dpnp/tensor/libtensor/include/kernels/elementwise_functions/expm1.hpp index e9e2a704cf0b..735cb1fada58 100644 --- a/dpnp/tensor/libtensor/include/kernels/elementwise_functions/expm1.hpp +++ b/dpnp/tensor/libtensor/include/kernels/elementwise_functions/expm1.hpp @@ -121,6 +121,10 @@ struct Expm1Functor } } + if (x == realT(0) && y == realT(0)) { + return resT{realT(0), y}; + } + // x, y finite numbers const realT cosY_val = sycl::cos(y); const realT sinY_val = (y == 0) ? y : sycl::sin(y); diff --git a/dpnp/tests/tensor/elementwise/test_expm1.py b/dpnp/tests/tensor/elementwise/test_expm1.py index bb665c424564..0e1a8205dab7 100644 --- a/dpnp/tests/tensor/elementwise/test_expm1.py +++ b/dpnp/tests/tensor/elementwise/test_expm1.py @@ -185,3 +185,31 @@ def test_expm1_special_cases(): tol = dpt.finfo(X.dtype).resolution with np.errstate(invalid="ignore"): assert_allclose(dpt.asnumpy(dpt.expm1(X)), res, atol=tol, rtol=tol) + + +@pytest.mark.parametrize("dtype", ["c8", "c16"]) +def test_expm1_zero_special_cases(dtype): + q = get_queue_or_skip() + skip_if_dtype_not_supported(dtype, q) + + x = [ + complex(+0.0, +0.0), + complex(-0.0, +0.0), + complex(+0.0, -0.0), + complex(-0.0, -0.0), + ] + expected = [ + complex(0.0, 0.0), + complex(0.0, 0.0), + complex(0.0, -0.0), + complex(0.0, -0.0), + ] + + xf = dpt.asarray(x, dtype=dtype, sycl_queue=q) + Y = dpt.asnumpy(dpt.expm1(xf)) + + for i in range(len(x)): + assert Y[i].real == expected[i].real + assert Y[i].imag == expected[i].imag + assert not np.signbit(Y[i].real) + assert np.signbit(Y[i].imag) == np.signbit(expected[i].imag)