Skip to content

Commit 57ff795

Browse files
vtavanavtavana
authored
Merge branch 'master' into impl_fft
2 parents 456674b + 0b7c230 commit 57ff795

27 files changed

+1777
-416
lines changed

‎doc/reference/logic.rst

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -26,6 +26,8 @@ Infinities and NaNs
2626
dpnp.isfinite
2727
dpnp.isinf
2828
dpnp.isnan
29+
dpnp.isneginf
30+
dpnp.isposinf
2931

3032

3133
Array type testing

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

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -26,6 +26,7 @@
2626
set(_elementwise_sources
2727
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/common.cpp
2828
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fabs.cpp
29+
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fmod.cpp
2930
)
3031

3132
set(python_module_name _ufunc_impl)

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

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -26,6 +26,7 @@
2626
#include <pybind11/pybind11.h>
2727

2828
#include "fabs.hpp"
29+
#include "fmod.hpp"
2930

3031
namespace py = pybind11;
3132

@@ -37,5 +38,6 @@ namespace dpnp::extensions::ufunc
3738
void init_elementwise_functions(py::module_ m)
3839
{
3940
init_fabs(m);
41+
init_fmod(m);
4042
}
4143
} // namespace dpnp::extensions::ufunc

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

Lines changed: 193 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,193 @@
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+
#include <sycl/sycl.hpp>
27+
28+
#include "dpctl4pybind11.hpp"
29+
30+
#include "fmod.hpp"
31+
#include "kernels/elementwise_functions/fmod.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 py = pybind11;
44+
45+
namespace dpnp::extensions::ufunc
46+
{
47+
namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
48+
namespace py_int = dpnp::extensions::py_internal;
49+
namespace td_ns = dpctl::tensor::type_dispatch;
50+
51+
using ew_cmn_ns::unary_contig_impl_fn_ptr_t;
52+
using ew_cmn_ns::unary_strided_impl_fn_ptr_t;
53+
54+
namespace impl
55+
{
56+
/**
57+
* @brief A factory to define pairs of supported types for which
58+
* sycl::fmod<T> function is available.
59+
*
60+
* @tparam T1 Type of input vectors `a`
61+
* @tparam T2 Type of input vectors `b`
62+
*/
63+
template <typename T1, typename T2>
64+
struct OutputType
65+
{
66+
using value_type = typename std::disjunction<
67+
td_ns::BinaryTypeMapResultEntry<T1, bool, T2, bool, std::int8_t>,
68+
td_ns::BinaryTypeMapResultEntry<T1,
69+
std::uint8_t,
70+
T2,
71+
std::uint8_t,
72+
std::uint8_t>,
73+
td_ns::BinaryTypeMapResultEntry<T1,
74+
std::int8_t,
75+
T2,
76+
std::int8_t,
77+
std::int8_t>,
78+
td_ns::BinaryTypeMapResultEntry<T1,
79+
std::uint16_t,
80+
T2,
81+
std::uint16_t,
82+
std::uint16_t>,
83+
td_ns::BinaryTypeMapResultEntry<T1,
84+
std::int16_t,
85+
T2,
86+
std::int16_t,
87+
std::int16_t>,
88+
td_ns::BinaryTypeMapResultEntry<T1,
89+
std::uint32_t,
90+
T2,
91+
std::uint32_t,
92+
std::uint32_t>,
93+
td_ns::BinaryTypeMapResultEntry<T1,
94+
std::int32_t,
95+
T2,
96+
std::int32_t,
97+
std::int32_t>,
98+
td_ns::BinaryTypeMapResultEntry<T1,
99+
std::uint64_t,
100+
T2,
101+
std::uint64_t,
102+
std::uint64_t>,
103+
td_ns::BinaryTypeMapResultEntry<T1,
104+
std::int64_t,
105+
T2,
106+
std::int64_t,
107+
std::int64_t>,
108+
td_ns::BinaryTypeMapResultEntry<T1,
109+
sycl::half,
110+
T2,
111+
sycl::half,
112+
sycl::half>,
113+
td_ns::BinaryTypeMapResultEntry<T1, float, T2, float, float>,
114+
td_ns::BinaryTypeMapResultEntry<T1, double, T2, double, double>,
115+
td_ns::DefaultResultEntry<void>>::result_type;
116+
};
117+
118+
using dpnp::kernels::fmod::FmodFunctor;
119+
120+
template <typename argT1,
121+
typename argT2,
122+
typename resT,
123+
unsigned int vec_sz = 4,
124+
unsigned int n_vecs = 2,
125+
bool enable_sg_loadstore = true>
126+
using ContigFunctor =
127+
ew_cmn_ns::BinaryContigFunctor<argT1,
128+
argT2,
129+
resT,
130+
FmodFunctor<argT1, argT2, resT>,
131+
vec_sz,
132+
n_vecs,
133+
enable_sg_loadstore>;
134+
135+
template <typename argT1, typename argT2, typename resT, typename IndexerT>
136+
using StridedFunctor =
137+
ew_cmn_ns::BinaryStridedFunctor<argT1,
138+
argT2,
139+
resT,
140+
IndexerT,
141+
FmodFunctor<argT1, argT2, resT>>;
142+
143+
using ew_cmn_ns::binary_contig_impl_fn_ptr_t;
144+
using ew_cmn_ns::binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t;
145+
using ew_cmn_ns::binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t;
146+
using ew_cmn_ns::binary_strided_impl_fn_ptr_t;
147+
148+
static binary_contig_impl_fn_ptr_t fmod_contig_dispatch_table[td_ns::num_types]
149+
[td_ns::num_types];
150+
static int fmod_output_typeid_table[td_ns::num_types][td_ns::num_types];
151+
static binary_strided_impl_fn_ptr_t
152+
fmod_strided_dispatch_table[td_ns::num_types][td_ns::num_types];
153+
154+
MACRO_POPULATE_DISPATCH_TABLES(fmod);
155+
} // namespace impl
156+
157+
void init_fmod(py::module_ m)
158+
{
159+
using arrayT = dpctl::tensor::usm_ndarray;
160+
using event_vecT = std::vector<sycl::event>;
161+
{
162+
impl::populate_fmod_dispatch_tables();
163+
using impl::fmod_contig_dispatch_table;
164+
using impl::fmod_output_typeid_table;
165+
using impl::fmod_strided_dispatch_table;
166+
167+
auto fmod_pyapi = [&](const arrayT &src1, const arrayT &src2,
168+
const arrayT &dst, sycl::queue &exec_q,
169+
const event_vecT &depends = {}) {
170+
return py_int::py_binary_ufunc(
171+
src1, src2, dst, exec_q, depends, fmod_output_typeid_table,
172+
fmod_contig_dispatch_table, fmod_strided_dispatch_table,
173+
// no support of C-contig row with broadcasting in OneMKL
174+
td_ns::NullPtrTable<
175+
impl::
176+
binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t>{},
177+
td_ns::NullPtrTable<
178+
impl::
179+
binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t>{});
180+
};
181+
m.def("_fmod", fmod_pyapi, "", py::arg("src1"), py::arg("src2"),
182+
py::arg("dst"), py::arg("sycl_queue"),
183+
py::arg("depends") = py::list());
184+
185+
auto fmod_result_type_pyapi = [&](const py::dtype &dtype1,
186+
const py::dtype &dtype2) {
187+
return py_int::py_binary_ufunc_result_type(
188+
dtype1, dtype2, fmod_output_typeid_table);
189+
};
190+
m.def("_fmod_result_type", fmod_result_type_pyapi);
191+
}
192+
}
193+
} // namespace dpnp::extensions::ufunc

