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 3 commits
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
43 changes: 29 additions & 14 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,28 +354,41 @@ 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_result_shape(input1_shape, input1_shape_ndim, \
input2_shape, input2_shape_ndim); \
\
DPNPC_id<_DataType_input1>* input1_it; \
const size_t input1_it_size_in_bytes = sizeof(DPNPC_id<_DataType_input1>); \
input1_it = reinterpret_cast<DPNPC_id<_DataType_input1>*>(dpnp_memory_alloc_c(input1_it_size_in_bytes)); \
new (input1_it) DPNPC_id<_DataType_input1>(input1_data, input1_shape, input1_shape_ndim); \
\
input1_it->broadcast_to_shape(result_shape); \
\
DPNPC_id<_DataType_input2>* input2_it; \
const size_t input2_it_size_in_bytes = sizeof(DPNPC_id<_DataType_input2>); \
input2_it = reinterpret_cast<DPNPC_id<_DataType_input2>*>(dpnp_memory_alloc_c(input2_it_size_in_bytes)); \
new (input2_it) DPNPC_id<_DataType_input2>(input2_data, input2_shape, input2_shape_ndim); \
\
input2_it->broadcast_to_shape(result_shape); \
\
const size_t result_size = input1_it->get_output_size(); \
\
cl::sycl::range<1> gws(result_size); \
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]; \
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) { \
Expand All @@ -390,9 +404,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 All @@ -405,6 +417,9 @@ static void func_map_init_elemwise_1arg_1type(func_map_t& fmap)
} \
\
event.wait(); \
\
input1_it->~DPNPC_id(); \
input2_it->~DPNPC_id(); \
Comment on lines +421 to +422

@samir-nasibli samir-nasibli Apr 1, 2021

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

No needed here.

@densmirn densmirn Apr 1, 2021

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Needed, we are responsible for destructing placed objects.

@samir-nasibli samir-nasibli Apr 2, 2021 •
edited
Loading

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

These are obvious things that you are saying, but a little bit not about that.

The comment was not about freeing resources in general, it is about that we should avoid such explicitly call the destructor.
I see that current iterator implementation has flaws, updates for to the iterator interface are required for this.
In anyway this is not the current PR problem.

}

#include <dpnp_gen_2arg_3type_tbl.hpp>
Expand Down
27 changes: 21 additions & 6 deletions dpnp/backend/src/dpnp_iterator.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -257,27 +257,32 @@ class DPNPC_id final

if (broadcastable(input_shape, input_shape_size, __shape))
{
broadcast_axes.clear();
free_broadcast_axes_memory();
free_output_memory();

std::vector<size_type> valid_axes;
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));

