Skip to content

Commit 99a243b

Browse files
antonwolfyantonwolfy
authored
Merge branch 'master' into update_backend
2 parents 77d387d + af601c6 commit 99a243b

39 files changed

+959
-511
lines changed

‎.github/workflows/build-sphinx.yml

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -91,7 +91,7 @@ jobs:
9191
sudo apt-get install -y nvidia-cuda-toolkit clinfo
9292
9393
- name: Checkout repo
94-
uses: actions/checkout@a5ac7e51b41094c92402da3b24376905380afc29 # v4.1.6
94+
uses: actions/checkout@692973e3d937129bcbf40652eb9f2f61becf3332 # v4.1.7
9595
with:
9696
fetch-depth: 0
9797

@@ -221,7 +221,7 @@ jobs:
221221
runs-on: ubuntu-20.04
222222

223223
steps:
224-
- uses: actions/checkout@a5ac7e51b41094c92402da3b24376905380afc29 # v4.1.6
224+
- uses: actions/checkout@692973e3d937129bcbf40652eb9f2f61becf3332 # v4.1.7
225225
with:
226226
fetch-depth: 0
227227

‎.github/workflows/conda-package.yml

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -90,7 +90,7 @@ jobs:
9090
access_token: ${{ github.token }}
9191

9292
- name: Checkout DPNP repo
93-
uses: actions/checkout@a5ac7e51b41094c92402da3b24376905380afc29 # v4.1.6
93+
uses: actions/checkout@692973e3d937129bcbf40652eb9f2f61becf3332 # v4.1.7
9494
with:
9595
fetch-depth: 0
9696

@@ -515,7 +515,7 @@ jobs:
515515
run: mamba install anaconda-client
516516

517517
- name: Checkout repo
518-
uses: actions/checkout@a5ac7e51b41094c92402da3b24376905380afc29 # v4.1.6
518+
uses: actions/checkout@692973e3d937129bcbf40652eb9f2f61becf3332 # v4.1.7
519519
with:
520520
repository: IntelPython/devops-tools
521521
fetch-depth: 0

‎.github/workflows/generate_coverage.yaml

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -32,7 +32,7 @@ jobs:
3232
access_token: ${{ github.token }}
3333

3434
- name: Checkout repo
35-
uses: actions/checkout@a5ac7e51b41094c92402da3b24376905380afc29 # v4.1.6
35+
uses: actions/checkout@692973e3d937129bcbf40652eb9f2f61becf3332 # v4.1.7
3636
with:
3737
fetch-depth: 0
3838

‎.github/workflows/openssf-scorecard.yml

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -33,7 +33,7 @@ jobs:
3333

3434
steps:
3535
- name: "Checkout code"
36-
uses: actions/checkout@a5ac7e51b41094c92402da3b24376905380afc29 # v4.1.6
36+
uses: actions/checkout@692973e3d937129bcbf40652eb9f2f61becf3332 # v4.1.7
3737
with:
3838
persist-credentials: false
3939

@@ -68,6 +68,6 @@ jobs:
6868

6969
# Upload the results to GitHub's code scanning dashboard.
7070
- name: "Upload to code-scanning"
71-
uses: github/codeql-action/upload-sarif@2e230e8fe0ad3a14a340ad0815ddb96d599d2aff # v3.25.8
71+
uses: github/codeql-action/upload-sarif@23acc5c183826b7a8a97bce3cecc52db901f8251 # v3.25.10
7272
with:
7373
sarif_file: results.sarif

‎.github/workflows/pre-commit.yml

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -26,7 +26,7 @@ jobs:
2626
pylint
2727
2828
- name: Checkout DPNP repo
29-
uses: actions/checkout@a5ac7e51b41094c92402da3b24376905380afc29 # v4.1.6
29+
uses: actions/checkout@692973e3d937129bcbf40652eb9f2f61becf3332 # v4.1.7
3030

3131
- name: Set up python
3232
uses: actions/setup-python@82c7e631bb3cdc910f68e0081d67478d79c6982d # v5.1.0

‎dpnp/CMakeLists.txt

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -60,6 +60,7 @@ add_subdirectory(backend/extensions/blas)
6060
add_subdirectory(backend/extensions/lapack)
6161
add_subdirectory(backend/extensions/vm)
6262
add_subdirectory(backend/extensions/sycl_ext)
63+
add_subdirectory(backend/extensions/ufunc)
6364

6465
add_subdirectory(dpnp_algo)
6566
add_subdirectory(dpnp_utils)

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

