Skip to content

MULTIPLY enable broadcasting #655

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
shssf merged 10 commits into from
Apr 6, 2021
Merged

MULTIPLY enable broadcasting #655

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
10 commits
Select commit Hold shift + click to select a range
6887fb3
MULTIPLY enable broadcasting
densmirn Mar 9, 2021
b66a4e8
Add gtests for broadcast iterator
densmirn Mar 22, 2021
052a307
Expand memory free before broadcasting
densmirn Mar 25, 2021
6f6f9ed
Disable broadcasting in element-wise operations
densmirn Mar 25, 2021
bc0c1f7
Change docs for 2 DPNPC_id methods
densmirn Mar 25, 2021
257700d
Merge branch 'master' of https://github.com/IntelPython/dpnp into fea…
densmirn Mar 25, 2021
7f2e8e2
Change memory free call before broadcasting
densmirn Mar 25, 2021
0731675
Enable gtests on broadcasting on GPU
densmirn Apr 1, 2021
da181e2
Enable broadcasting in element-wise operations
densmirn Apr 1, 2021
3b30fb6
Merge branch 'master' of https://github.com/IntelPython/dpnp into fea…
densmirn Apr 1, 2021
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
Next Next commit
MULTIPLY enable broadcasting
  • Loading branch information
@densmirn
densmirn committed Mar 19, 2021
commit 6887fb39411f1b431d465b4562a575b40c80ccba
38 changes: 23 additions & 15 deletions dpnp/backend/kernels/dpnp_krnl_elemwise.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,6 +28,7 @@

#include <dpnp_iface.hpp>
#include "dpnp_fptr.hpp"
#include "dpnp_iterator.hpp"
#include "dpnp_utils.hpp"
#include "queue_sycl.hpp"

Expand Down Expand Up @@ -353,29 +354,38 @@ static void func_map_init_elemwise_1arg_1type(func_map_t& fmap)
const size_t* where) \
{ \
/* avoid warning unused variable*/ \
(void)input1_shape; \
(void)input1_shape_ndim; \
(void)input2_shape; \
(void)input2_shape_ndim; \
(void)where; \
\
if (!input1_size || !input2_size) \
{ \
return; \
} \
\
const size_t result_size = (input2_size > input1_size) ? input2_size : input1_size; \
\
const _DataType_input1* input1_data = reinterpret_cast<const _DataType_input1*>(input1_in); \
const _DataType_input2* input2_data = reinterpret_cast<const _DataType_input2*>(input2_in); \
_DataType_input1* input1_data = reinterpret_cast<_DataType_input1*>(const_cast<void*>(input1_in)); \
_DataType_input2* input2_data = reinterpret_cast<_DataType_input2*>(const_cast<void*>(input2_in)); \
_DataType_output* result = reinterpret_cast<_DataType_output*>(result_out); \
\
std::vector<size_t> result_shape = get_common_shape(input1_shape, input1_shape_ndim, \

@shssf shssf Mar 19, 2021

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

maybe "get_result_shape" is better for this function name? not sure...

input2_shape, input2_shape_ndim); \
\
DPNPC_id<_DataType_input1> input1(input1_data, input1_shape, input1_shape_ndim); \
input1.broadcast(result_shape); \

@shssf shssf Mar 19, 2021

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

"broadcast_to_shape" func name?

\
DPNPC_id<_DataType_input2> input2(input2_data, input2_shape, input2_shape_ndim); \
input2.broadcast(result_shape); \
\
const size_t result_size = input1.get_output_size(); \
\
cl::sycl::range<1> gws(result_size); \
const DPNPC_id<_DataType_input1>* input1_it = &input1; \
const DPNPC_id<_DataType_input2>* input2_it = &input2; \
auto kernel_parallel_for_func = [=](cl::sycl::id<1> global_id) { \
size_t i = global_id[0]; /*for (size_t i = 0; i < result_size; ++i)*/ \
const _DataType_output input1_elem = (input1_size == 1) ? input1_data[0] : input1_data[i]; \
const _DataType_output input2_elem = (input2_size == 1) ? input2_data[0] : input2_data[i]; \
result[i] = __operation1__; \
const size_t i = global_id[0]; /*for (size_t i = 0; i < result_size; ++i)*/ \
{ \
const _DataType_output input1_elem = (*input1_it)[i]; \
const _DataType_output input2_elem = (*input2_it)[i]; \
result[i] = __operation1__; \
} \
}; \
auto kernel_func = [&](cl::sycl::handler& cgh) { \
cgh.parallel_for<class __name__##_kernel<_DataType_output, _DataType_input1, \
Expand All @@ -390,9 +400,7 @@ static void func_map_init_elemwise_1arg_1type(func_map_t& fmap)
std::is_same<_DataType_input1, float>::value) && \
std::is_same<_DataType_input2, _DataType_input1>::value) \
{ \
_DataType_input1* input1 = const_cast<_DataType_input1*>(input1_data); \
_DataType_input2* input2 = const_cast<_DataType_input2*>(input2_data); \
event = __operation2__(DPNP_QUEUE, result_size, input1, input2, result); \
event = __operation2__(DPNP_QUEUE, result_size, input1_data, input2_data, result); \
} \
else \
{ \
Expand Down
85 changes: 84 additions & 1 deletion dpnp/backend/src/dpnp_iterator.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -238,6 +238,58 @@ class DPNPC_id final
return output_size;
}

