Skip to content

Commit c4677e9

Browse files
antonwolfyantonwolfy
authored
Add implementation of dpnp.gcd and dpnp.lcm functions (#2091)
* Add gcd and lcm implementations to ufunc * Add gcd and lcm functions to mathematical interface * Add CFD tests * Add third party tests * Add more tests * Update description of input arrays
1 parent 48515c8 commit c4677e9

File tree

15 files changed

+763
-1
lines changed

15 files changed

+763
-1
lines changed

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

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -32,7 +32,9 @@ set(_elementwise_sources
3232
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fmax.cpp
3333
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fmin.cpp
3434
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fmod.cpp
35+
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/gcd.cpp
3536
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/heaviside.cpp
37+
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/lcm.cpp
3638
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/logaddexp2.cpp
3739
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/radians.cpp
3840
)

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

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -32,7 +32,9 @@
3232
#include "fmax.hpp"
3333
#include "fmin.hpp"
3434
#include "fmod.hpp"
35+
#include "gcd.hpp"
3536
#include "heaviside.hpp"
37+
#include "lcm.hpp"
3638
#include "logaddexp2.hpp"
3739
#include "radians.hpp"
3840

@@ -52,7 +54,9 @@ void init_elementwise_functions(py::module_ m)
5254
init_fmax(m);
5355
init_fmin(m);
5456
init_fmod(m);
57+
init_gcd(m);
5558
init_heaviside(m);
59+
init_lcm(m);
5660
init_logaddexp2(m);
5761
init_radians(m);
5862
}

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