for (int irit = input_shape_size - 1, orit = output_shape_size - 1; orit >= 0; --irit, --orit)
{
output_shape[orit] = __shape[orit];

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

broadcast_axes_size = valid_axes.size();
const size_type broadcast_axes_size_in_bytes = broadcast_axes_size * sizeof(size_type);
broadcast_axes = reinterpret_cast<size_type*>(dpnp_memory_alloc_c(broadcast_axes_size_in_bytes));
std::copy(valid_axes.begin(), valid_axes.end(), broadcast_axes);

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

Expand Down Expand Up @@ -505,7 +510,8 @@ class DPNPC_id final

for (int irit = input_shape_size - 1, orit = output_shape_size - 1; irit >= 0; --irit, --orit)
{
if (std::find(broadcast_axes.begin(), broadcast_axes.end(), orit) == broadcast_axes.end())
size_type* broadcast_axes_end = broadcast_axes + broadcast_axes_size;
if (std::find(broadcast_axes, broadcast_axes_end, orit) == broadcast_axes_end)
{
input_global_id += (sycl_output_xyz_thread[orit] * input_shape_strides[irit]);
}
Expand All @@ -528,6 +534,13 @@ class DPNPC_id final
axes_shape_strides = nullptr;
}

void free_broadcast_axes_memory()
{
broadcast_axes_size = size_type{};
dpnp_memory_free_c(broadcast_axes);
broadcast_axes = nullptr;
}

void free_input_memory()
{
input_size = size_type{};
Expand Down Expand Up @@ -561,6 +574,7 @@ class DPNPC_id final
void free_memory()
{
free_axes_memory();
free_broadcast_axes_memory();
free_input_memory();
free_iteration_memory();
free_output_memory();
Expand All @@ -575,7 +589,8 @@ 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 */
size_type* broadcast_axes = nullptr; /**< input shape broadcast axes */
size_type broadcast_axes_size = size_type{}; /**< input shape broadcast axes size */
bool broadcast_use = false;

size_type output_size = size_type{}; /**< output array size. Expected is same as GWS */
Expand Down
30 changes: 30 additions & 0 deletions dpnp/backend/tests/dpnp_test_utils.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,30 @@
#include <iostream>
#include <vector>

#include "dpnp_iterator.hpp"

using namespace std;
using dpnpc_it_t = DPNPC_id<size_t>::iterator;
using dpnpc_value_t = dpnpc_it_t::value_type;
using dpnpc_index_t = dpnpc_it_t::size_type;

template <typename _DataType>
vector<_DataType> get_input_data(const vector<dpnpc_index_t>& shape)
{
const dpnpc_index_t size = accumulate(shape.begin(), shape.end(), dpnpc_index_t(1), multiplies<dpnpc_index_t>());

vector<_DataType> input_data(size);
iota(input_data.begin(), input_data.end(), 1); // let's start from 1 to avoid cleaned memory comparison

return input_data;
}

template <typename _DataType>
_DataType* get_shared_data(const vector<_DataType>& input_data)
{
const size_t data_size_in_bytes = input_data.size() * sizeof(_DataType);
_DataType* shared_data = reinterpret_cast<_DataType*>(dpnp_memory_alloc_c(data_size_in_bytes));
copy(input_data.begin(), input_data.end(), shared_data);

return shared_data;
}
97 changes: 24 additions & 73 deletions dpnp/backend/tests/test_broadcast_iterator.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,64 +28,11 @@
#include <vector>

#include "dpnp_iterator.hpp"
#include "dpnp_test_utils.hpp"

#define DPNP_LOCAL_QUEUE 1 // TODO need to fix build procedure and remove this workaround. Issue #551

@shssf shssf Mar 26, 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.

@samir-nasibli it looks like it requires your attention

#include "queue_sycl.hpp"

using namespace std;
using dpnpc_it_t = DPNPC_id<size_t>::iterator;
using dpnpc_value_t = dpnpc_it_t::value_type;
using dpnpc_index_t = dpnpc_it_t::size_type;

template <typename _DataType>
vector<_DataType> get_input_data(const vector<dpnpc_index_t>& shape)
{
const dpnpc_index_t size = accumulate(shape.begin(), shape.end(), dpnpc_index_t(1), multiplies<dpnpc_index_t>());

vector<_DataType> input_data(size);
iota(input_data.begin(), input_data.end(), 1); // let's start from 1 to avoid cleaned memory comparison

return input_data;
}

TEST(TestBroadcastIterator, take_value_broadcast_loop_2D)
{
vector<dpnpc_value_t> expected_data{1, 1, 1, 1, 2, 2, 2, 2, 3, 3, 3, 3};
// expected input data 1, 2, 3
vector<dpnpc_value_t> input_data = get_input_data<dpnpc_value_t>({3, 1});
DPNPC_id<dpnpc_value_t> result_obj(input_data.data(), {3, 1});
result_obj.broadcast_to_shape({3, 4});

for (dpnpc_index_t output_id = 0; output_id < result_obj.get_output_size(); ++output_id)
{
EXPECT_EQ(result_obj[output_id], expected_data[output_id]);
}
}

TEST(TestBroadcastIterator, take_value_broadcast_loop_3D)
{
vector<dpnpc_value_t> expected_data{1, 2, 3, 4, 5, 6, 1, 2, 3, 4, 5, 6};
// expected input data 1, 2, 3, 4, 5, 6
vector<dpnpc_value_t> input_data = get_input_data<dpnpc_value_t>({2, 3});
DPNPC_id<dpnpc_value_t> result_obj(input_data.data(), {2, 3});
result_obj.broadcast_to_shape({2, 2, 3});

for (dpnpc_index_t output_id = 0; output_id < result_obj.get_output_size(); ++output_id)
{
EXPECT_EQ(result_obj[output_id], expected_data[output_id]);
}
}

TEST(TestBroadcastIterator, output_size_broadcast)
{
vector<dpnpc_value_t> input_data = get_input_data<dpnpc_value_t>({3, 4});
DPNPC_id<dpnpc_value_t> result_obj(input_data.data(), {3, 4});
result_obj.broadcast_to_shape({2, 3, 4});

const dpnpc_index_t output_size = result_obj.get_output_size();
EXPECT_EQ(output_size, 24);
}

struct IteratorParameters
{
vector<dpnpc_it_t::size_type> input_shape;
Expand All @@ -108,22 +55,19 @@ class IteratorBroadcasting : public ::testing::TestWithParam<IteratorParameters>

TEST_P(IteratorBroadcasting, loop_broadcast)
{
using data_type = double;

const IteratorParameters& param = GetParam();
const dpnpc_index_t result_size = param.result.size();
std::vector<data_type> input_data = get_input_data<data_type>(param.input_shape);

vector<dpnpc_value_t> input_data = get_input_data<dpnpc_value_t>(param.input_shape);
DPNPC_id<dpnpc_value_t> input(input_data.data(), param.input_shape);
DPNPC_id<data_type> input(input_data.data(), param.input_shape);
input.broadcast_to_shape(param.output_shape);

ASSERT_EQ(input.get_output_size(), result_size);
ASSERT_EQ(input.get_output_size(), param.result.size());

vector<dpnpc_value_t> test_result(result_size, 42);
for (dpnpc_index_t output_id = 0; output_id < result_size; ++output_id)
for (dpnpc_index_t output_id = 0; output_id < input.get_output_size(); ++output_id)
{
test_result[output_id] = input[output_id];
std::cout << "test_result[" << output_id << "] = " << test_result[output_id] << std::endl;
std::cout << "param.result.at(" << output_id << ") = " << param.result.at(output_id) << std::endl;
EXPECT_EQ(test_result[output_id], param.result.at(output_id));
EXPECT_EQ(input[output_id], param.result.at(output_id));
}
}

Expand All @@ -133,20 +77,23 @@ TEST_P(IteratorBroadcasting, sycl_broadcast)

const IteratorParameters& param = GetParam();
const dpnpc_index_t result_size = param.result.size();
vector<data_type> result(result_size, 42);
data_type* result_ptr = result.data();
data_type* result = reinterpret_cast<data_type*>(dpnp_memory_alloc_c(result_size * sizeof(data_type)));

vector<data_type> input_data = get_input_data<data_type>(param.input_shape);
DPNPC_id<data_type> input(input_data.data(), param.input_shape);
input.broadcast_to_shape(param.output_shape);
std::vector<data_type> input_data = get_input_data<data_type>(param.input_shape);
data_type* shared_data = get_shared_data<data_type>(input_data);

ASSERT_EQ(input.get_output_size(), result_size);
DPNPC_id<data_type>* input_it;
input_it = reinterpret_cast<DPNPC_id<data_type>*>(dpnp_memory_alloc_c(sizeof(DPNPC_id<data_type>)));
new (input_it) DPNPC_id<data_type>(shared_data, param.input_shape);

input_it->broadcast_to_shape(param.output_shape);

ASSERT_EQ(input_it->get_output_size(), result_size);

cl::sycl::range<1> gws(result_size);
const DPNPC_id<data_type>* input_it = &input;
auto kernel_parallel_for_func = [=](cl::sycl::id<1> global_id) {
const size_t idx = global_id[0];
result_ptr[idx] = (*input_it)[idx];
result[idx] = (*input_it)[idx];
};

auto kernel_func = [&](cl::sycl::handler& cgh) {
Expand All @@ -158,8 +105,12 @@ TEST_P(IteratorBroadcasting, sycl_broadcast)

for (dpnpc_index_t i = 0; i < result_size; ++i)
{
EXPECT_EQ(result.at(i), param.result.at(i));
EXPECT_EQ(result[i], param.result.at(i));
}

input_it->~DPNPC_id();

@samir-nasibli samir-nasibli Apr 1, 2021

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

no need to explicitly call the destructor.

@densmirn densmirn Apr 1, 2021

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Needed for placement new.

dpnp_memory_free_c(shared_data);
dpnp_memory_free_c(result);
}

/**
Expand Down
15 changes: 1 addition & 14 deletions dpnp/backend/tests/test_utils_iterator.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,25 +28,12 @@
#include <vector>

#include "dpnp_iterator.hpp"
#include "dpnp_test_utils.hpp"

#define DPNP_LOCAL_QUEUE 1 // TODO need to fix build procedure and remove this workaround. Issue #551
#include "queue_sycl.hpp"

using namespace std;
using dpnpc_it_t = DPNPC_id<size_t>::iterator;
using dpnpc_value_t = dpnpc_it_t::value_type;
using dpnpc_index_t = dpnpc_it_t::size_type;

template <typename _DataType>
vector<_DataType> get_input_data(const vector<dpnpc_index_t>& shape)
{
const dpnpc_index_t size = accumulate(shape.begin(), shape.end(), dpnpc_index_t(1), multiplies<dpnpc_index_t>());

vector<_DataType> input_data(size);
iota(input_data.begin(), input_data.end(), 1); // let's start from 1 to avoid cleaned memory comparison

return input_data;
}

TEST(TestUtilsIterator, begin_prefix_postfix)
{
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