Merge branch 'master' into remove_dot_and_multiply_from_backend

Author: AlexanderKalistratov

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

dpnp/backend/extensions/statistics/bincount.hpp

@@ -25,10 +25,11 @@
 
 #pragma once
 
-#include <dpctl4pybind11.hpp>
+#include <pybind11/pybind11.h>
 #include <sycl/sycl.hpp>
 
 #include "dispatch_table.hpp"
+#include "dpctl4pybind11.hpp"
 
 namespace dpctl_td_ns = dpctl::tensor::type_dispatch;
 

dpnp/backend/extensions/statistics/common.hpp

@@ -28,13 +28,12 @@
 #include <complex>
 #include <pybind11/numpy.h>
 #include <pybind11/pybind11.h>
-
-// clang-format off
-// math_utils.hpp doesn't include sycl header but uses sycl types
-// so sycl.hpp must be included before math_utils.hpp
 #include <sycl/sycl.hpp>
+
 #include "utils/math_utils.hpp"
-// clang-format on
+#include "utils/type_utils.hpp"
+
+namespace type_utils = dpctl::tensor::type_utils;
 
 namespace statistics::common
 {
@@ -54,24 +53,20 @@ constexpr auto Align(N n, D d)
 template <typename T, sycl::memory_order Order, sycl::memory_scope Scope>
 struct AtomicOp
 {
-    static void add(T &lhs, const T value)
+    static void add(T &lhs, const T &value)
     {
-        sycl::atomic_ref<T, Order, Scope> lh(lhs);
-        lh += value;
-    }
-};
+        if constexpr (type_utils::is_complex_v<T>) {
+            using vT = typename T::value_type;
+            vT *_lhs = reinterpret_cast<vT(&)[2]>(lhs);
+            const vT *_val = reinterpret_cast<const vT(&)[2]>(value);
 
-template <typename T, sycl::memory_order Order, sycl::memory_scope Scope>
-struct AtomicOp<std::complex<T>, Order, Scope>
-{
-    static void add(std::complex<T> &lhs, const std::complex<T> value)
-    {
-        T *_lhs = reinterpret_cast<T(&)[2]>(lhs);
-        const T *_val = reinterpret_cast<const T(&)[2]>(value);
-        sycl::atomic_ref<T, Order, Scope> lh0(_lhs[0]);
-        lh0 += _val[0];
-        sycl::atomic_ref<T, Order, Scope> lh1(_lhs[1]);
-        lh1 += _val[1];
+            AtomicOp<vT, Order, Scope>::add(_lhs[0], _val[0]);
+            AtomicOp<vT, Order, Scope>::add(_lhs[1], _val[1]);
+        }
+        else {
+            sycl::atomic_ref<T, Order, Scope> lh(lhs);
+            lh += value;
+        }
     }
 };
 
@@ -80,17 +75,12 @@ struct Less
 {
     bool operator()(const T &lhs, const T &rhs) const
     {
-        return std::less{}(lhs, rhs);
-    }
-};
-
-template <typename T>
-struct Less<std::complex<T>>
-{
-    bool operator()(const std::complex<T> &lhs,
-                    const std::complex<T> &rhs) const
-    {
-        return dpctl::tensor::math_utils::less_complex(lhs, rhs);
+        if constexpr (type_utils::is_complex_v<T>) {
+            return dpctl::tensor::math_utils::less_complex(lhs, rhs);
+        }
+        else {
+            return std::less{}(lhs, rhs);
+        }
     }
 };
 
@@ -99,26 +89,23 @@ struct IsNan
 {
     static bool isnan(const T &v)
     {
-        if constexpr (std::is_floating_point_v<T> ||
-                      std::is_same_v<T, sycl::half>) {
+        if constexpr (type_utils::is_complex_v<T>) {
+            using vT = typename T::value_type;
+
+            const vT real1 = std::real(v);
+            const vT imag1 = std::imag(v);
+
+            return IsNan<vT>::isnan(real1) || IsNan<vT>::isnan(imag1);
+        }
+        else if constexpr (std::is_floating_point_v<T> ||
+                           std::is_same_v<T, sycl::half>) {
             return sycl::isnan(v);
         }
 
         return false;
     }
 };
 
-template <typename T>
-struct IsNan<std::complex<T>>
-{
-    static bool isnan(const std::complex<T> &v)
-    {
-        T real1 = std::real(v);
-        T imag1 = std::imag(v);
-        return sycl::isnan(real1) || sycl::isnan(imag1);
-    }
-};
-
 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/statistics/histogram.hpp

@@ -25,11 +25,11 @@
 
 #pragma once
 
+#include <pybind11/pybind11.h>
 #include <sycl/sycl.hpp>
 
 #include "dispatch_table.hpp"
-
-// namespace dpctl_td_ns = dpctl::tensor::type_dispatch;
+#include "dpctl4pybind11.hpp"
 
 namespace statistics::histogram
 {

dpnp/backend/extensions/statistics/histogram_common.cpp

@@ -88,16 +88,14 @@ void validate(const usm_ndarray &sample,
                   {&histogram}, names);
 
     check_size_at_least(bins_ptr, 2, names);
-
     check_size_at_least(&histogram, 1, names);
-    check_num_dims(&histogram, 1, names);
 
     if (weights_ptr) {
         check_num_dims(weights_ptr, 1, names);
 
-        auto sample_size = sample.get_size();
+        auto sample_size = sample.get_shape(0);
         auto weights_size = weights_ptr->get_size();
-        if (sample.get_size() != weights_ptr->get_size()) {
+        if (sample_size != weights_ptr->get_size()) {
             throw py::value_error(name_of(&sample, names) + " size (" +
                                   std::to_string(sample_size) + ") and " +
                                   name_of(weights_ptr, names) + " size (" +
@@ -110,61 +108,74 @@ void validate(const usm_ndarray &sample,
 
     if (sample.get_ndim() == 1) {
         check_num_dims(bins_ptr, 1, names);
+
+        if (bins_ptr && histogram.get_size() != bins_ptr->get_size() - 1) {
+            auto hist_size = histogram.get_size();
+            auto bins_size = bins_ptr->get_size();
+            throw py::value_error(
+                name_of(&histogram, names) + " parameter and " +
+                name_of(bins_ptr, names) + " parameters shape mismatch. " +
+                name_of(&histogram, names) + " size is " +
+                std::to_string(hist_size) + name_of(bins_ptr, names) +
+                " must have size " + std::to_string(hist_size + 1) +
+                " but have " + std::to_string(bins_size));
+        }
     }
     else if (sample.get_ndim() == 2) {
         auto sample_count = sample.get_shape(0);
         auto expected_dims = sample.get_shape(1);
 
-        if (bins_ptr != nullptr && bins_ptr->get_ndim() != expected_dims) {
+        if (histogram.get_ndim() != expected_dims) {
             throw py::value_error(
-                name_of(&sample, names) + " parameter has shape {" +
-                std::to_string(sample_count) + "x" +
-                std::to_string(expected_dims) + "}" + ", so " +
-                name_of(bins_ptr, names) + " parameter expected to be " +
+                name_of(&sample, names) + " parameter has shape (" +
+                std::to_string(sample_count) + ", " +
+                std::to_string(expected_dims) + ")" + ", so " +
+                name_of(&histogram, names) + " parameter expected to be " +
                 std::to_string(expected_dims) +
                 "d. "
                 "Actual " +
-                std::to_string(bins->get_ndim()) + "d");
+                std::to_string(histogram.get_ndim()) + "d");
         }
-    }
 
-    if (bins_ptr != nullptr) {
-        py::ssize_t expected_hist_size = 1;
-        for (int i = 0; i < bins_ptr->get_ndim(); ++i) {
-            expected_hist_size *= (bins_ptr->get_shape(i) - 1);
+        if (bins_ptr != nullptr) {
+            py::ssize_t expected_bins_size = 0;
+            for (int i = 0; i < histogram.get_ndim(); ++i) {
+                expected_bins_size += histogram.get_shape(i) + 1;
+            }
+
+            auto actual_bins_size = bins_ptr->get_size();
+            if (actual_bins_size != expected_bins_size) {
+                throw py::value_error(
+                    name_of(&histogram, names) + " and " +
+                    name_of(bins_ptr, names) + " shape mismatch. " +
+                    name_of(bins_ptr, names) + " expected to have size = " +
+                    std::to_string(expected_bins_size) + ". Actual " +
+                    std::to_string(actual_bins_size));
+            }
         }
 
-        if (histogram.get_size() != expected_hist_size) {
-            throw py::value_error(
-                name_of(&histogram, names) + " and " +
-                name_of(bins_ptr, names) + " shape mismatch. " +
-                name_of(&histogram, names) + " expected to have size = " +
-                std::to_string(expected_hist_size) + ". Actual " +
-                std::to_string(histogram.get_size()));
+        int64_t max_hist_size = std::numeric_limits<uint32_t>::max() - 1;
+        if (histogram.get_size() > max_hist_size) {
+            throw py::value_error(name_of(&histogram, names) +
+                                  " parameter size expected to be less than " +
+                                  std::to_string(max_hist_size) + ". Actual " +
+                                  std::to_string(histogram.get_size()));
         }
-    }
-
-    int64_t max_hist_size = std::numeric_limits<uint32_t>::max() - 1;
-    if (histogram.get_size() > max_hist_size) {
-        throw py::value_error(name_of(&histogram, names) +
-                              " parameter size expected to be less than " +
-                              std::to_string(max_hist_size) + ". Actual " +
-                              std::to_string(histogram.get_size()));
-    }
 
-    auto array_types = dpctl_td_ns::usm_ndarray_types();
-    auto hist_type = static_cast<typenum_t>(
-        array_types.typenum_to_lookup_id(histogram.get_typenum()));
-    if (histogram.get_elemsize() == 8 && hist_type != typenum_t::CFLOAT) {
-        auto device = exec_q.get_device();
-        bool _64bit_atomics = device.has(sycl::aspect::atomic64);
-
-        if (!_64bit_atomics) {
-            auto device_name = device.get_info<sycl::info::device::name>();
-            throw py::value_error(
-                name_of(&histogram, names) +
-                " parameter has 64-bit type, but 64-bit atomics " +
-                " are not supported for " + device_name);
+        auto array_types = dpctl_td_ns::usm_ndarray_types();
+        auto hist_type = static_cast<typenum_t>(
+            array_types.typenum_to_lookup_id(histogram.get_typenum()));
+        if (histogram.get_elemsize() == 8 && hist_type != typenum_t::CFLOAT) {
+            auto device = exec_q.get_device();
+            bool _64bit_atomics = device.has(sycl::aspect::atomic64);
+
+            if (!_64bit_atomics) {
+                auto device_name = device.get_info<sycl::info::device::name>();
+                throw py::value_error(
+                    name_of(&histogram, names) +
+                    " parameter has 64-bit type, but 64-bit atomics " +
+                    " are not supported for " + device_name);
+            }
         }
     }
 }