/**
* @ingroup BACKEND_UTILS
* @brief Set output shape for the data object to use in computation.
*
* Set shape of output array to use in computation of input array index by index of output array.
*
* @note this function is designed for non-SYCL environment execution
*
* @param [in] __shape Output shape.
*/
inline void broadcast(const std::vector<size_type>& __shape)
{
if (axis_use || broadcast_use)
{
return;
}

if (broadcastable(input_shape, input_shape_size, __shape))
{
broadcast_use = true;

output_shape_size = __shape.size();
const size_type output_shape_size_in_bytes = output_shape_size * sizeof(size_type);

output_shape = reinterpret_cast<size_type*>(dpnp_memory_alloc_c(output_shape_size_in_bytes));
int in_idx = input_shape_size - 1;
int out_idx = output_shape_size - 1;
for (; out_idx >= 0; --in_idx, --out_idx)
{
output_shape[out_idx] = __shape[out_idx];

// ex: input_shape = {7, 1, 5}, output_shape = {8, 7, 6, 5} => broadcast_axes = {0, 2}
if (in_idx < 0 || input_shape[in_idx] != output_shape[out_idx])
{
broadcast_axes.insert(broadcast_axes.begin(), out_idx);
}
}

output_size = std::accumulate(
output_shape, output_shape + output_shape_size, size_type(1), std::multiplies<size_type>());

output_shape_strides = reinterpret_cast<size_type*>(dpnp_memory_alloc_c(output_shape_size_in_bytes));
get_shape_offsets_inkernel<size_type>(output_shape, output_shape_size, output_shape_strides);

iteration_size = 1;

// make thread private storage for each shape by multiplying memory
sycl_output_xyz =
reinterpret_cast<size_type*>(dpnp_memory_alloc_c(output_size * output_shape_size_in_bytes));
}
}

/**
* @ingroup BACKEND_UTILS
* @brief Set axis for the data object to use in computation.
Expand Down Expand Up @@ -268,7 +320,7 @@ class DPNPC_id final

/**
* @ingroup BACKEND_UTILS
* @brief Set axes for the data object to use in computation.
* @brief Broadcast input data to specified shape.
*
* Set axes of the shape of input array to use in iteration.
* Axes might be negative to indicate axes as reverse iterator
Expand All @@ -285,6 +337,11 @@ class DPNPC_id final
*/
inline void set_axes(const std::vector<long>& __axes)
{
if (axis_use || broadcast_use)
{
return;
}

if (!__axes.empty() && input_shape_size)
{
axes = get_validated_axes(__axes, input_shape_size);
Expand Down Expand Up @@ -364,6 +421,11 @@ class DPNPC_id final
/// this function is designed for SYCL environment execution
inline reference operator[](size_type __n) const
{
if (broadcast_use)
{
return *begin(__n);
}

const iterator it = begin();
return it[__n];
}
Expand Down Expand Up @@ -426,6 +488,24 @@ class DPNPC_id final
}
}
}
else if (broadcast_use)
{
assert(output_global_id < output_size);

// use thread private storage
size_type* sycl_output_xyz_thread = sycl_output_xyz + (output_global_id * output_shape_size);

get_xyz_by_id_inkernel(output_global_id, output_shape_strides, output_shape_size, sycl_output_xyz_thread);

for (size_t oit = 0, iit = 0; oit < output_shape_size; ++oit)
{
if (std::find(broadcast_axes.begin(), broadcast_axes.end(), oit) == broadcast_axes.end())
{
input_global_id += (sycl_output_xyz_thread[oit] * input_shape_strides[iit]);
++iit;
}
}
}

