From 704c22f1373e1276acb43c71e7e17048271bbc03 Mon Sep 17 00:00:00 2001 From: Adnan AlSinan Date: Tue, 24 Oct 2023 11:05:56 +0100 Subject: [GPU] Update Reverse layer to allow negative axis and reversed axis order - Adds option to use negative axis and inverted axis. - Adds validation tests for the above. Resolves COMPMID-6459 Change-Id: I88afd845d078f92c82ec8529ce7241fccd4c417e Signed-off-by: Adnan AlSinan Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/10523 Tested-by: Arm Jenkins Comments-Addressed: Arm Jenkins Reviewed-by: Viet-Hoa Do Benchmark: Arm Jenkins --- arm_compute/runtime/CL/functions/CLReverse.h | 41 +++++++++++--------- docs/user_guide/operator_list.dox | 4 +- docs/user_guide/release_version_and_change_log.dox | 4 +- src/core/CL/cl_kernels/common/reverse.cl | 28 ++++++++------ src/core/CL/kernels/CLReverseKernel.cpp | 29 ++++++++++---- src/core/CL/kernels/CLReverseKernel.h | 45 ++++++++++++++-------- .../CL/functions/CLDirectDeconvolutionLayer.cpp | 2 +- src/runtime/CL/functions/CLFFTConvolutionLayer.cpp | 5 ++- src/runtime/CL/functions/CLReverse.cpp | 18 +++++---- tests/validation/CL/Reverse.cpp | 29 +++++++------- tests/validation/fixtures/ReverseFixture.h | 31 +-------------- 11 files changed, 124 insertions(+), 112 deletions(-) diff --git a/arm_compute/runtime/CL/functions/CLReverse.h b/arm_compute/runtime/CL/functions/CLReverse.h index 0defc3f28b..46229540b4 100644 --- a/arm_compute/runtime/CL/functions/CLReverse.h +++ b/arm_compute/runtime/CL/functions/CLReverse.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2021 Arm Limited. + * Copyright (c) 2018-2021, 2023 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -21,8 +21,8 @@ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE * SOFTWARE. */ -#ifndef ARM_COMPUTE_CLREVERSE_H -#define ARM_COMPUTE_CLREVERSE_H +#ifndef ACL_ARM_COMPUTE_RUNTIME_CL_FUNCTIONS_CLREVERSE_H +#define ACL_ARM_COMPUTE_RUNTIME_CL_FUNCTIONS_CLREVERSE_H #include "arm_compute/core/Error.h" #include "arm_compute/runtime/CL/ICLSimpleFunction.h" @@ -45,33 +45,38 @@ public: * Valid data type configurations: * |src0 |src1 |dst | * |:--------------|:--------------|:--------------| - * |All |U32 |All | + * |All |U32, S32 |All | * - * @param[in] input Input tensor. Data types supported: All. - * @param[out] output Output tensor. Data type supported: Same as @p input - * @param[in] axis Axis tensor. Contains the indices of the dimensions to reverse. Data type supported: U32 + * @param[in] input Input tensor. Data types supported: All. + * @param[out] output Output tensor. Data type supported: Same as @p input + * @param[in] axis Axis tensor. Contains the indices of the dimensions to reverse. Data type supported: U32/S32 + * @param[in] use_inverted_axis Reverse ACL axis indices convention, if true, (inverted)axis = (tensor_rank - 1) - axis */ - void configure(const ICLTensor *input, ICLTensor *output, const ICLTensor *axis); + void configure(const ICLTensor *input, ICLTensor *output, const ICLTensor *axis, bool use_inverted_axis); /** Initialize the function * - * @param[in] compile_context The compile context to be used. - * @param[in] input Input tensor. Data types supported: All. - * @param[out] output Output tensor. Data type supported: Same as @p input - * @param[in] axis Axis tensor. Contains the indices of the dimensions to reverse. Data type supported: U32 + * @param[in] compile_context The compile context to be used. + * @param[in] input Input tensor. Data types supported: All. + * @param[out] output Output tensor. Data type supported: Same as @p input + * @param[in] axis Axis tensor. Contains the indices of the dimensions to reverse. Data type supported: U32/S32 + * @param[in] use_inverted_axis Reverse ACL axis indices convention, if true, (inverted)axis = (tensor_rank - 1) - axis */ void configure(const CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output, - const ICLTensor *axis); + const ICLTensor *axis, + bool use_inverted_axis); /** Static function to check if given info will lead to a valid configuration of @ref CLReverseKernel * - * @param[in] input Input tensor info. Data types supported: All. - * @param[in] output Output tensor info. Data type supported: Same as @p input - * @param[in] axis Axis tensor info. Contains the indices of the dimensions to reverse. Data type supported: U32 + * @param[in] input Input tensor info. Data types supported: All. + * @param[in] output Output tensor info. Data type supported: Same as @p input + * @param[in] axis Axis tensor info. Contains the indices of the dimensions to reverse. Data type supported: U32/S32 + * @param[in] use_inverted_axis Reverse ACL axis indices convention, if true, (inverted)axis = (tensor_rank - 1) - axis * * @return a status */ - static Status validate(const ITensorInfo *input, const ITensorInfo *output, const ITensorInfo *axis); + static Status + validate(const ITensorInfo *input, const ITensorInfo *output, const ITensorInfo *axis, bool use_inverted_axis); }; } // namespace arm_compute -#endif /* ARM_COMPUTE_CLREVERSE_H */ +#endif // ACL_ARM_COMPUTE_RUNTIME_CL_FUNCTIONS_CLREVERSE_H diff --git a/docs/user_guide/operator_list.dox b/docs/user_guide/operator_list.dox index 498b925cd3..0c74214c73 100644 --- a/docs/user_guide/operator_list.dox +++ b/docs/user_guide/operator_list.dox @@ -1,5 +1,5 @@ /// -/// Copyright (c) 2021-2022 Arm Limited. +/// Copyright (c) 2021-2023 Arm Limited. /// /// SPDX-License-Identifier: MIT /// @@ -2666,7 +2666,7 @@ where N = batches, C = channels, H = height, W = width, D = depth
src0src1dst -
AllU32All +
AllU32, S32All
RNNLayer diff --git a/docs/user_guide/release_version_and_change_log.dox b/docs/user_guide/release_version_and_change_log.dox index 9e67d7e6b4..dbd00eb00a 100644 --- a/docs/user_guide/release_version_and_change_log.dox +++ b/docs/user_guide/release_version_and_change_log.dox @@ -62,8 +62,8 @@ v23.11 Public major release - Add new OpenCL™ kernels: - @ref opencl::kernels::ClMatMulLowpNativeMMULKernel support for QASYMM8 and QASYMM8_SIGNED, with batch support - Deprecate support for Bfloat16 in @ref cpu::CpuCast. - - Add support for negative axis values and inverted axis values in @ref arm_compute::NEReverse. - - Support for U32 axis in @ref arm_compute::NEReverse will be deprecated in 24.02. + - Add support for negative axis values and inverted axis values in @ref arm_compute::NEReverse and @ref arm_compute::CLReverse. + - Support for U32 axis in @ref arm_compute::NEReverse and @ref arm_compute::CLReverse will be deprecated in 24.02. v23.08 Public major release - Deprecate the legacy 'libarm_compute_core' library. This library is an artifact of Compute Library's legacy library architecture and no longer serves any purpose. diff --git a/src/core/CL/cl_kernels/common/reverse.cl b/src/core/CL/cl_kernels/common/reverse.cl index 6b0afb9c2c..f94bfb6640 100644 --- a/src/core/CL/cl_kernels/common/reverse.cl +++ b/src/core/CL/cl_kernels/common/reverse.cl @@ -1,5 +1,5 @@ /* -* Copyright (c) 2018-2021 Arm Limited. +* Copyright (c) 2018-2021, 2023 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -33,6 +33,8 @@ * * @note The data type must be given as a preprocessor argument using -DDATA_TYPE=num. e.g. -DDATA_TYPE=uint * @note The number of dimensions to reverse must be given as a preprocessor argument using -DNUM_REVERSE_DIMS=num, e.g. -DNUM_REVERSE_DIMS=3 + * @note The number of dimensions of the source tensor must be given as a preprocessor argument using -DRANK=num, e.g. -DRANK=3 + * @note The values in axis_tensor must be within [-rank, rank-1]. * * @param[in] src_ptr Pointer to the source tensor. Supported data types: All * @param[in] src_stride_x Stride of the first source tensor in X dimension (in bytes) @@ -78,20 +80,24 @@ __kernel void reverse(TENSOR4D_DECLARATION(src), const uint4 dims = (uint4)(0, 1, 2, 3); int4 to_reverse = (int4)(0, 0, 0, 0); + + VEC_DATA_TYPE(int, NUM_REVERSE_DIMS) indices = VLOAD(NUM_REVERSE_DIMS)(0,(__global int *)axis.ptr); +#if defined(USE_INVERTED_AXIS) + indices = select((VEC_DATA_TYPE(int, NUM_REVERSE_DIMS)) RANK - 1, -1, indices < 0) - indices; +#else /* defined(USE_INVERTED_AXIS) */ + indices = select(indices, indices + RANK, indices < 0); +#endif /* defined(USE_INVERTED_AXIS) */ + #if NUM_REVERSE_DIMS == 1 - const uint index = *((__global uint *)axis.ptr); - to_reverse = (uint4)index == dims; + to_reverse = ((uint4)indices == dims); #elif NUM_REVERSE_DIMS == 2 - const uint2 indices = vload2(0, (__global uint *)axis.ptr); - to_reverse = ((uint4)indices.s0 == dims) || ((uint4)indices.s1 == dims); + to_reverse = ((uint4)indices.s0 == dims) || ((uint4)indices.s1 == dims); #elif NUM_REVERSE_DIMS == 3 - const uint2 indices01 = vload2(0, (__global uint *)axis.ptr); - const uint index2 = *((__global uint *)axis.ptr + 2); - to_reverse = ((uint4)indices01.s0 == dims) || ((uint4)indices01.s1 == dims) || ((uint4)index2 == dims); -#else /* NUM_REVERSE_DIMS == 3 */ - const uint4 indices = vload4(0, (__global uint *)axis.ptr); - to_reverse = ((uint4)indices.s0 == dims) || ((uint4)indices.s1 == dims) || ((uint4)indices.s2 == dims) || ((uint4)indices.s3 == dims); + to_reverse = ((uint4)indices.s0 == dims) || ((uint4)indices.s1 == dims) || ((uint4)indices.s2 == dims); +#else /* NUM_REVERSE_DIMS == 1 */ + to_reverse = ((uint4)indices.s0 == dims) || ((uint4)indices.s1 == dims) || ((uint4)indices.s2 == dims) || ((uint4)indices.s3 == dims); #endif /* NUM_REVERSE_DIMS == 1 */ + const uint x_out = to_reverse.s0 ? width - x_in - 1 : x_in; const uint y_out = to_reverse.s1 ? height - y_in - 1 : y_in; const uint z_out = to_reverse.s2 ? depth - z_in - 1 : z_in; diff --git a/src/core/CL/kernels/CLReverseKernel.cpp b/src/core/CL/kernels/CLReverseKernel.cpp index 79a0f03b1e..00241b161b 100644 --- a/src/core/CL/kernels/CLReverseKernel.cpp +++ b/src/core/CL/kernels/CLReverseKernel.cpp @@ -40,13 +40,17 @@ namespace arm_compute { namespace { -Status validate_arguments(const ITensorInfo *input, const ITensorInfo *output, const ITensorInfo *axis) +Status +validate_arguments(const ITensorInfo *input, const ITensorInfo *output, const ITensorInfo *axis, bool use_inverted_axis) { + ARM_COMPUTE_UNUSED(use_inverted_axis); ARM_COMPUTE_ERROR_ON_NULLPTR(input, output, axis); ARM_COMPUTE_RETURN_ERROR_ON_F16_UNSUPPORTED(input); ARM_COMPUTE_RETURN_ERROR_ON(input->data_type() == DataType::UNKNOWN); - ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(axis, 1, DataType::U32); + ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(axis, 1, DataType::U32, DataType::S32); ARM_COMPUTE_RETURN_ERROR_ON_MSG(axis->num_dimensions() > 1, "Axis must be a 1D tensor"); + ARM_COMPUTE_RETURN_ERROR_ON_MSG(input->num_dimensions() > 4, + "Current implementation only supports up to 4 dimensions."); ARM_COMPUTE_RETURN_ERROR_ON_MSG(axis->dimension(0) > 4, "Only up to 4 dimensions can be reversed"); // Checks performed when output is configured @@ -66,15 +70,19 @@ CLReverseKernel::CLReverseKernel() : _input(nullptr), _output(nullptr), _axis(nu _type = CLKernelType::ELEMENTWISE; } -void CLReverseKernel::configure(const ICLTensor *input, ICLTensor *output, const ICLTensor *axis) +void CLReverseKernel::configure(const ICLTensor *input, + ICLTensor *output, + const ICLTensor *axis, + bool use_inverted_axis) { - configure(CLKernelLibrary::get().get_compile_context(), input, output, axis); + configure(CLKernelLibrary::get().get_compile_context(), input, output, axis, use_inverted_axis); } void CLReverseKernel::configure(const CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output, - const ICLTensor *axis) + const ICLTensor *axis, + bool use_inverted_axis) { ARM_COMPUTE_ERROR_ON_NULLPTR(input, output, axis); auto padding_info = get_padding_info({input, output, axis}); @@ -86,12 +94,14 @@ void CLReverseKernel::configure(const CLCompileContext &compile_context, // Output tensor auto initialization if not yet initialized auto_init_if_empty(*output->info(), *input->info()->clone()); - ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(input->info(), output->info(), axis->info())); + ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(input->info(), output->info(), axis->info(), use_inverted_axis)); // Set kernel build options CLBuildOptions build_opts; build_opts.add_option("-DNUM_REVERSE_DIMS=" + support::cpp11::to_string(axis->info()->dimension(0))); build_opts.add_option("-DDATA_TYPE=" + get_cl_unsigned_type_from_element_size(input->info()->element_size())); + build_opts.add_option("-DRANK=" + support::cpp11::to_string(input->info()->num_dimensions())); + build_opts.add_option_if(use_inverted_axis, "-DUSE_INVERTED_AXIS"); // Create kernel _kernel = create_kernel(compile_context, "reverse", build_opts.options()); @@ -119,9 +129,12 @@ void CLReverseKernel::configure(const CLCompileContext &compile_context, ARM_COMPUTE_ERROR_ON(has_padding_changed(padding_info)); } -Status CLReverseKernel::validate(const ITensorInfo *input, const ITensorInfo *output, const ITensorInfo *axis) +Status CLReverseKernel::validate(const ITensorInfo *input, + const ITensorInfo *output, + const ITensorInfo *axis, + bool use_inverted_axis) { - ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(input, output, axis)); + ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(input, output, axis, use_inverted_axis)); return Status{}; } diff --git a/src/core/CL/kernels/CLReverseKernel.h b/src/core/CL/kernels/CLReverseKernel.h index fbd99dc883..a630aec15a 100644 --- a/src/core/CL/kernels/CLReverseKernel.h +++ b/src/core/CL/kernels/CLReverseKernel.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2020 Arm Limited. + * Copyright (c) 2018-2020, 2023 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -21,8 +21,8 @@ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE * SOFTWARE. */ -#ifndef ARM_COMPUTE_CLREVERSEKERNEL_H -#define ARM_COMPUTE_CLREVERSEKERNEL_H +#ifndef ACL_SRC_CORE_CL_KERNELS_CLREVERSEKERNEL_H +#define ACL_SRC_CORE_CL_KERNELS_CLREVERSEKERNEL_H #include "src/core/CL/ICLKernel.h" @@ -48,32 +48,43 @@ public: ~CLReverseKernel() = default; /** Initialise the kernel's inputis and output * - * @param[in] input Input tensor. Data types supported: All. - * @param[out] output Output tensor. Data type supported: Same as @p input - * @param[in] axis Axis tensor. Contains the indices of the dimensions to reverse. Data type supported: U32 + * @param[in] input Input tensor. Data types supported: All. + * @param[out] output Output tensor. Data type supported: Same as @p input + * @param[in] axis Axis tensor. Contains the indices of the dimensions to reverse. Data type supported: U32/S32 + * @param[in] use_inverted_axis Reverse ACL axis indices convention i.e. acl.dim(0) = tensor_rank -1 + * + * @note The value of each axis should be between [-rank, rank) + * @note If there are duplicate values in the tensor, the subsequent axis values are ignored. e.g. an array of [2, 2] has the same effects as [2]. + * + * @deprecated Support for U32 in axis tensor will be removed in 24.02 release + * */ - void configure(const ICLTensor *input, ICLTensor *output, const ICLTensor *axis); + void configure(const ICLTensor *input, ICLTensor *output, const ICLTensor *axis, bool use_inverted_axis); /** Initialise the kernel's inputis and output * - * @param[in] compile_context The compile context to be used. - * @param[in] input Input tensor. Data types supported: All. - * @param[out] output Output tensor. Data type supported: Same as @p input - * @param[in] axis Axis tensor. Contains the indices of the dimensions to reverse. Data type supported: U32 + * @param[in] compile_context The compile context to be used. + * @param[in] input Input tensor. Data types supported: All. + * @param[out] output Output tensor. Data type supported: Same as @p input + * @param[in] axis Axis tensor. Contains the indices of the dimensions to reverse. Data type supported: U32/S32 + * @param[in] use_inverted_axis Reverse ACL axis indices convention i.e. acl.dim(0) = tensor_rank -1 */ void configure(const CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output, - const ICLTensor *axis); + const ICLTensor *axis, + bool use_inverted_axis); /** Static function to check if given info will lead to a valid configuration of @ref CLReverseKernel * - * @param[in] input Input tensor info. Data types supported: All. - * @param[in] output Output tensor info. Data type supported: Same as @p input - * @param[in] axis Axis tensor info. Contains the indices of the dimensions to reverse. Data type supported: U32 + * @param[in] input Input tensor info. Data types supported: All. + * @param[in] output Output tensor info. Data type supported: Same as @p input + * @param[in] axis Axis tensor info. Contains the indices of the dimensions to reverse. Data type supported: U32/S32 + * @param[in] use_inverted_axis Reverse ACL axis indices convention i.e. acl.dim(0) = tensor_rank -1 * * @return a status */ - static Status validate(const ITensorInfo *input, const ITensorInfo *output, const ITensorInfo *axis); + static Status + validate(const ITensorInfo *input, const ITensorInfo *output, const ITensorInfo *axis, bool use_inverted_axis); // Inherited methods overridden: void run(const Window &window, cl::CommandQueue &queue) override; @@ -84,4 +95,4 @@ public: const ICLTensor *_axis; }; } // namespace arm_compute -#endif /*ARM_COMPUTE_CLREVERSEKERNEL_H */ +#endif // ACL_SRC_CORE_CL_KERNELS_CLREVERSEKERNEL_H diff --git a/src/runtime/CL/functions/CLDirectDeconvolutionLayer.cpp b/src/runtime/CL/functions/CLDirectDeconvolutionLayer.cpp index 3717f30ae1..7cd268ab0b 100644 --- a/src/runtime/CL/functions/CLDirectDeconvolutionLayer.cpp +++ b/src/runtime/CL/functions/CLDirectDeconvolutionLayer.cpp @@ -161,7 +161,7 @@ void CLDirectDeconvolutionLayer::configure(const CLCompileContext &compile_conte _original_weights = weights; _flip_axis.allocator()->init(TensorInfo(TensorShape(2U), 1, DataType::U32)); _weights_flipped.allocator()->init(weights->info()->clone()->set_data_layout(data_layout)); - _flip_weights.configure(compile_context, weights, &_weights_flipped, &_flip_axis); + _flip_weights.configure(compile_context, weights, &_weights_flipped, &_flip_axis, /* use_inverted_axis */ false); auto out_dims = deconvolution_output_dimensions(input->info()->dimension(idx_w), input->info()->dimension(idx_h), diff --git a/src/runtime/CL/functions/CLFFTConvolutionLayer.cpp b/src/runtime/CL/functions/CLFFTConvolutionLayer.cpp index 3894b10785..2a73517549 100644 --- a/src/runtime/CL/functions/CLFFTConvolutionLayer.cpp +++ b/src/runtime/CL/functions/CLFFTConvolutionLayer.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2021 Arm Limited. + * Copyright (c) 2019-2021, 2023 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -181,7 +181,8 @@ void CLFFTConvolutionLayer::configure(const CLCompileContext &compile_context // Flip weights _flipped_weights.allocator()->init(weights_to_use->info()->clone()->set_is_resizable(true).reset_padding()); _flip_axis.allocator()->init(TensorInfo(TensorShape(2U), 1, DataType::U32)); - _flip_weights_func.configure(compile_context, weights_to_use, &_flipped_weights, &_flip_axis); + _flip_weights_func.configure(compile_context, weights_to_use, &_flipped_weights, &_flip_axis, + /* use_inverted_axis */ false); // Pad weights const PaddingList padding_w = {{0, input_dims.x() + pad_valid.x() - 1}, {0, input_dims.y() + pad_valid.y() - 1}}; diff --git a/src/runtime/CL/functions/CLReverse.cpp b/src/runtime/CL/functions/CLReverse.cpp index 415de52e64..a20be2335a 100644 --- a/src/runtime/CL/functions/CLReverse.cpp +++ b/src/runtime/CL/functions/CLReverse.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2021 Arm Limited. + * Copyright (c) 2018-2021, 2023 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -30,24 +30,28 @@ namespace arm_compute { -void CLReverse::configure(const ICLTensor *input, ICLTensor *output, const ICLTensor *axis) +void CLReverse::configure(const ICLTensor *input, ICLTensor *output, const ICLTensor *axis, bool use_inverted_axis) { - configure(CLKernelLibrary::get().get_compile_context(), input, output, axis); + configure(CLKernelLibrary::get().get_compile_context(), input, output, axis, use_inverted_axis); } void CLReverse::configure(const CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output, - const ICLTensor *axis) + const ICLTensor *axis, + bool use_inverted_axis) { ARM_COMPUTE_LOG_PARAMS(input, output, axis); auto k = std::make_unique(); - k->configure(compile_context, input, output, axis); + k->configure(compile_context, input, output, axis, use_inverted_axis); _kernel = std::move(k); } -Status CLReverse::validate(const ITensorInfo *input, const ITensorInfo *output, const ITensorInfo *axis) +Status CLReverse::validate(const ITensorInfo *input, + const ITensorInfo *output, + const ITensorInfo *axis, + bool use_inverted_axis) { - return CLReverseKernel::validate(input, output, axis); + return CLReverseKernel::validate(input, output, axis, use_inverted_axis); } } // namespace arm_compute diff --git a/tests/validation/CL/Reverse.cpp b/tests/validation/CL/Reverse.cpp index ff46ba64ad..82effc2136 100644 --- a/tests/validation/CL/Reverse.cpp +++ b/tests/validation/CL/Reverse.cpp @@ -44,7 +44,7 @@ namespace validation using framework::dataset::make; namespace { -auto run_small_dataset = combine(datasets::SmallShapes(), datasets::Tiny1DShapes()); +auto run_small_dataset = combine(datasets::Small3DShapes(), datasets::Tiny1DShapes()); auto run_large_dataset = combine(datasets::LargeShapes(), datasets::Tiny1DShapes()); } // namespace @@ -80,7 +80,8 @@ DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip( { Status s = CLReverse::validate(&src_info.clone()->set_is_resizable(false), &dst_info.clone()->set_is_resizable(false), - &axis_info.clone()->set_is_resizable(false)); + &axis_info.clone()->set_is_resizable(false), + false); ARM_COMPUTE_EXPECT(bool(s) == expected, framework::LogLevel::ERRORS); } // clang-format on @@ -97,8 +98,8 @@ FIXTURE_DATA_TEST_CASE(RunSmall, combine( run_small_dataset, make("DataType", DataType::F16), - make("use_negative_axis", { false }), - make("use_inverted_axis", { false }))) + make("use_negative_axis", { true, false }), + make("use_inverted_axis", { true, false }))) { // Validate output validate(CLAccessor(_target), _reference); @@ -110,8 +111,8 @@ FIXTURE_DATA_TEST_CASE(RunLarge, combine( run_large_dataset, make("DataType", DataType::F16), - make("use_negative_axis", { false }), - make("use_inverted_axis", { false }))) + make("use_negative_axis", { true, false }), + make("use_inverted_axis", { true, false }))) { // Validate output validate(CLAccessor(_target), _reference); @@ -125,8 +126,8 @@ FIXTURE_DATA_TEST_CASE(RunSmall, combine( run_small_dataset, make("DataType", DataType::F32), - make("use_negative_axis", { false }), - make("use_inverted_axis", { false }))) + make("use_negative_axis", { true, false }), + make("use_inverted_axis", { true, false }))) { // Validate output validate(CLAccessor(_target), _reference); @@ -138,8 +139,8 @@ FIXTURE_DATA_TEST_CASE(RunLarge, combine( run_large_dataset, make("DataType", DataType::F32), - make("use_negative_axis", { false }), - make("use_inverted_axis", { false }))) + make("use_negative_axis", { true, false }), + make("use_inverted_axis", { true, false }))) { // Validate output validate(CLAccessor(_target), _reference); @@ -155,8 +156,8 @@ FIXTURE_DATA_TEST_CASE(RunSmall, combine( run_small_dataset, make("DataType", DataType::QASYMM8), - make("use_negative_axis", { false }), - make("use_inverted_axis", { false }))) + make("use_negative_axis", { true, false }), + make("use_inverted_axis", { true, false }))) { // Validate output validate(CLAccessor(_target), _reference); @@ -168,8 +169,8 @@ FIXTURE_DATA_TEST_CASE(RunLarge, combine( run_large_dataset, make("DataType", DataType::QASYMM8), - make("use_negative_axis", { false }), - make("use_inverted_axis", { false }))) + make("use_negative_axis", { true, false }), + make("use_inverted_axis", { true, false }))) { // Validate output validate(CLAccessor(_target), _reference); diff --git a/tests/validation/fixtures/ReverseFixture.h b/tests/validation/fixtures/ReverseFixture.h index d53945191e..856bff7b12 100644 --- a/tests/validation/fixtures/ReverseFixture.h +++ b/tests/validation/fixtures/ReverseFixture.h @@ -27,9 +27,6 @@ #include "arm_compute/core/Helpers.h" #include "arm_compute/core/TensorShape.h" #include "arm_compute/core/Types.h" -#ifdef ARM_COMPUTE_OPENCL_ENABLED -#include "arm_compute/runtime/CL/functions/CLReverse.h" -#endif // ARM_COMPUTE_OPENCL_ENABLED #include "arm_compute/runtime/Tensor.h" #include "tests/AssetsLibrary.h" #include "tests/Globals.h" @@ -44,31 +41,6 @@ namespace test { namespace validation { -namespace -{ -template -#ifdef ARM_COMPUTE_OPENCL_ENABLED -std::enable_if_t < !std::is_same::value, void > -#else // ARM_COMPUTE_OPENCL_ENABLED -void -#endif // ARM_COMPUTE_OPENCL_ENABLED -configureReverse(ReverseFunction &func, TensorType &src, TensorType &axis, TensorType &dst, bool use_inverted_axis) -{ - func.configure(&src, &dst, &axis, use_inverted_axis); -} - -#ifdef ARM_COMPUTE_OPENCL_ENABLED -template -std::enable_if_t::value, void> -configureReverse(ReverseFunction &func, TensorType &src, TensorType &axis, TensorType &dst, bool use_inverted_axis) -{ - ARM_COMPUTE_UNUSED(use_inverted_axis); - func.configure(&src, &dst, &axis); -} - -#endif // ARM_COMPUTE_OPENCL_ENABLED -} //namespace - template class ReverseValidationFixture : public framework::Fixture { @@ -113,8 +85,7 @@ protected: // Create and configure function FunctionType reverse_func; - - configureReverse(reverse_func, src, axis, dst, use_inverted_axis); + reverse_func.configure(&src, &dst, &axis, use_inverted_axis); ARM_COMPUTE_ASSERT(src.info()->is_resizable()); ARM_COMPUTE_ASSERT(axis.info()->is_resizable()); -- cgit v1.2.1