• Home
  • Features
  • Pricing
  • Docs
  • Announcements
  • Sign In

IntelPython / dpnp / 9797394893

04 Jul 2024 04:14PM UTC coverage: 54.375% (+0.4%) from 54.011%
9797394893

push

github

web-flow
Rework implementation of `dpnp.fmax` and `dpnp.fmin` functions (#1905)

* Implement dpnp.fmax and dpnp.fmin functions

* Updated existing tests and added new ones

* Removed unused code from cython backend

* Removed a reference to original descriptor

3573 of 11085 branches covered (32.23%)

Branch coverage included in aggregate %.

181 of 228 new or added lines in 11 files covered. (79.39%)

307 existing lines in 5 files now uncovered.

12993 of 19381 relevant lines covered (67.04%)

17127.72 hits per line

Source File
Press 'n' to go to next uncovered line, 'b' for previous

88.24
/dpnp/backend/extensions/vm/fmax.cpp
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 <oneapi/mkl.hpp>
27
#include <sycl/sycl.hpp>
28

29
#include "dpctl4pybind11.hpp"
30

31
#include "common.hpp"
32
#include "fmax.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
#include "utils/type_utils.hpp"
43

44
namespace dpnp::extensions::vm
45
{
46
namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
47
namespace py = pybind11;
48
namespace py_int = dpnp::extensions::py_internal;
49
namespace td_ns = dpctl::tensor::type_dispatch;
50
namespace tu_ns = dpctl::tensor::type_utils;
51

52
namespace impl
53
{
54
// OneMKL namespace with VM functions
55
namespace mkl_vm = oneapi::mkl::vm;
56

57
/**
58
 * @brief A factory to define pairs of supported types for which
59
 * MKL VM library provides support in oneapi::mkl::vm::fmax<T> function.
60
 *
61
 * @tparam T Type of input vectors `a` and `b` and of result vector `y`.
62
 */
63
template <typename T1, typename T2>
64
struct OutputType
65
{
66
    using value_type = typename std::disjunction<
67
        td_ns::BinaryTypeMapResultEntry<T1, double, T2, double, double>,
68
        td_ns::BinaryTypeMapResultEntry<T1, float, T2, float, float>,
69
        td_ns::DefaultResultEntry<void>>::result_type;
70
};
71

72
template <typename T1, typename T2>
73
static sycl::event fmax_contig_impl(sycl::queue &exec_q,
74
                                    std::size_t in_n,
75
                                    const char *in_a,
76
                                    py::ssize_t a_offset,
77
                                    const char *in_b,
78
                                    py::ssize_t b_offset,
79
                                    char *out_y,
80
                                    py::ssize_t out_offset,
81
                                    const std::vector<sycl::event> &depends)
82
{
39✔
83
    tu_ns::validate_type_for_device<T1>(exec_q);
39✔
84
    tu_ns::validate_type_for_device<T2>(exec_q);
39✔
85

86
    if ((a_offset != 0) || (b_offset != 0) || (out_offset != 0)) {
39!
NEW
87
        throw std::runtime_error("Arrays offsets have to be equals to 0");
×
NEW
88
    }
×
89

90
    std::int64_t n = static_cast<std::int64_t>(in_n);
39✔
91
    const T1 *a = reinterpret_cast<const T1 *>(in_a);
39✔
92
    const T2 *b = reinterpret_cast<const T2 *>(in_b);
39✔
93

94
    using resTy = typename OutputType<T1, T2>::value_type;
39✔
95
    resTy *y = reinterpret_cast<resTy *>(out_y);
39✔
96

97
    return mkl_vm::fmax(exec_q,
39✔
98
                        n, // number of elements to be calculated
39✔
99
                        a, // pointer `a` containing 1st input vector of size n
39✔
100
                        b, // pointer `b` containing 2nd input vector of size n
39✔
101
                        y, // pointer `y` to the output vector of size n
39✔
102
                        depends);
39✔
103
}
39✔
104

105
using ew_cmn_ns::binary_contig_impl_fn_ptr_t;
106
using ew_cmn_ns::binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t;
107
using ew_cmn_ns::binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t;
108
using ew_cmn_ns::binary_strided_impl_fn_ptr_t;
109

110
static int output_typeid_vector[td_ns::num_types][td_ns::num_types];
111
static binary_contig_impl_fn_ptr_t contig_dispatch_vector[td_ns::num_types]
112
                                                         [td_ns::num_types];
113

114
MACRO_POPULATE_DISPATCH_TABLES(fmax);
115
} // namespace impl
116

117
void init_fmax(py::module_ m)
118
{
1✔
119
    using arrayT = dpctl::tensor::usm_ndarray;
1✔
120
    using event_vecT = std::vector<sycl::event>;
1✔
121

122
    impl::populate_dispatch_tables();
1✔
123
    using impl::contig_dispatch_vector;
1✔
124
    using impl::output_typeid_vector;
1✔
125

126
    auto fmax_pyapi = [&](sycl::queue &exec_q, const arrayT &src1,
1✔
127
                          const arrayT &src2, const arrayT &dst,
1✔
128
                          const event_vecT &depends = {}) {
39✔
129
        return py_int::py_binary_ufunc(
39✔
130
            src1, src2, dst, exec_q, depends, output_typeid_vector,
39✔
131
            contig_dispatch_vector,
39✔
132
            // no support of strided implementation in OneMKL
133
            td_ns::NullPtrTable<impl::binary_strided_impl_fn_ptr_t>{},
39✔
134
            // no support of C-contig row with broadcasting in OneMKL
135
            td_ns::NullPtrTable<
39✔
136
                impl::
39✔
137
                    binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t>{},
39✔
138
            td_ns::NullPtrTable<
39✔
139
                impl::
39✔
140
                    binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t>{});
39✔
141
    };
39✔
142
    m.def("_fmax", fmax_pyapi,
1✔
143
          "Call `fmax` function from OneMKL VM library to performs element "
1✔
144
          "by element computation of the modulus function of vector `src1` "
1✔
145
          "with respect to vector `src2` to resulting vector `dst`",
1✔
146
          py::arg("sycl_queue"), py::arg("src1"), py::arg("src2"),
1✔
147
          py::arg("dst"), py::arg("depends") = py::list());
1✔
148

149
    auto fmax_need_to_call_pyapi = [&](sycl::queue &exec_q, const arrayT &src1,
1✔
150
                                       const arrayT &src2, const arrayT &dst) {
101✔
151
        return py_internal::need_to_call_binary_ufunc(exec_q, src1, src2, dst,
101✔
152
                                                      output_typeid_vector,
101✔
153
                                                      contig_dispatch_vector);
101✔
154
    };
101✔
155
    m.def("_mkl_fmax_to_call", fmax_need_to_call_pyapi,
1✔
156
          "Check input arguments to answer if `fmax` function from "
1✔
157
          "OneMKL VM library can be used",
1✔
158
          py::arg("sycl_queue"), py::arg("src1"), py::arg("src2"),
1✔
159
          py::arg("dst"));
1✔
160
}
1✔
161
} // namespace dpnp::extensions::vm
STATUS · Troubleshooting · Open an Issue · Sales · Support · CAREERS · ENTERPRISE · START FREE · SCHEDULE DEMO
ANNOUNCEMENTS · TWITTER · TOS & SLA · Supported CI Services · What's a CI service? · Automated Testing

© 2026 Coveralls, Inc