Fix review comments

Author: AlexanderKalistratov

dpnp/backend/extensions/statistics/common.hpp

@@ -101,7 +101,8 @@ struct IsNan
 {
     static bool isnan(const T &v)
     {
-        if constexpr (std::is_floating_point_v<T>) {
+        if constexpr (std::is_floating_point_v<T> ||
+                      std::is_same_v<T, sycl::half>) {
             return sycl::isnan(v);
         }
 

dpnp/backend/extensions/statistics/histogram.cpp

@@ -41,8 +41,6 @@
 #include "histogram.hpp"
 #include "histogram_common.hpp"
 
-#include <iostream>
-
 namespace dpctl_td_ns = dpctl::tensor::type_dispatch;
 using dpctl::tensor::usm_ndarray;
 
@@ -139,7 +137,7 @@ struct histogram_kernel
             cgh.depends_on(depends);
             constexpr uint32_t dims = 1;
 
-            auto dispatch_edges = [&](uint32_t local_mem, auto &weights,
+            auto dispatch_edges = [&](uint32_t local_mem, const auto &weights,
                                       auto &hist) {
                 if (device.is_gpu() && (local_mem >= bins_count + 1)) {
                     auto edges = CachedEdges(bins_edges, bins_count + 1, cgh);
@@ -153,7 +151,7 @@ struct histogram_kernel
                 }
             };
 
-            auto dispatch_bins = [&](auto &weights) {
+            auto dispatch_bins = [&](const auto &weights) {
                 const auto local_mem_size =
                     get_local_mem_size_in_items<T>(device);
                 if (local_mem_size >= bins_count) {

dpnp/backend/extensions/statistics/histogram_common.cpp

@@ -102,9 +102,13 @@ void validate(const usm_ndarray &sample,
 
     auto check_overlaping = [&](const array_ptr &first,
                                 const array_ptr &second) {
+        if (first == nullptr || second == nullptr) {
+            return;
+        }
+
         const auto &overlap = dpctl::tensor::overlap::MemoryOverlap();
 
-        if (overlap(sample, histogram)) {
+        if (overlap(*first, *second)) {
             throw py::value_error(get_name(first) +
                                   " has overlapping memory segments with " +
                                   get_name(second));

dpnp/dpnp_iface_histograms.py

@@ -42,6 +42,7 @@
 
 import dpctl.utils as dpu
 import numpy
+from dpctl.tensor._type_utils import _can_cast
 
 import dpnp
 
@@ -304,26 +305,10 @@ def _result_type_for_device(dtype1, dtype2, device):
     return map_dtype_to_device(rt, device)
 
 
-def _supported_by_device(dtype, device):
-    if dtype in [dpnp.float64, dpnp.complex128]:
-        return device.has_aspect_fp64
-
-    if dtype == dpnp.float16:
-        return device.has_aspect_fp16
-
-    return True
-
-
-def _can_cast(dtype1, dtype2, device):
-    if not _supported_by_device(dtype1, device) or not _supported_by_device(
-        dtype2, device
-    ):
-        return False
-
-    return dpnp.can_cast(dtype1, dtype2)
-
-
 def _align_dtypes(a_dtype, bins_dtype, ntype, supported_types, device):
+    has_fp64 = device.has_aspect_fp64
+    has_fp16 = device.has_aspect_fp16
+
     a_bin_dtype = _result_type_for_device(a_dtype, bins_dtype, device)
 
     # histogram implementation doesn't support uint64 as histogram type
@@ -335,11 +320,12 @@ def _align_dtypes(a_dtype, bins_dtype, ntype, supported_types, device):
         return a_bin_dtype, ntype
 
     for sample_type, hist_type in supported_types:
-        if _can_cast(a_bin_dtype, sample_type, device) and _can_cast(
-            ntype, hist_type, device
-        ):
+        if _can_cast(
+            a_bin_dtype, sample_type, has_fp16, has_fp64
+        ) and _can_cast(ntype, hist_type, has_fp16, has_fp64):
             return sample_type, hist_type
 
+    # should not happen
     return None, None
 
 
@@ -468,6 +454,7 @@ def histogram(a, bins=10, range=None, density=None, weights=None):
     # host usm memory
     n_usm_type = "device" if usm_type == "host" else usm_type
 
+    # histogram implementation requires output array to be filled with zeros
     n_casted = dpnp.zeros(
         bin_edges.size - 1,
         dtype=hist_dtype,