Merge branch 'master' into feature/dlpack-kdlcpu-support

Author: oleksandr-pavlyk

.github/workflows/CODEOWNERS

@@ -0,0 +1 @@
+*       @oleksandr-pavlyk @ndgrigorian

.github/workflows/conda-package.yml

@@ -62,12 +62,12 @@ jobs:
             $CHANNELS \
             conda-recipe
       - name: Upload artifact
-        uses: actions/[email protected].4
+        uses: actions/[email protected].5
         with:
           name: ${{ env.PACKAGE_NAME }} ${{ runner.os }} Python ${{ matrix.python }}
           path: /usr/share/miniconda/conda-bld/linux-64/${{ env.PACKAGE_NAME }}-*.tar.bz2
       - name: Upload wheels artifact
-        uses: actions/[email protected].4
+        uses: actions/[email protected].5
         with:
           name: ${{ env.PACKAGE_NAME }} ${{ runner.os }} Wheels Python ${{ matrix.python }}
           path: ${{ env.WHEELS_OUTPUT_FOLDER }}${{ env.PACKAGE_NAME }}-*.whl
@@ -111,12 +111,12 @@ jobs:
           OVERRIDE_INTEL_IPO: 1   # IPO requires more resources that GH actions VM provides
         run: conda build --no-test --python ${{ matrix.python }} -c dppy/label/bootstrap -c ${{ env.INTEL_CHANNEL }} -c conda-forge --override-channels conda-recipe
       - name: Upload artifact
-        uses: actions/[email protected].4
+        uses: actions/[email protected].5
         with:
           name: ${{ env.PACKAGE_NAME }} ${{ runner.os }} Python ${{ matrix.python }}
           path: ${{ env.conda-bld }}${{ env.PACKAGE_NAME }}-*.tar.bz2
       - name: Upload wheels artifact
-        uses: actions/[email protected].4
+        uses: actions/[email protected].5
         with:
           name: ${{ env.PACKAGE_NAME }} ${{ runner.os }} Wheels Python ${{ matrix.python }}
           path: ${{ env.WHEELS_OUTPUT_FOLDER }}${{ env.PACKAGE_NAME }}-*.whl

.github/workflows/generate-docs.yml

@@ -104,7 +104,7 @@ jobs:
           git push tokened_docs gh-pages
       - name: Save built docs as an artifact
         if: ${{ github.event.pull_request && github.event.pull_request.head.repo.fork && github.event.action != 'closed'}}
-        uses: actions/[email protected].4
+        uses: actions/[email protected].5
         with:
           name: ${{ env.PACKAGE_NAME }} rendered documentation
           path: ~/docs

.github/workflows/openssf-scorecard.yml

@@ -60,7 +60,7 @@ jobs:
       # Upload the results as artifacts (optional). Commenting out will disable uploads of run results in SARIF
       # format to the repository Actions tab.
       - name: "Upload artifact"
-        uses: actions/upload-artifact@0b2256b8c012f0828dc542b3febcab082c67f72b # v4.3.4
+        uses: actions/upload-artifact@89ef406dd8d7e03cfd12d9e0a4a378f454709029 # v4.3.5
         with:
           name: SARIF file
           path: results.sarif

conda-recipe/build.sh

@@ -27,26 +27,6 @@ CMAKE_ARGS="${CMAKE_ARGS} -DDPCTL_LEVEL_ZERO_INCLUDE_DIR=${PREFIX}/include/level
 # -wnx flags mean: --wheel --no-isolation --skip-dependency-check
 ${PYTHON} -m build -w -n -x
 
-pushd dist
-${PYTHON} -m wheel unpack -d dpctl_wheel dpctl*.whl
-export lib_name=libDPCTLSyclInterface
-export so_full_path=$(find dpctl_wheel -regextype posix-extended -regex "^.*${lib_name}\.so")
-export sox_full_path=$(find dpctl_wheel -regextype posix-extended -regex "^.*${lib_name}\.so\.[0-9]*$")
-export soxxx_full_path=$(find dpctl_wheel -regextype posix-extended -regex "^.*${lib_name}\.so\.[0-9]*\.[0-9]*$")
-
-rm -rf ${so_full_path} ${soxxx_full_path}
-
-export so_name=$(basename ${so_full_path})
-export sox_name=$(basename ${sox_full_path})
-export soxxx_name=$(basename ${soxxx_full_path})
-export wheel_path=$(dirname $(dirname ${so_full_path}))
-
-# deal with hard copies
-${PYTHON} -m wheel pack ${wheel_path}
-
-rm -rf dpctl_wheel
-popd
-
 ${PYTHON} -m wheel tags --remove --build "$GIT_DESCRIBE_NUMBER" \
     --platform-tag "manylinux_${GLIBC_MAJOR}_${GLIBC_MINOR}_x86_64" \
     dist/dpctl*.whl