Lines changed: 186 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,186 @@
1+
//*****************************************************************************
2+
// Copyright (c) 2024, Intel Corporation
3+
// All rights reserved.
4+
//
5+
// Redistribution and use in source and binary forms, with or without
6+
// maxification, 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 "gcd.hpp"
31+
#include "kernels/elementwise_functions/gcd.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+
namespace td_ns = dpctl::tensor::type_dispatch;
48+
49+
namespace impl
50+
{
51+
namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
52+
53+
template <typename T1, typename T2>
54+
struct OutputType
55+
{
56+
using value_type = typename std::disjunction< // disjunction is C++17
57+
// feature, supported by DPC++
58+
td_ns::BinaryTypeMapResultEntry<T1,
59+
std::uint8_t,
60+
T2,
61+
std::uint8_t,
62+
std::uint8_t>,
63+
td_ns::BinaryTypeMapResultEntry<T1,
64+
std::int8_t,
65+
T2,
66+
std::int8_t,
67+
std::int8_t>,
68+
td_ns::BinaryTypeMapResultEntry<T1,
69+
std::uint16_t,
70+
T2,
71+
std::uint16_t,
72+
std::uint16_t>,
73+
td_ns::BinaryTypeMapResultEntry<T1,
74+
std::int16_t,
75+
T2,
76+
std::int16_t,
77+
std::int16_t>,
78+
td_ns::BinaryTypeMapResultEntry<T1,
79+
std::uint32_t,
80+
T2,
81+
std::uint32_t,
82+
std::uint32_t>,
83+
td_ns::BinaryTypeMapResultEntry<T1,
84+
std::int32_t,
85+
T2,
86+
std::int32_t,
87+
std::int32_t>,
88+
td_ns::BinaryTypeMapResultEntry<T1,
89+
std::uint64_t,
90+
T2,
91+
std::uint64_t,
92+
std::uint64_t>,
93+
td_ns::BinaryTypeMapResultEntry<T1,
94+
std::int64_t,
95+
T2,
96+
std::int64_t,
97+
std::int64_t>,
98+
td_ns::BinaryTypeMapResultEntry<T1,
99+
std::uint64_t,
100+
T2,
101+
std::int64_t,
102+
std::uint64_t>,
103+
td_ns::BinaryTypeMapResultEntry<T1,
104+
std::int64_t,
105+
T2,
106+
std::uint64_t,
107+
std::int64_t>,
108+
td_ns::DefaultResultEntry<void>>::result_type;
109+
};
110+
111+
using dpnp::kernels::gcd::GcdFunctor;
112+
113+
template <typename argT1,
114+
typename argT2,
115+
typename resT,
116+
unsigned int vec_sz = 4,
117+
unsigned int n_vecs = 2,
118+
bool enable_sg_loadstore = true>
119+
using ContigFunctor =
120+
ew_cmn_ns::BinaryContigFunctor<argT1,
121+
argT2,
122+
resT,
123+
GcdFunctor<argT1, argT2, resT>,
124+
vec_sz,
125+
n_vecs,
126+
enable_sg_loadstore>;
127+
128+
template <typename argT1, typename argT2, typename resT, typename IndexerT>
129+
using StridedFunctor =
130+
ew_cmn_ns::BinaryStridedFunctor<argT1,
131+
argT2,
132+
resT,
133+
IndexerT,
134+
GcdFunctor<argT1, argT2, resT>>;
135+
136+
using ew_cmn_ns::binary_contig_impl_fn_ptr_t;
137+
using ew_cmn_ns::binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t;
138+
using ew_cmn_ns::binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t;
139+
using ew_cmn_ns::binary_strided_impl_fn_ptr_t;
140+
141+
static binary_contig_impl_fn_ptr_t gcd_contig_dispatch_table[td_ns::num_types]
142+
[td_ns::num_types];
143+
static int gcd_output_typeid_table[td_ns::num_types][td_ns::num_types];
144+
static binary_strided_impl_fn_ptr_t
145+
gcd_strided_dispatch_table[td_ns::num_types][td_ns::num_types];
146+
147+
MACRO_POPULATE_DISPATCH_TABLES(gcd);
148+
} // namespace impl
149+
150+
void init_gcd(py::module_ m)
151+
{
152+
using arrayT = dpctl::tensor::usm_ndarray;
153+
using event_vecT = std::vector<sycl::event>;
154+
{
155+
impl::populate_gcd_dispatch_tables();
156+
using impl::gcd_contig_dispatch_table;
157+
using impl::gcd_output_typeid_table;
158+
using impl::gcd_strided_dispatch_table;
159+
160+
auto gcd_pyapi = [&](const arrayT &src1, const arrayT &src2,
161+
const arrayT &dst, sycl::queue &exec_q,
162+
const event_vecT &depends = {}) {
163+
return py_int::py_binary_ufunc(
164+
src1, src2, dst, exec_q, depends, gcd_output_typeid_table,
165+
gcd_contig_dispatch_table, gcd_strided_dispatch_table,
166+
// no dedicated kernel for C-contig row with broadcasting
167+
td_ns::NullPtrTable<
168+
impl::
169+
binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t>{},
170+
td_ns::NullPtrTable<
171+
impl::
172+
binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t>{});
173+
};
174+
m.def("_gcd", gcd_pyapi, "", py::arg("src1"), py::arg("src2"),
175+
py::arg("dst"), py::arg("sycl_queue"),
176+
py::arg("depends") = py::list());
177+
178+
auto gcd_result_type_pyapi = [&](const py::dtype &dtype1,
179+
const py::dtype &dtype2) {
180+
return py_int::py_binary_ufunc_result_type(dtype1, dtype2,
181+
gcd_output_typeid_table);
182+
};
183+
m.def("_gcd_result_type", gcd_result_type_pyapi);
184+
}
185+
}
186+
} // namespace dpnp::extensions::ufunc

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

Lines changed: 35 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,35 @@
1+
//*****************************************************************************
2+
// Copyright (c) 2024, 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_gcd(py::module_ m);
35+
} // namespace dpnp::extensions::ufunc

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

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -38,7 +38,6 @@
3838

3939
// dpctl tensor headers
4040
#include "kernels/elementwise_functions/common.hpp"
41-
#include "kernels/elementwise_functions/maximum.hpp"
4241
#include "utils/type_dispatch.hpp"
4342

4443
namespace dpnp::extensions::ufunc

0 commit comments

Comments
 (0)