DPNP C++ backend kernel library 0.21.0dev3
Data Parallel Extension for NumPy*
Loading...
Searching...
No Matches
common.hpp
1//*****************************************************************************
2// Copyright (c) 2023, 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#pragma once
30
31#include <cstddef>
32#include <type_traits>
33#include <utility>
34#include <vector>
35
36#include <oneapi/mkl.hpp>
37#include <pybind11/pybind11.h>
38#include <sycl/sycl.hpp>
39
40#include "dpnp4pybind11.hpp"
41
42// utils extension header
43#include "ext/common.hpp"
44
45// dpnp tensor headers
46#include "utils/memory_overlap.hpp"
47#include "utils/type_dispatch.hpp"
48
49namespace ext_ns = ext::common;
50namespace py = pybind11;
51namespace td_ns = dpnp::tensor::type_dispatch;
52
53namespace dpnp::extensions::vm::py_internal
54{
55template <typename output_typesT, typename contig_dispatchT>
56bool need_to_call_unary_ufunc(sycl::queue &exec_q,
57 const dpnp::tensor::usm_ndarray &src,
58 const dpnp::tensor::usm_ndarray &dst,
59 const output_typesT &output_type_vec,
60 const contig_dispatchT &contig_dispatch_vector)
61{
62 // check type_nums
63 int src_typenum = src.get_typenum();
64 int dst_typenum = dst.get_typenum();
65
66 auto array_types = td_ns::usm_ndarray_types();
67 int src_typeid = array_types.typenum_to_lookup_id(src_typenum);
68 int dst_typeid = array_types.typenum_to_lookup_id(dst_typenum);
69
70 // check that types are supported
71 int func_output_typeid = output_type_vec[src_typeid];
72 if (dst_typeid != func_output_typeid) {
73 return false;
74 }
75
76 // OneMKL VM functions perform a copy on host if no double type support
77 if (!exec_q.get_device().has(sycl::aspect::fp64)) {
78 return false;
79 }
80
81 // check that queues are compatible
82 if (!dpnp::utils::queues_are_compatible(exec_q, {src, dst})) {
83 return false;
84 }
85
86 // dimensions must be the same
87 int dst_nd = dst.get_ndim();
88 if (dst_nd != src.get_ndim()) {
89 return false;
90 }
91 else if (dst_nd == 0) {
92 // don't call OneMKL for 0d arrays
93 return false;
94 }
95
96 // shapes must be the same
97 const py::ssize_t *src_shape = src.get_shape_raw();
98 const py::ssize_t *dst_shape = dst.get_shape_raw();
99 bool shapes_equal(true);
100 size_t src_nelems(1);
101
102 for (int i = 0; i < dst_nd; ++i) {
103 src_nelems *= static_cast<size_t>(src_shape[i]);
104 shapes_equal = shapes_equal && (src_shape[i] == dst_shape[i]);
105 }
106 if (!shapes_equal) {
107 return false;
108 }
109
110 // if nelems is zero, return false
111 if (src_nelems == 0) {
112 return false;
113 }
114
115 // ensure that output is ample enough to accommodate all elements
116 auto dst_offsets = dst.get_minmax_offsets();
117 // destination must be ample enough to accommodate all elements
118 {
119 size_t range =
120 static_cast<size_t>(dst_offsets.second - dst_offsets.first);
121 if (range + 1 < src_nelems) {
122 return false;
123 }
124 }
125
126 // check memory overlap
127 auto const &overlap = dpnp::tensor::overlap::MemoryOverlap();
128 if (overlap(src, dst)) {
129 return false;
130 }
131
132 // support only contiguous inputs
133 bool is_src_c_contig = src.is_c_contiguous();
134 bool is_dst_c_contig = dst.is_c_contiguous();
135
136 bool all_c_contig = (is_src_c_contig && is_dst_c_contig);
137 if (!all_c_contig) {
138 return false;
139 }
140
141 // MKL function is not defined for the type
142 if (contig_dispatch_vector[src_typeid] == nullptr) {
143 return false;
144 }
145 return true;
146}
147
148template <typename output_typesT, typename contig_dispatchT>
149bool need_to_call_unary_two_outputs_ufunc(
150 sycl::queue &exec_q,
151 const dpnp::tensor::usm_ndarray &src,
152 const dpnp::tensor::usm_ndarray &dst1,
153 const dpnp::tensor::usm_ndarray &dst2,
154 const output_typesT &output_type_vec,
155 const contig_dispatchT &contig_dispatch_vector)
156{
157 // check type_nums
158 int src_typenum = src.get_typenum();
159 int dst1_typenum = dst1.get_typenum();
160 int dst2_typenum = dst2.get_typenum();
161
162 const auto &array_types = td_ns::usm_ndarray_types();
163 int src_typeid = array_types.typenum_to_lookup_id(src_typenum);
164 int dst1_typeid = array_types.typenum_to_lookup_id(dst1_typenum);
165 int dst2_typeid = array_types.typenum_to_lookup_id(dst2_typenum);
166
167 std::pair<int, int> func_output_typeids = output_type_vec[src_typeid];
168
169 // check that types are supported
170 if (dst1_typeid != func_output_typeids.first ||
171 dst2_typeid != func_output_typeids.second) {
172 return false;
173 }
174
175 // OneMKL VM functions perform a copy on host if no double type support
176 if (!exec_q.get_device().has(sycl::aspect::fp64)) {
177 return false;
178 }
179
180 // check that queues are compatible
181 if (!dpnp::utils::queues_are_compatible(exec_q, {src, dst1, dst2})) {
182 return false;
183 }
184
185 // dimensions must be the same
186 int src_nd = src.get_ndim();
187 int dst1_nd = dst1.get_ndim();
188 int dst2_nd = dst2.get_ndim();
189 if (src_nd != dst1_nd || src_nd != dst2_nd) {
190 return false;
191 }
192 else if (dst1_nd == 0 || dst2_nd == 0) {
193 // don't call OneMKL for 0d arrays
194 return false;
195 }
196
197 // shapes must be the same
198 const py::ssize_t *src_shape = src.get_shape_raw();
199 const py::ssize_t *dst1_shape = dst1.get_shape_raw();
200 const py::ssize_t *dst2_shape = dst2.get_shape_raw();
201 bool shapes_equal(true);
202 size_t src_nelems(1);
203
204 for (int i = 0; i < src_nd; ++i) {
205 src_nelems *= static_cast<std::size_t>(src_shape[i]);
206 shapes_equal = shapes_equal && (src_shape[i] == dst1_shape[i]) &&
207 (src_shape[i] == dst2_shape[i]);
208 }
209 if (!shapes_equal) {
210 return false;
211 }
212
213 // if nelems is zero, return false
214 if (src_nelems == 0) {
215 return false;
216 }
217
218 // ensure that outputs are ample enough to accommodate all elements
219 auto dst1_offsets = dst1.get_minmax_offsets();
220 auto dst2_offsets = dst2.get_minmax_offsets();
221 // destinations must be ample enough to accommodate all elements
222 {
223 size_t range1 =
224 static_cast<size_t>(dst1_offsets.second - dst1_offsets.first);
225 size_t range2 =
226 static_cast<size_t>(dst2_offsets.second - dst2_offsets.first);
227 if ((range1 + 1 < src_nelems) || (range2 + 1 < src_nelems)) {
228 return false;
229 }
230 }
231
232 // check memory overlap
233 auto const &overlap = dpnp::tensor::overlap::MemoryOverlap();
234 if (overlap(src, dst1) || overlap(src, dst2) || overlap(dst1, dst2)) {
235 return false;
236 }
237
238 // support only contiguous inputs
239 bool is_src_c_contig = src.is_c_contiguous();
240 bool is_dst1_c_contig = dst1.is_c_contiguous();
241 bool is_dst2_c_contig = dst2.is_c_contiguous();
242
243 bool all_c_contig =
244 (is_src_c_contig && is_dst1_c_contig && is_dst2_c_contig);
245 if (!all_c_contig) {
246 return false;
247 }
248
249 // MKL function is not defined for the type
250 if (contig_dispatch_vector[src_typeid] == nullptr) {
251 return false;
252 }
253 return true;
254}
255
256template <typename output_typesT, typename contig_dispatchT>
257bool need_to_call_binary_ufunc(sycl::queue &exec_q,
258 const dpnp::tensor::usm_ndarray &src1,
259 const dpnp::tensor::usm_ndarray &src2,
260 const dpnp::tensor::usm_ndarray &dst,
261 const output_typesT &output_type_table,
262 const contig_dispatchT &contig_dispatch_table)
263{
264 // check type_nums
265 int src1_typenum = src1.get_typenum();
266 int src2_typenum = src2.get_typenum();
267 int dst_typenum = dst.get_typenum();
268
269 auto array_types = td_ns::usm_ndarray_types();
270 int src1_typeid = array_types.typenum_to_lookup_id(src1_typenum);
271 int src2_typeid = array_types.typenum_to_lookup_id(src2_typenum);
272 int dst_typeid = array_types.typenum_to_lookup_id(dst_typenum);
273
274 // check that types are supported
275 int output_typeid = output_type_table[src1_typeid][src2_typeid];
276 if (output_typeid != dst_typeid) {
277 return false;
278 }
279
280 // types must be the same
281 if (src1_typeid != src2_typeid) {
282 return false;
283 }
284
285 // OneMKL VM functions perform a copy on host if no double type support
286 if (!exec_q.get_device().has(sycl::aspect::fp64)) {
287 return false;
288 }
289
290 // check that queues are compatible
291 if (!dpnp::utils::queues_are_compatible(exec_q, {src1, src2, dst})) {
292 return false;
293 }
294
295 // dimensions must be the same
296 int dst_nd = dst.get_ndim();
297 if (dst_nd != src1.get_ndim() || dst_nd != src2.get_ndim()) {
298 return false;
299 }
300 else if (dst_nd == 0) {
301 // don't call OneMKL for 0d arrays
302 return false;
303 }
304
305 // shapes must be the same
306 const py::ssize_t *src1_shape = src1.get_shape_raw();
307 const py::ssize_t *src2_shape = src2.get_shape_raw();
308 const py::ssize_t *dst_shape = dst.get_shape_raw();
309 bool shapes_equal(true);
310 size_t src_nelems(1);
311
312 for (int i = 0; i < dst_nd; ++i) {
313 src_nelems *= static_cast<size_t>(src1_shape[i]);
314 shapes_equal = shapes_equal && (src1_shape[i] == dst_shape[i] &&
315 src2_shape[i] == dst_shape[i]);
316 }
317 if (!shapes_equal) {
318 return false;
319 }
320
321 // if nelems is zero, return false
322 if (src_nelems == 0) {
323 return false;
324 }
325
326 // ensure that output is ample enough to accommodate all elements
327 auto dst_offsets = dst.get_minmax_offsets();
328 // destination must be ample enough to accommodate all elements
329 {
330 size_t range =
331 static_cast<size_t>(dst_offsets.second - dst_offsets.first);
332 if (range + 1 < src_nelems) {
333 return false;
334 }
335 }
336
337 // check memory overlap
338 auto const &overlap = dpnp::tensor::overlap::MemoryOverlap();
339 if (overlap(src1, dst) || overlap(src2, dst)) {
340 return false;
341 }
342
343 // support only contiguous inputs
344 bool is_src1_c_contig = src1.is_c_contiguous();
345 bool is_src2_c_contig = src2.is_c_contiguous();
346 bool is_dst_c_contig = dst.is_c_contiguous();
347
348 bool all_c_contig =
349 (is_src1_c_contig && is_src2_c_contig && is_dst_c_contig);
350 if (!all_c_contig) {
351 return false;
352 }
353
354 // MKL function is not defined for the type
355 if (contig_dispatch_table[src1_typeid] == nullptr) {
356 return false;
357 }
358 return true;
359}
360
366#define MACRO_POPULATE_DISPATCH_VECTORS(__name__) \
367 template <typename fnT, typename T> \
368 struct ContigFactory \
369 { \
370 fnT get() \
371 { \
372 if constexpr (std::is_same_v<typename OutputType<T>::value_type, \
373 void>) { \
374 return nullptr; \
375 } \
376 else { \
377 return __name__##_contig_impl<T>; \
378 } \
379 } \
380 }; \
381 \
382 template <typename fnT, typename T> \
383 struct TypeMapFactory \
384 { \
385 std::enable_if_t<std::is_same<fnT, int>::value, int> get() \
386 { \
387 using rT = typename OutputType<T>::value_type; \
388 return td_ns::GetTypeid<rT>{}.get(); \
389 } \
390 }; \
391 \
392 static void populate_dispatch_vectors(void) \
393 { \
394 ext_ns::init_dispatch_vector<int, TypeMapFactory>( \
395 output_typeid_vector); \
396 ext_ns::init_dispatch_vector<unary_contig_impl_fn_ptr_t, \
397 ContigFactory>(contig_dispatch_vector); \
398 };
399
405#define MACRO_POPULATE_DISPATCH_2OUTS_VECTORS(__name__) \
406 template <typename fnT, typename T> \
407 struct ContigFactory \
408 { \
409 fnT get() \
410 { \
411 if constexpr (std::is_same_v<typename OutputType<T>::value_type1, \
412 void> || \
413 std::is_same_v<typename OutputType<T>::value_type2, \
414 void>) { \
415 fnT fn = nullptr; \
416 return fn; \
417 } \
418 else { \
419 fnT fn = __name__##_contig_impl<T>; \
420 return fn; \
421 } \
422 } \
423 }; \
424 \
425 template <typename fnT, typename T> \
426 struct TypeMapFactory \
427 { \
428 std::enable_if_t<std::is_same<fnT, std::pair<int, int>>::value, \
429 std::pair<int, int>> \
430 get() \
431 { \
432 using rT1 = typename OutputType<T>::value_type1; \
433 using rT2 = typename OutputType<T>::value_type2; \
434 return std::make_pair(td_ns::GetTypeid<rT1>{}.get(), \
435 td_ns::GetTypeid<rT2>{}.get()); \
436 } \
437 }; \
438 \
439 static void populate_dispatch_vectors(void) \
440 { \
441 ext_ns::init_dispatch_vector<std::pair<int, int>, TypeMapFactory>( \
442 output_typeid_vector); \
443 ext_ns::init_dispatch_vector<unary_two_outputs_contig_impl_fn_ptr_t, \
444 ContigFactory>(contig_dispatch_vector); \
445 };
446
452#define MACRO_POPULATE_DISPATCH_TABLES(__name__) \
453 template <typename fnT, typename T1, typename T2> \
454 struct ContigFactory \
455 { \
456 fnT get() \
457 { \
458 if constexpr (std::is_same_v< \
459 typename OutputType<T1, T2>::value_type, \
460 void>) { \
461 return nullptr; \
462 } \
463 else { \
464 return __name__##_contig_impl<T1, T2>; \
465 } \
466 } \
467 }; \
468 \
469 template <typename fnT, typename T1, typename T2> \
470 struct TypeMapFactory \
471 { \
472 std::enable_if_t<std::is_same<fnT, int>::value, int> get() \
473 { \
474 using rT = typename OutputType<T1, T2>::value_type; \
475 return td_ns::GetTypeid<rT>{}.get(); \
476 } \
477 }; \
478 \
479 static void populate_dispatch_tables(void) \
480 { \
481 ext_ns::init_dispatch_table<int, TypeMapFactory>( \
482 output_typeid_vector); \
483 ext_ns::init_dispatch_table<binary_contig_impl_fn_ptr_t, \
484 ContigFactory>(contig_dispatch_vector); \
485 };
486} // namespace dpnp::extensions::vm::py_internal