DPNP C++ backend kernel library 0.21.0dev7
Data Parallel Extension for NumPy*
Loading...
Searching...
No Matches
populate.hpp
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// - 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 <type_traits>
32#include <utility>
33#include <vector>
34
35#include <pybind11/pybind11.h>
36
37// utils extension header
38#include "ext/common.hpp"
39
40namespace ext_ns = ext::common;
41
46#define MACRO_POPULATE_DISPATCH_VECTORS(__name__) \
47 template <typename T1, typename T2, unsigned int vec_sz, \
48 unsigned int n_vecs> \
49 class __name__##_contig_kernel; \
50 \
51 template <typename argTy> \
52 sycl::event __name__##_contig_impl( \
53 sycl::queue &exec_q, size_t nelems, const char *arg_p, char *res_p, \
54 const std::vector<sycl::event> &depends = {}) \
55 { \
56 return ew_cmn_ns::unary_contig_impl<argTy, OutputType, ContigFunctor, \
57 __name__##_contig_kernel>( \
58 exec_q, nelems, arg_p, res_p, depends); \
59 } \
60 \
61 template <typename fnT, typename T> \
62 struct __name__##_ContigFactory \
63 { \
64 fnT get() \
65 { \
66 if constexpr (std::is_same_v<typename OutputType<T>::value_type, \
67 void>) { \
68 fnT fn = nullptr; \
69 return fn; \
70 } \
71 else { \
72 fnT fn = __name__##_contig_impl<T>; \
73 return fn; \
74 } \
75 } \
76 }; \
77 \
78 template <typename fnT, typename T> \
79 struct __name__##_TypeMapFactory \
80 { \
81 std::enable_if_t<std::is_same<fnT, int>::value, int> get() \
82 { \
83 using rT = typename OutputType<T>::value_type; \
84 return td_ns::GetTypeid<rT>{}.get(); \
85 } \
86 }; \
87 \
88 template <typename T1, typename T2, typename T3> \
89 class __name__##_strided_kernel; \
90 \
91 template <typename argTy> \
92 sycl::event __name__##_strided_impl( \
93 sycl::queue &exec_q, size_t nelems, int nd, \
94 const py::ssize_t *shape_and_strides, const char *arg_p, \
95 py::ssize_t arg_offset, char *res_p, py::ssize_t res_offset, \
96 const std::vector<sycl::event> &depends, \
97 const std::vector<sycl::event> &additional_depends) \
98 { \
99 return ew_cmn_ns::unary_strided_impl< \
100 argTy, OutputType, StridedFunctor, __name__##_strided_kernel>( \
101 exec_q, nelems, nd, shape_and_strides, arg_p, arg_offset, res_p, \
102 res_offset, depends, additional_depends); \
103 } \
104 \
105 template <typename fnT, typename T> \
106 struct __name__##_StridedFactory \
107 { \
108 fnT get() \
109 { \
110 if constexpr (std::is_same_v<typename OutputType<T>::value_type, \
111 void>) { \
112 fnT fn = nullptr; \
113 return fn; \
114 } \
115 else { \
116 fnT fn = __name__##_strided_impl<T>; \
117 return fn; \
118 } \
119 } \
120 }; \
121 \
122 void populate_##__name__##_dispatch_vectors(void) \
123 { \
124 ext_ns::init_dispatch_vector<unary_contig_impl_fn_ptr_t, \
125 __name__##_ContigFactory>( \
126 __name__##_contig_dispatch_vector); \
127 ext_ns::init_dispatch_vector<unary_strided_impl_fn_ptr_t, \
128 __name__##_StridedFactory>( \
129 __name__##_strided_dispatch_vector); \
130 ext_ns::init_dispatch_vector<int, __name__##_TypeMapFactory>( \
131 __name__##_output_typeid_vector); \
132 };
133
138#define MACRO_POPULATE_DISPATCH_2OUTS_VECTORS(__name__) \
139 template <typename T1, typename T2, typename T3, unsigned int vec_sz, \
140 unsigned int n_vecs> \
141 class __name__##_contig_kernel; \
142 \
143 template <typename argTy> \
144 sycl::event __name__##_contig_impl( \
145 sycl::queue &exec_q, size_t nelems, const char *arg_p, char *res1_p, \
146 char *res2_p, const std::vector<sycl::event> &depends = {}) \
147 { \
148 return ew_cmn_ns::unary_two_outputs_contig_impl< \
149 argTy, OutputType, ContigFunctor, __name__##_contig_kernel>( \
150 exec_q, nelems, arg_p, res1_p, res2_p, depends); \
151 } \
152 \
153 template <typename fnT, typename T> \
154 struct __name__##_ContigFactory \
155 { \
156 fnT get() \
157 { \
158 if constexpr (std::is_same_v<typename OutputType<T>::value_type1, \
159 void> || \
160 std::is_same_v<typename OutputType<T>::value_type2, \
161 void>) { \
162 fnT fn = nullptr; \
163 return fn; \
164 } \
165 else { \
166 fnT fn = __name__##_contig_impl<T>; \
167 return fn; \
168 } \
169 } \
170 }; \
171 \
172 template <typename fnT, typename T> \
173 struct __name__##_TypeMapFactory \
174 { \
175 std::enable_if_t<std::is_same<fnT, std::pair<int, int>>::value, \
176 std::pair<int, int>> \
177 get() \
178 { \
179 using rT1 = typename OutputType<T>::value_type1; \
180 using rT2 = typename OutputType<T>::value_type2; \
181 return std::make_pair(td_ns::GetTypeid<rT1>{}.get(), \
182 td_ns::GetTypeid<rT2>{}.get()); \
183 } \
184 }; \
185 \
186 template <typename T1, typename T2, typename T3, typename T4> \
187 class __name__##_strided_kernel; \
188 \
189 template <typename argTy> \
190 sycl::event __name__##_strided_impl( \
191 sycl::queue &exec_q, size_t nelems, int nd, \
192 const py::ssize_t *shape_and_strides, const char *arg_p, \
193 py::ssize_t arg_offset, char *res1_p, py::ssize_t res1_offset, \
194 char *res2_p, py::ssize_t res2_offset, \
195 const std::vector<sycl::event> &depends, \
196 const std::vector<sycl::event> &additional_depends) \
197 { \
198 return ew_cmn_ns::unary_two_outputs_strided_impl< \
199 argTy, OutputType, StridedFunctor, __name__##_strided_kernel>( \
200 exec_q, nelems, nd, shape_and_strides, arg_p, arg_offset, res1_p, \
201 res1_offset, res2_p, res2_offset, depends, additional_depends); \
202 } \
203 \
204 template <typename fnT, typename T> \
205 struct __name__##_StridedFactory \
206 { \
207 fnT get() \
208 { \
209 if constexpr (std::is_same_v<typename OutputType<T>::value_type1, \
210 void> || \
211 std::is_same_v<typename OutputType<T>::value_type2, \
212 void>) { \
213 fnT fn = nullptr; \
214 return fn; \
215 } \
216 else { \
217 fnT fn = __name__##_strided_impl<T>; \
218 return fn; \
219 } \
220 } \
221 }; \
222 \
223 void populate_##__name__##_dispatch_vectors(void) \
224 { \
225 ext_ns::init_dispatch_vector<unary_two_outputs_contig_impl_fn_ptr_t, \
226 __name__##_ContigFactory>( \
227 __name__##_contig_dispatch_vector); \
228 ext_ns::init_dispatch_vector<unary_two_outputs_strided_impl_fn_ptr_t, \
229 __name__##_StridedFactory>( \
230 __name__##_strided_dispatch_vector); \
231 ext_ns::init_dispatch_vector<std::pair<int, int>, \
232 __name__##_TypeMapFactory>( \
233 __name__##_output_typeid_vector); \
234 };
235
240#define MACRO_POPULATE_DISPATCH_TABLES(__name__) \
241 template <typename argT1, typename argT2, typename resT, \
242 unsigned int vec_sz, unsigned int n_vecs> \
243 class __name__##_contig_kernel; \
244 \
245 template <typename argTy1, typename argTy2> \
246 sycl::event __name__##_contig_impl( \
247 sycl::queue &exec_q, size_t nelems, const char *arg1_p, \
248 py::ssize_t arg1_offset, const char *arg2_p, py::ssize_t arg2_offset, \
249 char *res_p, py::ssize_t res_offset, \
250 const std::vector<sycl::event> &depends = {}) \
251 { \
252 return ew_cmn_ns::binary_contig_impl<argTy1, argTy2, OutputType, \
253 ContigFunctor, \
254 __name__##_contig_kernel>( \
255 exec_q, nelems, arg1_p, arg1_offset, arg2_p, arg2_offset, res_p, \
256 res_offset, depends); \
257 } \
258 \
259 template <typename fnT, typename T1, typename T2> \
260 struct __name__##_ContigFactory \
261 { \
262 fnT get() \
263 { \
264 if constexpr (std::is_same_v< \
265 typename OutputType<T1, T2>::value_type, \
266 void>) { \
267 \
268 fnT fn = nullptr; \
269 return fn; \
270 } \
271 else { \
272 fnT fn = __name__##_contig_impl<T1, T2>; \
273 return fn; \
274 } \
275 } \
276 }; \
277 \
278 template <typename fnT, typename T1, typename T2> \
279 struct __name__##_TypeMapFactory \
280 { \
281 std::enable_if_t<std::is_same<fnT, int>::value, int> get() \
282 { \
283 using rT = typename OutputType<T1, T2>::value_type; \
284 return td_ns::GetTypeid<rT>{}.get(); \
285 } \
286 }; \
287 \
288 template <typename T1, typename T2, typename resT, typename IndexerT> \
289 class __name__##_strided_kernel; \
290 \
291 template <typename argTy1, typename argTy2> \
292 sycl::event __name__##_strided_impl( \
293 sycl::queue &exec_q, size_t nelems, int nd, \
294 const py::ssize_t *shape_and_strides, const char *arg1_p, \
295 py::ssize_t arg1_offset, const char *arg2_p, py::ssize_t arg2_offset, \
296 char *res_p, py::ssize_t res_offset, \
297 const std::vector<sycl::event> &depends, \
298 const std::vector<sycl::event> &additional_depends) \
299 { \
300 return ew_cmn_ns::binary_strided_impl<argTy1, argTy2, OutputType, \
301 StridedFunctor, \
302 __name__##_strided_kernel>( \
303 exec_q, nelems, nd, shape_and_strides, arg1_p, arg1_offset, \
304 arg2_p, arg2_offset, res_p, res_offset, depends, \
305 additional_depends); \
306 } \
307 \
308 template <typename fnT, typename T1, typename T2> \
309 struct __name__##_StridedFactory \
310 { \
311 fnT get() \
312 { \
313 if constexpr (std::is_same_v< \
314 typename OutputType<T1, T2>::value_type, \
315 void>) { \
316 fnT fn = nullptr; \
317 return fn; \
318 } \
319 else { \
320 fnT fn = __name__##_strided_impl<T1, T2>; \
321 return fn; \
322 } \
323 } \
324 }; \
325 \
326 void populate_##__name__##_dispatch_tables(void) \
327 { \
328 ext_ns::init_dispatch_table<binary_contig_impl_fn_ptr_t, \
329 __name__##_ContigFactory>( \
330 __name__##_contig_dispatch_table); \
331 ext_ns::init_dispatch_table<binary_strided_impl_fn_ptr_t, \
332 __name__##_StridedFactory>( \
333 __name__##_strided_dispatch_table); \
334 ext_ns::init_dispatch_table<int, __name__##_TypeMapFactory>( \
335 __name__##_output_typeid_table); \
336 };
337
342#define MACRO_POPULATE_DISPATCH_2OUTS_TABLES(__name__) \
343 template <typename argT1, typename argT2, typename resT1, typename resT2, \
344 unsigned int vec_sz, unsigned int n_vecs> \
345 class __name__##_contig_kernel; \
346 \
347 template <typename argTy1, typename argTy2> \
348 sycl::event __name__##_contig_impl( \
349 sycl::queue &exec_q, size_t nelems, const char *arg1_p, \
350 py::ssize_t arg1_offset, const char *arg2_p, py::ssize_t arg2_offset, \
351 char *res1_p, py::ssize_t res1_offset, char *res2_p, \
352 py::ssize_t res2_offset, const std::vector<sycl::event> &depends = {}) \
353 { \
354 return ew_cmn_ns::binary_two_outputs_contig_impl< \
355 argTy1, argTy2, OutputType, ContigFunctor, \
356 __name__##_contig_kernel>( \
357 exec_q, nelems, arg1_p, arg1_offset, arg2_p, arg2_offset, res1_p, \
358 res1_offset, res2_p, res2_offset, depends); \
359 } \
360 \
361 template <typename fnT, typename T1, typename T2> \
362 struct __name__##_ContigFactory \
363 { \
364 fnT get() \
365 { \
366 if constexpr (std::is_same_v< \
367 typename OutputType<T1, T2>::value_type1, \
368 void> || \
369 std::is_same_v< \
370 typename OutputType<T1, T2>::value_type2, \
371 void>) { \
372 \
373 fnT fn = nullptr; \
374 return fn; \
375 } \
376 else { \
377 fnT fn = __name__##_contig_impl<T1, T2>; \
378 return fn; \
379 } \
380 } \
381 }; \
382 \
383 template <typename fnT, typename T1, typename T2> \
384 struct __name__##_TypeMapFactory \
385 { \
386 std::enable_if_t<std::is_same<fnT, std::pair<int, int>>::value, \
387 std::pair<int, int>> \
388 get() \
389 { \
390 using rT1 = typename OutputType<T1, T2>::value_type1; \
391 using rT2 = typename OutputType<T1, T2>::value_type2; \
392 return std::make_pair(td_ns::GetTypeid<rT1>{}.get(), \
393 td_ns::GetTypeid<rT2>{}.get()); \
394 } \
395 }; \
396 \
397 template <typename T1, typename T2, typename resT1, typename resT2, \
398 typename IndexerT> \
399 class __name__##_strided_kernel; \
400 \
401 template <typename argTy1, typename argTy2> \
402 sycl::event __name__##_strided_impl( \
403 sycl::queue &exec_q, size_t nelems, int nd, \
404 const py::ssize_t *shape_and_strides, const char *arg1_p, \
405 py::ssize_t arg1_offset, const char *arg2_p, py::ssize_t arg2_offset, \
406 char *res1_p, py::ssize_t res1_offset, char *res2_p, \
407 py::ssize_t res2_offset, const std::vector<sycl::event> &depends, \
408 const std::vector<sycl::event> &additional_depends) \
409 { \
410 return ew_cmn_ns::binary_two_outputs_strided_impl< \
411 argTy1, argTy2, OutputType, StridedFunctor, \
412 __name__##_strided_kernel>( \
413 exec_q, nelems, nd, shape_and_strides, arg1_p, arg1_offset, \
414 arg2_p, arg2_offset, res1_p, res1_offset, res2_p, res2_offset, \
415 depends, additional_depends); \
416 } \
417 \
418 template <typename fnT, typename T1, typename T2> \
419 struct __name__##_StridedFactory \
420 { \
421 fnT get() \
422 { \
423 if constexpr (std::is_same_v< \
424 typename OutputType<T1, T2>::value_type1, \
425 void> || \
426 std::is_same_v< \
427 typename OutputType<T1, T2>::value_type2, \
428 void>) { \
429 fnT fn = nullptr; \
430 return fn; \
431 } \
432 else { \
433 fnT fn = __name__##_strided_impl<T1, T2>; \
434 return fn; \
435 } \
436 } \
437 }; \
438 \
439 void populate_##__name__##_dispatch_tables(void) \
440 { \
441 ext_ns::init_dispatch_table<binary_two_outputs_contig_impl_fn_ptr_t, \
442 __name__##_ContigFactory>( \
443 __name__##_contig_dispatch_table); \
444 ext_ns::init_dispatch_table<binary_two_outputs_strided_impl_fn_ptr_t, \
445 __name__##_StridedFactory>( \
446 __name__##_strided_dispatch_table); \
447 ext_ns::init_dispatch_table<std::pair<int, int>, \
448 __name__##_TypeMapFactory>( \
449 __name__##_output_typeid_table); \
450 };