diff options
21 files changed, 286 insertions, 519 deletions
diff --git a/arm_compute/core/Helpers.h b/arm_compute/core/Helpers.h index 48ac38b170..90dd6082e1 100644 --- a/arm_compute/core/Helpers.h +++ b/arm_compute/core/Helpers.h @@ -801,16 +801,6 @@ inline T wrap_around(T x, T m) return x >= 0 ? x % m : (x % m + m) % m; } -/** Convert a dimension axis to the number of dimensions in the range [0, @p dim_axis] - * Handle negative axis, negative axis is used to specify axis from the end (e.g. -1 for the last axis). - * - * @param[in] dim_axis The last axis (inclusive) in the range [0, @p dim_axis] - * @param[in] num_dims The total number of dimensions - * - * @return The number of dimensions in the range [0, @p dim_axis] - */ -inline size_t dim_index_2_num_dims(int32_t dim_axis, int32_t num_dims); - /** Convert negative coordinates to positive in the range [0, num_dims_input] * * @param[out] coords Array of coordinates to be converted. @@ -852,6 +842,21 @@ inline unsigned int get_next_power_two(unsigned int x) return x; } + +/** Given a softmax axis, this function returns the permutation vector required to put the axis to the front + * + * @note This function assumes a tensor rank <= 4 + * + * Axis selects the dimension on which softmax is performed. + * E.g. For input of shape 4x5x6 and axis=1, softmax will be applied to 4x6=24 vectors of size 5. + * Interally softmax kernels is always performed on the first dimension (front dimension), therefore permutation is + * required to put the dimension specified by @p axis to the first dimension. + * + * @param[in] axis Axis on which to perform softmax. Supported: 1, 2, 3 (0 implies no permutation needed) + * + * @return the permutation vector + */ +PermutationVector get_permutation_vector_from_softmax_axis(size_t axis); } // namespace arm_compute #include "arm_compute/core/Helpers.inl" diff --git a/arm_compute/core/Helpers.inl b/arm_compute/core/Helpers.inl index df0c929372..5613e8c74e 100644 --- a/arm_compute/core/Helpers.inl +++ b/arm_compute/core/Helpers.inl @@ -29,11 +29,6 @@ namespace arm_compute { -inline size_t dim_index_2_num_dims(int32_t dim_axis, int32_t num_dims) -{ - return static_cast<size_t>(wrap_around(dim_axis, num_dims)) + 1; -} - inline uint8_t pixel_area_c1u8_clamp(const uint8_t *first_pixel_ptr, size_t stride, size_t width, size_t height, float wr, float hr, int x, int y) { ARM_COMPUTE_ERROR_ON(first_pixel_ptr == nullptr); diff --git a/arm_compute/runtime/CL/functions/CLSoftmaxLayer.h b/arm_compute/runtime/CL/functions/CLSoftmaxLayer.h index bb01584ff4..fd71f3ed4d 100644 --- a/arm_compute/runtime/CL/functions/CLSoftmaxLayer.h +++ b/arm_compute/runtime/CL/functions/CLSoftmaxLayer.h @@ -26,8 +26,7 @@ #include "arm_compute/core/CL/kernels/CLSoftmaxLayerKernel.h" #include "arm_compute/runtime/CL/CLTensor.h" -#include "arm_compute/runtime/CL/functions/CLFlattenLayer.h" -#include "arm_compute/runtime/CL/functions/CLReshapeLayer.h" +#include "arm_compute/runtime/CL/functions/CLPermute.h" #include "arm_compute/runtime/IFunction.h" #include "arm_compute/runtime/IMemoryManager.h" #include "arm_compute/runtime/MemoryGroup.h" @@ -47,7 +46,10 @@ class ICLTensor; * @f[ out = (x - max(x) * beta) - log(\sum{e^{x - max(x) * beta}}) @f] * * This function runs the following kernels: + * -# If axis is not 0: + * -# @ref CLPermute * -# @ref CLLogits1DNormKernel + * -# @ref CLLogits1DMaxShiftExpSumKernel */ template <bool IS_LOG = false> class CLSoftmaxLayerGeneric : public IFunction @@ -60,70 +62,47 @@ public: * @param[in] input Source tensor. Data types supported: QASYMM8/QASYMM8_SIGNED/F16/F32 for Softmax and F16/F32 for Log Softmax * @param[out] output Destination tensor. Data types supported: same as @p input * @param[in] beta (Optional) A scaling factor for the exponent. Defaults to 1.f - * @param[in] axis (Optional) The last axis of the first n dimensions (inclusive)to reduce. Only supports axis 0. + * @param[in] axis (Optional) The dimension in which to apply the function. E.g. for input of shape 4x5x6 and + * axis=1, softmax will be applied to 4x6=24 vectors of size 5. Defaults to 0 */ - void configure(const ICLTensor *input, ICLTensor *output, float beta = 1.0f, size_t axis = 0); + void configure(const ICLTensor *input, ICLTensor *output, float beta = 1.0f, int32_t axis = 0); /** Set the input and output tensors. * * @param[in] compile_context The compile context to be used. * @param[in] input Source tensor. Data types supported: QASYMM8/QASYMM8_SIGNED/F16/F32 for Softmax and F16/F32 for Log Softmax * @param[out] output Destination tensor. Data types supported: same as @p input * @param[in] beta (Optional) A scaling factor for the exponent. Defaults to 1.f - * @param[in] axis (Optional) The last axis of the first n dimensions (inclusive)to reduce. Only supports axis 0. + * @param[in] axis (Optional) The dimension in which to apply the function. E.g. for input of shape 4x5x6 and + * axis=1, softmax will be applied to 4x6=24 vectors of size 5. Defaults to 0 */ - void configure(const CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output, float beta = 1.0f, size_t axis = 0); + void configure(const CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output, float beta = 1.0f, int32_t axis = 0); /** Static function to check if given info will lead to a valid configuration of @ref CLSoftmaxLayer * * @param[in] input Source tensor. Data types supported: QASYMM8/QASYMM8_SIGNED/F16/F32 for Softmax and F16/F32 for Log Softmax * @param[in] output Destination tensor. Data types supported: same as @p input * @param[in] beta (Optional) A scaling factor for the exponent. Defaults to 1.f - * @param[in] axis (Optional) The last axis of the first n dimensions (inclusive)to reduce. Only supports axis 0. + * @param[in] axis (Optional) The dimension in which to apply the function. E.g. for input of shape 4x5x6 and + * axis=1, softmax will be applied to 4x6=24 vectors of size 5. Defaults to 0 * * @return a status */ - static Status validate(const ITensorInfo *input, const ITensorInfo *output, float beta = 1.0f, size_t axis = 0); + static Status validate(const ITensorInfo *input, const ITensorInfo *output, float beta = 1.0f, int32_t axis = 0); // Inherited methods overridden: void run() override; private: - /** Utility method to configure the kernels needed to flatten the input - * tensor. - * - * @note This function changes the internal state of this class. In particular, - * it initializes the kernel @p _flatten_kernel and the tensors @p _input_flat and - * @p _output_flat - * - * @param[in] input Original source tensor. - * @param[in] output Original destination tensor. - * @param[in] axis (Optional) The last axis of the first n dimensions (inclusive)to reduce. Only supports axis 0. - */ - void configure_reshape_input_kernel(const ICLTensor *input, const ICLTensor *output, size_t axis); - /** Utility method to configure the kernels needed to flatten the input - * tensor. - * - * @note This function changes the internal state of this class. In particular, - * it initializes the kernel @p _flatten_kernel and the tensors @p _input_flat and - * @p _output_flat - * - * @param[in] compile_context The compile context to be used. - * @param[in] input Original source tensor. - * @param[in] output Original destination tensor. - * @param[in] axis (Optional) The last axis of the first n dimensions (inclusive)to reduce. Only supports axis 0. - */ - void configure_reshape_input_kernel(const CLCompileContext &compile_context, const ICLTensor *input, const ICLTensor *output, size_t axis); - MemoryGroup _memory_group; + CLPermute _permute_input; + CLPermute _permute_output; CLLogits1DMaxShiftExpSumKernel _max_shift_exp_sum_kernel; CLLogits1DNormKernel _norm_kernel; - std::unique_ptr<IFunction> _flatten_ptr; - CLReshapeLayer _reshape; CLTensor _max; CLTensor _sum; CLTensor _tmp; - CLTensor _input_flattened; - CLTensor _output_flattened; - bool _needs_flattening; + CLTensor _input_permuted; + CLTensor _output_permuted; + bool _needs_permute; }; using CLSoftmaxLayer = CLSoftmaxLayerGeneric<false>; diff --git a/arm_compute/runtime/GLES_COMPUTE/functions/GCSoftmaxLayer.h b/arm_compute/runtime/GLES_COMPUTE/functions/GCSoftmaxLayer.h index 4ccfe2684e..0279edf63d 100644 --- a/arm_compute/runtime/GLES_COMPUTE/functions/GCSoftmaxLayer.h +++ b/arm_compute/runtime/GLES_COMPUTE/functions/GCSoftmaxLayer.h @@ -50,17 +50,15 @@ public: GCSoftmaxLayer(std::shared_ptr<IMemoryManager> memory_manager = nullptr); /** Set the input and output tensors. * - * @param[in] input Source tensor. Data types supported: F16/F32 - * @param[out] output Destination tensor. Data types supported: same as @p input - * @param[in] beta (Optional) A scaling factor for the exponent. Only beta = 1 is supported - * @param[in] reduce_end_axis (Optional) The last axis of the first n dimensions (inclusive)to reduce. Defaults to 0. - * It has the purpose of squashing together the first n dimensions till (including) the @p reduce_end_axis. For instance, given a [2x3x4x5] image, - * when @p reduce_end_axis is 1, the reduction will be applied to axes 0 and 1, and the Softmax op will be applied on each of the [2x3] planes of the input image. - * Must be in range [0, input_num_dimensions). + * @param[in] input Source tensor. Data types supported: F16/F32 + * @param[out] output Destination tensor. Data types supported: same as @p input + * @param[in] beta (Optional) A scaling factor for the exponent. Only beta = 1 is supported + * @param[in] axis (Optional) The dimension in which to apply the function. E.g. for input of shape 4x5x6 and + * axis=1, softmax will be applied to 4x6=24 vectors of size 5. Defaults to 0 * - * @note The value of @p reduce_end_axis must be always 0 for GLES + * @note The value of @p axis must be always 0 for GLES */ - void configure(const IGCTensor *input, IGCTensor *output, float beta = 1.0f, size_t reduce_end_axis = 0); + void configure(const IGCTensor *input, IGCTensor *output, float beta = 1.0f, int32_t axis = 0); // Inherited methods overridden: void run() override; diff --git a/arm_compute/runtime/NEON/functions/NESoftmaxLayer.h b/arm_compute/runtime/NEON/functions/NESoftmaxLayer.h index 9fb4d85262..20b20201d5 100644 --- a/arm_compute/runtime/NEON/functions/NESoftmaxLayer.h +++ b/arm_compute/runtime/NEON/functions/NESoftmaxLayer.h @@ -28,8 +28,7 @@ #include "arm_compute/core/NEON/kernels/NESoftmaxLayerKernel.h" #include "arm_compute/runtime/IFunction.h" #include "arm_compute/runtime/MemoryGroup.h" -#include "arm_compute/runtime/NEON/functions/NEFlattenLayer.h" -#include "arm_compute/runtime/NEON/functions/NEReshapeLayer.h" +#include "arm_compute/runtime/NEON/functions/NEPermute.h" #include "arm_compute/runtime/Tensor.h" namespace arm_compute @@ -44,7 +43,9 @@ class ITensor; * Log Softmax is calculated by : * @f[ out = (x - max(x) * beta) - log(\sum{e^{x - max(x) * beta}}) @f] * - * This function runs the following kernels: + * This function runs the following function/kernels: + * -# If axis is not 0: + * -# @ref NEPermute * -# @ref NEFillBorderKernel * -# @ref NELogits1DMaxKernel * -# @ref NELogits1DSoftmaxKernel @@ -70,7 +71,8 @@ public: * last value of each row to the nearest multiple. * @param[out] output Destination tensor. Data types supported: same as @p input. * @param[in] beta (Optional) A scaling factor for the exponent. - * @param[in] axis (Optional) The last axis of the first n dimensions (inclusive)to reduce. Only supports axis 0. + * @param[in] axis (Optional) The dimension in which to apply the function. E.g. for input of shape 4x5x6 and + * axis=1, softmax will be applied to 4x6=24 vectors of size 5. Defaults to 0 */ void configure(ITensor *input, ITensor *output, float beta = 1.0f, int32_t axis = 0); /** Static function to check if given info will lead to a valid configuration of @ref NESoftmaxLayer @@ -78,7 +80,8 @@ public: * @param[in] input Source tensor info. Data types supported: QASYMM8/QASYMM8_SIGNED/F16/F32. * @param[in] output Destination tensor info. Data types supported: same as @p input * @param[in] beta (Optional) A scaling factor for the exponent. - * @param[in] axis (Optional) The last axis of the first n dimensions (inclusive)to reduce. Only supports axis 0. + * @param[in] axis (Optional) The dimension in which to apply the function. E.g. for input of shape 4x5x6 and + * axis=1, softmax will be applied to 4x6=24 vectors of size 5. Defaults to 0 * * @return a status */ @@ -88,30 +91,17 @@ public: void run() override; private: - /** Utility method to configure the kernels needed to flatten the input - * tensor. - * - * @note This function changes the internal state of this class. In particular, - * it initializes the kernel @p _flatten_kernel and the tensors @p _input_flat and - * @p _output_flat - * - * @param[in] input Original source tensor. - * @param[in] output Original destination tensor. - * @param[in] axis (Optional) The last axis of the first n dimensions (inclusive)to reduce. Only supports axis 0. - */ - void configure_reshape_input_kernel(const ITensor *input, const ITensor *output, int32_t axis); - MemoryGroup _memory_group; + NEPermute _permute_input; + NEPermute _permute_output; NELogits1DMaxKernel _max_kernel; NELogits1DSoftmaxKernel<IS_LOG> _softmax_kernel; - std::unique_ptr<IFunction> _flat_or_reshape_ptr; NEFillBorderKernel _fill_border_kernel; - NEReshapeLayer _reshape; Tensor _max; Tensor _tmp; - Tensor _input_flattened; - Tensor _output_flattened; - bool _needs_flattening; + Tensor _input_permuted; + Tensor _output_permuted; + bool _needs_permute; }; using NESoftmaxLayer = NESoftmaxLayerGeneric<false>; diff --git a/docs/00_introduction.dox b/docs/00_introduction.dox index bb1dfec69e..37c39f50a4 100644 --- a/docs/00_introduction.dox +++ b/docs/00_introduction.dox @@ -237,6 +237,18 @@ If there is more than one release in a month then an extra sequential number is @subsection S2_2_changelog Changelog +v20.11 Public major release + - Interface change + - Properly support softmax axis to have the same meaning as other major frameworks. That is, axis now defines the dimension + on which Softmax/Logsoftmax is performed. E.g. for input of shape 4x5x6 and axis=1, softmax will be applied to 4x6=24 vectors of size 5. + The supported value range of axis is [-rank, rank). + This change applies to the following functions: + - @ref NESoftmaxLayer + - @ref NELogSoftmaxLayer + - @ref CLSoftmaxLayer + - @ref CLLogSoftmaxLayer + - @ref GCSoftmaxLayer + v20.08 Public major release - Various bug fixes. - Various optimisations. diff --git a/src/core/CL/kernels/CLPermuteKernel.cpp b/src/core/CL/kernels/CLPermuteKernel.cpp index 1636e5a1bc..dc2d6fe4b4 100644 --- a/src/core/CL/kernels/CLPermuteKernel.cpp +++ b/src/core/CL/kernels/CLPermuteKernel.cpp @@ -75,16 +75,16 @@ void CLPermuteKernel::configure(const ICLTensor *input, ICLTensor *output, const void CLPermuteKernel::configure(const CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output, const PermutationVector &perm) { ARM_COMPUTE_ERROR_ON_NULLPTR(input, output); + const TensorShape output_shape = get_output_shape(input->info(), perm); + // Output auto inizialitation if not yet initialized + auto_init_if_empty(*output->info(), input->info()->clone()->set_tensor_shape(output_shape)); + ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(input->info(), output->info(), perm)); _input = input; _output = output; _perm = perm; - const TensorShape output_shape = get_output_shape(input->info(), perm); - // Output auto inizialitation if not yet initialized - auto_init_if_empty(*output->info(), input->info()->clone()->set_tensor_shape(output_shape)); - // Create kernel CLBuildOptions build_opts; build_opts.add_option("-DDATA_TYPE=" + get_cl_unsigned_type_from_element_size(data_size_from_type(input->info()->data_type()))); diff --git a/src/core/Helpers.cpp b/src/core/Helpers.cpp index bfc4a8d101..5c7200b35c 100644 --- a/src/core/Helpers.cpp +++ b/src/core/Helpers.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2018 Arm Limited. + * Copyright (c) 2016-2020 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -23,9 +23,9 @@ */ #include "arm_compute/core/Helpers.h" -using namespace arm_compute; - -Window arm_compute::calculate_max_window(const ValidRegion &valid_region, const Steps &steps, bool skip_border, BorderSize border_size) +namespace arm_compute +{ +Window calculate_max_window(const ValidRegion &valid_region, const Steps &steps, bool skip_border, BorderSize border_size) { if(!skip_border) { @@ -79,7 +79,7 @@ Window arm_compute::calculate_max_window(const ValidRegion &valid_region, const return window; } -Window arm_compute::calculate_max_enlarged_window(const ValidRegion &valid_region, const Steps &steps, BorderSize border_size) +Window calculate_max_enlarged_window(const ValidRegion &valid_region, const Steps &steps, BorderSize border_size) { const Coordinates &anchor = valid_region.anchor; const TensorShape &shape = valid_region.shape; @@ -128,7 +128,7 @@ Window arm_compute::calculate_max_enlarged_window(const ValidRegion &valid_regio return window; } -Window arm_compute::calculate_max_window_horizontal(const ValidRegion &valid_region, const Steps &steps, bool skip_border, BorderSize border_size) +Window calculate_max_window_horizontal(const ValidRegion &valid_region, const Steps &steps, bool skip_border, BorderSize border_size) { if(skip_border) { @@ -181,8 +181,8 @@ Window arm_compute::calculate_max_window_horizontal(const ValidRegion &valid_reg return window; } -ValidRegion arm_compute::calculate_valid_region_scale(const ITensorInfo &src_info, const TensorShape &dst_shape, - InterpolationPolicy interpolate_policy, SamplingPolicy sampling_policy, bool border_undefined) +ValidRegion calculate_valid_region_scale(const ITensorInfo &src_info, const TensorShape &dst_shape, + InterpolationPolicy interpolate_policy, SamplingPolicy sampling_policy, bool border_undefined) { const DataLayout data_layout = src_info.data_layout(); const int idx_width = get_data_layout_dimension_index(data_layout, DataLayoutDimension::WIDTH); @@ -255,4 +255,20 @@ ValidRegion arm_compute::calculate_valid_region_scale(const ITensorInfo &src_inf valid_region.shape.set(idx_height, std::min<size_t>(valid_end_out_y - valid_start_out_y, dst_shape[idx_height])); return valid_region; -}
\ No newline at end of file +} + +PermutationVector get_permutation_vector_from_softmax_axis(size_t actual_axis) +{ + switch(actual_axis) + { + case 1: + return PermutationVector(1U, 0U, 2U, 3U); + case 2: + return PermutationVector(2U, 1U, 0U, 3U); + case 3: + return PermutationVector(3U, 1U, 2U, 0U); + default: + ARM_COMPUTE_ERROR("Axis not supported"); + } +} +} // namespace arm_compute
\ No newline at end of file diff --git a/src/runtime/CL/functions/CLSoftmaxLayer.cpp b/src/runtime/CL/functions/CLSoftmaxLayer.cpp index f7b2935622..720f9111a5 100644 --- a/src/runtime/CL/functions/CLSoftmaxLayer.cpp +++ b/src/runtime/CL/functions/CLSoftmaxLayer.cpp @@ -36,96 +36,45 @@ namespace arm_compute { template <bool IS_LOG> CLSoftmaxLayerGeneric<IS_LOG>::CLSoftmaxLayerGeneric(std::shared_ptr<IMemoryManager> memory_manager) - : _memory_group(std::move(memory_manager)), _max_shift_exp_sum_kernel(), _norm_kernel(), _flatten_ptr(), _reshape(), _max(), _sum(), _tmp(), _input_flattened(), _output_flattened(), - _needs_flattening(false) + : _memory_group(std::move(memory_manager)), _permute_input(), _permute_output(), _max_shift_exp_sum_kernel(), _norm_kernel(), _max(), _sum(), _tmp(), _input_permuted(), _output_permuted(), + _needs_permute() { } template <bool IS_LOG> -void CLSoftmaxLayerGeneric<IS_LOG>::configure_reshape_input_kernel(const ICLTensor *input, const ICLTensor *output, size_t first_n_reduce_axes) -{ - configure_reshape_input_kernel(CLKernelLibrary::get().get_compile_context(), input, output, first_n_reduce_axes); -} - -template <bool IS_LOG> -void CLSoftmaxLayerGeneric<IS_LOG>::configure_reshape_input_kernel(const CLCompileContext &compile_context, const ICLTensor *input, const ICLTensor *output, size_t first_n_reduce_axes) -{ - // Flatten the input - const TensorShape shape_flatten = misc::shape_calculator::compute_softmax_shape(input->info(), first_n_reduce_axes); - - // Initialize the flat input - _input_flattened.allocator()->init(input->info()->clone()->set_is_resizable(true).reset_padding().set_tensor_shape(shape_flatten)); - - // If we need to flatten the input, we can use CLFlattenKernel or CLReshapeKernel - // If the number of reduced axes is 3 (max dimension), which means collapsing all axes except the batch axis, we use CLFlattenKernel. - // In all other cases we have to use CLReshapeKernel - // Note that the "other cases" include both: - // 1. first_n_reduce_axes < 3: Reduce the first 1 (no need to reduce) or 2 dimensions (inclusive) - // 2. first_n_reduce_axes == 4: Reduce all 4 dimensions. This can only be handled by CLReshapeKernel instead of CLFlattenKernel. - if(first_n_reduce_axes == 3) - { - auto flatten = support::cpp14::make_unique<CLFlattenLayer>(); - flatten->configure(compile_context, input, &_input_flattened); - _flatten_ptr = std::move(flatten); - } - else - { - auto reshape_ptr = support::cpp14::make_unique<CLReshapeLayer>(); - reshape_ptr->configure(compile_context, input, &_input_flattened); - _flatten_ptr = std::move(reshape_ptr); - } - - // We need to init the output tensor here. Indeed, the reshape kernel expects - // both tensors to be already initialized - auto_init_if_empty(*output->info(), *input->info()->clone()); -} - -template <bool IS_LOG> -void CLSoftmaxLayerGeneric<IS_LOG>::configure(const ICLTensor *input, ICLTensor *output, float beta, size_t axis) +void CLSoftmaxLayerGeneric<IS_LOG>::configure(const ICLTensor *input, ICLTensor *output, float beta, int32_t axis) { configure(CLKernelLibrary::get().get_compile_context(), input, output, beta, axis); } template <bool IS_LOG> -void CLSoftmaxLayerGeneric<IS_LOG>::configure(const CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output, float beta, size_t axis) +void CLSoftmaxLayerGeneric<IS_LOG>::configure(const CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output, float beta, int32_t axis) { // Perform validation step ARM_COMPUTE_ERROR_ON_NULLPTR(input, output); ARM_COMPUTE_ERROR_THROW_ON(CLSoftmaxLayerGeneric<IS_LOG>::validate(input->info(), output->info(), beta, axis)); - // Convert reduce-before axis (inclusive) to first n axes to reduce - size_t first_n_reduce_axes = dim_index_2_num_dims(axis, input->info()->num_dimensions()); - - // We only need flattening when the number of axes to reduce is greater than 1 - _needs_flattening = first_n_reduce_axes > 1; + const size_t actual_axis = static_cast<size_t>(wrap_around(axis, static_cast<int32_t>(input->info()->num_dimensions()))); - // If we are dealing with a 4D tensor, we will: - // - Flatten the input, so that we end up with a [width*height*depth] * batches 2D tensor - // - Execute all the pipeline (reduction + normalization) on the flattened tensor - // - Reshape the flattened output into the real output - if(_needs_flattening) + _needs_permute = actual_axis != 0; + ICLTensor *tmp_output = output; + const ICLTensor *tmp_input = _needs_permute ? &_input_permuted : input; + if(_needs_permute) { - // Add to the memory manager _input_flattened - _memory_group.manage(&_input_flattened); - - // Cofigure _flatten_kernel and _input_flattened - configure_reshape_input_kernel(input, output, first_n_reduce_axes); + _memory_group.manage(&_input_permuted); + _memory_group.manage(&_output_permuted); + _permute_input.configure(compile_context, input, &_input_permuted, get_permutation_vector_from_softmax_axis(actual_axis)); + tmp_output = &_output_permuted; } - // We want to deal with a 2D input. Either it is the flattened version of the original input (4D case) - // or it is the original input case (2D case) - const ICLTensor *input_2D = (_needs_flattening ? &_input_flattened : input); - - // Create intermediate tensors shapes - TensorInfo input_info = input_2D->info()->clone()->reset_padding().set_is_resizable(true); - DataType tmp_data_type = is_data_type_quantized_asymmetric(input_2D->info()->data_type()) ? DataType::S32 : input_2D->info()->data_type(); - TensorInfo tensor_info_tmp(input_info.clone()->set_data_type(tmp_data_type)); + // Create intermediate tensors + DataType tmp_data_type = is_data_type_quantized_asymmetric(tmp_input->info()->data_type()) ? DataType::S32 : tmp_input->info()->data_type(); + TensorInfo tensor_info_tmp(tmp_input->info()->clone()->set_data_type(tmp_data_type)); _tmp.allocator()->init(tensor_info_tmp); - - TensorShape max_sum_shape = input_2D->info()->tensor_shape(); + TensorShape max_sum_shape = tmp_input->info()->tensor_shape(); max_sum_shape.set(0, 1); - _max.allocator()->init(input_info.clone()->set_tensor_shape(max_sum_shape)); - _sum.allocator()->init(input_info.clone()->set_tensor_shape(max_sum_shape).set_data_type(tmp_data_type)); + _max.allocator()->init(tmp_input->info()->clone()->set_tensor_shape(max_sum_shape)); + _sum.allocator()->init(tmp_input->info()->clone()->set_tensor_shape(max_sum_shape).set_data_type(tmp_data_type)); // Set GPU target to kernels _max_shift_exp_sum_kernel.set_target(CLScheduler::get().target()); @@ -138,49 +87,43 @@ void CLSoftmaxLayerGeneric<IS_LOG>::configure(const CLCompileContext &compile_co SoftmaxKernelInfo softmax_info; softmax_info.beta = beta; softmax_info.is_log = IS_LOG; - softmax_info.input_data_type = input_2D->info()->data_type(); + softmax_info.input_data_type = tmp_input->info()->data_type(); // Configure kernels - _max_shift_exp_sum_kernel.configure(compile_context, input_2D, &_max, &_tmp, &_sum, softmax_info); - - if(_needs_flattening) - { - // Add to the memory manager _output_flattened - _memory_group.manage(&_output_flattened); - - // The normalization kernel stores the result in a flat output tensor - _norm_kernel.configure(compile_context, &_tmp, &_sum, &_output_flattened, softmax_info); - - // Reshape the flat output into a the requested (4D) output - _reshape.configure(compile_context, &_output_flattened, output); - - // Allocate the intermediate flat tensors - _input_flattened.allocator()->allocate(); - _output_flattened.allocator()->allocate(); - } - else - { - // Softmax 2D case - _norm_kernel.configure(compile_context, &_tmp, &_sum, output, softmax_info); - } + _max_shift_exp_sum_kernel.configure(compile_context, tmp_input, &_max, &_tmp, &_sum, softmax_info); + _norm_kernel.configure(compile_context, &_tmp, &_sum, tmp_output, softmax_info); // Allocate intermediate buffers _tmp.allocator()->allocate(); _max.allocator()->allocate(); _sum.allocator()->allocate(); + if(_needs_permute) + { + _permute_output.configure(compile_context, &_output_permuted, output, get_permutation_vector_from_softmax_axis(actual_axis)); + _input_permuted.allocator()->allocate(); + _output_permuted.allocator()->allocate(); + } } template <bool IS_LOG> -Status CLSoftmaxLayerGeneric<IS_LOG>::validate(const ITensorInfo *input, const ITensorInfo *output, float beta, size_t axis) +Status CLSoftmaxLayerGeneric<IS_LOG>::validate(const ITensorInfo *input, const ITensorInfo *output, float beta, int32_t axis) { ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input, output); ARM_COMPUTE_RETURN_ERROR_ON_MSG(input->num_dimensions() > 4, "Only up to 4 dimensions are supported"); - ARM_COMPUTE_RETURN_ERROR_ON_MSG(axis != 0, "Only axis 0 supported in tensors"); ARM_COMPUTE_UNUSED(beta); - ARM_COMPUTE_RETURN_ERROR_ON(input->num_dimensions() <= axis); + ARM_COMPUTE_RETURN_ERROR_ON(axis < static_cast<int32_t>(-input->num_dimensions()) || static_cast<int32_t>(input->num_dimensions()) <= axis); - // Convert reduce-before axis (inclusive) to first n axes to reduce - size_t first_n_reduce_axes = dim_index_2_num_dims(axis, input->num_dimensions()); + const size_t actual_axis = static_cast<size_t>(wrap_around(axis, static_cast<int32_t>(input->num_dimensions()))); + const bool needs_permute = actual_axis != 0; + if(needs_permute) + { + const PermutationVector permutation_vector = get_permutation_vector_from_softmax_axis(actual_axis); + const TensorShape permuted_shape = misc::shape_calculator::compute_permutation_output_shape(*input, permutation_vector); + TensorInfo input_permuted(input->clone()->set_tensor_shape(permuted_shape)); + ARM_COMPUTE_RETURN_ON_ERROR(CLPermute::validate(input, &input_permuted, permutation_vector)); + TensorInfo output_permuted(output->clone()->set_tensor_shape(permuted_shape)); + ARM_COMPUTE_RETURN_ON_ERROR(CLPermute::validate(&output_permuted, output, permutation_vector)); + } // Create intermediate tensor info DataType tmp_data_type = is_data_type_quantized_asymmetric(input->data_type()) ? DataType::S32 : input->data_type(); @@ -191,23 +134,6 @@ Status CLSoftmaxLayerGeneric<IS_LOG>::validate(const ITensorInfo *input, const I TensorInfo tensor_info_max(input->clone()->set_tensor_shape(max_sum_shape).set_is_resizable(true)); TensorInfo tensor_info_sum(input->clone()->set_tensor_shape(max_sum_shape).set_data_type(tmp_data_type).set_quantization_info(QuantizationInfo()).set_is_resizable(true)); - const bool needs_flattening = (first_n_reduce_axes > 1); - - if(needs_flattening) - { - const TensorShape shape_flatten = misc::shape_calculator::compute_softmax_shape(input, first_n_reduce_axes); - TensorInfo tensor_info_flat(input->clone()->set_tensor_shape(shape_flatten).set_is_resizable(true)); - - if(first_n_reduce_axes == 3) - { - ARM_COMPUTE_RETURN_ON_ERROR(CLFlattenLayer::validate(input, &tensor_info_flat)); - } - else - { - ARM_COMPUTE_RETURN_ON_ERROR(CLReshapeLayer::validate(input, &tensor_info_flat)); - } - } - SoftmaxKernelInfo softmax_info; softmax_info.beta = beta; softmax_info.is_log = IS_LOG; @@ -216,12 +142,6 @@ Status CLSoftmaxLayerGeneric<IS_LOG>::validate(const ITensorInfo *input, const I ARM_COMPUTE_RETURN_ON_ERROR(CLLogits1DMaxShiftExpSumKernel::validate(input, &tensor_info_max, &tensor_info_tmp, &tensor_info_sum)); ARM_COMPUTE_RETURN_ON_ERROR(CLLogits1DNormKernel::validate(&tensor_info_tmp, &tensor_info_sum, output, softmax_info)); - if(needs_flattening) - { - const TensorShape shape_flatten = misc::shape_calculator::compute_softmax_shape(input); - TensorInfo tensor_info_flat(input->clone()->set_tensor_shape(shape_flatten).set_is_resizable(true)); - } - return Status{}; } @@ -230,17 +150,17 @@ void CLSoftmaxLayerGeneric<IS_LOG>::run() { MemoryGroupResourceScope scope_mg(_memory_group); - if(_needs_flattening) + if(_needs_permute) { - _flatten_ptr->run(); + _permute_input.run(); } CLScheduler::get().enqueue(_max_shift_exp_sum_kernel, false); - CLScheduler::get().enqueue(_norm_kernel, !_needs_flattening); + CLScheduler::get().enqueue(_norm_kernel, !_needs_permute); - if(_needs_flattening) + if(_needs_permute) { - _reshape.run(); + _permute_output.run(); } } diff --git a/src/runtime/GLES_COMPUTE/functions/GCSoftmaxLayer.cpp b/src/runtime/GLES_COMPUTE/functions/GCSoftmaxLayer.cpp index 48d8cb576b..fdb9a42f13 100644 --- a/src/runtime/GLES_COMPUTE/functions/GCSoftmaxLayer.cpp +++ b/src/runtime/GLES_COMPUTE/functions/GCSoftmaxLayer.cpp @@ -34,13 +34,13 @@ GCSoftmaxLayer::GCSoftmaxLayer(std::shared_ptr<IMemoryManager> memory_manager) { } -void GCSoftmaxLayer::configure(const IGCTensor *input, IGCTensor *output, float beta, size_t reduce_end_axis) +void GCSoftmaxLayer::configure(const IGCTensor *input, IGCTensor *output, float beta, int32_t axis) { - ARM_COMPUTE_UNUSED(beta, reduce_end_axis); + ARM_COMPUTE_UNUSED(beta, axis); ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F16, DataType::F32); ARM_COMPUTE_ERROR_ON(beta != 1.0f); - ARM_COMPUTE_ERROR_ON_MSG(reduce_end_axis != 0, "Reduce_end_axis must be 0 for GLES"); + ARM_COMPUTE_ERROR_ON_MSG(axis != 0, "axis must be 0 for GLES"); // Create intermediate tensors shapes _tmp.allocator()->init(TensorInfo(input->info()->tensor_shape(), input->info()->num_channels(), input->info()->data_type())); diff --git a/src/runtime/NEON/functions/NESoftmaxLayer.cpp b/src/runtime/NEON/functions/NESoftmaxLayer.cpp index 750992fca6..e763caa3a3 100644 --- a/src/runtime/NEON/functions/NESoftmaxLayer.cpp +++ b/src/runtime/NEON/functions/NESoftmaxLayer.cpp @@ -32,78 +32,41 @@ namespace arm_compute { template <bool IS_LOG> NESoftmaxLayerGeneric<IS_LOG>::NESoftmaxLayerGeneric(std::shared_ptr<IMemoryManager> memory_manager) - : _memory_group(std::move(memory_manager)), _max_kernel(), _softmax_kernel(), _flat_or_reshape_ptr(nullptr), _fill_border_kernel(), _reshape(), _max(), _tmp(), _input_flattened(), _output_flattened(), - _needs_flattening(false) + : _memory_group(std::move(memory_manager)), _permute_input(), _permute_output(), _max_kernel(), _softmax_kernel(), _fill_border_kernel(), _max(), _tmp(), _input_permuted(), _output_permuted(), + _needs_permute(false) { } template <bool IS_LOG> -void NESoftmaxLayerGeneric<IS_LOG>::configure_reshape_input_kernel(const ITensor *input, const ITensor *output, int32_t first_n_reduce_axes) -{ - // Flatten the input - const TensorShape shape_flatten = misc::shape_calculator::compute_softmax_shape(input->info(), first_n_reduce_axes); - - // Initialize the flat input - _input_flattened.allocator()->init(input->info()->clone()->set_is_resizable(true).reset_padding().set_tensor_shape(shape_flatten)); - - // Note that the "other cases" include both: - // 1. first_n_reduce_axes < 3: Reduce the first 1 (no need to reduce) or 2 dimensions (inclusive) - // 2. first_n_reduce_axes == 4: Reduce all 4 dimensions. This can only be handled by NEReshapeKernel instead of NEFlattenKernel. - if(first_n_reduce_axes == 3) - { - auto flatten_kernel_ptr = support::cpp14::make_unique<NEFlattenLayer>(); - flatten_kernel_ptr->configure(input, &_input_flattened); - _flat_or_reshape_ptr = std::move(flatten_kernel_ptr); - } - else - { - auto reshape_kernel_ptr = support::cpp14::make_unique<NEReshapeLayer>(); - reshape_kernel_ptr->configure(input, &_input_flattened); - _flat_or_reshape_ptr = std::move(reshape_kernel_ptr); - } - - // We need to init the output tensor here. Indeed, the reshape kernel expects - // both tensors to be already initialized - auto_init_if_empty(*output->info(), *input->info()->clone()); -} - -template <bool IS_LOG> void NESoftmaxLayerGeneric<IS_LOG>::configure(ITensor *input, ITensor *output, float beta, int32_t axis) { // Perform validation step ARM_COMPUTE_ERROR_ON_NULLPTR(input, output); ARM_COMPUTE_ERROR_THROW_ON(NESoftmaxLayerGeneric::validate(input->info(), output->info(), beta, axis)); - // Convert reduce-before axis (inclusive) to first n axes to reduce - size_t first_n_reduce_axes = dim_index_2_num_dims(axis, static_cast<int32_t>(input->info()->num_dimensions())); + const unsigned int actual_axis = static_cast<unsigned int>(wrap_around(axis, static_cast<int32_t>(input->info()->num_dimensions()))); - // We only need flattening when the number of axes to reduce is greater than 1 - _needs_flattening = first_n_reduce_axes > 1; + _needs_permute = actual_axis > 0; - // If we are dealing with a 4D tensor, we will: - // - Flatten the input, so that we end up with a [width*height*depth] * batches 2D tensor - // - Execute all the pipeline (reduction + normalization) on the flattened tensor - // - Reshape the flattened output into the real output - if(_needs_flattening) + if(_needs_permute) { - // Add to the memory manager _input_flattened - _memory_group.manage(&_input_flattened); + // Add to the memory manager _input_permuted + _memory_group.manage(&_input_permuted); - // Configure _flatten_kernel and _input_flattened - configure_reshape_input_kernel(input, output, first_n_reduce_axes); + _permute_input.configure(input, &_input_permuted, get_permutation_vector_from_softmax_axis(actual_axis)); } - // We want to deal with a 2D input. Either it is the flattened version of the original input (4D case) + // We want to deal with a 2D input. Either it is the permuted version of the original input (4D case) // or it is the original input case (2D case) - ITensor *input_2D = (_needs_flattening ? &_input_flattened : input); + ITensor *tmp_input = (_needs_permute ? &_input_permuted : input); // Create intermediate tensors shapes - const TensorInfo input_info = input_2D->info()->clone()->reset_padding().set_is_resizable(true); - DataType tmp_data_type = is_data_type_quantized_asymmetric(input_2D->info()->data_type()) ? DataType::F32 : input_2D->info()->data_type(); + const TensorInfo input_info = tmp_input->info()->clone()->reset_padding().set_is_resizable(true); + DataType tmp_data_type = is_data_type_quantized_asymmetric(tmp_input->info()->data_type()) ? DataType::F32 : tmp_input->info()->data_type(); TensorInfo tensor_info_tmp(input_info.clone()->set_data_type(tmp_data_type)); // Init intermediate tensors - TensorShape max_sum_shape = input_2D->info()->tensor_shape(); + TensorShape max_sum_shape = tmp_input->info()->tensor_shape(); max_sum_shape.set(0, 1); _max.allocator()->init(input_info.clone()->set_tensor_shape(max_sum_shape)); _tmp.allocator()->init(tensor_info_tmp); @@ -113,27 +76,27 @@ void NESoftmaxLayerGeneric<IS_LOG>::configure(ITensor *input, ITensor *output, f _memory_group.manage(&_tmp); // Configure Kernels - _max_kernel.configure(input_2D, &_max); - if(_needs_flattening) + _max_kernel.configure(tmp_input, &_max); + if(_needs_permute) { - // Add to the memory manager _output_flattened - _memory_group.manage(&_output_flattened); + // Add to the memory manager _output_permuted + _memory_group.manage(&_output_permuted); - // The normalization kernel stores the result in a flat output tensor - _softmax_kernel.configure(input_2D, &_max, &_output_flattened, beta, &_tmp); - _input_flattened.allocator()->allocate(); + // The normalization kernel stores the result in a permuted output tensor + _softmax_kernel.configure(tmp_input, &_max, &_output_permuted, beta, &_tmp); + _input_permuted.allocator()->allocate(); - // Reshape the flat output into the requested (4D) output - _reshape.configure(&_output_flattened, output); + // Re-permute the permuted output into the requested (4D) output + _permute_output.configure(&_output_permuted, output, get_permutation_vector_from_softmax_axis(actual_axis)); - // Allocate the intermediate flat tensors - _output_flattened.allocator()->allocate(); + // Allocate the intermediate permuted tensors + _output_permuted.allocator()->allocate(); } else { // Softmax 2D case - _fill_border_kernel.configure(input_2D, _max_kernel.border_size(), BorderMode::REPLICATE); - _softmax_kernel.configure(input_2D, &_max, output, beta, &_tmp); + _fill_border_kernel.configure(tmp_input, _max_kernel.border_size(), BorderMode::REPLICATE); + _softmax_kernel.configure(tmp_input, &_max, output, beta, &_tmp); } // Allocate intermediate buffers @@ -148,12 +111,8 @@ Status NESoftmaxLayerGeneric<IS_LOG>::validate(const ITensorInfo *input, const I ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input, output); ARM_COMPUTE_RETURN_ERROR_ON_MSG(input->num_dimensions() > 4, "Only up to 4 dimensions are supported"); ARM_COMPUTE_UNUSED(beta); - ARM_COMPUTE_RETURN_ERROR_ON_MSG(axis != 0, "Only axis 0 supported"); ARM_COMPUTE_RETURN_ERROR_ON(axis < static_cast<int32_t>(-input->num_dimensions()) || static_cast<int32_t>(input->num_dimensions()) <= axis); - // Convert reduce-before axis (inclusive) to first n axes to reduce - size_t first_n_reduce_axes = dim_index_2_num_dims(axis, static_cast<int32_t>(input->num_dimensions())); - // Create intermediate tensor info DataType tmp_data_type = input->data_type(); const TensorInfo tensor_info_tmp(input->clone()->set_data_type(tmp_data_type).set_is_resizable(true)); @@ -163,21 +122,18 @@ Status NESoftmaxLayerGeneric<IS_LOG>::validate(const ITensorInfo *input, const I const TensorInfo tensor_info_max_sum(input->clone()->set_tensor_shape(max_sum_shape).set_data_type(tmp_data_type).set_quantization_info(input->quantization_info()).set_is_resizable(true)); const TensorInfo dont_care; - const bool needs_flattening = (first_n_reduce_axes > 1); + const unsigned int actual_axis = static_cast<unsigned int>(wrap_around(axis, static_cast<int32_t>(input->num_dimensions()))); + + const bool needs_permute = actual_axis > 0; - if(needs_flattening) + if(needs_permute) { - const TensorShape shape_flatten = misc::shape_calculator::compute_softmax_shape(input, first_n_reduce_axes); - TensorInfo tensor_info_flat(input->clone()->set_tensor_shape(shape_flatten).set_is_resizable(true)); - - if(first_n_reduce_axes == 3) - { - ARM_COMPUTE_RETURN_ON_ERROR(NEFlattenLayer::validate(input, &tensor_info_flat)); - } - else - { - ARM_COMPUTE_RETURN_ON_ERROR(NEReshapeLayer::validate(input, &tensor_info_flat)); - } + const PermutationVector permutation_vector = get_permutation_vector_from_softmax_axis(actual_axis); + const TensorShape permuted_shape = misc::shape_calculator::compute_permutation_output_shape(*input, permutation_vector); + TensorInfo input_permuted(input->clone()->set_tensor_shape(permuted_shape)); + ARM_COMPUTE_RETURN_ON_ERROR(NEPermute::validate(input, &input_permuted, permutation_vector)); + TensorInfo output_permuted(output->clone()->set_tensor_shape(permuted_shape)); + ARM_COMPUTE_RETURN_ON_ERROR(NEPermute::validate(&output_permuted, output, permutation_vector)); } ARM_COMPUTE_RETURN_ON_ERROR(NELogits1DMaxKernel::validate(input, &tensor_info_max_sum)); @@ -191,18 +147,18 @@ void NESoftmaxLayerGeneric<IS_LOG>::run() { MemoryGroupResourceScope scope_mg(_memory_group); - if(_needs_flattening) + if(_needs_permute) { - _flat_or_reshape_ptr->run(); + _permute_input.run(); } NEScheduler::get().schedule(&_fill_border_kernel, Window::DimY); NEScheduler::get().schedule(&_max_kernel, Window::DimY); NEScheduler::get().schedule(&_softmax_kernel, Window::DimY); - if(_needs_flattening) + if(_needs_permute) { - _reshape.run(); + _permute_output.run(); } } diff --git a/tests/validation/CL/LogSoftmaxLayer.cpp b/tests/validation/CL/LogSoftmaxLayer.cpp index 15466affc4..8fdc745d13 100644 --- a/tests/validation/CL/LogSoftmaxLayer.cpp +++ b/tests/validation/CL/LogSoftmaxLayer.cpp @@ -59,7 +59,7 @@ TEST_SUITE(FP16) FIXTURE_DATA_TEST_CASE(RunSmall, CLLogSoftmaxLayerFixture<half>, framework::DatasetMode::ALL, combine(combine(combine(datasets::SoftmaxLayerSmallShapes(), framework::dataset::make("DataType", DataType::F16)), framework::dataset::make("Beta", { 1.0f, 2.0f })), - framework::dataset::make("Axis", { 0 }))) + framework::dataset::make("Axis", { 0, -1 }))) { // Validate output validate(CLAccessor(_target), _reference, tolerance_f16); @@ -75,7 +75,7 @@ FIXTURE_DATA_TEST_CASE(RunLarge, CLLogSoftmaxLayerFixture<half>, framework::Data FIXTURE_DATA_TEST_CASE(Run4D, CLLogSoftmaxLayerFixture<half>, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::SoftmaxLayer4DShapes(), framework::dataset::make("DataType", DataType::F16)), framework::dataset::make("Beta", { 1.0f, 2.0f })), - framework::dataset::make("Axis", { 0 }))) + framework::dataset::make("Axis", { 0, -3, 2 }))) { // Validate output validate(CLAccessor(_target), _reference, tolerance_f16); @@ -86,7 +86,7 @@ TEST_SUITE(FP32) FIXTURE_DATA_TEST_CASE(RunSmall, CLLogSoftmaxLayerFixture<float>, framework::DatasetMode::ALL, combine(combine(combine(datasets::SoftmaxLayerSmallShapes(), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("Beta", { 1.0f, 2.0f })), - framework::dataset::make("Axis", { 0 }))) + framework::dataset::make("Axis", { 0, 1 }))) { // Validate output validate(CLAccessor(_target), _reference, tolerance_f32); @@ -99,10 +99,10 @@ FIXTURE_DATA_TEST_CASE(RunLarge, CLLogSoftmaxLayerFixture<float>, framework::Dat // Validate output validate(CLAccessor(_target), _reference, tolerance_f32); } -FIXTURE_DATA_TEST_CASE(Run4D, CLLogSoftmaxLayerFixture<float>, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::SoftmaxLayer4DShapes(), - framework::dataset::make("DataType", DataType::F32)), - framework::dataset::make("Beta", { 1.0f, 2.0f })), - framework::dataset::make("Axis", { 0 }))) +FIXTURE_DATA_TEST_CASE(Run4D, CLLogSoftmaxLayerFixture<float>, framework::DatasetMode::ALL, combine(combine(combine(datasets::SoftmaxLayer4DShapes(), + framework::dataset::make("DataType", DataType::F32)), + framework::dataset::make("Beta", { 1.0f, 2.0f })), + framework::dataset::make("Axis", { 0, -4, 3 }))) { // Validate output validate(CLAccessor(_target), _reference, tolerance_f32); diff --git a/tests/validation/CL/SoftmaxLayer.cpp b/tests/validation/CL/SoftmaxLayer.cpp index 90c3058c5d..0e4952b02a 100644 --- a/tests/validation/CL/SoftmaxLayer.cpp +++ b/tests/validation/CL/SoftmaxLayer.cpp @@ -129,9 +129,9 @@ DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip(zip( QuantizationInfo(1.f/256, 0)), TensorInfo(TensorShape(32U, 13U), 1, DataType::QASYMM8_SIGNED, QuantizationInfo(1.f/256, -128)), - TensorInfo(TensorShape(32U, 13U), 1, DataType::QASYMM8_SIGNED, // Invalid axis high + TensorInfo(TensorShape(32U, 13U), 1, DataType::QASYMM8_SIGNED, QuantizationInfo(1.f/256, -128)), - TensorInfo(TensorShape(32U, 13U), 1, DataType::QASYMM8_SIGNED, // Invalid axis low + TensorInfo(TensorShape(32U, 13U), 1, DataType::QASYMM8_SIGNED, QuantizationInfo(1.f/256, -128)), })), framework::dataset::make("beta", { 1.0, @@ -151,11 +151,11 @@ DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip(zip( 0, 0, 0, + 1, 0, - 0, - 0, - 2, -1, + 2, + -3, })), framework::dataset::make("Expected", { false, false, false, false, false, true, true, true, false, false })), input_info, output_info, beta, axis, expected) @@ -173,7 +173,7 @@ TEST_SUITE(FP16) FIXTURE_DATA_TEST_CASE(RunSmall, CLSoftmaxLayerFixture<half>, framework::DatasetMode::ALL, combine(combine(combine(datasets::SoftmaxLayerSmallShapes(), framework::dataset::make("DataType", DataType::F16)), framework::dataset::make("Beta", { 1.0f, 2.0f })), - framework::dataset::make("Axis", { 0 }))) + framework::dataset::make("Axis", { 0, -1 }))) { // Validate output validate(CLAccessor(_target), _reference, tolerance_f16); @@ -189,7 +189,7 @@ FIXTURE_DATA_TEST_CASE(RunLarge, CLSoftmaxLayerFixture<half>, framework::Dataset FIXTURE_DATA_TEST_CASE(Run4D, CLSoftmaxLayerFixture<half>, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::SoftmaxLayer4DShapes(), framework::dataset::make("DataType", DataType::F16)), framework::dataset::make("Beta", { 1.0f, 2.0f })), - framework::dataset::make("Axis", { 0 }))) + framework::dataset::make("Axis", { 0, -1, 2 }))) { // Validate output validate(CLAccessor(_target), _reference, tolerance_f16); @@ -200,7 +200,7 @@ TEST_SUITE(FP32) FIXTURE_DATA_TEST_CASE(RunSmall, CLSoftmaxLayerFixture<float>, framework::DatasetMode::ALL, combine(combine(combine(datasets::SoftmaxLayerSmallShapes(), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("Beta", { 1.0f, 2.0f })), - framework::dataset::make("Axis", { 0 }))) + framework::dataset::make("Axis", { 0, 1 }))) { // Validate output validate(CLAccessor(_target), _reference, tolerance_f32); @@ -216,7 +216,7 @@ FIXTURE_DATA_TEST_CASE(RunLarge, CLSoftmaxLayerFixture<float>, framework::Datase FIXTURE_DATA_TEST_CASE(Run4D, CLSoftmaxLayerFixture<float>, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::SoftmaxLayer4DShapes(), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("Beta", { 1.0f, 2.0f })), - framework::dataset::make("Axis", { 0 }))) + framework::dataset::make("Axis", { 0, -2, 3 }))) { // Validate output validate(CLAccessor(_target), _reference, tolerance_f32); @@ -233,7 +233,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLSoftmaxLayerQuantizedFixture<uint8_t>, framew framework::dataset::make("DataType", DataType::QASYMM8)), combine(framework::dataset::make("QuantizationInfo", { QuantizationInfo(0.5f, -10) }), framework::dataset::make("Beta", { 1.0f, 2.f }))), - framework::dataset::make("Axis", { 0 }))) + framework::dataset::make("Axis", { 0, 1 }))) { // Validate output validate(CLAccessor(_target), _reference, tolerance_qasymm8); @@ -251,7 +251,7 @@ FIXTURE_DATA_TEST_CASE(Run4D, CLSoftmaxLayerQuantizedFixture<uint8_t>, framework framework::dataset::make("DataType", DataType::QASYMM8)), combine(framework::dataset::make("QuantizationInfo", { QuantizationInfo(0.5f, -10) }), framework::dataset::make("Beta", { 1.0f, 2.0f }))), - framework::dataset::make("Axis", { 0 }))) + framework::dataset::make("Axis", { 0, -4, 1 }))) { // Validate output validate(CLAccessor(_target), _reference, tolerance_qasymm8); @@ -265,7 +265,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLSoftmaxLayerQuantizedFixture<int8_t>, framewo framework::dataset::make("DataType", DataType::QASYMM8_SIGNED)), combine(framework::dataset::make("QuantizationInfo", { QuantizationInfo(0.5f, -10) }), framework::dataset::make("Beta", { 1.0f, 2.f }))), - framework::dataset::make("Axis", { 0 }))) + framework::dataset::make("Axis", { 0, 1 }))) { // Validate output validate(CLAccessor(_target), _reference, tolerance_qasymm8_signed); diff --git a/tests/validation/GLES_COMPUTE/SoftmaxLayer.cpp b/tests/validation/GLES_COMPUTE/SoftmaxLayer.cpp index af92cff813..2a8a33eaf7 100644 --- a/tests/validation/GLES_COMPUTE/SoftmaxLayer.cpp +++ b/tests/validation/GLES_COMPUTE/SoftmaxLayer.cpp @@ -89,7 +89,7 @@ TEST_SUITE(FP16) FIXTURE_DATA_TEST_CASE(RunSmall, GCSoftmaxLayerFixture<half_float::half>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SoftmaxLayerSmallShapes(), framework::dataset::make("DataType", DataType::F16)), framework::dataset::make("Beta", 1.0f)), - framework::dataset::make("ReduceEndAxis", 0))) + framework::dataset::make("Axis", 0))) { // Validate output validate(GCAccessor(_target), _reference, tolerance_f16); @@ -97,18 +97,18 @@ FIXTURE_DATA_TEST_CASE(RunSmall, GCSoftmaxLayerFixture<half_float::half>, framew FIXTURE_DATA_TEST_CASE(RunLarge, GCSoftmaxLayerFixture<half_float::half>, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::SoftmaxLayerLargeShapes(), framework::dataset::make("DataType", DataType::F16)), framework::dataset::make("Beta", 1.0f)), - framework::dataset::make("ReduceEndAxis", 0))) + framework::dataset::make("Axis", 0))) { // Validate output validate(GCAccessor(_target), _reference, tolerance_f16); } -TEST_SUITE_END() +TEST_SUITE_END() // FP16 TEST_SUITE(FP32) FIXTURE_DATA_TEST_CASE(RunSmall, GCSoftmaxLayerFixture<float>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SoftmaxLayerSmallShapes(), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("Beta", 1.0f)), - framework::dataset::make("ReduceEndAxis", 0))) + framework::dataset::make("Axis", 0))) { // Validate output validate(GCAccessor(_target), _reference, tolerance_f32); @@ -116,16 +116,16 @@ FIXTURE_DATA_TEST_CASE(RunSmall, GCSoftmaxLayerFixture<float>, framework::Datase FIXTURE_DATA_TEST_CASE(RunLarge, GCSoftmaxLayerFixture<float>, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::SoftmaxLayerLargeShapes(), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("Beta", 1.0f)), - framework::dataset::make("ReduceEndAxis", 0))) + framework::dataset::make("Axis", 0))) { // Validate output validate(GCAccessor(_target), _reference, tolerance_f32); } -TEST_SUITE_END() -TEST_SUITE_END() +TEST_SUITE_END() // FP32 +TEST_SUITE_END() // Float -TEST_SUITE_END() -TEST_SUITE_END() +TEST_SUITE_END() // SoftmaxLayer +TEST_SUITE_END() // GC } // namespace validation } // namespace test } // namespace arm_compute diff --git a/tests/validation/NEON/LogSoftmaxLayer.cpp b/tests/validation/NEON/LogSoftmaxLayer.cpp index 3f85e3f7a2..a7ab033359 100644 --- a/tests/validation/NEON/LogSoftmaxLayer.cpp +++ b/tests/validation/NEON/LogSoftmaxLayer.cpp @@ -71,7 +71,7 @@ TEST_SUITE(FP16) FIXTURE_DATA_TEST_CASE(RunSmall, NELogSoftmaxLayerFixture<half>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::Small4DShapes(), framework::dataset::make("DataType", DataType::F16)), framework::dataset::make("Beta", { 1.0f, 2.0f })), - framework::dataset::make("Axis", { 0 }))) + framework::dataset::make("Axis", { 0, -1 }))) { // Validate output validate(Accessor(_target), _reference, tolerance_f16); @@ -79,7 +79,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NELogSoftmaxLayerFixture<half>, framework::Data FIXTURE_DATA_TEST_CASE(RunSmall4D, NELogSoftmaxLayerFixture<half>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::Small4DShapes(), framework::dataset::make("DataType", DataType::F16)), framework::dataset::make("Beta", { 1.0f, 2.0f })), - framework::dataset::make("Axis", { 0 }))) + framework::dataset::make("Axis", { 0, -3, 2 }))) { // Validate output validate(Accessor(_target), _reference, tolerance_f16); @@ -99,7 +99,7 @@ TEST_SUITE(FP32) FIXTURE_DATA_TEST_CASE(RunSmall2D, NELogSoftmaxLayerFixture<float>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SoftmaxLayerSmallShapes(), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("Beta", { 1.0f, 2.0f })), - framework::dataset::make("Axis", { 0 }))) + framework::dataset::make("Axis", { 0, 1 }))) { // Validate output validate(Accessor(_target), _reference, tolerance_f32); @@ -107,7 +107,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall2D, NELogSoftmaxLayerFixture<float>, framework::D FIXTURE_DATA_TEST_CASE(RunSmall4D, NELogSoftmaxLayerFixture<float>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::Small4DShapes(), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("Beta", { 1.0f, 2.0f })), - framework::dataset::make("Axis", { 0 }))) + framework::dataset::make("Axis", { 0, 2, -1 }))) { // Validate output validate(Accessor(_target), _reference, tolerance_f32); @@ -132,7 +132,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall2D, NELogSoftmaxLayerQuantizedFixture<uint8_t>, f framework::dataset::make("DataType", DataType::QASYMM8)), combine(framework::dataset::make("QuantizationInfo", { QuantizationInfo(0.5f, -10) }), framework::dataset::make("Beta", { 1.0f, 2.f }))), - framework::dataset::make("Axis", { 0 }))) + framework::dataset::make("Axis", { 0, 1 }))) { // Validate output validate(Accessor(_target), _reference, tolerance_qasymm8); @@ -141,7 +141,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall4D, NELogSoftmaxLayerQuantizedFixture<uint8_t>, f framework::dataset::make("DataType", DataType::QASYMM8)), combine(framework::dataset::make("QuantizationInfo", { QuantizationInfo(0.5f, -10) }), framework::dataset::make("Beta", { 1.0f, 2.f }))), - framework::dataset::make("Axis", { 0 }))) + framework::dataset::make("Axis", { 0, -1, 1 }))) { // Validate output validate(Accessor(_target), _reference, tolerance_qasymm8); diff --git a/tests/validation/NEON/SoftmaxLayer.cpp b/tests/validation/NEON/SoftmaxLayer.cpp index 70203d9ce9..2a9e30604e 100644 --- a/tests/validation/NEON/SoftmaxLayer.cpp +++ b/tests/validation/NEON/SoftmaxLayer.cpp @@ -73,6 +73,7 @@ DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip(zip( TensorInfo(TensorShape(32U, 13U), 1, DataType::F32), TensorInfo(TensorShape(32U, 13U), 1, DataType::QASYMM8, QuantizationInfo(1.f/256, 12)), + TensorInfo(TensorShape(32U, 13U), 1, DataType::F32), TensorInfo(TensorShape(32U, 13U), 1, DataType::QASYMM8, //Invalid axis high QuantizationInfo(1.f/256, 12)), TensorInfo(TensorShape(32U, 13U), 1, DataType::QASYMM8, //Invalid axis low @@ -85,6 +86,7 @@ DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip(zip( TensorInfo(TensorShape(32U, 13U), 1, DataType::F32), TensorInfo(TensorShape(32U, 13U), 1, DataType::QASYMM8, QuantizationInfo(1.f/256, 0)), + TensorInfo(TensorShape(32U, 13U), 1, DataType::F32), TensorInfo(TensorShape(32U, 13U), 1, DataType::QASYMM8, QuantizationInfo(1.f/256, 0)), TensorInfo(TensorShape(32U, 13U), 1, DataType::QASYMM8, @@ -95,18 +97,20 @@ DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip(zip( 1.0, 2.0, 1.0, + 1.0, 2.0, 1.0, })), framework::dataset::make("axis", { 0, 0, 0, + 1, 0, - 0, + -1, 2, -3, })), - framework::dataset::make("Expected", { false, false, false, true, true, false, false })), + framework::dataset::make("Expected", { false, false, false, true, true, true, false, false })), input_info, output_info, beta, axis, expected) { ARM_COMPUTE_EXPECT(bool(NESoftmaxLayer::validate(&input_info.clone()->set_is_resizable(false), &output_info.clone()->set_is_resizable(false), beta, axis)) == expected, framework::LogLevel::ERRORS); @@ -123,7 +127,7 @@ TEST_SUITE(FP16) FIXTURE_DATA_TEST_CASE(RunSmall, NESoftmaxLayerFixture<half>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::Small4DShapes(), framework::dataset::make("DataType", DataType::F16)), framework::dataset::make("Beta", { 1.0f, 2.0f })), - framework::dataset::make("Axis", { 0 }))) + framework::dataset::make("Axis", { 0, 1 }))) { // Validate output validate(Accessor(_target), _reference, tolerance_f16); @@ -131,7 +135,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NESoftmaxLayerFixture<half>, framework::Dataset FIXTURE_DATA_TEST_CASE(RunSmall4D, NESoftmaxLayerFixture<half>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::Small4DShapes(), framework::dataset::make("DataType", DataType::F16)), framework::dataset::make("Beta", { 1.0f, 2.0f })), - framework::dataset::make("Axis", { 0 }))) + framework::dataset::make("Axis", { 0, 2, -1 }))) { // Validate output validate(Accessor(_target), _reference, tolerance_f16); @@ -151,7 +155,7 @@ TEST_SUITE(FP32) FIXTURE_DATA_TEST_CASE(RunSmall2D, NESoftmaxLayerFixture<float>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SoftmaxLayerSmallShapes(), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("Beta", { 1.0f, 2.0f })), - framework::dataset::make("Axis", { 0 }))) + framework::dataset::make("Axis", { 0, -1 }))) { // Validate output validate(Accessor(_target), _reference, tolerance_f32); @@ -159,7 +163,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall2D, NESoftmaxLayerFixture<float>, framework::Data FIXTURE_DATA_TEST_CASE(RunSmall4D, NESoftmaxLayerFixture<float>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::Small4DShapes(), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("Beta", { 1.0f, 2.0f })), - framework::dataset::make("Axis", { 0 }))) + framework::dataset::make("Axis", { 0, -2, 3 }))) { // Validate output validate(Accessor(_target), _reference, tolerance_f32); @@ -184,7 +188,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall2D, NESoftmaxLayerQuantizedFixture<uint8_t>, fram framework::dataset::make("DataType", DataType::QASYMM8)), combine(framework::dataset::make("QuantizationInfo", { QuantizationInfo(0.5f, -10) }), framework::dataset::make("Beta", { 1.0f, 2.f }))), - framework::dataset::make("Axis", { 0 }))) + framework::dataset::make("Axis", { 0, -1 }))) { // Validate output validate(Accessor(_target), _reference, tolerance_qasymm8); @@ -193,7 +197,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall4D, NESoftmaxLayerQuantizedFixture<uint8_t>, fram framework::dataset::make("DataType", DataType::QASYMM8)), combine(framework::dataset::make("QuantizationInfo", { QuantizationInfo(0.5f, -10) }), framework::dataset::make("Beta", { 1.0f, 2.f }))), - framework::dataset::make("Axis", { 0 }))) + framework::dataset::make("Axis", { 0, 1, -2 }))) { // Validate output validate(Accessor(_target), _reference, tolerance_qasymm8); @@ -214,7 +218,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall2D, NESoftmaxLayerQuantizedFixture<int8_t>, frame framework::dataset::make("DataType", DataType::QASYMM8_SIGNED)), combine(framework::dataset::make("QuantizationInfo", { QuantizationInfo(0.5f, -10) }), framework::dataset::make("Beta", { 1.0f, 2.f }))), - framework::dataset::make("Axis", { 0 }))) + framework::dataset::make("Axis", { 0, -1 }))) { // Validate output validate(Accessor(_target), _reference, tolerance_qasymm8_signed); @@ -223,7 +227,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall4D, NESoftmaxLayerQuantizedFixture<int8_t>, frame framework::dataset::make("DataType", DataType::QASYMM8_SIGNED)), combine(framework::dataset::make("QuantizationInfo", { QuantizationInfo(0.5f, -10) }), framework::dataset::make("Beta", { 1.0f, 2.f }))), - framework::dataset::make("Axis", { 0 }))) + framework::dataset::make("Axis", { 0, 1, -1 }))) { // Validate output validate(Accessor(_target), _reference, tolerance_qasymm8_signed); diff --git a/tests/validation/fixtures/SoftmaxLayerFixture.h b/tests/validation/fixtures/SoftmaxLayerFixture.h index 29a3ed2cd0..30356d648d 100644 --- a/tests/validation/fixtures/SoftmaxLayerFixture.h +++ b/tests/validation/fixtures/SoftmaxLayerFixture.h @@ -32,7 +32,6 @@ #include "tests/IAccessor.h" #include "tests/framework/Asserts.h" #include "tests/framework/Fixture.h" -#include "tests/validation/reference/LogSoftmaxLayer.h" #include "tests/validation/reference/SoftmaxLayer.h" #include <random> @@ -52,8 +51,8 @@ public: { _quantization_info = quantization_info; - _target = compute_target(shape, data_type, quantization_info, beta, axis); _reference = compute_reference(shape, data_type, quantization_info, beta, axis); + _target = compute_target(shape, data_type, quantization_info, beta, axis); } protected: @@ -62,7 +61,7 @@ protected: { if(!is_data_type_quantized(tensor.data_type())) { - std::uniform_real_distribution<> distribution(-1000.f, 1000.f); + std::uniform_real_distribution<> distribution(-10.f, 10.f); library->fill(tensor, distribution, 0); } else // data type is quantized_asymmetric (signed or unsigned) @@ -111,14 +110,7 @@ protected: // Fill reference fill(src); - if(IS_LOG) - { - return reference::log_softmax_layer<T>(src, beta, axis); - } - else - { - return reference::softmax_layer<T>(src, beta, axis); - } + return reference::softmax_layer<T>(src, beta, axis, IS_LOG); } TensorType _target{}; @@ -155,6 +147,7 @@ public: axis); } }; + } // namespace validation } // namespace test } // namespace arm_compute diff --git a/tests/validation/reference/LogSoftmaxLayer.cpp b/tests/validation/reference/LogSoftmaxLayer.cpp deleted file mode 100644 index 8d3b8f7579..0000000000 --- a/tests/validation/reference/LogSoftmaxLayer.cpp +++ /dev/null @@ -1,61 +0,0 @@ -/* - * Copyright (c) 2019-2020 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ -#include "LogSoftmaxLayer.h" -#include "SoftmaxLayer.h" - -#include "arm_compute/core/Types.h" - -namespace arm_compute -{ -namespace test -{ -namespace validation -{ -namespace reference -{ -template <typename T, typename std::enable_if<is_floating_point<T>::value, int>::type> -SimpleTensor<T> log_softmax_layer(const SimpleTensor<T> &src, float beta, int32_t reduce_end_axis) -{ - return softmax_layer_generic<T>(src, beta, reduce_end_axis, true); -} - -template < typename T, typename std::enable_if < std::is_same<T, uint8_t>::value || std::is_same<T, int8_t>::value, int >::type > -SimpleTensor<T> log_softmax_layer(const SimpleTensor<T> &src, float beta, int32_t reduce_end_axis) -{ - const QuantizationInfo output_quantization_info = arm_compute::get_softmax_output_quantization_info(src.data_type(), true); - - SimpleTensor<float> src_tmp = convert_from_asymmetric(src); - SimpleTensor<float> dst_tmp = log_softmax_layer<float>(src_tmp, beta, reduce_end_axis); - SimpleTensor<T> dst = convert_to_asymmetric<T>(dst_tmp, output_quantization_info); - return dst; -} - -template SimpleTensor<float> log_softmax_layer(const SimpleTensor<float> &src, float beta, int32_t reduce_end_axis); -template SimpleTensor<half> log_softmax_layer(const SimpleTensor<half> &src, float beta, int32_t reduce_end_axis); -template SimpleTensor<uint8_t> log_softmax_layer(const SimpleTensor<uint8_t> &src, float beta, int32_t reduce_end_axis); -template SimpleTensor<int8_t> log_softmax_layer(const SimpleTensor<int8_t> &src, float beta, int32_t reduce_end_axis); -} // namespace reference -} // namespace validation -} // namespace test -} // namespace arm_compute diff --git a/tests/validation/reference/LogSoftmaxLayer.h b/tests/validation/reference/LogSoftmaxLayer.h deleted file mode 100644 index db945074a2..0000000000 --- a/tests/validation/reference/LogSoftmaxLayer.h +++ /dev/null @@ -1,47 +0,0 @@ -/* - * Copyright (c) 2019-2020 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ -#ifndef ARM_COMPUTE_TEST_LOG_SOFTMAX_LAYER_H -#define ARM_COMPUTE_TEST_LOG_SOFTMAX_LAYER_H - -#include "tests/SimpleTensor.h" -#include "tests/validation/Helpers.h" - -namespace arm_compute -{ -namespace test -{ -namespace validation -{ -namespace reference -{ -template <typename T, typename std::enable_if<is_floating_point<T>::value, int>::type = 0> -SimpleTensor<T> log_softmax_layer(const SimpleTensor<T> &src, float beta, int32_t reduce_end_axis = 0); - -template < typename T, typename std::enable_if < std::is_same<T, uint8_t>::value || std::is_same<T, int8_t>::value, int >::type = 0 > -SimpleTensor<T> log_softmax_layer(const SimpleTensor<T> &src, float beta, int32_t reduce_end_axis = 0); -} // namespace reference -} // namespace validation -} // namespace test -} // namespace arm_compute -#endif /* ARM_COMPUTE_TEST_SOFTMAX_LAYER_H */ diff --git a/tests/validation/reference/SoftmaxLayer.cpp b/tests/validation/reference/SoftmaxLayer.cpp index 00206766f8..3fbac32a9b 100644 --- a/tests/validation/reference/SoftmaxLayer.cpp +++ b/tests/validation/reference/SoftmaxLayer.cpp @@ -25,6 +25,7 @@ #include "arm_compute/core/Helpers.h" #include "arm_compute/core/Types.h" +#include "utils/TypePrinter.h" namespace arm_compute { @@ -35,39 +36,43 @@ namespace validation namespace reference { template <typename T, typename std::enable_if<is_floating_point<T>::value, int>::type> -SimpleTensor<T> softmax_layer_generic(const SimpleTensor<T> &src, float beta, int32_t reduce_end_axis, bool is_log) +SimpleTensor<T> softmax_layer_generic(const SimpleTensor<T> &src, float beta, int32_t axis, bool is_log) { // Create reference SimpleTensor<T> dst{ src.shape(), src.data_type(), 1 }; - // Convert reduce-before axis (inclusive) to first n axes to reduce - const size_t first_n_reduce_axes = dim_index_2_num_dims(reduce_end_axis, static_cast<int32_t>(src.shape().num_dimensions())); + const int32_t n_dims = static_cast<int32_t>(src.shape().num_dimensions()); + ARM_COMPUTE_ERROR_ON(axis < -n_dims || axis >= n_dims); - // Compute reference. Lower dims are the collapsing of the first axis - // dimensions (i.e., the flattened dimension of each batch). The upper dims are - // instead the batches we want to normalize + const unsigned int actual_axis = static_cast<unsigned int>(wrap_around(axis, n_dims)); + Window window; + window.use_tensor_dimensions(src.shape()); + const unsigned int axis_dimension = src.shape()[actual_axis]; + window.set(actual_axis, Window::Dimension(0, 1, 1)); - const int lower_dims = src.shape().total_size_lower(first_n_reduce_axes); - - const int upper_dims = src.shape().total_size_upper(first_n_reduce_axes); - -#if defined(_OPENMP) - #pragma omp parallel for -#endif /* _OPENMP */ - for(int r = 0; r < upper_dims; ++r) + execute_window_loop(window, [&](const Coordinates & id) { - const T *src_row_ptr = src.data() + r * lower_dims; - T *dst_row_ptr = dst.data() + r * lower_dims; - - // Find max - const T max = *std::max_element(src_row_ptr, src_row_ptr + lower_dims); + // Find max along axis + Coordinates offset(id); + offset.set(actual_axis, 0); + T max = *reinterpret_cast<const T *>(src(offset)); + for(unsigned int axis_id = 1; axis_id < axis_dimension; ++axis_id) + { + offset.set(actual_axis, axis_id); + const T val = *reinterpret_cast<const T *>(src(offset)); + if(val > max) + { + max = val; + } + } // Regularize T sum(0.f); - std::transform(src_row_ptr, src_row_ptr + lower_dims, dst_row_ptr, [&sum, max, beta, is_log](T val) + for(unsigned int axis_id = 0; axis_id < axis_dimension; ++axis_id) { - T res{ (val - max) *beta }; - + offset.set(actual_axis, axis_id); + const T val = *reinterpret_cast<const T *>(src(offset)); + T res{ (val - max) *beta }; if(is_log) { sum += std::exp(res); @@ -77,50 +82,52 @@ SimpleTensor<T> softmax_layer_generic(const SimpleTensor<T> &src, float beta, in res = std::exp(res); sum += res; } - return res; - }); + *reinterpret_cast<T *>(dst(offset)) = res; + } // Normalize - std::transform(dst_row_ptr, dst_row_ptr + lower_dims, dst_row_ptr, [sum, is_log](T val) + for(unsigned int axis_id = 0; axis_id < axis_dimension; ++axis_id) { + offset.set(actual_axis, axis_id); + const T val = *reinterpret_cast<const T *>(dst(offset)); if(is_log) { - return val - static_cast<T>(std::log(sum)); + *reinterpret_cast<T *>(dst(offset)) = val - static_cast<T>(std::log(sum)); } else { - return val / sum; + *reinterpret_cast<T *>(dst(offset)) = val / sum; } - }); - } - + } + }); return dst; } -template SimpleTensor<float> softmax_layer_generic(const SimpleTensor<float> &src, float beta, int32_t reduce_end_axis, bool is_log); -template SimpleTensor<half> softmax_layer_generic(const SimpleTensor<half> &src, float beta, int32_t reduce_end_axis, bool is_log); +template SimpleTensor<float> softmax_layer_generic(const SimpleTensor<float> &src, float beta, int32_t axis, bool is_log); +template SimpleTensor<half> softmax_layer_generic(const SimpleTensor<half> &src, float beta, int32_t axis, bool is_log); template <typename T, typename std::enable_if<is_floating_point<T>::value, int>::type> -SimpleTensor<T> softmax_layer(const SimpleTensor<T> &src, float beta, int32_t reduce_end_axis) +SimpleTensor<T> softmax_layer(const SimpleTensor<T> &src, float beta, int32_t axis, bool is_log) { - return softmax_layer_generic<T>(src, beta, reduce_end_axis, false); + return softmax_layer_generic<T>(src, beta, axis, is_log); } template < typename T, typename std::enable_if < std::is_same<T, uint8_t>::value || std::is_same<T, int8_t>::value, int >::type > -SimpleTensor<T> softmax_layer(const SimpleTensor<T> &src, float beta, int32_t reduce_end_axis) +SimpleTensor<T> softmax_layer(const SimpleTensor<T> &src, float beta, int32_t axis, bool is_log) { - const QuantizationInfo output_quantization_info = arm_compute::get_softmax_output_quantization_info(src.data_type(), false); + const QuantizationInfo output_quantization_info = arm_compute::get_softmax_output_quantization_info(src.data_type(), is_log); SimpleTensor<float> src_tmp = convert_from_asymmetric(src); - SimpleTensor<float> dst_tmp = softmax_layer<float>(src_tmp, beta, reduce_end_axis); + SimpleTensor<float> dst_tmp = softmax_layer<float>(src_tmp, beta, axis, is_log); SimpleTensor<T> dst = convert_to_asymmetric<T>(dst_tmp, output_quantization_info); return dst; } -template SimpleTensor<float> softmax_layer(const SimpleTensor<float> &src, float beta, int32_t reduce_end_axis); -template SimpleTensor<half> softmax_layer(const SimpleTensor<half> &src, float beta, int32_t reduce_end_axis); -template SimpleTensor<uint8_t> softmax_layer(const SimpleTensor<uint8_t> &src, float beta, int32_t reduce_end_axis); -template SimpleTensor<int8_t> softmax_layer(const SimpleTensor<int8_t> &src, float beta, int32_t reduce_end_axis); +template SimpleTensor<float> softmax_layer(const SimpleTensor<float> &src, float beta, int32_t axis, bool is_log); +template SimpleTensor<half> softmax_layer(const SimpleTensor<half> &src, float beta, int32_t axis, bool is_log); +template SimpleTensor<uint8_t> softmax_layer(const SimpleTensor<uint8_t> &src, float beta, int32_t axis, bool is_log); +template SimpleTensor<int8_t> softmax_layer(const SimpleTensor<int8_t> &src, float beta, int32_t axis, bool is_log); + } // namespace reference } // namespace validation } // namespace test diff --git a/tests/validation/reference/SoftmaxLayer.h b/tests/validation/reference/SoftmaxLayer.h index 2af0b6d36a..3362f195c9 100644 --- a/tests/validation/reference/SoftmaxLayer.h +++ b/tests/validation/reference/SoftmaxLayer.h @@ -36,13 +36,13 @@ namespace validation namespace reference { template <typename T, typename std::enable_if<is_floating_point<T>::value, int>::type = 0> -SimpleTensor<T> softmax_layer_generic(const SimpleTensor<T> &src, float beta, int32_t reduce_end_axis, bool is_log = false); +SimpleTensor<T> softmax_layer_generic(const SimpleTensor<T> &src, float beta, int32_t axis, bool is_log = false); template <typename T, typename std::enable_if<is_floating_point<T>::value, int>::type = 0> -SimpleTensor<T> softmax_layer(const SimpleTensor<T> &src, float beta, int32_t reduce_end_axis = 0); +SimpleTensor<T> softmax_layer(const SimpleTensor<T> &src, float beta, int32_t axis = 0, bool is_log = false); template < typename T, typename std::enable_if < std::is_same<T, uint8_t>::value || std::is_same<T, int8_t>::value, int >::type = 0 > -SimpleTensor<T> softmax_layer(const SimpleTensor<T> &src, float beta, int32_t reduce_end_axis = 0); +SimpleTensor<T> softmax_layer(const SimpleTensor<T> &src, float beta, int32_t axis = 0, bool is_log = false); } // namespace reference } // namespace validation } // namespace test |