Merge branch 'master' into correlate_fft

Author: AlexanderKalistratov

.github/workflows/conda-package.yml

@@ -83,13 +83,13 @@ jobs:
           MAX_BUILD_CMPL_MKL_VERSION: '2025.1a0'
 
       - name: Upload artifact
-        uses: actions/upload-artifact@b4b15b8c7c6ac21ea08fcf65892d2ee8f75cf882 # v4.4.3
+        uses: actions/upload-artifact@6f51ac03b9356f520e9adb1b1b7802705f340c2b # v4.5.0
         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/upload-artifact@b4b15b8c7c6ac21ea08fcf65892d2ee8f75cf882 # v4.4.3
+        uses: actions/upload-artifact@6f51ac03b9356f520e9adb1b1b7802705f340c2b # v4.5.0
         with:
           name: ${{ env.PACKAGE_NAME }} ${{ runner.os }} Wheels Python ${{ matrix.python }}
           path: ${{ env.WHEELS_OUTPUT_FOLDER }}${{ env.PACKAGE_NAME }}-*.whl
@@ -148,7 +148,7 @@ jobs:
 
       - name: Test conda channel
         run: |
-          mamba search ${{ env.PACKAGE_NAME }} -c ${{ env.channel-path }} --override-channels --info --json > ${{ env.ver-json-path }}
+          conda search ${{ env.PACKAGE_NAME }} -c ${{ env.channel-path }} --override-channels --info --json > ${{ env.ver-json-path }}
           cat ${{ env.ver-json-path }}
 
       - name: Get package version
@@ -182,7 +182,7 @@ jobs:
         id: run_tests_linux
         uses: nick-fields/retry@7152eba30c6575329ac0576536151aca5a72780e # v3.0.0
         with:
-          timeout_minutes: 10
+          timeout_minutes: 12
           max_attempts: ${{ env.RUN_TESTS_MAX_ATTEMPTS }}
           retry_on: any
           command: |
@@ -264,7 +264,7 @@ jobs:
       - name: Test conda channel
         run: |
           @echo on
-          mamba search ${{ env.PACKAGE_NAME }} -c ${{ env.channel-path }} --override-channels --info --json > ${{ env.ver-json-path }}
+          conda search ${{ env.PACKAGE_NAME }} -c ${{ env.channel-path }} --override-channels --info --json > ${{ env.ver-json-path }}
 
       - name: Dump version.json
         run: more ${{ env.ver-json-path }}

.github/workflows/openssf-scorecard.yml

@@ -60,14 +60,14 @@ 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@b4b15b8c7c6ac21ea08fcf65892d2ee8f75cf882 # v4.4.3
+        uses: actions/upload-artifact@6f51ac03b9356f520e9adb1b1b7802705f340c2b # v4.5.0
         with:
           name: SARIF file
           path: results.sarif
           retention-days: 14
 
       # Upload the results to GitHub's code scanning dashboard.
       - name: "Upload to code-scanning"
-        uses: github/codeql-action/upload-sarif@aa578102511db1f4524ed59b8cc2bae4f6e88195 # v3.27.6
+        uses: github/codeql-action/upload-sarif@48ab28a6f5dbc2a99bf1e0131198dd8f1df78169 # v3.28.0
         with:
           sarif_file: results.sarif

CMakeLists.txt

@@ -112,8 +112,8 @@ if(_use_onemkl_interfaces)
         set(ENABLE_CUSOLVER_BACKEND True)
         set(ENABLE_CUFFT_BACKEND True)
         # set(ENABLE_CURAND_BACKEND True)
-        set(ENABLE_MKLGPU_BACKEND False)
-        set(ENABLE_MKLCPU_BACKEND False)
+        set(ENABLE_MKLGPU_BACKEND True)
+        set(ENABLE_MKLCPU_BACKEND True)
     endif()
 
     if(DPNP_ONEMKL_INTERFACES_DIR)

doc/conf.py

@@ -218,9 +218,10 @@ def _can_document_member(member, *args, **kwargs):
 
 intersphinx_mapping = {
     "python": ("https://docs.python.org/3/", None),
-    "numpy": ("https://docs.scipy.org/doc/numpy/", None),
+    "numpy": ("https://numpy.org/doc/stable/", None),
     "scipy": ("https://docs.scipy.org/doc/scipy/reference/", None),
     "dpctl": ("https://intelpython.github.io/dpctl/latest/", None),
+    "cupy": ("https://docs.cupy.dev/en/stable/", None),
 }
 
 # If true, `todo` and `todoList` produce output, else they produce nothing.

doc/known_words.txt

@@ -1,11 +1,13 @@
 al
 ary
 backend
+bandlimited
 bitwise
 boolean
 broadcastable
 broadcasted
 byteorder
