|
| 1 | +//***************************************************************************** |
| 2 | +// Copyright (c) 2026, 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 | +// - Neither the name of the copyright holder nor the names of its contributors |
| 13 | +// may be used to endorse or promote products derived from this software |
| 14 | +// without specific prior written permission. |
| 15 | +// |
| 16 | +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" |
| 17 | +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE |
| 18 | +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE |
| 19 | +// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE |
| 20 | +// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR |
| 21 | +// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF |
| 22 | +// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS |
| 23 | +// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN |
| 24 | +// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) |
| 25 | +// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF |
| 26 | +// THE POSSIBILITY OF SUCH DAMAGE. |
| 27 | +//***************************************************************************** |
| 28 | +// |
| 29 | +//===---------------------------------------------------------------------===// |
| 30 | +/// |
| 31 | +/// \file |
| 32 | +/// This file defines functions of dpctl.tensor._tensor_elementwise_impl |
| 33 | +/// extension, specifically functions for elementwise operations. |
| 34 | +//===---------------------------------------------------------------------===// |
| 35 | + |
| 36 | +#include <vector> |
| 37 | + |
| 38 | +#include <sycl/sycl.hpp> |
| 39 | + |
| 40 | +#include "dpnp4pybind11.hpp" |
| 41 | +#include <pybind11/numpy.h> |
| 42 | +#include <pybind11/pybind11.h> |
| 43 | +#include <pybind11/stl.h> |
| 44 | + |
| 45 | +#include "elementwise_functions.hpp" |
| 46 | +#include "logaddexp.hpp" |
| 47 | +#include "utils/type_dispatch.hpp" |
| 48 | + |
| 49 | +#include "kernels/elementwise_functions/common.hpp" |
| 50 | +#include "kernels/elementwise_functions/logaddexp.hpp" |
| 51 | + |
| 52 | +namespace dpctl::tensor::py_internal |
| 53 | +{ |
| 54 | + |
| 55 | +namespace py = pybind11; |
| 56 | +namespace td_ns = dpctl::tensor::type_dispatch; |
| 57 | + |
| 58 | +namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common; |
| 59 | +using ew_cmn_ns::binary_contig_impl_fn_ptr_t; |
| 60 | +using ew_cmn_ns::binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t; |
| 61 | +using ew_cmn_ns::binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t; |
| 62 | +using ew_cmn_ns::binary_strided_impl_fn_ptr_t; |
| 63 | + |
| 64 | +// B15: ===== LOGADDEXP (x1, x2) |
| 65 | +namespace impl |
| 66 | +{ |
| 67 | +namespace logaddexp_fn_ns = dpctl::tensor::kernels::logaddexp; |
| 68 | + |
| 69 | +static binary_contig_impl_fn_ptr_t |
| 70 | + logaddexp_contig_dispatch_table[td_ns::num_types][td_ns::num_types]; |
| 71 | +static int logaddexp_output_id_table[td_ns::num_types][td_ns::num_types]; |
| 72 | + |
| 73 | +static binary_strided_impl_fn_ptr_t |
| 74 | + logaddexp_strided_dispatch_table[td_ns::num_types][td_ns::num_types]; |
| 75 | + |
| 76 | +void populate_logaddexp_dispatch_tables(void) |
| 77 | +{ |
| 78 | + using namespace td_ns; |
| 79 | + namespace fn_ns = logaddexp_fn_ns; |
| 80 | + |
| 81 | + // which input types are supported, and what is the type of the result |
| 82 | + using fn_ns::LogAddExpTypeMapFactory; |
| 83 | + DispatchTableBuilder<int, LogAddExpTypeMapFactory, num_types> dtb1; |
| 84 | + dtb1.populate_dispatch_table(logaddexp_output_id_table); |
| 85 | + |
| 86 | + // function pointers for operation on general strided arrays |
| 87 | + using fn_ns::LogAddExpStridedFactory; |
| 88 | + DispatchTableBuilder<binary_strided_impl_fn_ptr_t, LogAddExpStridedFactory, |
| 89 | + num_types> |
| 90 | + dtb2; |
| 91 | + dtb2.populate_dispatch_table(logaddexp_strided_dispatch_table); |
| 92 | + |
| 93 | + // function pointers for operation on contiguous inputs and output |
| 94 | + using fn_ns::LogAddExpContigFactory; |
| 95 | + DispatchTableBuilder<binary_contig_impl_fn_ptr_t, LogAddExpContigFactory, |
| 96 | + num_types> |
| 97 | + dtb3; |
| 98 | + dtb3.populate_dispatch_table(logaddexp_contig_dispatch_table); |
| 99 | +}; |
| 100 | + |
| 101 | +} // namespace impl |
| 102 | + |
| 103 | +void init_logaddexp(py::module_ m) |
| 104 | +{ |
| 105 | + using arrayT = dpctl::tensor::usm_ndarray; |
| 106 | + using event_vecT = std::vector<sycl::event>; |
| 107 | + { |
| 108 | + impl::populate_logaddexp_dispatch_tables(); |
| 109 | + using impl::logaddexp_contig_dispatch_table; |
| 110 | + using impl::logaddexp_output_id_table; |
| 111 | + using impl::logaddexp_strided_dispatch_table; |
| 112 | + |
| 113 | + auto logaddexp_pyapi = [&](const arrayT &src1, const arrayT &src2, |
| 114 | + const arrayT &dst, sycl::queue &exec_q, |
| 115 | + const event_vecT &depends = {}) { |
| 116 | + return py_binary_ufunc( |
| 117 | + src1, src2, dst, exec_q, depends, logaddexp_output_id_table, |
| 118 | + // function pointers to handle operation on contiguous arrays |
| 119 | + // (pointers may be nullptr) |
| 120 | + logaddexp_contig_dispatch_table, |
| 121 | + // function pointers to handle operation on strided arrays (most |
| 122 | + // general case) |
| 123 | + logaddexp_strided_dispatch_table, |
| 124 | + // function pointers to handle operation of c-contig matrix and |
| 125 | + // c-contig row with broadcasting (may be nullptr) |
| 126 | + td_ns::NullPtrTable< |
| 127 | + binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t>{}, |
| 128 | + // function pointers to handle operation of c-contig matrix and |
| 129 | + // c-contig row with broadcasting (may be nullptr) |
| 130 | + td_ns::NullPtrTable< |
| 131 | + binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t>{}); |
| 132 | + }; |
| 133 | + auto logaddexp_result_type_pyapi = [&](const py::dtype &dtype1, |
| 134 | + const py::dtype &dtype2) { |
| 135 | + return py_binary_ufunc_result_type(dtype1, dtype2, |
| 136 | + logaddexp_output_id_table); |
| 137 | + }; |
| 138 | + m.def("_logaddexp", logaddexp_pyapi, "", py::arg("src1"), |
| 139 | + py::arg("src2"), py::arg("dst"), py::arg("sycl_queue"), |
| 140 | + py::arg("depends") = py::list()); |
| 141 | + m.def("_logaddexp_result_type", logaddexp_result_type_pyapi, ""); |
| 142 | + } |
| 143 | +} |
| 144 | + |
| 145 | +} // namespace dpctl::tensor::py_internal |
0 commit comments