Merge branch 'master' into impl_pad

Author: vtavana

.github/workflows/openssf-scorecard.yml

@@ -68,6 +68,6 @@ jobs:
 
       # Upload the results to GitHub's code scanning dashboard.
       - name: "Upload to code-scanning"
-        uses: github/codeql-action/upload-sarif@c36620d31ac7c881962c3d9dd939c40ec9434f2b # v3.26.12
+        uses: github/codeql-action/upload-sarif@f779452ac5af1c261dce0346a8f964149f49322b # v3.26.13
         with:
           sarif_file: results.sarif

doc/reference/linalg.rst

@@ -69,7 +69,6 @@ Norms and other numbers
    dpnp.trace
    dpnp.linalg.trace (Array API compatible)
 
-
 Solving linear equations
 --------------------------
 

doc/reference/ndarray.rst

@@ -92,6 +92,7 @@ Other attributes
    :nosignatures:
 
    dpnp.ndarray.T
+   dpnp.ndarray.mT
    dpnp.ndarray.real
    dpnp.ndarray.imag
    dpnp.ndarray.flat

dpnp/backend/CMakeLists.txt

@@ -30,8 +30,6 @@ set(DPNP_SRC
     kernels/dpnp_krnl_indexing.cpp
     kernels/dpnp_krnl_mathematical.cpp
     kernels/dpnp_krnl_random.cpp
-    kernels/dpnp_krnl_reduction.cpp
-    kernels/dpnp_krnl_searching.cpp
     kernels/dpnp_krnl_sorting.cpp
     kernels/dpnp_krnl_statistics.cpp
     src/constants.cpp

dpnp/backend/extensions/ufunc/CMakeLists.txt

@@ -35,6 +35,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/lcm.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/ldexp.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/logaddexp2.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/radians.cpp
 )

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

@@ -35,6 +35,7 @@
 #include "gcd.hpp"
 #include "heaviside.hpp"
 #include "lcm.hpp"
+#include "ldexp.hpp"
 #include "logaddexp2.hpp"
 #include "radians.hpp"
 
@@ -57,6 +58,7 @@ void init_elementwise_functions(py::module_ m)
     init_gcd(m);
     init_heaviside(m);
     init_lcm(m);
+    init_ldexp(m);
     init_logaddexp2(m);
     init_radians(m);
 }

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

