diff --git a/CHANGELOG.md b/CHANGELOG.md index f370cb27b673..060e57825c50 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -14,6 +14,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 * Adjusted the `pre-commit` configuration to run autoupdate weekly [#2479](https://github.com/IntelPython/dpnp/pull/2479) * Improved validation of `--target-hip` build option to only accept a gfx-prefixed value [#2481](https://github.com/IntelPython/dpnp/pull/2481) +* Cleaned up backend code to remove obsolete and unused parts of functionality [#2485](https://github.com/IntelPython/dpnp/pull/2485) ### Fixed diff --git a/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp b/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp index 793b66341ecf..3e0d0bcfee4d 100644 --- a/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp @@ -30,7 +30,6 @@ #include #include "dpnp_fptr.hpp" -#include "dpnp_iterator.hpp" #include "dpnp_utils.hpp" #include "dpnpc_memory_adapter.hpp" #include "queue_sycl.hpp" diff --git a/dpnp/backend/src/dpnp_iterator.hpp b/dpnp/backend/src/dpnp_iterator.hpp deleted file mode 100644 index e6008a97c052..000000000000 --- a/dpnp/backend/src/dpnp_iterator.hpp +++ /dev/null @@ -1,707 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016-2025, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -#pragma once -#ifndef DPNP_ITERATOR_H // Cython compatibility -#define DPNP_ITERATOR_H - -#include -#include -#include -#include -#include -#include - -#include - -/** - * @ingroup BACKEND_UTILS - * @brief Iterator for @ref DPNPC_id type - * - * This type should be used to simplify data iteration over input with - * parameters - * "[axis|axes]" It is designed to be used in SYCL environment - * - */ -template -class DPNP_USM_iterator final -{ -public: - using value_type = _Tp; - using difference_type = std::ptrdiff_t; - using iterator_category = std::random_access_iterator_tag; - using pointer = value_type *; - using reference = value_type &; - using size_type = shape_elem_type; - - DPNP_USM_iterator(pointer __base_ptr, - size_type __id, - const size_type *__shape_stride = nullptr, - const size_type *__axes_stride = nullptr, - size_type __shape_size = 0) - : base(__base_ptr), iter_id(__id), iteration_shape_size(__shape_size), - iteration_shape_strides(__shape_stride), - axes_shape_strides(__axes_stride) - { - } - - DPNP_USM_iterator() = delete; - - inline reference operator*() const - { - return *ptr(); - } - - inline pointer operator->() const - { - return ptr(); - } - - /// prefix increment - inline DPNP_USM_iterator &operator++() - { - ++iter_id; - - return *this; - } - - /// postfix increment - inline DPNP_USM_iterator operator++(int) - { - DPNP_USM_iterator tmp = *this; - ++(*this); // call prefix increment - - return tmp; - } - - inline bool operator==(const DPNP_USM_iterator &__rhs) const - { - assert(base == __rhs.base); // iterators are incomparable if base - // pointers are different - return (iter_id == __rhs.iter_id); - }; - - inline bool operator!=(const DPNP_USM_iterator &__rhs) const - { - return !(*this == __rhs); - }; - - inline bool operator<(const DPNP_USM_iterator &__rhs) const - { - return iter_id < __rhs.iter_id; - }; - - // TODO need more operators - - // Random access iterator requirements - inline reference operator[](size_type __n) const - { - return *ptr(__n); - } - - inline difference_type operator-(const DPNP_USM_iterator &__rhs) const - { - difference_type diff = - difference_type(iter_id) - difference_type(__rhs.iter_id); - - return diff; - } - - /// Print this container in human readable form in error reporting - friend std::ostream &operator<<(std::ostream &__out, - const DPNP_USM_iterator &__it) - { - const std::vector it_strides(__it.iteration_shape_strides, - __it.iteration_shape_strides + - __it.iteration_shape_size); - const std::vector it_axes_strides( - __it.axes_shape_strides, - __it.axes_shape_strides + __it.iteration_shape_size); - - __out << "DPNP_USM_iterator(base=" << __it.base; - __out << ", iter_id=" << __it.iter_id; - __out << ", iteration_shape_size=" << __it.iteration_shape_size; - __out << ", iteration_shape_strides=" << it_strides; - __out << ", axes_shape_strides=" << it_axes_strides; - __out << ")"; - - return __out; - } - -private: - inline pointer ptr() const - { - return ptr(iter_id); - } - - inline pointer ptr(size_type iteration_id) const - { - size_type offset = 0; - - if (iteration_shape_size > 0) { - long reminder = iteration_id; - for (size_t it = 0; it < static_cast(iteration_shape_size); - ++it) { - const size_type axis_val = iteration_shape_strides[it]; - size_type xyz_id = reminder / axis_val; - offset += (xyz_id * axes_shape_strides[it]); - - reminder = reminder % axis_val; - } - } - else { - offset = iteration_id; - } - - return base + offset; - } - - const pointer base = nullptr; - size_type iter_id = - size_type{}; /**< Iterator logical ID over iteration shape */ - const size_type iteration_shape_size = - size_type{}; /**< Number of elements in @ref iteration_shape_strides - array */ - const size_type *iteration_shape_strides = nullptr; - const size_type *axes_shape_strides = nullptr; -}; - -/** - * @ingroup BACKEND_UTILS - * @brief Type to keep USM array pointers used in kernels - * - * This type should be used in host part of the code to provide pre-calculated - * data. The @ref DPNP_USM_iterator will be used later in SYCL environment - * - */ -template -class DPNPC_id final -{ -public: - using value_type = _Tp; - using iterator = DPNP_USM_iterator; - using pointer = value_type *; - using reference = value_type &; - using size_type = shape_elem_type; - - DPNPC_id(DPCTLSyclQueueRef q_ref, - pointer __ptr, - const size_type *__shape, - const size_type __shape_size) - { - queue_ref = q_ref; - std::vector shape(__shape, __shape + __shape_size); - init_container(__ptr, shape); - } - - DPNPC_id(DPCTLSyclQueueRef q_ref, - pointer __ptr, - const size_type *__shape, - const size_type *__strides, - const size_type __ndim) - { - queue_ref = q_ref; - std::vector shape(__shape, __shape + __ndim); - std::vector strides(__strides, __strides + __ndim); - init_container(__ptr, shape, strides); - } - - /** - * @ingroup BACKEND_UTILS - * @brief Main container for reduction iterator - * - * Construct object to hold @ref __ptr data with shape @ref __shape. - * It is needed to provide reduction iterator over the data. - * - * @note this function is designed for non-SYCL environment execution - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] __ptr Pointer to input data. Used to get values only. - * @param [in] __shape Shape of data provided by @ref __ptr. - * Empty container means scalar value pointed by @ref - * __ptr. - */ - DPNPC_id(DPCTLSyclQueueRef q_ref, - pointer __ptr, - const std::vector &__shape) - { - queue_ref = q_ref; - init_container(__ptr, __shape); - } - - /** - * @ingroup BACKEND_UTILS - * @brief Main container for reduction/broadcasting iterator - * - * Construct object to hold @ref __ptr data with shape @ref __shape and - * strides @ref __strides. - * - * @note this function is designed for non-SYCL environment execution - * - * @param [in] __ptr Pointer to input data. Used to get values only. - * @param [in] __shape Shape of data provided by @ref __ptr. - * Empty container means scalar value pointed by - * @ref __ptr. - * @param [in] __strides Strides of data provided by @ref __ptr. - */ - DPNPC_id(pointer __ptr, - const std::vector &__shape, - const std::vector &__strides) - { - init_container(__ptr, __shape, __strides); - } - - DPNPC_id() = delete; - - ~DPNPC_id() - { - free_memory(); - } - - /// this function return number of elements in output - inline size_type get_output_size() const - { - return output_size; - } - - inline void broadcast_to_shape(const size_type *__shape, - const size_type __shape_size) - { - std::vector shape(__shape, __shape + __shape_size); - broadcast_to_shape(shape); - } - - /** - * @ingroup BACKEND_UTILS - * @brief Broadcast input data to specified shape. - * - * Set output shape to use in computation of input index by output index. - * - * @note this function is designed for non-SYCL environment execution - * - * @param [in] __shape Output shape. - */ - inline void broadcast_to_shape(const std::vector &__shape) - { - if (axis_use) { - return; - } - - if (broadcastable(input_shape, input_shape_size, __shape)) { - free_broadcast_axes_memory(); - free_output_memory(); - - std::vector 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( - dpnp_memory_alloc_c(queue_ref, 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} => - // valid_axes = {0, 2} - if (irit < 0 || input_shape[irit] != output_shape[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( - dpnp_memory_alloc_c(queue_ref, 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()); - - output_shape_strides = reinterpret_cast( - dpnp_memory_alloc_c(queue_ref, output_shape_size_in_bytes)); - get_shape_offsets_inkernel( - output_shape, output_shape_size, output_shape_strides); - - iteration_size = 1; - } - } - - /** - * @ingroup BACKEND_UTILS - * @brief Set axis for the data object to use in computation. - * - * Set axis of the shape of input array to use in iteration. - * Axis might be negative to indicate right to left axes indexing in a shape - * - * Indexing goes from left to right. - * Reduction example: - * Input shape A[6, 7, 8, 9] - * set_axis(1) // same as -3 in this example - * output shape will be C[6, 8, 9] - * - * @note this function is designed for non-SYCL environment execution - * - * @param [in] __axis Axis in a shape of input array. - */ - inline void set_axis(shape_elem_type __axis) - { - set_axes({__axis}); - } - - inline void set_axes(const shape_elem_type *__axes, const size_t axes_ndim) - { - const std::vector axes_vec(__axes, __axes + axes_ndim); - set_axes(axes_vec); - } - - /** - * @ingroup BACKEND_UTILS - * @brief Set axes for the data object to use in computation. - * - * Set axes of the shape of input array to use in iteration. - * Axes might be negative to indicate axes as reverse iterator - * - * Indexing goes from left to right. - * Reduction example: - * Input shape A[2, 3, 4, 5] - * set_axes({1, 2}) // same as {-3, -2} - * output shape will be C[2, 5] - * - * @note this function is designed for non-SYCL environment execution - * - * @param [in] __axes Vector of axes of a shape of input array. - */ - inline void set_axes(const std::vector &__axes) - { - if (broadcast_use) { - return; - } - - if (!__axes.empty() && input_shape_size) { - free_axes_memory(); - free_iteration_memory(); - free_output_memory(); - - axes = get_validated_axes(__axes, input_shape_size); - axis_use = true; - - 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(); - const size_type iteration_shape_size_in_bytes = - iteration_shape_size * sizeof(size_type); - std::vector iteration_shape; - - output_shape = reinterpret_cast( - dpnp_memory_alloc_c(queue_ref, output_shape_size_in_bytes)); - 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()) { - *output_shape_it = input_shape[i]; - ++output_shape_it; - } - } - - output_size = - std::accumulate(output_shape, output_shape + output_shape_size, - size_type(1), std::multiplies()); - - output_shape_strides = reinterpret_cast( - dpnp_memory_alloc_c(queue_ref, output_shape_size_in_bytes)); - get_shape_offsets_inkernel( - output_shape, output_shape_size, output_shape_strides); - - iteration_size = 1; - iteration_shape.reserve(iteration_shape_size); - for (const auto &axis : axes) { - const size_type axis_dim = input_shape[axis]; - iteration_shape.push_back(axis_dim); - iteration_size *= axis_dim; - } - - iteration_shape_strides = reinterpret_cast( - dpnp_memory_alloc_c(queue_ref, iteration_shape_size_in_bytes)); - get_shape_offsets_inkernel(iteration_shape.data(), - iteration_shape.size(), - iteration_shape_strides); - - axes_shape_strides = reinterpret_cast( - dpnp_memory_alloc_c(queue_ref, iteration_shape_size_in_bytes)); - for (size_t i = 0; i < static_cast(iteration_shape_size); - ++i) { - axes_shape_strides[i] = input_shape_strides[axes[i]]; - } - } - } - - /// this function is designed for SYCL environment execution - inline iterator begin(size_type output_global_id = 0) const - { - return iterator(data + get_input_begin_offset(output_global_id), 0, - iteration_shape_strides, axes_shape_strides, - iteration_shape_size); - } - - /// this function is designed for SYCL environment execution - inline iterator end(size_type output_global_id = 0) const - { - // TODO it is better to get begin() iterator as a parameter - - return iterator(data + get_input_begin_offset(output_global_id), - get_iteration_size(), iteration_shape_strides, - axes_shape_strides, iteration_shape_size); - } - - /// 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]; - } - -private: - void init_container(pointer __ptr, const std::vector &__shape) - { - // TODO needs to address negative values in __shape with exception - if ((__ptr == nullptr) && __shape.empty()) { - return; - } - - if (__ptr != nullptr) { - data = __ptr; - input_size = 1; // means scalar at this stage - output_size = 1; // if input size is not zero it means we have - // scalar as output - iteration_size = 1; - } - - if (!__shape.empty()) { - input_size = - std::accumulate(__shape.begin(), __shape.end(), size_type(1), - std::multiplies()); - if (input_size == 0) - { // shape might be shape[3, 4, 0, 6]. This means no input memory - // and no output expected - output_size = 0; // depends on axes. zero at this stage only - } - - input_shape_size = __shape.size(); - input_shape = reinterpret_cast(dpnp_memory_alloc_c( - queue_ref, input_shape_size * sizeof(size_type))); - std::copy(__shape.begin(), __shape.end(), input_shape); - - input_shape_strides = - reinterpret_cast(dpnp_memory_alloc_c( - queue_ref, input_shape_size * sizeof(size_type))); - get_shape_offsets_inkernel(input_shape, input_shape_size, - input_shape_strides); - } - iteration_size = input_size; - } - - void init_container(pointer __ptr, - const std::vector &__shape, - const std::vector &__strides) - { - // TODO needs to address negative values in __shape with exception - if ((__ptr == nullptr) && __shape.empty()) { - return; - } - - if (__ptr != nullptr) { - data = __ptr; - input_size = 1; // means scalar at this stage - output_size = 1; // if input size is not zero it means we have - // scalar as output - iteration_size = 1; - } - - if (!__shape.empty()) { - input_size = - std::accumulate(__shape.begin(), __shape.end(), size_type(1), - std::multiplies()); - if (input_size == 0) - { // shape might be shape[3, 4, 0, 6]. This means no input memory - // and no output expected - output_size = 0; // depends on axes. zero at this stage only - } - - input_shape_size = __shape.size(); - input_shape = reinterpret_cast(dpnp_memory_alloc_c( - queue_ref, input_shape_size * sizeof(size_type))); - std::copy(__shape.begin(), __shape.end(), input_shape); - - input_shape_strides = - reinterpret_cast(dpnp_memory_alloc_c( - queue_ref, input_shape_size * sizeof(size_type))); - std::copy(__strides.begin(), __strides.end(), input_shape_strides); - } - iteration_size = input_size; - } - - /// this function is designed for SYCL environment execution - size_type get_input_begin_offset(size_type output_global_id) const - { - size_type input_global_id = 0; - if (axis_use) { - assert(output_global_id < output_size); - - for (size_t iit = 0, oit = 0; - iit < static_cast(input_shape_size); ++iit) - { - if (std::find(axes.begin(), axes.end(), iit) == axes.end()) { - const size_type output_xyz_id = get_xyz_id_by_id_inkernel( - output_global_id, output_shape_strides, - output_shape_size, oit); - input_global_id += - (output_xyz_id * input_shape_strides[iit]); - ++oit; - } - } - } - else if (broadcast_use) { - assert(output_global_id < output_size); - assert(input_shape_size <= output_shape_size); - - for (int irit = input_shape_size - 1, orit = output_shape_size - 1; - irit >= 0; --irit, --orit) - { - size_type *broadcast_axes_end = - broadcast_axes + broadcast_axes_size; - if (std::find(broadcast_axes, broadcast_axes_end, orit) == - broadcast_axes_end) { - const size_type output_xyz_id = get_xyz_id_by_id_inkernel( - output_global_id, output_shape_strides, - output_shape_size, orit); - input_global_id += - (output_xyz_id * input_shape_strides[irit]); - } - } - } - - return input_global_id; - } - - /// this function is designed for SYCL environment execution - size_type get_iteration_size() const - { - return iteration_size; - } - - void free_axes_memory() - { - axes.clear(); - dpnp_memory_free_c(queue_ref, axes_shape_strides); - axes_shape_strides = nullptr; - } - - void free_broadcast_axes_memory() - { - broadcast_axes_size = size_type{}; - dpnp_memory_free_c(queue_ref, broadcast_axes); - broadcast_axes = nullptr; - } - - void free_input_memory() - { - input_size = size_type{}; - input_shape_size = size_type{}; - dpnp_memory_free_c(queue_ref, input_shape); - dpnp_memory_free_c(queue_ref, input_shape_strides); - input_shape = nullptr; - input_shape_strides = nullptr; - } - - void free_iteration_memory() - { - iteration_size = size_type{}; - iteration_shape_size = size_type{}; - dpnp_memory_free_c(queue_ref, iteration_shape_strides); - iteration_shape_strides = nullptr; - } - - void free_output_memory() - { - output_size = size_type{}; - output_shape_size = size_type{}; - dpnp_memory_free_c(queue_ref, output_shape); - dpnp_memory_free_c(queue_ref, output_shape_strides); - output_shape = nullptr; - output_shape_strides = nullptr; - } - - void free_memory() - { - free_axes_memory(); - free_broadcast_axes_memory(); - free_input_memory(); - free_iteration_memory(); - free_output_memory(); - } - - DPCTLSyclQueueRef queue_ref = nullptr; /**< reference to SYCL queue */ - - pointer data = nullptr; /**< input array begin pointer */ - size_type input_size = size_type{}; /**< input array size */ - size_type *input_shape = nullptr; /**< input array shape */ - 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 */ - bool axis_use = false; - - 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 */ - size_type *output_shape = nullptr; /**< output array shape */ - size_type output_shape_size = size_type{}; /**< output array shape size */ - size_type *output_shape_strides = - nullptr; /**< output array shape strides (same size as output_shape) */ - - size_type iteration_size = - size_type{}; /**< iteration array size in elements */ - size_type iteration_shape_size = size_type{}; - size_type *iteration_shape_strides = nullptr; - size_type *axes_shape_strides = nullptr; -}; - -#endif // DPNP_ITERATOR_H diff --git a/dpnp/backend/src/dpnp_utils.hpp b/dpnp/backend/src/dpnp_utils.hpp index cec94a46ded2..9c27b98d78d6 100644 --- a/dpnp/backend/src/dpnp_utils.hpp +++ b/dpnp/backend/src/dpnp_utils.hpp @@ -38,14 +38,6 @@ #include -#define LIBSYCL_VERSION_GREATER(major, minor, patch) \ - (__LIBSYCL_MAJOR_VERSION > major) || \ - (__LIBSYCL_MAJOR_VERSION == major and \ - __LIBSYCL_MINOR_VERSION > minor) || \ - (__LIBSYCL_MAJOR_VERSION == major and \ - __LIBSYCL_MINOR_VERSION == minor and \ - __LIBSYCL_PATCH_VERSION >= patch) - /** * Version of SYCL DPC++ 2023 compiler where a return type of sycl::abs() is * changed from unsigned integer to signed one of input vector. @@ -144,47 +136,6 @@ _DataType get_xyz_id_by_id_inkernel(size_t global_id, return xyz_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 &input_shape, - const std::vector &output_shape) -{ - if (input_shape.size() > output_shape.size()) { - return false; - } - - std::vector::const_reverse_iterator irit = - input_shape.rbegin(); - std::vector::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 shape_elem_type *input_shape, - const size_t input_shape_size, - const std::vector &output_shape) -{ - const std::vector input_shape_vec( - input_shape, input_shape + input_shape_size); - return broadcastable(input_shape_vec, output_shape); -} - /** * @ingroup BACKEND_UTILS * @brief Check arrays are equal. @@ -212,78 +163,6 @@ static inline bool array_equal(const _DataType *input1, std::begin(input2_vec)); } -/** - * @ingroup BACKEND_UTILS - * @brief Normalizes an axes into a non-negative integer axes. - * - * Return vector of normalized axes with a non-negative integer axes. - * - * By default, this forbids axes from being specified multiple times. - * - * @param [in] __axes Array with positive or negative indexes. - * @param [in] __shape_size The number of dimensions of the array that - * @ref __axes should be normalized against. - * @param [in] __allow_duplicate Disallow an axis from being specified twice. - * Default: false - * - * @exception std::range_error Particular axis is out of range or other - * error. - * @return The normalized axes indexes, such that `0 <= - * result < __shape_size` - */ -static inline std::vector - get_validated_axes(const std::vector &__axes, - const size_t __shape_size, - const bool __allow_duplicate = false) -{ - std::vector result; - - if (__axes.empty()) { - goto out; - } - - if (__axes.size() > __shape_size) { - goto err; - } - - result.reserve(__axes.size()); - for (std::vector::const_iterator it = __axes.cbegin(); - it != __axes.cend(); ++it) - { - const shape_elem_type _axis = *it; - const shape_elem_type input_shape_size_signed = - static_cast(__shape_size); - if (_axis >= input_shape_size_signed) { // positive axis range check - goto err; - } - - if (_axis < -input_shape_size_signed) { // negative axis range check - goto err; - } - - const shape_elem_type positive_axis = - _axis < 0 ? (_axis + input_shape_size_signed) : _axis; - - if (!__allow_duplicate) { - if (std::find(result.begin(), result.end(), positive_axis) != - result.end()) { // find axis duplication - goto err; - } - } - - result.push_back(positive_axis); - } - -out: - return result; - -err: - // TODO exception if wrong axis? need common function for throwing - // exceptions - throw std::range_error( - "DPNP Error: validate_axes() failed with axis check"); -} - /** * @ingroup BACKEND_UTILS * @brief check support of type T by SYCL device. diff --git a/dpnp/backend/tests/CMakeLists.txt b/dpnp/backend/tests/CMakeLists.txt index b3deef8ff584..bc220eaeee80 100644 --- a/dpnp/backend/tests/CMakeLists.txt +++ b/dpnp/backend/tests/CMakeLists.txt @@ -46,11 +46,8 @@ link_directories(${GTEST_LIB_DIR}) # TODO split add_executable(dpnpc_tests - test_broadcast_iterator.cpp test_main.cpp - test_random.cpp - test_utils.cpp - test_utils_iterator.cpp) + test_random.cpp) target_link_libraries(dpnpc_tests GTest::GTest GTest::Main pthread dpnp_backend_library) # TODO split diff --git a/dpnp/backend/tests/dpnp_test_utils.hpp b/dpnp/backend/tests/dpnp_test_utils.hpp deleted file mode 100644 index cfae5b3f6c0e..000000000000 --- a/dpnp/backend/tests/dpnp_test_utils.hpp +++ /dev/null @@ -1,34 +0,0 @@ -#include -#include - -#include "dpnp_iterator.hpp" - -using namespace std; -using dpnpc_it_t = DPNPC_id::iterator; -using dpnpc_value_t = dpnpc_it_t::value_type; -using dpnpc_index_t = dpnpc_it_t::size_type; - -template -vector<_DataType> get_input_data(const vector &shape) -{ - const dpnpc_index_t size = - accumulate(shape.begin(), shape.end(), dpnpc_index_t(1), - multiplies()); - - 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 -_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; -} diff --git a/dpnp/backend/tests/test_broadcast_iterator.cpp b/dpnp/backend/tests/test_broadcast_iterator.cpp deleted file mode 100644 index 9b97e82bc681..000000000000 --- a/dpnp/backend/tests/test_broadcast_iterator.cpp +++ /dev/null @@ -1,176 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016-2025, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -#include -#include -#include - -#include "dpnp_iterator.hpp" -#include "dpnp_test_utils.hpp" - -#include "queue_sycl.hpp" - -struct IteratorParameters -{ - vector input_shape; - vector output_shape; - vector result; - - /// Operator needs to print this container in human readable form in error - /// reporting - friend std::ostream &operator<<(std::ostream &out, - const IteratorParameters &data) - { - out << "IteratorParameters(input_shape=" << data.input_shape - << ", output_shape=" << data.output_shape - << ", result=" << data.result << ")"; - - return out; - } -}; - -class IteratorBroadcasting : public ::testing::TestWithParam -{ -}; - -TEST_P(IteratorBroadcasting, loop_broadcast) -{ - using data_type = double; - - const IteratorParameters ¶m = GetParam(); - std::vector input_data = - get_input_data(param.input_shape); - - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - - DPNPC_id input(q_ref, input_data.data(), param.input_shape); - input.broadcast_to_shape(param.output_shape); - - ASSERT_EQ(input.get_output_size(), param.result.size()); - - for (dpnpc_index_t output_id = 0; output_id < input.get_output_size(); - ++output_id) - { - EXPECT_EQ(input[output_id], param.result.at(output_id)); - } -} - -TEST_P(IteratorBroadcasting, sycl_broadcast) -{ - using data_type = double; - - const IteratorParameters ¶m = GetParam(); - const dpnpc_index_t result_size = param.result.size(); - data_type *result = reinterpret_cast( - dpnp_memory_alloc_c(result_size * sizeof(data_type))); - - std::vector input_data = - get_input_data(param.input_shape); - data_type *shared_data = get_shared_data(input_data); - - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - - DPNPC_id *input_it; - input_it = reinterpret_cast *>( - dpnp_memory_alloc_c(q_ref, sizeof(DPNPC_id))); - new (input_it) DPNPC_id(q_ref, shared_data, param.input_shape); - - input_it->broadcast_to_shape(param.output_shape); - - ASSERT_EQ(input_it->get_output_size(), result_size); - - sycl::range<1> gws(result_size); - auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { - const size_t idx = global_id[0]; - result[idx] = (*input_it)[idx]; - }; - - auto kernel_func = [&](sycl::handler &cgh) { - cgh.parallel_for( - gws, kernel_parallel_for_func); - }; - - sycl::event event = DPNP_QUEUE.submit(kernel_func); - event.wait(); - - for (dpnpc_index_t i = 0; i < result_size; ++i) { - EXPECT_EQ(result[i], param.result.at(i)); - } - - input_it->~DPNPC_id(); - dpnp_memory_free_c(shared_data); - dpnp_memory_free_c(result); -} - -/** - * Expected values produced by following script: - * - * import numpy as np - * - * input_size = 12 - * input_shape = [3, 4] - * input = np.arange(1, input_size + 1, dtype=np.int64).reshape(input_shape) - * print(f"input shape={input.shape}") - * print(f"input:\n{input}\n") - * - * output_shape = [2, 3, 4] - * output = np.ones(output_shape, dtype=np.int64) - * print(f"output shape={output.shape}") - * - * result = input * output - * print(f"result={np.array2string(result.reshape(result.size), separator=', - * ')}\n", sep=", ") - */ -INSTANTIATE_TEST_SUITE_P( - TestBroadcastIterator, - IteratorBroadcasting, - testing::Values( - IteratorParameters{{1}, {1}, {1}}, - IteratorParameters{{1}, {4}, {1, 1, 1, 1}}, - IteratorParameters{{1}, {3, 4}, {1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1}}, - IteratorParameters{{1}, {2, 3, 4}, {1, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1}}, - IteratorParameters{{4}, {4}, {1, 2, 3, 4}}, - IteratorParameters{{4}, {3, 4}, {1, 2, 3, 4, 1, 2, 3, 4, 1, 2, 3, 4}}, - IteratorParameters{{4}, {2, 3, 4}, {1, 2, 3, 4, 1, 2, 3, 4, - 1, 2, 3, 4, 1, 2, 3, 4, - 1, 2, 3, 4, 1, 2, 3, 4}}, - IteratorParameters{{3, 4}, - {3, 4}, - {1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12}}, - IteratorParameters{{3, 4}, {2, 3, 4}, {1, 2, 3, 4, 5, 6, 7, 8, - 9, 10, 11, 12, 1, 2, 3, 4, - 5, 6, 7, 8, 9, 10, 11, 12}}, - IteratorParameters{{2, 3, 4}, - {2, 3, 4}, - {1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, - 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24}}, - IteratorParameters{{2, 3, 1}, {2, 3, 4}, {1, 1, 1, 1, 2, 2, 2, 2, - 3, 3, 3, 3, 4, 4, 4, 4, - 5, 5, 5, 5, 6, 6, 6, 6}}, - IteratorParameters{{2, 1, 4}, {2, 3, 4}, {1, 2, 3, 4, 1, 2, 3, 4, - 1, 2, 3, 4, 5, 6, 7, 8, - 5, 6, 7, 8, 5, 6, 7, 8}})); diff --git a/dpnp/backend/tests/test_utils.cpp b/dpnp/backend/tests/test_utils.cpp deleted file mode 100644 index dc3a94b2d921..000000000000 --- a/dpnp/backend/tests/test_utils.cpp +++ /dev/null @@ -1,97 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016-2025, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -#include -#include - -#include "dpnp_utils.hpp" - -using namespace std; - -struct AxesParameters -{ - vector axes; - size_t shape_size = size_t{}; - bool duplications = bool{}; - vector result; - - /// Operator needs to print this container in human readable form in error - /// reporting - friend std::ostream &operator<<(std::ostream &out, - const AxesParameters &data) - { - out << "AxesParameters(axes=" << data.axes - << ", shape_size=" << data.shape_size - << ", duplications=" << data.duplications - << ", result=" << data.result << ")"; - - return out; - } -}; - -struct AxesNormalization : public ::testing::TestWithParam -{ - // need to change test name string - struct PrintToStringParamName - { - template - string operator()(const testing::TestParamInfo &info) const - { - const AxesParameters ¶m = - static_cast(info.param); - stringstream ss; - ss << "axes=" << param.axes << ", shape_size=" << param.shape_size - << ", duplications=" << param.duplications - << ", result=" << param.result; - return ss.str(); - } - }; -}; - -TEST_P(AxesNormalization, get_validated_axes) -{ - const AxesParameters ¶m = GetParam(); - vector result = - get_validated_axes(param.axes, param.shape_size, param.duplications); - EXPECT_EQ(result, param.result); -} - -INSTANTIATE_TEST_SUITE_P( - TestUtilsAxesNormalization, - AxesNormalization, - testing::Values(AxesParameters{{0}, 1, true, {0}}, - AxesParameters{{1}, 4, false, {1}}, - AxesParameters{{-1}, 4, false, {3}}, - AxesParameters{{0, 1, 2, 3}, 4, false, {0, 1, 2, 3}}, - /* AxesParameters{{0, 1, 1, 3}, 4, false, {0, 1, 3}}, */ - AxesParameters{{0, 1, 1, 3}, 4, true, {0, 1, 1, 3}}, - AxesParameters{{-2, 1, -4, 3}, 4, false, {2, 1, 0, 3}}, - AxesParameters{{-4, -3, -2, -1}, 4, false, {0, 1, 2, 3}}, - AxesParameters{{-1, -2, -3, -4}, 4, false, {3, 2, 1, 0}}, - AxesParameters{{}, 0, true, {}}, - AxesParameters{{}, 0, false, {}}, - AxesParameters{{}, 2, true, {}}) - /*, AxesNormalization::PrintToStringParamName()*/ -); diff --git a/dpnp/backend/tests/test_utils_iterator.cpp b/dpnp/backend/tests/test_utils_iterator.cpp deleted file mode 100644 index c625199fac4b..000000000000 --- a/dpnp/backend/tests/test_utils_iterator.cpp +++ /dev/null @@ -1,556 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016-2025, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -#include -#include -#include - -#include "dpnp_iterator.hpp" -#include "dpnp_test_utils.hpp" - -#include "queue_sycl.hpp" - -using namespace std; - -TEST(TestUtilsIterator, begin_prefix_postfix) -{ - using test_it = dpnpc_it_t; - - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - - vector input_data = get_input_data({2}); - DPNPC_id result_obj(q_ref, input_data.data(), {2}); - - test_it begin = result_obj.begin(); - test_it end = result_obj.end(); - - EXPECT_NE(begin, end); - - test_it begin0 = begin; - EXPECT_EQ(begin0, begin); - - test_it begin1 = begin0++; - EXPECT_NE(begin1, begin0); - EXPECT_EQ(begin1, begin); - - begin1++; - EXPECT_EQ(begin1, begin0); - - test_it begin_1 = ++begin0; - EXPECT_EQ(begin_1, begin0); - EXPECT_EQ(begin0, end); -} - -TEST(TestUtilsIterator, take_value) -{ - using test_it = dpnpc_it_t; - - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - - // expected data 1, 2 - vector input_data = get_input_data({2}); - DPNPC_id result_obj(q_ref, input_data.data(), {2}); - - test_it begin = result_obj.begin(); - EXPECT_EQ(*begin, 1); - - ++begin; - EXPECT_EQ(*begin, 2); - - EXPECT_EQ(result_obj[1], 2); -} - -TEST(TestUtilsIterator, take_value_loop) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - - // expected data 1, 2 ,3, 4 - vector input_data = get_input_data({4}); - DPNPC_id result_obj(q_ref, input_data.data(), {4}); - - dpnpc_it_t begin = result_obj.begin(); - for (size_t i = 0; i < input_data.size(); ++i, ++begin) { - EXPECT_EQ(result_obj[i], i + 1); - } -} - -TEST(TestUtilsIterator, take_value_loop_3D) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - - // expected input data 1, 2 ,3, 4...24 - vector input_data = get_input_data({2, 3, 4}); - DPNPC_id result_obj(q_ref, input_data.data(), {2, 3, 4}); - - dpnpc_it_t begin = result_obj.begin(); - for (size_t i = 0; i < input_data.size(); ++i, ++begin) { - EXPECT_EQ(result_obj[i], i + 1); - } -} - -TEST(TestUtilsIterator, take_value_axes_loop_3D) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - - vector expected_data{1, 2, 3, 4, 13, 14, 15, 16}; - // expected input data 1, 2 ,3, 4...24 - vector input_data = get_input_data({2, 3, 4}); - DPNPC_id result_obj(q_ref, input_data.data(), {2, 3, 4}); - result_obj.set_axes({0, 2}); - - vector::iterator expected_it = expected_data.begin(); - DPNPC_id::iterator end = result_obj.end(); - for (DPNPC_id::iterator it = result_obj.begin(); it != end; - ++it, ++expected_it) - { - EXPECT_EQ(*it, *expected_it); - } -} - -TEST(TestUtilsIterator, take_value_axis_0_0) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - - vector input_data = get_input_data({4}); - DPNPC_id result_obj(q_ref, input_data.data(), {2, 2}); - result_obj.set_axis(0); // expected data {{1, 3}, {2 ,4}} with shape {2, 2} - - dpnpc_it_t begin = result_obj.begin(); - dpnpc_it_t end = result_obj.end(); - EXPECT_NE(begin, end); - EXPECT_EQ(*begin, 1); - - ++begin; - EXPECT_EQ(*begin, 3); -} - -TEST(TestUtilsIterator, take_value_axis_0_1) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - - vector input_data = get_input_data({4}); - DPNPC_id result_obj(q_ref, input_data.data(), {2, 2}); - result_obj.set_axis(0); // expected data {{1, 3}, {2 ,4}} with shape {2, 2} - - dpnpc_it_t begin = result_obj.begin(1); - dpnpc_it_t end = result_obj.end(1); - EXPECT_NE(begin, end); - EXPECT_EQ(*begin, 2); - - ++begin; - EXPECT_EQ(*begin, 4); -} - -TEST(TestUtilsIterator, take_value_axis_1) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - - vector input_data = get_input_data({4}); - DPNPC_id result_obj(q_ref, input_data.data(), {2, 2}); - result_obj.set_axis(1); // expected data {{1, 2}, {3 ,4}} - - dpnpc_it_t begin = result_obj.begin(); - dpnpc_it_t end = result_obj.end(); - EXPECT_NE(begin, end); - EXPECT_EQ(*begin, 1); - EXPECT_EQ(*end, 3); // linear data space - - ++begin; - EXPECT_EQ(*begin, 2); -} - -TEST(TestUtilsIterator, full_reduction_with_input_shape) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - - vector input_data = get_input_data({2, 3}); - DPNPC_id result_obj(q_ref, input_data.data(), {2, 3}); - - dpnpc_value_t result = 0; - for (dpnpc_it_t data_it = result_obj.begin(0); data_it != result_obj.end(0); - ++data_it) - { - result += *data_it; - } - - EXPECT_EQ(result, 21); -} - -TEST(TestUtilsIterator, output_size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - - // expected data 1, 2 - vector input_data = get_input_data({2, 3}); - DPNPC_id result_obj(q_ref, input_data.data(), {2, 3}); - - const dpnpc_index_t output_size = result_obj.get_output_size(); - - EXPECT_EQ(output_size, 1); -} - -TEST(TestUtilsIterator, output_size_empty) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - - vector input_data = get_input_data({}); - DPNPC_id result_obj(q_ref, input_data.data(), {}); - - const dpnpc_index_t output_size = result_obj.get_output_size(); - - EXPECT_EQ(output_size, 1); -} - -TEST(TestUtilsIterator, output_size_nullptr) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - - DPNPC_id result_obj(q_ref, nullptr, {}); - - const dpnpc_index_t output_size = result_obj.get_output_size(); - - EXPECT_EQ(output_size, 0); -} - -TEST(TestUtilsIterator, output_size_axis) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - - // expected data 1, 2 - vector input_data = get_input_data({2, 3, 4}); - DPNPC_id result_obj(q_ref, input_data.data(), {2, 3, 4}); - - result_obj.set_axis(1); - const dpnpc_index_t output_size = result_obj.get_output_size(); - EXPECT_EQ(output_size, 8); - - result_obj.set_axis(-2); - const dpnpc_index_t output_size_1 = result_obj.get_output_size(); - EXPECT_EQ(output_size_1, 8); -} - -TEST(TestUtilsIterator, output_size_axis_2D) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - - // expected data 1, 2 - vector input_data = get_input_data({2, 3, 4}); - DPNPC_id result_obj(q_ref, input_data.data(), {2, 3, 4}); - - result_obj.set_axes({0, 2}); - const dpnpc_index_t output_size = result_obj.get_output_size(); - EXPECT_EQ(output_size, 3); - - result_obj.set_axes({-3, 2}); - const dpnpc_index_t output_size_1 = result_obj.get_output_size(); - EXPECT_EQ(output_size_1, 3); -} - -TEST(TestUtilsIterator, iterator_loop) -{ - const dpnpc_it_t::size_type size = 10; - - vector expected = get_input_data({size}); - - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - - dpnpc_value_t input_data[size]; - DPNPC_id result(q_ref, input_data, {size}); - iota(result.begin(), result.end(), 1); - - vector::iterator it_expected = expected.begin(); - dpnpc_it_t it_result = result.begin(); - - for (; it_expected != expected.end(); ++it_expected, ++it_result) { - EXPECT_EQ(*it_expected, *it_result); - } -} - -TEST(TestUtilsIterator, operator_minus) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - - vector input_data = get_input_data({3, 4}); - DPNPC_id obj(q_ref, input_data.data(), {3, 4}); - - EXPECT_EQ(obj.begin() - obj.end(), -12); - EXPECT_EQ(obj.end() - obj.begin(), 12); - - obj.set_axis(0); - EXPECT_EQ(obj.begin() - obj.end(), -3); - EXPECT_EQ(obj.end() - obj.begin(), 3); - - EXPECT_EQ(obj.begin(1) - obj.end(1), -3); - EXPECT_EQ(obj.end(1) - obj.begin(1), 3); - - obj.set_axis(1); - EXPECT_EQ(obj.begin() - obj.end(), -4); - EXPECT_EQ(obj.end() - obj.begin(), 4); - - EXPECT_EQ(obj.begin(1) - obj.end(1), -4); - EXPECT_EQ(obj.end(1) - obj.begin(1), 4); -} - -TEST(TestUtilsIterator, iterator_distance) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - - vector input_data = get_input_data({3, 4}); - DPNPC_id obj(q_ref, input_data.data(), {3, 4}); - - dpnpc_it_t::difference_type default_diff_distance = - std::distance(obj.begin(), obj.end()); - EXPECT_EQ(default_diff_distance, 12); - - obj.set_axis(0); - dpnpc_it_t::difference_type axis_0_diff_distance = - std::distance(obj.begin(), obj.end()); - EXPECT_EQ(axis_0_diff_distance, 3); - - dpnpc_it_t::difference_type axis_0_1_diff_distance = - std::distance(obj.begin(1), obj.end(1)); - EXPECT_EQ(axis_0_1_diff_distance, 3); - - obj.set_axis(1); - dpnpc_it_t::difference_type axis_1_diff_distance = - std::distance(obj.begin(), obj.end()); - EXPECT_EQ(axis_1_diff_distance, 4); - - dpnpc_it_t::difference_type axis_1_1_diff_distance = - std::distance(obj.begin(1), obj.end(1)); - EXPECT_EQ(axis_1_1_diff_distance, 4); -} - -struct IteratorParameters -{ - vector input_shape; - vector axes; - vector result; - - /// Operator needs to print this container in human readable form in error - /// reporting - friend std::ostream &operator<<(std::ostream &out, - const IteratorParameters &data) - { - out << "IteratorParameters(input_shape=" << data.input_shape - << ", axis=" << data.axes << ", result=" << data.result << ")"; - - return out; - } -}; - -class IteratorReduction : public ::testing::TestWithParam -{ -}; - -TEST_P(IteratorReduction, loop_reduce_axis) -{ - const IteratorParameters ¶m = GetParam(); - const dpnpc_index_t result_size = param.result.size(); - - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - - vector input_data = - get_input_data(param.input_shape); - DPNPC_id input(q_ref, input_data.data(), param.input_shape); - input.set_axes(param.axes); - - ASSERT_EQ(input.get_output_size(), result_size); - - vector test_result(result_size, 42); - for (dpnpc_index_t output_id = 0; output_id < result_size; ++output_id) { - test_result[output_id] = 0; - for (dpnpc_it_t data_it = input.begin(output_id); - data_it != input.end(output_id); ++data_it) - { - test_result[output_id] += *data_it; - } - - EXPECT_EQ(test_result[output_id], param.result.at(output_id)); - } -} - -TEST_P(IteratorReduction, pstl_reduce_axis) -{ - using data_type = double; - - const IteratorParameters ¶m = GetParam(); - const dpnpc_index_t result_size = param.result.size(); - - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - - vector input_data = get_input_data(param.input_shape); - DPNPC_id input(q_ref, input_data.data(), param.input_shape); - input.set_axes(param.axes); - - ASSERT_EQ(input.get_output_size(), result_size); - - vector result(result_size, 42); - for (dpnpc_index_t output_id = 0; output_id < result_size; ++output_id) { - auto policy = oneapi::dpl::execution::make_device_policy< - class test_pstl_reduce_axis_kernel>(DPNP_QUEUE); - result[output_id] = - std::reduce(policy, input.begin(output_id), input.end(output_id), - data_type(0), std::plus()); - policy.queue().wait(); - - EXPECT_EQ(result[output_id], param.result.at(output_id)); - } -} - -TEST_P(IteratorReduction, sycl_reduce_axis) -{ - using data_type = double; - - const IteratorParameters ¶m = GetParam(); - const dpnpc_index_t result_size = param.result.size(); - vector result(result_size, 42); - data_type *result_ptr = result.data(); - - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - - vector input_data = get_input_data(param.input_shape); - DPNPC_id input(q_ref, input_data.data(), param.input_shape); - input.set_axes(param.axes); - - ASSERT_EQ(input.get_output_size(), result_size); - - sycl::range<1> gws(result_size); - const DPNPC_id *input_it = &input; - auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { - const size_t idx = global_id[0]; - - data_type accumulator = 0; - for (DPNPC_id::iterator data_it = input_it->begin(idx); - data_it != input_it->end(idx); ++data_it) - { - accumulator += *data_it; - } - result_ptr[idx] = accumulator; - }; - - auto kernel_func = [&](sycl::handler &cgh) { - cgh.parallel_for( - gws, kernel_parallel_for_func); - }; - - 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), param.result.at(i)); - } -} - -/** - * Expected values produced by following script: - * - * import numpy as np - * - * shape = [2, 3, 4] - * size = 24 - * axis=1 - * input = np.arange(1, size + 1).reshape(shape) - * print(f"axis={axis}") - * print(f"input.dtype={input.dtype}") - * print(f"input shape={input.shape}") - * print(f"input:\n{input}\n") - * - * result = np.sum(input, axis=axis) - * print(f"result.dtype={result.dtype}") - * print(f"result shape={result.shape}") - * - * print(f"result={np.array2string(result.reshape(result.size), - * separator=',')}\n", sep=",") - */ -INSTANTIATE_TEST_SUITE_P( - TestUtilsIterator, - IteratorReduction, - testing::Values( - IteratorParameters{{2, 3, 4}, - {0}, - {14, 16, 18, 20, 22, 24, 26, 28, 30, 32, 34, 36}}, - IteratorParameters{{2, 3, 4}, {1}, {15, 18, 21, 24, 51, 54, 57, 60}}, - IteratorParameters{{2, 3, 4}, {2}, {10, 26, 42, 58, 74, 90}}, - IteratorParameters{{1, 1, 1}, {0}, {1}}, - IteratorParameters{{1, 1, 1}, {1}, {1}}, - IteratorParameters{{1, 1, 1}, {2}, {1}}, - IteratorParameters{{2, 3, 4, 2}, {0}, {26, 28, 30, 32, 34, 36, 38, 40, - 42, 44, 46, 48, 50, 52, 54, 56, - 58, 60, 62, 64, 66, 68, 70, 72}}, - IteratorParameters{{2, 3, 4, 2}, - {1}, - {27, 30, 33, 36, 39, 42, 45, 48, 99, 102, 105, 108, - 111, 114, 117, 120}}, - IteratorParameters{ - {2, 3, 4, 2}, - {2}, - {16, 20, 48, 52, 80, 84, 112, 116, 144, 148, 176, 180}}, - IteratorParameters{{2, 3, 4, 2}, {3}, {3, 7, 11, 15, 19, 23, 27, 31, - 35, 39, 43, 47, 51, 55, 59, 63, - 67, 71, 75, 79, 83, 87, 91, 95}}, - IteratorParameters{{2, 3, 4, 2}, - {0, 1}, - {126, 132, 138, 144, 150, 156, 162, 168}}, - IteratorParameters{{2, 3, 4, 2}, {2, 3}, {36, 100, 164, 228, 292, 356}}, - IteratorParameters{ - {2, 3, 4, 2}, - {0, 3}, - {54, 62, 70, 78, 86, 94, 102, 110, 118, 126, 134, 142}}, - IteratorParameters{{2, 3, 4, 2}, {0, 1, 2}, {576, 600}}, - IteratorParameters{{2, 3, 4, 2}, {0, 2, 3}, {264, 392, 520}}, - IteratorParameters{{2, 3, 4, 2}, {0, -2, -1}, {264, 392, 520}}, - IteratorParameters{{3, 4}, {0}, {15, 18, 21, 24}}, - IteratorParameters{{3, 4}, {1}, {10, 26, 42}}, - IteratorParameters{{2, 3, 4, 5, 6}, {0, 1, 2, 3, 4}, {259560}}, - IteratorParameters{{2, 3, 4, 5, 6}, - {0, 1, 3, 4}, - {56790, 62190, 67590, 72990}}, - IteratorParameters{{2, 3, 4, 5, 6}, - {1, 2, 3}, - {10680, 10740, 10800, 10860, 10920, 10980, 32280, - 32340, 32400, 32460, 32520, 32580}}, - IteratorParameters{{2, 3, 4, 5, 6}, - {3, 1, 2}, - {10680, 10740, 10800, 10860, 10920, 10980, 32280, - 32340, 32400, 32460, 32520, 32580}}, - IteratorParameters{{2, 3, 4, 5, 6}, - {0, 3, 1, 2}, - {42960, 43080, 43200, 43320, 43440, 43560}}, - IteratorParameters{{2, 3, 4, 5}, - {1, 3}, - {345, 420, 495, 570, 1245, 1320, 1395, 1470}}, - IteratorParameters{{2, 3, 0, 5}, {1, 3}, {}}, - IteratorParameters{{2, 0, 4, 5}, {1, 3}, {0, 0, 0, 0, 0, 0, 0, 0}}, - // IteratorParameters{{2, 3, -4, 5}, {1, 3}, {}}, - // IteratorParameters{{2, -3, 4, 5}, {1, 3}, {0,0,0,0,0,0,0,0}}, - IteratorParameters{{}, {}, {1}}, - IteratorParameters{{0}, {}, {}}, - IteratorParameters{{}, {0}, {1}}, - IteratorParameters{{0}, {0}, {0}}, - IteratorParameters{ - {1}, - {0}, - {1}}) /*TODO , testing::PrintToStringParamName() */);