return input_global_id;
}
Expand Down Expand Up @@ -456,6 +536,9 @@ class DPNPC_id final
std::vector<size_type> axes; /**< input shape reduction axes */
bool axis_use = false;

std::vector<size_type> broadcast_axes; /**< input shape broadcast axes */
bool broadcast_use = false;

size_type output_size = size_type{}; /**< output array size. Expected is same as GWS */
size_type* output_shape = nullptr; /**< output array shape */
size_type output_shape_size = size_type{}; /**< output array shape size */
Expand Down
86 changes: 86 additions & 0 deletions dpnp/backend/src/dpnp_utils.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -127,6 +127,92 @@ size_t get_id_by_xyz_inkernel(const _DataType* xyz, size_t xyz_size, const _Data
return global_id;
}

/**
* @ingroup BACKEND_UTILS
* @brief Check input shape is broadcastable to output one.
*
* @param [in] input_shape Input shape.
* @param [in] output_shape Output shape.
*
* @return Input shape is broadcastable to output one or not.
*/
static inline bool
broadcastable(const std::vector<size_t>& input_shape, const std::vector<size_t>& output_shape)
{
if (input_shape.size() > output_shape.size())
{
return false;
}

std::vector<size_t>::const_reverse_iterator irit = input_shape.rbegin();
std::vector<size_t>::const_reverse_iterator orit = output_shape.rbegin();
for (; irit != input_shape.rend(); ++irit, ++orit)
{
if (*irit != 1 && *irit != *orit)
{
return false;
}
}

return true;
}

static inline bool
broadcastable(const size_t* input_shape, const size_t input_shape_size, const std::vector<size_t>& output_shape)
{
const std::vector<size_t> input_shape_vec(input_shape, input_shape + input_shape_size);
return broadcastable(input_shape_vec, output_shape);
}

/**
* @ingroup BACKEND_UTILS
* @brief Get common shape based on input shapes.
*
* Example:
* Input1 shape A[8, 1, 6, 1]
* Input2 shape B[7, 1, 5]
* Output shape will be C[8, 7, 6, 5]
*
* @param [in] input1_shape Input1 shape.
* @param [in] input1_shape_size Input1 shape size.
* @param [in] input2_shape Input2 shape.
* @param [in] input2_shape_size Input2 shape size.
*
* @exception std::domain_error Input shapes are not broadcastable.
* @return Common shape.
*/
static inline std::vector<size_t>
get_common_shape(const size_t* input1_shape, const size_t input1_shape_size,
const size_t* input2_shape, const size_t input2_shape_size)
{
const size_t result_shape_size = (input2_shape_size > input1_shape_size) ? input2_shape_size : input1_shape_size;
std::vector<size_t> result_shape;
result_shape.reserve(result_shape_size);

int in1_idx = input1_shape_size - 1;
int in2_idx = input2_shape_size - 1;
for (; in1_idx >= 0 || in2_idx >= 0; --in1_idx, --in2_idx)
{
size_t input1_val = (in1_idx >= 0) ? input1_shape[in1_idx] : 1;
size_t input2_val = (in2_idx >= 0) ? input2_shape[in2_idx] : 1;

if (input1_val == input2_val || input1_val == 1)
{
result_shape.insert(result_shape.begin(), input2_val);
}
else if (input2_val == 1)
{
result_shape.insert(result_shape.begin(), input1_val);
}
else
{
throw std::domain_error("DPNP Error: get_common_shape() failed with input shapes check");
}
}

return result_shape;
}

/**
* @ingroup BACKEND_UTILS
* @brief Normalizes an axes into a non-negative integer axes.
Expand Down
4 changes: 0 additions & 4 deletions dpnp/dpnp_iface_mathematical.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1071,10 +1071,6 @@ def multiply(x1, x2, dtype=None, out=None, where=True, **kwargs):
pass
elif x2_is_dparray and x2.ndim == 0:
pass
elif x1_is_dparray and x2_is_dparray and x1.size != x2.size:
pass
elif x1_is_dparray and x2_is_dparray and x1.shape != x2.shape:
pass
elif dtype is not None:
pass
elif out is not None:
Expand Down