+cardinalis
 Cholesky
 combinatorially
 conda
@@ -34,11 +36,16 @@ fs
 getter
 Golub
 Hadamard
+histogrammed
 Hypergeometric
+kwargs
+iaxis
 iinfo
 Infs
 intp
+ints
 iterable
+Lanczos
 Lomax
 Mersenne
 meshgrid
@@ -47,13 +54,17 @@ multinomial
 multivalued
 NaN
 NaT
+nd
 ndarray
 ndarrays
 ndim
 Nj
 Nk
 normed
 nuc
+numpy
+nx
+ny
 Nyquist
 oneAPI
 ord
@@ -62,12 +73,15 @@ Penrose
 Polyutils
 pre
 prepend
+prepended
 prepending
 representable
+resampling
 runtimes
 scikit
 signbit
 signum
+sinc
 subarray
 subarrays
 subclasses
@@ -84,7 +98,9 @@ ufunc
 ufuncs
 Unary
 unscaled
+unstacked
 unicode
+Upcasting
 usm
 Vandermonde
 vectorized

dpnp/__init__.py

@@ -25,6 +25,7 @@
 # *****************************************************************************
 
 import os
+import sys
 
 mypath = os.path.dirname(os.path.realpath(__file__))
 
@@ -45,10 +46,22 @@
     if hasattr(os, "add_dll_directory"):
         os.add_dll_directory(mypath)
         os.add_dll_directory(dpctlpath)
+
     os.environ["PATH"] = os.pathsep.join(
         [os.getenv("PATH", ""), mypath, dpctlpath]
     )
 
+    # For virtual environments on Windows, add folder with DPC++ libraries
+    # to the DLL search path
+    if sys.base_exec_prefix != sys.exec_prefix and os.path.isfile(
+        os.path.join(sys.exec_prefix, "pyvenv.cfg")
+    ):
+        dll_path = os.path.join(sys.exec_prefix, "Library", "bin")
+        if os.path.isdir(dll_path):
+            os.environ["PATH"] = os.pathsep.join(
+                [os.getenv("PATH", ""), dll_path]
+            )
+
 # Borrowed from DPCTL
 from dpctl.tensor import DLDeviceType
 

dpnp/backend/extensions/blas/gemm.cpp

@@ -55,7 +55,9 @@ typedef sycl::event (*gemm_impl_fn_ptr_t)(sycl::queue &,
                                           const std::int64_t,
                                           char *,
                                           const std::int64_t,
+#if !defined(USE_ONEMKL_CUBLAS)
                                           const bool,
+#endif // !USE_ONEMKL_CUBLAS
                                           const std::vector<sycl::event> &);
 
 static gemm_impl_fn_ptr_t gemm_dispatch_table[dpctl_td_ns::num_types]