‎dpnp/backend/extensions/ufunc/elementwise_functions/fmod.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_fmod(py::module_ m);
35+
} // namespace dpnp::extensions::ufunc

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

Lines changed: 109 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -26,7 +26,8 @@
2626
#pragma once
2727

2828
/**
29-
* @brief A macro used to define factories and a populating universal functions.
29+
* @brief A macro used to define factories and a populating unary universal
30+
* functions.
3031
*/
3132
#define MACRO_POPULATE_DISPATCH_VECTORS(__name__) \
3233
template <typename T1, typename T2, unsigned int vec_sz, \
@@ -120,3 +121,110 @@
120121
dvb3; \
121122
dvb3.populate_dispatch_vector(__name__##_output_typeid_vector); \
122123
};
124+
125+
/**
126+
* @brief A macro used to define factories and a populating binary universal
127+
* functions.
128+
*/
129+
#define MACRO_POPULATE_DISPATCH_TABLES(__name__) \
130+
template <typename argT1, typename argT2, typename resT, \
131+
unsigned int vec_sz, unsigned int n_vecs> \
132+
class __name__##_contig_kernel; \
133+
\
134+
template <typename argTy1, typename argTy2> \
135+
sycl::event __name__##_contig_impl( \
136+
sycl::queue &exec_q, size_t nelems, const char *arg1_p, \
137+
py::ssize_t arg1_offset, const char *arg2_p, py::ssize_t arg2_offset, \
138+
char *res_p, py::ssize_t res_offset, \
139+
const std::vector<sycl::event> &depends = {}) \
140+
{ \
141+
return ew_cmn_ns::binary_contig_impl<argTy1, argTy2, OutputType, \
142+
ContigFunctor, \
143+
__name__##_contig_kernel>( \
144+
exec_q, nelems, arg1_p, arg1_offset, arg2_p, arg2_offset, res_p, \
145+
res_offset, depends); \
146+
} \
147+
\
148+
template <typename fnT, typename T1, typename T2> \
149+
struct ContigFactory \
150+
{ \
151+
fnT get() \
152+
{ \
153+
if constexpr (std::is_same_v< \
154+
typename OutputType<T1, T2>::value_type, void>) \
155+
{ \
156+
\
157+
fnT fn = nullptr; \
158+
return fn; \
159+
} \
160+
else { \
161+
fnT fn = __name__##_contig_impl<T1, T2>; \
162+
return fn; \
163+
} \
164+
} \
165+
}; \
166+
\
167+
template <typename fnT, typename T1, typename T2> \
168+
struct TypeMapFactory \
169+
{ \
170+
std::enable_if_t<std::is_same<fnT, int>::value, int> get() \
171+
{ \
172+
using rT = typename OutputType<T1, T2>::value_type; \
173+
return td_ns::GetTypeid<rT>{}.get(); \
174+
} \
175+
}; \
176+
\
177+
template <typename T1, typename T2, typename resT, typename IndexerT> \
178+
class __name__##_strided_kernel; \
179+
\
180+
template <typename argTy1, typename argTy2> \
181+
sycl::event __name__##_strided_impl( \
182+
sycl::queue &exec_q, size_t nelems, int nd, \
183+
const py::ssize_t *shape_and_strides, const char *arg1_p, \
184+
py::ssize_t arg1_offset, const char *arg2_p, py::ssize_t arg2_offset, \
185+
char *res_p, py::ssize_t res_offset, \
186+
const std::vector<sycl::event> &depends, \
187+
const std::vector<sycl::event> &additional_depends) \
188+
{ \
189+
return ew_cmn_ns::binary_strided_impl<argTy1, argTy2, OutputType, \
190+
StridedFunctor, \
191+
__name__##_strided_kernel>( \
192+
exec_q, nelems, nd, shape_and_strides, arg1_p, arg1_offset, \
193+
arg2_p, arg2_offset, res_p, res_offset, depends, \
194+
additional_depends); \
195+
} \
196+
\
197+
template <typename fnT, typename T1, typename T2> \
198+
struct StridedFactory \
199+
{ \
200+
fnT get() \
201+
{ \
202+
if constexpr (std::is_same_v< \
203+
typename OutputType<T1, T2>::value_type, void>) \
204+
{ \
205+
fnT fn = nullptr; \
206+
return fn; \
207+
} \
208+
else { \
209+
fnT fn = __name__##_strided_impl<T1, T2>; \
210+
return fn; \
211+
} \
212+
} \
213+
}; \
214+
\
215+
void populate_##__name__##_dispatch_tables(void) \
216+
{ \
217+
td_ns::DispatchTableBuilder<binary_contig_impl_fn_ptr_t, \
218+
ContigFactory, td_ns::num_types> \
219+
dvb1; \
220+
dvb1.populate_dispatch_table(__name__##_contig_dispatch_table); \
221+
\
222+
td_ns::DispatchTableBuilder<binary_strided_impl_fn_ptr_t, \
223+
StridedFactory, td_ns::num_types> \
224+
dvb2; \
225+
dvb2.populate_dispatch_table(__name__##_strided_dispatch_table); \
226+
\
227+
td_ns::DispatchTableBuilder<int, TypeMapFactory, td_ns::num_types> \
228+
dvb3; \
229+
dvb3.populate_dispatch_table(__name__##_output_typeid_table); \
230+
};

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

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -43,6 +43,7 @@ set(_elementwise_sources
4343
${CMAKE_CURRENT_SOURCE_DIR}/exp2.cpp
4444
${CMAKE_CURRENT_SOURCE_DIR}/expm1.cpp
4545
${CMAKE_CURRENT_SOURCE_DIR}/floor.cpp
46+
${CMAKE_CURRENT_SOURCE_DIR}/fmod.cpp
4647
${CMAKE_CURRENT_SOURCE_DIR}/hypot.cpp
4748
${CMAKE_CURRENT_SOURCE_DIR}/ln.cpp
4849
${CMAKE_CURRENT_SOURCE_DIR}/log10.cpp

0 commit comments

Comments
 (0)