Lines changed: 79 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,79 @@
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+
set(_elementwise_sources
27+
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/common.cpp
28+
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fabs.cpp
29+
)
30+
31+
set(python_module_name _ufunc_impl)
32+
33+
set(_module_src
34+
# TODO: remove sources from `elementwise_functions` folder
35+
${CMAKE_CURRENT_SOURCE_DIR}/../elementwise_functions/elementwise_functions_type_utils.cpp
36+
${CMAKE_CURRENT_SOURCE_DIR}/../elementwise_functions/simplify_iteration_space.cpp
37+
${CMAKE_CURRENT_SOURCE_DIR}/ufunc_py.cpp
38+
${_elementwise_sources}
39+
)
40+
41+
pybind11_add_module(${python_module_name} MODULE ${_module_src})
42+
add_sycl_to_target(TARGET ${python_module_name} SOURCES ${_module_src})
43+
44+
if (WIN32)
45+
if (${CMAKE_VERSION} VERSION_LESS "3.27")
46+
# this is a work-around for target_link_options inserting option after -link option, cause
47+
# linker to ignore it.
48+
set(CMAKE_CXX_LINK_FLAGS "${CMAKE_CXX_LINK_FLAGS} -fsycl-device-code-split=per_kernel")
49+
endif()
50+
endif()
51+
52+
set_target_properties(${python_module_name} PROPERTIES CMAKE_POSITION_INDEPENDENT_CODE ON)
53+
54+
target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../)
55+
56+
target_include_directories(${python_module_name} PUBLIC ${Dpctl_INCLUDE_DIR})
57+
target_include_directories(${python_module_name} PUBLIC ${Dpctl_TENSOR_INCLUDE_DIR})
58+
59+
if (WIN32)
60+
target_compile_options(${python_module_name} PRIVATE
61+
/clang:-fno-approx-func
62+
/clang:-fno-finite-math-only
63+
)
64+
else()
65+
target_compile_options(${python_module_name} PRIVATE
66+
-fno-approx-func
67+
-fno-finite-math-only
68+
)
69+
endif()
70+
71+
target_link_options(${python_module_name} PUBLIC -fsycl-device-code-split=per_kernel)
72+
73+
if (DPNP_GENERATE_COVERAGE)
74+
target_link_options(${python_module_name} PRIVATE -fprofile-instr-generate -fcoverage-mapping)
75+
endif()
76+
77+
install(TARGETS ${python_module_name}
78+
DESTINATION "dpnp/backend/extensions/ufunc"
79+
)

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

Lines changed: 41 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,41 @@
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 <pybind11/pybind11.h>
27+
28+
#include "fabs.hpp"
29+
30+
namespace py = pybind11;
31+
32+
namespace dpnp::extensions::ufunc
33+
{
34+
/**
35+
* @brief Add elementwise functions to Python module
36+
*/
37+
void init_elementwise_functions(py::module_ m)
38+
{
39+
init_fabs(m);
40+
}
41+
} // namespace dpnp::extensions::ufunc

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

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

Lines changed: 128 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,128 @@
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 "fabs.hpp"
31+
#include "kernels/elementwise_functions/fabs.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::fabs<T> function is available.
59+
*
60+
* @tparam T Type of input vector `a` and of result vector `y`.
61+
*/
62+
template <typename T>
63+
struct OutputType
64+
{
65+
using value_type =
66+
typename std::disjunction<td_ns::TypeMapResultEntry<T, sycl::half>,
67+
td_ns::TypeMapResultEntry<T, float>,
68+
td_ns::TypeMapResultEntry<T, double>,
69+
td_ns::DefaultResultEntry<void>>::result_type;
70+
};
71+
72+
using dpnp::kernels::fabs::FabsFunctor;
73+
74+
template <typename argT,
75+
typename resT = argT,
76+
unsigned int vec_sz = 4,
77+
unsigned int n_vecs = 2,
78+
bool enable_sg_loadstore = true>
79+
using ContigFunctor = ew_cmn_ns::UnaryContigFunctor<argT,
80+
resT,
81+
FabsFunctor<argT, resT>,
82+
vec_sz,
83+
n_vecs,
84+
enable_sg_loadstore>;
85+
86+
template <typename argTy, typename resTy, typename IndexerT>
87+
using StridedFunctor = ew_cmn_ns::
88+
UnaryStridedFunctor<argTy, resTy, IndexerT, FabsFunctor<argTy, resTy>>;
89+
90+
using ew_cmn_ns::unary_contig_impl_fn_ptr_t;
91+
using ew_cmn_ns::unary_strided_impl_fn_ptr_t;
92+
93+
static unary_contig_impl_fn_ptr_t fabs_contig_dispatch_vector[td_ns::num_types];
94+
static int fabs_output_typeid_vector[td_ns::num_types];
95+
static unary_strided_impl_fn_ptr_t
96+
fabs_strided_dispatch_vector[td_ns::num_types];
97+
98+
MACRO_POPULATE_DISPATCH_VECTORS(fabs);
99+
} // namespace impl
100+
101+
void init_fabs(py::module_ m)
102+
{
103+
using arrayT = dpctl::tensor::usm_ndarray;
104+
using event_vecT = std::vector<sycl::event>;
105+
{
106+
impl::populate_fabs_dispatch_vectors();
107+
using impl::fabs_contig_dispatch_vector;
108+
using impl::fabs_output_typeid_vector;
109+
using impl::fabs_strided_dispatch_vector;
110+
111+
auto fabs_pyapi = [&](const arrayT &src, const arrayT &dst,
112+
sycl::queue &exec_q,
113+
const event_vecT &depends = {}) {
114+
return py_int::py_unary_ufunc(
115+
src, dst, exec_q, depends, fabs_output_typeid_vector,
116+
fabs_contig_dispatch_vector, fabs_strided_dispatch_vector);
117+
};
118+
m.def("_fabs", fabs_pyapi, "", py::arg("src"), py::arg("dst"),
119+
py::arg("sycl_queue"), py::arg("depends") = py::list());
120+
121+
auto fabs_result_type_pyapi = [&](const py::dtype &dtype) {
122+
return py_int::py_unary_ufunc_result_type(
123+
dtype, fabs_output_typeid_vector);
124+
};
125+
m.def("_fabs_result_type", fabs_result_type_pyapi);
126+
}
127+
}
128+
} // namespace dpnp::extensions::ufunc

0 commit comments

Comments
 (0)