@@ -0,0 +1,169 @@
+//*****************************************************************************
+// Copyright (c) 2024, Intel Corporation
+// All rights reserved.
+//
+// Redistribution and use in source and binary forms, with or without
+// maxification, 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.
+//
+// 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.
+//*****************************************************************************
+
+#include <sycl/sycl.hpp>
+
+#include "dpctl4pybind11.hpp"
+
+#include "kernels/elementwise_functions/ldexp.hpp"
+#include "ldexp.hpp"
+#include "populate.hpp"
+
+// include a local copy of elementwise common header from dpctl tensor:
+// dpctl/tensor/libtensor/source/elementwise_functions/elementwise_functions.hpp
+// TODO: replace by including dpctl header once available
+#include "../../elementwise_functions/elementwise_functions.hpp"
+
+// dpctl tensor headers
+#include "kernels/elementwise_functions/common.hpp"
+#include "kernels/elementwise_functions/maximum.hpp"
+#include "utils/type_dispatch.hpp"
+
+namespace dpnp::extensions::ufunc
+{
+namespace py = pybind11;
+namespace py_int = dpnp::extensions::py_internal;
+namespace td_ns = dpctl::tensor::type_dispatch;
+
+namespace impl
+{
+namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
+namespace max_ns = dpctl::tensor::kernels::maximum;
+
+// Supports the same types table as for maximum function in dpctl
+// template <typename T1, typename T2>
+// using OutputType = max_ns::MaximumOutputType<T1, T2>;
+template <typename T1, typename T2>
+struct OutputType
+{
+    using value_type = typename std::disjunction< // disjunction is C++17
+                                                  // feature, supported by DPC++
+        td_ns::BinaryTypeMapResultEntry<T1,
+                                        sycl::half,
+                                        T2,
+                                        std::int8_t,
+                                        sycl::half>,
+        td_ns::BinaryTypeMapResultEntry<T1,
+                                        sycl::half,
+                                        T2,
+                                        std::int16_t,
+                                        sycl::half>,
+        td_ns::BinaryTypeMapResultEntry<T1,
+                                        sycl::half,
+                                        T2,
+                                        std::int32_t,
+                                        sycl::half>,
+        td_ns::BinaryTypeMapResultEntry<T1,
+                                        sycl::half,
+                                        T2,
+                                        std::int64_t,
+                                        sycl::half>,
+        td_ns::BinaryTypeMapResultEntry<T1, float, T2, std::int8_t, float>,
+        td_ns::BinaryTypeMapResultEntry<T1, float, T2, std::int16_t, float>,
+        td_ns::BinaryTypeMapResultEntry<T1, float, T2, std::int32_t, float>,
+        td_ns::BinaryTypeMapResultEntry<T1, float, T2, std::int64_t, float>,
+        td_ns::BinaryTypeMapResultEntry<T1, double, T2, std::int8_t, double>,
+        td_ns::BinaryTypeMapResultEntry<T1, double, T2, std::int16_t, double>,
+        td_ns::BinaryTypeMapResultEntry<T1, double, T2, std::int32_t, double>,
+        td_ns::BinaryTypeMapResultEntry<T1, double, T2, std::int64_t, double>,
+        td_ns::DefaultResultEntry<void>>::result_type;
+};
+
+using dpnp::kernels::ldexp::LdexpFunctor;
+
+template <typename argT1,
+          typename argT2,
+          typename resT,
+          unsigned int vec_sz = 4,
+          unsigned int n_vecs = 2,
+          bool enable_sg_loadstore = true>
+using ContigFunctor =
+    ew_cmn_ns::BinaryContigFunctor<argT1,
+                                   argT2,
+                                   resT,
+                                   LdexpFunctor<argT1, argT2, resT>,
+                                   vec_sz,
+                                   n_vecs,
+                                   enable_sg_loadstore>;
+
+template <typename argT1, typename argT2, typename resT, typename IndexerT>
+using StridedFunctor =
+    ew_cmn_ns::BinaryStridedFunctor<argT1,
+                                    argT2,
+                                    resT,
+                                    IndexerT,
+                                    LdexpFunctor<argT1, argT2, resT>>;
+
+using ew_cmn_ns::binary_contig_impl_fn_ptr_t;
+using ew_cmn_ns::binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t;
+using ew_cmn_ns::binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t;
+using ew_cmn_ns::binary_strided_impl_fn_ptr_t;
+
+static binary_contig_impl_fn_ptr_t
+    ldexp_contig_dispatch_table[td_ns::num_types][td_ns::num_types];
+static int ldexp_output_typeid_table[td_ns::num_types][td_ns::num_types];
+static binary_strided_impl_fn_ptr_t
+    ldexp_strided_dispatch_table[td_ns::num_types][td_ns::num_types];
+
+MACRO_POPULATE_DISPATCH_TABLES(ldexp);
+} // namespace impl
+
+void init_ldexp(py::module_ m)
+{
+    using arrayT = dpctl::tensor::usm_ndarray;
+    using event_vecT = std::vector<sycl::event>;
+    {
+        impl::populate_ldexp_dispatch_tables();
+        using impl::ldexp_contig_dispatch_table;
+        using impl::ldexp_output_typeid_table;
+        using impl::ldexp_strided_dispatch_table;
+
+        auto ldexp_pyapi = [&](const arrayT &src1, const arrayT &src2,
+                               const arrayT &dst, sycl::queue &exec_q,
+                               const event_vecT &depends = {}) {
+            return py_int::py_binary_ufunc(
+                src1, src2, dst, exec_q, depends, ldexp_output_typeid_table,
+                ldexp_contig_dispatch_table, ldexp_strided_dispatch_table,
+                // no support of C-contig row with broadcasting in OneMKL
+                td_ns::NullPtrTable<
+                    impl::
+                        binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t>{},
+                td_ns::NullPtrTable<
+                    impl::
+                        binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t>{});
+        };
+        m.def("_ldexp", ldexp_pyapi, "", py::arg("src1"), py::arg("src2"),
+              py::arg("dst"), py::arg("sycl_queue"),
+              py::arg("depends") = py::list());
+
+        auto ldexp_result_type_pyapi = [&](const py::dtype &dtype1,
+                                           const py::dtype &dtype2) {
+            return py_int::py_binary_ufunc_result_type(
+                dtype1, dtype2, ldexp_output_typeid_table);
+        };
+        m.def("_ldexp_result_type", ldexp_result_type_pyapi);
+    }
+}
+} // namespace dpnp::extensions::ufunc

dpnp/backend/examples/example8.cpp renamed to dpnp/backend/extensions/ufunc/elementwise_functions/ldexp.hpp

@@ -1,5 +1,5 @@
 //*****************************************************************************
-// Copyright (c) 2016-2024, Intel Corporation
+// Copyright (c) 2024, Intel Corporation
 // All rights reserved.
 //
 // Redistribution and use in source and binary forms, with or without
@@ -23,45 +23,13 @@
 // THE POSSIBILITY OF SUCH DAMAGE.
 //*****************************************************************************
 
-/**
- * Example 9.
- *
- * TODO explanation of the example
- *
- * Possible compile line:
- * . /opt/intel/oneapi/setvars.sh
- * g++ -g dpnp/backend/examples/example8.cpp -Idpnp -Idpnp/backend/include
- * -Ldpnp -Wl,-rpath='$ORIGIN'/dpnp -ldpnp_backend_c -o example8
- *
- */
-#include <iostream>
+#pragma once
 
-#include "dpnp_iface.hpp"
+#include <pybind11/pybind11.h>
 
-int main(int, char **)
-{
-    const size_t size = 16;
-
-    double *array = (double *)dpnp_memory_alloc_c(size * sizeof(double));
-    long *result = (long *)dpnp_memory_alloc_c(size * sizeof(long));
-
-    std::cout << "array" << std::endl;
-    for (size_t i = 0; i < size; ++i) {
-        array[i] = (double)(size - i) / 2;
-        std::cout << array[i] << ", ";
-    }
-    std::cout << std::endl;
-
-    dpnp_argsort_c<double, long>(array, result, size);
+namespace py = pybind11;
 
-    std::cout << "array with 'sorted' indices" << std::endl;
-    for (size_t i = 0; i < size; ++i) {
-        std::cout << result[i] << ", ";
-    }
-    std::cout << std::endl;
-
-    dpnp_memory_free_c(result);
-    dpnp_memory_free_c(array);
-
-    return 0;
-}
+namespace dpnp::extensions::ufunc
+{
+void init_ldexp(py::module_ m);
+} // namespace dpnp::extensions::ufunc