Skip to content

Commit 2c6c697

Browse files
vtavanavtavana
authored
Merge 715d11a into 2700d92 
2 parents 2700d92 + 715d11a commit 2c6c697

File tree

10 files changed

+460
-158
lines changed

10 files changed

+460
-158
lines changed

‎dpnp/backend/extensions/ufunc/CMakeLists.txt

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -24,6 +24,7 @@
2424
# *****************************************************************************
2525

2626
set(_elementwise_sources
27+
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/bitwise_count.cpp
2728
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/common.cpp
2829
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/degrees.cpp
2930
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fabs.cpp

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

Lines changed: 137 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,137 @@
1+
//*****************************************************************************
2+
// Copyright (c) 2025, Intel Corporation
3+
// All rights reserved.
4+
//
5+
// Redistribution and use in source and binary forms, with or without
6+
// modification, are permitted provided that the following conditions are met:
7+
// - Redistributions of source code must retain the above copyright notice,
8+
// this list of conditions and the following disclaimer.
9+
// - Redistributions in binary form must reproduce the above copyright notice,
10+
// this list of conditions and the following disclaimer in the documentation
11+
// and/or other materials provided with the distribution.
12+
//
13+
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
14+
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
15+
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
16+
// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
17+
// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
18+
// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
19+
// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
20+
// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
21+
// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
22+
// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
23+
// THE POSSIBILITY OF SUCH DAMAGE.
24+
//*****************************************************************************
25+
26+
#include <sycl/sycl.hpp>
27+
28+
#include "dpctl4pybind11.hpp"
29+
30+
#include "bitwise_count.hpp"
31+
#include "kernels/elementwise_functions/bitwise_count.hpp"
32+
#include "populate.hpp"
33+
34+
// include a local copy of elementwise common header from dpctl tensor:
35+
// dpctl/tensor/libtensor/source/elementwise_functions/elementwise_functions.hpp
36+
// TODO: replace by including dpctl header once available
37+
#include "../../elementwise_functions/elementwise_functions.hpp"
38+
39+
// dpctl tensor headers
40+
#include "kernels/elementwise_functions/common.hpp"
41+
#include "utils/type_dispatch.hpp"
42+
43+
namespace dpnp::extensions::ufunc
44+
{
45+
namespace py = pybind11;
46+
namespace py_int = dpnp::extensions::py_internal;
47+
48+
namespace impl
49+
{
50+
namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
51+
namespace td_ns = dpctl::tensor::type_dispatch;
52+
53+
/**
54+
* @brief A factory to define pairs of supported types for which
55+
* sycl::bitwise_count<T> function is available.
56+
*
57+
* @tparam T Type of input vector `a` and of result vector `y`.
58+
*/
59+
template <typename T>
60+
struct OutputType
61+
{
62+
using value_type = typename std::disjunction<
63+
td_ns::TypeMapResultEntry<T, std::uint8_t, std::uint8_t>,
64+
td_ns::TypeMapResultEntry<T, std::int8_t, std::uint8_t>,
65+
td_ns::TypeMapResultEntry<T, std::uint16_t, std::uint8_t>,
66+
td_ns::TypeMapResultEntry<T, std::int16_t, std::uint8_t>,
67+
td_ns::TypeMapResultEntry<T, std::uint32_t, std::uint8_t>,
68+
td_ns::TypeMapResultEntry<T, std::int32_t, std::uint8_t>,
69+
td_ns::TypeMapResultEntry<T, std::uint64_t, std::uint8_t>,
70+
td_ns::TypeMapResultEntry<T, std::int64_t, std::uint8_t>,
71+
td_ns::DefaultResultEntry<void>>::result_type;
72+
};
73+
74+
using dpnp::kernels::bitwise_count::BitwiseCountFunctor;
75+
76+
template <typename argT,
77+
typename resT = argT,
78+
unsigned int vec_sz = 4,
79+
unsigned int n_vecs = 2,
80+
bool enable_sg_loadstore = true>
81+
using ContigFunctor =
82+
ew_cmn_ns::UnaryContigFunctor<argT,
83+
resT,
84+
BitwiseCountFunctor<argT, resT>,
85+
vec_sz,
86+
n_vecs,
87+
enable_sg_loadstore>;
88+
89+
template <typename argTy, typename resTy, typename IndexerT>
90+
using StridedFunctor =
91+
ew_cmn_ns::UnaryStridedFunctor<argTy,
92+
resTy,
93+
IndexerT,
94+
BitwiseCountFunctor<argTy, resTy>>;
95+
96+
using ew_cmn_ns::unary_contig_impl_fn_ptr_t;
97+
using ew_cmn_ns::unary_strided_impl_fn_ptr_t;
98+
99+
static unary_contig_impl_fn_ptr_t
100+
bitwise_count_contig_dispatch_vector[td_ns::num_types];
101+
static int bitwise_count_output_typeid_vector[td_ns::num_types];
102+
static unary_strided_impl_fn_ptr_t
103+
bitwise_count_strided_dispatch_vector[td_ns::num_types];
104+
105+
MACRO_POPULATE_DISPATCH_VECTORS(bitwise_count);
106+
} // namespace impl
107+
108+
void init_bitwise_count(py::module_ m)
109+
{
110+
using arrayT = dpctl::tensor::usm_ndarray;
111+
using event_vecT = std::vector<sycl::event>;
112+
{
113+
impl::populate_bitwise_count_dispatch_vectors();
114+
using impl::bitwise_count_contig_dispatch_vector;
115+
using impl::bitwise_count_output_typeid_vector;
116+
using impl::bitwise_count_strided_dispatch_vector;
117+
118+
auto bitwise_count_pyapi = [&](const arrayT &src, const arrayT &dst,
119+
sycl::queue &exec_q,
120+
const event_vecT &depends = {}) {
121+
return py_int::py_unary_ufunc(
122+
src, dst, exec_q, depends, bitwise_count_output_typeid_vector,
123+
bitwise_count_contig_dispatch_vector,
124+
bitwise_count_strided_dispatch_vector);
125+
};
126+
m.def("_bitwise_count", bitwise_count_pyapi, "", py::arg("src"),
127+
py::arg("dst"), py::arg("sycl_queue"),
128+
py::arg("depends") = py::list());
129+
130+
auto bitwise_count_result_type_pyapi = [&](const py::dtype &dtype) {
131+
return py_int::py_unary_ufunc_result_type(
132+
dtype, bitwise_count_output_typeid_vector);
133+
};
134+
m.def("_bitwise_count_result_type", bitwise_count_result_type_pyapi);
135+
}
136+
}
137+
} // namespace dpnp::extensions::ufunc

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

Lines changed: 35 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,35 @@
1+
//*****************************************************************************
2+
// Copyright (c) 2025, Intel Corporation
3+
// All rights reserved.
4+
//
5+
// Redistribution and use in source and binary forms, with or without
6+
// modification, are permitted provided that the following conditions are met:
7+
// - Redistributions of source code must retain the above copyright notice,
8+
// this list of conditions and the following disclaimer.
9+
// - Redistributions in binary form must reproduce the above copyright notice,
10+
// this list of conditions and the following disclaimer in the documentation
11+
// and/or other materials provided with the distribution.
12+
//
13+
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
14+
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
15+
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
16+
// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
17+
// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
18+
// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
19+
// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
20+
// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
21+
// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
22+
// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
23+
// THE POSSIBILITY OF SUCH DAMAGE.
24+
//*****************************************************************************
25+
26+
#pragma once
27+
28+
#include <pybind11/pybind11.h>
29+
30+
namespace py = pybind11;
31+
32+
namespace dpnp::extensions::ufunc
33+
{
34+
void init_bitwise_count(py::module_ m);
35+
} // namespace dpnp::extensions::ufunc

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

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -25,6 +25,7 @@
2525

2626
#include <pybind11/pybind11.h>
2727

28+
#include "bitwise_count.hpp"
2829
#include "degrees.hpp"
2930
#include "fabs.hpp"
3031
#include "fix.hpp"
@@ -52,6 +53,7 @@ namespace dpnp::extensions::ufunc
5253
*/
5354
void init_elementwise_functions(py::module_ m)
5455
{
56+
init_bitwise_count(m);
5557
init_degrees(m);
5658
init_fabs(m);
5759
init_fix(m);

‎dpnp/backend/kernels/elementwise_functions/bitwise_count.hpp

Lines changed: 90 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,90 @@
1+
//*****************************************************************************
2+
// Copyright (c) 2025, Intel Corporation
3+
// All rights reserved.
4+
//
5+
// Redistribution and use in source and binary forms, with or without
6+
// modification, are permitted provided that the following conditions are met:
7+
// - Redistributions of source code must retain the above copyright notice,
8+
// this list of conditions and the following disclaimer.
9+
// - Redistributions in binary form must reproduce the above copyright notice,
10+
// this list of conditions and the following disclaimer in the documentation
11+
// and/or other materials provided with the distribution.
12+
//
13+
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
14+
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
15+
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
16+
// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
17+
// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
18+
// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
19+
// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
20+
// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
21+
// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
22+
// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
23+
// THE POSSIBILITY OF SUCH DAMAGE.
24+
//*****************************************************************************
25+
26+
#pragma once
27+
28+
#include <sycl/sycl.hpp>
29+
30+
// dpctl tensor headers
31+
#include "utils/type_utils.hpp"
32+
33+
namespace dpnp::kernels::bitwise_count
34+
{
35+
namespace tu_ns = dpctl::tensor::type_utils;
36+
37+
template <typename argT, typename resT>
38+
struct BitwiseCountFunctor
39+
{
40+
// is function constant for given argT
41+
using is_constant = typename std::false_type;
42+
// constant value, if constant
43+
// constexpr resT constant_value = resT{};
44+
// is function defined for sycl::vec
45+
using supports_vec = typename std::true_type;
46+
// do both argT and resT support subgroup store/load operation
47+
using supports_sg_loadstore = typename std::true_type;
48+
49+
resT operator()(const argT &x) const
50+
{
51+
if constexpr (std::is_unsigned_v<argT>) {
52+
return sycl::popcount(x);
53+
}
54+
else {
55+
return sycl::popcount(sycl::abs(x));
56+
}
57+
}
58+
59+
template <int vec_sz>
60+
sycl::vec<resT, vec_sz> operator()(const sycl::vec<argT, vec_sz> &x) const
61+
{
62+
if constexpr (std::is_unsigned_v<argT>) {
63+
auto const &res_vec = sycl::popcount(x);
64+
65+
using deducedT = typename std::remove_cv_t<
66+
std::remove_reference_t<decltype(res_vec)>>::element_type;
67+
68+
if constexpr (std::is_same_v<resT, deducedT>) {
69+
return res_vec;
70+
}
71+
else {
72+
return tu_ns::vec_cast<std::uint8_t, deducedT, vec_sz>(res_vec);
73+
}
74+
}
75+
else {
76+
auto const &res_vec = sycl::popcount(sycl::abs(x));
77+
78+
using deducedT = typename std::remove_cv_t<
79+
std::remove_reference_t<decltype(res_vec)>>::element_type;
80+
81+
if constexpr (std::is_same_v<resT, deducedT>) {
82+
return res_vec;
83+
}
84+
else {
85+
return tu_ns::vec_cast<std::uint8_t, deducedT, vec_sz>(res_vec);
86+
}
87+
}
88+
}
89+
};
90+
} // namespace dpnp::kernels::bitwise_count

0 commit comments

Comments
 (0)