Skip to content

Commit 863f587

Browse files
antonwolfyantonwolfy
authored
dpnp.add() doesn't work properly with a scalar (#1288)
* dpnp.add() doesn't work properly with a scalar * get rid of dpctl.SyclQueue() call in tests with unsupported device keyword * Add a fix for crash on CPU device * USM type in operations with a scalar * Porting fix for crash to logic kernel
1 parent d1acfaf commit 863f587

17 files changed

+605
-475
lines changed

‎dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp

+79-14
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
//*****************************************************************************
2-
// Copyright (c) 2016-2020, Intel Corporation
2+
// Copyright (c) 2016-2023, Intel Corporation
33
// All rights reserved.
44
//
55
// Redistribution and use in source and binary forms, with or without
@@ -31,7 +31,10 @@
3131
* Parameters:
3232
* - public name of the function and kernel name
3333
* - operation used to calculate the result
34+
* - vector operation over SYCL group used to calculate the result
35+
* - list of types vector operation accepts
3436
* - mkl operation used to calculate the result
37+
* - list of types mkl operation accepts
3538
*
3639
*/
3740

@@ -41,11 +44,12 @@
4144

4245
#ifdef _SECTION_DOCUMENTATION_GENERATION_
4346

44-
#define MACRO_2ARG_3TYPES_OP(__name__, __operation1__, __operation2__) \
47+
#define MACRO_2ARG_3TYPES_OP( \
48+
__name__, __operation__, __vec_operation__, __vec_types__, __mkl_operation__, __mkl_types__) \
4549
/** @ingroup BACKEND_API */ \
4650
/** @brief Per element operation function __name__ */ \
4751
/** */ \
48-
/** Function "__name__" executes operator "__operation1__" over corresponding elements of input arrays */ \
52+
/** Function "__name__" executes operator "__operation__" over corresponding elements of input arrays */ \
4953
/** */ \
5054
/** @param[in] q_ref Reference to SYCL queue. */ \
5155
/** @param[out] result_out Output array. */ \
@@ -105,23 +109,84 @@
105109

106110
#endif
107111

108-
MACRO_2ARG_3TYPES_OP(dpnp_add_c, input1_elem + input2_elem, oneapi::mkl::vm::add)
109-
MACRO_2ARG_3TYPES_OP(dpnp_arctan2_c, sycl::atan2((double)input1_elem, (double)input2_elem), oneapi::mkl::vm::atan2)
112+
MACRO_2ARG_3TYPES_OP(dpnp_add_c,
113+
input1_elem + input2_elem,
114+
sycl::add_sat(x1, x2),
115+
MACRO_UNPACK_TYPES(int, long),
116+
oneapi::mkl::vm::add,
117+
MACRO_UNPACK_TYPES(float, double, std::complex<float>, std::complex<double>))
118+
119+
MACRO_2ARG_3TYPES_OP(dpnp_arctan2_c,
120+
sycl::atan2((double)input1_elem, (double)input2_elem),
121+
nullptr,
122+
std::false_type,
123+
oneapi::mkl::vm::atan2,
124+
MACRO_UNPACK_TYPES(float, double))
125+
110126
MACRO_2ARG_3TYPES_OP(dpnp_copysign_c,
111127
sycl::copysign((double)input1_elem, (double)input2_elem),
112-
oneapi::mkl::vm::copysign)
113-
MACRO_2ARG_3TYPES_OP(dpnp_divide_c, input1_elem / input2_elem, oneapi::mkl::vm::div)
114-
MACRO_2ARG_3TYPES_OP(dpnp_fmod_c, sycl::fmod((double)input1_elem, (double)input2_elem), oneapi::mkl::vm::fmod)
115-
MACRO_2ARG_3TYPES_OP(dpnp_hypot_c, sycl::hypot((double)input1_elem, (double)input2_elem), oneapi::mkl::vm::hypot)
116-
MACRO_2ARG_3TYPES_OP(dpnp_maximum_c, sycl::max(input1_elem, input2_elem), oneapi::mkl::vm::fmax)
117-
MACRO_2ARG_3TYPES_OP(dpnp_minimum_c, sycl::min(input1_elem, input2_elem), oneapi::mkl::vm::fmin)
128+
nullptr,
129+
std::false_type,
130+
oneapi::mkl::vm::copysign,
131+
MACRO_UNPACK_TYPES(float, double))
132+
133+
MACRO_2ARG_3TYPES_OP(dpnp_divide_c,
134+
input1_elem / input2_elem,
135+
nullptr,
136+
std::false_type,
137+
oneapi::mkl::vm::div,
138+
MACRO_UNPACK_TYPES(float, double))
139+
140+
MACRO_2ARG_3TYPES_OP(dpnp_fmod_c,
141+
sycl::fmod((double)input1_elem, (double)input2_elem),
142+
nullptr,
143+
std::false_type,
144+
oneapi::mkl::vm::fmod,
145+
MACRO_UNPACK_TYPES(float, double))
146+
147+
MACRO_2ARG_3TYPES_OP(dpnp_hypot_c,
148+
sycl::hypot((double)input1_elem, (double)input2_elem),
149+
nullptr,
150+
std::false_type,
151+
oneapi::mkl::vm::hypot,
152+
MACRO_UNPACK_TYPES(float, double))
153+
154+
MACRO_2ARG_3TYPES_OP(dpnp_maximum_c,
155+
sycl::max(input1_elem, input2_elem),
156+
nullptr,
157+
std::false_type,
158+
oneapi::mkl::vm::fmax,
159+
MACRO_UNPACK_TYPES(float, double))
160+
161+
MACRO_2ARG_3TYPES_OP(dpnp_minimum_c,
162+
sycl::min(input1_elem, input2_elem),
163+
nullptr,
164+
std::false_type,
165+
oneapi::mkl::vm::fmin,
166+
MACRO_UNPACK_TYPES(float, double))
118167

