implement dpnp.bitwise_count (#2308)

In this PR, `dpnp.bitwise_count` is implemented using [SYCL built-in
function](https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#api:popcount)
`sycl::popcount`.

Author: vtavana

dpnp/backend/extensions/ufunc/CMakeLists.txt

@@ -24,6 +24,7 @@
 # *****************************************************************************
 
 set(_elementwise_sources
+    ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/bitwise_count.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/common.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/degrees.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fabs.cpp

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

@@ -0,0 +1,137 @@
+//*****************************************************************************
+// Copyright (c) 2025, Intel Corporation
+// All rights reserved.
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, 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 "bitwise_count.hpp"
+#include "kernels/elementwise_functions/bitwise_count.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 "utils/type_dispatch.hpp"
+
+namespace dpnp::extensions::ufunc
+{
+namespace py = pybind11;
+namespace py_int = dpnp::extensions::py_internal;
+
+namespace impl
+{
+namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
+namespace td_ns = dpctl::tensor::type_dispatch;
+
+/**
+ * @brief A factory to define pairs of supported types for which
+ * sycl::bitwise_count<T> function is available.
+ *
+ * @tparam T Type of input vector `a` and of result vector `y`.
+ */
+template <typename T>
+struct OutputType
+{
+    using value_type = typename std::disjunction<
+        td_ns::TypeMapResultEntry<T, std::uint8_t, std::uint8_t>,
+        td_ns::TypeMapResultEntry<T, std::int8_t, std::uint8_t>,
+        td_ns::TypeMapResultEntry<T, std::uint16_t, std::uint8_t>,
+        td_ns::TypeMapResultEntry<T, std::int16_t, std::uint8_t>,
+        td_ns::TypeMapResultEntry<T, std::uint32_t, std::uint8_t>,
+        td_ns::TypeMapResultEntry<T, std::int32_t, std::uint8_t>,
+        td_ns::TypeMapResultEntry<T, std::uint64_t, std::uint8_t>,
+        td_ns::TypeMapResultEntry<T, std::int64_t, std::uint8_t>,
+        td_ns::DefaultResultEntry<void>>::result_type;
+};
+
+using dpnp::kernels::bitwise_count::BitwiseCountFunctor;
+
+template <typename argT,
+          typename resT = argT,
+          unsigned int vec_sz = 4,
+          unsigned int n_vecs = 2,
+          bool enable_sg_loadstore = true>
+using ContigFunctor =
+    ew_cmn_ns::UnaryContigFunctor<argT,
+                                  resT,
+                                  BitwiseCountFunctor<argT, resT>,
+                                  vec_sz,
+                                  n_vecs,
+                                  enable_sg_loadstore>;
+
+template <typename argTy, typename resTy, typename IndexerT>
+using StridedFunctor =
+    ew_cmn_ns::UnaryStridedFunctor<argTy,
+                                   resTy,
+                                   IndexerT,
+                                   BitwiseCountFunctor<argTy, resTy>>;
+
+using ew_cmn_ns::unary_contig_impl_fn_ptr_t;
+using ew_cmn_ns::unary_strided_impl_fn_ptr_t;
+
+static unary_contig_impl_fn_ptr_t
+    bitwise_count_contig_dispatch_vector[td_ns::num_types];
+static int bitwise_count_output_typeid_vector[td_ns::num_types];
+static unary_strided_impl_fn_ptr_t
+    bitwise_count_strided_dispatch_vector[td_ns::num_types];
+
+MACRO_POPULATE_DISPATCH_VECTORS(bitwise_count);
+} // namespace impl
+
+void init_bitwise_count(py::module_ m)
+{
+    using arrayT = dpctl::tensor::usm_ndarray;
+    using event_vecT = std::vector<sycl::event>;
+    {
+        impl::populate_bitwise_count_dispatch_vectors();
+        using impl::bitwise_count_contig_dispatch_vector;
+        using impl::bitwise_count_output_typeid_vector;
+        using impl::bitwise_count_strided_dispatch_vector;
+
+        auto bitwise_count_pyapi = [&](const arrayT &src, const arrayT &dst,
+                                       sycl::queue &exec_q,
+                                       const event_vecT &depends = {}) {
+            return py_int::py_unary_ufunc(
+                src, dst, exec_q, depends, bitwise_count_output_typeid_vector,
+                bitwise_count_contig_dispatch_vector,
+                bitwise_count_strided_dispatch_vector);
+        };
+        m.def("_bitwise_count", bitwise_count_pyapi, "", py::arg("src"),
+              py::arg("dst"), py::arg("sycl_queue"),
+              py::arg("depends") = py::list());
+
+        auto bitwise_count_result_type_pyapi = [&](const py::dtype &dtype) {
+            return py_int::py_unary_ufunc_result_type(
+                dtype, bitwise_count_output_typeid_vector);
+        };
+        m.def("_bitwise_count_result_type", bitwise_count_result_type_pyapi);
+    }
+}
+} // namespace dpnp::extensions::ufunc

dpnp/backend/extensions/ufunc/elementwise_functions/bitwise_count.hpp

@@ -0,0 +1,35 @@
+//*****************************************************************************
+// Copyright (c) 2025, Intel Corporation
+// All rights reserved.
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, 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.
+//*****************************************************************************
+
+#pragma once
+
+#include <pybind11/pybind11.h>
+
+namespace py = pybind11;
+
+namespace dpnp::extensions::ufunc
+{
+void init_bitwise_count(py::module_ m);
+} // namespace dpnp::extensions::ufunc

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

@@ -25,6 +25,7 @@
 
 #include <pybind11/pybind11.h>
 
+#include "bitwise_count.hpp"
 #include "degrees.hpp"
 #include "fabs.hpp"
 #include "fix.hpp"
@@ -52,6 +53,7 @@ namespace dpnp::extensions::ufunc
  */
 void init_elementwise_functions(py::module_ m)
 {
+    init_bitwise_count(m);
     init_degrees(m);
     init_fabs(m);
     init_fix(m);

dpnp/backend/kernels/elementwise_functions/bitwise_count.hpp

@@ -0,0 +1,59 @@
+//*****************************************************************************
+// Copyright (c) 2025, Intel Corporation
+// All rights reserved.
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, 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.
+//*****************************************************************************
+
+#pragma once
+
+#include <sycl/sycl.hpp>
+
+// dpctl tensor headers
+#include "utils/type_utils.hpp"
+
+namespace dpnp::kernels::bitwise_count
+{
+namespace tu_ns = dpctl::tensor::type_utils;
+
+template <typename argT, typename resT>
+struct BitwiseCountFunctor
+{
+    // is function constant for given argT
+    using is_constant = typename std::false_type;
+    // constant value, if constant
+    // constexpr resT constant_value = resT{};
+    // is function defined for sycl::vec
+    using supports_vec = typename std::false_type;
+    // do both argT and resT support subgroup store/load operation
+    using supports_sg_loadstore = typename std::true_type;
+
+    resT operator()(const argT &x) const
+    {
+        if constexpr (std::is_unsigned_v<argT>) {
+            return sycl::popcount(x);
+        }
+        else {
+            return sycl::popcount(sycl::abs(x));
+        }
+    }
+};
+} // namespace dpnp::kernels::bitwise_count

dpnp/dpnp_iface_bitwise.py

@@ -38,16 +38,18 @@
 """
 
 # pylint: disable=protected-access
-
+# pylint: disable=no-name-in-module
 
 import dpctl.tensor._tensor_elementwise_impl as ti
 import numpy
 
+import dpnp.backend.extensions.ufunc._ufunc_impl as ufi
 from dpnp.dpnp_algo.dpnp_elementwise_common import DPNPBinaryFunc, DPNPUnaryFunc
 
 __all__ = [
     "binary_repr",
     "bitwise_and",
+    "bitwise_count",
     "bitwise_invert",
     "bitwise_left_shift",
     "bitwise_not",
@@ -215,6 +217,59 @@ def binary_repr(num, width=None):
 )
 
 
+_BITWISE_COUNT_DOCSTRING = """
+Computes the number of 1-bits in the absolute value of `x`.
+
+For full documentation refer to :obj:`numpy.bitwise_count`.
+
+Parameters
+----------
+x : {dpnp.ndarray, usm_ndarray}
+    Input array, expected to have integer or boolean data type.
+out : {None, dpnp.ndarray, usm_ndarray}, optional
+    Output array to populate.
+    Array must have the correct shape and the expected data type.
+
+    Default: ``None``.
+order : {"C", "F", "A", "K"}, optional
+    Memory layout of the newly output array, if parameter `out` is ``None``.
+
+    Default: ``"K"``.
+
+Returns
+-------
+out : dpnp.ndarray
+    The corresponding number of 1-bits in the input. Returns ``uint8`` for all
+    integer types.
+
+Limitations
+-----------
+Parameters `where` and `subok` are supported with their default values.
+Keyword argument `kwargs` is currently unsupported.
+Otherwise ``NotImplementedError`` exception will be raised.
+
+Examples
+--------
+>>> import dpnp as np
+>>> a = np.array(1023)
+>>> np.bitwise_count(a)
+array(10, dtype=uint8)
+
+>>> a = np.array([2**i - 1 for i in range(16)])
+>>> np.bitwise_count(a)
+array([ 0,  1,  2,  3,  4,  5,  6,  7,  8,  9, 10, 11, 12, 13, 14, 15],
+      dtype=uint8)
+
+"""
+
+bitwise_count = DPNPUnaryFunc(
+    "bitwise_count",
+    ufi._bitwise_count_result_type,
+    ufi._bitwise_count,
+    _BITWISE_COUNT_DOCSTRING,
+)
+
+
 _BITWISE_OR_DOCSTRING = """
 Computes the bitwise OR of the underlying binary representation of each
 element `x1_i` of the input array `x1` with the respective element `x2_i`