"src/core/CL/kernels/CLThresholdKernel.cpp",
"src/core/CL/kernels/CLTileKernel.cpp",
"src/core/CL/kernels/CLTransposeKernel.cpp",
- "src/core/CL/kernels/CLUpsampleLayerKernel.cpp",
"src/core/CL/kernels/CLWarpAffineKernel.cpp",
"src/core/CL/kernels/CLWarpPerspectiveKernel.cpp",
"src/core/CL/kernels/CLWeightsReshapeKernel.cpp",
"src/core/NEON/kernels/NEThresholdKernel.cpp",
"src/core/NEON/kernels/NETileKernel.cpp",
"src/core/NEON/kernels/NETransposeKernel.cpp",
- "src/core/NEON/kernels/NEUpsampleLayerKernel.cpp",
"src/core/NEON/kernels/NEWarpKernel.cpp",
"src/core/NEON/kernels/NEWeightsReshapeKernel.cpp",
"src/core/NEON/kernels/NEWidthConcatenateLayerKernel.cpp",
"src/runtime/CL/functions/CLTile.cpp",
"src/runtime/CL/functions/CLTranspose.cpp",
"src/runtime/CL/functions/CLUnstack.cpp",
- "src/runtime/CL/functions/CLUpsampleLayer.cpp",
"src/runtime/CL/functions/CLWarpAffine.cpp",
"src/runtime/CL/functions/CLWarpPerspective.cpp",
"src/runtime/CL/functions/CLWinogradConvolutionLayer.cpp",
"src/runtime/NEON/functions/NETile.cpp",
"src/runtime/NEON/functions/NETranspose.cpp",
"src/runtime/NEON/functions/NEUnstack.cpp",
- "src/runtime/NEON/functions/NEUpsampleLayer.cpp",
"src/runtime/NEON/functions/NEWarpAffine.cpp",
"src/runtime/NEON/functions/NEWarpPerspective.cpp",
"src/runtime/NEON/functions/NEWinogradConvolutionLayer.cpp",
* @return Node ID of the created node, EmptyNodeID in case of error
*/
static NodeID add_strided_slice_node(Graph &g, NodeParams params, NodeIdxPair input, Coordinates &starts, Coordinates &ends, BiStrides &strides, StridedSliceLayerInfo info);
- /** Adds an upsample layer to the graph
- *
- * @param[in] g Graph to add the node to
- * @param[in] params Common node parameters
- * @param[in] input Input to the yolo layer node as a NodeID-Index pair
- * @param[in] info Upsample layer stride info
- * @param[in] upsampling_policy Upsampling policy used
- *
- * @return Node ID of the created node, EmptyNodeID in case of error
- */
- static NodeID add_upsample_node(Graph &g, NodeParams params, NodeIdxPair input, Size2D info, InterpolationPolicy upsampling_policy);
/** Adds a yolo layer to the graph
*
* @param[in] g Graph to add the node to
// Create and configure function
auto func = std::make_unique<ResizeLayerFunction>();
- func->configure(input, output, ScaleKernelInfo{ policy, BorderMode::CONSTANT });
+ func->configure(input, output, ScaleKernelInfo{ policy, BorderMode::CONSTANT, PixelValue(), SamplingPolicy::CENTER, false, false });
// Log info
ARM_COMPUTE_LOG_GRAPH_INFO("Instantiated "
return RETURN_UNIQUE_PTR(func);
}
-
-/** Create a backend Upsample layer function
- *
- * @tparam UpsampleLayerFunction Backend Upsample function
- * @tparam TargetInfo Target-specific information
- *
- * @param[in] node Node to create the backend function for
- * @param[in] ctx Graph context
- *
- * @return Backend Upsample layer function
- */
-template <typename UpsampleLayerFunction, typename TargetInfo>
-std::unique_ptr<IFunction> create_upsample_layer(UpsampleLayerNode &node, GraphContext &ctx)
-{
- ARM_COMPUTE_UNUSED(ctx);
- validate_node<TargetInfo>(node, 1 /* expected inputs */, 1 /* expected outputs */);
-
- // Extract IO and info
- typename TargetInfo::TensorType *input = get_backing_tensor<TargetInfo>(node.input(0));
- typename TargetInfo::TensorType *output = get_backing_tensor<TargetInfo>(node.output(0));
- const Size2D info = node.info();
- const InterpolationPolicy upsampling_policy = node.upsampling_policy();
- ARM_COMPUTE_ERROR_ON(upsampling_policy != InterpolationPolicy::NEAREST_NEIGHBOR);
- ARM_COMPUTE_ERROR_ON(info.x() != 2 || info.y() != 2);
- ARM_COMPUTE_ERROR_ON(input == nullptr);
- ARM_COMPUTE_ERROR_ON(output == nullptr);
-
- // Create and configure function
- auto func = std::make_unique<UpsampleLayerFunction>();
- func->configure(input, output, info, upsampling_policy);
-
- // Log info
- ARM_COMPUTE_LOG_GRAPH_INFO("Instantiated "
- << node.name()
- << " Type: " << node.type()
- << " Target: " << TargetInfo::TargetType
- << " Data Type: " << input->info()->data_type()
- << " Input shape: " << input->info()->tensor_shape()
- << " Output shape: " << output->info()->tensor_shape()
- << " Strides: " << info
- << " Upsampling policy: " << upsampling_policy
- << std::endl);
-
- return RETURN_UNIQUE_PTR(func);
-}
} // namespace detail
} // namespace backends
} // namespace graph
return StridedSliceLayer::validate(input, output, starts, ends, strides, info.begin_mask(), info.end_mask(), info.shrink_axis_mask());
}
-/** Validates a Upsample layer node
- *
- * @tparam UpsampleLayer Upsample layer type
- *
- * @param[in] node Node to validate
- *
- * @return Status
- */
-template <typename UpsampleLayer>
-Status validate_upsample_layer(UpsampleLayerNode &node)
-{
- ARM_COMPUTE_LOG_GRAPH_VERBOSE("Validating UpsampleLayer node with ID : " << node.id() << " and Name: " << node.name() << std::endl);
- ARM_COMPUTE_RETURN_ERROR_ON(node.num_inputs() != 1);
- ARM_COMPUTE_RETURN_ERROR_ON(node.num_outputs() != 1);
-
- // Extract input and output
- arm_compute::ITensorInfo *input = detail::get_backing_tensor_info(node.input(0));
- arm_compute::ITensorInfo *output = get_backing_tensor_info(node.output(0));
-
- // Validate function
- return UpsampleLayer::validate(input, output, node.info(), node.upsampling_policy());
-}
/** Validates a element-wise layer node
*
* @param[in] node Node to validate
StridedSliceLayerInfo _info;
};
-/** Upsample Layer */
-class UpsampleLayer final : public ILayer
-{
-public:
- /** Construct a Upsample layer.
- *
- * @param[in] info Stride info
- * @param[in] upsampling_policy Upsampling policy
- */
- UpsampleLayer(Size2D info, InterpolationPolicy upsampling_policy)
- : _info(info), _upsampling_policy(upsampling_policy)
- {
- }
-
- NodeID create_layer(IStream &s) override
- {
- NodeParams common_params = { name(), s.hints().target_hint };
- NodeIdxPair input = { s.tail_node(), 0 };
- return GraphBuilder::add_upsample_node(s.graph(), common_params, input, _info, _upsampling_policy);
- }
-
-private:
- Size2D _info;
- InterpolationPolicy _upsampling_policy;
-};
-
/** YOLO Layer */
class YOLOLayer final : public ILayer
{
#include "arm_compute/graph/nodes/SplitLayerNode.h"
#include "arm_compute/graph/nodes/StackLayerNode.h"
#include "arm_compute/graph/nodes/StridedSliceLayerNode.h"
-#include "arm_compute/graph/nodes/UpsampleLayerNode.h"
#endif /* ARM_COMPUTE_GRAPH_NODES_H */
class SplitLayerNode;
class StackLayerNode;
class StridedSliceLayerNode;
-class UpsampleLayerNode;
} // namespace graph
} // namespace arm_compute
#endif /* ARM_COMPUTE_GRAPH_NODES_FWD_H */
+++ /dev/null
-/*
- * Copyright (c) 2018-2019 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_GRAPH_UPSAMPLE_LAYER_NODE_H
-#define ARM_COMPUTE_GRAPH_UPSAMPLE_LAYER_NODE_H
-
-#include "arm_compute/graph/INode.h"
-
-namespace arm_compute
-{
-namespace graph
-{
-/** Upsample Layer node */
-class UpsampleLayerNode final : public INode
-{
-public:
- /** Constructor
- *
- * @param[in] info Stride info
- * @param[in] upsampling_policy Upsampling policy
- */
- UpsampleLayerNode(Size2D info, InterpolationPolicy upsampling_policy);
- /** Stride info metadata accessor
- *
- * @return The stride info of the layer
- */
- Size2D info() const;
- /** Upsampling policy metadata accessor
- *
- * @return The upsampling policy of the layer
- */
- InterpolationPolicy upsampling_policy() const;
- /** Computes upsample output descriptor
- *
- * @param[in] input_descriptor Input descriptor
- * @param[in] info Stride information
- *
- * @return Output descriptor
- */
- static TensorDescriptor compute_output_descriptor(const TensorDescriptor &input_descriptor, Size2D info);
-
- // Inherited overridden methods:
- NodeType type() const override;
- bool forward_descriptors() override;
- TensorDescriptor configure_output(size_t idx) const override;
- void accept(INodeVisitor &v) override;
-
-private:
- Size2D _info;
- InterpolationPolicy _upsampling_policy;
-};
-} // namespace graph
-} // namespace arm_compute
-#endif /* ARM_COMPUTE_GRAPH_UPSAMPLE_LAYER_NODE_H */
#include "arm_compute/runtime/CL/functions/CLTile.h"
#include "arm_compute/runtime/CL/functions/CLTranspose.h"
#include "arm_compute/runtime/CL/functions/CLUnstack.h"
-#include "arm_compute/runtime/CL/functions/CLUpsampleLayer.h"
#include "arm_compute/runtime/CL/functions/CLWarpAffine.h"
#include "arm_compute/runtime/CL/functions/CLWarpPerspective.h"
#include "arm_compute/runtime/CL/functions/CLWinogradConvolutionLayer.h"
+++ /dev/null
-/*
- * Copyright (c) 2018-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_CLUPSAMPLELAYER_H
-#define ARM_COMPUTE_CLUPSAMPLELAYER_H
-
-#include "arm_compute/runtime/IFunction.h"
-
-#include "arm_compute/core/Types.h"
-#include "arm_compute/runtime/IFunction.h"
-
-#include <memory>
-
-namespace arm_compute
-{
-class CLCompileContext;
-class CLUpsampleLayerKernel;
-class ICLTensor;
-class ITensorInfo;
-
-/** Basic function to run @ref CLUpsampleLayerKernel */
-class CLUpsampleLayer : public IFunction
-{
-public:
- /** Default constructor */
- CLUpsampleLayer();
- /** Prevent instances of this class from being copied (As this class contains pointers) */
- CLUpsampleLayer(const CLUpsampleLayer &) = delete;
- /** Prevent instances of this class from being copied (As this class contains pointers) */
- CLUpsampleLayer &operator=(const CLUpsampleLayer &) = delete;
- /** Allow instances of this class to be moved */
- CLUpsampleLayer(CLUpsampleLayer &&) = default;
- /** Allow instances of this class to be moved */
- CLUpsampleLayer &operator=(CLUpsampleLayer &&) = default;
- /** Default destructor */
- ~CLUpsampleLayer();
-
- /** Initialize the function's source, destination, interpolation type and border_mode.
- *
- * @param[in] input Source tensor. Data type supported: All.
- * @param[out] output Destination tensor. Data types supported: same as @p input.
- * @param[in] info Contains stride information described in @ref Size2D.
- * @param[in] upsampling_policy Defines the policy to fill the intermediate pixels.
- */
- void configure(ICLTensor *input, ICLTensor *output,
- const Size2D &info, const InterpolationPolicy upsampling_policy);
- /** Initialize the function's source, destination, interpolation type and border_mode.
- *
- * @param[in] compile_context The compile context to be used.
- * @param[in] input Source tensor. Data type supported: All.
- * @param[out] output Destination tensor. Data types supported: same as @p input.
- * @param[in] info Contains stride information described in @ref Size2D.
- * @param[in] upsampling_policy Defines the policy to fill the intermediate pixels.
- */
- void configure(const CLCompileContext &compile_context, ICLTensor *input, ICLTensor *output,
- const Size2D &info, const InterpolationPolicy upsampling_policy);
- /** Static function to check if given info will lead to a valid configuration of @ref CLUpsampleLayerKernel
- *
- * @param[in] input Source tensor info. Data types supported: All.
- * @param[in] output Destination tensor info. Data types supported: same as @p input.
- * @param[in] info Contains stride information described in @ref Size2D.
- * @param[in] upsampling_policy Defines the policy to fill the intermediate pixels.
- *
- * @return a status
- */
- static Status validate(const ITensorInfo *input, const ITensorInfo *output,
- const Size2D &info, const InterpolationPolicy upsampling_policy);
-
- // Inherited methods overridden:
- void run() override;
-
-private:
- std::unique_ptr<CLUpsampleLayerKernel> _upsample;
- ICLTensor *_output;
-};
-} // namespace arm_compute
-#endif /* ARM_COMPUTE_CLUPSAMPLELAYER_H */
#include "arm_compute/runtime/NEON/functions/NETile.h"
#include "arm_compute/runtime/NEON/functions/NETranspose.h"
#include "arm_compute/runtime/NEON/functions/NEUnstack.h"
-#include "arm_compute/runtime/NEON/functions/NEUpsampleLayer.h"
#include "arm_compute/runtime/NEON/functions/NEWarpAffine.h"
#include "arm_compute/runtime/NEON/functions/NEWarpPerspective.h"
#include "arm_compute/runtime/NEON/functions/NEWinogradConvolutionLayer.h"
+++ /dev/null
-/*
- * Copyright (c) 2018-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_NEUPSAMPLELAYER_H
-#define ARM_COMPUTE_NEUPSAMPLELAYER_H
-
-#include "arm_compute/core/Types.h"
-#include "arm_compute/runtime/IFunction.h"
-#include "arm_compute/runtime/NEON/NEScheduler.h"
-#include "arm_compute/runtime/Tensor.h"
-
-#include <memory>
-
-namespace arm_compute
-{
-class ITensor;
-class NEUpsampleLayerKernel;
-
-/** Function to run upsample layer */
-class NEUpsampleLayer : public IFunction
-{
-public:
- /** Constructor */
- NEUpsampleLayer();
- /** Prevent instances of this class from being copied (As this class contains pointers) */
- NEUpsampleLayer(const NEUpsampleLayer &) = delete;
- /** Prevent instances of this class from being copied (As this class contains pointers) */
- NEUpsampleLayer &operator=(const NEUpsampleLayer &) = delete;
- /** Prevent instances of this class from being moved (As this class contains non movable objects) */
- NEUpsampleLayer(NEUpsampleLayer &&) = delete;
- /** Prevent instances of this class from being moved (As this class contains non movable objects) */
- NEUpsampleLayer &operator=(NEUpsampleLayer &&) = delete;
- /** Default destructor */
- ~NEUpsampleLayer();
- /** Set the input output tensors.
- *
- * @param[in] input Source tensor. Data types supported: QASYMM8_SIGNED/QASYMM8/F16/F32.
- * @param[out] output Destination tensor. Data types supported: same as @p input.
- * @param[in] info Contains stride information described in @ref Size2D.
- * @param[in] policy Defines the policy to fill the intermediate pixels.
- *
- */
- void configure(const ITensor *input, ITensor *output, const Size2D &info,
- const InterpolationPolicy &policy);
- /** Static function to check if given info will lead to a valid configuration of @ref NEUpsampleLayer
- *
- * @param[in] input Source tensor info. Data types supported: QASYMM8_SIGNED/QASYMM8/F16/F32.
- * @param[out] output Destination tensor info. Data types supported: same as @p input.
- * @param[in] info Contains stride information described in @ref Size2D.
- * @param[in] policy Defines the policy to fill the intermediate pixels.
- *
- * @return a status
- */
- static Status validate(const ITensorInfo *input, const ITensorInfo *output, const Size2D &info,
- const InterpolationPolicy &policy);
-
- // Inherited methods overridden:
- void run() override;
-
-private:
- std::unique_ptr<NEUpsampleLayerKernel> _kernel;
- DataLayout _data_layout;
-};
-} // arm_compute
-#endif /* ARM_COMPUTE_NEUPSAMPLELAYER_H */
- NEGEMMTranspose1xW
- NEComputeAllAnchors / CLComputeAllAnchors
- NEGEMMAssemblyDispatch
+ - NEUpsampleLayer / CLUpsampleLayer
- Removed kernels:
- NEGEMMMatrixVectorMultiplyKernel
- NELocallyConnectedMatrixMultiplyKernel / CLLocallyConnectedMatrixMultiplyKernel
+ - NEUpsampleLayerKernel / CLUpsampleLayerKernel
v20.11 Public major release
- Various bug fixes.
- @ref NEROIPoolingLayerKernel
- @ref NEROIAlignLayerKernel
- NEYOLOLayerKernel
- - @ref NEUpsampleLayerKernel
+ - NEUpsampleLayerKernel
- @ref NEFloorKernel
- @ref NEWidthConcatenateLayerKernel
- @ref NEDepthConcatenateLayerKernel
- @ref CLReduceMean
- @ref NEScale
- @ref NEScaleKernel
- - @ref NEUpsampleLayer
+ - NEUpsampleLayer
- @ref NECast
- @ref NEReductionOperation
- @ref NEReduceMean
- @ref NEReduceMean
- @ref NEReorgLayer / @ref NEReorgLayerKernel
- @ref NEPriorBoxLayer / @ref NEPriorBoxLayerKernel
- - @ref NEUpsampleLayer / @ref NEUpsampleLayerKernel
+ - NEUpsampleLayer / NEUpsampleLayerKernel
- NEYOLOLayer / NEYOLOLayerKernel
- New OpenCL kernels / functions:
- @ref CLBatchToSpaceLayer / @ref CLBatchToSpaceLayerKernel
- @ref CLSlice
- @ref CLSplit
- @ref CLStridedSlice / @ref CLStridedSliceKernel
- - @ref CLUpsampleLayer / @ref CLUpsampleLayerKernel
+ - CLUpsampleLayer / CLUpsampleLayerKernel
- CLYOLOLayer / CLYOLOLayerKernel
- New CPP kernels / functions:
- @ref CPPBoxWithNonMaximaSuppressionLimit / @ref CPPBoxWithNonMaximaSuppressionLimitKernel
0.000001f)
.set_name("conv2d_59/BatchNorm")
<< ActivationLayer(ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::LEAKY_RELU, 0.1f)).set_name("conv2d_60/LeakyRelu")
- << UpsampleLayer(Size2D(2, 2), InterpolationPolicy::NEAREST_NEIGHBOR).set_name("Upsample_60");
+ << ResizeLayer(InterpolationPolicy::NEAREST_NEIGHBOR, 2, 2).set_name("Upsample_60");
SubStream concat_1(route_1);
concat_1 << ConcatLayer(std::move(route_1), std::move(intermediate_layers.second)).set_name("Route1")
<< ConvolutionLayer(
0.000001f)
.set_name("conv2d_66/BatchNorm")
<< ActivationLayer(ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::LEAKY_RELU, 0.1f)).set_name("conv2d_68/LeakyRelu")
- << UpsampleLayer(Size2D(2, 2), InterpolationPolicy::NEAREST_NEIGHBOR).set_name("Upsample_68");
+ << ResizeLayer(InterpolationPolicy::NEAREST_NEIGHBOR, 2, 2).set_name("Upsample_68");
SubStream concat_2(route_2);
concat_2 << ConcatLayer(std::move(route_2), std::move(intermediate_layers.first)).set_name("Route2")
<< ConvolutionLayer(
#include "src/core/CL/kernels/CLThresholdKernel.h"
#include "src/core/CL/kernels/CLTileKernel.h"
#include "src/core/CL/kernels/CLTransposeKernel.h"
-#include "src/core/CL/kernels/CLUpsampleLayerKernel.h"
#include "src/core/CL/kernels/CLWarpAffineKernel.h"
#include "src/core/CL/kernels/CLWarpPerspectiveKernel.h"
#include "src/core/CL/kernels/CLWeightsReshapeKernel.h"
+++ /dev/null
-/*
- * Copyright (c) 2018-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 "src/core/CL/kernels/CLUpsampleLayerKernel.h"
-
-#include "arm_compute/core/CL/CLHelpers.h"
-#include "arm_compute/core/CL/CLKernelLibrary.h"
-#include "arm_compute/core/CL/ICLTensor.h"
-#include "arm_compute/core/Error.h"
-#include "arm_compute/core/Helpers.h"
-#include "arm_compute/core/Validate.h"
-#include "arm_compute/core/Window.h"
-#include "arm_compute/core/utils/misc/ShapeCalculator.h"
-#include "src/core/AccessWindowStatic.h"
-#include "src/core/CL/CLValidate.h"
-#include "src/core/helpers/AutoConfiguration.h"
-#include "src/core/helpers/WindowHelpers.h"
-#include "support/StringSupport.h"
-
-namespace arm_compute
-{
-CLUpsampleLayerKernel::CLUpsampleLayerKernel()
- : _input(nullptr), _output(nullptr), _info(), _data_layout(DataLayout::UNKNOWN), _num_elems_processed_per_iteration_input_x()
-{
-}
-
-Status CLUpsampleLayerKernel::validate(const ITensorInfo *input, const ITensorInfo *output, const Size2D &info, const InterpolationPolicy upsampling_policy)
-{
- ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input, output);
- ARM_COMPUTE_UNUSED(upsampling_policy);
-
- DataLayout data_layout = input->data_layout();
- const int idx_width = get_data_layout_dimension_index(data_layout, DataLayoutDimension::WIDTH);
- const int idx_height = get_data_layout_dimension_index(data_layout, DataLayoutDimension::HEIGHT);
-
- ARM_COMPUTE_RETURN_ERROR_ON_F16_UNSUPPORTED(input);
- ARM_COMPUTE_RETURN_ERROR_ON(input->data_type() == DataType::UNKNOWN);
-
- ARM_COMPUTE_RETURN_ERROR_ON_MSG(info.x() != 2 || info.y() != 2, "Only stride 2 is supported");
- ARM_COMPUTE_RETURN_ERROR_ON_MSG(upsampling_policy != InterpolationPolicy::NEAREST_NEIGHBOR, "Only nearest neighbor policy supported");
-
- if(output->total_size() != 0)
- {
- ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input, output);
- ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_LAYOUT(input, output);
- ARM_COMPUTE_RETURN_ERROR_ON(output->dimension(idx_width) != info.x() * input->dimension(idx_width));
- ARM_COMPUTE_RETURN_ERROR_ON(output->dimension(idx_height) != info.y() * input->dimension(idx_height));
- ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_QUANTIZATION_INFO(input, output);
- }
-
- return Status{};
-}
-
-void CLUpsampleLayerKernel::configure(const ICLTensor *input, ICLTensor *output, const Size2D &info, const InterpolationPolicy upsampling_policy)
-{
- configure(CLKernelLibrary::get().get_compile_context(), input, output, info, upsampling_policy);
-}
-
-void CLUpsampleLayerKernel::configure(const CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output, const Size2D &info, const InterpolationPolicy upsampling_policy)
-{
- ARM_COMPUTE_ERROR_ON_NULLPTR(input, output);
- ARM_COMPUTE_UNUSED(upsampling_policy);
-
- _input = input;
- _output = output;
- _info = info;
- _data_layout = input->info()->data_layout();
- _num_elems_processed_per_iteration_input_x = 1;
-
- TensorShape output_shape = misc::shape_calculator::compute_upsample_shape(*input->info(), info);
- auto_init_if_empty(*output->info(), output_shape, 1, input->info()->data_type());
- output->info()->set_data_layout(_data_layout);
-
- unsigned int num_elems_processed_per_iteration_x = 16;
- const int output_width_x = output->info()->dimension(0);
- const bool multi_access_x = ((output_width_x / num_elems_processed_per_iteration_x) > 0);
-
- // Perform validation step
- ARM_COMPUTE_ERROR_THROW_ON(CLUpsampleLayerKernel::validate(input->info(), output->info(), info, upsampling_policy));
-
- Window win{};
-
- switch(_data_layout)
- {
- case DataLayout::NCHW:
- {
- win = calculate_max_window(*output->info());
- win.set(Window::DimY, Window::Dimension(win.y().start(), win.y().end(), info.y()));
- if(multi_access_x)
- {
- _num_elems_processed_per_iteration_input_x = num_elems_processed_per_iteration_x / info.x();
- win.set(Window::DimX, Window::Dimension(win.x().start(), ceil_to_multiple(win.x().end(), num_elems_processed_per_iteration_x), num_elems_processed_per_iteration_x));
- }
- break;
- }
- case DataLayout::NHWC:
- {
- win = calculate_max_window(*output->info());
- win.set(Window::DimY, Window::Dimension(win.y().start(), win.y().end(), info.x()));
- win.set(Window::DimZ, Window::Dimension(win.z().start(), win.z().end(), info.y()));
- if(multi_access_x)
- {
- _num_elems_processed_per_iteration_input_x = num_elems_processed_per_iteration_x;
- win.set(Window::DimX, Window::Dimension(win.x().start(), ceil_to_multiple(win.x().end(),
- num_elems_processed_per_iteration_x),
- num_elems_processed_per_iteration_x));
- }
- break;
- }
- default:
- ARM_COMPUTE_ERROR("Not implemented");
- }
-
- // Create kernel
- CLBuildOptions build_opts;
- build_opts.add_option(("-DDATA_TYPE=" + get_cl_unsigned_type_from_element_size(input->info()->element_size())));
- build_opts.add_option_if(multi_access_x, "-DVEC_SIZE_IN=" + support::cpp11::to_string(_num_elems_processed_per_iteration_input_x));
- build_opts.add_option_if(multi_access_x, "-DVEC_SIZE_OUT=" + support::cpp11::to_string(num_elems_processed_per_iteration_x));
- build_opts.add_option_if(multi_access_x, "-DLAST_ACCESSED_X_IN=" + support::cpp11::to_string(std::max<int>(_input->info()->dimension(0) - _num_elems_processed_per_iteration_input_x, 0)));
- build_opts.add_option_if(multi_access_x, "-DLAST_ACCESSED_X_OUT=" + support::cpp11::to_string(std::max<int>(output_width_x - num_elems_processed_per_iteration_x, 0)));
- _kernel = create_kernel(compile_context, "upsample_layer_" + lower_string(string_from_data_layout(input->info()->data_layout())), build_opts.options());
-
- ICLKernel::configure_internal(win);
-}
-
-void CLUpsampleLayerKernel::run(const Window &window, cl::CommandQueue &queue)
-{
- ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this);
- ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(ICLKernel::window(), window);
-
- Window collapsed_window = window.collapse_if_possible(ICLKernel::window(), Window::DimZ);
- Window slice_out = collapsed_window.first_slice_window_3D();
- Window slice_in = collapsed_window.first_slice_window_3D();
-
- switch(_data_layout)
- {
- case DataLayout::NCHW:
- slice_in.set(Window::DimX, Window::Dimension(0, _input->info()->dimension(0), _num_elems_processed_per_iteration_input_x));
- slice_in.set(Window::DimY, Window::Dimension(0, _input->info()->dimension(1), 1));
- break;
- case DataLayout::NHWC:
- slice_in.set(Window::DimY, Window::Dimension(0, _input->info()->dimension(1), 1));
- slice_in.set(Window::DimZ, Window::Dimension(0, _input->info()->dimension(2), 1));
- break;
- default:
- ARM_COMPUTE_ERROR("Not implemented");
- }
-
- do
- {
- unsigned int idx = 0;
- add_3D_tensor_argument(idx, _input, slice_in);
- add_3D_tensor_argument(idx, _output, slice_out);
- enqueue(queue, *this, slice_out, lws_hint());
- }
- while(collapsed_window.slide_window_slice_3D(slice_out) && collapsed_window.slide_window_slice_3D(slice_in));
-}
-} // namespace arm_compute
+++ /dev/null
-/*
- * Copyright (c) 2018-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_CLUPSAMPLELAYERKERNEL_H
-#define ARM_COMPUTE_CLUPSAMPLELAYERKERNEL_H
-
-#include "src/core/CL/ICLKernel.h"
-
-namespace arm_compute
-{
-class ICLTensor;
-
-/** Interface for the UpsampleLayer kernel on OpenCL. */
-class CLUpsampleLayerKernel : public ICLKernel
-{
-public:
- /** Constructor */
- CLUpsampleLayerKernel();
- /** Prevent instances of this class from being copied (As this class contains pointers) */
- CLUpsampleLayerKernel(const CLUpsampleLayerKernel &) = delete;
- /** Prevent instances of this class from being copied (As this class contains pointers) */
- CLUpsampleLayerKernel &operator=(const CLUpsampleLayerKernel &) = delete;
- /** Default Move Constructor. */
- CLUpsampleLayerKernel(CLUpsampleLayerKernel &&) = default;
- /** Default move assignment operator */
- CLUpsampleLayerKernel &operator=(CLUpsampleLayerKernel &&) = default;
- /** Default destructor */
- ~CLUpsampleLayerKernel() = default;
-
- /** Initialise the kernel's input and output.
- *
- * @param[in] input Source tensor. Data types supported: All.
- * @param[out] output Destination tensor. Data types supported: same as @p input.
- * @param[in] info Contains stride information described in @ref Size2D.
- * @param[in] upsampling_policy Defines the policy to fill the intermediate pixels.
- */
- void configure(const ICLTensor *input, ICLTensor *output, const Size2D &info, const InterpolationPolicy upsampling_policy);
- /** Initialise the kernel's input and output.
- *
- * @param[in] compile_context The compile context to be used.
- * @param[in] input Source tensor. Data types supported: All.
- * @param[out] output Destination tensor. Data types supported: same as @p input.
- * @param[in] info Contains stride information described in @ref Size2D.
- * @param[in] upsampling_policy Defines the policy to fill the intermediate pixels.
- */
- void configure(const CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output, const Size2D &info, const InterpolationPolicy upsampling_policy);
- /** Static function to check if given info will lead to a valid configuration of @ref CLUpsampleLayerKernel
- *
- * @param[in] input Source tensor info. Data types supported: All.
- * @param[in] output Destination tensor info. Data types supported: same as @p input.
- * @param[in] info Contains stride information described in @ref Size2D.
- * @param[in] upsampling_policy Defines the policy to fill the intermediate pixels.
- *
- * @return a status
- */
- static Status validate(const ITensorInfo *input, const ITensorInfo *output, const Size2D &info, const InterpolationPolicy upsampling_policy);
-
- // Inherited methods overridden:
- void run(const Window &window, cl::CommandQueue &queue) override;
-
-private:
- const ICLTensor *_input;
- ICLTensor *_output;
- Size2D _info;
- DataLayout _data_layout;
- unsigned int _num_elems_processed_per_iteration_input_x;
-};
-} // namespace arm_compute
-#endif /*ARM_COMPUTE_CLUPSAMPLELAYERKERNEL_H */
#include "src/core/NEON/kernels/NEThresholdKernel.h"
#include "src/core/NEON/kernels/NETileKernel.h"
#include "src/core/NEON/kernels/NETransposeKernel.h"
-#include "src/core/NEON/kernels/NEUpsampleLayerKernel.h"
#include "src/core/NEON/kernels/NEWarpKernel.h"
#include "src/core/NEON/kernels/NEWeightsReshapeKernel.h"
#include "src/core/NEON/kernels/NEWidthConcatenateLayerKernel.h"
+++ /dev/null
-/*
- * Copyright (c) 2018-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 "src/core/NEON/kernels/NEUpsampleLayerKernel.h"
-
-#include "arm_compute/core/Error.h"
-#include "arm_compute/core/Helpers.h"
-#include "arm_compute/core/ITensor.h"
-#include "arm_compute/core/TensorInfo.h"
-#include "arm_compute/core/Validate.h"
-#include "arm_compute/core/Window.h"
-#include "arm_compute/core/utils/misc/ShapeCalculator.h"
-#include "src/core/CPP/Validate.h"
-#include "src/core/NEON/wrapper/wrapper.h"
-#include "src/core/helpers/AutoConfiguration.h"
-#include "src/core/helpers/WindowHelpers.h"
-
-#include <arm_neon.h>
-
-namespace arm_compute
-{
-namespace
-{
-template <typename T, int S>
-inline T get_data_out(T data, int offset)
-{
- T out{ 0 };
- for(int i = 0; i < S / 2; ++i)
- {
- out[2 * i] = wrapper::vgetlane(data, i + offset);
- out[2 * i + 1] = wrapper::vgetlane(data, i + offset);
- }
- return out;
-}
-} // namespace
-NEUpsampleLayerKernel::NEUpsampleLayerKernel()
- : _func(nullptr), _input(nullptr), _output(nullptr), _info()
-{
-}
-
-Status NEUpsampleLayerKernel::validate(const ITensorInfo *input, const ITensorInfo *output, const Size2D &info, const InterpolationPolicy policy)
-{
- ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input, output);
- ARM_COMPUTE_UNUSED(policy);
-
- const DataLayout data_layout = input->data_layout();
- const int idx_width = get_data_layout_dimension_index(data_layout, DataLayoutDimension::WIDTH);
- const int idx_height = get_data_layout_dimension_index(data_layout, DataLayoutDimension::HEIGHT);
-
- ARM_COMPUTE_RETURN_ERROR_ON_CPU_F16_UNSUPPORTED(input);
- ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::QASYMM8_SIGNED, DataType::QASYMM8, DataType::F16, DataType::F32);
- ARM_COMPUTE_RETURN_ERROR_ON_MSG(info.x() != 2 || info.y() != 2, "Only stride 2 is supported");
- ARM_COMPUTE_RETURN_ERROR_ON_MSG(policy != InterpolationPolicy::NEAREST_NEIGHBOR, "Only nearest neighbor policy supported");
-
- // Check output if configured
- if(output->total_size() != 0)
- {
- ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input, output);
- ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_LAYOUT(input, output);
- ARM_COMPUTE_RETURN_ERROR_ON(output->dimension(idx_width) != info.x() * input->dimension(idx_width));
- ARM_COMPUTE_RETURN_ERROR_ON(output->dimension(idx_height) != info.y() * input->dimension(idx_height));
- ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_QUANTIZATION_INFO(input, output);
- }
- return Status{};
-}
-
-template <typename T, int S>
-void NEUpsampleLayerKernel::upsample_nchw(const arm_compute::Window &window)
-{
- using VectorType = typename wrapper::traits::neon_vector<T, S>::type;
-
- Window window_in(window);
- window_in.set(Window::DimX, Window::Dimension(0, 1, 1));
-
- Window window_out(window);
- window_out.set(Window::DimX, Window::Dimension(0, 1, 1));
- window_out.set(Window::DimY, Window::Dimension(0, _output->info()->dimension(1), _info.y()));
-
- const auto window_start_x = static_cast<int>(window.x().start());
- const auto window_end_x = static_cast<int>(window.x().end());
- const int window_step_x = S;
-
- Iterator input(_input, window_in);
- Iterator output(_output, window_out);
- const int offset_y_out = _output->info()->strides_in_bytes().y() / sizeof(T);
-
- execute_window_loop(window_out, [&](const Coordinates &)
- {
- const auto input_ptr = reinterpret_cast<const T *>(input.ptr());
- const auto output_ptr = reinterpret_cast<T *>(output.ptr());
-
- int x = window_start_x;
- for(; x <= (window_end_x - window_step_x); x += window_step_x)
- {
- const VectorType data = wrapper::vloadq(reinterpret_cast<const T *>(input_ptr + x));
- const VectorType data_out1 = get_data_out<VectorType, S>(data, 0);
- const VectorType data_out2 = get_data_out<VectorType, S>(data, S / 2);
-
- wrapper::vstore(output_ptr + 2 * x, data_out1);
- wrapper::vstore(output_ptr + 2 * x + S, data_out2);
- wrapper::vstore(output_ptr + 2 * x + offset_y_out, data_out1);
- wrapper::vstore(output_ptr + 2 * x + offset_y_out + S, data_out2);
- }
-
- // Compute left-over elements
- for(; x < window_end_x; ++x)
- {
- *(output_ptr + 2 * x) = *(input_ptr + x);
- *(output_ptr + 2 * x + 1) = *(input_ptr + x);
- *(output_ptr + 2 * x + offset_y_out) = *(input_ptr + x);
- *(output_ptr + 2 * x + offset_y_out + 1) = *(input_ptr + x);
- }
- },
- input, output);
-}
-
-template <typename T, int S>
-void NEUpsampleLayerKernel::upsample_nhwc(const arm_compute::Window &window)
-{
- using VectorType = typename wrapper::traits::neon_vector<T, S>::type;
-
- Window window_out(window);
- window_out.set(Window::DimX, Window::Dimension(0, 1, 1));
- window_out.set(Window::DimY, Window::Dimension(0, _output->info()->dimension(1), _info.x()));
- window_out.set(Window::DimZ, Window::Dimension(0, _output->info()->dimension(2), _info.y()));
-
- const auto window_start_x = static_cast<int>(window.x().start());
- const auto window_end_x = static_cast<int>(window.x().end());
- const int window_step_x = S;
-
- Window window_in{ window };
- window_in.set(Window::DimX, Window::Dimension(0, 1, 1));
-
- Iterator input(_input, window_in);
- Iterator output(_output, window_out);
-
- const int offset_y_out = _output->info()->strides_in_bytes().y() / sizeof(T);
- const int offset_z_out = _output->info()->strides_in_bytes().z() / sizeof(T);
-
- execute_window_loop(window_out, [&](const Coordinates &)
- {
- const auto input_ptr = reinterpret_cast<const T *>(input.ptr());
- const auto output_ptr = reinterpret_cast<T *>(output.ptr());
-
- int x = window_start_x;
- for(; x <= (window_end_x - window_step_x); x += window_step_x)
- {
- const VectorType data = wrapper::vloadq(reinterpret_cast<const T *>(input_ptr + x));
-
- wrapper::vstore(output_ptr + x, data);
- wrapper::vstore(output_ptr + x + offset_y_out, data);
- wrapper::vstore(output_ptr + x + offset_z_out, data);
- wrapper::vstore(output_ptr + x + offset_y_out + offset_z_out, data);
- }
-
- // Compute left-over elements
- for(; x < window_end_x; ++x)
- {
- *(output_ptr + x) = *(input_ptr + x);
- *(output_ptr + x + offset_y_out) = *(input_ptr + x);
- *(output_ptr + x + offset_z_out) = *(input_ptr + x);
- *(output_ptr + x + offset_y_out + offset_z_out) = *(input_ptr + x);
- }
- },
- input, output);
-}
-
-void NEUpsampleLayerKernel::configure(const ITensor *input, ITensor *output, const Size2D &info, const InterpolationPolicy policy)
-{
- ARM_COMPUTE_ERROR_ON_NULLPTR(input, output);
- ARM_COMPUTE_UNUSED(policy);
-
- _input = input;
- _output = output;
- _info = info;
-
- const DataLayout data_layout = input->info()->data_layout();
-
- TensorShape output_shape = misc::shape_calculator::compute_upsample_shape(*input->info(), info);
- auto_init_if_empty(*output->info(), output_shape, 1, input->info()->data_type());
- output->info()->set_data_layout(data_layout);
-
- // Perform validation step
- ARM_COMPUTE_ERROR_THROW_ON(NEUpsampleLayerKernel::validate(input->info(), output->info(), info, policy));
-
- switch(data_layout)
- {
- case DataLayout::NCHW:
- {
- switch(input->info()->data_type())
- {
- case DataType::QASYMM8_SIGNED:
- _func = &NEUpsampleLayerKernel::upsample_nchw<int8_t, 16>;
- break;
- case DataType::QASYMM8:
- _func = &NEUpsampleLayerKernel::upsample_nchw<uint8_t, 16>;
- break;
- case DataType::F32:
- _func = &NEUpsampleLayerKernel::upsample_nchw<float, 4>;
- break;
-#ifdef __ARM_FEATURE_FP16_VECTOR_ARITHMETIC
- case DataType::F16:
- _func = &NEUpsampleLayerKernel::upsample_nchw<float16_t, 8>;
- ;
- break;
-#endif /* __ARM_FEATURE_FP16_VECTOR_ARITHMETIC */
- default:
- ARM_COMPUTE_ERROR("Not implemented");
- }
- break;
- }
- case DataLayout::NHWC:
- {
- switch(input->info()->data_type())
- {
- case DataType::QASYMM8_SIGNED:
- _func = &NEUpsampleLayerKernel::upsample_nhwc<int8_t, 16>;
- break;
- case DataType::QASYMM8:
- _func = &NEUpsampleLayerKernel::upsample_nhwc<uint8_t, 16>;
- break;
- case DataType::F32:
- _func = &NEUpsampleLayerKernel::upsample_nhwc<float, 4>;
- break;
-#ifdef __ARM_FEATURE_FP16_VECTOR_ARITHMETIC
- case DataType::F16:
- _func = &NEUpsampleLayerKernel::upsample_nhwc<float16_t, 8>;
- break;
-#endif /* __ARM_FEATURE_FP16_VECTOR_ARITHMETIC */
- default:
- ARM_COMPUTE_ERROR("Not implemented");
- }
- break;
- }
- default:
- ARM_COMPUTE_ERROR("Not implemented");
- }
-
- // Configure window
- Window win = calculate_max_window(*input->info(), Steps());
- Coordinates coord;
- coord.set_num_dimensions(output->info()->num_dimensions());
- output->info()->set_valid_region(ValidRegion(coord, output->info()->tensor_shape()));
- INEKernel::configure(win);
-}
-
-void NEUpsampleLayerKernel::run(const Window &window, const ThreadInfo &info)
-{
- ARM_COMPUTE_UNUSED(info);
- ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this);
- ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(INEKernel::window(), window);
- ARM_COMPUTE_ERROR_ON(_func == nullptr);
-
- (this->*_func)(window);
-}
-} // namespace arm_compute
+++ /dev/null
-/*
- * Copyright (c) 2018-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_NEUPSAMPLELAYERKERNEL_H
-#define ARM_COMPUTE_NEUPSAMPLELAYERKERNEL_H
-
-#include "src/core/NEON/INEKernel.h"
-
-namespace arm_compute
-{
-class ITensor;
-
-/** Interface for the Upsample layer kernel.*/
-class NEUpsampleLayerKernel : public INEKernel
-{
-public:
- const char *name() const override
- {
- return "NEUpsampleLayerKernel";
- }
- /** Default constructor */
- NEUpsampleLayerKernel();
- /** Prevent instances of this class from being copied (As this class contains pointers) */
- NEUpsampleLayerKernel(const NEUpsampleLayerKernel &) = delete;
- /** Prevent instances of this class from being copied (As this class contains pointers) */
- NEUpsampleLayerKernel &operator=(const NEUpsampleLayerKernel &) = delete;
- /** Default Move Constructor. */
- NEUpsampleLayerKernel(NEUpsampleLayerKernel &&) = default;
- /** Default move assignment operator */
- NEUpsampleLayerKernel &operator=(NEUpsampleLayerKernel &&) = default;
- /** Default destructor */
- ~NEUpsampleLayerKernel() = default;
- /** Set the input output tensors.
- *
- * @param[in] input Source tensor. Data types supported: QASYMM8/QASYMM8_SIGNED/F16/F32.
- * @param[out] output Destination tensor. Data types supported: same as @p input.
- * @param[in] info Contains stride information described in @ref Size2D.
- * @param[in] policy Defines the policy to fill the intermediate pixels.
- *
- */
- void configure(const ITensor *input, ITensor *output, const Size2D &info, const InterpolationPolicy policy);
- /** Static function to check if given info will lead to a valid configuration of @ref NEUpsampleLayerKernel
- *
- * @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] info Contains stride information described in @ref Size2D.
- * @param[in] policy Defines the policy to fill the intermediate pixels.
- *
- * @return a status
- */
- static Status validate(const ITensorInfo *input, const ITensorInfo *output, const Size2D &info, const InterpolationPolicy policy);
-
- // Inherited methods overridden:
- void run(const Window &window, const ThreadInfo &info) override;
-
-private:
- /** Function to run upsample layer (NCHW)
- *
- * @param[in] window Region on which to execute the kernel. (Must be a valid region of the window returned by window()).
- */
- template <typename T, int S>
- void upsample_nchw(const Window &window);
- /** Function to run upsample layer (NHWC)
- *
- * @param[in] window Region on which to execute the kernel. (Must be a valid region of the window returned by window()).
- */
- template <typename T, int S>
- void upsample_nhwc(const Window &window);
-
- using UpsampleFunctionPtr = void (NEUpsampleLayerKernel::*)(const Window &window);
-
-private:
- UpsampleFunctionPtr _func;
- const ITensor *_input;
- ITensor *_output;
- Size2D _info;
-};
-} // namespace arm_compute
-#endif /*ARM_COMPUTE_NEUPSAMPLELAYERKERNEL_H */
return create_simple_multiple_input_single_output_node<StackLayerNode>(g, params, inputs, inputs.size(), axis);
}
-NodeID GraphBuilder::add_upsample_node(Graph &g, NodeParams params, NodeIdxPair input, Size2D info, InterpolationPolicy upsampling_policy)
-{
- return create_simple_single_input_output_node<UpsampleLayerNode>(g, params, input, info, upsampling_policy);
-}
-
NodeID GraphBuilder::add_yolo_node(Graph &g, NodeParams params, NodeIdxPair input, ActivationLayerInfo act_info)
{
check_nodeidx_pair(input, g);
return detail::create_stack_layer<CLStackLayer, CLTargetInfo>(*polymorphic_downcast<StackLayerNode *>(node));
case NodeType::StridedSliceLayer:
return detail::create_strided_slice_layer<CLStridedSlice, CLTargetInfo>(*polymorphic_downcast<StridedSliceLayerNode *>(node));
- case NodeType::UpsampleLayer:
- return detail::create_upsample_layer<CLUpsampleLayer, CLTargetInfo>(*polymorphic_downcast<UpsampleLayerNode *>(node), ctx);
default:
return nullptr;
}
return detail::validate_slice_layer<CLSlice>(*polymorphic_downcast<SliceLayerNode *>(node));
case NodeType::StridedSliceLayer:
return detail::validate_strided_slice_layer<CLStridedSlice>(*polymorphic_downcast<StridedSliceLayerNode *>(node));
- case NodeType::UpsampleLayer:
- return detail::validate_upsample_layer<CLUpsampleLayer>(*polymorphic_downcast<UpsampleLayerNode *>(node));
case NodeType::EltwiseLayer:
return detail::validate_eltwise_Layer<CLEltwiseLayerFunctions>(*polymorphic_downcast<EltwiseLayerNode *>(node));
case NodeType::UnaryEltwiseLayer:
return ARM_COMPUTE_CREATE_ERROR(arm_compute::ErrorCode::RUNTIME_ERROR, "Unsupported operation : ROIAlignLayer");
case NodeType::SliceLayer:
return ARM_COMPUTE_CREATE_ERROR(arm_compute::ErrorCode::RUNTIME_ERROR, "Unsupported operation : SliceLayer");
- case NodeType::UpsampleLayer:
- return ARM_COMPUTE_CREATE_ERROR(arm_compute::ErrorCode::RUNTIME_ERROR, "Unsupported operation : UpsampleLayer");
- case NodeType::YOLOLayer:
- return ARM_COMPUTE_CREATE_ERROR(arm_compute::ErrorCode::RUNTIME_ERROR, "Unsupported operation : YOLOLayer");
default:
return Status{};
}
return detail::create_stack_layer<NEStackLayer, NETargetInfo>(*polymorphic_downcast<StackLayerNode *>(node));
case NodeType::StridedSliceLayer:
return detail::create_strided_slice_layer<NEStridedSlice, NETargetInfo>(*polymorphic_downcast<StridedSliceLayerNode *>(node));
- case NodeType::UpsampleLayer:
- return detail::create_upsample_layer<NEUpsampleLayer, NETargetInfo>(*polymorphic_downcast<UpsampleLayerNode *>(node), ctx);
default:
return nullptr;
}
return detail::validate_slice_layer<NESlice>(*polymorphic_downcast<SliceLayerNode *>(node));
case NodeType::StridedSliceLayer:
return detail::validate_strided_slice_layer<NEStridedSlice>(*polymorphic_downcast<StridedSliceLayerNode *>(node));
- case NodeType::UpsampleLayer:
- return detail::validate_upsample_layer<NEUpsampleLayer>(*polymorphic_downcast<UpsampleLayerNode *>(node));
case NodeType::EltwiseLayer:
return detail::validate_eltwise_Layer<NEEltwiseLayerFunctions>(*polymorphic_downcast<EltwiseLayerNode *>(node));
case NodeType::UnaryEltwiseLayer:
+++ /dev/null
-/*
- * Copyright (c) 2018-2019 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 "arm_compute/graph/nodes/UpsampleLayerNode.h"
-
-#include "arm_compute/graph/Graph.h"
-#include "arm_compute/graph/INodeVisitor.h"
-#include "arm_compute/graph/Utils.h"
-
-namespace arm_compute
-{
-namespace graph
-{
-UpsampleLayerNode::UpsampleLayerNode(Size2D info, InterpolationPolicy upsampling_policy)
- : _info(info), _upsampling_policy(upsampling_policy)
-{
- _input_edges.resize(1, EmptyEdgeID);
- _outputs.resize(1, NullTensorID);
-}
-
-Size2D UpsampleLayerNode::info() const
-{
- return _info;
-}
-
-InterpolationPolicy UpsampleLayerNode::upsampling_policy() const
-{
- return _upsampling_policy;
-}
-
-TensorDescriptor UpsampleLayerNode::compute_output_descriptor(const TensorDescriptor &input_descriptor,
- Size2D info)
-{
- const unsigned int input_width = get_dimension_size(input_descriptor, DataLayoutDimension::WIDTH);
- const unsigned int input_height = get_dimension_size(input_descriptor, DataLayoutDimension::HEIGHT);
-
- const DataLayout data_layout = input_descriptor.layout;
- TensorDescriptor output_descriptor = input_descriptor;
- output_descriptor.shape.set(get_dimension_idx(data_layout, DataLayoutDimension::WIDTH), input_width * info.x());
- output_descriptor.shape.set(get_dimension_idx(data_layout, DataLayoutDimension::HEIGHT), input_height * info.y());
-
- return output_descriptor;
-}
-
-bool UpsampleLayerNode::forward_descriptors()
-{
- if((input_id(0) != NullTensorID) && (output_id(0) != NullTensorID))
- {
- Tensor *dst = output(0);
- ARM_COMPUTE_ERROR_ON(dst == nullptr);
- dst->desc() = configure_output(0);
- return true;
- }
- return false;
-}
-
-TensorDescriptor UpsampleLayerNode::configure_output(size_t idx) const
-{
- ARM_COMPUTE_UNUSED(idx);
- ARM_COMPUTE_ERROR_ON(idx >= _outputs.size());
-
- const Tensor *src = input(0);
- ARM_COMPUTE_ERROR_ON(src == nullptr);
-
- return compute_output_descriptor(src->desc(), _info);
-}
-
-NodeType UpsampleLayerNode::type() const
-{
- return NodeType::UpsampleLayer;
-}
-
-void UpsampleLayerNode::accept(INodeVisitor &v)
-{
- v.visit(*this);
-}
-} // namespace graph
-} // namespace arm_compute
\ No newline at end of file
+++ /dev/null
-/*
- * Copyright (c) 2018-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 "arm_compute/runtime/CL/functions/CLUpsampleLayer.h"
-
-#include "arm_compute/core/CL/OpenCL.h"
-#include "arm_compute/core/Utils.h"
-#include "arm_compute/runtime/CL/CLScheduler.h"
-#include "src/core/CL/kernels/CLUpsampleLayerKernel.h"
-
-namespace arm_compute
-{
-CLUpsampleLayer::CLUpsampleLayer() // NOLINT
- : _upsample(std::make_unique<CLUpsampleLayerKernel>()),
- _output(nullptr)
-{
-}
-
-CLUpsampleLayer::~CLUpsampleLayer() = default;
-
-Status CLUpsampleLayer::validate(const ITensorInfo *input, const ITensorInfo *output,
- const Size2D &info, const InterpolationPolicy upsampling_policy)
-{
- return CLUpsampleLayerKernel::validate(input, output, info, upsampling_policy);
-}
-
-void CLUpsampleLayer::configure(ICLTensor *input, ICLTensor *output,
- const Size2D &info, const InterpolationPolicy upsampling_policy)
-{
- configure(CLKernelLibrary::get().get_compile_context(), input, output, info, upsampling_policy);
-}
-
-void CLUpsampleLayer::configure(const CLCompileContext &compile_context, ICLTensor *input, ICLTensor *output,
- const Size2D &info, const InterpolationPolicy upsampling_policy)
-{
- ARM_COMPUTE_ERROR_ON_NULLPTR(input, output);
-
- _output = output;
- _upsample->configure(compile_context, input, _output, info, upsampling_policy);
-}
-
-void CLUpsampleLayer::run()
-{
- CLScheduler::get().enqueue(*_upsample, false);
-}
-} // namespace arm_compute
+++ /dev/null
-/*
- * Copyright (c) 2018-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 "arm_compute/runtime/NEON/functions/NEUpsampleLayer.h"
-
-#include "src/core/NEON/kernels/NEUpsampleLayerKernel.h"
-
-namespace arm_compute
-{
-NEUpsampleLayer::~NEUpsampleLayer() = default;
-
-NEUpsampleLayer::NEUpsampleLayer()
- : _kernel(), _data_layout()
-{
-}
-
-Status NEUpsampleLayer::validate(const ITensorInfo *input, const ITensorInfo *output, const Size2D &info,
- const InterpolationPolicy &policy)
-{
- return NEUpsampleLayerKernel::validate(input, output, info, policy);
-}
-
-void NEUpsampleLayer::configure(const ITensor *input, ITensor *output, const Size2D &info, const InterpolationPolicy &policy)
-{
- _data_layout = input->info()->data_layout();
- _kernel = std::make_unique<NEUpsampleLayerKernel>();
- _kernel->configure(input, output, info, policy);
-}
-
-void NEUpsampleLayer::run()
-{
- const auto win = (_data_layout == DataLayout::NCHW) ? Window::DimZ : Window::DimX;
- NEScheduler::get().schedule(_kernel.get(), win);
-}
-} // namespace arm_compute
+++ /dev/null
-/*
- * Copyright (c) 2018-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 "arm_compute/core/Types.h"
-#include "arm_compute/runtime/CL/CLTensor.h"
-#include "arm_compute/runtime/CL/CLTensorAllocator.h"
-#include "arm_compute/runtime/CL/functions/CLUpsampleLayer.h"
-#include "tests/CL/CLAccessor.h"
-#include "tests/PaddingCalculator.h"
-#include "tests/datasets/ShapeDatasets.h"
-#include "tests/framework/Asserts.h"
-#include "tests/framework/Macros.h"
-#include "tests/framework/datasets/Datasets.h"
-#include "tests/validation/Validation.h"
-#include "tests/validation/fixtures/UpsampleLayerFixture.h"
-
-namespace arm_compute
-{
-namespace test
-{
-namespace validation
-{
-namespace
-{
-constexpr AbsoluteTolerance<float> tolerance(0.001f);
-} // namespace
-
-TEST_SUITE(CL)
-TEST_SUITE(UpsampleLayer)
-
-// *INDENT-OFF*
-// clang-format off
-DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip(zip(
- framework::dataset::make("InputInfo", { TensorInfo(TensorShape(10U, 10U, 2U), 1, DataType::F32), // Mismatching data type
- TensorInfo(TensorShape(10U, 10U, 2U), 1, DataType::F32), // Invalid output shape
- TensorInfo(TensorShape(10U, 10U, 2U), 1, DataType::F32), // Invalid stride
- TensorInfo(TensorShape(10U, 10U, 2U), 1, DataType::F32), // Invalid policy
- TensorInfo(TensorShape(10U, 10U, 2U), 1, DataType::F32),
- }),
- framework::dataset::make("OutputInfo",{ TensorInfo(TensorShape(20U, 20U, 2U), 1, DataType::F16),
- TensorInfo(TensorShape(20U, 10U, 2U), 1, DataType::F32),
- TensorInfo(TensorShape(20U, 20U, 2U), 1, DataType::F32),
- TensorInfo(TensorShape(20U, 20U, 2U), 1, DataType::F32),
- TensorInfo(TensorShape(20U, 20U, 2U), 1, DataType::F32),
- })),
- framework::dataset::make("PadInfo", { Size2D(2, 2),
- Size2D(2, 2),
- Size2D(1, 1),
- Size2D(2, 2),
- Size2D(2, 2),
- })),
- framework::dataset::make("UpsamplingPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR,
- InterpolationPolicy::NEAREST_NEIGHBOR,
- InterpolationPolicy::NEAREST_NEIGHBOR,
- InterpolationPolicy::BILINEAR,
- InterpolationPolicy::NEAREST_NEIGHBOR,
- })),
- framework::dataset::make("Expected", { false, false, false, false, true })),
- input_info, output_info, pad_info, upsampling_policy, expected)
-{
- bool is_valid = bool(CLUpsampleLayer::validate(&input_info.clone()->set_is_resizable(false), &output_info.clone()->set_is_resizable(false), pad_info, upsampling_policy));
- ARM_COMPUTE_EXPECT(is_valid == expected, framework::LogLevel::ERRORS);
-}
-// clang-format on
-// *INDENT-ON*
-
-TEST_SUITE(Float)
-template <typename T>
-using CLUpsampleLayerFixture = UpsampleLayerFixture<CLTensor, CLAccessor, CLUpsampleLayer, T>;
-TEST_SUITE(FP32)
-FIXTURE_DATA_TEST_CASE(RunSmall, CLUpsampleLayerFixture<float>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(),
- framework::dataset::make("DataType", DataType::F32)),
- framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC })),
- framework::dataset::make("PadInfo", { Size2D(2, 2) })),
- framework::dataset::make("UpsamplingPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR })))
-{
- // Validate output
- validate(CLAccessor(_target), _reference, tolerance);
-}
-TEST_SUITE_END() // FP32
-
-TEST_SUITE(FP16)
-
-FIXTURE_DATA_TEST_CASE(RunSmall, CLUpsampleLayerFixture<half>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(),
- framework::dataset::make("DataType",
- DataType::F16)),
- framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC })),
- framework::dataset::make("PadInfo", { Size2D(2, 2) })),
- framework::dataset::make("UpsamplingPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR })))
-{
- // Validate output
- validate(CLAccessor(_target), _reference, tolerance);
-}
-
-TEST_SUITE_END() // FP16
-TEST_SUITE_END() // Float
-
-TEST_SUITE(Quantized)
-template <typename T>
-using CLUpsampleLayerQuantizedFixture = UpsampleLayerQuantizedFixture<CLTensor, CLAccessor, CLUpsampleLayer, T>;
-TEST_SUITE(QASYMM8)
-FIXTURE_DATA_TEST_CASE(RunSmall, CLUpsampleLayerQuantizedFixture<uint8_t>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(datasets::SmallShapes(),
- framework::dataset::make("DataType", DataType::QASYMM8)),
- framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC })),
- framework::dataset::make("PadInfo", { Size2D(2, 2) })),
- framework::dataset::make("UpsamplingPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR })),
- framework::dataset::make("QuantizationInfo", { QuantizationInfo(1.f / 255.f, 10) })))
-{
- // Validate output
- validate(CLAccessor(_target), _reference, tolerance);
-}
-TEST_SUITE_END() // QASYMM8
-TEST_SUITE(QASYMM8_SIGNED)
-FIXTURE_DATA_TEST_CASE(RunSmall, CLUpsampleLayerQuantizedFixture<int8_t>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(datasets::SmallShapes(),
- framework::dataset::make("DataType", DataType::QASYMM8_SIGNED)),
- framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC })),
- framework::dataset::make("PadInfo", { Size2D(2, 2) })),
- framework::dataset::make("UpsamplingPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR })),
- framework::dataset::make("QuantizationInfo", { QuantizationInfo(1.f / 255.f, 10) })))
-{
- // Validate output
- validate(CLAccessor(_target), _reference, tolerance);
-}
-TEST_SUITE_END() // QASYMM8_SIGNED
-TEST_SUITE_END() // Quantized
-
-TEST_SUITE_END() // UpsampleLayer
-TEST_SUITE_END() // CL
-} // namespace validation
-} // namespace test
-} // namespace arm_compute
+++ /dev/null
-/*
- * Copyright (c) 2018-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 "arm_compute/core/Types.h"
-#include "arm_compute/runtime/NEON/functions/NEUpsampleLayer.h"
-#include "tests/NEON/Accessor.h"
-#include "tests/PaddingCalculator.h"
-#include "tests/datasets/ShapeDatasets.h"
-#include "tests/framework/Asserts.h"
-#include "tests/framework/Macros.h"
-#include "tests/framework/datasets/Datasets.h"
-#include "tests/validation/Validation.h"
-#include "tests/validation/fixtures/UpsampleLayerFixture.h"
-
-namespace arm_compute
-{
-namespace test
-{
-namespace validation
-{
-TEST_SUITE(NEON)
-TEST_SUITE(UpsampleLayer)
-
-// *INDENT-OFF*
-// clang-format off
-DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip(zip(
- framework::dataset::make("InputInfo", { TensorInfo(TensorShape(10U, 10U, 2U), 1, DataType::F32), // Mismatching data type
- TensorInfo(TensorShape(10U, 10U, 2U), 1, DataType::F32), // Invalid output shape
- TensorInfo(TensorShape(10U, 10U, 2U), 1, DataType::F32), // Invalid stride
- TensorInfo(TensorShape(10U, 10U, 2U), 1, DataType::F32), // Invalid policy
- TensorInfo(TensorShape(32U, 32U), 1, DataType::F32),
- }),
- framework::dataset::make("OutputInfo",{ TensorInfo(TensorShape(20U, 20U, 2U), 1, DataType::F16),
- TensorInfo(TensorShape(20U, 10U, 2U), 1, DataType::F32),
- TensorInfo(TensorShape(20U, 20U, 2U), 1, DataType::F32),
- TensorInfo(TensorShape(20U, 20U, 2U), 1, DataType::F32),
- TensorInfo(TensorShape(64U, 64U), 1, DataType::F32),
- })),
- framework::dataset::make("PadInfo", { Size2D(2, 2),
- Size2D(2, 2),
- Size2D(1, 1),
- Size2D(2, 2),
- Size2D(2, 2),
- })),
- framework::dataset::make("UpsamplingPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR,
- InterpolationPolicy::NEAREST_NEIGHBOR,
- InterpolationPolicy::NEAREST_NEIGHBOR,
- InterpolationPolicy::BILINEAR,
- InterpolationPolicy::NEAREST_NEIGHBOR,
- })),
- framework::dataset::make("Expected", { false, false, false, false, true })),
- input_info, output_info, pad_info, policy, expected)
-{
- bool is_valid = bool(NEUpsampleLayer::validate(&input_info.clone()->set_is_resizable(false), &output_info.clone()->set_is_resizable(false), pad_info, policy));
- ARM_COMPUTE_EXPECT(is_valid == expected, framework::LogLevel::ERRORS);
-}
-// clang-format on
-// *INDENT-ON*
-
-template <typename T>
-using NEUpsampleLayerFixture = UpsampleLayerFixture<Tensor, Accessor, NEUpsampleLayer, T>;
-
-template <typename T>
-using NEUpsampleLayerQuantizedFixture = UpsampleLayerQuantizedFixture<Tensor, Accessor, NEUpsampleLayer, T>;
-
-TEST_SUITE(Float)
-TEST_SUITE(FP32)
-FIXTURE_DATA_TEST_CASE(RunSmall, NEUpsampleLayerFixture<float>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(),
- framework::dataset::make("DataType", DataType::F32)),
- framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC })),
- framework::dataset::make("PadInfo", { Size2D(2, 2) })),
- framework::dataset::make("UpsamplingPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR })))
-{
- // Validate output
- validate(Accessor(_target), _reference);
-}
-TEST_SUITE_END() // FP32
-
-#ifdef __ARM_FEATURE_FP16_VECTOR_ARITHMETIC
-TEST_SUITE(FP16)
-FIXTURE_DATA_TEST_CASE(RunSmall, NEUpsampleLayerFixture<half>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(),
- framework::dataset::make("DataType",
- DataType::F16)),
- framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC })),
- framework::dataset::make("PadInfo", { Size2D(2, 2) })),
- framework::dataset::make("UpsamplingPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR })))
-{
- // Validate output
- validate(Accessor(_target), _reference);
-}
-TEST_SUITE_END() // FP16
-#endif /* __ARM_FEATURE_FP16_VECTOR_ARITHMETIC */
-TEST_SUITE_END() // Float
-
-TEST_SUITE(Quantized)
-TEST_SUITE(QASYMM8)
-FIXTURE_DATA_TEST_CASE(RunSmall, NEUpsampleLayerQuantizedFixture<uint8_t>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(datasets::SmallShapes(),
- framework::dataset::make("DataType", DataType::QASYMM8)),
- framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC })),
- framework::dataset::make("PadInfo", { Size2D(2, 2) })),
- framework::dataset::make("UpsamplingPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR })),
- framework::dataset::make("QuantizationInfo", QuantizationInfo(2.f / 255.f, 10))))
-{
- // Validate output
- validate(Accessor(_target), _reference);
-}
-TEST_SUITE_END() // QASYMM8
-TEST_SUITE(QASYMM8_SIGNED)
-FIXTURE_DATA_TEST_CASE(RunSmall, NEUpsampleLayerQuantizedFixture<int8_t>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(datasets::SmallShapes(),
- framework::dataset::make("DataType", DataType::QASYMM8_SIGNED)),
- framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC })),
- framework::dataset::make("PadInfo", { Size2D(2, 2) })),
- framework::dataset::make("UpsamplingPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR })),
- framework::dataset::make("QuantizationInfo", QuantizationInfo(2.f / 255.f, 10))))
-{
- // Validate output
- validate(Accessor(_target), _reference);
-}
-TEST_SUITE_END() // QASYMM8_SIGNED
-TEST_SUITE_END() // Quantized
-
-TEST_SUITE_END() // UpsampleLayer
-TEST_SUITE_END() // NEON
-
-} // namespace validation
-} // namespace test
-} // namespace arm_compute
+++ /dev/null
-/*
- * Copyright (c) 2018-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 "arm_compute/core/TensorShape.h"
-#include "arm_compute/core/Types.h"
-#include "tests/AssetsLibrary.h"
-#include "tests/Globals.h"
-#include "tests/IAccessor.h"
-#include "tests/framework/Asserts.h"
-#include "tests/framework/Fixture.h"
-#include "tests/validation/Helpers.h"
-#include "tests/validation/reference/UpsampleLayer.h"
-
-#include <random>
-
-namespace arm_compute
-{
-namespace test
-{
-namespace validation
-{
-template <typename TensorType, typename AccessorType, typename FunctionType, typename T>
-class UpsampleLayerFixtureBase : public framework::Fixture
-{
-public:
- template <typename...>
- void setup(TensorShape input_shape, DataType data_type, DataLayout data_layout,
- Size2D info, const InterpolationPolicy &policy, QuantizationInfo quantization_info)
- {
- _data_type = data_type;
-
- _target = compute_target(input_shape, info, policy, data_type, data_layout, quantization_info);
- _reference = compute_reference(input_shape, info, policy, data_type, quantization_info);
- }
-
-protected:
- template <typename U>
- void fill(U &&tensor, int i)
- {
- library->fill_tensor_uniform(tensor, i);
- }
-
- TensorType compute_target(TensorShape input_shape, const Size2D &info, const InterpolationPolicy &policy,
- DataType data_type, DataLayout data_layout, QuantizationInfo quantization_info)
- {
- TensorShape output_shape(input_shape);
- output_shape.set(0, info.x() * input_shape[0]);
- output_shape.set(1, info.y() * input_shape[1]);
-
- if(data_layout == DataLayout::NHWC)
- {
- permute(input_shape, PermutationVector(2U, 0U, 1U));
- permute(output_shape, PermutationVector(2U, 0U, 1U));
- }
-
- // Create tensors
- TensorType src = create_tensor<TensorType>(input_shape, data_type, 1, quantization_info, data_layout);
- TensorType dst = create_tensor<TensorType>(output_shape, data_type, 1, quantization_info, data_layout);
-
- // Create and configure function
- FunctionType upsample;
- upsample.configure(&src, &dst, info, policy);
-
- ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS);
- ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS);
-
- // Allocate tensors
- src.allocator()->allocate();
- dst.allocator()->allocate();
-
- ARM_COMPUTE_EXPECT(!src.info()->is_resizable(), framework::LogLevel::ERRORS);
- ARM_COMPUTE_EXPECT(!dst.info()->is_resizable(), framework::LogLevel::ERRORS);
-
- // Fill tensors
- fill(AccessorType(src), 0);
-
- // Compute DeconvolutionLayer function
- upsample.run();
-
- return dst;
- }
-
- SimpleTensor<T> compute_reference(const TensorShape &input_shape, const Size2D &info, const InterpolationPolicy &policy,
- DataType data_type, QuantizationInfo quantization_info)
- {
- // Create reference
- SimpleTensor<T> src{ input_shape, data_type, 1, quantization_info };
-
- // Fill reference
- fill(src, 0);
-
- return reference::upsample_layer<T>(src, info, policy);
- }
-
- TensorType _target{};
- SimpleTensor<T> _reference{};
- DataType _data_type{};
-};
-
-template <typename TensorType, typename AccessorType, typename FunctionType, typename T>
-class UpsampleLayerFixture : public UpsampleLayerFixtureBase<TensorType, AccessorType, FunctionType, T>
-{
-public:
- template <typename...>
- void setup(TensorShape input_shape, DataType data_type, DataLayout data_layout,
- Size2D info, const InterpolationPolicy &policy)
- {
- UpsampleLayerFixtureBase<TensorType, AccessorType, FunctionType, T>::setup(input_shape, data_type, data_layout,
- info, policy, QuantizationInfo());
- }
-};
-
-template <typename TensorType, typename AccessorType, typename FunctionType, typename T>
-class UpsampleLayerQuantizedFixture : public UpsampleLayerFixtureBase<TensorType, AccessorType, FunctionType, T>
-{
-public:
- template <typename...>
- void setup(TensorShape input_shape, DataType data_type, DataLayout data_layout,
- Size2D info, const InterpolationPolicy &policy, QuantizationInfo quantization_info)
- {
- UpsampleLayerFixtureBase<TensorType, AccessorType, FunctionType, T>::setup(input_shape, data_type, data_layout,
- info, policy, quantization_info);
- }
-};
-
-} // namespace validation
-} // namespace test
-} // namespace arm_compute
+++ /dev/null
-/*
- * Copyright (c) 2018-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 "UpsampleLayer.h"
-
-#include "support/Requires.h"
-#include "tests/validation/Helpers.h"
-
-namespace arm_compute
-{
-namespace test
-{
-namespace validation
-{
-namespace reference
-{
-template <typename T>
-SimpleTensor<T> upsample_layer(const SimpleTensor<T> &src, const Size2D &info, const InterpolationPolicy policy)
-{
- ARM_COMPUTE_ERROR_ON(policy != InterpolationPolicy::NEAREST_NEIGHBOR);
- ARM_COMPUTE_UNUSED(policy);
-
- TensorShape output_shape = src.shape();
- output_shape.set(0, src.shape().x() * info.x());
- output_shape.set(1, src.shape().y() * info.y());
-
- // Create reference
- const int stride_x = info.x();
- const int stride_y = info.y();
- int width_out = output_shape.x();
- int height_out = output_shape.y();
- SimpleTensor<T> out{ output_shape, src.data_type(), 1, src.quantization_info() };
-
- const int width_in = src.shape().x();
- const int height_in = src.shape().y();
- const int num_2d_slices = src.shape().total_size() / (width_in * height_in);
-
- for(int slice = 0; slice < num_2d_slices; ++slice)
- {
- const int offset_slice_in = slice * width_in * height_in;
- const int offset_slice_out = slice * height_out * width_out;
- for(int y = 0; y < height_out; ++y)
- {
- for(int x = 0; x < width_out; ++x)
- {
- const int out_offset = y * width_out + x;
- const int in_offset = (y / stride_y) * width_in + x / stride_x;
-
- T *_out = out.data() + offset_slice_out + out_offset;
- const T *in = src.data() + offset_slice_in + in_offset;
- *_out = *in;
- }
- }
- }
- return out;
-}
-
-template SimpleTensor<float> upsample_layer(const SimpleTensor<float> &src,
- const Size2D &info, const InterpolationPolicy policy);
-template SimpleTensor<half> upsample_layer(const SimpleTensor<half> &src,
- const Size2D &info, const InterpolationPolicy policy);
-template SimpleTensor<uint8_t> upsample_layer(const SimpleTensor<uint8_t> &src,
- const Size2D &info, const InterpolationPolicy policy);
-template SimpleTensor<int8_t> upsample_layer(const SimpleTensor<int8_t> &src,
- const Size2D &info, const InterpolationPolicy policy);
-} // namespace reference
-} // namespace validation
-} // namespace test
-} // namespace arm_compute
+++ /dev/null
-/*
- * Copyright (c) 2018-2019 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_UPSAMPLE_LAYER_H
-#define ARM_COMPUTE_TEST_UPSAMPLE_LAYER_H
-
-#include "tests/SimpleTensor.h"
-#include "tests/validation/Helpers.h"
-
-namespace arm_compute
-{
-namespace test
-{
-namespace validation
-{
-namespace reference
-{
-template <typename T>
-SimpleTensor<T> upsample_layer(const SimpleTensor<T> &src,
- const Size2D &info, const InterpolationPolicy policy);
-} // namespace reference
-} // namespace validation
-} // namespace test
-} // namespace arm_compute
-#endif /* ARM_COMPUTE_TEST_UPSAMPLE_LAYER_H */
projects / platform / upstream / armcl.git / commitdiff