Skip to content

Commit 30466c5

Browse files
antonwolfyantonwolfy
authored
Merge branch 'master' into assume-mkl-2024.0
2 parents 5492cd9 + 0b7c230 commit 30466c5

30 files changed

+1802
-435
lines changed

‎conda-recipe/bld.bat

Lines changed: 18 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -1,14 +1,7 @@
11
REM A workaround for activate-dpcpp.bat issue to be addressed in 2021.4
2-
set "LIB=%BUILD_PREFIX%\Library\lib;%BUILD_PREFIX%\compiler\lib;%LIB%"
2+
SET "LIB=%BUILD_PREFIX%\Library\lib;%BUILD_PREFIX%\compiler\lib;%LIB%"
33
SET "INCLUDE=%BUILD_PREFIX%\include;%INCLUDE%"
44

5-
REM Since the 60.0.0 release, setuptools includes a local, vendored copy
6-
REM of distutils (from late copies of CPython) that is enabled by default.
7-
REM It breaks build for Windows, so use distutils from "stdlib" as before.
8-
REM @TODO: remove the setting, once transition to build backend on Windows
9-
REM to cmake is complete.
10-
SET "SETUPTOOLS_USE_DISTUTILS=stdlib"
11-
125
"%PYTHON%" setup.py clean --all
136

147
set "MKLROOT=%PREFIX%/Library"
@@ -18,10 +11,15 @@ set "DPL_ROOT_HINT=%PREFIX%/Library"
1811
set "SKBUILD_ARGS=-G Ninja -- -DCMAKE_C_COMPILER:PATH=icx -DCMAKE_CXX_COMPILER:PATH=icx -DCMAKE_VERBOSE_MAKEFILE:BOOL=ON"
1912
set "SKBUILD_ARGS=%SKBUILD_ARGS% -DCMAKE_VERBOSE_MAKEFILE:BOOL=ON"
2013

14+
REM Overriding IPO is useful for building in resources constrained VMs (public CI)
15+
if DEFINED OVERRIDE_INTEL_IPO (
16+
set "SKBUILD_ARGS=%SKBUILD_ARGS% -DCMAKE_INTERPROCEDURAL_OPTIMIZATION:BOOL=FALSE"
17+
)
18+
2119
FOR %%V IN (14.0.0 14 15.0.0 15 16.0.0 16 17.0.0 17) DO @(
2220
REM set DIR_HINT if directory exists
2321
IF EXIST "%BUILD_PREFIX%\Library\lib\clang\%%V\" (
24-
SET "SYCL_INCLUDE_DIR_HINT=%BUILD_PREFIX%\Library\lib\clang\%%V"
22+
SET "SYCL_INCLUDE_DIR_HINT=%BUILD_PREFIX%\Library\lib\clang\%%V"
2523
)
2624
)
2725

@@ -40,19 +38,20 @@ if EXIST "%PLATFORM_DIR%" (
4038
)
4139

4240
if NOT "%WHEELS_OUTPUT_FOLDER%"=="" (
43-
rem Install and assemble wheel package from the build bits
44-
"%PYTHON%" setup.py install bdist_wheel %SKBUILD_ARGS%
45-
if errorlevel 1 exit 1
46-
copy dist\dpnp*.whl %WHEELS_OUTPUT_FOLDER%
47-
if errorlevel 1 exit 1
41+
rem Install and assemble wheel package from the build bits
42+
"%PYTHON%" setup.py install bdist_wheel %SKBUILD_ARGS%
43+
if errorlevel 1 exit 1
44+
copy dist\dpnp*.whl %WHEELS_OUTPUT_FOLDER%
45+
if errorlevel 1 exit 1
4846
) ELSE (
49-
rem Only install
50-
"%PYTHON%" setup.py install %SKBUILD_ARGS%
51-
if errorlevel 1 exit 1
47+
rem Only install
48+
"%PYTHON%" setup.py install %SKBUILD_ARGS%
49+
if errorlevel 1 exit 1
5250
)
5351

5452
rem copy back
5553
if EXIST "%PLATFORM_DIR%" (
56-
copy /Y "%FN%" "%PLATFORM_DIR%\%FN%"
57-
if errorlevel 1 exit 1
54+
rem copy back
55+
copy /Y "%FN%" "%PLATFORM_DIR%\%FN%"
56+
if errorlevel 1 exit 1
5857
)

‎conda-recipe/meta.yaml

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -42,6 +42,7 @@ build:
4242
include_recipe: False
4343
script_env:
4444
- WHEELS_OUTPUT_FOLDER
45+
- OVERRIDE_INTEL_IPO # [win]
4546

4647
test:
4748
requires:

‎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

0 commit comments

Comments
 (0)