Skip to content

Add scalar support in bitwise functions (partial) #641

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Mar 6, 2021
Merged

Add scalar support in bitwise functions (partial) #641

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
7 changes: 4 additions & 3 deletions dpnp/backend/include/dpnp_gen_2arg_1type_tbl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -46,12 +46,13 @@
/** */ \
/** Function "__name__" executes operator "__operation__" over corresponding elements of input arrays */ \
/** */ \
/** @param[out] result1 Output array. */ \
/** @param[in] array1 Input array 1. */ \
/** @param[in] size1 Number of elements in @ref array1 */ \
/** @param[in] array2 Input array 2. */ \
/** @param[out] result1 Output array. */ \
/** @param[in] size Number of elements in the output array. */ \
/** @param[in] size2 Number of elements in @ref array2 */ \
template <typename _DataType> \
void __name__(void* array1, void* array2, void* result1, size_t size);
void __name__(void* result1, const void* array1, const size_t size1, const void* array2, const size_t size2);

#endif

Expand Down
3 changes: 2 additions & 1 deletion dpnp/backend/include/dpnp_iface.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -666,7 +666,8 @@ INP_DLLEXPORT void dpnp_invert_c(void* array1_in, void* result, size_t size);

#define MACRO_2ARG_1TYPE_OP(__name__, __operation__) \
template <typename _DataType> \
INP_DLLEXPORT void __name__(void* array1_in1, void* array2_in, void* result1, size_t size);
INP_DLLEXPORT void __name__( \
void* result1, const void* array1, const size_t size1, const void* array2, const size_t size2);

#include <dpnp_gen_2arg_1type_tbl.hpp>