@@ -59,15 +39,6 @@ ${PYTHON} -m pip install dist/dpctl*.whl \
     --prefix "${PREFIX}" \
     -vv
 
-export libdir=$(find $PREFIX -name 'libDPCTLSyclInterface*' -exec dirname \{\} \;)
-
-# Recover symbolic links
-# libDPCTLSyclInterface.so.0 -> libDPCTLSyclInterface.so.0.17
-# libDPCTLSyclInterface.so -> libDPCTLSyclInterface.so.0
-mv ${libdir}/${sox_name} ${libdir}/${soxxx_name}
-ln -s ${libdir}/${soxxx_name} ${libdir}/${sox_name}
-ln -s ${libdir}/${sox_name} ${libdir}/${so_name}
-
 # Copy wheel package
 if [[ -d "${WHEELS_OUTPUT_FOLDER}" ]]; then
     cp dist/dpctl*.whl "${WHEELS_OUTPUT_FOLDER[@]}"

docs/doc_sources/api_reference/dpctl/tensor.indexing_functions.rst

@@ -15,3 +15,4 @@ by either integral arrays of indices or boolean mask arrays.
     place
     put
     take
+    take_along_axis

dpctl/apis/include/dpctl4pybind11.hpp

@@ -27,6 +27,7 @@
 
 #include "dpctl_capi.h"
 #include <complex>
+#include <exception>
 #include <memory>
 #include <pybind11/pybind11.h>
 #include <sycl/sycl.hpp>
@@ -748,6 +749,53 @@ class usm_memory : public py::object
             throw py::error_already_set();
     }
 
+    /*! @brief Create usm_memory object from shared pointer that manages
+     *  lifetime of the USM allocation.
+     */
+    usm_memory(void *usm_ptr,
+               size_t nbytes,
+               const sycl::queue &q,
+               std::shared_ptr<void> shptr)
+    {
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
+        DPCTLSyclUSMRef usm_ref = reinterpret_cast<DPCTLSyclUSMRef>(usm_ptr);
+        auto q_uptr = std::make_unique<sycl::queue>(q);
+        DPCTLSyclQueueRef QRef =
+            reinterpret_cast<DPCTLSyclQueueRef>(q_uptr.get());
+
+        auto vacuous_destructor = []() {};
+        py::capsule mock_owner(vacuous_destructor);
+
+        // create memory object owned by mock_owner, it is a new reference
+        PyObject *_memory =
+            api.Memory_Make_(usm_ref, nbytes, QRef, mock_owner.ptr());
+        auto ref_count_decrementer = [](PyObject *o) noexcept { Py_DECREF(o); };
+
+        using py_uptrT =
+            std::unique_ptr<PyObject, decltype(ref_count_decrementer)>;
+
+        if (!_memory) {
+            throw py::error_already_set();
+        }
+
+        auto memory_uptr = py_uptrT(_memory, ref_count_decrementer);
+        std::shared_ptr<void> *opaque_ptr = new std::shared_ptr<void>(shptr);
+
+        Py_MemoryObject *memobj = reinterpret_cast<Py_MemoryObject *>(_memory);
+        // replace mock_owner capsule as the owner
+        memobj->refobj = Py_None;
+        // set opaque ptr field, usm_memory now knowns that USM is managed
+        // by smart pointer
+        memobj->_opaque_ptr = reinterpret_cast<void *>(opaque_ptr);
+
+        // _memory will delete created copies of sycl::queue, and
+        // std::shared_ptr and the deleter of the shared_ptr<void> is
+        // supposed to free the USM allocation
+        m_ptr = _memory;
+        q_uptr.release();
+        memory_uptr.release();
+    }
+
     sycl::queue get_queue() const
     {
         Py_MemoryObject *mem_obj = reinterpret_cast<Py_MemoryObject *>(m_ptr);

dpctl/memory/_opaque_smart_ptr.hpp

@@ -34,6 +34,7 @@
 #include "syclinterface/dpctl_sycl_types.h"
 #include <memory>
 #include <sycl/sycl.hpp>
+#include <utility>
 
 #include <exception>
 #include <iostream>

dpctl/tensor/__init__.py

@@ -60,7 +60,14 @@
 )
 from dpctl.tensor._device import Device
 from dpctl.tensor._dlpack import from_dlpack