119168
// "multiply" needs to be standalone kernel (not autogenerated) due to complex algorithm. This is not an element wise.
120169
// pytest "tests/third_party/cupy/creation_tests/test_ranges.py::TestMgrid::test_mgrid3"
121170
// requires multiplication shape1[10] with shape2[10,1] and result expected as shape[10,10]
122-
MACRO_2ARG_3TYPES_OP(dpnp_multiply_c, input1_elem* input2_elem, oneapi::mkl::vm::mul)
171+
MACRO_2ARG_3TYPES_OP(dpnp_multiply_c,
172+
input1_elem* input2_elem,
173+
nullptr,
174+
std::false_type,
175+
oneapi::mkl::vm::mul,
176+
MACRO_UNPACK_TYPES(float, double, std::complex<float>, std::complex<double>))
177+
178+
MACRO_2ARG_3TYPES_OP(dpnp_power_c,
179+
sycl::pow((double)input1_elem, (double)input2_elem),
180+
nullptr,
181+
std::false_type,
182+
oneapi::mkl::vm::pow,
183+
MACRO_UNPACK_TYPES(float, double))
123184

124-
MACRO_2ARG_3TYPES_OP(dpnp_power_c, sycl::pow((double)input1_elem, (double)input2_elem), oneapi::mkl::vm::pow)
125-
MACRO_2ARG_3TYPES_OP(dpnp_subtract_c, input1_elem - input2_elem, oneapi::mkl::vm::sub)
185+
MACRO_2ARG_3TYPES_OP(dpnp_subtract_c,
186+
input1_elem - input2_elem,
187+
nullptr,
188+
std::false_type,
189+
oneapi::mkl::vm::sub,
190+
MACRO_UNPACK_TYPES(float, double))
126191

127192
#undef MACRO_2ARG_3TYPES_OP

‎dpnp/backend/include/dpnp_iface.hpp

+2-1
Original file line numberDiff line numberDiff line change
@@ -1829,7 +1829,8 @@ INP_DLLEXPORT void dpnp_invert_c(void* array1_in, void* result, size_t size);
18291829

18301830
#include <dpnp_gen_2arg_2type_tbl.hpp>
18311831

1832-
#define MACRO_2ARG_3TYPES_OP(__name__, __operation1__, __operation2__) \
1832+
#define MACRO_2ARG_3TYPES_OP( \
1833+
__name__, __operation__, __vec_operation__, __vec_types__, __mkl_operation__, __mkl_types__) \
18331834
template <typename _DataType_output, typename _DataType_input1, typename _DataType_input2> \
18341835
INP_DLLEXPORT DPCTLSyclEventRef __name__(DPCTLSyclQueueRef q_ref, \
18351836
void* result_out, \

‎dpnp/backend/include/dpnp_iface_fptr.hpp

+2-2
Original file line numberDiff line numberDiff line change
@@ -394,13 +394,13 @@ enum class DPNPFuncName : size_t
394394
enum class DPNPFuncType : size_t
395395
{
396396
DPNP_FT_NONE, /**< Very first element of the enumeration */
397+
DPNP_FT_BOOL, /**< analog of numpy.bool_ or bool */
397398
DPNP_FT_INT, /**< analog of numpy.int32 or int */
398399
DPNP_FT_LONG, /**< analog of numpy.int64 or long */
399400
DPNP_FT_FLOAT, /**< analog of numpy.float32 or float */
400401
DPNP_FT_DOUBLE, /**< analog of numpy.float32 or double */
401402
DPNP_FT_CMPLX64, /**< analog of numpy.complex64 or std::complex<float> */
402-
DPNP_FT_CMPLX128, /**< analog of numpy.complex128 or std::complex<double> */
403-
DPNP_FT_BOOL /**< analog of numpy.bool_ or bool */
403+
DPNP_FT_CMPLX128 /**< analog of numpy.complex128 or std::complex<double> */
404404
};
405405

406406
/**

0 commit comments

Comments
 (0)