Expand Down
18 changes: 12 additions & 6 deletions dpnp/backend/kernels/dpnp_krnl_bitwise.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -70,19 +70,25 @@ static void func_map_init_bitwise_1arg_1type(func_map_t& fmap)
class __name__##_kernel; \
\
template <typename _DataType> \
void __name__(void* array1_in, void* array2_in, void* result1, size_t size) \
void __name__(void* result1, const void* array1_in, const size_t size1, const void* array2_in, const size_t size2) \
{ \
if (!size1 || !size2) \
{ \
return; \
} \
\
cl::sycl::event event; \
_DataType* array1 = reinterpret_cast<_DataType*>(array1_in); \
_DataType* array2 = reinterpret_cast<_DataType*>(array2_in); \
const _DataType* array1 = reinterpret_cast<const _DataType*>(array1_in); \
const _DataType* array2 = reinterpret_cast<const _DataType*>(array2_in); \
_DataType* result = reinterpret_cast<_DataType*>(result1); \
\
cl::sycl::range<1> gws(size); \
const size_t gws_size = std::max(size1, size2); \
cl::sycl::range<1> gws(gws_size); \
auto kernel_parallel_for_func = [=](cl::sycl::id<1> global_id) { \
size_t i = global_id[0]; /*for (size_t i = 0; i < size; ++i)*/ \
{ \
_DataType input_elem1 = array1[i]; \
_DataType input_elem2 = array2[i]; \
const _DataType input_elem1 = (size1 == 1) ? array1[0] : array1[i]; \
const _DataType input_elem2 = (size2 == 1) ? array2[0] : array2[i]; \
result[i] = __operation__; \
} \
}; \
Expand Down
3 changes: 2 additions & 1 deletion dpnp/dpnp_algo/dpnp_algo.pxd
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -211,12 +211,13 @@ cdef extern from "dpnp_iface.hpp":
ctypedef void(*fptr_1out_t)(void *, size_t)
ctypedef void(*fptr_1in_1out_t)(void * , void * , size_t)
ctypedef void(*fptr_2in_1out_t)(void * , void*, void*, size_t)
ctypedef void(*fptr_2in_1out_new_t)(void * , void*, size_t, void*, size_t) # to be fused with fptr_2in_1out_t
ctypedef void(*fptr_blas_gemm_2in_1out_t)(void * , void * , void * , size_t, size_t, size_t)
ctypedef void(*dpnp_reduction_c_t)(void * , const void * , const size_t*, const size_t, const long*, const size_t, const void * , const long*)

cdef dparray call_fptr_1out(DPNPFuncName fptr_name, result_shape, result_dtype)
cdef dparray call_fptr_1in_1out(DPNPFuncName fptr_name, dparray x1, dparray_shape_type result_shape)
cdef dparray call_fptr_2in_1out(DPNPFuncName fptr_name, dparray x1, dparray x2, dparray_shape_type result_shape)
cdef dparray call_fptr_2in_1out(DPNPFuncName fptr_name, dparray x1, dparray x2, dparray_shape_type result_shape, new_version=*)


cpdef dparray dpnp_astype(dparray array1, dtype_target)
Expand Down
11 changes: 8 additions & 3 deletions dpnp/dpnp_algo/dpnp_algo.pyx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -330,7 +330,7 @@ cdef dparray call_fptr_1in_1out(DPNPFuncName fptr_name, dparray x1, dparray_shap
return result


cdef dparray call_fptr_2in_1out(DPNPFuncName fptr_name, dparray x1, dparray x2, dparray_shape_type result_shape):
cdef dparray call_fptr_2in_1out(DPNPFuncName fptr_name, dparray x1, dparray x2, dparray_shape_type result_shape, new_version=False):

""" Convert string type names (dparray.dtype) to C enum DPNPFuncType """
cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(x1.dtype)
Expand All @@ -343,8 +343,13 @@ cdef dparray call_fptr_2in_1out(DPNPFuncName fptr_name, dparray x1, dparray x2,
""" Create result array with type given by FPTR data """
cdef dparray result = dparray(result_shape, dtype=result_type)

cdef fptr_2in_1out_t func = <fptr_2in_1out_t > kernel_data.ptr
""" Call FPTR function """
func(x1.get_data(), x2.get_data(), result.get_data(), x1.size)
# parameter 'new_version' must be removed in shortly
cdef fptr_2in_1out_t func_old = <fptr_2in_1out_t > kernel_data.ptr # can't define it inside 'if' due Cython limitation
cdef fptr_2in_1out_new_t func_new = <fptr_2in_1out_new_t > kernel_data.ptr
if (new_version):
func_new(result.get_data(), x1.get_data(), x1.size, x2.get_data(), x2.size)
else:
func_old(x1.get_data(), x2.get_data(), result.get_data(), x1.size)

return result
10 changes: 5 additions & 5 deletions dpnp/dpnp_algo/dpnp_algo_bitwise.pyx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -47,23 +47,23 @@ __all__ += [


cpdef dparray dpnp_bitwise_and(dparray array1, dparray array2):
return call_fptr_2in_1out(DPNP_FN_BITWISE_AND, array1, array2, array1.shape)
return call_fptr_2in_1out(DPNP_FN_BITWISE_AND, array1, array2, array1.shape, True)


cpdef dparray dpnp_bitwise_or(dparray array1, dparray array2):
return call_fptr_2in_1out(DPNP_FN_BITWISE_OR, array1, array2, array1.shape)
return call_fptr_2in_1out(DPNP_FN_BITWISE_OR, array1, array2, array1.shape, True)


cpdef dparray dpnp_bitwise_xor(dparray array1, dparray array2):
return call_fptr_2in_1out(DPNP_FN_BITWISE_XOR, array1, array2, array1.shape)
return call_fptr_2in_1out(DPNP_FN_BITWISE_XOR, array1, array2, array1.shape, True)


cpdef dparray dpnp_invert(dparray arr):
return call_fptr_1in_1out(DPNP_FN_INVERT, arr, arr.shape)


cpdef dparray dpnp_left_shift(dparray array1, dparray array2):
return call_fptr_2in_1out(DPNP_FN_LEFT_SHIFT, array1, array2, array1.shape)
return call_fptr_2in_1out(DPNP_FN_LEFT_SHIFT, array1, array2, array1.shape, True)

cpdef dparray dpnp_right_shift(dparray array1, dparray array2):
return call_fptr_2in_1out(DPNP_FN_RIGHT_SHIFT, array1, array2, array1.shape)
return call_fptr_2in_1out(DPNP_FN_RIGHT_SHIFT, array1, array2, array1.shape, True)