-from dpctl.tensor._indexing_functions import extract, nonzero, place, put, take
+from dpctl.tensor._indexing_functions import (
+    extract,
+    nonzero,
+    place,
+    put,
+    take,
+    take_along_axis,
+)
 from dpctl.tensor._linear_algebra_functions import (
     matmul,
     matrix_transpose,
@@ -377,4 +384,5 @@
     "diff",
     "count_nonzero",
     "DLDeviceType",
+    "take_along_axis",
 ]

dpctl/tensor/_copy_utils.py

@@ -795,13 +795,18 @@ def _nonzero_impl(ary):
     return res
 
 
-def _take_multi_index(ary, inds, p):
+def _take_multi_index(ary, inds, p, mode=0):
     if not isinstance(ary, dpt.usm_ndarray):
         raise TypeError(
             f"Expecting type dpctl.tensor.usm_ndarray, got {type(ary)}"
         )
     ary_nd = ary.ndim
     p = normalize_axis_index(operator.index(p), ary_nd)
+    mode = operator.index(mode)
+    if mode not in [0, 1]:
+        raise ValueError(
+            "Invalid value for mode keyword, only 0 or 1 is supported"
+        )
     queues_ = [
         ary.sycl_queue,
     ]
@@ -860,7 +865,7 @@ def _take_multi_index(ary, inds, p):
         ind=inds,
         dst=res,
         axis_start=p,
-        mode=0,
+        mode=mode,
         sycl_queue=exec_q,
         depends=dep_ev,
     )

dpctl/tensor/_indexing_functions.py

@@ -21,7 +21,7 @@
 import dpctl.tensor._tensor_impl as ti
 import dpctl.utils
 
-from ._copy_utils import _extract_impl, _nonzero_impl
+from ._copy_utils import _extract_impl, _nonzero_impl, _take_multi_index
 from ._numpy_helper import normalize_axis_index
 
 
@@ -423,3 +423,82 @@ def nonzero(arr):
     if arr.ndim == 0:
         raise ValueError("Array of positive rank is expected")
     return _nonzero_impl(arr)
+
+
+def _range(sh_i, i, nd, q, usm_t, dt):
+    ind = dpt.arange(sh_i, dtype=dt, usm_type=usm_t, sycl_queue=q)
+    ind.shape = tuple(sh_i if i == j else 1 for j in range(nd))
+    return ind
+
+
+def take_along_axis(x, indices, /, *, axis=-1, mode="wrap"):
+    """
+    Returns elements from an array at the one-dimensional indices specified
+    by ``indices`` along a provided ``axis``.
+
+    Args:
+        x (usm_ndarray):
+            input array. Must be compatible with ``indices``, except for the
+            axis (dimension) specified by ``axis``.
+        indices (usm_ndarray):
+            array indices. Must have the same rank (i.e., number of dimensions)
+            as ``x``.
+        axis: int
+            axis along which to select values. If ``axis`` is negative, the
+            function determines the axis along which to select values by
+            counting from the last dimension. Default: ``-1``.
+        mode (str, optional):
+            How out-of-bounds indices will be handled. Possible values
+            are:
+
+            - ``"wrap"``: clamps indices to (``-n <= i < n``), then wraps
+              negative indices.
+            - ``"clip"``: clips indices to (``0 <= i < n``).
+
+            Default: ``"wrap"``.
+
+    Returns:
+        usm_ndarray:
+            an array having the same data type as ``x``. The returned array has
+            the same rank (i.e., number of dimensions) as ``x`` and a shape
+            determined according to :ref:`broadcasting`, except for the axis
+            (dimension) specified by ``axis`` whose size must equal the size
+            of the corresponding axis (dimension) in ``indices``.
+
+    Note:
+        Treatment of the out-of-bound indices in ``indices`` array is controlled
+        by the value of ``mode`` keyword.
+    """
+    if not isinstance(x, dpt.usm_ndarray):
+        raise TypeError(f"Expected dpctl.tensor.usm_ndarray, got {type(x)}")
+    if not isinstance(indices, dpt.usm_ndarray):
+        raise TypeError(
+            f"Expected dpctl.tensor.usm_ndarray, got {type(indices)}"
+        )
+    x_nd = x.ndim
+    if x_nd != indices.ndim:
+        raise ValueError(
+            "Number of dimensions in the first and the second "
+            "argument arrays must be equal"
+        )
+    pp = normalize_axis_index(operator.index(axis), x_nd)
+    out_usm_type = dpctl.utils.get_coerced_usm_type(
+        (x.usm_type, indices.usm_type)
+    )
+    exec_q = dpctl.utils.get_execution_queue((x.sycl_queue, indices.sycl_queue))
+    if exec_q is None:
+        raise dpctl.utils.ExecutionPlacementError(
+            "Execution placement can not be unambiguously inferred "
+            "from input arguments. "
+        )
+    mode_i = _get_indexing_mode(mode)
+    indexes_dt = ti.default_device_index_type(exec_q.sycl_device)
+    _ind = tuple(
+        (
+            indices
+            if i == pp
+            else _range(x.shape[i], i, x_nd, exec_q, out_usm_type, indexes_dt)
+        )
+        for i in range(x_nd)
+    )
+    return _take_multi_index(x, _ind, 0, mode=mode_i)

