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
Prev Previous commit
Next Next commit
Enable gtests on broadcasting on GPU
  • Loading branch information
@densmirn
densmirn committed Apr 1, 2021
commit 0731675949c1ee7552619a5f30d595c96c84e84f
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