Add support of uint64 dtype in dpnp.bincount

CHANGELOG.md

@@ -14,6 +14,8 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0
 
 ### Changed
 
+* Allowed input array of `uint64` dtype in `dpnp.bincount` [#2361](https://github.com/IntelPython/dpnp/pull/2361)
+
 ### Fixed
 
 

dpnp/backend/extensions/statistics/bincount.cpp

@@ -72,7 +72,7 @@ struct BincountEdges
     template <typename dT>
     bool in_bounds(const dT *val, const boundsT &bounds) const
     {
-        return check_in_bounds(val[0], std::get<0>(bounds),
+        return check_in_bounds(static_cast<T>(val[0]), std::get<0>(bounds),
                                std::get<1>(bounds));
     }
 
@@ -81,16 +81,17 @@ struct BincountEdges
     T max;
 };
 
-template <typename T, typename HistType = size_t>
+using DefaultHistType = int64_t;
+
+template <typename T, typename HistType = DefaultHistType>
 struct BincountF
 {
     static sycl::event impl(sycl::queue &exec_q,
                             const void *vin,
-                            const int64_t min,
-                            const int64_t max,
+                            const uint64_t min,
+                            const uint64_t max,
                             const void *vweights,
                             void *vout,
-                            const size_t,
                             const size_t size,
                             const std::vector<sycl::event> &depends)
     {
@@ -145,9 +146,12 @@ struct BincountF
     }
 };
 
-using SupportedTypes = std::tuple<std::tuple<int64_t, int64_t>,
+using SupportedTypes = std::tuple<std::tuple<int64_t, DefaultHistType>,
+                                  std::tuple<uint64_t, DefaultHistType>,
                                   std::tuple<int64_t, float>,
-                                  std::tuple<int64_t, double>>;
+                                  std::tuple<uint64_t, float>,
+                                  std::tuple<int64_t, double>,
+                                  std::tuple<uint64_t, double>>;
 
 } // namespace
 
@@ -158,8 +162,8 @@ Bincount::Bincount() : dispatch_table("sample", "histogram")
 
 std::tuple<sycl::event, sycl::event> Bincount::call(
     const dpctl::tensor::usm_ndarray &sample,
-    const int64_t min,
-    const int64_t max,
+    const uint64_t min,
+    const uint64_t max,
     const std::optional<const dpctl::tensor::usm_ndarray> &weights,
     dpctl::tensor::usm_ndarray &histogram,
     const std::vector<sycl::event> &depends)