dpctl/tensor/_usmarray.pyx

@@ -621,6 +621,22 @@ cdef class usm_ndarray:
 
     @shape.setter
     def shape(self, new_shape):
+        """
+        Modifies usm_ndarray instance in-place by changing its metadata
+        about the shape and the strides of the array, or raises
+        `AttributeError` exception if in-place change is not possible.
+
+        Args:
+            new_shape: (tuple, int)
+                New shape. Only non-negative values are supported.
+                The new shape may not lead to the change in the
+                number of elements in the array.
+
+        Whether the array can be reshape in-place depends on its
+        strides. Use :func:`dpctl.tensor.reshape` function which
+        always succeeds to reshape the array by performing a copy
+        if necessary.
+        """
         cdef int new_nd = -1
         cdef Py_ssize_t nelems = -1
         cdef int err = 0
@@ -634,7 +650,11 @@ cdef class usm_ndarray:
 
         from ._reshape import reshaped_strides
 
-        new_nd = len(new_shape)
+        try:
+            new_nd = len(new_shape)
+        except TypeError:
+            new_nd = 1
+            new_shape = (new_shape,)
         try:
             new_shape = tuple(operator.index(dim) for dim in new_shape)
         except TypeError:

dpctl/tensor/_utility_functions.py

@@ -453,7 +453,8 @@ def diff(x, /, *, axis=-1, n=1, prepend=None, append=None):
     x_nd = x.ndim
     axis = normalize_axis_index(operator.index(axis), x_nd)
     n = operator.index(n)
-
+    if n < 0:
+        raise ValueError(f"`n` must be positive, got {n}")
     arr = _concat_diff_input(x, axis, prepend, append)
     if n == 0:
         return arr

dpctl/tests/test_service.py

@@ -163,9 +163,11 @@ def test_syclinterface(should_skip_syclinterface):
         if should_skip_syclinterface:
             pass
         else:
-            assert len(paths) > 1 and any(
-                [os.path.islink(fn) for fn in paths]
-            ), "All library instances are hard links"
+            len_p = len(paths)
+            n_symlinks = sum([int(os.path.islink(fn)) for fn in paths])
+            assert len_p == 1 + n_symlinks, (
+                "Multiple library instances are hard links",
+            )
     elif sys.platform in ["win32", "cygwin"]:
         exts = []
         for fn in paths:

dpctl/tests/test_tensor_diff.py

@@ -313,3 +313,17 @@ def test_diff_input_validation():
         dpt.diff,
         bad_in,
     )
+
+
+def test_diff_positive_order():
+    get_queue_or_skip()
+
+    x = dpt.ones(1, dtype="i4")
+    n = -1
+    assert_raises_regex(
+        ValueError,
+        ".*must be positive.*",
+        dpt.diff,
+        x,
+        n=n,
+    )