From 2656d7130472f09357b6ee6dd765da5ed76c6404 Mon Sep 17 00:00:00 2001 From: Denis Smirnov Date: Thu, 25 Mar 2021 14:51:57 +0300 Subject: [PATCH 1/6] Run gtests on GPU --- dpnp/backend/tests/test_main.cpp | 2 ++ dpnp/backend/tests/test_utils_iterator.cpp | 34 ++++++++++++++++++++++ 2 files changed, 36 insertions(+) diff --git a/dpnp/backend/tests/test_main.cpp b/dpnp/backend/tests/test_main.cpp index a28463c31c7c..daa2c4b12679 100644 --- a/dpnp/backend/tests/test_main.cpp +++ b/dpnp/backend/tests/test_main.cpp @@ -57,6 +57,8 @@ int main(int argc, char** argv) // currently using global queue + dpnp_queue_initialize_c(QueueOptions::GPU_SELECTOR); + // It returns 0 if all tests are successful, or 1 otherwise. return RUN_ALL_TESTS(); } diff --git a/dpnp/backend/tests/test_utils_iterator.cpp b/dpnp/backend/tests/test_utils_iterator.cpp index ffc2e0658659..014aa7e749b9 100644 --- a/dpnp/backend/tests/test_utils_iterator.cpp +++ b/dpnp/backend/tests/test_utils_iterator.cpp @@ -429,6 +429,40 @@ TEST_P(IteratorReduction, sycl_reduce_axis) } } +TEST(TestUtilsIterator, sycl_get_first) +{ + using data_type = double; + + const dpnpc_index_t result_size = 1; + vector result(result_size, 42); + data_type* result_ptr = result.data(); + + vector input_data = get_input_data({1}); + data_type* input_ptr = input_data.data(); + DPNPC_id input(input_ptr, {1}); + + ASSERT_EQ(input.get_output_size(), result_size); + + cl::sycl::range<1> gws(result_size); + const DPNPC_id* 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->begin()); + }; + + auto kernel_func = [&](cl::sycl::handler& cgh) { + cgh.parallel_for(gws, kernel_parallel_for_func); + }; + + cl::sycl::event event = DPNP_QUEUE.submit(kernel_func); + event.wait(); + + for (dpnpc_index_t i = 0; i < result_size; ++i) + { + EXPECT_EQ(result.at(i), input_data[0]); + } +} + /** * Expected values produced by following script: * From be5c2b031eee2e67ccf8ce54e3087a39e4a0f53f Mon Sep 17 00:00:00 2001 From: Denis Smirnov Date: Fri, 26 Mar 2021 09:40:16 +0300 Subject: [PATCH 2/6] Fix usage of iterator in SYCL --- dpnp/backend/tests/test_utils_iterator.cpp | 15 ++++++++++----- 1 file changed, 10 insertions(+), 5 deletions(-) diff --git a/dpnp/backend/tests/test_utils_iterator.cpp b/dpnp/backend/tests/test_utils_iterator.cpp index 014aa7e749b9..39a7332cdf4c 100644 --- a/dpnp/backend/tests/test_utils_iterator.cpp +++ b/dpnp/backend/tests/test_utils_iterator.cpp @@ -24,6 +24,7 @@ //***************************************************************************** #include +#include #include #include @@ -439,15 +440,17 @@ TEST(TestUtilsIterator, sycl_get_first) vector input_data = get_input_data({1}); data_type* input_ptr = input_data.data(); - DPNPC_id input(input_ptr, {1}); - ASSERT_EQ(input.get_output_size(), result_size); + // DPNPC_id* input_it = reinterpret_cast*>(dpnp_memory_alloc_c(sizeof(DPNPC_id))); + DPNPC_id* input_it = sycl::malloc_shared>(1, DPNP_QUEUE); + new (input_it) DPNPC_id(input_ptr, {1}); + + ASSERT_EQ(input_it->get_output_size(), result_size); cl::sycl::range<1> gws(result_size); - const DPNPC_id* 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->begin()); + result_ptr[idx] = input_it->get_output_size(); }; auto kernel_func = [&](cl::sycl::handler& cgh) { @@ -459,8 +462,10 @@ TEST(TestUtilsIterator, sycl_get_first) for (dpnpc_index_t i = 0; i < result_size; ++i) { - EXPECT_EQ(result.at(i), input_data[0]); + EXPECT_EQ(result.at(i), result_size); } + + input_it->~DPNPC_id(); } /** From 127875be5de3fecd5bbeabd8beaf198699881804 Mon Sep 17 00:00:00 2001 From: Denis Smirnov Date: Fri, 26 Mar 2021 09:40:16 +0300 Subject: [PATCH 3/6] Fix usage of iterator in SYCL part 2 --- dpnp/backend/tests/test_main.cpp | 2 - dpnp/backend/tests/test_utils_iterator.cpp | 103 ++++++++++++--------- 2 files changed, 58 insertions(+), 47 deletions(-) diff --git a/dpnp/backend/tests/test_main.cpp b/dpnp/backend/tests/test_main.cpp index daa2c4b12679..a28463c31c7c 100644 --- a/dpnp/backend/tests/test_main.cpp +++ b/dpnp/backend/tests/test_main.cpp @@ -57,8 +57,6 @@ int main(int argc, char** argv) // currently using global queue - dpnp_queue_initialize_c(QueueOptions::GPU_SELECTOR); - // It returns 0 if all tests are successful, or 1 otherwise. return RUN_ALL_TESTS(); } diff --git a/dpnp/backend/tests/test_utils_iterator.cpp b/dpnp/backend/tests/test_utils_iterator.cpp index 39a7332cdf4c..d0a2b8c5cab6 100644 --- a/dpnp/backend/tests/test_utils_iterator.cpp +++ b/dpnp/backend/tests/test_utils_iterator.cpp @@ -49,6 +49,15 @@ vector<_DataType> get_input_data(const vector& shape) return input_data; } +template +_DataType* get_shared_data(const vector& shape) +{ + vector<_DataType> input_data = get_input_data<_DataType>(shape); + _DataType* shared_data = reinterpret_cast<_DataType*>(dpnp_memory_alloc_c(input_data.size() * sizeof(_DataType))); + + return shared_data; +} + TEST(TestUtilsIterator, begin_prefix_postfix) { using test_it = dpnpc_it_t; @@ -320,6 +329,43 @@ TEST(TestUtilsIterator, iterator_distance) EXPECT_EQ(axis_1_1_diff_distance, 4); } +TEST(TestUtilsIterator, sycl_getitem) +{ + using data_type = double; + + const dpnpc_index_t result_size = 12; + data_type* result = reinterpret_cast(dpnp_memory_alloc_c(result_size * sizeof(data_type))); + data_type* input_data = get_shared_data({3, 4}); + + DPNPC_id* input_it; + input_it = reinterpret_cast*>(dpnp_memory_alloc_c(sizeof(DPNPC_id))); + new (input_it) DPNPC_id(input_data, {3, 4}); + + ASSERT_EQ(input_it->get_output_size(), result_size); + + cl::sycl::range<1> gws(result_size); + auto kernel_parallel_for_func = [=](cl::sycl::id<1> global_id) { + const size_t idx = global_id[0]; + result[idx] = (*input_it)[idx]; + }; + + auto kernel_func = [&](cl::sycl::handler& cgh) { + cgh.parallel_for(gws, kernel_parallel_for_func); + }; + + cl::sycl::event event = DPNP_QUEUE.submit(kernel_func); + event.wait(); + + for (dpnpc_index_t i = 0; i < result_size; ++i) + { + EXPECT_EQ(result[i], input_data[i]); + } + + input_it->~DPNPC_id(); + dpnp_memory_free_c(input_data); + dpnp_memory_free_c(result); +} + struct IteratorParameters { vector input_shape; @@ -395,17 +441,18 @@ TEST_P(IteratorReduction, sycl_reduce_axis) const IteratorParameters& param = GetParam(); const dpnpc_index_t result_size = param.result.size(); - vector result(result_size, 42); - data_type* result_ptr = result.data(); + data_type* result = reinterpret_cast(dpnp_memory_alloc_c(result_size * sizeof(data_type))); + data_type* input_data = get_shared_data(param.input_shape); - vector input_data = get_input_data(param.input_shape); - DPNPC_id input(input_data.data(), param.input_shape); - input.set_axes(param.axes); + DPNPC_id* input_it; + input_it = reinterpret_cast*>(dpnp_memory_alloc_c(sizeof(DPNPC_id))); + new (input_it) DPNPC_id(input_data, param.input_shape); - ASSERT_EQ(input.get_output_size(), result_size); + input_it->set_axes(param.axes); + + ASSERT_EQ(input_it->get_output_size(), result_size); cl::sycl::range<1> gws(result_size); - const DPNPC_id* input_it = &input; auto kernel_parallel_for_func = [=](cl::sycl::id<1> global_id) { const size_t idx = global_id[0]; @@ -414,7 +461,7 @@ TEST_P(IteratorReduction, sycl_reduce_axis) { accumulator += *data_it; } - result_ptr[idx] = accumulator; + result[idx] = accumulator; }; auto kernel_func = [&](cl::sycl::handler& cgh) { @@ -426,46 +473,12 @@ TEST_P(IteratorReduction, sycl_reduce_axis) for (dpnpc_index_t i = 0; i < result_size; ++i) { - EXPECT_EQ(result.at(i), param.result.at(i)); - } -} - -TEST(TestUtilsIterator, sycl_get_first) -{ - using data_type = double; - - const dpnpc_index_t result_size = 1; - vector result(result_size, 42); - data_type* result_ptr = result.data(); - - vector input_data = get_input_data({1}); - data_type* input_ptr = input_data.data(); - - // DPNPC_id* input_it = reinterpret_cast*>(dpnp_memory_alloc_c(sizeof(DPNPC_id))); - DPNPC_id* input_it = sycl::malloc_shared>(1, DPNP_QUEUE); - new (input_it) DPNPC_id(input_ptr, {1}); - - ASSERT_EQ(input_it->get_output_size(), result_size); - - cl::sycl::range<1> gws(result_size); - auto kernel_parallel_for_func = [=](cl::sycl::id<1> global_id) { - const size_t idx = global_id[0]; - result_ptr[idx] = input_it->get_output_size(); - }; - - auto kernel_func = [&](cl::sycl::handler& cgh) { - cgh.parallel_for(gws, kernel_parallel_for_func); - }; - - cl::sycl::event event = DPNP_QUEUE.submit(kernel_func); - event.wait(); - - for (dpnpc_index_t i = 0; i < result_size; ++i) - { - EXPECT_EQ(result.at(i), result_size); + EXPECT_EQ(result[i], param.result.at(i)); } input_it->~DPNPC_id(); + dpnp_memory_free_c(input_data); + dpnp_memory_free_c(result); } /** From 7f771de66d513a4307e5016f446fb4c87cc25b57 Mon Sep 17 00:00:00 2001 From: Denis Smirnov Date: Tue, 30 Mar 2021 14:43:39 +0300 Subject: [PATCH 4/6] Fix usage of iterator in SYCL part 3 --- dpnp/backend/tests/test_utils_iterator.cpp | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/dpnp/backend/tests/test_utils_iterator.cpp b/dpnp/backend/tests/test_utils_iterator.cpp index d0a2b8c5cab6..12a4bfa19811 100644 --- a/dpnp/backend/tests/test_utils_iterator.cpp +++ b/dpnp/backend/tests/test_utils_iterator.cpp @@ -55,6 +55,11 @@ _DataType* get_shared_data(const vector& shape) vector<_DataType> input_data = get_input_data<_DataType>(shape); _DataType* shared_data = reinterpret_cast<_DataType*>(dpnp_memory_alloc_c(input_data.size() * sizeof(_DataType))); + for (size_t i = 0; i < input_data.size(); ++i) + { + shared_data[i] = input_data[i]; + } + return shared_data; } @@ -341,8 +346,6 @@ TEST(TestUtilsIterator, sycl_getitem) input_it = reinterpret_cast*>(dpnp_memory_alloc_c(sizeof(DPNPC_id))); new (input_it) DPNPC_id(input_data, {3, 4}); - ASSERT_EQ(input_it->get_output_size(), result_size); - cl::sycl::range<1> gws(result_size); auto kernel_parallel_for_func = [=](cl::sycl::id<1> global_id) { const size_t idx = global_id[0]; From 957cd385a2887d07fe7cd2fabe9ddb281ed2246f Mon Sep 17 00:00:00 2001 From: Denis Smirnov Date: Tue, 30 Mar 2021 15:38:04 +0300 Subject: [PATCH 5/6] Fix usage of iterator in SYCL part 4 --- dpnp/backend/src/dpnp_iterator.hpp | 27 +++++++++++++++------- dpnp/backend/tests/test_main.cpp | 2 ++ dpnp/backend/tests/test_utils_iterator.cpp | 25 +++++++++++--------- 3 files changed, 35 insertions(+), 19 deletions(-) diff --git a/dpnp/backend/src/dpnp_iterator.hpp b/dpnp/backend/src/dpnp_iterator.hpp index 6c6d14d4d041..02c39c55ed29 100644 --- a/dpnp/backend/src/dpnp_iterator.hpp +++ b/dpnp/backend/src/dpnp_iterator.hpp @@ -291,13 +291,21 @@ class DPNPC_id final free_iteration_memory(); free_output_memory(); - axes = get_validated_axes(__axes, input_shape_size); + std::vector valid_axes = get_validated_axes(__axes, input_shape_size); axis_use = true; - output_shape_size = input_shape_size - axes.size(); + axes_size = valid_axes.size(); + const size_type axes_size_in_bytes = axes_size * sizeof(size_type); + axes = reinterpret_cast(dpnp_memory_alloc_c(axes_size_in_bytes)); + for (size_type i = 0; i < axes_size; ++i) + { + axes[i] = valid_axes[i]; + } + + output_shape_size = input_shape_size - axes_size; const size_type output_shape_size_in_bytes = output_shape_size * sizeof(size_type); - iteration_shape_size = axes.size(); + iteration_shape_size = axes_size; const size_type iteration_shape_size_in_bytes = iteration_shape_size * sizeof(size_type); std::vector iteration_shape; @@ -305,7 +313,7 @@ class DPNPC_id final size_type* output_shape_it = output_shape; for (size_type i = 0; i < input_shape_size; ++i) { - if (std::find(axes.begin(), axes.end(), i) == axes.end()) + if (std::find(valid_axes.begin(), valid_axes.end(), i) == valid_axes.end()) { *output_shape_it = input_shape[i]; ++output_shape_it; @@ -320,7 +328,7 @@ class DPNPC_id final iteration_size = 1; iteration_shape.reserve(iteration_shape_size); - for (const auto& axis : axes) + for (const auto& axis : valid_axes) { const size_type axis_dim = input_shape[axis]; iteration_shape.push_back(axis_dim); @@ -423,7 +431,7 @@ class DPNPC_id final for (size_t iit = 0, oit = 0; iit < input_shape_size; ++iit) { - if (std::find(axes.begin(), axes.end(), iit) == axes.end()) + if (std::find(axes, axes + axes_size, iit) == axes + axes_size) { input_global_id += (sycl_output_xyz_thread[oit] * input_shape_strides[iit]); ++oit; @@ -442,8 +450,10 @@ class DPNPC_id final void free_axes_memory() { - axes.clear(); + axes_size = size_type{}; + dpnp_memory_free_c(axes); dpnp_memory_free_c(axes_shape_strides); + axes = nullptr; axes_shape_strides = nullptr; } @@ -491,7 +501,8 @@ class DPNPC_id final size_type input_shape_size = size_type{}; /**< input array shape size */ size_type* input_shape_strides = nullptr; /**< input array shape strides (same size as input_shape) */ - std::vector axes; /**< input shape reduction axes */ + size_type* axes = nullptr; /**< input shape reduction axes */ + size_type axes_size = size_type{}; /**< input shape reduction axes size */ bool axis_use = false; size_type output_size = size_type{}; /**< output array size. Expected is same as GWS */ diff --git a/dpnp/backend/tests/test_main.cpp b/dpnp/backend/tests/test_main.cpp index a28463c31c7c..daa2c4b12679 100644 --- a/dpnp/backend/tests/test_main.cpp +++ b/dpnp/backend/tests/test_main.cpp @@ -57,6 +57,8 @@ int main(int argc, char** argv) // currently using global queue + dpnp_queue_initialize_c(QueueOptions::GPU_SELECTOR); + // It returns 0 if all tests are successful, or 1 otherwise. return RUN_ALL_TESTS(); } diff --git a/dpnp/backend/tests/test_utils_iterator.cpp b/dpnp/backend/tests/test_utils_iterator.cpp index 12a4bfa19811..f5d3d98148fe 100644 --- a/dpnp/backend/tests/test_utils_iterator.cpp +++ b/dpnp/backend/tests/test_utils_iterator.cpp @@ -50,11 +50,10 @@ vector<_DataType> get_input_data(const vector& shape) } template -_DataType* get_shared_data(const vector& shape) +_DataType* get_shared_data(const vector<_DataType>& input_data) { - vector<_DataType> input_data = get_input_data<_DataType>(shape); - _DataType* shared_data = reinterpret_cast<_DataType*>(dpnp_memory_alloc_c(input_data.size() * sizeof(_DataType))); - + 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)); for (size_t i = 0; i < input_data.size(); ++i) { shared_data[i] = input_data[i]; @@ -340,11 +339,13 @@ TEST(TestUtilsIterator, sycl_getitem) const dpnpc_index_t result_size = 12; data_type* result = reinterpret_cast(dpnp_memory_alloc_c(result_size * sizeof(data_type))); - data_type* input_data = get_shared_data({3, 4}); + + vector input_data = get_input_data({3, 4}); + data_type* shared_data = get_shared_data(input_data); DPNPC_id* input_it; input_it = reinterpret_cast*>(dpnp_memory_alloc_c(sizeof(DPNPC_id))); - new (input_it) DPNPC_id(input_data, {3, 4}); + new (input_it) DPNPC_id(shared_data, {3, 4}); cl::sycl::range<1> gws(result_size); auto kernel_parallel_for_func = [=](cl::sycl::id<1> global_id) { @@ -361,11 +362,11 @@ TEST(TestUtilsIterator, sycl_getitem) for (dpnpc_index_t i = 0; i < result_size; ++i) { - EXPECT_EQ(result[i], input_data[i]); + EXPECT_EQ(result[i], shared_data[i]); } input_it->~DPNPC_id(); - dpnp_memory_free_c(input_data); + dpnp_memory_free_c(shared_data); dpnp_memory_free_c(result); } @@ -445,11 +446,13 @@ TEST_P(IteratorReduction, sycl_reduce_axis) const IteratorParameters& param = GetParam(); const dpnpc_index_t result_size = param.result.size(); data_type* result = reinterpret_cast(dpnp_memory_alloc_c(result_size * sizeof(data_type))); - data_type* input_data = get_shared_data(param.input_shape); + + vector input_data = get_input_data(param.input_shape); + data_type* shared_data = get_shared_data(input_data); DPNPC_id* input_it; input_it = reinterpret_cast*>(dpnp_memory_alloc_c(sizeof(DPNPC_id))); - new (input_it) DPNPC_id(input_data, param.input_shape); + new (input_it) DPNPC_id(shared_data, param.input_shape); input_it->set_axes(param.axes); @@ -480,7 +483,7 @@ TEST_P(IteratorReduction, sycl_reduce_axis) } input_it->~DPNPC_id(); - dpnp_memory_free_c(input_data); + dpnp_memory_free_c(shared_data); dpnp_memory_free_c(result); } From 9b4825861b651bcf9a7f74b0760f6b83c3f14d6c Mon Sep 17 00:00:00 2001 From: Denis Smirnov Date: Tue, 30 Mar 2021 18:58:36 +0300 Subject: [PATCH 6/6] Avoid forced testing on GPU --- dpnp/backend/tests/test_main.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/dpnp/backend/tests/test_main.cpp b/dpnp/backend/tests/test_main.cpp index daa2c4b12679..a28463c31c7c 100644 --- a/dpnp/backend/tests/test_main.cpp +++ b/dpnp/backend/tests/test_main.cpp @@ -57,8 +57,6 @@ int main(int argc, char** argv) // currently using global queue - dpnp_queue_initialize_c(QueueOptions::GPU_SELECTOR); - // It returns 0 if all tests are successful, or 1 otherwise. return RUN_ALL_TESTS(); }