@@ -182,8 +186,7 @@ std::tuple<sycl::event, sycl::event> Bincount::call(
         weights.has_value() ? weights.value().get_data() : nullptr;
 
     auto ev = bincount_func(exec_q, sample.get_data(), min, max, weights_ptr,
-                            histogram.get_data(), histogram.get_shape(0),
-                            sample.get_shape(0), depends);
+                            histogram.get_data(), sample.get_size(), depends);
 
     sycl::event args_ev;
     if (weights.has_value()) {

dpnp/backend/extensions/statistics/bincount.hpp

@@ -39,12 +39,11 @@ struct Bincount
 {
     using FnT = sycl::event (*)(sycl::queue &,
                                 const void *,
-                                const int64_t,
-                                const int64_t,
+                                const uint64_t,
+                                const uint64_t,
                                 const void *,
                                 void *,
                                 const size_t,
-                                const size_t,
                                 const std::vector<sycl::event> &);
 
     common::DispatchTable2<FnT> dispatch_table;
@@ -53,8 +52,8 @@ struct Bincount
 
     std::tuple<sycl::event, sycl::event>
         call(const dpctl::tensor::usm_ndarray &input,
-             const int64_t min,
-             const int64_t max,
+             const uint64_t min,
+             const uint64_t max,
              const std::optional<const dpctl::tensor::usm_ndarray> &weights,
              dpctl::tensor::usm_ndarray &output,
              const std::vector<sycl::event> &depends);

dpnp/backend/extensions/statistics/histogram.cpp

@@ -106,7 +106,9 @@ using CachedEdges = HistogramEdges<T, CachedData<const T, 1>>;
 template <typename T>
 using UncachedEdges = HistogramEdges<T, UncachedData<const T, 1>>;
 
-template <typename T, typename BinsT, typename HistType = size_t>
+using DefaultHistType = int64_t;
+
+template <typename T, typename BinsT, typename HistType = DefaultHistType>
 struct HistogramF
 {
     static sycl::event impl(sycl::queue &exec_q,
@@ -186,26 +188,27 @@ using HistogramF_ = HistogramF<SampleType, SampleType, HistType>;
 
 } // namespace
 
-using SupportedTypes = std::tuple<std::tuple<uint64_t, int64_t>,
-                                  std::tuple<int64_t, int64_t>,
-                                  std::tuple<uint64_t, float>,
-                                  std::tuple<int64_t, float>,
-                                  std::tuple<uint64_t, double>,
-                                  std::tuple<int64_t, double>,
-                                  std::tuple<uint64_t, std::complex<float>>,
-                                  std::tuple<int64_t, std::complex<float>>,
-                                  std::tuple<uint64_t, std::complex<double>>,
-                                  std::tuple<int64_t, std::complex<double>>,
-                                  std::tuple<float, int64_t>,
-                                  std::tuple<double, int64_t>,
-                                  std::tuple<float, float>,
-                                  std::tuple<double, double>,
-                                  std::tuple<float, std::complex<float>>,
-                                  std::tuple<double, std::complex<double>>,
-                                  std::tuple<std::complex<float>, int64_t>,
-                                  std::tuple<std::complex<double>, int64_t>,
-                                  std::tuple<std::complex<float>, float>,
-                                  std::tuple<std::complex<double>, double>>;
+using SupportedTypes =
+    std::tuple<std::tuple<uint64_t, DefaultHistType>,
+               std::tuple<int64_t, DefaultHistType>,
+               std::tuple<uint64_t, float>,
+               std::tuple<int64_t, float>,
+               std::tuple<uint64_t, double>,
+               std::tuple<int64_t, double>,
+               std::tuple<uint64_t, std::complex<float>>,
+               std::tuple<int64_t, std::complex<float>>,
+               std::tuple<uint64_t, std::complex<double>>,
+               std::tuple<int64_t, std::complex<double>>,
+               std::tuple<float, DefaultHistType>,
+               std::tuple<double, DefaultHistType>,
+               std::tuple<float, float>,
+               std::tuple<double, double>,
+               std::tuple<float, std::complex<float>>,
+               std::tuple<double, std::complex<double>>,
+               std::tuple<std::complex<float>, DefaultHistType>,
+               std::tuple<std::complex<double>, DefaultHistType>,
+               std::tuple<std::complex<float>, float>,
+               std::tuple<std::complex<double>, double>>;
 
 Histogram::Histogram() : dispatch_table("sample", "histogram")
 {

dpnp/dpnp_iface_histograms.py

@@ -293,8 +293,8 @@ def _bincount_run_native(
 
     mem_ev, bc_ev = statistics_ext.bincount(
         x_usm,
-        min_v,
-        max_v,
+        min_v.item(),
+        max_v.item(),
         weights_usm,
         n_usm,
         depends=_manager.submitted_events,
@@ -313,6 +313,11 @@ def bincount(x, weights=None, minlength=0):
 
     For full documentation refer to :obj:`numpy.bincount`.
 
+    Warning
+    -------
+    This function synchronizes in order to calculate binning edges.
+    This may harm performance in some applications.
+
     Parameters
     ----------
     x : {dpnp.ndarray, usm_ndarray}
@@ -391,10 +396,8 @@ def bincount(x, weights=None, minlength=0):
 
     if x_casted_dtype is None or ntype_casted is None:  # pragma: no cover
         raise ValueError(
-            f"function '{bincount}' does not support input types "
-            f"({x.dtype}, {ntype}), "
-            "and the inputs could not be coerced to any "
-            "supported types"
+            f"Input types ({x.dtype}, {ntype}) are not supported, "
+            "and the inputs could not be coerced to any supported types"
         )
 
     x_casted = dpnp.asarray(x, dtype=x_casted_dtype, order="C")
@@ -508,6 +511,11 @@ def histogram(a, bins=10, range=None, density=None, weights=None):
 
     For full documentation refer to :obj:`numpy.histogram`.
 
+    Warning
+    -------
+    This function may synchronize in order to check a monotonically increasing
+    array of bin edges. This may harm performance in some applications.
+
     Parameters
     ----------
     a : {dpnp.ndarray, usm_ndarray}
@@ -611,9 +619,8 @@ def histogram(a, bins=10, range=None, density=None, weights=None):
 
     if a_bin_dtype is None or hist_dtype is None:  # pragma: no cover
         raise ValueError(
-            f"function '{histogram}' does not support input types "
-            f"({a.dtype}, {bin_edges.dtype}, {ntype}), "
-            "and the inputs could not be coerced to any "
+            f"Input types ({a.dtype}, {bin_edges.dtype}, {ntype}) "
+            "are not supported, and the inputs could not be coerced to any "
             "supported types"
         )
 
@@ -675,6 +682,11 @@ def histogram_bin_edges(a, bins=10, range=None, weights=None):
 
     For full documentation refer to :obj:`numpy.histogram_bin_edges`.
 
+    Warning
+    -------
+    This function may synchronize in order to check a monotonically increasing
+    array of bin edges. This may harm performance in some applications.
+
     Parameters
     ----------
     a : {dpnp.ndarray, usm_ndarray}
@@ -760,6 +772,13 @@ def histogram2d(x, y, bins=10, range=None, density=None, weights=None):
     """
     Compute the bi-dimensional histogram of two data samples.
 
+    For full documentation refer to :obj:`numpy.histogram2d`.
+
+    Warning
+    -------
+    This function may synchronize in order to check a monotonically increasing
+    array of bin edges. This may harm performance in some applications.
+
     Parameters
     ----------
     x : {dpnp.ndarray, usm_ndarray} of shape (N,)
@@ -1088,6 +1107,11 @@ def histogramdd(sample, bins=10, range=None, density=None, weights=None):
 
     For full documentation refer to :obj:`numpy.histogramdd`.
 
+    Warning
+    -------
+    This function may synchronize in order to check a monotonically increasing
+    array of bin edges. This may harm performance in some applications.
+
     Parameters
     ----------
     sample : {dpnp.ndarray, usm_ndarray}

dpnp/tests/test_histogram.py

@@ -12,8 +12,11 @@
 
 from .helper import (
     assert_dtype_allclose,
+    generate_random_numpy_array,
     get_abs_array,
     get_all_dtypes,
+    get_complex_dtypes,
+    get_float_complex_dtypes,
     get_float_dtypes,
     get_integer_dtypes,
     has_support_aspect64,
@@ -532,34 +535,29 @@ def test_range(self, range, dtype):
 
 
 class TestBincount:
-    @pytest.mark.parametrize("dtype", get_integer_dtypes())
-    def test_rand_data(self, dtype):
-        n = 100
-        upper_bound = 10 if dtype != dpnp.bool_ else 1
-        v = numpy.random.randint(0, upper_bound, size=n, dtype=dtype)
+    @pytest.mark.parametrize("dt", get_integer_dtypes() + [numpy.bool_])
+    def test_rand_data(self, dt):
+        v = generate_random_numpy_array(100, dtype=dt, low=0)
         iv = dpnp.array(v)
 
-        if numpy.issubdtype(dtype, numpy.uint64):
-            # discussed in numpy issue 17760
-            assert_raises(TypeError, numpy.bincount, v)
-            assert_raises(ValueError, dpnp.bincount, iv)
-        else:
-            expected_hist = numpy.bincount(v)
-            result_hist = dpnp.bincount(iv)
-            assert_array_equal(result_hist, expected_hist)
+        if numpy.issubdtype(dt, numpy.uint64) and numpy_version() < "2.2.4":
+            v = v.astype(numpy.int64)
 
-    @pytest.mark.parametrize("dtype", get_integer_dtypes())
-    def test_arange_data(self, dtype):
-        v = numpy.arange(100).astype(dtype)
+        expected_hist = numpy.bincount(v)
+        result_hist = dpnp.bincount(iv)
+        assert_array_equal(result_hist, expected_hist)
+
+    @pytest.mark.parametrize("dt", get_integer_dtypes())
+    def test_arange_data(self, dt):
+        v = numpy.arange(100, dtype=dt)
         iv = dpnp.array(v)
 
-        if numpy.issubdtype(dtype, numpy.uint64):
-            assert_raises(TypeError, numpy.bincount, v)
-            assert_raises(ValueError, dpnp.bincount, iv)
-        else:
-            expected_hist = numpy.bincount(v)
-            result_hist = dpnp.bincount(iv)
-            assert_array_equal(result_hist, expected_hist)
+        if numpy.issubdtype(dt, numpy.uint64) and numpy_version() < "2.2.4":
+            v = v.astype(numpy.int64)
+
+        expected_hist = numpy.bincount(v)
+        result_hist = dpnp.bincount(iv)
+        assert_array_equal(result_hist, expected_hist)
 
     @pytest.mark.parametrize("xp", [numpy, dpnp])
     def test_negative_values(self, xp):
@@ -581,11 +579,17 @@ def test_weights_another_sycl_queue(self):
             dpnp.bincount(v, weights=w)
 
     @pytest.mark.parametrize("xp", [numpy, dpnp])
-    def test_weights_unsupported_dtype(self, xp):
-        v = dpnp.arange(5)
-        w = dpnp.arange(5, dtype=dpnp.complex64)
-        with assert_raises(ValueError):
-            dpnp.bincount(v, weights=w)
+    @pytest.mark.parametrize("dt", get_float_complex_dtypes())
+    def test_data_unsupported_dtype(self, xp, dt):
+        v = xp.arange(5, dtype=dt)
+        assert_raises(TypeError, xp.bincount, v)
+
+    @pytest.mark.parametrize("xp", [numpy, dpnp])
+    @pytest.mark.parametrize("dt", get_complex_dtypes())
+    def test_weights_unsupported_dtype(self, xp, dt):
+        v = xp.arange(5)
+        w = xp.arange(5, dtype=dt)
+        assert_raises((TypeError, ValueError), xp.bincount, v, weights=w)
 
     @pytest.mark.parametrize(
         "bins_count",
@@ -606,11 +610,11 @@ def test_different_bins_amount(self, bins_count):
     )
     @pytest.mark.parametrize("minlength", [0, 1, 3, 5])
     def test_minlength(self, array, minlength):
-        np_a = numpy.array(array)
-        dpnp_a = dpnp.array(array)
+        a = numpy.array(array)
+        ia = dpnp.array(a)
 
-        expected = numpy.bincount(np_a, minlength=minlength)
-        result = dpnp.bincount(dpnp_a, minlength=minlength)
+        expected = numpy.bincount(a, minlength=minlength)
+        result = dpnp.bincount(ia, minlength=minlength)
         assert_allclose(result, expected)
 
     @pytest.mark.filterwarnings("ignore::DeprecationWarning")
@@ -639,21 +643,23 @@ def test_minlength_none(self, xp):
         )
 
     @pytest.mark.parametrize(
-        "array", [[1, 2, 2, 1, 2, 4]], ids=["[1, 2, 2, 1, 2, 4]"]
+        "weights",
+        [None, [0.3, 0.5, 0, 0.7, 1.0, -0.6], [2, 2, 2, 2, 2, 2]],
+        ids=["None", "float_data", "int_data"],
     )
     @pytest.mark.parametrize(
-        "weights",
-        [None, [0.3, 0.5, 0.2, 0.7, 1.0, -0.6], [2, 2, 2, 2, 2, 2]],
-        ids=["None", "[0.3, 0.5, 0.2, 0.7, 1., -0.6]", "[2, 2, 2, 2, 2, 2]"],
+        "dt", get_all_dtypes(no_none=True, no_complex=True)
     )
-    def test_weights(self, array, weights):
-        np_a = numpy.array(array)
-        np_weights = numpy.array(weights) if weights is not None else weights
-        dpnp_a = dpnp.array(array)
-        dpnp_weights = dpnp.array(weights) if weights is not None else weights
-
-        expected = numpy.bincount(np_a, weights=np_weights)
-        result = dpnp.bincount(dpnp_a, weights=dpnp_weights)
+    def test_weights(self, weights, dt):
+        a = numpy.array([1, 2, 2, 1, 2, 4])
+        ia = dpnp.array(a)
+        w = iw = None
+        if weights is not None:
+            w = numpy.array(weights, dtype=dt)
+            iw = dpnp.array(w)
+
+        expected = numpy.bincount(a, weights=w)
+        result = dpnp.bincount(ia, weights=iw)
         assert_allclose(result, expected)
 
     @pytest.mark.parametrize(