From a387e271b1e02ffd5c2993702b9a21c1ed5c95fa Mon Sep 17 00:00:00 2001 From: Sheri Zhang Date: Tue, 29 Jun 2021 17:34:06 +0100 Subject: Add in-place calculation support for CL elementwise arithmetic kernels - Add in-place calculation support in ClArithmeticKernel, ClSaturatedArithmeticKernel and ClMulKernel - Add in-place test cases Resolves: COMPMID-4431 Signed-off-by: Sheri Zhang Change-Id: Id484bdb76b74478a33fedb471ae0c7f799c599f6 Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/5885 Comments-Addressed: Arm Jenkins Reviewed-by: SiCong Li Tested-by: Arm Jenkins --- src/core/CL/cl_kernels/elementwise_operation.cl | 21 ++- .../cl_kernels/elementwise_operation_quantized.cl | 21 ++- src/core/CL/cl_kernels/pixelwise_mul_float.cl | 13 +- src/core/CL/cl_kernels/pixelwise_mul_int.cl | 26 ++- src/core/gpu/cl/kernels/ClElementwiseKernel.cpp | 50 ++++- src/core/gpu/cl/kernels/ClElementwiseKernel.h | 2 + src/core/gpu/cl/kernels/ClMulKernel.cpp | 32 +++- src/core/gpu/cl/kernels/ClMulKernel.h | 6 +- src/graph/mutators/InPlaceOperationMutator.cpp | 8 +- tests/datasets/ShapeDatasets.h | 19 ++ tests/validation/CL/ArithmeticAddition.cpp | 109 +++++++---- tests/validation/CL/ArithmeticDivision.cpp | 52 ++++-- tests/validation/CL/ArithmeticSubtraction.cpp | 23 ++- tests/validation/CL/ElementwiseMax.cpp | 52 ++++-- tests/validation/CL/ElementwiseMin.cpp | 52 ++++-- tests/validation/CL/ElementwisePower.cpp | 44 +++-- tests/validation/CL/ElementwiseSquaredDiff.cpp | 48 +++-- tests/validation/CL/PixelWiseMultiplication.cpp | 47 ++++- tests/validation/NEON/ArithmeticAddition.cpp | 91 ++++++---- tests/validation/NEON/ArithmeticSubtraction.cpp | 18 +- tests/validation/NEON/ElementwiseDivision.cpp | 31 +++- tests/validation/NEON/ElementwiseMax.cpp | 59 ++++-- tests/validation/NEON/ElementwiseMin.cpp | 62 +++++-- tests/validation/NEON/ElementwisePower.cpp | 29 ++- tests/validation/NEON/ElementwiseSquareDiff.cpp | 56 ++++-- tests/validation/NEON/PixelWiseMultiplication.cpp | 61 +++++-- .../fixtures/ArithmeticOperationsFixture.h | 117 +++++++----- .../fixtures/ElementwiseOperationsFixture.h | 201 ++++++++++++--------- .../fixtures/PixelWiseMultiplicationFixture.h | 44 +++-- 29 files changed, 966 insertions(+), 428 deletions(-) diff --git a/src/core/CL/cl_kernels/elementwise_operation.cl b/src/core/CL/cl_kernels/elementwise_operation.cl index 99f725645d..45dcbfc6e2 100644 --- a/src/core/CL/cl_kernels/elementwise_operation.cl +++ b/src/core/CL/cl_kernels/elementwise_operation.cl @@ -97,8 +97,12 @@ */ __kernel void OP_FUN_NAME(OP)( TENSOR3D_DECLARATION(in1), - TENSOR3D_DECLARATION(in2), - TENSOR3D_DECLARATION(out)) + TENSOR3D_DECLARATION(in2) +#if !defined(IN_PLACE) + , + TENSOR3D_DECLARATION(out) +#endif // !defined(IN_PLACE) +) { #if VEC_SIZE_IN1 == 1 uint in1_x_offs = 0; @@ -110,12 +114,23 @@ __kernel void OP_FUN_NAME(OP)( #else // VEC_SIZE_IN2 == 1 uint in2_x_offs = max((int)(get_global_id(0) * VEC_SIZE_IN2 - (VEC_SIZE_IN2 - VEC_SIZE_LEFTOVER) % VEC_SIZE_IN2), 0); #endif // VEC_SIZE_IN2 == 1 +#if !defined(IN_PLACE) uint out_x_offs = max((int)(get_global_id(0) * VEC_SIZE_OUT - (VEC_SIZE_OUT - VEC_SIZE_LEFTOVER) % VEC_SIZE_OUT), 0); +#endif // !defined(IN_PLACE) // Get pixels pointer __global uchar *in1_addr = in1_ptr + in1_offset_first_element_in_bytes + in1_x_offs * sizeof(DATA_TYPE) + get_global_id(1) * in1_step_y + get_global_id(2) * in1_step_z; __global uchar *in2_addr = in2_ptr + in2_offset_first_element_in_bytes + in2_x_offs * sizeof(DATA_TYPE) + get_global_id(1) * in2_step_y + get_global_id(2) * in2_step_z; - __global uchar *out_addr = out_ptr + out_offset_first_element_in_bytes + out_x_offs * sizeof(DATA_TYPE) + get_global_id(1) * out_step_y + get_global_id(2) * out_step_z; + __global uchar * +#if !defined(IN_PLACE) + out_addr = out_ptr + out_offset_first_element_in_bytes + out_x_offs * sizeof(DATA_TYPE) + get_global_id(1) * out_step_y + get_global_id(2) * out_step_z; +#else // !defined(IN_PLACE) +#if defined(SRC1_IN_PLACE) + out_addr = in1_addr; +#else //defined(SRC1_IN_PLACE) + out_addr = in2_addr; +#endif //defined(SRC1_IN_PLACE) +#endif // !defined(IN_PLACE) // Load values VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE_OUT) diff --git a/src/core/CL/cl_kernels/elementwise_operation_quantized.cl b/src/core/CL/cl_kernels/elementwise_operation_quantized.cl index 0051babf03..a11be80875 100644 --- a/src/core/CL/cl_kernels/elementwise_operation_quantized.cl +++ b/src/core/CL/cl_kernels/elementwise_operation_quantized.cl @@ -86,8 +86,12 @@ */ __kernel void OP_FUN_NAME(OP)( TENSOR3D_DECLARATION(in1), - TENSOR3D_DECLARATION(in2), - TENSOR3D_DECLARATION(out)) + TENSOR3D_DECLARATION(in2) +#if !defined(IN_PLACE) + , + TENSOR3D_DECLARATION(out) +#endif // !defined(IN_PLACE) +) { #if VEC_SIZE_IN1 == 1 uint in1_x_offs = 0; @@ -99,12 +103,23 @@ __kernel void OP_FUN_NAME(OP)( #else // VEC_SIZE_IN2 == 1 uint in2_x_offs = max((int)(get_global_id(0) * VEC_SIZE_IN2 - (VEC_SIZE_IN2 - VEC_SIZE_LEFTOVER) % VEC_SIZE_IN2), 0); #endif // VEC_SIZE_IN2 == 1 +#if !defined(IN_PLACE) uint out_x_offs = max((int)(get_global_id(0) * VEC_SIZE_OUT - (VEC_SIZE_OUT - VEC_SIZE_LEFTOVER) % VEC_SIZE_OUT), 0); +#endif // !defined(IN_PLACE) // Get pixels pointer __global uchar *in1_addr = in1_ptr + in1_offset_first_element_in_bytes + in1_x_offs * sizeof(DATA_TYPE) + get_global_id(1) * in1_step_y + get_global_id(2) * in1_step_z; __global uchar *in2_addr = in2_ptr + in2_offset_first_element_in_bytes + in2_x_offs * sizeof(DATA_TYPE) + get_global_id(1) * in2_step_y + get_global_id(2) * in2_step_z; - __global uchar *out_addr = out_ptr + out_offset_first_element_in_bytes + out_x_offs * sizeof(DATA_TYPE) + get_global_id(1) * out_step_y + get_global_id(2) * out_step_z; + __global uchar * +#if !defined(IN_PLACE) + out_addr = out_ptr + out_offset_first_element_in_bytes + out_x_offs * sizeof(DATA_TYPE) + get_global_id(1) * out_step_y + get_global_id(2) * out_step_z; +#else // !defined(IN_PLACE) +#if defined(SRC1_IN_PLACE) + out_addr = in1_addr; +#else //defined(SRC1_IN_PLACE) + out_addr = in2_addr; +#endif //defined(SRC1_IN_PLACE) +#endif // !defined(IN_PLACE) VEC_INT in_a = CONVERT((VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE_OUT))(VLOAD(VEC_SIZE_IN1)(0, (__global DATA_TYPE *)in1_addr)), VEC_INT); VEC_INT in_b = CONVERT((VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE_OUT))(VLOAD(VEC_SIZE_IN2)(0, (__global DATA_TYPE *)in2_addr)), VEC_INT); diff --git a/src/core/CL/cl_kernels/pixelwise_mul_float.cl b/src/core/CL/cl_kernels/pixelwise_mul_float.cl index 0016775893..10875293a9 100644 --- a/src/core/CL/cl_kernels/pixelwise_mul_float.cl +++ b/src/core/CL/cl_kernels/pixelwise_mul_float.cl @@ -77,7 +77,9 @@ __kernel void pixelwise_mul_float( TENSOR3D_DECLARATION(in1), TENSOR3D_DECLARATION(in2), +#if !defined(IN_PLACE) TENSOR3D_DECLARATION(out), +#endif // !defined(IN_PLACE) const float scale) { // Get pixels pointer @@ -87,7 +89,16 @@ __kernel void pixelwise_mul_float( __global uchar *in1_addr = in1_ptr + in1_offset_first_element_in_bytes + x * in1_stride_x + y * in1_stride_y + z * in1_stride_z; __global uchar *in2_addr = in2_ptr + in2_offset_first_element_in_bytes + x * in2_stride_x + y * in2_stride_y + z * in2_stride_z; - __global uchar *out_addr = out_ptr + out_offset_first_element_in_bytes + x * out_stride_x + y * out_stride_y + z * out_stride_z; + __global uchar * +#if !defined(IN_PLACE) + out_addr = out_ptr + out_offset_first_element_in_bytes + x * out_stride_x + y * out_stride_y + z * out_stride_z; +#else // !defined(IN_PLACE) +#if defined(SRC1_IN_PLACE) + out_addr = in1_addr; +#else //defined(SRC1_IN_PLACE) + out_addr = in2_addr; +#endif //defined(SRC1_IN_PLACE) +#endif // !defined(IN_PLACE) // Load data VEC_ACC_TYPE in1_data = CONVERT((VEC_DATA_TYPE(DATA_TYPE_IN1, VEC_SIZE_OUT))(VLOAD(VEC_SIZE_IN1)(0, (__global DATA_TYPE_IN1 *)in1_addr)), VEC_ACC_TYPE); diff --git a/src/core/CL/cl_kernels/pixelwise_mul_int.cl b/src/core/CL/cl_kernels/pixelwise_mul_int.cl index ac5cabcb8c..6d1c2d0c79 100644 --- a/src/core/CL/cl_kernels/pixelwise_mul_int.cl +++ b/src/core/CL/cl_kernels/pixelwise_mul_int.cl @@ -76,7 +76,9 @@ __kernel void pixelwise_mul_int( TENSOR3D_DECLARATION(in1), TENSOR3D_DECLARATION(in2), +#if !defined(IN_PLACE) TENSOR3D_DECLARATION(out), +#endif // !defined(IN_PLACE) const uint scale) { size_t x = max((int)(get_global_id(0) * VEC_SIZE_OUT - (VEC_SIZE_OUT - VEC_SIZE_LEFTOVER) % VEC_SIZE_OUT), 0); @@ -85,7 +87,16 @@ __kernel void pixelwise_mul_int( __global uchar *in1_addr = in1_ptr + in1_offset_first_element_in_bytes + x * in1_stride_x + y * in1_stride_y + z * in1_stride_z; __global uchar *in2_addr = in2_ptr + in2_offset_first_element_in_bytes + x * in2_stride_x + y * in2_stride_y + z * in2_stride_z; - __global uchar *out_addr = out_ptr + out_offset_first_element_in_bytes + x * out_stride_x + y * out_stride_y + z * out_stride_z; + __global uchar * +#if !defined(IN_PLACE) + out_addr = out_ptr + out_offset_first_element_in_bytes + x * out_stride_x + y * out_stride_y + z * out_stride_z; +#else // !defined(IN_PLACE) +#if defined(SRC1_IN_PLACE) + out_addr = in1_addr; +#else //defined(SRC1_IN_PLACE) + out_addr = in2_addr; +#endif //defined(SRC1_IN_PLACE) +#endif // !defined(IN_PLACE) // Load data VEC_ACC_TYPE in1_data = CONVERT((VEC_DATA_TYPE(DATA_TYPE_IN1, VEC_SIZE_OUT))VLOAD(VEC_SIZE_IN1)(0, (__global DATA_TYPE_IN1 *)in1_addr), VEC_ACC_TYPE); @@ -143,7 +154,9 @@ __kernel void pixelwise_mul_int( __kernel void pixelwise_mul_quantized( TENSOR3D_DECLARATION(in1), TENSOR3D_DECLARATION(in2), +#if !defined(IN_PLACE) TENSOR3D_DECLARATION(out), +#endif // !defined(IN_PLACE) const float scale) { size_t x = max((int)(get_global_id(0) * VEC_SIZE_OUT - (VEC_SIZE_OUT - VEC_SIZE_LEFTOVER) % VEC_SIZE_OUT), 0); @@ -152,7 +165,16 @@ __kernel void pixelwise_mul_quantized( __global uchar *in1_addr = in1_ptr + in1_offset_first_element_in_bytes + x * in1_stride_x + y * in1_stride_y + z * in1_stride_z; __global uchar *in2_addr = in2_ptr + in2_offset_first_element_in_bytes + x * in2_stride_x + y * in2_stride_y + z * in2_stride_z; - __global uchar *out_addr = out_ptr + out_offset_first_element_in_bytes + x * out_stride_x + y * out_stride_y + z * out_stride_z; + __global uchar * +#if !defined(IN_PLACE) + out_addr = out_ptr + out_offset_first_element_in_bytes + x * out_stride_x + y * out_stride_y + z * out_stride_z; +#else // !defined(IN_PLACE) +#if defined(SRC1_IN_PLACE) + out_addr = in1_addr; +#else //defined(SRC1_IN_PLACE) + out_addr = in2_addr; +#endif //defined(SRC1_IN_PLACE) +#endif // !defined(IN_PLACE) // Load data VEC_INT in_a = CONVERT((VEC_TYPE)(VLOAD(VEC_SIZE_IN1)(0, (__global DATA_TYPE_OUT *)in1_addr)), VEC_INT); diff --git a/src/core/gpu/cl/kernels/ClElementwiseKernel.cpp b/src/core/gpu/cl/kernels/ClElementwiseKernel.cpp index f005e9226e..3d9f0b6fcf 100644 --- a/src/core/gpu/cl/kernels/ClElementwiseKernel.cpp +++ b/src/core/gpu/cl/kernels/ClElementwiseKernel.cpp @@ -75,6 +75,21 @@ std::string generate_id_for_tuning_common(const std::string &kernel_name, const return config_id; } +Status validate_in_place_output_shape(const bool in_place, const bool src1_in_place, const ITensorInfo &src1, const ITensorInfo &src2, const ITensorInfo &dst, const TensorShape &out_shape) +{ + if(in_place) + { + ARM_COMPUTE_RETURN_ERROR_ON_MSG(detail::have_different_dimensions(out_shape, src1_in_place ? src1.tensor_shape() : src2.tensor_shape(), 0), + "Wrong shape for dst, cannot do in_place calculation"); + } + else + { + ARM_COMPUTE_RETURN_ERROR_ON_MSG(detail::have_different_dimensions(out_shape, dst.tensor_shape(), 0), + "Wrong shape for dst"); + } + return Status{}; +} + Status validate_arguments_with_float_only_supported_rules(const ITensorInfo &src1, const ITensorInfo &src2, const ITensorInfo &dst) { ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(&src1, &src2, &dst); @@ -82,6 +97,10 @@ Status validate_arguments_with_float_only_supported_rules(const ITensorInfo &src ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(&src1, 1, DataType::F16, DataType::F32); ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(&src1, &src2); + // Check whether it is in_place calculation + const bool in_place = (&src1 == &dst) || (&src2 == &dst); + const bool src1_in_place = in_place && (&src1 == &dst); + const TensorShape out_shape = TensorShape::broadcast_shape(src1.tensor_shape(), src2.tensor_shape()); ARM_COMPUTE_RETURN_ERROR_ON_MSG(out_shape.total_size() == 0, "Inputs are not broadcast compatible"); @@ -91,8 +110,7 @@ Status validate_arguments_with_float_only_supported_rules(const ITensorInfo &src { ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(&dst, 1, DataType::F16, DataType::F32); ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(&src1, &dst); - ARM_COMPUTE_RETURN_ERROR_ON_MSG(detail::have_different_dimensions(out_shape, dst.tensor_shape(), 0), - "Wrong shape for dst"); + ARM_COMPUTE_RETURN_ON_ERROR(validate_in_place_output_shape(in_place, src1_in_place, src1, src2, dst, out_shape)); } return Status{}; @@ -105,6 +123,10 @@ Status validate_arguments_divide_operation(const ITensorInfo *src1, const ITenso ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(src1, 1, DataType::F16, DataType::F32, DataType::S32); ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(src1, src2); + // Check whether it is in_place calculation + const bool in_place = (src1 == dst) || (src2 == dst); + const bool src1_in_place = in_place && (src1 == dst); + const TensorShape out_shape = TensorShape::broadcast_shape(src1->tensor_shape(), src2->tensor_shape()); ARM_COMPUTE_RETURN_ERROR_ON_MSG(out_shape.total_size() == 0, "Inputs are not broadcast compatible"); @@ -114,8 +136,7 @@ Status validate_arguments_divide_operation(const ITensorInfo *src1, const ITenso { ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(dst, 1, DataType::F16, DataType::F32, DataType::S32); ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(src1, dst); - ARM_COMPUTE_RETURN_ERROR_ON_MSG(detail::have_different_dimensions(out_shape, dst->tensor_shape(), 0), - "Wrong shape for dst"); + ARM_COMPUTE_RETURN_ON_ERROR(validate_in_place_output_shape(in_place, src1_in_place, *src1, *src2, *dst, out_shape)); } return Status{}; @@ -137,6 +158,10 @@ Status validate_arguments_with_arithmetic_rules(const ITensorInfo &src1, const I ARM_COMPUTE_RETURN_ERROR_ON_MSG(in2_offset != 0, "For quantized symmetric, offset must be zero"); } + // Check whether it is in_place calculation + const bool in_place = (&src1 == &dst) || (&src2 == &dst); + const bool src1_in_place = in_place && (&src1 == &dst); + const TensorShape out_shape = TensorShape::broadcast_shape(src1.tensor_shape(), src2.tensor_shape()); ARM_COMPUTE_RETURN_ERROR_ON_MSG(out_shape.total_size() == 0, "Inputs are not broadcast compatible"); @@ -145,6 +170,7 @@ Status validate_arguments_with_arithmetic_rules(const ITensorInfo &src1, const I { ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(&src1, &dst); ARM_COMPUTE_RETURN_ERROR_ON_MSG(detail::have_different_dimensions(out_shape, dst.tensor_shape(), 0), "Wrong shape for dst"); + ARM_COMPUTE_RETURN_ON_ERROR(validate_in_place_output_shape(in_place, src1_in_place, src1, src2, dst, out_shape)); if(is_data_type_quantized_symmetric(dst.data_type())) { @@ -182,6 +208,12 @@ CLBuildOptions generate_build_options_with_arithmetic_rules(const ITensorInfo &s } build_opts.add_option_if(src1.data_type() == DataType::S32, "-DS32"); + // Check whether it is in_place calculation + const bool in_place = (&src1 == &dst) || (&src2 == &dst); + const bool src1_in_place = in_place && (&src1 == &dst); + build_opts.add_option_if(in_place, "-DIN_PLACE"); + build_opts.add_option_if(src1_in_place, "-DSRC1_IN_PLACE"); + return build_opts; } @@ -267,6 +299,8 @@ void ClElementwiseKernel::run_op(ITensorPack &tensors, const Window &window, ::c const auto src_1 = utils::cast::polymorphic_downcast(tensors.get_const_tensor(TensorType::ACL_SRC_1)); auto dst = utils::cast::polymorphic_downcast(tensors.get_tensor(TensorType::ACL_DST)); + ARM_COMPUTE_ERROR_ON_NULLPTR(src_0, src_1, dst); + const TensorShape &in_shape1 = src_0->info()->tensor_shape(); const TensorShape &in_shape2 = src_1->info()->tensor_shape(); const TensorShape &out_shape = dst->info()->tensor_shape(); @@ -291,12 +325,18 @@ void ClElementwiseKernel::run_op(ITensorPack &tensors, const Window &window, ::c Window slice = collapsed.first_slice_window_3D(); Window slice_src1 = slice.broadcast_if_dimension_le_one(in_shape1_collapsed); Window slice_src2 = slice.broadcast_if_dimension_le_one(in_shape2_collapsed); + + // Check whether it is in_place calculation + const bool in_place = (src_0 == dst) || (src_1 == dst); do { unsigned int idx = 0; add_3D_tensor_argument(idx, src_0, slice_src1); add_3D_tensor_argument(idx, src_1, slice_src2); - add_3D_tensor_argument(idx, dst, slice); + if(!in_place) + { + add_3D_tensor_argument(idx, dst, slice); + } enqueue(queue, *this, slice, lws_hint()); ARM_COMPUTE_UNUSED(collapsed.slide_window_slice_3D(slice_src1)); diff --git a/src/core/gpu/cl/kernels/ClElementwiseKernel.h b/src/core/gpu/cl/kernels/ClElementwiseKernel.h index ab5c777ae6..4525cec55b 100644 --- a/src/core/gpu/cl/kernels/ClElementwiseKernel.h +++ b/src/core/gpu/cl/kernels/ClElementwiseKernel.h @@ -40,6 +40,8 @@ namespace kernels * Element-wise operation is computed by: * @f[ dst(x,y) = OP(src1(x,y), src2(x,y))@f] * + * For binary elementwise ops in-place cannot be enabled by passing nullptr to dst, it can only be enabled by passing either src1 or src2 to dst instead. + * */ class ClElementwiseKernel : public IClKernel { diff --git a/src/core/gpu/cl/kernels/ClMulKernel.cpp b/src/core/gpu/cl/kernels/ClMulKernel.cpp index 65f3bec099..7c4dddc20e 100644 --- a/src/core/gpu/cl/kernels/ClMulKernel.cpp +++ b/src/core/gpu/cl/kernels/ClMulKernel.cpp @@ -63,6 +63,10 @@ Status validate_arguments(const ITensorInfo *src1, const ITensorInfo *src2, cons ARM_COMPUTE_RETURN_ERROR_ON_MSG(scale < 0, "Scale cannot be negative."); ARM_COMPUTE_RETURN_ERROR_ON(act_info.enabled() && !is_data_type_float(dst->data_type())); + // Check whether it is in_place calculation + const bool in_place = (src1 == dst) || (src2 == dst); + const bool src1_in_place = in_place && (src1 == dst); + const TensorShape &out_shape = TensorShape::broadcast_shape(src1->tensor_shape(), src2->tensor_shape()); ARM_COMPUTE_RETURN_ERROR_ON_MSG(out_shape.total_size() == 0, "Inputs are not broadcast compatible"); @@ -85,7 +89,16 @@ Status validate_arguments(const ITensorInfo *src1, const ITensorInfo *src2, cons "Dst can only be QSYMM16 if both src are QSYMM16"); ARM_COMPUTE_RETURN_ERROR_ON_MSG((src1->data_type() == DataType::S32 || src2->data_type() == DataType::S32) && (dst->data_type() != DataType::S32), "Dst must be S32 if source tensors are S32"); - ARM_COMPUTE_RETURN_ERROR_ON_MSG(detail::have_different_dimensions(out_shape, dst->tensor_shape(), 0), "Wrong shape for dst"); + if(in_place) + { + ARM_COMPUTE_RETURN_ERROR_ON_MSG(detail::have_different_dimensions(out_shape, src1_in_place ? src1->tensor_shape() : src2->tensor_shape(), 0), + "Wrong shape for dst, cannot do in_place calculation"); + } + else + { + ARM_COMPUTE_RETURN_ERROR_ON_MSG(detail::have_different_dimensions(out_shape, dst->tensor_shape(), 0), + "Wrong shape for dst"); + } } return Status{}; @@ -194,11 +207,17 @@ void ClMulKernel::configure(const CLCompileContext &compile_context, ITensorInfo } } + // Check whether it is in_place calculation + const bool in_place = (src1 == dst) || (src2 == dst); + const bool src1_in_place = in_place && (src1 == dst); + build_opts.add_option_if(in_place, "-DIN_PLACE"); + build_opts.add_option_if(src1_in_place, "-DSRC1_IN_PLACE"); + // Create kernel _kernel = create_kernel(compile_context, kernel_name, build_opts.options()); // Set scale argument - unsigned int idx = 3 * num_arguments_per_3D_tensor(); // Skip the src and dst parameters + unsigned int idx = (in_place ? 2 : 3) * num_arguments_per_3D_tensor(); // Skip the src and dst parameters if(scale_int >= 0 && !is_quantized) { @@ -256,6 +275,8 @@ void ClMulKernel::run_op(ITensorPack &tensors, const Window &window, cl::Command const auto src_1 = utils::cast::polymorphic_downcast(tensors.get_const_tensor(TensorType::ACL_SRC_1)); auto dst = utils::cast::polymorphic_downcast(tensors.get_tensor(TensorType::ACL_DST)); + ARM_COMPUTE_ERROR_ON_NULLPTR(src_0, src_1, dst); + const TensorShape &in_shape1 = src_0->info()->tensor_shape(); const TensorShape &in_shape2 = src_1->info()->tensor_shape(); const TensorShape &out_shape = dst->info()->tensor_shape(); @@ -280,12 +301,17 @@ void ClMulKernel::run_op(ITensorPack &tensors, const Window &window, cl::Command Window slice_input1 = slice.broadcast_if_dimension_le_one(in_shape1_collapsed); Window slice_input2 = slice.broadcast_if_dimension_le_one(in_shape2_collapsed); + // Check whether it is in_place calculation + const bool in_place = (src_0 == dst) || (src_1 == dst); do { unsigned int idx = 0; add_3D_tensor_argument(idx, src_0, slice_input1); add_3D_tensor_argument(idx, src_1, slice_input2); - add_3D_tensor_argument(idx, dst, slice); + if(!in_place) + { + add_3D_tensor_argument(idx, dst, slice); + } enqueue(queue, *this, slice, lws_hint()); ARM_COMPUTE_UNUSED(collapsed.slide_window_slice_3D(slice_input1)); diff --git a/src/core/gpu/cl/kernels/ClMulKernel.h b/src/core/gpu/cl/kernels/ClMulKernel.h index dec8dba61c..2ee182b932 100644 --- a/src/core/gpu/cl/kernels/ClMulKernel.h +++ b/src/core/gpu/cl/kernels/ClMulKernel.h @@ -34,7 +34,11 @@ namespace opencl { namespace kernels { -/** Interface for the pixelwise multiplication kernel. */ +/** Interface for the pixelwise multiplication kernel. + * + * For binary elementwise ops in-place cannot be enabled by passing nullptr to dst, it can only be enabled by passing either src1 or src2 to dst instead. + * +*/ class ClMulKernel : public IClKernel { public: diff --git a/src/graph/mutators/InPlaceOperationMutator.cpp b/src/graph/mutators/InPlaceOperationMutator.cpp index 86236e8854..d3ea940895 100644 --- a/src/graph/mutators/InPlaceOperationMutator.cpp +++ b/src/graph/mutators/InPlaceOperationMutator.cpp @@ -180,9 +180,11 @@ void try_in_place_elementwise(std::unique_ptr &node) ARM_COMPUTE_ERROR_ON(current_output_tensor == nullptr); const auto qinfo_out = current_output_tensor->desc().quant_info; - // Can do in place, if the input has same shape as output, has same quntisation info as output, and input doesn't have accessor. - bool input0_can_in_place = !arm_compute::detail::have_different_dimensions(out_shape, shape0, 0) && (qinfo0 == qinfo_out) && (input0_tensor->accessor() == nullptr); - bool input1_can_in_place = !arm_compute::detail::have_different_dimensions(out_shape, shape1, 0) && (qinfo1 == qinfo_out) && (input1_tensor->accessor() == nullptr); + // Can do in place, if the input has same shape as output, has same quntisation info as output, has same data type as output and input doesn't have accessor. + bool input0_can_in_place = !arm_compute::detail::have_different_dimensions(out_shape, shape0, 0) && (qinfo0 == qinfo_out) + && (input0_tensor->desc().data_type == current_output_tensor->desc().data_type) && (input0_tensor->accessor() == nullptr); + bool input1_can_in_place = !arm_compute::detail::have_different_dimensions(out_shape, shape1, 0) && (qinfo1 == qinfo_out) + && (input1_tensor->desc().data_type == current_output_tensor->desc().data_type) && (input1_tensor->accessor() == nullptr); if(input0_can_in_place) { diff --git a/tests/datasets/ShapeDatasets.h b/tests/datasets/ShapeDatasets.h index 37c5f1626d..f3ef2c2985 100644 --- a/tests/datasets/ShapeDatasets.h +++ b/tests/datasets/ShapeDatasets.h @@ -211,6 +211,25 @@ public: { } }; +/** Data set containing pairs of tiny tensor shapes that are broadcast compatible and can do in_place calculation. */ +class TinyShapesBroadcastInplace final : public framework::dataset::ZipDataset +{ +public: + TinyShapesBroadcastInplace() + : ZipDataset( + ShapeDataset("Shape0", + { + TensorShape{ 9U }, + TensorShape{ 10U, 2U, 14U, 2U }, + }), + ShapeDataset("Shape1", + { + TensorShape{ 9U, 1U, 9U }, + TensorShape{ 10U }, + })) + { + } +}; /** Data set containing pairs of small tensor shapes that are broadcast compatible. */ class SmallShapesBroadcast final : public framework::dataset::ZipDataset { diff --git a/tests/validation/CL/ArithmeticAddition.cpp b/tests/validation/CL/ArithmeticAddition.cpp index 9e3d9afc36..45632dc7e2 100644 --- a/tests/validation/CL/ArithmeticAddition.cpp +++ b/tests/validation/CL/ArithmeticAddition.cpp @@ -51,6 +51,8 @@ const auto ActivationFunctionsDataset = framework::dataset::make("ActivationInfo ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::BOUNDED_RELU, 0.75f, 0.25f), ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::LOGISTIC, 0.75f, 0.25f) }); +const auto InPlaceDataSet = framework::dataset::make("InPlace", { false, true }); +const auto OutOfPlaceDataSet = framework::dataset::make("InPlace", { false }); } // namespace TEST_SUITE(CL) @@ -109,8 +111,10 @@ using CLArithmeticAdditionFixture = ArithmeticAdditionValidationFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::U8)), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP }))) +FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticAdditionFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", + DataType::U8)), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); @@ -118,15 +122,19 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticAdditionFixture, framework TEST_SUITE_END() // U8 TEST_SUITE(S16) -FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticAdditionFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::S16)), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP }))) +FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticAdditionFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", + DataType::S16)), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); } -FIXTURE_DATA_TEST_CASE(RunLarge, CLArithmeticAdditionFixture, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapes(), framework::dataset::make("DataType", DataType::S16)), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP }))) +FIXTURE_DATA_TEST_CASE(RunLarge, CLArithmeticAdditionFixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::LargeShapes(), framework::dataset::make("DataType", + DataType::S16)), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); @@ -139,12 +147,13 @@ using CLArithmeticAdditionQuantizedFixture = ArithmeticAdditionValidationQuantiz TEST_SUITE(Quantized) TEST_SUITE(QASYMM8) -FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticAdditionQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(datasets::SmallShapes(), +FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticAdditionQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::QASYMM8)), framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), framework::dataset::make("Src0QInfo", { QuantizationInfo(5.f / 255.f, 20) })), framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 255.f, 10) })), - framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 255.f, 5) }))) + framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 255.f, 5) })), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); @@ -152,36 +161,51 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticAdditionQuantizedFixture, template using CLArithmeticAdditionBroadcastQuantizedFixture = ArithmeticAdditionValidationQuantizedBroadcastFixture; FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLArithmeticAdditionBroadcastQuantizedFixture, framework::DatasetMode::PRECOMMIT, - combine(combine(combine(combine(combine(datasets::SmallShapesBroadcast(), - framework::dataset::make("DataType", DataType::QASYMM8)), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), - framework::dataset::make("Src0QInfo", { QuantizationInfo(5.f / 255.f, 20) })), - framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 255.f, 10) })), - framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 255.f, 5) }))) + combine(combine(combine(combine(combine(combine(datasets::SmallShapesBroadcast(), + framework::dataset::make("DataType", DataType::QASYMM8)), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), + framework::dataset::make("Src0QInfo", { QuantizationInfo(5.f / 255.f, 20) })), + framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 255.f, 10) })), + framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 255.f, 5) })), + OutOfPlaceDataSet)) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +FIXTURE_DATA_TEST_CASE(RunTinyBroadcastInPlace, CLArithmeticAdditionBroadcastQuantizedFixture, framework::DatasetMode::PRECOMMIT, + combine(combine(combine(combine(combine(combine(datasets::TinyShapesBroadcastInplace(), + framework::dataset::make("DataType", DataType::QASYMM8)), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), + framework::dataset::make("Src0QInfo", { QuantizationInfo(1.f / 255.f, 10) })), + framework::dataset::make("Src1QInfo", { QuantizationInfo(1.f / 255.f, 10) })), + framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 255.f, 10) })), + InPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); } TEST_SUITE_END() // QASYMM8 TEST_SUITE(QASYMM8_SIGNED) -FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticAdditionQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(datasets::SmallShapes(), +FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticAdditionQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::QASYMM8_SIGNED)), framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), framework::dataset::make("Src0QInfo", { QuantizationInfo(5.f / 255.f, 10) })), framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 255.f, 10) })), - framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 255.f, 5) }))) + framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 255.f, 5) })), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); } TEST_SUITE_END() // QASYMM8_SIGNED TEST_SUITE(QSYMM16) -FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticAdditionQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(datasets::SmallShapes(), +FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticAdditionQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::QSYMM16)), framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), framework::dataset::make("Src0QInfo", { QuantizationInfo(1.f / 32768.f, 0), QuantizationInfo(5.f / 32768.f, 0) })), framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 32768.f, 0), QuantizationInfo(5.f / 32768.f, 0) })), - framework::dataset::make("OutQInfo", { QuantizationInfo(5.f / 32768.f, 0) }))) + framework::dataset::make("OutQInfo", { QuantizationInfo(5.f / 32768.f, 0) })), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); @@ -194,18 +218,21 @@ using CLArithmeticAdditionFloatFixture = ArithmeticAdditionValidationFloatFixtur TEST_SUITE(Float) TEST_SUITE(FP16) -FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticAdditionFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", +FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticAdditionFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::F16)), framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), - EmptyActivationFunctionsDataset)) + EmptyActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); } -FIXTURE_DATA_TEST_CASE(RunWithActivation, CLArithmeticAdditionFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::TinyShapes(), framework::dataset::make("DataType", - DataType::F16)), +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLArithmeticAdditionFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(combine(datasets::TinyShapes(), + framework::dataset::make("DataType", + DataType::F16)), framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), - ActivationFunctionsDataset)) + ActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); @@ -213,27 +240,32 @@ FIXTURE_DATA_TEST_CASE(RunWithActivation, CLArithmeticAdditionFloatFixture TEST_SUITE_END() // FP16 TEST_SUITE(FP32) -FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticAdditionFloatFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", - DataType::F32)), +FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticAdditionFloatFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(), + framework::dataset::make("DataType", + DataType::F32)), framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), - EmptyActivationFunctionsDataset)) + EmptyActivationFunctionsDataset), + InPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); } -FIXTURE_DATA_TEST_CASE(RunWithActivation, CLArithmeticAdditionFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::TinyShapes(), framework::dataset::make("DataType", - DataType::F32)), +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLArithmeticAdditionFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(combine(datasets::TinyShapes(), + framework::dataset::make("DataType", + DataType::F32)), framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), - ActivationFunctionsDataset)) + ActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); } -FIXTURE_DATA_TEST_CASE(RunLarge, CLArithmeticAdditionFloatFixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::LargeShapes(), framework::dataset::make("DataType", +FIXTURE_DATA_TEST_CASE(RunLarge, CLArithmeticAdditionFloatFixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(combine(datasets::LargeShapes(), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), - EmptyActivationFunctionsDataset)) + EmptyActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); @@ -242,27 +274,30 @@ FIXTURE_DATA_TEST_CASE(RunLarge, CLArithmeticAdditionFloatFixture, framew template using CLArithmeticAdditionBroadcastFloatFixture = ArithmeticAdditionBroadcastValidationFloatFixture; -FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLArithmeticAdditionBroadcastFloatFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallShapesBroadcast(), +FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLArithmeticAdditionBroadcastFloatFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapesBroadcast(), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), - EmptyActivationFunctionsDataset)) + EmptyActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); } -FIXTURE_DATA_TEST_CASE(RunWithActivationBroadcast, CLArithmeticAdditionBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::TinyShapesBroadcast(), +FIXTURE_DATA_TEST_CASE(RunWithActivationBroadcast, CLArithmeticAdditionBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(combine(datasets::TinyShapesBroadcast(), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), - ActivationFunctionsDataset)) + ActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); } -FIXTURE_DATA_TEST_CASE(RunLargeBroadcast, CLArithmeticAdditionBroadcastFloatFixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::LargeShapesBroadcast(), +FIXTURE_DATA_TEST_CASE(RunLargeBroadcast, CLArithmeticAdditionBroadcastFloatFixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(combine(datasets::LargeShapesBroadcast(), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), - EmptyActivationFunctionsDataset)) + EmptyActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); diff --git a/tests/validation/CL/ArithmeticDivision.cpp b/tests/validation/CL/ArithmeticDivision.cpp index 9dcdfb83e1..94bacba7e5 100644 --- a/tests/validation/CL/ArithmeticDivision.cpp +++ b/tests/validation/CL/ArithmeticDivision.cpp @@ -59,6 +59,8 @@ const auto ActivationFunctionsDataset = framework::dataset::make("Activatio ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::BOUNDED_RELU, 0.75f, 0.25f), ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::LOGISTIC, 0.75f, 0.25f) }); +const auto InPlaceDataSet = framework::dataset::make("InPlace", { false, true }); +const auto OutOfPlaceDataSet = framework::dataset::make("InPlace", { false }); } // namespace TEST_SUITE(CL) @@ -95,14 +97,16 @@ using CLArithmeticDivisionIntegerFixture = ArithmeticDivisionValidationIntegerFi TEST_SUITE(Integer) TEST_SUITE(S32) -FIXTURE_DATA_TEST_CASE(RunSmallInteger, CLArithmeticDivisionIntegerFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), ArithmeticDivisionS32Dataset), - EmptyActivationFunctionsDataset)) +FIXTURE_DATA_TEST_CASE(RunSmallInteger, CLArithmeticDivisionIntegerFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallShapes(), ArithmeticDivisionS32Dataset), + EmptyActivationFunctionsDataset), + InPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); } -FIXTURE_DATA_TEST_CASE(RunIntegerWithActivation, CLArithmeticDivisionIntegerFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ArithmeticDivisionS32Dataset), - ActivationFunctionsDataset)) +FIXTURE_DATA_TEST_CASE(RunIntegerWithActivation, CLArithmeticDivisionIntegerFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::SmallShapes(), ArithmeticDivisionS32Dataset), + ActivationFunctionsDataset), + InPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); @@ -116,14 +120,16 @@ using CLArithmeticDivisionFloatFixture = ArithmeticDivisionValidationFloatFixtur TEST_SUITE(Float) TEST_SUITE(FP16) -FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticDivisionFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ArithmeticDivisionFP16Dataset), - EmptyActivationFunctionsDataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticDivisionFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::SmallShapes(), ArithmeticDivisionFP16Dataset), + EmptyActivationFunctionsDataset), + InPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp16, 0.01); } -FIXTURE_DATA_TEST_CASE(RunWithActivation, CLArithmeticDivisionFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapes(), ArithmeticDivisionFP16Dataset), - ActivationFunctionsDataset)) +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLArithmeticDivisionFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::TinyShapes(), ArithmeticDivisionFP16Dataset), + ActivationFunctionsDataset), + InPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp16, 0.01); @@ -131,21 +137,24 @@ FIXTURE_DATA_TEST_CASE(RunWithActivation, CLArithmeticDivisionFloatFixture TEST_SUITE_END() TEST_SUITE(FP32) -FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticDivisionFloatFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), ArithmeticDivisionFP32Dataset), - EmptyActivationFunctionsDataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticDivisionFloatFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallShapes(), ArithmeticDivisionFP32Dataset), + EmptyActivationFunctionsDataset), + InPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); } -FIXTURE_DATA_TEST_CASE(RunWithActivation, CLArithmeticDivisionFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapes(), ArithmeticDivisionFP32Dataset), - ActivationFunctionsDataset)) +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLArithmeticDivisionFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::TinyShapes(), ArithmeticDivisionFP32Dataset), + ActivationFunctionsDataset), + InPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); } -FIXTURE_DATA_TEST_CASE(RunLarge, CLArithmeticDivisionFloatFixture, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapes(), ArithmeticDivisionFP32Dataset), - EmptyActivationFunctionsDataset)) +FIXTURE_DATA_TEST_CASE(RunLarge, CLArithmeticDivisionFloatFixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::LargeShapes(), ArithmeticDivisionFP32Dataset), + EmptyActivationFunctionsDataset), + InPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); @@ -154,24 +163,27 @@ FIXTURE_DATA_TEST_CASE(RunLarge, CLArithmeticDivisionFloatFixture, framew template using CLArithmeticDivisionBroadcastFloatFixture = ArithmeticDivisionBroadcastValidationFloatFixture; -FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLArithmeticDivisionBroadcastFloatFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapesBroadcast(), +FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLArithmeticDivisionBroadcastFloatFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallShapesBroadcast(), ArithmeticDivisionFP32Dataset), - EmptyActivationFunctionsDataset)) + EmptyActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); } -FIXTURE_DATA_TEST_CASE(RunWithActivationBroadcast, CLArithmeticDivisionBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapesBroadcast(), +FIXTURE_DATA_TEST_CASE(RunWithActivationBroadcast, CLArithmeticDivisionBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::TinyShapesBroadcast(), ArithmeticDivisionFP32Dataset), - ActivationFunctionsDataset)) + ActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); } -FIXTURE_DATA_TEST_CASE(RunLargeBroadcast, CLArithmeticDivisionBroadcastFloatFixture, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapesBroadcast(), +FIXTURE_DATA_TEST_CASE(RunLargeBroadcast, CLArithmeticDivisionBroadcastFloatFixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::LargeShapesBroadcast(), ArithmeticDivisionFP32Dataset), - EmptyActivationFunctionsDataset)) + EmptyActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); diff --git a/tests/validation/CL/ArithmeticSubtraction.cpp b/tests/validation/CL/ArithmeticSubtraction.cpp index 00eba7f92a..6a82471dfa 100644 --- a/tests/validation/CL/ArithmeticSubtraction.cpp +++ b/tests/validation/CL/ArithmeticSubtraction.cpp @@ -180,6 +180,17 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticSubtractionQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(combine(datasets::TinyShapes(), + framework::dataset::make("DataType", DataType::QASYMM8)), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), + framework::dataset::make("Src0QInfo", { QuantizationInfo(5.f / 255.f, 20) })), + framework::dataset::make("Src1QInfo", { QuantizationInfo(5.f / 255.f, 20) })), + framework::dataset::make("OutQInfo", { QuantizationInfo(5.f / 255.f, 20) })), InPlaceDataSet)) { // Validate output @@ -193,7 +204,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticSubtractionQuantizedFixture framework::dataset::make("Src0QInfo", { QuantizationInfo(5.f / 255.f, 10) })), framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 255.f, 10) })), framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 255.f, 5) })), - InPlaceDataSet)) + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); @@ -281,6 +292,16 @@ FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLArithmeticSubtractionBroadcastFloatF // Validate output validate(CLAccessor(_target), _reference); } +FIXTURE_DATA_TEST_CASE(RunTinyBroadcastInplace, CLArithmeticSubtractionBroadcastFloatFixture, framework::DatasetMode::PRECOMMIT, + combine(combine(combine(combine(datasets::TinyShapesBroadcastInplace(), + framework::dataset::make("DataType", DataType::F32)), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + EmptyActivationFunctionsDataset), + InPlaceDataSet)) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} FIXTURE_DATA_TEST_CASE(RunWithActivationBroadcast, CLArithmeticSubtractionBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(combine(datasets::TinyShapesBroadcast(), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), diff --git a/tests/validation/CL/ElementwiseMax.cpp b/tests/validation/CL/ElementwiseMax.cpp index 225968efd1..bd47c23256 100644 --- a/tests/validation/CL/ElementwiseMax.cpp +++ b/tests/validation/CL/ElementwiseMax.cpp @@ -71,6 +71,8 @@ const auto ActivationFunctionsDataset = framework::dataset::make("ActivationInfo ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::BOUNDED_RELU, 0.75f, 0.25f), ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::LOGISTIC, 0.75f, 0.25f) }); +const auto InPlaceDataSet = framework::dataset::make("InPlace", { false, true }); +const auto OutOfPlaceDataSet = framework::dataset::make("InPlace", { false }); } // namespace TEST_SUITE(CL) @@ -104,7 +106,8 @@ using CLElementwiseMaxFixture = ElementwiseMaxValidationFixture, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallShapes(), ElementwiseMaxU8Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMaxFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), ElementwiseMaxU8Dataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); @@ -112,7 +115,8 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMaxFixture, framework::Da TEST_SUITE_END() TEST_SUITE(S16) -FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMaxFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwiseMaxS16Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMaxFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseMaxS16Dataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); @@ -125,33 +129,36 @@ using CLElementwiseMaxQuantizedFixture = ElementwiseMaxValidationQuantizedFixtur TEST_SUITE(Quantized) TEST_SUITE(QASYMM8) -FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMaxQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(), +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMaxQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(datasets::SmallShapes(), ElementwiseMaxQASYMM8Dataset), framework::dataset::make("Src0QInfo", { QuantizationInfo(5.f / 255.f, 20) })), framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 255.f, 10) })), - framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 255.f, 5) }))) + framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 255.f, 5) })), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32, 0.01); } TEST_SUITE_END() TEST_SUITE(QASYMM8_SIGNED) -FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMaxQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(), +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMaxQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(datasets::SmallShapes(), ElementwiseMaxQASYMM8SignedDataset), framework::dataset::make("Src0QInfo", { QuantizationInfo(5.f / 255.f, 20) })), framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 255.f, 10) })), - framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 255.f, 5) }))) + framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 255.f, 5) })), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); } TEST_SUITE_END() TEST_SUITE(QSYMM16) -FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMaxQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(), +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMaxQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(datasets::SmallShapes(), ElementwiseMaxQSYMM16Dataset), framework::dataset::make("Src0QInfo", { QuantizationInfo(1.f / 32768.f, 0), QuantizationInfo(5.f / 32768.f, 0) })), framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 32768.f, 0), QuantizationInfo(5.f / 32768.f, 0) })), - framework::dataset::make("OutQInfo", { QuantizationInfo(5.f / 32768.f, 0) }))) + framework::dataset::make("OutQInfo", { QuantizationInfo(5.f / 32768.f, 0) })), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); @@ -164,13 +171,16 @@ using CLElementwiseMaxFloatFixture = ElementwiseMaxValidationFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseMaxFP16Dataset), EmptyActivationFunctionsDataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMaxFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::SmallShapes(), ElementwiseMaxFP16Dataset), + EmptyActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp16, 0.01); } -FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwiseMaxFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapes(), ElementwiseMaxFP16Dataset), - ActivationFunctionsDataset)) +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwiseMaxFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::TinyShapes(), ElementwiseMaxFP16Dataset), + ActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp16, 0.01); @@ -178,14 +188,16 @@ FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwiseMaxFloatFixture, fr TEST_SUITE_END() TEST_SUITE(FP32) -FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMaxFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseMaxFP32Dataset), - EmptyActivationFunctionsDataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMaxFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::SmallShapes(), ElementwiseMaxFP32Dataset), + EmptyActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); } -FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwiseMaxFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapes(), ElementwiseMaxFP32Dataset), - ActivationFunctionsDataset)) +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwiseMaxFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::TinyShapes(), ElementwiseMaxFP32Dataset), + ActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); @@ -194,16 +206,18 @@ FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwiseMaxFloatFixture, f template using CLElementwiseMaxBroadcastFloatFixture = ElementwiseMaxBroadcastValidationFloatFixture; -FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLElementwiseMaxBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapesBroadcast(), +FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLElementwiseMaxBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::SmallShapesBroadcast(), ElementwiseMaxFP32Dataset), - EmptyActivationFunctionsDataset)) + EmptyActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); } -FIXTURE_DATA_TEST_CASE(RunWithActivationBroadcast, CLElementwiseMaxBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapesBroadcast(), +FIXTURE_DATA_TEST_CASE(RunWithActivationBroadcast, CLElementwiseMaxBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::TinyShapesBroadcast(), ElementwiseMaxFP32Dataset), - ActivationFunctionsDataset)) + ActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); diff --git a/tests/validation/CL/ElementwiseMin.cpp b/tests/validation/CL/ElementwiseMin.cpp index 2a066908fa..ee229a0941 100644 --- a/tests/validation/CL/ElementwiseMin.cpp +++ b/tests/validation/CL/ElementwiseMin.cpp @@ -71,6 +71,8 @@ const auto ActivationFunctionsDataset = framework::dataset::make("ActivationInfo ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::BOUNDED_RELU, 0.75f, 0.25f), ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::LOGISTIC, 0.75f, 0.25f) }); +const auto InPlaceDataSet = framework::dataset::make("InPlace", { false, true }); +const auto OutOfPlaceDataSet = framework::dataset::make("InPlace", { false }); } // namespace TEST_SUITE(CL) @@ -104,7 +106,8 @@ using CLElementwiseMinFixture = ElementwiseMinValidationFixture, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallShapes(), ElementwiseMinU8Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMinFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), ElementwiseMinU8Dataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); @@ -112,7 +115,8 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMinFixture, framework::Da TEST_SUITE_END() TEST_SUITE(S16) -FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMinFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwiseMinS16Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMinFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseMinS16Dataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); @@ -125,33 +129,36 @@ using CLElementwiseMinQuantizedFixture = ElementwiseMinValidationQuantizedFixtur TEST_SUITE(Quantized) TEST_SUITE(QASYMM8) -FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMinQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(), +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMinQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(datasets::SmallShapes(), ElementwiseMinQASYMM8Dataset), framework::dataset::make("Src0QInfo", { QuantizationInfo(5.f / 255.f, 20) })), framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 255.f, 10) })), - framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 255.f, 5) }))) + framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 255.f, 5) })), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32, 0.01); } TEST_SUITE_END() TEST_SUITE(QASYMM8_SIGNED) -FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMinQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(), +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMinQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(datasets::SmallShapes(), ElementwiseMinQASYMM8SignedDataset), framework::dataset::make("Src0QInfo", { QuantizationInfo(5.f / 255.f, 20) })), framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 255.f, 10) })), - framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 255.f, 5) }))) + framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 255.f, 5) })), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); } TEST_SUITE_END() TEST_SUITE(QSYMM16) -FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMinQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(), +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMinQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(datasets::SmallShapes(), ElementwiseMinQSYMM16Dataset), framework::dataset::make("SrcQInfo0", { QuantizationInfo(1.f / 32768.f, 0), QuantizationInfo(5.f / 32768.f, 0) })), framework::dataset::make("SrcQInfo1", { QuantizationInfo(2.f / 32768.f, 0), QuantizationInfo(5.f / 32768.f, 0) })), - framework::dataset::make("OutQInfo", { QuantizationInfo(5.f / 32768.f, 0) }))) + framework::dataset::make("OutQInfo", { QuantizationInfo(5.f / 32768.f, 0) })), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); @@ -164,13 +171,16 @@ using CLElementwiseMinFloatFixture = ElementwiseMinValidationFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseMinFP16Dataset), EmptyActivationFunctionsDataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMinFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::SmallShapes(), ElementwiseMinFP16Dataset), + EmptyActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp16, 0.01); } -FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwiseMinFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapes(), ElementwiseMinFP16Dataset), - ActivationFunctionsDataset)) +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwiseMinFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::TinyShapes(), ElementwiseMinFP16Dataset), + ActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp16, 0.01); @@ -178,14 +188,16 @@ FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwiseMinFloatFixture, fr TEST_SUITE_END() TEST_SUITE(FP32) -FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMinFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseMinFP32Dataset), - EmptyActivationFunctionsDataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMinFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::SmallShapes(), ElementwiseMinFP32Dataset), + EmptyActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); } -FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwiseMinFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapes(), ElementwiseMinFP32Dataset), - ActivationFunctionsDataset)) +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwiseMinFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::TinyShapes(), ElementwiseMinFP32Dataset), + ActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); @@ -193,16 +205,18 @@ FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwiseMinFloatFixture, f template using CLElementwiseMinBroadcastFloatFixture = ElementwiseMinBroadcastValidationFloatFixture; -FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLElementwiseMinBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapesBroadcast(), +FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLElementwiseMinBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::SmallShapesBroadcast(), ElementwiseMinFP32Dataset), - EmptyActivationFunctionsDataset)) + EmptyActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); } -FIXTURE_DATA_TEST_CASE(RunWithActivationBroadcast, CLElementwiseMinBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapesBroadcast(), +FIXTURE_DATA_TEST_CASE(RunWithActivationBroadcast, CLElementwiseMinBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::TinyShapesBroadcast(), ElementwiseMinFP32Dataset), - ActivationFunctionsDataset)) + ActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); diff --git a/tests/validation/CL/ElementwisePower.cpp b/tests/validation/CL/ElementwisePower.cpp index a2d3ba6c09..c2aeb6e045 100644 --- a/tests/validation/CL/ElementwisePower.cpp +++ b/tests/validation/CL/ElementwisePower.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020 Arm Limited. + * Copyright (c) 2019-2021 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -57,6 +57,8 @@ const auto ActivationFunctionsDataset = framework::dataset::make("ActivationInfo ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::BOUNDED_RELU, 0.75f, 0.25f), ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::LOGISTIC, 0.75f, 0.25f) }); +const auto InPlaceDataSet = framework::dataset::make("InPlace", { false, true }); +const auto OutOfPlaceDataSet = framework::dataset::make("InPlace", { false }); } // namespace TEST_SUITE(CL) @@ -96,29 +98,33 @@ using CLElementwisePowerBroadcastFloatFixture = ElementwisePowerBroadcastValidat TEST_SUITE(Float) TEST_SUITE(FP16) -FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwisePowerFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwisePowerFP16Dataset), - EmptyActivationFunctionsDataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwisePowerFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::SmallShapes(), ElementwisePowerFP16Dataset), + EmptyActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp16, 0.01); } -FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwisePowerFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapes(), ElementwisePowerFP16Dataset), - ActivationFunctionsDataset)) +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwisePowerFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::TinyShapes(), ElementwisePowerFP16Dataset), + ActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp16, 0.01); } -FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLElementwisePowerBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapesBroadcast(), +FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLElementwisePowerBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::SmallShapesBroadcast(), ElementwisePowerFP16Dataset), - EmptyActivationFunctionsDataset)) + EmptyActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp16, 0.01); } -FIXTURE_DATA_TEST_CASE(RunWithActivationBroadcast, CLElementwisePowerBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapesBroadcast(), +FIXTURE_DATA_TEST_CASE(RunWithActivationBroadcast, CLElementwisePowerBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::TinyShapesBroadcast(), ElementwisePowerFP16Dataset), - ActivationFunctionsDataset)) + ActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp16, 0.01); @@ -126,29 +132,33 @@ FIXTURE_DATA_TEST_CASE(RunWithActivationBroadcast, CLElementwisePowerBroadcastFl TEST_SUITE_END() //FP16 TEST_SUITE(FP32) -FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwisePowerFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwisePowerFP32Dataset), - EmptyActivationFunctionsDataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwisePowerFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::SmallShapes(), ElementwisePowerFP32Dataset), + EmptyActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); } -FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwisePowerFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapes(), ElementwisePowerFP32Dataset), - ActivationFunctionsDataset)) +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwisePowerFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::TinyShapes(), ElementwisePowerFP32Dataset), + ActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); } -FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLElementwisePowerBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapesBroadcast(), +FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLElementwisePowerBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::SmallShapesBroadcast(), ElementwisePowerFP32Dataset), - EmptyActivationFunctionsDataset)) + EmptyActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); } -FIXTURE_DATA_TEST_CASE(RunWithActivationBroadcast, CLElementwisePowerBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapesBroadcast(), +FIXTURE_DATA_TEST_CASE(RunWithActivationBroadcast, CLElementwisePowerBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::TinyShapesBroadcast(), ElementwisePowerFP32Dataset), - ActivationFunctionsDataset)) + ActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); diff --git a/tests/validation/CL/ElementwiseSquaredDiff.cpp b/tests/validation/CL/ElementwiseSquaredDiff.cpp index 4c732b0885..ee0279df33 100644 --- a/tests/validation/CL/ElementwiseSquaredDiff.cpp +++ b/tests/validation/CL/ElementwiseSquaredDiff.cpp @@ -70,6 +70,8 @@ const auto ActivationFunctionsDataset = framework::dataset::make("ActivationInfo ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::BOUNDED_RELU, 0.75f, 0.25f), ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::LOGISTIC, 0.75f, 0.25f) }); +const auto InPlaceDataSet = framework::dataset::make("InPlace", { false, true }); +const auto OutOfPlaceDataSet = framework::dataset::make("InPlace", { false }); } // namespace TEST_SUITE(CL) @@ -103,7 +105,8 @@ using CLElementwiseSquaredDiffFixture = ElementwiseSquaredDiffValidationFixture< TEST_SUITE(Integer) TEST_SUITE(U8) -FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseSquaredDiffFixture, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallShapes(), ElementwiseSquaredDiffU8Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseSquaredDiffFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), ElementwiseSquaredDiffU8Dataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); @@ -111,7 +114,8 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseSquaredDiffFixture, frame TEST_SUITE_END() TEST_SUITE(S16) -FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseSquaredDiffFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwiseSquaredDiffS16Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseSquaredDiffFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseSquaredDiffS16Dataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference); @@ -124,22 +128,24 @@ using CLElementwiseSquaredDiffQuantizedFixture = ElementwiseSquaredDiffValidatio TEST_SUITE(Quantized) TEST_SUITE(QASYMM8) -FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseSquaredDiffQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(), +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseSquaredDiffQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(datasets::SmallShapes(), ElementwiseSquaredDiffQASYMM8Dataset), framework::dataset::make("Src0QInfo", { QuantizationInfo(5.f / 255.f, 20) })), framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 255.f, 10) })), - framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 255.f, 5) }))) + framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 255.f, 5) })), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32, 0.01); } TEST_SUITE_END() TEST_SUITE(QSYMM16) -FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseSquaredDiffQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(), +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseSquaredDiffQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(datasets::SmallShapes(), ElementwiseSquaredDiffQSYMM16Dataset), framework::dataset::make("Src0QInfo", { QuantizationInfo(1.f / 32768.f, 0), QuantizationInfo(5.f / 32768.f, 0) })), framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 32768.f, 0), QuantizationInfo(5.f / 32768.f, 0) })), - framework::dataset::make("OutQInfo", { QuantizationInfo(5.f / 32768.f, 0) }))) + framework::dataset::make("OutQInfo", { QuantizationInfo(5.f / 32768.f, 0) })), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_qsymm16); @@ -152,14 +158,16 @@ using CLElementwiseSquaredDiffFloatFixture = ElementwiseSquaredDiffValidationFlo TEST_SUITE(Float) TEST_SUITE(FP16) -FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseSquaredDiffFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseSquaredDiffFP16Dataset), - EmptyActivationFunctionsDataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseSquaredDiffFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::SmallShapes(), ElementwiseSquaredDiffFP16Dataset), + EmptyActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp16, 0.01); } -FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwiseSquaredDiffFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapes(), ElementwiseSquaredDiffFP16Dataset), - ActivationFunctionsDataset)) +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwiseSquaredDiffFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::TinyShapes(), ElementwiseSquaredDiffFP16Dataset), + ActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp16, 0.01); @@ -167,14 +175,16 @@ FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwiseSquaredDiffFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseSquaredDiffFP32Dataset), - EmptyActivationFunctionsDataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseSquaredDiffFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::SmallShapes(), ElementwiseSquaredDiffFP32Dataset), + EmptyActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); } -FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwiseSquaredDiffFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapes(), ElementwiseSquaredDiffFP32Dataset), - ActivationFunctionsDataset)) +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwiseSquaredDiffFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::TinyShapes(), ElementwiseSquaredDiffFP32Dataset), + ActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); @@ -182,16 +192,18 @@ FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwiseSquaredDiffFloatFixture using CLElementwiseSquaredDiffBroadcastFloatFixture = ElementwiseSquaredDiffBroadcastValidationFloatFixture; -FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLElementwiseSquaredDiffBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapesBroadcast(), +FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLElementwiseSquaredDiffBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::SmallShapesBroadcast(), ElementwiseSquaredDiffFP32Dataset), - EmptyActivationFunctionsDataset)) + EmptyActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); } -FIXTURE_DATA_TEST_CASE(RunWithActivationBroadcast, CLElementwiseSquaredDiffBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapesBroadcast(), +FIXTURE_DATA_TEST_CASE(RunWithActivationBroadcast, CLElementwiseSquaredDiffBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::TinyShapesBroadcast(), ElementwiseSquaredDiffFP32Dataset), - ActivationFunctionsDataset)) + ActivationFunctionsDataset), + OutOfPlaceDataSet)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); diff --git a/tests/validation/CL/PixelWiseMultiplication.cpp b/tests/validation/CL/PixelWiseMultiplication.cpp index 9e0a6243d7..84aa2e7ee6 100644 --- a/tests/validation/CL/PixelWiseMultiplication.cpp +++ b/tests/validation/CL/PixelWiseMultiplication.cpp @@ -50,9 +50,6 @@ const auto ActivationFunctionsDataset = framework::dataset::make("ActivationInfo ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::BOUNDED_RELU, 0.75f, 0.25f), ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::LOGISTIC, 0.75f, 0.25f) }); -// Since in-place computation on CL-side hasn't been intended to be implemented, they are not tested. -// However, this dataset is required for the shared fixture and it would make extension easier when -// CL-side also starts supporting in-place computation. const auto InPlaceDataSet = framework::dataset::make("InPlace", { false }); } //namespace // *INDENT-OFF* @@ -124,7 +121,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLPixelWiseMultiplicationIntegerFixture, f datasets::SmallShapes(), framework::dataset::make("DataType1", DataType::S32)), framework::dataset::make("DataType2", DataType::S32)), - framework::dataset::make("Scale", {1.f})), + framework::dataset::make("Scale", { 1.f })), datasets::ConvertPolicies()), framework::dataset::make("RoundingPolicy", RoundingPolicy::TO_NEAREST_UP)), EmptyActivationFunctionsDataset), @@ -132,6 +129,18 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLPixelWiseMultiplicationIntegerFixture, f { validate(CLAccessor(_target), _reference); } +FIXTURE_DATA_TEST_CASE(RunInplace, CLPixelWiseMultiplicationIntegerFixture, framework::DatasetMode::PRECOMMIT, + combine(combine(combine(combine(combine(combine(combine(datasets::TinyShapes(), + framework::dataset::make("DataType1", DataType::S32)), + framework::dataset::make("DataType2", DataType::S32)), + framework::dataset::make("Scale", { 1.f })), + datasets::ConvertPolicies()), + framework::dataset::make("RoundingPolicy", RoundingPolicy::TO_NEAREST_UP)), + EmptyActivationFunctionsDataset), + framework::dataset::make("InPlace", { true }))) +{ + validate(CLAccessor(_target), _reference); +} TEST_SUITE_END() TEST_SUITE(F16toF16) @@ -147,6 +156,19 @@ TEST_SUITE(F32toF32) TEST_SUITE(Scale255) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmall, ToF32Fixture, PRECOMMIT, SmallShapes(), F32, F32, scale_255, TO_NEAREST_UP, EmptyActivationFunctionsDataset, VALIDATE(float, 1.f)) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunWithActivation, ToF32Fixture, ALL, TinyShapes(), F32, F32, scale_255, TO_NEAREST_UP, ActivationFunctionsDataset, VALIDATE(float, 1.f)) +FIXTURE_DATA_TEST_CASE(RunInplace, CLPixelWiseMultiplicationToF32Fixture, framework::DatasetMode::PRECOMMIT, + combine(combine(combine(combine(combine(combine(combine(datasets::TinyShapes(), + framework::dataset::make("DataTypeIn1", DataType::F32)), + framework::dataset::make("DataTypeIn2", DataType::F32)), + framework::dataset::make("Scale", { scale_255 })), + datasets::ConvertPolicies()), + framework::dataset::make("RoundingPolicy", RoundingPolicy::TO_NEAREST_UP)), + EmptyActivationFunctionsDataset), + framework::dataset::make("InPlace", { true }))) +{ + // Validate output + VALIDATE(float, 1.f) +} TEST_SUITE_END() // Scale255 TEST_SUITE_END() // F32toF32 @@ -200,6 +222,23 @@ FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLPixelWiseMultiplicationQuantizedBroa validate(CLAccessor(_target), _reference, tolerance_qasymm8); } +FIXTURE_DATA_TEST_CASE(RunInplace, CLPixelWiseMultiplicationQuantizedBroadcastFixture, framework::DatasetMode::PRECOMMIT, + combine(combine(combine(combine(combine(combine(combine(combine(combine(combine(datasets::TinyShapesBroadcastInplace(), + framework::dataset::make("DataTypeIn1", DataType::QASYMM8)), + framework::dataset::make("DataTypeIn2", DataType::QASYMM8)), + framework::dataset::make("DataTypeOut", DataType::QASYMM8)), + framework::dataset::make("Scale", { 1.f, 2.f })), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), + framework::dataset::make("RoundingPolicy", RoundingPolicy::TO_NEAREST_EVEN)), + framework::dataset::make("Src0QInfo", { QuantizationInfo(2.f / 255.f, 10) })), + framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 255.f, 10) })), + framework::dataset::make("OUtQInfo", { QuantizationInfo(2.f / 255.f, 10) })), + framework::dataset::make("InPlace", { true }))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_qasymm8); +} + TEST_SUITE_END() // QASYMM8 TEST_SUITE(QASYMM8_SIGNED) diff --git a/tests/validation/NEON/ArithmeticAddition.cpp b/tests/validation/NEON/ArithmeticAddition.cpp index 213dbc1f5e..f3e4dfc6e5 100644 --- a/tests/validation/NEON/ArithmeticAddition.cpp +++ b/tests/validation/NEON/ArithmeticAddition.cpp @@ -48,6 +48,8 @@ constexpr AbsoluteTolerance tolerance_quant(1); /**< Tolerance value for #else // !defined(__aarch64__) || defined(ENABLE_SVE) constexpr AbsoluteTolerance tolerance_quant(0); #endif // !defined(__aarch64__) || defined(ENABLE_SVE) +const auto InPlaceDataSet = framework::dataset::make("InPlace", { false, true }); +const auto OutOfPlaceDataSet = framework::dataset::make("InPlace", { false }); } // namespace TEST_SUITE(NEON) @@ -104,8 +106,10 @@ TEST_CASE(NoPaddingAdded, framework::DatasetMode::PRECOMMIT) TEST_SUITE(Integer) TEST_SUITE(U8) -FIXTURE_DATA_TEST_CASE(RunSmall, NEArithmeticAdditionFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::U8)), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP }))) +FIXTURE_DATA_TEST_CASE(RunSmall, NEArithmeticAdditionFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", + DataType::U8)), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); @@ -113,15 +117,19 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEArithmeticAdditionFixture, framework TEST_SUITE_END() // U8 TEST_SUITE(S16) -FIXTURE_DATA_TEST_CASE(RunSmall, NEArithmeticAdditionFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::S16)), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP }))) +FIXTURE_DATA_TEST_CASE(RunSmall, NEArithmeticAdditionFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", + DataType::S16)), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); } -FIXTURE_DATA_TEST_CASE(RunLarge, NEArithmeticAdditionFixture, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapes(), framework::dataset::make("DataType", DataType::S16)), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP }))) +FIXTURE_DATA_TEST_CASE(RunLarge, NEArithmeticAdditionFixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::LargeShapes(), framework::dataset::make("DataType", + DataType::S16)), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); @@ -129,8 +137,10 @@ FIXTURE_DATA_TEST_CASE(RunLarge, NEArithmeticAdditionFixture, framework TEST_SUITE_END() // S16 TEST_SUITE(S32) -FIXTURE_DATA_TEST_CASE(RunSmall, NEArithmeticAdditionFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::S32)), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP }))) +FIXTURE_DATA_TEST_CASE(RunSmall, NEArithmeticAdditionFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", + DataType::S32)), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); @@ -141,8 +151,9 @@ TEST_SUITE_END() // Integer TEST_SUITE(Float) #ifdef __ARM_FEATURE_FP16_VECTOR_ARITHMETIC TEST_SUITE(F16) -FIXTURE_DATA_TEST_CASE(RunSmall, NEArithmeticAdditionFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::F16)), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP }))) +FIXTURE_DATA_TEST_CASE(RunSmall, NEArithmeticAdditionFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::F16)), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); @@ -151,15 +162,19 @@ TEST_SUITE_END() // F16 #endif /* __ARM_FEATURE_FP16_VECTOR_ARITHMETIC */ TEST_SUITE(F32) -FIXTURE_DATA_TEST_CASE(RunSmall, NEArithmeticAdditionFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::F32)), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP }))) +FIXTURE_DATA_TEST_CASE(RunSmall, NEArithmeticAdditionFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", + DataType::F32)), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); } -FIXTURE_DATA_TEST_CASE(RunLarge, NEArithmeticAdditionFixture, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapes(), framework::dataset::make("DataType", DataType::F32)), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP }))) +FIXTURE_DATA_TEST_CASE(RunLarge, NEArithmeticAdditionFixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::LargeShapes(), framework::dataset::make("DataType", + DataType::F32)), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); @@ -168,17 +183,19 @@ FIXTURE_DATA_TEST_CASE(RunLarge, NEArithmeticAdditionFixture, framework:: template using NEArithmeticAdditionBroadcastFixture = ArithmeticAdditionBroadcastValidationFixture; -FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, NEArithmeticAdditionBroadcastFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapesBroadcast(), +FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, NEArithmeticAdditionBroadcastFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallShapesBroadcast(), framework::dataset::make("DataType", DataType::F32)), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP }))) + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); } -FIXTURE_DATA_TEST_CASE(RunLargeBroadcast, NEArithmeticAdditionBroadcastFixture, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapesBroadcast(), +FIXTURE_DATA_TEST_CASE(RunLargeBroadcast, NEArithmeticAdditionBroadcastFixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::LargeShapesBroadcast(), framework::dataset::make("DataType", DataType::F32)), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP }))) + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); @@ -197,11 +214,12 @@ TEST_SUITE(QASYMM8) FIXTURE_DATA_TEST_CASE(RunSmall, NEArithmeticAdditionQuantizedFixture, framework::DatasetMode::PRECOMMIT, - combine(combine(combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::QASYMM8)), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), - framework::dataset::make("Src0QInfo", { QuantizationInfo(5.f / 255.f, 20) })), - framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 255.f, 10) })), - framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 255.f, 5) }))) + combine(combine(combine(combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::QASYMM8)), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), + framework::dataset::make("Src0QInfo", { QuantizationInfo(5.f / 255.f, 20) })), + framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 255.f, 10) })), + framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 255.f, 5) })), + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_quant); @@ -212,22 +230,24 @@ TEST_SUITE(QASYMM8_SIGNED) FIXTURE_DATA_TEST_CASE(RunSmall, NEArithmeticAdditionQuantizedFixture, framework::DatasetMode::ALL, - combine(combine(combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::QASYMM8_SIGNED)), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), - framework::dataset::make("Src0QInfo", { QuantizationInfo(0.5f, 20) })), - framework::dataset::make("Src1QInfo", { QuantizationInfo(0.5f, 10) })), - framework::dataset::make("OutQInfo", { QuantizationInfo(0.5f, 5) }))) + combine(combine(combine(combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::QASYMM8_SIGNED)), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), + framework::dataset::make("Src0QInfo", { QuantizationInfo(0.5f, 20) })), + framework::dataset::make("Src1QInfo", { QuantizationInfo(0.5f, 10) })), + framework::dataset::make("OutQInfo", { QuantizationInfo(0.5f, 5) })), + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_quant); } -FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, NEArithmeticAdditionQuantizedBroadcastFixture, framework::DatasetMode::ALL, combine(combine(combine(combine(combine( +FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, NEArithmeticAdditionQuantizedBroadcastFixture, framework::DatasetMode::ALL, combine(combine(combine(combine(combine(combine( datasets::SmallShapesBroadcast(), framework::dataset::make("DataType", DataType::QASYMM8_SIGNED)), framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), framework::dataset::make("Src0QInfo", { QuantizationInfo(0.5f, 20) })), framework::dataset::make("Src1QInfo", { QuantizationInfo(0.5f, 10) })), - framework::dataset::make("OutQInfo", { QuantizationInfo(0.5f, 5) }))) + framework::dataset::make("OutQInfo", { QuantizationInfo(0.5f, 5) })), + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_quant); @@ -238,11 +258,12 @@ TEST_SUITE(QSYMM16) FIXTURE_DATA_TEST_CASE(RunSmall, NEArithmeticAdditionQuantizedFixture, framework::DatasetMode::PRECOMMIT, - combine(combine(combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::QSYMM16)), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), - framework::dataset::make("Src0QInfo", { QuantizationInfo(1.f / 32768.f, 0), QuantizationInfo(5.f / 32768.f, 0) })), - framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 32768.f, 0), QuantizationInfo(5.f / 32768.f, 0) })), - framework::dataset::make("OutQInfo", { QuantizationInfo(5.f / 32768.f, 0) }))) + combine(combine(combine(combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::QSYMM16)), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), + framework::dataset::make("Src0QInfo", { QuantizationInfo(1.f / 32768.f, 0), QuantizationInfo(5.f / 32768.f, 0) })), + framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 32768.f, 0), QuantizationInfo(5.f / 32768.f, 0) })), + framework::dataset::make("OutQInfo", { QuantizationInfo(5.f / 32768.f, 0) })), + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_quant); diff --git a/tests/validation/NEON/ArithmeticSubtraction.cpp b/tests/validation/NEON/ArithmeticSubtraction.cpp index 68213fb51f..8886ca2db5 100644 --- a/tests/validation/NEON/ArithmeticSubtraction.cpp +++ b/tests/validation/NEON/ArithmeticSubtraction.cpp @@ -57,6 +57,9 @@ const auto ArithmeticSubtractionQuantizationInfoDataset = combine(combine(framew const auto ArithmeticSubtractionQuantizationInfoSignedDataset = combine(combine(framework::dataset::make("QuantizationInfoIn1", { QuantizationInfo(0.5f, 10) }), framework::dataset::make("QuantizationInfoIn2", { QuantizationInfo(0.5f, 20) })), framework::dataset::make("QuantizationInfoOut", { QuantizationInfo(0.5f, 50) })); +const auto ArithmeticSubtractionQuantizationInfoSignedInPlaceDataset = combine(combine(framework::dataset::make("QuantizationInfoIn1", { QuantizationInfo(0.8f, 10) }), + framework::dataset::make("QuantizationInfoIn2", { QuantizationInfo(0.8f, 10) })), + framework::dataset::make("QuantizationInfoOut", { QuantizationInfo(0.8f, 10) })); const auto ArithmeticSubtractionQuantizationInfoSymmetric = combine(combine(framework::dataset::make("QuantizationInfoIn1", { QuantizationInfo(0.3f, 0) }), framework::dataset::make("QuantizationInfoIn2", { QuantizationInfo(0.7f, 0) })), framework::dataset::make("QuantizationInfoOut", { QuantizationInfo(0.2f, 0) })); @@ -179,7 +182,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEArithmeticSubtractionQASYMM8Fixture, framewor DataType::QASYMM8)), framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), ArithmeticSubtractionQuantizationInfoDataset), - InPlaceDataSet)) + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_qasymm8); @@ -191,12 +194,11 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEArithmeticSubtractionQASYMM8SignedFixture, fr datasets::SmallShapes(), framework::dataset::make("DataType", DataType::QASYMM8_SIGNED)), framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), ArithmeticSubtractionQuantizationInfoSignedDataset), - InPlaceDataSet)) + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_qasymm8); } - FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, NEArithmeticSubtractionQASYMM8SignedBroadcastFixture, framework::DatasetMode::ALL, combine(combine(combine(combine( datasets::SmallShapesBroadcast(), framework::dataset::make("DataType", DataType::QASYMM8_SIGNED)), @@ -207,6 +209,16 @@ FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, NEArithmeticSubtractionQASYMM8SignedBr // Validate output validate(Accessor(_target), _reference, tolerance_qasymm8); } +FIXTURE_DATA_TEST_CASE(RunTinyBroadcastInPlace, NEArithmeticSubtractionQASYMM8SignedBroadcastFixture, framework::DatasetMode::ALL, combine(combine(combine(combine( + datasets::TinyShapesBroadcastInplace(), + framework::dataset::make("DataType", DataType::QASYMM8_SIGNED)), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), + ArithmeticSubtractionQuantizationInfoSignedInPlaceDataset), + InPlaceDataSet)) +{ + // Validate output + validate(Accessor(_target), _reference, tolerance_qasymm8); +} TEST_SUITE_END() // QASYMM8_SIGNED TEST_SUITE(QSYMM16) diff --git a/tests/validation/NEON/ElementwiseDivision.cpp b/tests/validation/NEON/ElementwiseDivision.cpp index 8abccb2ed6..5f0224c91d 100644 --- a/tests/validation/NEON/ElementwiseDivision.cpp +++ b/tests/validation/NEON/ElementwiseDivision.cpp @@ -56,6 +56,8 @@ const auto ElementwiseDivisionFP16Dataset = combine(combine(framewo #endif /* __ARM_FEATURE_FP16_VECTOR_ARITHMETIC */ const auto ElementwiseDivisionFP32Dataset = combine(combine(framework::dataset::make("DataType", DataType::F32), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("DataType", DataType::F32)); +const auto InPlaceDataSet = framework::dataset::make("InPlace", { false, true }); +const auto OutOfPlaceDataSet = framework::dataset::make("InPlace", { false }); } // namespace TEST_SUITE(NEON) @@ -105,14 +107,16 @@ using CpuElementwiseDivisionBroadcastDynamicShapeFixture = ArithmeticDivisionBro TEST_SUITE(F32) -FIXTURE_DATA_TEST_CASE(RunSmall, CpuElementwiseDivisionDynamicShapeFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwiseDivisionFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CpuElementwiseDivisionDynamicShapeFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseDivisionFP32Dataset), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_fp32, 0.01); } -FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CpuElementwiseDivisionBroadcastDynamicShapeFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapesBroadcast(), - ElementwiseDivisionFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CpuElementwiseDivisionBroadcastDynamicShapeFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapesBroadcast(), + ElementwiseDivisionFP32Dataset), + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_fp32, 0.01); @@ -124,7 +128,8 @@ TEST_SUITE_END() // DynamicShape TEST_SUITE(Float) #ifdef __ARM_FEATURE_FP16_VECTOR_ARITHMETIC TEST_SUITE(F16) -FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseDivisionFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwiseDivisionFP16Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseDivisionFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseDivisionFP16Dataset), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_fp16, 0.01); @@ -133,7 +138,8 @@ TEST_SUITE_END() // F16 #endif /* __ARM_FEATURE_FP16_VECTOR_ARITHMETIC */ TEST_SUITE(F32) -FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseDivisionFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwiseDivisionFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseDivisionFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseDivisionFP32Dataset), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_fp32, 0.01); @@ -142,8 +148,16 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseDivisionFixture, framework: template using NEElementwiseDivisionBroadcastFixture = ArithmeticDivisionBroadcastValidationFixture; -FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, NEElementwiseDivisionBroadcastFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapesBroadcast(), - ElementwiseDivisionFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, NEElementwiseDivisionBroadcastFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapesBroadcast(), + ElementwiseDivisionFP32Dataset), + OutOfPlaceDataSet)) +{ + // Validate output + validate(Accessor(_target), _reference, tolerance_fp32, 0.01); +} +FIXTURE_DATA_TEST_CASE(RunTinyBroadcastInPlace, NEElementwiseDivisionBroadcastFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapesBroadcastInplace(), + ElementwiseDivisionFP32Dataset), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_fp32, 0.01); @@ -153,7 +167,8 @@ TEST_SUITE_END() // Float TEST_SUITE(Integer) TEST_SUITE(S32) -FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseDivisionFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwiseDivisionS32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseDivisionFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseDivisionS32Dataset), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_zero_s32); diff --git a/tests/validation/NEON/ElementwiseMax.cpp b/tests/validation/NEON/ElementwiseMax.cpp index 4bc263184e..61421ab3e5 100644 --- a/tests/validation/NEON/ElementwiseMax.cpp +++ b/tests/validation/NEON/ElementwiseMax.cpp @@ -62,6 +62,8 @@ const auto ElementwiseMaxFP16Dataset = combine(combine(framework::dataset::make( #endif /* __ARM_FEATURE_FP16_VECTOR_ARITHMETIC */ const auto ElementwiseMaxFP32Dataset = combine(combine(framework::dataset::make("DataType", DataType::F32), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("DataType", DataType::F32)); +const auto InPlaceDataSet = framework::dataset::make("InPlace", { false, true }); +const auto OutOfPlaceDataSet = framework::dataset::make("InPlace", { false }); } // namespace TEST_SUITE(NEON) @@ -111,7 +113,8 @@ DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip( // *INDENT-ON* TEST_SUITE(S32) -FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMaxFixture, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallShapes(), ElementwiseMaxS32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMaxFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), ElementwiseMaxS32Dataset), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); @@ -119,7 +122,8 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMaxFixture, framework::Da TEST_SUITE_END() // S32 TEST_SUITE(S16) -FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMaxFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwiseMaxS16Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMaxFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseMaxS16Dataset), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); @@ -131,11 +135,12 @@ using NEElementwiseMaxQuantizedFixture = ElementwiseMaxValidationQuantizedFixtur TEST_SUITE(Quantized) TEST_SUITE(QASYMM8) -FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMaxQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(), +FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMaxQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(datasets::SmallShapes(), ElementwiseMaxQASYMM8Dataset), framework::dataset::make("QuantizationInfo", { QuantizationInfo(5.f / 255.f, 20) })), framework::dataset::make("QuantizationInfo", { QuantizationInfo(2.f / 255.f, 10) })), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(1.f / 255.f, 5) }))) + framework::dataset::make("QuantizationInfo", { QuantizationInfo(1.f / 255.f, 5) })), + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_fp32, 0.01); @@ -144,11 +149,13 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMaxQuantizedFixture, fram template using NEElementwiseMaxQuantizedBroadcastFixture = ElementwiseMaxQuantizedBroadcastValidationFixture; -FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, NEElementwiseMaxQuantizedBroadcastFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapesBroadcast(), - ElementwiseMaxQASYMM8Dataset), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(5.f / 255.f, 20) })), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(2.f / 255.f, 10) })), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(1.f / 255.f, 5) }))) +FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, NEElementwiseMaxQuantizedBroadcastFixture, framework::DatasetMode::PRECOMMIT, + combine(combine(combine(combine(combine(datasets::SmallShapesBroadcast(), + ElementwiseMaxQASYMM8Dataset), + framework::dataset::make("QuantizationInfo", { QuantizationInfo(5.f / 255.f, 20) })), + framework::dataset::make("QuantizationInfo", { QuantizationInfo(2.f / 255.f, 10) })), + framework::dataset::make("QuantizationInfo", { QuantizationInfo(1.f / 255.f, 5) })), + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); @@ -156,16 +163,26 @@ FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, NEElementwiseMaxQuantizedBroadcastFixt TEST_SUITE_END() TEST_SUITE(QASYMM8_SIGNED) -FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMaxQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(), +FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMaxQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(datasets::SmallShapes(), ElementwiseMaxQASYMM8SignedDataset), framework::dataset::make("QuantizationInfo", { QuantizationInfo(10.f, 20) })), framework::dataset::make("QuantizationInfo", { QuantizationInfo(1.f, 0) })), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(2.f, -27) }))) + framework::dataset::make("QuantizationInfo", { QuantizationInfo(2.f, -27) })), + OutOfPlaceDataSet)) +{ + // Validate output + validate(Accessor(_target), _reference, tolerance_qasymm8_signed); +} +FIXTURE_DATA_TEST_CASE(RunSmallInPlace, NEElementwiseMaxQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(datasets::SmallShapes(), + ElementwiseMaxQASYMM8SignedDataset), + framework::dataset::make("QuantizationInfo", { QuantizationInfo(10.f, -20) })), + framework::dataset::make("QuantizationInfo", { QuantizationInfo(10.f, -20) })), + framework::dataset::make("QuantizationInfo", { QuantizationInfo(10.f, -20) })), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_qasymm8_signed); } - TEST_SUITE_END() TEST_SUITE_END() @@ -173,7 +190,8 @@ TEST_SUITE_END() TEST_SUITE(Float) #ifdef __ARM_FEATURE_FP16_VECTOR_ARITHMETIC TEST_SUITE(F16) -FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMaxFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwiseMaxFP16Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMaxFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseMaxFP16Dataset), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); @@ -182,7 +200,8 @@ TEST_SUITE_END() // F16 #endif /* __ARM_FEATURE_FP16_VECTOR_ARITHMETIC */ TEST_SUITE(F32) -FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMaxFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwiseMaxFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMaxFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseMaxFP32Dataset), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); @@ -190,8 +209,16 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMaxFixture, framework::Data template using NEElementwiseMaxBroadcastFixture = ElementwiseMaxBroadcastValidationFixture; -FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, NEElementwiseMaxBroadcastFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapesBroadcast(), - ElementwiseMaxFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, NEElementwiseMaxBroadcastFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapesBroadcast(), + ElementwiseMaxFP32Dataset), + OutOfPlaceDataSet)) +{ + // Validate output + validate(Accessor(_target), _reference); +} +FIXTURE_DATA_TEST_CASE(RunTinyBroadcastInPlace, NEElementwiseMaxBroadcastFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapesBroadcastInplace(), + ElementwiseMaxFP32Dataset), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); diff --git a/tests/validation/NEON/ElementwiseMin.cpp b/tests/validation/NEON/ElementwiseMin.cpp index 3836b90308..a134eb354d 100644 --- a/tests/validation/NEON/ElementwiseMin.cpp +++ b/tests/validation/NEON/ElementwiseMin.cpp @@ -62,6 +62,8 @@ const auto ElementwiseMinFP16Dataset = combine(combine(framework::dataset::make( #endif /* __ARM_FEATURE_FP16_VECTOR_ARITHMETIC */ const auto ElementwiseMinFP32Dataset = combine(combine(framework::dataset::make("DataType", DataType::F32), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("DataType", DataType::F32)); +const auto InPlaceDataSet = framework::dataset::make("InPlace", { false, true }); +const auto OutOfPlaceDataSet = framework::dataset::make("InPlace", { false }); } // namespace TEST_SUITE(NEON) @@ -110,7 +112,8 @@ DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip( // *INDENT-ON* TEST_SUITE(S32) -FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMinFixture, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallShapes(), ElementwiseMinS32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMinFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), ElementwiseMinS32Dataset), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); @@ -118,7 +121,8 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMinFixture, framework::Da TEST_SUITE_END() // S32 TEST_SUITE(S16) -FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMinFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwiseMinS16Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMinFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseMinS16Dataset), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); @@ -133,23 +137,34 @@ TEST_SUITE(QASYMM8) template using NEElementwiseMinQuantizedBroadcastFixture = ElementwiseMinQuantizedBroadcastValidationFixture; -FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, NEElementwiseMinQuantizedBroadcastFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapesBroadcast(), - ElementwiseMinQASYMM8Dataset), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(5.f / 255.f, 20) })), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(2.f / 255.f, 10) })), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(1.f / 255.f, 5) }))) +FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, NEElementwiseMinQuantizedBroadcastFixture, framework::DatasetMode::PRECOMMIT, + combine(combine(combine(combine(combine(datasets::SmallShapesBroadcast(), + ElementwiseMinQASYMM8Dataset), + framework::dataset::make("QuantizationInfo", { QuantizationInfo(5.f / 255.f, 20) })), + framework::dataset::make("QuantizationInfo", { QuantizationInfo(2.f / 255.f, 10) })), + framework::dataset::make("QuantizationInfo", { QuantizationInfo(1.f / 255.f, 5) })), + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); } - -FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMinQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(), +FIXTURE_DATA_TEST_CASE(RunTinyBroadcastInPlace, NEElementwiseMinQuantizedBroadcastFixture, framework::DatasetMode::PRECOMMIT, + combine(combine(combine(combine(combine(datasets::TinyShapesBroadcastInplace(), + ElementwiseMinQASYMM8Dataset), + framework::dataset::make("QuantizationInfo", { QuantizationInfo(2.f / 255.f, 20) })), + framework::dataset::make("QuantizationInfo", { QuantizationInfo(2.f / 255.f, 20) })), + framework::dataset::make("QuantizationInfo", { QuantizationInfo(2.f / 255.f, 20) })), + InPlaceDataSet)) +{ + // Validate output + validate(Accessor(_target), _reference); +} +FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMinQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(datasets::SmallShapes(), ElementwiseMinQASYMM8Dataset), framework::dataset::make("QuantizationInfo", { QuantizationInfo(5.f / 255.f, 20) })), framework::dataset::make("QuantizationInfo", { QuantizationInfo(2.f / 255.f, 10) })), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(1.f / 255.f, 5) })) - - ) + framework::dataset::make("QuantizationInfo", { QuantizationInfo(1.f / 255.f, 5) })), + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_fp32, 0.01); @@ -157,11 +172,12 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMinQuantizedFixture, fram TEST_SUITE_END() TEST_SUITE(QASYMM8_SIGNED) -FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMinQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(), +FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMinQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(datasets::SmallShapes(), ElementwiseMaxQASYMM8SignedDataset), framework::dataset::make("QuantizationInfo", { QuantizationInfo(10.f, 20) })), framework::dataset::make("QuantizationInfo", { QuantizationInfo(1.f, 0) })), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(2.f, -27) }))) + framework::dataset::make("QuantizationInfo", { QuantizationInfo(2.f, -27) })), + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_qasymm8_signed); @@ -174,7 +190,8 @@ TEST_SUITE_END() TEST_SUITE(Float) #ifdef __ARM_FEATURE_FP16_VECTOR_ARITHMETIC TEST_SUITE(F16) -FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMinFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwiseMinFP16Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMinFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseMinFP16Dataset), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); @@ -183,7 +200,8 @@ TEST_SUITE_END() // F16 #endif /* __ARM_FEATURE_FP16_VECTOR_ARITHMETIC */ TEST_SUITE(F32) -FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMinFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwiseMinFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMinFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseMinFP32Dataset), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); @@ -192,8 +210,16 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseMinFixture, framework::Data template using NEElementwiseMinBroadcastFixture = ElementwiseMinBroadcastValidationFixture; -FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, NEElementwiseMinBroadcastFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapesBroadcast(), - ElementwiseMinFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, NEElementwiseMinBroadcastFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapesBroadcast(), + ElementwiseMinFP32Dataset), + OutOfPlaceDataSet)) +{ + // Validate output + validate(Accessor(_target), _reference); +} +FIXTURE_DATA_TEST_CASE(RunTinyBroadcastInPlace, NEElementwiseMinBroadcastFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapesBroadcastInplace(), + ElementwiseMinFP32Dataset), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); diff --git a/tests/validation/NEON/ElementwisePower.cpp b/tests/validation/NEON/ElementwisePower.cpp index 4305387c5f..9ac9eec280 100644 --- a/tests/validation/NEON/ElementwisePower.cpp +++ b/tests/validation/NEON/ElementwisePower.cpp @@ -51,6 +51,8 @@ const auto ElementwisePowerFP16Dataset = combine(combine(framework: #endif /* __ARM_FEATURE_FP16_VECTOR_ARITHMETIC */ const auto ElementwisePowerFP32Dataset = combine(combine(framework::dataset::make("DataType", DataType::F32), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("DataType", DataType::F32)); +const auto InPlaceDataSet = framework::dataset::make("InPlace", { false, true }); +const auto OutOfPlaceDataSet = framework::dataset::make("InPlace", { false }); } // namespace TEST_SUITE(NEON) @@ -91,7 +93,8 @@ DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip( TEST_SUITE(Float) #ifdef __ARM_FEATURE_FP16_VECTOR_ARITHMETIC TEST_SUITE(F16) -FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwisePowerFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwisePowerFP16Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwisePowerFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwisePowerFP16Dataset), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_fp16, 0.01); @@ -101,13 +104,15 @@ TEST_SUITE_END() // F16 TEST_SUITE(F32) -FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwisePowerFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwisePowerFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwisePowerFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwisePowerFP32Dataset), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_fp32, 0.01); } -FIXTURE_DATA_TEST_CASE(RunLarge, NEElementwisePowerFixture, framework::DatasetMode::NIGHTLY, combine(datasets::LargeShapes(), ElementwisePowerFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunLarge, NEElementwisePowerFixture, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapes(), ElementwisePowerFP32Dataset), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_fp32, 0.01); @@ -116,15 +121,23 @@ FIXTURE_DATA_TEST_CASE(RunLarge, NEElementwisePowerFixture, framework::Da template using NEElementwisePowerBroadcastFixture = ElementwisePowerBroadcastValidationFixture; -FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, NEElementwisePowerBroadcastFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapesBroadcast(), - ElementwisePowerFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, NEElementwisePowerBroadcastFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapesBroadcast(), + ElementwisePowerFP32Dataset), + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_fp32, 0.01); } - -FIXTURE_DATA_TEST_CASE(RunLargeBroadcast, NEElementwisePowerBroadcastFixture, framework::DatasetMode::NIGHTLY, combine(datasets::LargeShapesBroadcast(), - ElementwisePowerFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunTinyBroadcastInPlace, NEElementwisePowerBroadcastFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapesBroadcastInplace(), + ElementwisePowerFP32Dataset), + InPlaceDataSet)) +{ + // Validate output + validate(Accessor(_target), _reference, tolerance_fp32, 0.01); +} +FIXTURE_DATA_TEST_CASE(RunLargeBroadcast, NEElementwisePowerBroadcastFixture, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapesBroadcast(), + ElementwisePowerFP32Dataset), + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_fp32, 0.01); diff --git a/tests/validation/NEON/ElementwiseSquareDiff.cpp b/tests/validation/NEON/ElementwiseSquareDiff.cpp index 069cbbd7fa..9a86b541de 100644 --- a/tests/validation/NEON/ElementwiseSquareDiff.cpp +++ b/tests/validation/NEON/ElementwiseSquareDiff.cpp @@ -68,6 +68,8 @@ const auto ElementwiseSquaredDiffFP16Dataset = combine(combine(framework::datase #endif /* __ARM_FEATURE_FP16_VECTOR_ARITHMETIC */ const auto ElementwiseSquaredDiffFP32Dataset = combine(combine(framework::dataset::make("DataType", DataType::F32), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("DataType", DataType::F32)); +const auto InPlaceDataSet = framework::dataset::make("InPlace", { false, true }); +const auto OutOfPlaceDataSet = framework::dataset::make("InPlace", { false }); } // namespace TEST_SUITE(NEON) @@ -109,7 +111,8 @@ DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip( // *INDENT-ON* TEST_SUITE(S32) -FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseSquaredDiffFixture, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallShapes(), ElementwiseSquaredDiffS32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseSquaredDiffFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), ElementwiseSquaredDiffS32Dataset), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); @@ -117,7 +120,8 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseSquaredDiffFixture, frame TEST_SUITE_END() // S32 TEST_SUITE(S16) -FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseSquaredDiffFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwiseSquaredDiffS16Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseSquaredDiffFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseSquaredDiffS16Dataset), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); @@ -129,13 +133,12 @@ using NEElementwiseSquaredDiffQuantizedFixture = ElementwiseSquaredDiffValidatio TEST_SUITE(Quantized) TEST_SUITE(QASYMM8) -FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseSquaredDiffQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(), +FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseSquaredDiffQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(datasets::SmallShapes(), ElementwiseSquaredDiffQASYMM8Dataset), framework::dataset::make("QuantizationInfo", { QuantizationInfo(5.f / 255.f, 20) })), framework::dataset::make("QuantizationInfo", { QuantizationInfo(2.f / 255.f, 10) })), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(1.f / 255.f, 5) })) - - ) + framework::dataset::make("QuantizationInfo", { QuantizationInfo(1.f / 255.f, 5) })), + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_fp32, 0.01); @@ -144,11 +147,23 @@ template using NEElementwiseSquaredDiffQuantizedBroadcastFixture = ElementwiseSquaredDiffQuantizedBroadcastValidationFixture; FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, NEElementwiseSquaredDiffQuantizedBroadcastFixture, framework::DatasetMode::PRECOMMIT, - combine(combine(combine(combine(datasets::SmallShapesBroadcast(), - ElementwiseSquaredDiffQASYMM8Dataset), + combine(combine(combine(combine(combine(datasets::SmallShapesBroadcast(), + ElementwiseSquaredDiffQASYMM8Dataset), + framework::dataset::make("QuantizationInfo", { QuantizationInfo(5.f / 255.f, 20) })), + framework::dataset::make("QuantizationInfo", { QuantizationInfo(2.f / 255.f, 10) })), + framework::dataset::make("QuantizationInfo", { QuantizationInfo(1.f / 255.f, 5) })), + OutOfPlaceDataSet)) +{ + // Validate output + validate(Accessor(_target), _reference); +} +FIXTURE_DATA_TEST_CASE(RunTinyBroadcastInPlace, NEElementwiseSquaredDiffQuantizedBroadcastFixture, framework::DatasetMode::ALL, + combine(combine(combine(combine(combine(datasets::TinyShapesBroadcastInplace(), + ElementwiseSquaredDiffQASYMM8Dataset), + framework::dataset::make("QuantizationInfo", { QuantizationInfo(5.f / 255.f, 20) })), framework::dataset::make("QuantizationInfo", { QuantizationInfo(5.f / 255.f, 20) })), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(2.f / 255.f, 10) })), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(1.f / 255.f, 5) }))) + framework::dataset::make("QuantizationInfo", { QuantizationInfo(5.f / 255.f, 20) })), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); @@ -156,11 +171,12 @@ FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, NEElementwiseSquaredDiffQuantizedBroad TEST_SUITE_END() TEST_SUITE(QASYMM8_SIGNED) -FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseSquaredDiffQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(), +FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseSquaredDiffQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(datasets::SmallShapes(), ElementwiseSquaredDiffQASYMM8SignedDataset), framework::dataset::make("QuantizationInfo", { QuantizationInfo(1.f, 5) })), framework::dataset::make("QuantizationInfo", { QuantizationInfo(.5f, 5) })), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(.2f, 5) }))) + framework::dataset::make("QuantizationInfo", { QuantizationInfo(.2f, 5) })), + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); @@ -171,7 +187,8 @@ TEST_SUITE_END() TEST_SUITE(Float) #ifdef __ARM_FEATURE_FP16_VECTOR_ARITHMETIC TEST_SUITE(F16) -FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseSquaredDiffFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwiseSquaredDiffFP16Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseSquaredDiffFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseSquaredDiffFP16Dataset), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_fp16, 0.01); @@ -180,7 +197,8 @@ TEST_SUITE_END() // F16 #endif /* __ARM_FEATURE_FP16_VECTOR_ARITHMETIC */ TEST_SUITE(F32) -FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseSquaredDiffFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwiseSquaredDiffFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseSquaredDiffFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseSquaredDiffFP32Dataset), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); @@ -188,15 +206,17 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseSquaredDiffFixture, framewo template using NEElementwiseSquaredDiffBroadcastFixture = ElementwiseSquaredDiffBroadcastValidationFixture; -FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, NEElementwiseSquaredDiffBroadcastFixture, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallShapesBroadcast(), - ElementwiseSquaredDiffFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, NEElementwiseSquaredDiffBroadcastFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapesBroadcast(), + ElementwiseSquaredDiffFP32Dataset), + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); } -FIXTURE_DATA_TEST_CASE(RunLargeBroadcast, NEElementwiseSquaredDiffBroadcastFixture, framework::DatasetMode::NIGHTLY, combine(datasets::LargeShapesBroadcast(), - ElementwiseSquaredDiffFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunLargeBroadcast, NEElementwiseSquaredDiffBroadcastFixture, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapesBroadcast(), + ElementwiseSquaredDiffFP32Dataset), + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); diff --git a/tests/validation/NEON/PixelWiseMultiplication.cpp b/tests/validation/NEON/PixelWiseMultiplication.cpp index 1bb0588919..964d1c5deb 100644 --- a/tests/validation/NEON/PixelWiseMultiplication.cpp +++ b/tests/validation/NEON/PixelWiseMultiplication.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2020 Arm Limited. + * Copyright (c) 2017-2021 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -56,6 +56,11 @@ const auto PixelWiseMultiplicationQASYMM8QuantDataset = combine(combine( framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 32768.f, 0) })), framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 32768.f, 0) })); +const auto PixelWiseMultiplicationQASYMM8QuantInPlaceDataset = combine(combine( + framework::dataset::make("Src0QInfo", { QuantizationInfo(5.f / 32768.f, 10) }), + framework::dataset::make("Src1QInfo", { QuantizationInfo(5.f / 32768.f, 10) })), + framework::dataset::make("OutQInfo", { QuantizationInfo(5.f / 32768.f, 10) })); + const auto PixelWiseMultiplicationPolicySTNUDataset = combine( framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE }), framework::dataset::make("RoundingPolicy", { RoundingPolicy::TO_NEAREST_UP })); @@ -75,7 +80,8 @@ const auto PixelWiseMultiplicationPolicySTZDataset = combine( * expected to have either different quantization information, data type * or different shape we are not testing in-place computation. */ -const auto InPlaceDataSet = framework::dataset::make("InPlace", { false, true }); +const auto InPlaceDataSet = framework::dataset::make("InPlace", { false, true }); +const auto OutOfPlaceDataSet = framework::dataset::make("InPlace", { false }); #define DEFAULT_VALIDATE validate(Accessor(_target), _reference); #define VALIDATE(TYPE, TOLERANCE) validate(Accessor(_target), _reference, AbsoluteTolerance(TOLERANCE), 0.f); @@ -275,7 +281,19 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQASYMM8SignedFixture, framework::dataset::make("Scale", { scale_unity })), PixelWiseMultiplicationPolicySTZDataset), PixelWiseMultiplicationQASYMM8QuantDataset), - InPlaceDataSet)) + OutOfPlaceDataSet)) +{ + // Validate output + validate(Accessor(_target), _reference, tolerance_qasymm8); +} +FIXTURE_DATA_TEST_CASE(RunSmallInPlace, NEPixelWiseMultiplicationQASYMM8SignedFixture, framework::DatasetMode::ALL, combine(combine(combine(combine(combine(combine(combine(datasets::SmallShapes(), + framework::dataset::make("DataTypeIn1", DataType::QASYMM8_SIGNED)), + framework::dataset::make("DataTypeIn2", DataType::QASYMM8_SIGNED)), + framework::dataset::make("DataTypeOut", DataType::QASYMM8_SIGNED)), + framework::dataset::make("Scale", { scale_unity })), + PixelWiseMultiplicationPolicySTZDataset), + PixelWiseMultiplicationQASYMM8QuantInPlaceDataset), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_qasymm8); @@ -292,7 +310,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQASYMM8Fixture, framew framework::dataset::make("Scale", { scale_255 })), PixelWiseMultiplicationPolicySTNUDataset), PixelWiseMultiplicationQASYMM8QuantDataset), - InPlaceDataSet)) + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_qasymm8); @@ -306,7 +324,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQASYMM8Fixture, framew framework::dataset::make("Scale", { scale_unity })), PixelWiseMultiplicationPolicySTZDataset), PixelWiseMultiplicationQASYMM8QuantDataset), - InPlaceDataSet)) + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_qasymm8); @@ -320,7 +338,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQASYMM8Fixture, framew framework::dataset::make("Scale", { scale_other })), PixelWiseMultiplicationPolicySTZDataset), PixelWiseMultiplicationQASYMM8QuantDataset), - InPlaceDataSet)) + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_qasymm8); @@ -335,7 +353,20 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationBroadcastQASYMM8Fixtur framework::dataset::make("Scale", { scale_other })), PixelWiseMultiplicationPolicySTZDataset), PixelWiseMultiplicationQASYMM8QuantDataset), - framework::dataset::make("InPlace", { false }))) + OutOfPlaceDataSet)) +{ + // Validate output + validate(Accessor(_target), _reference, tolerance_qasymm8); +} +FIXTURE_DATA_TEST_CASE(RunTinyInPlace, NEPixelWiseMultiplicationBroadcastQASYMM8Fixture, framework::DatasetMode::ALL, + combine(combine(combine(combine(combine(combine(combine(datasets::TinyShapesBroadcastInplace(), + framework::dataset::make("DataTypeIn1", DataType::QASYMM8)), + framework::dataset::make("DataTypeIn2", DataType::QASYMM8)), + framework::dataset::make("DataTypeOut", DataType::QASYMM8)), + framework::dataset::make("Scale", { scale_other })), + PixelWiseMultiplicationPolicySTZDataset), + PixelWiseMultiplicationQASYMM8QuantInPlaceDataset), + InPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_qasymm8); @@ -351,7 +382,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQSYMM16Fixture, framew framework::dataset::make("Scale", { scale_255 })), PixelWiseMultiplicationPolicySTNUDataset), PixelWiseMultiplicationQSYMM16QuantDataset), - InPlaceDataSet)) + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_qsymm16); @@ -365,7 +396,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQSYMM16Fixture, framew framework::dataset::make("Scale", { scale_unity })), PixelWiseMultiplicationPolicySTZDataset), PixelWiseMultiplicationQSYMM16QuantDataset), - InPlaceDataSet)) + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_qsymm16); @@ -379,7 +410,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQSYMM16Fixture, framew framework::dataset::make("Scale", { scale_other })), PixelWiseMultiplicationPolicySTZDataset), PixelWiseMultiplicationQSYMM16QuantDataset), - InPlaceDataSet)) + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference, tolerance_qsymm16); @@ -394,7 +425,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQSYMM16ToS32Fixture, f framework::dataset::make("Scale", { scale_unity })), PixelWiseMultiplicationPolicySTZDataset), PixelWiseMultiplicationQSYMM16QuantDataset), - framework::dataset::make("InPlace", { false }))) + OutOfPlaceDataSet)) { // Validate output validate(Accessor(_target), _reference); @@ -411,7 +442,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationU8U8ToS16Fixture, fram framework::dataset::make("Scale", { scale_255 })), datasets::ConvertPolicies()), framework::dataset::make("RoundingPolicy", RoundingPolicy::TO_NEAREST_UP)), - framework::dataset::make("InPlace", { false }))) + OutOfPlaceDataSet)) { // Validate output validate_wrap(Accessor(_target), _reference, AbsoluteTolerance(1), 0.f); @@ -451,17 +482,17 @@ TEST_SUITE_END() // U8toU8 TEST_SUITE(U8toS16) TEST_SUITE(Scale255) -PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmall, ToS16Fixture, ALL, SmallShapes(), U8, S16, S16, scale_255, TO_NEAREST_UP, framework::dataset::make("InPlace", { false }), +PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmall, ToS16Fixture, ALL, SmallShapes(), U8, S16, S16, scale_255, TO_NEAREST_UP, OutOfPlaceDataSet, WRAP_VALIDATE(int16_t, 2)) TEST_SUITE_END() // Scale255 TEST_SUITE(ScaleUnity) -PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmall, ToS16Fixture, ALL, SmallShapes(), U8, S16, S16, scale_unity, TO_ZERO, framework::dataset::make("InPlace", { false }), +PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmall, ToS16Fixture, ALL, SmallShapes(), U8, S16, S16, scale_unity, TO_ZERO, OutOfPlaceDataSet, DEFAULT_VALIDATE) TEST_SUITE_END() // ScaleUnity TEST_SUITE(ScaleOther) -PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmall, ToS16Fixture, ALL, SmallShapes(), U8, S16, S16, scale_other, TO_ZERO, framework::dataset::make("InPlace", { false }), +PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmall, ToS16Fixture, ALL, SmallShapes(), U8, S16, S16, scale_other, TO_ZERO, OutOfPlaceDataSet, DEFAULT_VALIDATE) TEST_SUITE_END() // ScaleOther diff --git a/tests/validation/fixtures/ArithmeticOperationsFixture.h b/tests/validation/fixtures/ArithmeticOperationsFixture.h index 7aa716d676..0f7e44e588 100644 --- a/tests/validation/fixtures/ArithmeticOperationsFixture.h +++ b/tests/validation/fixtures/ArithmeticOperationsFixture.h @@ -47,13 +47,13 @@ class ArithmeticOperationGenericFixture : public framework::Fixture public: template void setup(reference::ArithmeticOperation op, const TensorShape &shape0, const TensorShape &shape1, DataType data_type, ConvertPolicy convert_policy, - QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out, ActivationLayerInfo act_info, bool in_place) + QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out, ActivationLayerInfo act_info, bool is_inplace) { - _op = op; - _act_info = act_info; - _in_place = in_place; - _target = compute_target(shape0, shape1, data_type, convert_policy, qinfo0, qinfo1, qinfo_out); - _reference = compute_reference(shape0, shape1, data_type, convert_policy, qinfo0, qinfo1, qinfo_out); + _op = op; + _act_info = act_info; + _is_inplace = is_inplace; + _target = compute_target(shape0, shape1, data_type, convert_policy, qinfo0, qinfo1, qinfo_out); + _reference = compute_reference(shape0, shape1, data_type, convert_policy, qinfo0, qinfo1, qinfo_out); } protected: @@ -67,27 +67,51 @@ protected: QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out) { // Create tensors - TensorType ref_src1 = create_tensor(shape0, data_type, 1, qinfo0); - TensorType ref_src2 = create_tensor(shape1, data_type, 1, qinfo1); - TensorType dst = create_tensor(TensorShape::broadcast_shape(shape0, shape1), data_type, 1, qinfo_out); - TensorType *dst_to_use = _in_place ? &ref_src1 : &dst; + const TensorShape out_shape = TensorShape::broadcast_shape(shape0, shape1); + TensorType ref_src1 = create_tensor(shape0, data_type, 1, qinfo0); + TensorType ref_src2 = create_tensor(shape1, data_type, 1, qinfo1); + TensorType dst = create_tensor(out_shape, data_type, 1, qinfo_out); + + // Check whether do in-place computation and whether inputs are broadcast compatible + TensorType *actual_dst = &dst; + if(_is_inplace) + { + bool src1_is_inplace = !arm_compute::detail::have_different_dimensions(out_shape, shape0, 0) && (qinfo0 == qinfo_out); + bool src2_is_inplace = !arm_compute::detail::have_different_dimensions(out_shape, shape1, 0) && (qinfo1 == qinfo_out); + bool do_in_place = out_shape.total_size() != 0 && (src1_is_inplace || src2_is_inplace); + ARM_COMPUTE_ASSERT(do_in_place); + + if(src1_is_inplace) + { + actual_dst = &ref_src1; + } + else + { + actual_dst = &ref_src2; + } + } // Create and configure function FunctionType arith_op; - arith_op.configure(&ref_src1, &ref_src2, dst_to_use, convert_policy, _act_info); + arith_op.configure(&ref_src1, &ref_src2, actual_dst, convert_policy, _act_info); ARM_COMPUTE_ASSERT(ref_src1.info()->is_resizable()); ARM_COMPUTE_ASSERT(ref_src2.info()->is_resizable()); - ARM_COMPUTE_ASSERT(dst_to_use->info()->is_resizable()); // Allocate tensors ref_src1.allocator()->allocate(); ref_src2.allocator()->allocate(); - dst_to_use->allocator()->allocate(); ARM_COMPUTE_ASSERT(!ref_src1.info()->is_resizable()); ARM_COMPUTE_ASSERT(!ref_src2.info()->is_resizable()); - ARM_COMPUTE_ASSERT(!dst_to_use->info()->is_resizable()); + + // If don't do in-place computation, still need to allocate original dst + if(!_is_inplace) + { + ARM_COMPUTE_ASSERT(dst.info()->is_resizable()); + dst.allocator()->allocate(); + ARM_COMPUTE_ASSERT(!dst.info()->is_resizable()); + } // Fill tensors fill(AccessorType(ref_src1), 0); @@ -96,38 +120,30 @@ protected: // Compute function arith_op.run(); - if(_in_place) - { - return ref_src1; - } - return dst; + return std::move(*actual_dst); } SimpleTensor compute_reference(const TensorShape &shape0, const TensorShape &shape1, DataType data_type, ConvertPolicy convert_policy, QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out) { - // current in-place implementation only supports same metadata of input and output tensors. - // By ignoring output quantization information here, we can make test cases implementation much simpler. - QuantizationInfo output_qinfo = _in_place ? qinfo0 : qinfo_out; - // Create reference SimpleTensor ref_src1{ shape0, data_type, 1, qinfo0 }; SimpleTensor ref_src2{ shape1, data_type, 1, qinfo1 }; - SimpleTensor ref_dst{ TensorShape::broadcast_shape(shape0, shape1), data_type, 1, output_qinfo }; + SimpleTensor ref_dst{ TensorShape::broadcast_shape(shape0, shape1), data_type, 1, qinfo_out }; // Fill reference fill(ref_src1, 0); fill(ref_src2, 1); auto result = reference::arithmetic_operation(_op, ref_src1, ref_src2, ref_dst, convert_policy); - return _act_info.enabled() ? reference::activation_layer(result, _act_info, output_qinfo) : result; + return _act_info.enabled() ? reference::activation_layer(result, _act_info, qinfo_out) : result; } TensorType _target{}; SimpleTensor _reference{}; reference::ArithmeticOperation _op{ reference::ArithmeticOperation::ADD }; ActivationLayerInfo _act_info{}; - bool _in_place{}; + bool _is_inplace{}; }; template @@ -135,10 +151,10 @@ class ArithmeticAdditionBroadcastValidationFixture : public ArithmeticOperationG { public: template - void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type, ConvertPolicy convert_policy) + void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type, ConvertPolicy convert_policy, bool is_inplace) { ArithmeticOperationGenericFixture::setup(reference::ArithmeticOperation::ADD, shape0, shape1, data_type, convert_policy, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), ActivationLayerInfo(), false); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), ActivationLayerInfo(), is_inplace); } }; @@ -147,10 +163,10 @@ class ArithmeticAdditionValidationFixture : public ArithmeticOperationGenericFix { public: template - void setup(const TensorShape &shape, DataType data_type, ConvertPolicy convert_policy) + void setup(const TensorShape &shape, DataType data_type, ConvertPolicy convert_policy, bool is_inplace) { ArithmeticOperationGenericFixture::setup(reference::ArithmeticOperation::ADD, shape, shape, data_type, convert_policy, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), ActivationLayerInfo(), false); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), ActivationLayerInfo(), is_inplace); } }; @@ -159,10 +175,10 @@ class ArithmeticAdditionBroadcastValidationFloatFixture : public ArithmeticOpera { public: template - void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type, ConvertPolicy convert_policy, ActivationLayerInfo act_info) + void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type, ConvertPolicy convert_policy, ActivationLayerInfo act_info, bool is_inplace) { ArithmeticOperationGenericFixture::setup(reference::ArithmeticOperation::ADD, shape0, shape1, data_type, convert_policy, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info, false); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info, is_inplace); } }; @@ -171,10 +187,10 @@ class ArithmeticAdditionValidationFloatFixture : public ArithmeticOperationGener { public: template - void setup(const TensorShape &shape, DataType data_type, ConvertPolicy convert_policy, ActivationLayerInfo act_info) + void setup(const TensorShape &shape, DataType data_type, ConvertPolicy convert_policy, ActivationLayerInfo act_info, bool is_inplace) { ArithmeticOperationGenericFixture::setup(reference::ArithmeticOperation::ADD, shape, shape, data_type, convert_policy, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info, false); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info, is_inplace); } }; @@ -183,11 +199,11 @@ class ArithmeticAdditionValidationQuantizedFixture : public ArithmeticOperationG { public: template - void setup(const TensorShape &shape, DataType data_type, ConvertPolicy convert_policy, QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out) + void setup(const TensorShape &shape, DataType data_type, ConvertPolicy convert_policy, QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out, bool is_inplace) { ArithmeticOperationGenericFixture::setup(reference::ArithmeticOperation::ADD, shape, shape, data_type, convert_policy, - qinfo0, qinfo1, qinfo_out, ActivationLayerInfo(), false); + qinfo0, qinfo1, qinfo_out, ActivationLayerInfo(), is_inplace); } }; @@ -196,10 +212,11 @@ class ArithmeticAdditionValidationQuantizedBroadcastFixture : public ArithmeticO { public: template - void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type, ConvertPolicy convert_policy, QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out) + void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type, ConvertPolicy convert_policy, QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out, + bool is_inplace) { ArithmeticOperationGenericFixture::setup(reference::ArithmeticOperation::ADD, shape0, shape1, data_type, convert_policy, - qinfo0, qinfo1, qinfo_out, ActivationLayerInfo(), false); + qinfo0, qinfo1, qinfo_out, ActivationLayerInfo(), is_inplace); } }; @@ -208,10 +225,10 @@ class ArithmeticSubtractionBroadcastValidationFixture : public ArithmeticOperati { public: template - void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type, ConvertPolicy convert_policy, bool in_place) + void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type, ConvertPolicy convert_policy, bool is_inplace) { ArithmeticOperationGenericFixture::setup(reference::ArithmeticOperation::SUB, shape0, shape1, data_type, convert_policy, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), ActivationLayerInfo(), in_place); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), ActivationLayerInfo(), is_inplace); } }; @@ -221,10 +238,10 @@ class ArithmeticSubtractionBroadcastValidationFloatFixture : public ArithmeticOp public: template void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type, ConvertPolicy convert_policy, ActivationLayerInfo act_info, - bool in_place) + bool is_inplace) { ArithmeticOperationGenericFixture::setup(reference::ArithmeticOperation::SUB, shape0, shape1, data_type, convert_policy, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info, in_place); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info, is_inplace); } }; @@ -233,10 +250,10 @@ class ArithmeticSubtractionValidationFixture : public ArithmeticOperationGeneric { public: template - void setup(const TensorShape &shape, DataType data_type, ConvertPolicy convert_policy, bool in_place) + void setup(const TensorShape &shape, DataType data_type, ConvertPolicy convert_policy, bool is_inplace) { ArithmeticOperationGenericFixture::setup(reference::ArithmeticOperation::SUB, shape, shape, data_type, convert_policy, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), ActivationLayerInfo(), in_place); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), ActivationLayerInfo(), is_inplace); } }; @@ -245,10 +262,10 @@ class ArithmeticSubtractionValidationFloatFixture : public ArithmeticOperationGe { public: template - void setup(const TensorShape &shape, DataType data_type, ConvertPolicy convert_policy, ActivationLayerInfo act_info, bool in_place) + void setup(const TensorShape &shape, DataType data_type, ConvertPolicy convert_policy, ActivationLayerInfo act_info, bool is_inplace) { ArithmeticOperationGenericFixture::setup(reference::ArithmeticOperation::SUB, shape, shape, data_type, convert_policy, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info, in_place); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info, is_inplace); } }; @@ -257,11 +274,11 @@ class ArithmeticSubtractionValidationQuantizedFixture : public ArithmeticOperati { public: template - void setup(const TensorShape &shape, DataType data_type, ConvertPolicy convert_policy, QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out, bool in_place) + void setup(const TensorShape &shape, DataType data_type, ConvertPolicy convert_policy, QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out, bool is_inplace) { ArithmeticOperationGenericFixture::setup(reference::ArithmeticOperation::SUB, shape, shape, data_type, convert_policy, - qinfo0, qinfo1, qinfo_out, ActivationLayerInfo(), in_place); + qinfo0, qinfo1, qinfo_out, ActivationLayerInfo(), is_inplace); } }; @@ -271,10 +288,10 @@ class ArithmeticSubtractionValidationQuantizedBroadcastFixture : public Arithmet public: template void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type, ConvertPolicy convert_policy, QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out, - bool in_place) + bool is_inplace) { ArithmeticOperationGenericFixture::setup(reference::ArithmeticOperation::SUB, shape0, shape1, data_type, convert_policy, - qinfo0, qinfo1, qinfo_out, ActivationLayerInfo(), in_place); + qinfo0, qinfo1, qinfo_out, ActivationLayerInfo(), is_inplace); } }; } // namespace validation diff --git a/tests/validation/fixtures/ElementwiseOperationsFixture.h b/tests/validation/fixtures/ElementwiseOperationsFixture.h index 6661862342..6f31a730ae 100644 --- a/tests/validation/fixtures/ElementwiseOperationsFixture.h +++ b/tests/validation/fixtures/ElementwiseOperationsFixture.h @@ -49,12 +49,13 @@ public: template void setup(ArithmeticOperation op, const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, - QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out, bool in_place = false, bool use_dynamic_shape = false) + QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out, bool is_inplace = false, bool use_dynamic_shape = false) { _op = op; _use_dynamic_shape = use_dynamic_shape; + _is_inplace = is_inplace; - _target = compute_target(shape0, shape1, data_type0, data_type1, output_data_type, qinfo0, qinfo1, qinfo_out, in_place); + _target = compute_target(shape0, shape1, data_type0, data_type1, output_data_type, qinfo0, qinfo1, qinfo_out); _reference = compute_reference(shape0, shape1, data_type0, data_type1, output_data_type, qinfo0, qinfo1, qinfo_out); } @@ -83,7 +84,7 @@ protected: } TensorType compute_target(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, - QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out, bool in_place = false) + QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out) { // Create tensors const TensorShape out_shape = TensorShape::broadcast_shape(shape0, shape1); @@ -92,17 +93,19 @@ protected: TensorType dst = create_tensor(out_shape, output_data_type, 1, qinfo_out); // Check whether do in-place computation and whether inputs are broadcast compatible - TensorType *actual_dst = &dst; - bool src1_can_in_place = !arm_compute::detail::have_different_dimensions(out_shape, shape0, 0) && (qinfo0 == qinfo_out); - bool src2_can_in_place = !arm_compute::detail::have_different_dimensions(out_shape, shape1, 0) && (qinfo1 == qinfo_out); - bool do_in_place = in_place && out_shape.total_size() != 0 && (src1_can_in_place || src2_can_in_place); - if(do_in_place) + TensorType *actual_dst = &dst; + if(_is_inplace) { - if(src1_can_in_place) + bool src1_is_inplace = !arm_compute::detail::have_different_dimensions(out_shape, shape0, 0) && (qinfo0 == qinfo_out) && (data_type0 == output_data_type); + bool src2_is_inplace = !arm_compute::detail::have_different_dimensions(out_shape, shape1, 0) && (qinfo1 == qinfo_out) && (data_type1 == output_data_type); + bool do_in_place = out_shape.total_size() != 0 && (src1_is_inplace || src2_is_inplace); + ARM_COMPUTE_ASSERT(do_in_place); + + if(src1_is_inplace) { actual_dst = &ref_src1; } - else if(src2_can_in_place) + else { actual_dst = &ref_src2; } @@ -135,8 +138,8 @@ protected: ref_src1.allocator()->allocate(); ref_src2.allocator()->allocate(); - // If in-place computation is not supported, still need to allocate original dst - if(!do_in_place) + // If don't do in-place computation, still need to allocate original dst + if(!_is_inplace) { ARM_COMPUTE_ASSERT(dst.info()->is_resizable()); dst.allocator()->allocate(); @@ -176,6 +179,7 @@ protected: SimpleTensor _reference{}; ArithmeticOperation _op{ ArithmeticOperation::ADD }; bool _use_dynamic_shape{ false }; + bool _is_inplace{ false }; }; // Arithmetic operation fused with activation function @@ -186,12 +190,13 @@ public: template void setup(ArithmeticOperation op, const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, - QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out, ActivationLayerInfo act_info, bool in_place = false) + QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out, ActivationLayerInfo act_info, bool is_inplace = true) { ArithmeticOperationsGenericFixture::setup(op, shape0, shape1, data_type0, data_type1, output_data_type, - qinfo0, qinfo1, qinfo_out, in_place); - _act_info = act_info; + qinfo0, qinfo1, qinfo_out, is_inplace); + _act_info = act_info; + _is_inplace = is_inplace; } protected: @@ -199,26 +204,51 @@ protected: QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out) { // Create tensors - TensorType ref_src1 = create_tensor(shape0, data_type0, 1, qinfo0); - TensorType ref_src2 = create_tensor(shape1, data_type1, 1, qinfo1); - TensorType dst = create_tensor(TensorShape::broadcast_shape(shape0, shape1), output_data_type, 1, qinfo_out); + const TensorShape out_shape = TensorShape::broadcast_shape(shape0, shape1); + TensorType ref_src1 = create_tensor(shape0, data_type0, 1, qinfo0); + TensorType ref_src2 = create_tensor(shape1, data_type1, 1, qinfo1); + TensorType dst = create_tensor(out_shape, output_data_type, 1, qinfo_out); + + // Check whether do in-place computation and whether inputs are broadcast compatible + TensorType *actual_dst = &dst; + if(_is_inplace) + { + bool src1_is_inplace = !arm_compute::detail::have_different_dimensions(out_shape, shape0, 0) && (qinfo0 == qinfo_out) && (data_type0 == output_data_type); + bool src2_is_inplace = !arm_compute::detail::have_different_dimensions(out_shape, shape1, 0) && (qinfo1 == qinfo_out) && (data_type1 == output_data_type); + bool do_in_place = out_shape.total_size() != 0 && (src1_is_inplace || src2_is_inplace); + ARM_COMPUTE_ASSERT(do_in_place); + + if(src1_is_inplace) + { + actual_dst = &ref_src1; + } + else + { + actual_dst = &ref_src2; + } + } // Create and configure function FunctionType elem_op; - elem_op.configure(&ref_src1, &ref_src2, &dst, _act_info); + elem_op.configure(&ref_src1, &ref_src2, actual_dst, _act_info); ARM_COMPUTE_ASSERT(ref_src1.info()->is_resizable()); ARM_COMPUTE_ASSERT(ref_src2.info()->is_resizable()); - ARM_COMPUTE_ASSERT(dst.info()->is_resizable()); // Allocate tensors ref_src1.allocator()->allocate(); ref_src2.allocator()->allocate(); - dst.allocator()->allocate(); + + // If don't do in-place computation, still need to allocate original dst + if(!_is_inplace) + { + ARM_COMPUTE_ASSERT(dst.info()->is_resizable()); + dst.allocator()->allocate(); + ARM_COMPUTE_ASSERT(!dst.info()->is_resizable()); + } ARM_COMPUTE_ASSERT(!ref_src1.info()->is_resizable()); ARM_COMPUTE_ASSERT(!ref_src2.info()->is_resizable()); - ARM_COMPUTE_ASSERT(!dst.info()->is_resizable()); // Fill tensors fill(AccessorType(ref_src1), 0); @@ -227,7 +257,7 @@ protected: // Compute function elem_op.run(); - return dst; + return std::move(*actual_dst); } SimpleTensor compute_reference(const TensorShape &shape0, const TensorShape &shape1, @@ -240,6 +270,7 @@ protected: } ActivationLayerInfo _act_info{}; + bool _is_inplace{ false }; }; template @@ -247,11 +278,11 @@ class ArithmeticDivisionBroadcastValidationFixture : public ArithmeticOperations { public: template - void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type) + void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, bool is_inplace) { ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::DIV, shape0, shape1, data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), true); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), is_inplace); } }; @@ -260,11 +291,11 @@ class ArithmeticDivisionValidationFixture : public ArithmeticOperationsGenericFi { public: template - void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type) + void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, bool is_inplace) { ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::DIV, shape, shape, data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo()); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), is_inplace); } }; @@ -273,11 +304,11 @@ class ArithmeticDivisionBroadcastDynamicShapeValidationFixture : public Arithmet { public: template - void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type) + void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, bool is_inplace) { ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::DIV, shape0, shape1, data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), true, true); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), is_inplace, true); } }; @@ -286,11 +317,11 @@ class ArithmeticDivisionDynamicShapeValidationFixture : public ArithmeticOperati { public: template - void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type) + void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, bool is_inplace) { ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::DIV, shape, shape, data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), true); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), is_inplace); } }; @@ -299,11 +330,11 @@ class ArithmeticDivisionBroadcastValidationFloatFixture : public ArithmeticOpera { public: template - void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info) + void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info, bool is_inplace) { ArithmeticOperationsFuseActivationFixture::setup(ArithmeticOperation::DIV, shape0, shape1, data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info, true); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info, is_inplace); } }; @@ -312,11 +343,11 @@ class ArithmeticDivisionValidationFloatFixture : public ArithmeticOperationsFuse { public: template - void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info) + void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info, bool is_inplace) { ArithmeticOperationsFuseActivationFixture::setup(ArithmeticOperation::DIV, shape, shape, data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info, is_inplace); } }; @@ -325,11 +356,11 @@ class ArithmeticDivisionValidationIntegerFixture : public ArithmeticOperationsFu { public: template - void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info) + void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info, bool is_inplace) { ArithmeticOperationsFuseActivationFixture::setup(ArithmeticOperation::DIV, shape, shape, data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info, is_inplace); } }; @@ -339,12 +370,12 @@ class ArithmeticDivisionValidationQuantizedFixture : public ArithmeticOperations public: template void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, - QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out) + QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out, bool is_inplace) { ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::DIV, shape, shape, data_type0, data_type1, output_data_type, - qinfo0, qinfo1, qinfo_out); + qinfo0, qinfo1, qinfo_out, is_inplace); } }; @@ -353,11 +384,11 @@ class ElementwiseMaxBroadcastValidationFixture : public ArithmeticOperationsGene { public: template - void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type) + void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, bool is_inplace) { ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::MAX, shape0, shape1, data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), true); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), is_inplace); } }; @@ -366,11 +397,11 @@ class ElementwiseMaxValidationFixture : public ArithmeticOperationsGenericFixtur { public: template - void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type) + void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, bool is_inplace) { ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::MAX, shape, shape, data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo()); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), is_inplace); } }; @@ -379,11 +410,11 @@ class ElementwiseMaxBroadcastValidationFloatFixture : public ArithmeticOperation { public: template - void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info) + void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info, bool is_inplace) { ArithmeticOperationsFuseActivationFixture::setup(ArithmeticOperation::MAX, shape0, shape1, data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info, true); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info, is_inplace); } }; @@ -392,11 +423,11 @@ class ElementwiseMaxValidationFloatFixture : public ArithmeticOperationsFuseActi { public: template - void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info) + void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info, bool is_inplace) { ArithmeticOperationsFuseActivationFixture::setup(ArithmeticOperation::MAX, shape, shape, data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info, is_inplace); } }; @@ -406,12 +437,12 @@ class ElementwiseMaxValidationQuantizedFixture : public ArithmeticOperationsGene public: template void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, - QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out) + QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out, bool is_inplace) { ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::MAX, shape, shape, data_type0, data_type1, output_data_type, - qinfo0, qinfo1, qinfo_out); + qinfo0, qinfo1, qinfo_out, is_inplace); } }; @@ -421,12 +452,12 @@ class ElementwiseMaxQuantizedBroadcastValidationFixture : public ArithmeticOpera public: template void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, - QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out) + QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out, bool is_inplace) { ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::MAX, shape0, shape1, data_type0, data_type1, output_data_type, - qinfo0, qinfo1, qinfo_out, true); + qinfo0, qinfo1, qinfo_out, is_inplace); } }; @@ -435,11 +466,11 @@ class ElementwiseMinBroadcastValidationFixture : public ArithmeticOperationsGene { public: template - void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type) + void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, bool is_inplace) { ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::MIN, shape0, shape1, data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), true); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), is_inplace); } }; @@ -448,11 +479,11 @@ class ElementwiseMinValidationFixture : public ArithmeticOperationsGenericFixtur { public: template - void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type) + void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, bool is_inplace) { ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::MIN, shape, shape, data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo()); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), is_inplace); } }; @@ -461,11 +492,11 @@ class ElementwiseMinBroadcastValidationFloatFixture : public ArithmeticOperation { public: template - void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info) + void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info, bool is_inplace) { ArithmeticOperationsFuseActivationFixture::setup(ArithmeticOperation::MIN, shape0, shape1, data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info, true); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info, is_inplace); } }; @@ -474,11 +505,11 @@ class ElementwiseMinValidationFloatFixture : public ArithmeticOperationsFuseActi { public: template - void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info) + void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info, bool is_inplace) { ArithmeticOperationsFuseActivationFixture::setup(ArithmeticOperation::MIN, shape, shape, data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info, is_inplace); } }; @@ -488,12 +519,12 @@ class ElementwiseMinValidationQuantizedFixture : public ArithmeticOperationsGene public: template void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, - QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out) + QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out, bool is_inplace) { ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::MIN, shape, shape, data_type0, data_type1, output_data_type, - qinfo0, qinfo1, qinfo_out); + qinfo0, qinfo1, qinfo_out, is_inplace); } }; @@ -503,12 +534,12 @@ class ElementwiseMinQuantizedBroadcastValidationFixture : public ArithmeticOpera public: template void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, - QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out) + QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out, bool is_inplace) { ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::MIN, shape0, shape1, data_type0, data_type1, output_data_type, - qinfo0, qinfo1, qinfo_out, true); + qinfo0, qinfo1, qinfo_out, is_inplace); } }; @@ -517,11 +548,11 @@ class ElementwiseSquaredDiffBroadcastValidationFixture : public ArithmeticOperat { public: template - void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type) + void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, bool is_inplace) { ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::SQUARED_DIFF, shape0, shape1, data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), true); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), is_inplace); } }; @@ -530,11 +561,11 @@ class ElementwiseSquaredDiffValidationFixture : public ArithmeticOperationsGener { public: template - void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type) + void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, bool is_inplace) { ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::SQUARED_DIFF, shape, shape, data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo()); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), is_inplace); } }; @@ -543,11 +574,11 @@ class ElementwiseSquaredDiffBroadcastValidationFloatFixture : public ArithmeticO { public: template - void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info) + void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info, bool is_inplace) { ArithmeticOperationsFuseActivationFixture::setup(ArithmeticOperation::SQUARED_DIFF, shape0, shape1, data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info, true); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info, is_inplace); } }; @@ -556,11 +587,11 @@ class ElementwiseSquaredDiffValidationFloatFixture : public ArithmeticOperations { public: template - void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info) + void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info, bool is_inplace) { ArithmeticOperationsFuseActivationFixture::setup(ArithmeticOperation::SQUARED_DIFF, shape, shape, data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info, is_inplace); } }; @@ -570,12 +601,12 @@ class ElementwiseSquaredDiffValidationQuantizedFixture : public ArithmeticOperat public: template void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, - QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out) + QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out, bool is_inplace) { ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::SQUARED_DIFF, shape, shape, data_type0, data_type1, output_data_type, - qinfo0, qinfo1, qinfo_out); + qinfo0, qinfo1, qinfo_out, is_inplace); } }; @@ -585,12 +616,12 @@ class ElementwiseSquaredDiffQuantizedBroadcastValidationFixture : public Arithme public: template void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, - QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out) + QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out, bool is_inplace) { ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::SQUARED_DIFF, shape0, shape1, data_type0, data_type1, output_data_type, - qinfo0, qinfo1, qinfo_out, true); + qinfo0, qinfo1, qinfo_out, is_inplace); } }; @@ -603,7 +634,7 @@ public: { ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::PRELU, shape0, shape1, data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), true); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo()); } }; @@ -646,7 +677,7 @@ public: { ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::PRELU, shape0, shape1, data_type0, data_type1, output_data_type, - qinfo0, qinfo1, qinfo_out, true); + qinfo0, qinfo1, qinfo_out); } }; @@ -655,11 +686,11 @@ class ElementwisePowerBroadcastValidationFixture : public ArithmeticOperationsGe { public: template - void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type) + void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, bool is_inplace) { ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::POWER, shape0, shape1, data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), true); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), is_inplace); } }; @@ -668,11 +699,11 @@ class ElementwisePowerValidationFixture : public ArithmeticOperationsGenericFixt { public: template - void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type) + void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, bool is_inplace) { ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::POWER, shape, shape, data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo()); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), is_inplace); } }; @@ -681,11 +712,11 @@ class ElementwisePowerBroadcastValidationFloatFixture : public ArithmeticOperati { public: template - void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info) + void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info, bool is_inplace) { ArithmeticOperationsFuseActivationFixture::setup(ArithmeticOperation::POWER, shape0, shape1, data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info, true); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info, is_inplace); } }; @@ -694,11 +725,11 @@ class ElementwisePowerValidationFloatFixture : public ArithmeticOperationsFuseAc { public: template - void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info) + void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info, bool is_inplace) { ArithmeticOperationsFuseActivationFixture::setup(ArithmeticOperation::POWER, shape, shape, data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info, is_inplace); } }; diff --git a/tests/validation/fixtures/PixelWiseMultiplicationFixture.h b/tests/validation/fixtures/PixelWiseMultiplicationFixture.h index 7c643bd726..8dc5179109 100644 --- a/tests/validation/fixtures/PixelWiseMultiplicationFixture.h +++ b/tests/validation/fixtures/PixelWiseMultiplicationFixture.h @@ -76,9 +76,29 @@ protected: QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out, ActivationLayerInfo act_info) { // Create tensors - TensorType src1 = create_tensor(shape0, dt_in1, 1, qinfo0); - TensorType src2 = create_tensor(shape1, dt_in2, 1, qinfo1); - TensorType dst = create_tensor(TensorShape::broadcast_shape(shape0, shape1), dt_out, 1, qinfo_out); + const TensorShape out_shape = TensorShape::broadcast_shape(shape0, shape1); + TensorType src1 = create_tensor(shape0, dt_in1, 1, qinfo0); + TensorType src2 = create_tensor(shape1, dt_in2, 1, qinfo1); + TensorType dst = create_tensor(out_shape, dt_out, 1, qinfo_out); + + // Check whether do in-place computation and whether inputs are broadcast compatible + TensorType *actual_dst = &dst; + if(_is_inplace) + { + bool src1_is_inplace = !arm_compute::detail::have_different_dimensions(out_shape, shape0, 0) && (qinfo0 == qinfo_out) && (dt_in1 == dt_out); + bool src2_is_inplace = !arm_compute::detail::have_different_dimensions(out_shape, shape1, 0) && (qinfo1 == qinfo_out) && (dt_in2 == dt_out); + bool do_in_place = out_shape.total_size() != 0 && (src1_is_inplace || src2_is_inplace); + ARM_COMPUTE_ASSERT(do_in_place); + + if(src1_is_inplace) + { + actual_dst = &src1; + } + else + { + actual_dst = &src2; + } + } auto allocate_tensor = [](TensorType & t) { @@ -89,11 +109,12 @@ protected: // Create and configure function FunctionType multiply; - multiply.configure(&src1, &src2, (_is_inplace ? &src1 : &dst), scale, convert_policy, rounding_policy, act_info); + multiply.configure(&src1, &src2, actual_dst, scale, convert_policy, rounding_policy, act_info); allocate_tensor(src1); allocate_tensor(src2); + // If don't do in-place computation, still need to allocate original dst if(!_is_inplace) { allocate_tensor(dst); @@ -106,12 +127,7 @@ protected: // Compute function multiply.run(); - if(_is_inplace) - { - return src1; - } - - return dst; + return std::move(*actual_dst); } SimpleTensor compute_reference(const TensorShape &shape0, const TensorShape &shape1, DataType dt_in1, DataType dt_in2, DataType dt_out, @@ -122,16 +138,12 @@ protected: SimpleTensor src1{ shape0, dt_in1, 1, qinfo0 }; SimpleTensor src2{ shape1, dt_in2, 1, qinfo1 }; - // current in-place implementation only supports same metadata of input and output tensors. - // By ignoring output quantization information here, we can make test cases implementation much simpler. - QuantizationInfo output_qinfo = _is_inplace ? qinfo0 : qinfo_out; - // Fill reference fill(src1, 0); fill(src2, 1); - auto result = reference::pixel_wise_multiplication(src1, src2, scale, convert_policy, rounding_policy, dt_out, output_qinfo); - return act_info.enabled() ? reference::activation_layer(result, act_info, output_qinfo) : result; + auto result = reference::pixel_wise_multiplication(src1, src2, scale, convert_policy, rounding_policy, dt_out, qinfo_out); + return act_info.enabled() ? reference::activation_layer(result, act_info, qinfo_out) : result; } TensorType _target{}; -- cgit v1.2.1