Merge branch 'master' into update-ndarray-docstrings

Author: antonwolfy

.pre-commit-config.yaml

@@ -66,7 +66,7 @@ repos:
         name: isort (pyi)
         types: [pyi]
 -   repo: https://github.com/pycqa/flake8
-    rev: 7.1.2
+    rev: 7.2.0
     hooks:
     -   id: flake8
         args: ["--config=.flake8"]
@@ -79,7 +79,7 @@ repos:
     -   id: clang-format
         args: ["-i"]
 -   repo: https://github.com/gitleaks/gitleaks
-    rev: v8.24.0
+    rev: v8.26.0
     hooks:
     -   id: gitleaks
 -   repo: https://github.com/jumanjihouse/pre-commit-hooks

CHANGELOG.md

@@ -16,6 +16,8 @@ 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)
+* Added implementation of `dpnp.interp` [#2417](https://github.com/IntelPython/dpnp/pull/2417)
+* Added support to build `dpnp` for specified AMD GPU architecture using [CodePlay oneAPI plug-in](https://developer.codeplay.com/products/oneapi/amd/home/) [#2302](https://github.com/IntelPython/dpnp/pull/2302)
 
 ### Changed
 
@@ -27,7 +29,16 @@ This release achieves 100% compliance with Python Array API specification (revis
 * Updated `dpnp.einsum` to add support for `order=None` [#2411](https://github.com/IntelPython/dpnp/pull/2411)
 * Updated Python Array API specification version supported to `2024.12` [#2416](https://github.com/IntelPython/dpnp/pull/2416)
 * Removed `einsum_call` keyword from `dpnp.einsum_path` signature [#2421](https://github.com/IntelPython/dpnp/pull/2421)
+* Updated `dpnp.vdot` to return a 0-D array when one of the inputs is a scalar [#2295](https://github.com/IntelPython/dpnp/pull/2295)
+* Updated `dpnp.outer` to return the same dtype as NumPy when multiplying an array with a scalar [#2295](https://github.com/IntelPython/dpnp/pull/2295)
 * 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

@@ -75,27 +75,60 @@ option(DPNP_USE_ONEMKL_INTERFACES
     "Build DPNP with oneMKL Interfaces"
     OFF
 )
+set(HIP_TARGETS "" CACHE STRING "HIP architecture for target")
+
 set(_dpnp_sycl_targets)
+set(_use_onemkl_interfaces OFF)
 set(_use_onemkl_interfaces_cuda OFF)
+set(_use_onemkl_interfaces_hip OFF)
+
+set(_dpnp_sycl_target_compile_options)
+set(_dpnp_sycl_target_link_options)
+
 if ("x${DPNP_SYCL_TARGETS}" STREQUAL "x")
-   if(DPNP_TARGET_CUDA)
-      set(_dpnp_sycl_targets "nvptx64-nvidia-cuda,spir64-unknown-unknown")
-      set(_use_onemkl_interfaces_cuda ON)
-   else()
-      if(DEFINED ENV{DPNP_TARGET_CUDA})
-          set(_dpnp_sycl_targets "nvptx64-nvidia-cuda,spir64-unknown-unknown")
-          set(_use_onemkl_interfaces_cuda ON)
-      endif()
-   endif()
+    if(DPNP_TARGET_CUDA)
+        set(_dpnp_sycl_targets "nvptx64-nvidia-cuda,spir64-unknown-unknown")
+        set(_use_onemkl_interfaces_cuda ON)
+    endif()
+
+    if (NOT "x${HIP_TARGETS}" STREQUAL "x")
+        set(_use_onemkl_interfaces_hip ON)
+
+        if ("x${_dpnp_sycl_targets}" STREQUAL "x")
+            set(_dpnp_sycl_targets "amd_gpu_${HIP_TARGETS},spir64-unknown-unknown")
+        else()
+            set(_dpnp_sycl_targets "amd_gpu_${HIP_TARGETS},${_dpnp_sycl_targets}")
+        endif()
+    endif()
 else()
-   set(_dpnp_sycl_targets ${DPNP_SYCL_TARGETS})
+    set(_dpnp_sycl_targets ${DPNP_SYCL_TARGETS})
+
+    if ("${DPNP_SYCL_TARGETS}" MATCHES "nvptx64-nvidia-cuda")
+        set(_use_onemkl_interfaces_cuda ON)
+    endif()
+
+    if ("${DPNP_SYCL_TARGETS}" MATCHES "amd_gpu_")
+        set(_use_onemkl_interfaces_hip ON)
+
+        if ("x${HIP_TARGETS}" STREQUAL "x")
+            message(FATAL_ERROR "HIP_TARGETS must be specified when using HIP backend")
+        endif()
+    endif()
+
+    if ("${DPNP_SYCL_TARGETS}" MATCHES "amdgcn-amd-amdhsa")
+        message(FATAL_ERROR
+            "Legacy target 'amdgcn-amd-amdhsa' is not supported. "
+            "Use alias form 'amd_gpu_<arch>' instead"
+        )
+    endif()
 endif()
 
-if(_dpnp_sycl_targets)
+if (_dpnp_sycl_targets)
    message(STATUS "Compiling for -fsycl-targets=${_dpnp_sycl_targets}")
+   list(APPEND _dpnp_sycl_target_compile_options -fsycl-targets=${_dpnp_sycl_targets})
+   list(APPEND _dpnp_sycl_target_link_options -fsycl-targets=${_dpnp_sycl_targets})
 endif()
 
-set(_use_onemkl_interfaces OFF)
 if(DPNP_USE_ONEMKL_INTERFACES)
     set(_use_onemkl_interfaces ON)
 else()
@@ -107,26 +140,33 @@ endif()
 if(_use_onemkl_interfaces)
     set(BUILD_FUNCTIONAL_TESTS False)
     set(BUILD_EXAMPLES False)
+    set(ENABLE_MKLGPU_BACKEND True)
+    set(ENABLE_MKLCPU_BACKEND True)
+
     if(_use_onemkl_interfaces_cuda)
         set(ENABLE_CUBLAS_BACKEND True)
         set(ENABLE_CUSOLVER_BACKEND True)
         set(ENABLE_CUFFT_BACKEND True)
         # set(ENABLE_CURAND_BACKEND True)
-        set(ENABLE_MKLGPU_BACKEND True)
-        set(ENABLE_MKLCPU_BACKEND True)
+    endif()
+    if(_use_onemkl_interfaces_hip)
+        set(ENABLE_ROCBLAS_BACKEND True)
+        set(ENABLE_ROCSOLVER_BACKEND True)
+        set(ENABLE_ROCFFT_BACKEND True)
+        # set(ENABLE_ROCRAND_BACKEND True)
     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)
@@ -136,6 +176,13 @@ if(_use_onemkl_interfaces)
     endif()
     message(STATUS "MKL interfaces lib target used: ${MKL_INTERFACES_LIB}")
     set(CMAKE_INSTALL_RPATH "${CMAKE_BINARY_DIR}/lib")
+else()
+    if(_use_onemkl_interfaces_cuda OR _use_onemkl_interfaces_hip)
+        message(FATAL_ERROR
+            "CUDA or HIP targets are enabled, but oneMKL Interfaces are not. "
+            "Please set DPNP_USE_ONEMKL_INTERFACES=ON to enable them."
+        )
+    endif()
 endif()
 
 if(WIN32)

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:
@@ -130,18 +136,52 @@ Building ``dpnp`` for these targets requires that these CodePlay plugins be inst
 installation layout of compatible version. The following plugins from CodePlay are supported:
 
     - `oneAPI for NVIDIA(R) GPUs <codeplay_nv_plugin_>`_
+    - `oneAPI for AMD GPUs <codeplay_amd_plugin_>`_
 
 .. _codeplay_nv_plugin: https://developer.codeplay.com/products/oneapi/nvidia/
+.. _codeplay_amd_plugin: https://developer.codeplay.com/products/oneapi/amd/
 
 Building ``dpnp`` also requires `building Data Parallel Control Library for custom SYCL targets.
 <https://intelpython.github.io/dpctl/latest/beginners_guides/installation.html#building-for-custom-sycl-targets>`_
 
-Build ``dpnp`` as follows:
+``dpnp`` can be built for CUDA devices as follows:
 
 .. code-block:: bash
 
     python scripts/build_locally.py --target=cuda
 
+And for AMD devices:
+
+.. code-block:: bash
+
+    python scripts/build_locally.py --target-hip=<arch>
+
+Note that the *oneAPI for AMD GPUs* plugin requires the architecture be specified and only
+one architecture can be specified at a time.
+
+To determine the architecture code (``<arch>``) for your AMD GPU, run:
+
+.. code-block:: bash
+
+    rocminfo | grep 'Name: *gfx.*'
+
+This will print names like ``gfx90a``, ``gfx1030``, etc.
+You can then use one of them as the argument to ``--target-hip``.
+
+For example:
+
+.. code-block:: bash
+    python scripts/build_locally.py --target-hip=gfx90a
+
+
+It is, however, possible to build for Intel devices, CUDA devices, and an AMD device
+architecture all at once:
+
+.. code-block:: bash
+
+    python scripts/build_locally.py --target=cuda --target-hip=gfx90a
+
+
 Testing
 =======
 

dpnp/backend/extensions/blas/CMakeLists.txt

@@ -40,12 +40,12 @@ if(_dpnp_sycl_targets)
     target_compile_options(
         ${python_module_name}
         PRIVATE
-        -fsycl-targets=${_dpnp_sycl_targets}
+        ${_dpnp_sycl_target_compile_options}
     )
     target_link_options(
         ${python_module_name}
         PRIVATE
-        -fsycl-targets=${_dpnp_sycl_targets}
+        ${_dpnp_sycl_target_link_options}
     )
 endif()
 

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,