Merge remote-tracking branch 'origin/master' into enable_amd_build

Author: vlad-perevezentsev

CHANGELOG.md

@@ -16,7 +16,7 @@ This release achieves 100% compliance with Python Array API specification (revis
 * Added implementation of `dpnp.bartlett` [#2366](https://github.com/IntelPython/dpnp/pull/2366)
 * Added implementation of `dpnp.convolve` [#2205](https://github.com/IntelPython/dpnp/pull/2205)
 * Added implementation of `dpnp.kaiser` [#2387](https://github.com/IntelPython/dpnp/pull/2387)
-* Support for compiling `dpnp` for specified AMD GPU architecture with use of [CodePlay oneAPI plug-in](https://developer.codeplay.com/products/oneapi/amd/home/) [#2302](https://github.com/IntelPython/dpnp/pull/2302)
+* Added implementation of `dpnp.interp` [#2417](https://github.com/IntelPython/dpnp/pull/2417)
 
 ### Changed
 
@@ -30,6 +30,12 @@ This release achieves 100% compliance with Python Array API specification (revis
 * Removed `einsum_call` keyword from `dpnp.einsum_path` signature [#2421](https://github.com/IntelPython/dpnp/pull/2421)
 * Changed `"max dimensions"` to `None` in array API capabilities [#2432](https://github.com/IntelPython/dpnp/pull/2432)
 * Updated kernel header `i0.hpp` to expose `cyl_bessel_i0` function depending on build target [#2440](https://github.com/IntelPython/dpnp/pull/2440)
+* Added MKL functions `arg`, `copysign`, `i0`, and `inv` from VM namespace to be used by implementation of the appropriate element-wise functions [#2445](https://github.com/IntelPython/dpnp/pull/2445)
+* Clarified details about conda install instructions in `Quick start quide` and `README` [#2446](https://github.com/IntelPython/dpnp/pull/2446)
+* Bumped oneMKL version up to `0.7` [#2448](https://github.com/IntelPython/dpnp/pull/2448)
+* The parameter `axis` in `dpnp.take_along_axis` function has now a default value of `-1` [#2442](https://github.com/IntelPython/dpnp/pull/2442)
+* Updates the list of required python versions documented in `Quick Start Guide` [#2449](https://github.com/IntelPython/dpnp/pull/2449)
+* Updated FFT module to ensure an input array is Hermitian before calling complex-to-real FFT [#2444](https://github.com/IntelPython/dpnp/pull/2444)
 
 ### Fixed
 

CMakeLists.txt

@@ -157,16 +157,16 @@ if(_use_onemkl_interfaces)
     endif()
 
     if(DPNP_ONEMKL_INTERFACES_DIR)
-        FetchContent_Declare(onemkl_interfaces_library SOURCE_DIR "${DPNP_ONEMKL_INTERFACES_DIR}")
+        FetchContent_Declare(onemath_library SOURCE_DIR "${DPNP_ONEMKL_INTERFACES_DIR}")
     else()
         FetchContent_Declare(
-                onemkl_interfaces_library
+            onemath_library
                 GIT_REPOSITORY https://github.com/uxlfoundation/oneMath.git
-                GIT_TAG 8f4312ef966420b9b8b4b82b9d5c22e2c91a1fe7  # v0.6
+                GIT_TAG 20ba6fd7ae4af6ed693246cfd22c343e6522edbe  # v0.7
         )
     endif()
 
-    FetchContent_MakeAvailable(onemkl_interfaces_library)
+    FetchContent_MakeAvailable(onemath_library)
     if(TARGET onemath)
         set(MKL_INTERFACES_LIB "onemath" CACHE INTERNAL "OneMath lib target")
     elseif(TARGET onemkl)

README.md

@@ -41,7 +41,7 @@ To get the library from the latest release, follow the instructions from
 To install `dpnp` from the Intel(R) conda channel, use the following command:
 
 ```bash
-conda install dpnp -c https://software.repos.intel.com/python/conda/ -c conda-forge
+conda install dpnp -c https://software.repos.intel.com/python/conda/ -c conda-forge --override-channels
 ```
 
 ## Pip
@@ -60,7 +60,7 @@ To try out the latest features, install `dpnp` using our development channel on
 Anaconda cloud:
 
 ```bash
-conda install dpnp -c dppy/label/dev -c https://software.repos.intel.com/python/conda/ -c conda-forge
+conda install dpnp -c dppy/label/dev -c https://software.repos.intel.com/python/conda/ -c conda-forge --override-channels
 ```
 
 

doc/quick_start_guide.rst

@@ -24,7 +24,7 @@ Follow device driver installation instructions to complete the step.
 Python Interpreter
 ==================
 
-You will need Python 3.8, 3.9, or 3.10 installed on your system. If you
+You will need Python 3.9, 3.10, 3.11 or 3.12 installed on your system. If you
 do not have one yet the easiest way to do that is to install
 `Intel Distribution for Python*`_. It installs all essential Python numerical
 and machine learning packages optimized for the Intel hardware, including
@@ -42,14 +42,20 @@ Install Package from Intel(R) channel
 
 You will need one of the commands below:
 
-* Conda: ``conda install dpnp -c https://software.repos.intel.com/python/conda/ -c conda-forge``
+* Conda: ``conda install dpnp -c https://software.repos.intel.com/python/conda/ -c conda-forge --override-channels``
 
 * Pip: ``python -m pip install --index-url https://software.repos.intel.com/python/pypi dpnp``
 
 These commands install dpnp package along with its dependencies, including
 ``dpctl`` package with `Data Parallel Control Library`_ and all required
 compiler runtimes and OneMKL.
 
+.. warning::
+    Packages from the Intel channel are meant to be used together with dependencies from the **conda-forge** channel, and might not
+    work correctly when used in an environment where packages from the ``anaconda`` default channel have been installed. It is
+    advisable to use the `miniforge <https://github.com/conda-forge/miniforge>`__ installer for ``conda``/``mamba``, as it comes with
+    ``conda-forge`` as the only default channel.
+
 .. note::
    Before installing with conda or pip it is strongly advised to update ``conda`` and ``pip`` to latest versions
 
@@ -68,7 +74,7 @@ And to build dpnp package from the sources:
 
 .. code-block:: bash
 
-    conda build conda-recipe -c https://software.repos.intel.com/python/conda/ -c conda-forge
+    conda build conda-recipe -c https://software.repos.intel.com/python/conda/ -c conda-forge --override-channels
 
 Finally, to install the result package:
 
@@ -90,7 +96,7 @@ On Linux:
 
     conda create -n build-env dpctl cython dpcpp_linux-64 mkl-devel-dpcpp tbb-devel \
           onedpl-devel cmake scikit-build ninja pytest intel-gpu-ocl-icd-system     \
-          -c dppy/label/dev -c https://software.repos.intel.com/python/conda/ -c conda-forge
+          -c dppy/label/dev -c https://software.repos.intel.com/python/conda/ -c conda-forge --override-channels
     conda activate build-env
 
 On Windows:
@@ -99,7 +105,7 @@ On Windows:
 
     conda create -n build-env dpctl cython dpcpp_win-64 mkl-devel-dpcpp tbb-devel \
           onedpl-devel cmake scikit-build ninja pytest intel-gpu-ocl-icd-system   \
-          -c dppy/label/dev -c https://software.repos.intel.com/python/conda/ -c conda-forge
+          -c dppy/label/dev -c https://software.repos.intel.com/python/conda/ -c conda-forge --override-channels
     conda activate build-env
 
 To build and install the package on Linux OS, run:

dpnp/backend/extensions/common/ext/common.hpp

@@ -106,6 +106,27 @@ struct IsNan
     }
 };
 
+template <typename T, bool hasValueType>
+struct value_type_of_impl;
+
+template <typename T>
+struct value_type_of_impl<T, false>
+{
+    using type = T;
+};
+
+template <typename T>
+struct value_type_of_impl<T, true>
+{
+    using type = typename T::value_type;
+};
+
+template <typename T>
+using value_type_of = value_type_of_impl<T, type_utils::is_complex_v<T>>;
+
+template <typename T>
+using value_type_of_t = typename value_type_of<T>::type;
+
 size_t get_max_local_size(const sycl::device &device);
 size_t get_max_local_size(const sycl::device &device,
                           int cpu_local_size_limit,

dpnp/backend/extensions/common/ext/details/validation_utils_internal.hpp

@@ -23,9 +23,17 @@
 // THE POSSIBILITY OF SUCH DAMAGE.
 //*****************************************************************************
 
+#include <pybind11/numpy.h>
+#include <pybind11/pybind11.h>
+
+#include "ext/common.hpp"
+
 #include "ext/validation_utils.hpp"
 #include "utils/memory_overlap.hpp"
 
+namespace td_ns = dpctl::tensor::type_dispatch;
+namespace common = ext::common;
+
 namespace ext::validation
 {
 inline sycl::queue get_queue(const std::vector<array_ptr> &inputs,
@@ -137,6 +145,15 @@ inline void check_num_dims(const array_ptr &arr,
     }
 }
 
+inline void check_num_dims(const std::vector<array_ptr> &arrays,
+                           const size_t ndim,
+                           const array_names &names)
+{
+    for (const auto &arr : arrays) {
+        check_num_dims(arr, ndim, names);
+    }
+}
+
 inline void check_max_dims(const array_ptr &arr,
                            const size_t max_ndim,
                            const array_names &names)
@@ -163,6 +180,103 @@ inline void check_size_at_least(const array_ptr &arr,
     }
 }
 
+inline void check_has_dtype(const array_ptr &arr,
+                            const typenum_t dtype,
+                            const array_names &names)
+{
+    if (arr == nullptr) {
+        return;
+    }
+
+    auto array_types = td_ns::usm_ndarray_types();
+    int array_type_id = array_types.typenum_to_lookup_id(arr->get_typenum());
+    int expected_type_id = static_cast<int>(dtype);
+
+    if (array_type_id != expected_type_id) {
+        py::dtype actual_dtype = common::dtype_from_typenum(array_type_id);
+        py::dtype dtype_py = common::dtype_from_typenum(expected_type_id);
+
+        std::string msg = "Array " + name_of(arr, names) + " must have dtype " +
+                          std::string(py::str(dtype_py)) + ", but got " +
+                          std::string(py::str(actual_dtype));
+
+        throw py::value_error(msg);
+    }
+}
+
+inline void check_same_dtype(const array_ptr &arr1,
+                             const array_ptr &arr2,
+                             const array_names &names)
+{
+    if (arr1 == nullptr || arr2 == nullptr) {
+        return;
+    }
+
+    auto array_types = td_ns::usm_ndarray_types();
+    int first_type_id = array_types.typenum_to_lookup_id(arr1->get_typenum());
+    int second_type_id = array_types.typenum_to_lookup_id(arr2->get_typenum());
+
+    if (first_type_id != second_type_id) {
+        py::dtype first_dtype = common::dtype_from_typenum(first_type_id);
+        py::dtype second_dtype = common::dtype_from_typenum(second_type_id);
+
+        std::string msg = "Arrays " + name_of(arr1, names) + " and " +
+                          name_of(arr2, names) +
+                          " must have the same dtype, but got " +
+                          std::string(py::str(first_dtype)) + " and " +
+                          std::string(py::str(second_dtype));
+
+        throw py::value_error(msg);
+    }
+}
+
+inline void check_same_dtype(const std::vector<array_ptr> &arrays,
+                             const array_names &names)
+{
+    if (arrays.empty()) {
+        return;
+    }
+
+    const auto *first = arrays[0];
+    for (size_t i = 1; i < arrays.size(); ++i) {
+        check_same_dtype(first, arrays[i], names);
+    }
+}
+
+inline void check_same_size(const array_ptr &arr1,
+                            const array_ptr &arr2,
+                            const array_names &names)
+{
+    if (arr1 == nullptr || arr2 == nullptr) {
+        return;
+    }
+
+    auto size1 = arr1->get_size();
+    auto size2 = arr2->get_size();
+
+    if (size1 != size2) {
+        std::string msg =
+            "Arrays " + name_of(arr1, names) + " and " + name_of(arr2, names) +
+            " must have the same size, but got " + std::to_string(size1) +
+            " and " + std::to_string(size2);
+
+        throw py::value_error(msg);
+    }
+}
+
+inline void check_same_size(const std::vector<array_ptr> &arrays,
+                            const array_names &names)
+{
+    if (arrays.empty()) {
+        return;
+    }
+
+    auto first = arrays[0];
+    for (size_t i = 1; i < arrays.size(); ++i) {
+        check_same_size(first, arrays[i], names);
+    }
+}
+
 inline void common_checks(const std::vector<array_ptr> &inputs,
                           const std::vector<array_ptr> &outputs,
                           const array_names &names)

dpnp/backend/extensions/common/ext/validation_utils.hpp

@@ -35,6 +35,7 @@ namespace ext::validation
 {
 using array_ptr = const dpctl::tensor::usm_ndarray *;
 using array_names = std::unordered_map<array_ptr, std::string>;
+using dpctl::tensor::type_dispatch::typenum_t;
 
 std::string name_of(const array_ptr &arr, const array_names &names);
 
@@ -56,6 +57,9 @@ void check_no_overlap(const std::vector<array_ptr> &inputs,
 void check_num_dims(const array_ptr &arr,
                     const size_t ndim,
                     const array_names &names);
+void check_num_dims(const std::vector<array_ptr> &arrays,
+                    const size_t ndim,
+                    const array_names &names);
 void check_max_dims(const array_ptr &arr,
                     const size_t max_ndim,
                     const array_names &names);
@@ -64,6 +68,20 @@ void check_size_at_least(const array_ptr &arr,
                          const size_t size,
                          const array_names &names);
 
+void check_has_dtype(const array_ptr &arr,
+                     const typenum_t dtype,
+                     const array_names &names);
+
+void check_same_dtype(const array_ptr &arr1,
+                      const array_ptr &arr2,
+                      const array_names &names);
+
+void check_same_size(const array_ptr &arr1,
+                     const array_ptr &arr2,
+                     const array_names &names);
+void check_same_size(const std::vector<array_ptr> &arrays,
+                     const array_names &names);
+
 void common_checks(const std::vector<array_ptr> &inputs,
                    const std::vector<array_ptr> &outputs,
                    const array_names &names);

dpnp/backend/extensions/ufunc/CMakeLists.txt

@@ -36,6 +36,7 @@ set(_elementwise_sources
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/gcd.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/heaviside.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/i0.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/interpolate.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/lcm.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/ldexp.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/logaddexp2.cpp
@@ -69,6 +70,7 @@ endif()
 set_target_properties(${python_module_name} PROPERTIES CMAKE_POSITION_INDEPENDENT_CODE ON)
 
 target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../)
+target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../common)
 
 target_include_directories(${python_module_name} PUBLIC ${Dpctl_INCLUDE_DIR})
 target_include_directories(${python_module_name} PUBLIC ${Dpctl_TENSOR_INCLUDE_DIR})

dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp

@@ -36,6 +36,7 @@
 #include "gcd.hpp"
 #include "heaviside.hpp"
 #include "i0.hpp"
+#include "interpolate.hpp"
 #include "lcm.hpp"
 #include "ldexp.hpp"
 #include "logaddexp2.hpp"
@@ -64,6 +65,7 @@ void init_elementwise_functions(py::module_ m)
     init_gcd(m);
     init_heaviside(m);
     init_i0(m);
+    init_interpolate(m);
     init_lcm(m);
     init_ldexp(m);
     init_logaddexp2(m);