@@ -74,7 +76,9 @@ static sycl::event gemm_impl(sycl::queue &exec_q,
                              const std::int64_t ldb,
                              char *resultC,
                              const std::int64_t ldc,
+#if !defined(USE_ONEMKL_CUBLAS)
                              const bool is_row_major,
+#endif // !USE_ONEMKL_CUBLAS
                              const std::vector<sycl::event> &depends)
 {
     type_utils::validate_type_for_device<Tab>(exec_q);
@@ -236,6 +240,7 @@ std::tuple<sycl::event, sycl::event, bool>
     std::int64_t lda;
     std::int64_t ldb;
 
+// cuBLAS supports only column-major storage
 #if defined(USE_ONEMKL_CUBLAS)
     const bool is_row_major = false;
 
@@ -315,9 +320,15 @@ std::tuple<sycl::event, sycl::event, bool>
     const char *b_typeless_ptr = matrixB.get_data();
     char *r_typeless_ptr = resultC.get_data();
 
+#if defined(USE_ONEMKL_CUBLAS)
+    sycl::event gemm_ev =
+        gemm_fn(exec_q, transA, transB, m, n, k, a_typeless_ptr, lda,
+                b_typeless_ptr, ldb, r_typeless_ptr, ldc, depends);
+#else
     sycl::event gemm_ev = gemm_fn(exec_q, transA, transB, m, n, k,
                                   a_typeless_ptr, lda, b_typeless_ptr, ldb,
                                   r_typeless_ptr, ldc, is_row_major, depends);
+#endif // USE_ONEMKL_CUBLAS
 
     sycl::event args_ev = dpctl::utils::keep_args_alive(
         exec_q, {matrixA, matrixB, resultC}, {gemm_ev});

dpnp/backend/extensions/blas/gemm_batch.cpp

@@ -60,7 +60,9 @@ typedef sycl::event (*gemm_batch_impl_fn_ptr_t)(
     const char *,
     const char *,
     char *,
+#if !defined(USE_ONEMKL_CUBLAS)
     const bool,
+#endif // !USE_ONEMKL_CUBLAS
     const std::vector<sycl::event> &);
 
 static gemm_batch_impl_fn_ptr_t
@@ -83,7 +85,9 @@ static sycl::event gemm_batch_impl(sycl::queue &exec_q,
                                    const char *matrixA,
                                    const char *matrixB,
                                    char *resultC,
+#if !defined(USE_ONEMKL_CUBLAS)
                                    const bool is_row_major,
+#endif // !USE_ONEMKL_CUBLAS
                                    const std::vector<sycl::event> &depends)
 {
     type_utils::validate_type_for_device<Tab>(exec_q);
@@ -311,6 +315,7 @@ std::tuple<sycl::event, sycl::event, bool>
     std::int64_t lda;
     std::int64_t ldb;
 
+// cuBLAS supports only column-major storage
 #if defined(USE_ONEMKL_CUBLAS)
     const bool is_row_major = false;
 
@@ -391,10 +396,17 @@ std::tuple<sycl::event, sycl::event, bool>
     const char *b_typeless_ptr = matrixB.get_data();
     char *r_typeless_ptr = resultC.get_data();
 
+#if defined(USE_ONEMKL_CUBLAS)
+    sycl::event gemm_batch_ev =
+        gemm_batch_fn(exec_q, m, n, k, batch_size, lda, ldb, ldc, stridea,
+                      strideb, stridec, transA, transB, a_typeless_ptr,
+                      b_typeless_ptr, r_typeless_ptr, depends);
+#else
     sycl::event gemm_batch_ev =
         gemm_batch_fn(exec_q, m, n, k, batch_size, lda, ldb, ldc, stridea,
                       strideb, stridec, transA, transB, a_typeless_ptr,
                       b_typeless_ptr, r_typeless_ptr, is_row_major, depends);
+#endif // USE_ONEMKL_CUBLAS
 
     sycl::event args_ev = dpctl::utils::keep_args_alive(
         exec_q, {matrixA, matrixB, resultC}, {gemm_batch_ev});

dpnp/backend/extensions/blas/gemv.cpp

@@ -53,7 +53,9 @@ typedef sycl::event (*gemv_impl_fn_ptr_t)(sycl::queue &,
                                           const std::int64_t,
                                           char *,
                                           const std::int64_t,
+#if !defined(USE_ONEMKL_CUBLAS)
                                           const bool,
+#endif // !USE_ONEMKL_CUBLAS
                                           const std::vector<sycl::event> &);
 
 static gemv_impl_fn_ptr_t gemv_dispatch_vector[dpctl_td_ns::num_types];
@@ -69,7 +71,9 @@ static sycl::event gemv_impl(sycl::queue &exec_q,
                              const std::int64_t incx,
                              char *vectorY,
                              const std::int64_t incy,
+#if !defined(USE_ONEMKL_CUBLAS)
                              const bool is_row_major,
+#endif // !USE_ONEMKL_CUBLAS
                              const std::vector<sycl::event> &depends)
 {
     type_utils::validate_type_for_device<T>(exec_q);
@@ -190,6 +194,7 @@ std::pair<sycl::event, sycl::event>
     oneapi::mkl::transpose transA;
     std::size_t src_nelems;
 
+// cuBLAS supports only column-major storage
 #if defined(USE_ONEMKL_CUBLAS)
     const bool is_row_major = false;
     std::int64_t m;
@@ -299,9 +304,15 @@ std::pair<sycl::event, sycl::event>
         y_typeless_ptr -= (y_shape[0] - 1) * std::abs(incy) * y_elemsize;
     }
 
+#if defined(USE_ONEMKL_CUBLAS)
+    sycl::event gemv_ev =
+        gemv_fn(exec_q, transA, m, n, a_typeless_ptr, lda, x_typeless_ptr, incx,
+                y_typeless_ptr, incy, depends);
+#else
     sycl::event gemv_ev =
         gemv_fn(exec_q, transA, m, n, a_typeless_ptr, lda, x_typeless_ptr, incx,
                 y_typeless_ptr, incy, is_row_major, depends);
+#endif // USE_ONEMKL_CUBLAS
 
     sycl::event args_ev = dpctl::utils::keep_args_alive(
         exec_q, {matrixA, vectorX, vectorY}, {gemv_ev});

dpnp/backend/extensions/statistics/CMakeLists.txt

@@ -29,6 +29,7 @@ set(_module_src
     ${CMAKE_CURRENT_SOURCE_DIR}/bincount.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/common.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/histogram.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/histogramdd.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/histogram_common.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/sliding_dot_product1d.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/sliding_window1d.cpp