aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorMichalis Spyrou <michalis.spyrou@arm.com>2018-08-29 16:28:11 +0100
committerMichalis Spyrou <michalis.spyrou@arm.com>2018-11-14 15:42:37 +0000
commit6c7c38e70c795077ba727aadeefc670888bec089 (patch)
tree7969d7976b22b881f2c08e97a772a4537e203629
parentd8e340d2fef39c78a9afda93b210a8640145df87 (diff)
downloadComputeLibrary-6c7c38e70c795077ba727aadeefc670888bec089.tar.gz
COMPMID-1462 SSD support: Create CL PriorBox
Change-Id: I5bf5d751ec7c02d96c26a769f49d03ea23a248b7
-rw-r--r--arm_compute/core/CL/CLKernels.h1
-rw-r--r--arm_compute/core/CL/kernels/CLPriorBoxLayerKernel.h86
-rw-r--r--arm_compute/core/Types.h122
-rw-r--r--arm_compute/core/utils/misc/ShapeCalculator.h14
-rw-r--r--arm_compute/runtime/CL/CLFunctions.h1
-rw-r--r--arm_compute/runtime/CL/functions/CLPriorBoxLayer.h68
-rw-r--r--src/core/CL/CLKernelLibrary.cpp6
-rw-r--r--src/core/CL/cl_kernels/prior_box_layer.cl256
-rw-r--r--src/core/CL/kernels/CLPriorBoxLayerKernel.cpp275
-rw-r--r--src/runtime/CL/functions/CLPriorBoxLayer.cpp58
-rw-r--r--tests/datasets/PriorBoxLayerDataset.h134
-rw-r--r--tests/validation/CL/PriorBoxLayer.cpp103
-rw-r--r--tests/validation/fixtures/PriorBoxLayerFixture.h112
-rw-r--r--tests/validation/reference/PriorBoxLayer.cpp158
-rw-r--r--tests/validation/reference/PriorBoxLayer.h44
-rw-r--r--utils/TypePrinter.h47
16 files changed, 1485 insertions, 0 deletions
diff --git a/arm_compute/core/CL/CLKernels.h b/arm_compute/core/CL/CLKernels.h
index df76366a4b..12700192e6 100644
--- a/arm_compute/core/CL/CLKernels.h
+++ b/arm_compute/core/CL/CLKernels.h
@@ -105,6 +105,7 @@
#include "arm_compute/core/CL/kernels/CLPermuteKernel.h"
#include "arm_compute/core/CL/kernels/CLPixelWiseMultiplicationKernel.h"
#include "arm_compute/core/CL/kernels/CLPoolingLayerKernel.h"
+#include "arm_compute/core/CL/kernels/CLPriorBoxLayerKernel.h"
#include "arm_compute/core/CL/kernels/CLQuantizationLayerKernel.h"
#include "arm_compute/core/CL/kernels/CLROIAlignLayerKernel.h"
#include "arm_compute/core/CL/kernels/CLROIPoolingLayerKernel.h"
diff --git a/arm_compute/core/CL/kernels/CLPriorBoxLayerKernel.h b/arm_compute/core/CL/kernels/CLPriorBoxLayerKernel.h
new file mode 100644
index 0000000000..a5423e4af0
--- /dev/null
+++ b/arm_compute/core/CL/kernels/CLPriorBoxLayerKernel.h
@@ -0,0 +1,86 @@
+/*
+ * Copyright (c) 2018 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_CLPRIORBOXLAYERKERNEL_H__
+#define __ARM_COMPUTE_CLPRIORBOXLAYERKERNEL_H__
+
+#include "arm_compute/core/CL/ICLKernel.h"
+
+namespace arm_compute
+{
+class ICLTensor;
+
+/** Interface for the PriorBox layer kernel. */
+class CLPriorBoxLayerKernel : public ICLKernel
+{
+public:
+ /** Constructor */
+ CLPriorBoxLayerKernel();
+ /** Prevent instances of this class from being copied (As this class contains pointers) */
+ CLPriorBoxLayerKernel(const CLPriorBoxLayerKernel &) = delete;
+ /** Prevent instances of this class from being copied (As this class contains pointers) */
+ CLPriorBoxLayerKernel &operator=(const CLPriorBoxLayerKernel &) = delete;
+ /** Default Move Constructor. */
+ CLPriorBoxLayerKernel(CLPriorBoxLayerKernel &&) = default;
+ /** Default move assignment operator */
+ CLPriorBoxLayerKernel &operator=(CLPriorBoxLayerKernel &&) = default;
+ /** Default destructor */
+ ~CLPriorBoxLayerKernel() = default;
+
+ /** Set the input and output tensors.
+ *
+ * @param[in] input1 First source tensor. Data types supported: F32. Data layouts supported: NCHW/NHWC.
+ * @param[in] input2 Second source tensor. Data types and layouts supported: same as @p input1
+ * @param[out] output Destination tensor. Output dimensions are [W * H * num_priors * 4, 2]. Data types and layouts supported: same as @p input1
+ * @param[in] info Prior box layer info.
+ * @param[in] min Minimum prior box values
+ * @param[in] max Maximum prior box values
+ * @param[in] aspect_ratios Aspect ratio values
+ */
+ void configure(const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output, const PriorBoxLayerInfo &info, cl::Buffer *min, cl::Buffer *max, cl::Buffer *aspect_ratios);
+ /** Static function to check if given info will lead to a valid configuration of @ref CLPriorBoxLayerKernel
+ *
+ * @param[in] input1 First source tensor info. Data types supported: F32. Data layouts supported: NCHW/NHWC.
+ * @param[in] input2 Second source tensor info. Data types and layouts supported: same as @p input1
+ * @param[in] output Destination tensor info. Output dimensions are [W * H * num_priors * 4, 2]. Data type supported: same as @p input1
+ * @param[in] info Prior box layer info.
+ *
+ * @return a status
+ */
+ static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const PriorBoxLayerInfo &info);
+
+ // Inherited methods overridden:
+ void run(const Window &window, cl::CommandQueue &queue) override;
+
+private:
+ const ICLTensor *_input1;
+ const ICLTensor *_input2;
+ ICLTensor *_output;
+ PriorBoxLayerInfo _info;
+ int _num_priors;
+ cl::Buffer *_min;
+ cl::Buffer *_max;
+ cl::Buffer *_aspect_ratios;
+};
+} // namespace arm_compute
+#endif /*__ARM_COMPUTE_CLPRIORBOXLAYERKERNEL_H__ */
diff --git a/arm_compute/core/Types.h b/arm_compute/core/Types.h
index 8df5c65e1e..03f195f7da 100644
--- a/arm_compute/core/Types.h
+++ b/arm_compute/core/Types.h
@@ -799,6 +799,128 @@ struct FullyConnectedLayerInfo
}
};
+/** PriorBox layer info */
+class PriorBoxLayerInfo final
+{
+public:
+ /** Default Constructor */
+ PriorBoxLayerInfo()
+ : _min_sizes(),
+ _variances(),
+ _offset(),
+ _flip(true),
+ _clip(false),
+ _max_sizes(),
+ _aspect_ratios(),
+ _img_size(),
+ _steps()
+ {
+ }
+ /** Constructor
+ *
+ * @param[in] min_sizes Min sizes vector.
+ * @param[in] variances Variances vector. Size must be equal to 4.
+ * @param[in] offset Offset value.
+ * @param[in] flip (Optional) Flip the aspect ratios.
+ * @param[in] clip (Optional) Clip coordinates so that they're within [0,1].
+ * @param[in] max_sizes (Optional) Max sizes vector.
+ * @param[in] aspect_ratios (Optional) Aspect ratios of the boxes.
+ * @param[in] img_size (Optional) Image size.
+ * @param[in] steps (Optional) Step values.
+ */
+ PriorBoxLayerInfo(const std::vector<float> &min_sizes, const std::vector<float> &variances, float offset, bool flip = true, bool clip = false,
+ const std::vector<float> &max_sizes = {}, const std::vector<float> &aspect_ratios = {}, const Coordinates2D &img_size = Coordinates2D{ 0, 0 }, const std::array<float, 2> &steps = { { 0.f, 0.f } })
+ : _min_sizes(min_sizes),
+ _variances(variances),
+ _offset(offset),
+ _flip(flip),
+ _clip(clip),
+ _max_sizes(max_sizes),
+ _aspect_ratios(aspect_ratios),
+ _img_size(img_size),
+ _steps(steps)
+ {
+ _aspect_ratios.push_back(1.);
+ for(unsigned int i = 0; i < aspect_ratios.size(); ++i)
+ {
+ float ar = aspect_ratios[i];
+ bool already_exist = false;
+ for(auto ar_new : _aspect_ratios)
+ {
+ if(fabs(ar - ar_new) < 1e-6)
+ {
+ already_exist = true;
+ break;
+ }
+ }
+ if(!already_exist)
+ {
+ _aspect_ratios.push_back(ar);
+ if(flip)
+ {
+ _aspect_ratios.push_back(1.f / ar);
+ }
+ }
+ }
+ }
+ /** Get min sizes. */
+ std::vector<float> min_sizes() const
+ {
+ return _min_sizes;
+ }
+ /** Get min variances. */
+ std::vector<float> variances() const
+ {
+ return _variances;
+ }
+ /** Get the step coordinates */
+ std::array<float, 2> steps() const
+ {
+ return _steps;
+ }
+ /** Get the image size coordinates */
+ Coordinates2D img_size() const
+ {
+ return _img_size;
+ }
+ /** Get the offset */
+ float offset() const
+ {
+ return _offset;
+ }
+ /** Get the flip value */
+ bool flip() const
+ {
+ return _flip;
+ }
+ /** Get the clip value */
+ bool clip() const
+ {
+ return _clip;
+ }
+ /** Get max sizes. */
+ std::vector<float> max_sizes() const
+ {
+ return _max_sizes;
+ }
+ /** Get aspect ratios. */
+ std::vector<float> aspect_ratios() const
+ {
+ return _aspect_ratios;
+ }
+
+private:
+ std::vector<float> _min_sizes;
+ std::vector<float> _variances;
+ float _offset;
+ bool _flip;
+ bool _clip;
+ std::vector<float> _max_sizes;
+ std::vector<float> _aspect_ratios;
+ Coordinates2D _img_size;
+ std::array<float, 2> _steps;
+};
+
/** Pooling Layer Information class */
class PoolingLayerInfo
{
diff --git a/arm_compute/core/utils/misc/ShapeCalculator.h b/arm_compute/core/utils/misc/ShapeCalculator.h
index 1f532ca31e..5c9457ed6b 100644
--- a/arm_compute/core/utils/misc/ShapeCalculator.h
+++ b/arm_compute/core/utils/misc/ShapeCalculator.h
@@ -646,6 +646,20 @@ inline TensorShape compute_space_to_batch_shape(const ITensorInfo *input, const
return output_shape;
}
+inline TensorShape compute_prior_box_shape(const ITensorInfo &input, const PriorBoxLayerInfo &info)
+{
+ DataLayout data_layout = input.data_layout();
+ const size_t idx_w = get_data_layout_dimension_index(data_layout, DataLayoutDimension::WIDTH);
+ const size_t idx_h = get_data_layout_dimension_index(data_layout, DataLayoutDimension::HEIGHT);
+
+ const int num_priors = info.aspect_ratios().size() * info.min_sizes().size() + info.max_sizes().size();
+
+ TensorShape output_shape{};
+ output_shape.set(0, input.dimension(idx_w) * input.dimension(idx_h) * num_priors * 4);
+ output_shape.set(1, 2);
+
+ return output_shape;
+}
inline TensorShape compute_padded_shape(const TensorShape &input_shape, const PaddingList &padding)
{
diff --git a/arm_compute/runtime/CL/CLFunctions.h b/arm_compute/runtime/CL/CLFunctions.h
index 7e36afd6f1..694e818788 100644
--- a/arm_compute/runtime/CL/CLFunctions.h
+++ b/arm_compute/runtime/CL/CLFunctions.h
@@ -102,6 +102,7 @@
#include "arm_compute/runtime/CL/functions/CLPhase.h"
#include "arm_compute/runtime/CL/functions/CLPixelWiseMultiplication.h"
#include "arm_compute/runtime/CL/functions/CLPoolingLayer.h"
+#include "arm_compute/runtime/CL/functions/CLPriorBoxLayer.h"
#include "arm_compute/runtime/CL/functions/CLQuantizationLayer.h"
#include "arm_compute/runtime/CL/functions/CLRNNLayer.h"
#include "arm_compute/runtime/CL/functions/CLROIAlignLayer.h"
diff --git a/arm_compute/runtime/CL/functions/CLPriorBoxLayer.h b/arm_compute/runtime/CL/functions/CLPriorBoxLayer.h
new file mode 100644
index 0000000000..2376cd3fc6
--- /dev/null
+++ b/arm_compute/runtime/CL/functions/CLPriorBoxLayer.h
@@ -0,0 +1,68 @@
+/*
+ * Copyright (c) 2018 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_CLPRIORBOXLAYER_H__
+#define __ARM_COMPUTE_CLPRIORBOXLAYER_H__
+
+#include "arm_compute/runtime/IFunction.h"
+
+#include "arm_compute/core/CL/kernels/CLPriorBoxLayerKernel.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/CL/ICLSimpleFunction.h"
+
+namespace arm_compute
+{
+class ICLTensor;
+
+/** Basic function to run @ref CLPriorBoxLayerKernel. */
+class CLPriorBoxLayer : public ICLSimpleFunction
+{
+public:
+ /** Constructor */
+ CLPriorBoxLayer();
+ /** Set the input and output tensors.
+ *
+ * @param[in] input1 First source tensor. Data types supported: F32. Data layouts supported: NCHW/NHWC.
+ * @param[in] input2 Second source tensor. Data types and layouts supported: same as @p input1
+ * @param[out] output Destination tensor. Output dimensions are [W * H * num_priors * 4, 2]. Data types and layouts supported: same as @p input1
+ * @param[in] info Prior box layer info.
+ */
+ void configure(const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output, const PriorBoxLayerInfo &info);
+ /** Static function to check if given info will lead to a valid configuration of @ref CLPriorBoxLayer
+ *
+ * @param[in] input1 First source tensor info. Data types supported: F32. Data layouts supported: NCHW/NHWC.
+ * @param[in] input2 Second source tensor info. Data types and layouts supported: same as @p input1
+ * @param[in] output Destination tensor info. Output dimensions are [W * H * num_priors * 4, 2]. Data types and layouts supported: same as @p input1
+ * @param[in] info Prior box layer info.
+ *
+ * @return a status
+ */
+ static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const PriorBoxLayerInfo &info);
+
+private:
+ cl::Buffer _min;
+ cl::Buffer _max;
+ cl::Buffer _aspect_ratios;
+};
+} // arm_compute
+#endif /* __ARM_COMPUTE_CLPRIORBOXLAYER_H__ */
diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp
index 3c2528f358..ccc9aec0d8 100644
--- a/src/core/CL/CLKernelLibrary.cpp
+++ b/src/core/CL/CLKernelLibrary.cpp
@@ -351,6 +351,8 @@ const std::map<std::string, std::string> CLKernelLibrary::_kernel_program_map =
{ "pooling_layer_MxN_nhwc", "pooling_layer.cl" },
{ "pooling_layer_MxN_quantized_nhwc", "pooling_layer_quantized.cl" },
{ "pooling_layer_MxN_quantized_nchw", "pooling_layer_quantized.cl" },
+ { "prior_box_layer_nchw", "prior_box_layer.cl" },
+ { "prior_box_layer_nhwc", "prior_box_layer.cl" },
{ "quantization_layer", "quantization_layer.cl" },
{ "reduction_operation_x", "reduction_operation.cl" },
{ "reduction_operation_quantized_x", "reduction_operation.cl" },
@@ -761,6 +763,10 @@ const std::map<std::string, std::string> CLKernelLibrary::_program_source_map =
#include "./cl_kernels/pooling_layer_quantized.clembed"
},
{
+ "prior_box_layer.cl",
+#include "./cl_kernels/prior_box_layer.clembed"
+ },
+ {
"quantization_layer.cl",
#include "./cl_kernels/quantization_layer.clembed"
},
diff --git a/src/core/CL/cl_kernels/prior_box_layer.cl b/src/core/CL/cl_kernels/prior_box_layer.cl
new file mode 100644
index 0000000000..2a0329d239
--- /dev/null
+++ b/src/core/CL/cl_kernels/prior_box_layer.cl
@@ -0,0 +1,256 @@
+/*
+ * Copyright (c) 2018 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 "helpers.h"
+
+#if defined(DATA_TYPE) && defined(WIDTH) && defined(HEIGHT) && defined(LAYER_WIDTH) && defined(LAYER_HEIGHT) && defined(OFFSET) && defined(STEP_X) && defined(STEP_Y) && defined(NUM_PRIORS) && defined(VARIANCE_0) && defined(VARIANCE_1) && defined(VARIANCE_2) && defined(VARIANCE_3)
+
+/** Compute prior boxes and clip (NCHW)
+ *
+ * @param[out] output_ptr Pointer to the destination tensor. Supported data types: F32
+ * @param[in] output_stride_x Stride of the destination tensor in X dimension (in bytes)
+ * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in] output_stride_y Stride of the destination tensor in Y dimension (in bytes)
+ * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes)
+ * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination tensor
+ * @param[in] idx Index to write to
+ * @param[in] center_x Center value of the x axis
+ * @param[in] center_y Center value of the y axis
+ * @param[in] box_width Prior box width
+ * @param[in] box_height Prior box height
+ *
+ */
+inline void calculate_xy_min_max_nchw(Image *out, int idx, float center_x, float center_y, float box_width, float box_height)
+{
+ float xmin = (center_x - box_width / 2.f) / WIDTH;
+ float ymin = (center_y - box_height / 2.f) / HEIGHT;
+ float xmax = (center_x + box_width / 2.f) / WIDTH;
+ float ymax = (center_y + box_height / 2.f) / HEIGHT;
+
+#if defined(CLIP)
+ xmin = clamp(xmin, 0.f, 1.f);
+ ymin = clamp(ymin, 0.f, 1.f);
+ xmax = clamp(xmax, 0.f, 1.f);
+ ymax = clamp(ymax, 0.f, 1.f);
+#endif // defined(CLIP)
+
+ // Store result
+ vstore4((VEC_DATA_TYPE(DATA_TYPE, 4))(xmin, ymin, xmax, ymax), 0, ((__global DATA_TYPE *)offset(out, idx + 0, 0)));
+}
+
+/** Compute prior boxes (NCHW)
+ *
+ * @param[out] output_ptr Pointer to the destination tensor. Supported data types: F32
+ * @param[in] output_stride_x Stride of the destination tensor in X dimension (in bytes)
+ * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in] output_stride_y Stride of the destination tensor in Y dimension (in bytes)
+ * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes)
+ * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination tensor
+ * @param[in] min_size Prior box min size
+ * @param[in] min_idx Index of the min vector
+ * @param[in] idx Index to write to
+ *
+ * @return The updated index
+ */
+inline int calculate_min_nchw(Image *out, float *max, float *aspect_ratios, int max_size, int aspect_ratios_size, float min_size, int min_idx, int idx)
+{
+ const float center_x = ((float)(get_global_id(0) % LAYER_WIDTH) + OFFSET) * STEP_X;
+ const float center_y = ((float)(get_global_id(0) / LAYER_WIDTH) + OFFSET) * STEP_Y;
+
+ float box_width = min_size;
+ float box_height = min_size;
+ calculate_xy_min_max_nchw(out, idx, center_x, center_y, box_width, box_height);
+ idx += 4;
+
+ if(max_size > 0)
+ {
+ box_width = sqrt(min_size * max[min_idx]);
+ box_height = box_width;
+ calculate_xy_min_max_nchw(out, idx, center_x, center_y, box_width, box_height);
+ idx += 4;
+ }
+ for(unsigned int i = 0; i < aspect_ratios_size; ++i)
+ {
+ if(fabs(aspect_ratios[i] - 1.f) < 1e-6f)
+ {
+ continue;
+ }
+ box_width = min_size * sqrt(aspect_ratios[i]);
+ box_height = min_size * rsqrt(aspect_ratios[i]);
+
+ calculate_xy_min_max_nchw(out, idx, center_x, center_y, box_width, box_height);
+ idx += 4;
+ }
+
+ return idx;
+}
+
+/** Compute prior boxes and clip (NHWC)
+ *
+ * @param[out] output_ptr Pointer to the destination tensor. Supported data types: F32
+ * @param[in] output_stride_x Stride of the destination tensor in X dimension (in bytes)
+ * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in] output_stride_y Stride of the destination tensor in Y dimension (in bytes)
+ * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes)
+ * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination tensor
+ * @param[in] idx Index to write to
+ * @param[in] center_x Center value of the x axis
+ * @param[in] center_y Center value of the y axis
+ * @param[in] box_width Prior box width
+ * @param[in] box_height Prior box height
+ *
+ */
+inline void calculate_xy_min_max_nhwc(Tensor3D *out, int idx, float center_x, float center_y, float box_width, float box_height)
+{
+ float xmin = (center_x - box_width / 2.f) / WIDTH;
+ float ymin = (center_y - box_height / 2.f) / HEIGHT;
+ float xmax = (center_x + box_width / 2.f) / WIDTH;
+ float ymax = (center_y + box_height / 2.f) / HEIGHT;
+
+#if defined(CLIP)
+ xmin = clamp(xmin, 0.f, 1.f);
+ ymin = clamp(ymin, 0.f, 1.f);
+ xmax = clamp(xmax, 0.f, 1.f);
+ ymax = clamp(ymax, 0.f, 1.f);
+#endif // defined(CLIP)
+
+ *((__global DATA_TYPE *)tensor3D_offset(out, 0, idx + 0, 0)) = xmin;
+ *((__global DATA_TYPE *)tensor3D_offset(out, 0, idx + 1, 0)) = ymin;
+ *((__global DATA_TYPE *)tensor3D_offset(out, 0, idx + 2, 0)) = xmax;
+ *((__global DATA_TYPE *)tensor3D_offset(out, 0, idx + 3, 0)) = ymax;
+}
+
+/** Compute prior boxes (NHWC)
+ *
+ * @param[in,out] out Tensor output
+ * @param[in] max The maximum values
+ * @param[in] aspect_ratios The aspect ratio values
+ * @param[in] max_size The maximum values values size
+ * @param[in] aspect_ratios_size The aspect ratio values size
+ * @param[in] min_size The minimum values size
+ * @param[in] min_idx Index of the min vector
+ * @param[in] idx Index to write to
+ *
+ * @return The updated index
+ */
+inline int calculate_min_nhwc(Image *out, float *max, float *aspect_ratios, int max_size, int aspect_ratios_size, float min_size, int min_idx, int idx)
+{
+ const float center_x = ((float)(get_global_id(1) % LAYER_WIDTH) + OFFSET) * STEP_X;
+ const float center_y = ((float)(get_global_id(1) / LAYER_WIDTH) + OFFSET) * STEP_Y;
+
+ float box_width = min_size;
+ float box_height = min_size;
+
+ calculate_xy_min_max_nhwc(out, idx, center_x, center_y, box_width, box_height);
+ idx += 4;
+ if(max_size > 0)
+ {
+ box_width = sqrt(min_size * max[min_idx]);
+ box_height = box_width;
+ calculate_xy_min_max_nhwc(out, idx, center_x, center_y, box_width, box_height);
+ idx += 4;
+ }
+ for(unsigned int i = 0; i < aspect_ratios_size; ++i)
+ {
+ if(fabs(aspect_ratios[i] - 1.f) < 1e-6f)
+ {
+ continue;
+ }
+ box_width = min_size * sqrt(aspect_ratios[i]);
+ box_height = min_size * rsqrt(aspect_ratios[i]);
+
+ calculate_xy_min_max_nhwc(out, idx, center_x, center_y, box_width, box_height);
+ idx += 4;
+ }
+
+ return idx;
+}
+
+/** Calculate prior boxes with NCHW format.
+ *
+ * @param[out] output_ptr Pointer to the destination tensor. Supported data types: F32
+ * @param[in] output_stride_x Stride of the destination tensor in X dimension (in bytes)
+ * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in] output_stride_y Stride of the destination tensor in Y dimension (in bytes)
+ * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes)
+ * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination tensor
+ * @param[in] min The minimum values
+ * @param[in] max The maximum_values
+ * @param[in] aspect_ratios The aspect ratio values
+ * @param[in] min_size The minimum values size
+ * @param[in] max_size The maximum_values values size
+ * @param[in] aspect_ratios_size The aspect ratio values size
+ */
+__kernel void prior_box_layer_nchw(IMAGE_DECLARATION(output), __global float *min, __global float *max, __global float *aspect_ratios, unsigned int min_size, unsigned int max_size,
+ unsigned int aspect_ratios_size)
+{
+ Image out = CONVERT_TO_IMAGE_STRUCT(output);
+
+ int idx = 0;
+ for(unsigned int i = 0; i < min_size; ++i)
+ {
+ idx = calculate_min_nchw(&out, max, aspect_ratios, max_size, aspect_ratios_size, min[i], i, idx);
+ }
+
+ // Store variances
+ for(int i = 0; i < (NUM_PRIORS * 4); i += 4)
+ {
+ vstore4((VEC_DATA_TYPE(DATA_TYPE, 4))(VARIANCE_0, VARIANCE_1, VARIANCE_2, VARIANCE_3), 0, ((__global DATA_TYPE *)offset(&out, i, 1)));
+ }
+}
+
+/** Calculate prior boxes with NHWC format.
+ *
+ * @param[out] output_ptr Pointer to the destination tensor. Supported data types: F32
+ * @param[in] output_stride_x Stride of the destination tensor in X dimension (in bytes)
+ * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in] output_stride_y Stride of the destination tensor in Y dimension (in bytes)
+ * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes)
+ * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination tensor
+ * @param[in] min The minimum values
+ * @param[in] max The maximum_values
+ * @param[in] aspect_ratios The aspect ratio values
+ * @param[in] min_size The minimum values size
+ * @param[in] max_size The maximum_values values size
+ * @param[in] aspect_ratios_size The aspect ratio values size
+ */
+__kernel void prior_box_layer_nhwc(TENSOR3D_DECLARATION(output), __global float *min, __global float *max, __global float *aspect_ratios, unsigned int min_size, unsigned int max_size,
+ unsigned int aspect_ratios_size)
+{
+ Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(output);
+
+ int idx = 0;
+ for(unsigned int i = 0; i < min_size; ++i)
+ {
+ idx = calculate_min_nhwc(&out, max, aspect_ratios, max_size, aspect_ratios_size, min[i], i, idx);
+ }
+
+ for(int i = 0; i < (NUM_PRIORS * 4); i += 4)
+ {
+ *((__global DATA_TYPE *)tensor3D_offset(&out, 0, i + 0, 1)) = VARIANCE_0;
+ *((__global DATA_TYPE *)tensor3D_offset(&out, 0, i + 1, 1)) = VARIANCE_1;
+ *((__global DATA_TYPE *)tensor3D_offset(&out, 0, i + 2, 1)) = VARIANCE_2;
+ *((__global DATA_TYPE *)tensor3D_offset(&out, 0, i + 3, 1)) = VARIANCE_3;
+ }
+}
+#endif /* defined(DATA_TYPE) && defined(WIDTH) && defined(HEIGHT) && defined(LAYER_WIDTH) && defined(LAYER_HEIGHT) && defined(OFFSET) && defined(STEP_X) && defined(STEP_Y) && defined(NUM_PRIORS) && defined(VARIANCE_0) && defined(VARIANCE_1) && defined(VARIANCE_2) && defined(VARIANCE_3) */
diff --git a/src/core/CL/kernels/CLPriorBoxLayerKernel.cpp b/src/core/CL/kernels/CLPriorBoxLayerKernel.cpp
new file mode 100644
index 0000000000..63e745ed10
--- /dev/null
+++ b/src/core/CL/kernels/CLPriorBoxLayerKernel.cpp
@@ -0,0 +1,275 @@
+/*
+ * Copyright (c) 2018 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/CL/kernels/CLPriorBoxLayerKernel.h"
+
+#include "arm_compute/core/CL/CLHelpers.h"
+#include "arm_compute/core/CL/CLKernelLibrary.h"
+#include "arm_compute/core/CL/CLValidate.h"
+#include "arm_compute/core/CL/ICLTensor.h"
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/TensorInfo.h"
+#include "arm_compute/core/Utils.h"
+#include "arm_compute/core/Window.h"
+#include "arm_compute/core/utils/misc/ShapeCalculator.h"
+
+#include "support/ToolchainSupport.h"
+
+using namespace arm_compute::misc::shape_calculator;
+
+namespace arm_compute
+{
+namespace
+{
+Status validate_arguments(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const PriorBoxLayerInfo &info)
+{
+ ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input1, input2, output);
+ ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input1, 1, DataType::F32);
+ ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_LAYOUT(input1, input2);
+ ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input1, input2);
+
+ // Check variances
+ const int var_size = info.variances().size();
+ if(var_size > 1)
+ {
+ ARM_COMPUTE_RETURN_ERROR_ON_MSG(var_size != 4, "Must provide 4 variance values");
+ for(int i = 0; i < var_size; ++i)
+ {
+ ARM_COMPUTE_RETURN_ERROR_ON_MSG(var_size <= 0, "Must be greater than 0");
+ }
+ }
+ ARM_COMPUTE_RETURN_ERROR_ON_MSG(info.steps()[0] < 0.f, "Step x should be greater or equal to 0");
+ ARM_COMPUTE_RETURN_ERROR_ON_MSG(info.steps()[1] < 0.f, "Step y should be greater or equal to 0");
+
+ if(!info.max_sizes().empty())
+ {
+ ARM_COMPUTE_RETURN_ERROR_ON_MSG(info.max_sizes().size() != info.min_sizes().size(), "Max and min sizes dimensions should match");
+ }
+
+ for(unsigned int i = 0; i < info.max_sizes().size(); ++i)
+ {
+ ARM_COMPUTE_RETURN_ERROR_ON_MSG(info.max_sizes()[i] < info.min_sizes()[i], "Max size should be greater than min size");
+ }
+
+ if(output != nullptr && output->total_size() != 0)
+ {
+ ARM_COMPUTE_RETURN_ERROR_ON(output->dimension(get_data_layout_dimension_index(input1->data_layout(), DataLayoutDimension::HEIGHT)) != 2);
+ ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_LAYOUT(input1, output);
+ }
+
+ return Status{};
+}
+
+std::pair<Status, Window> validate_and_configure_window(const ITensorInfo *input1, const ITensorInfo *input2, ITensorInfo *output, const PriorBoxLayerInfo &info, int num_priors)
+{
+ ARM_COMPUTE_UNUSED(input2);
+ // Output tensor auto initialization if not yet initialized
+ TensorShape output_shape = compute_prior_box_shape(*input1, info);
+ auto_init_if_empty(*output, output_shape, 1, input1->data_type());
+
+ Window win{};
+ bool window_changed = false;
+
+ switch(input1->data_layout())
+ {
+ case DataLayout::NCHW:
+ {
+ const unsigned int num_elems_processed_per_iteration = 4 * num_priors;
+
+ win = calculate_max_window(*output, Steps(num_elems_processed_per_iteration));
+ AccessWindowHorizontal output_access(output, 0, num_elems_processed_per_iteration);
+ window_changed = update_window_and_padding(win, output_access);
+ break;
+ }
+ case DataLayout::NHWC:
+ {
+ win = calculate_max_window(*output, Steps());
+ break;
+ }
+ default:
+ ARM_COMPUTE_ERROR("Not implemented");
+ };
+ Status err = (window_changed) ? ARM_COMPUTE_CREATE_ERROR(ErrorCode::RUNTIME_ERROR, "Insufficient Padding!") : Status{};
+ return std::make_pair(err, win);
+}
+} // namespace
+
+CLPriorBoxLayerKernel::CLPriorBoxLayerKernel()
+ : _input1(nullptr), _input2(nullptr), _output(nullptr), _info(), _num_priors(), _min(), _max(), _aspect_ratios()
+{
+}
+
+void CLPriorBoxLayerKernel::configure(const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output, const PriorBoxLayerInfo &info, cl::Buffer *min, cl::Buffer *max, cl::Buffer *aspect_ratios)
+{
+ ARM_COMPUTE_ERROR_ON_NULLPTR(input1, input2, output);
+
+ _input1 = input1;
+ _input2 = input2;
+ _output = output;
+ _info = info;
+ _min = min;
+ _max = max;
+ _aspect_ratios = aspect_ratios;
+
+ ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(input1->info(), input2->info(), output->info(), info));
+
+ // Calculate number of aspect ratios
+ _num_priors = info.aspect_ratios().size() * info.min_sizes().size() + info.max_sizes().size();
+
+ const DataLayout data_layout = input1->info()->data_layout();
+
+ const int width_idx = get_data_layout_dimension_index(data_layout, DataLayoutDimension::WIDTH);
+ const int height_idx = get_data_layout_dimension_index(data_layout, DataLayoutDimension::HEIGHT);
+
+ const int layer_width = input1->info()->dimension(width_idx);
+ const int layer_height = input1->info()->dimension(height_idx);
+
+ int img_width = info.img_size().x;
+ int img_height = info.img_size().y;
+ if(img_width == 0 || img_height == 0)
+ {
+ img_width = input2->info()->dimension(width_idx);
+ img_height = input2->info()->dimension(height_idx);
+ }
+
+ float step_x = info.steps()[0];
+ float step_y = info.steps()[0];
+ if(step_x == 0.f || step_y == 0.f)
+ {
+ step_x = static_cast<float>(img_width) / layer_width;
+ step_y = static_cast<float>(img_height) / layer_height;
+ }
+
+ // Set build options
+ CLBuildOptions build_opts;
+ build_opts.add_option("-DDATA_TYPE=" + get_cl_type_from_data_type(input1->info()->data_type()));
+ build_opts.add_option("-DWIDTH=" + support::cpp11::to_string(img_width));
+ build_opts.add_option("-DHEIGHT=" + support::cpp11::to_string(img_height));
+ build_opts.add_option("-DLAYER_WIDTH=" + support::cpp11::to_string(layer_width));
+ build_opts.add_option("-DLAYER_HEIGHT=" + support::cpp11::to_string(layer_height));
+ build_opts.add_option("-DSTEP_X=" + support::cpp11::to_string(step_x));
+ build_opts.add_option("-DSTEP_Y=" + support::cpp11::to_string(step_y));
+ build_opts.add_option("-DNUM_PRIORS=" + support::cpp11::to_string(_num_priors));
+ build_opts.add_option("-DOFFSET=" + support::cpp11::to_string(info.offset()));
+ build_opts.add_option_if(info.clip(), "-DIN_PLACE");
+
+ if(info.variances().size() > 1)
+ {
+ for(unsigned int i = 0; i < info.variances().size(); ++i)
+ {
+ build_opts.add_option("-DVARIANCE_" + support::cpp11::to_string(i) + "=" + support::cpp11::to_string(info.variances().at(i)));
+ }
+ }
+ else
+ {
+ for(unsigned int i = 0; i < 4; ++i)
+ {
+ build_opts.add_option("-DVARIANCE_" + support::cpp11::to_string(i) + "=" + support::cpp11::to_string(info.variances().at(0)));
+ }
+ }
+
+ unsigned int idx = 0;
+ // Create kernel
+ switch(data_layout)
+ {
+ case DataLayout::NCHW:
+ {
+ idx = num_arguments_per_2D_tensor();
+ _kernel = static_cast<cl::Kernel>(CLKernelLibrary::get().create_kernel("prior_box_layer_nchw", build_opts.options()));
+ break;
+ }
+ case DataLayout::NHWC:
+ {
+ idx = num_arguments_per_3D_tensor();
+ _kernel = static_cast<cl::Kernel>(CLKernelLibrary::get().create_kernel("prior_box_layer_nhwc", build_opts.options()));
+ break;
+ }
+ default:
+ ARM_COMPUTE_ERROR("Not implemented");
+ }
+
+ _kernel.setArg(idx++, *_min);
+ _kernel.setArg(idx++, *_max);
+ _kernel.setArg(idx++, *_aspect_ratios);
+ _kernel.setArg<unsigned int>(idx++, info.min_sizes().size());
+ _kernel.setArg<unsigned int>(idx++, info.max_sizes().size());
+ _kernel.setArg<unsigned int>(idx++, info.aspect_ratios().size());
+
+ // Configure kernel window
+ auto win_config = validate_and_configure_window(input1->info(), input2->info(), output->info(), info, _num_priors);
+
+ ARM_COMPUTE_ERROR_THROW_ON(win_config.first);
+ ICLKernel::configure_internal(win_config.second);
+}
+
+Status CLPriorBoxLayerKernel::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const PriorBoxLayerInfo &info)
+{
+ ARM_COMPUTE_ERROR_ON_NULLPTR(input1, input2, output);
+ ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(input1, input2, output, info));
+ const int num_priors = info.aspect_ratios().size() * info.min_sizes().size() + info.max_sizes().size();
+ ARM_COMPUTE_RETURN_ON_ERROR(validate_and_configure_window(input1->clone().get(), input2->clone().get(), output->clone().get(), info, num_priors)
+ .first);
+
+ return Status{};
+}
+
+void CLPriorBoxLayerKernel::run(const Window &window, cl::CommandQueue &queue)
+{
+ ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this);
+ ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(IKernel::window(), window);
+
+ queue.enqueueWriteBuffer(*_min, CL_TRUE, 0, _info.min_sizes().size() * sizeof(float), _info.min_sizes().data());
+ queue.enqueueWriteBuffer(*_aspect_ratios, CL_TRUE, 0, _info.aspect_ratios().size() * sizeof(float), _info.aspect_ratios().data());
+ if(!_info.max_sizes().empty())
+ {
+ queue.enqueueWriteBuffer(*_max, CL_TRUE, 0, _info.max_sizes().size() * sizeof(float), _info.max_sizes().data());
+ }
+
+ switch(_input1->info()->data_layout())
+ {
+ case DataLayout::NCHW:
+ {
+ Window slice = window.first_slice_window_2D();
+ slice.set(Window::DimY, Window::Dimension(0, _output->info()->dimension(1), 2));
+
+ unsigned int idx = 0;
+ add_2D_tensor_argument(idx, _output, slice);
+ enqueue(queue, *this, slice);
+ break;
+ }
+ case DataLayout::NHWC:
+ {
+ Window slice = window.first_slice_window_3D();
+ slice.set(Window::DimY, Window::Dimension(0, _output->info()->dimension(1), 4 * _num_priors));
+ slice.set(Window::DimZ, Window::Dimension(0, _output->info()->dimension(2), 2));
+
+ unsigned int idx = 0;
+ add_3D_tensor_argument(idx, _output, slice);
+ enqueue(queue, *this, slice);
+ break;
+ }
+ default:
+ ARM_COMPUTE_ERROR("Not implemented");
+ }
+}
+} // namespace arm_compute
diff --git a/src/runtime/CL/functions/CLPriorBoxLayer.cpp b/src/runtime/CL/functions/CLPriorBoxLayer.cpp
new file mode 100644
index 0000000000..4f6c969a92
--- /dev/null
+++ b/src/runtime/CL/functions/CLPriorBoxLayer.cpp
@@ -0,0 +1,58 @@
+/*
+ * Copyright (c) 2018 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/CLPriorBoxLayer.h"
+
+#include "arm_compute/core/CL/kernels/CLPriorBoxLayerKernel.h"
+#include "arm_compute/core/Error.h"
+#include "arm_compute/core/TensorInfo.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/core/Validate.h"
+#include "arm_compute/runtime/CL/CLScheduler.h"
+
+using namespace arm_compute;
+
+CLPriorBoxLayer::CLPriorBoxLayer()
+ : _min(nullptr), _max(nullptr), _aspect_ratios(nullptr)
+{
+}
+
+void CLPriorBoxLayer::configure(const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output, const PriorBoxLayerInfo &info)
+{
+ _min = cl::Buffer(CLScheduler::get().context(), CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_WRITE, info.min_sizes().size() * sizeof(float));
+ _aspect_ratios = cl::Buffer(CLScheduler::get().context(), CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_WRITE, info.aspect_ratios().size() * sizeof(float));
+ if(!info.max_sizes().empty())
+ {
+ _max = cl::Buffer(CLScheduler::get().context(), CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_WRITE, info.max_sizes().size() * sizeof(float));
+ }
+
+ auto k = arm_compute::support::cpp14::make_unique<CLPriorBoxLayerKernel>();
+ k->configure(input1, input2, output, info, &_min, &_max, &_aspect_ratios);
+ _kernel = std::move(k);
+}
+
+Status CLPriorBoxLayer::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const PriorBoxLayerInfo &info)
+{
+ return CLPriorBoxLayerKernel::validate(input1, input2, output, info);
+} \ No newline at end of file
diff --git a/tests/datasets/PriorBoxLayerDataset.h b/tests/datasets/PriorBoxLayerDataset.h
new file mode 100644
index 0000000000..22a0d97057
--- /dev/null
+++ b/tests/datasets/PriorBoxLayerDataset.h
@@ -0,0 +1,134 @@
+/*
+ * Copyright (c) 2018 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_PRIORBOX_LAYER_DATASET
+#define ARM_COMPUTE_TEST_PRIORBOX_LAYER_DATASET
+
+#include "utils/TypePrinter.h"
+
+#include "arm_compute/core/TensorShape.h"
+#include "arm_compute/core/Types.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace datasets
+{
+class PriorBoxLayerDataset
+{
+public:
+ using type = std::tuple<TensorShape, PriorBoxLayerInfo>;
+
+ struct iterator
+ {
+ iterator(std::vector<TensorShape>::const_iterator src_it,
+ std::vector<PriorBoxLayerInfo>::const_iterator infos_it)
+ : _src_it{ std::move(src_it) },
+ _infos_it{ std::move(infos_it) }
+ {
+ }
+
+ std::string description() const
+ {
+ std::stringstream description;
+ description << "In=" << *_src_it << ":";
+ description << "Info=" << *_infos_it << ":";
+ return description.str();
+ }
+
+ PriorBoxLayerDataset::type operator*() const
+ {
+ return std::make_tuple(*_src_it, *_infos_it);
+ }
+
+ iterator &operator++()
+ {
+ ++_src_it;
+ ++_infos_it;
+
+ return *this;
+ }
+
+ private:
+ std::vector<TensorShape>::const_iterator _src_it;
+ std::vector<PriorBoxLayerInfo>::const_iterator _infos_it;
+ };
+
+ iterator begin() const
+ {
+ return iterator(_src_shapes.begin(), _infos.begin());
+ }
+
+ int size() const
+ {
+ return std::min(_src_shapes.size(), _infos.size());
+ }
+
+ void add_config(TensorShape src, PriorBoxLayerInfo info)
+ {
+ _src_shapes.emplace_back(std::move(src));
+ _infos.emplace_back(std::move(info));
+ }
+
+protected:
+ PriorBoxLayerDataset() = default;
+ PriorBoxLayerDataset(PriorBoxLayerDataset &&) = default;
+
+private:
+ std::vector<TensorShape> _src_shapes{};
+ std::vector<PriorBoxLayerInfo> _infos{};
+};
+
+class SmallPriorBoxLayerDataset final : public PriorBoxLayerDataset
+{
+public:
+ SmallPriorBoxLayerDataset()
+ {
+ std::vector<float> min_val = { 30.f };
+ std::vector<float> var = { 0.1, 0.1, 0.2, 0.2 };
+ std::vector<float> max_val = { 60.f };
+ std::vector<float> aspect_ratio = { 2.f };
+ std::array<float, 2> steps = { 8.f, 8.f };
+
+ add_config(TensorShape(4U, 4U), PriorBoxLayerInfo(min_val, var, 0.5f, true, false, max_val, aspect_ratio, Coordinates2D{ 8, 8 }, steps));
+ }
+};
+
+class LargePriorBoxLayerDataset final : public PriorBoxLayerDataset
+{
+public:
+ LargePriorBoxLayerDataset()
+ {
+ std::vector<float> min_val = { 30.f };
+ std::vector<float> var = { 0.1, 0.1, 0.2, 0.2 };
+ std::vector<float> max_val = { 60.f };
+ std::vector<float> aspect_ratio = { 2.f };
+ std::array<float, 2> steps = { 8.f, 8.f };
+ add_config(TensorShape(150U, 245U, 4U, 12U), PriorBoxLayerInfo(min_val, var, 0.5f, true, false, max_val, aspect_ratio, Coordinates2D{ 8, 8 }, steps));
+ }
+};
+} // namespace datasets
+} // namespace test
+} // namespace arm_compute
+#endif /* ARM_COMPUTE_TEST_PRIORBOX_LAYER_DATASET */
diff --git a/tests/validation/CL/PriorBoxLayer.cpp b/tests/validation/CL/PriorBoxLayer.cpp
new file mode 100644
index 0000000000..79776b5bcd
--- /dev/null
+++ b/tests/validation/CL/PriorBoxLayer.cpp
@@ -0,0 +1,103 @@
+/*
+ * Copyright (c) 2018 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/CLPriorBoxLayer.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/PriorBoxLayerDataset.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/Helpers.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/PriorBoxLayerFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+constexpr AbsoluteTolerance<float> tolerance_f32(0.00001f); /**< Tolerance value for comparing reference's output against implementation's output for DataType::F32 */
+} // namespace
+
+TEST_SUITE(CL)
+TEST_SUITE(PriorBoxLayer)
+
+template <typename T>
+using CLPriorBoxLayerFixture = PriorBoxLayerValidationFixture<CLTensor, CLAccessor, CLPriorBoxLayer, T>;
+
+// *INDENT-OFF*
+// clang-format off
+DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip(zip(
+ framework::dataset::make("Input1Info", { TensorInfo(TensorShape(10U, 10U, 2U), 1, DataType::F32),
+ TensorInfo(TensorShape(10U, 10U, 2U), 1, DataType::F32), // Window shrink
+ }),
+ framework::dataset::make("Input2Info", { TensorInfo(TensorShape(10U, 10U, 2U), 1, DataType::F32),
+ TensorInfo(TensorShape(10U, 10U, 2U), 1, DataType::F32),
+ })),
+ framework::dataset::make("OutputInfo",{ TensorInfo(TensorShape(1200U, 2U), 1, DataType::F32),
+ TensorInfo(TensorShape(1000U, 2U), 1, DataType::F32),
+ })),
+ framework::dataset::make("PriorBoxInfo",{ PriorBoxLayerInfo(std::vector<float>(1), std::vector<float>(1), 0, true, true, std::vector<float>(1), std::vector<float>(1), Coordinates2D{8, 8}, std::array<float, 2>()),
+ PriorBoxLayerInfo(std::vector<float>(1), std::vector<float>(1), 0, true, true, std::vector<float>(1), std::vector<float>(1), Coordinates2D{8, 8}, std::array<float, 2>()),
+ })),
+ framework::dataset::make("Expected", { true, false})),
+ input1_info, input2_info, output_info, info, expected)
+{
+ bool has_error = bool(CLPriorBoxLayer::validate(&input1_info.clone()->set_is_resizable(false), &input2_info.clone()->set_is_resizable(false), &output_info.clone()->set_is_resizable(false), info));
+ ARM_COMPUTE_EXPECT(has_error == expected, framework::LogLevel::ERRORS);
+}
+// clang-format on
+// *INDENT-ON*
+
+TEST_SUITE(Float)
+TEST_SUITE(FP32)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLPriorBoxLayerFixture<float>, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallPriorBoxLayerDataset(),
+ framework::dataset::make("DataType", DataType::F32)),
+ framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC })))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f32, 0);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLPriorBoxLayerFixture<float>, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargePriorBoxLayerDataset(),
+ framework::dataset::make("DataType", DataType::F32)),
+ framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC })))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f32, 0);
+}
+TEST_SUITE_END() // Float
+TEST_SUITE_END() // FP32
+
+TEST_SUITE_END() // PriorBoxLayer
+TEST_SUITE_END() // CL
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/fixtures/PriorBoxLayerFixture.h b/tests/validation/fixtures/PriorBoxLayerFixture.h
new file mode 100644
index 0000000000..dd7a49ee1f
--- /dev/null
+++ b/tests/validation/fixtures/PriorBoxLayerFixture.h
@@ -0,0 +1,112 @@
+/*
+ * Copyright (c) 2018 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_PRIOR_BOX_LAYER_FIXTURE
+#define ARM_COMPUTE_TEST_PRIOR_BOX_LAYER_FIXTURE
+
+#include "arm_compute/core/TensorShape.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/core/utils/misc/ShapeCalculator.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/PriorBoxLayer.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+template <typename TensorType, typename AccessorType, typename FunctionType, typename T>
+class PriorBoxLayerValidationFixture : public framework::Fixture
+{
+public:
+ template <typename...>
+ void setup(TensorShape input_shape, PriorBoxLayerInfo info, DataType data_type, DataLayout data_layout)
+ {
+ TensorInfo input_info(input_shape, 1, data_type);
+ const TensorShape output_shape = misc::shape_calculator::compute_prior_box_shape(input_info, info);
+ _data_type = data_type;
+
+ _target = compute_target(input_shape, output_shape, info, data_type, data_layout);
+ _reference = compute_reference(input_shape, output_shape, info, data_type);
+ }
+
+protected:
+ TensorType compute_target(TensorShape input_shape, TensorShape output_shape, PriorBoxLayerInfo info, DataType data_type, DataLayout data_layout)
+ {
+ if(data_layout == DataLayout::NHWC)
+ {
+ permute(input_shape, PermutationVector(2U, 0U, 1U));
+ permute(output_shape, PermutationVector(2U, 0U, 1U));
+ }
+
+ // Create tensors
+ TensorType src1 = create_tensor<TensorType>(input_shape, data_type, 1, QuantizationInfo(), data_layout);
+ TensorType src2 = create_tensor<TensorType>(input_shape, data_type, 1, QuantizationInfo(), data_layout);
+ TensorType dst = create_tensor<TensorType>(output_shape, data_type, 1, QuantizationInfo(), data_layout);
+
+ // Create and configure function
+ FunctionType prior_box;
+ prior_box.configure(&src1, &src2, &dst, info);
+
+ ARM_COMPUTE_EXPECT(src1.info()->is_resizable(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(src2.info()->is_resizable(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+ // Allocate tensors
+ src1.allocator()->allocate();
+ src2.allocator()->allocate();
+ dst.allocator()->allocate();
+
+ ARM_COMPUTE_EXPECT(!src1.info()->is_resizable(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(!src2.info()->is_resizable(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(!dst.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+ // Compute function
+ prior_box.run();
+
+ return dst;
+ }
+
+ SimpleTensor<T> compute_reference(const TensorShape &input_shape, const TensorShape &output_shape, PriorBoxLayerInfo info, DataType data_type)
+ {
+ // Create reference
+ SimpleTensor<T> input1{ input_shape, data_type, 1 };
+ SimpleTensor<T> input2{ input_shape, data_type, 1 };
+
+ return reference::prior_box_layer(input1, input2, info, output_shape);
+ }
+
+ TensorType _target{};
+ SimpleTensor<T> _reference{};
+ DataType _data_type{};
+};
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* ARM_COMPUTE_TEST_PRIOR_BOX_LAYER_FIXTURE */
diff --git a/tests/validation/reference/PriorBoxLayer.cpp b/tests/validation/reference/PriorBoxLayer.cpp
new file mode 100644
index 0000000000..0fd4a8aa48
--- /dev/null
+++ b/tests/validation/reference/PriorBoxLayer.cpp
@@ -0,0 +1,158 @@
+/*
+ * Copyright (c) 2018 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 "PriorBoxLayer.h"
+
+#include "ActivationLayer.h"
+
+#include "tests/validation/Helpers.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace reference
+{
+template <typename T>
+SimpleTensor<T> prior_box_layer(const SimpleTensor<T> &src1, const SimpleTensor<T> &src2, const PriorBoxLayerInfo &info, const TensorShape &output_shape)
+{
+ const auto layer_width = static_cast<int>(src1.shape()[0]);
+ const auto layer_height = static_cast<int>(src1.shape()[1]);
+
+ int img_width = info.img_size().x;
+ int img_height = info.img_size().y;
+ if(img_width == 0 || img_height == 0)
+ {
+ img_width = static_cast<int>(src2.shape()[0]);
+ img_height = static_cast<int>(src2.shape()[1]);
+ }
+
+ float step_x = info.steps()[0];
+ float step_y = info.steps()[1];
+ if(step_x == 0.f || step_y == 0.f)
+ {
+ step_x = static_cast<float>(img_width) / layer_width;
+ step_x = static_cast<float>(img_height) / layer_height;
+ }
+
+ // Calculate number of aspect ratios
+ const int num_priors = info.aspect_ratios().size() * info.min_sizes().size() + info.max_sizes().size();
+ const int total_elements = layer_width * layer_height * num_priors * 4;
+
+ SimpleTensor<T> result(output_shape, src1.data_type());
+
+ int idx = 0;
+ for(int y = 0; y < layer_height; ++y)
+ {
+ for(int x = 0; x < layer_width; ++x)
+ {
+ const float center_x = (x + info.offset()) * step_x;
+ const float center_y = (y + info.offset()) * step_y;
+ float box_width;
+ float box_height;
+ for(unsigned int i = 0; i < info.min_sizes().size(); ++i)
+ {
+ const float min_size = info.min_sizes().at(i);
+ box_width = min_size;
+ box_height = min_size;
+ // (xmin, ymin, xmax, ymax)
+ result[idx++] = (center_x - box_width / 2.f) / img_width;
+ result[idx++] = (center_y - box_height / 2.f) / img_height;
+ result[idx++] = (center_x + box_width / 2.f) / img_width;
+ result[idx++] = (center_y + box_height / 2.f) / img_height;
+
+ if(!info.max_sizes().empty())
+ {
+ const float max_size = info.max_sizes().at(i);
+ box_width = sqrt(min_size * max_size);
+ box_height = box_width;
+
+ // (xmin, ymin, xmax, ymax)
+ result[idx++] = (center_x - box_width / 2.f) / img_width;
+ result[idx++] = (center_y - box_height / 2.f) / img_height;
+ result[idx++] = (center_x + box_width / 2.f) / img_width;
+ result[idx++] = (center_y + box_height / 2.f) / img_height;
+ }
+
+ // rest of priors
+ for(auto ar : info.aspect_ratios())
+ {
+ if(fabs(ar - 1.) < 1e-6)
+ {
+ continue;
+ }
+
+ box_width = min_size * sqrt(ar);
+ box_height = min_size / sqrt(ar);
+
+ // (xmin, ymin, xmax, ymax)
+ result[idx++] = (center_x - box_width / 2.f) / img_width;
+ result[idx++] = (center_y - box_height / 2.f) / img_height;
+ result[idx++] = (center_x + box_width / 2.f) / img_width;
+ result[idx++] = (center_y + box_height / 2.f) / img_height;
+ }
+ }
+ }
+ }
+
+ // clip the coordinates
+ if(info.clip())
+ {
+ for(int i = 0; i < total_elements; ++i)
+ {
+ result[i] = std::min<T>(std::max<T>(result[i], 0.f), 1.f);
+ }
+ }
+
+ // set the variance.
+ if(info.variances().size() == 1)
+ {
+ std::fill_n(result.data() + idx, total_elements, info.variances().at(0));
+ }
+ else
+ {
+ for(int h = 0; h < layer_height; ++h)
+ {
+ for(int w = 0; w < layer_width; ++w)
+ {
+ for(int i = 0; i < num_priors; ++i)
+ {
+ for(int j = 0; j < 4; ++j)
+ {
+ result[idx++] = info.variances().at(j);
+ }
+ }
+ }
+ }
+ }
+
+ return result;
+}
+template SimpleTensor<float> prior_box_layer(const SimpleTensor<float> &src1, const SimpleTensor<float> &src2, const PriorBoxLayerInfo &info, const TensorShape &output_shape);
+
+} // namespace reference
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/reference/PriorBoxLayer.h b/tests/validation/reference/PriorBoxLayer.h
new file mode 100644
index 0000000000..25e567f59a
--- /dev/null
+++ b/tests/validation/reference/PriorBoxLayer.h
@@ -0,0 +1,44 @@
+/*
+ * Copyright (c) 2018 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_PRIOR_BOX_LAYER_H__
+#define __ARM_COMPUTE_TEST_PRIOR_BOX_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> prior_box_layer(const SimpleTensor<T> &src1, const SimpleTensor<T> &src2, const PriorBoxLayerInfo &info, const TensorShape &output_shape);
+} // namespace reference
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* __ARM_COMPUTE_TEST_PRIOR_BOX_LAYER_H__ */
diff --git a/utils/TypePrinter.h b/utils/TypePrinter.h
index 571de9a4ef..df16cba9b5 100644
--- a/utils/TypePrinter.h
+++ b/utils/TypePrinter.h
@@ -1324,6 +1324,30 @@ inline std::string to_string(const PoolingLayerInfo &info)
return str.str();
}
+/** Formatted output of the PriorBoxLayerInfo.
+ *
+ * @param[in] info Type to output.
+ *
+ * @return Formatted string.
+ */
+inline std::string to_string(const PriorBoxLayerInfo &info)
+{
+ std::stringstream str;
+ str << "{";
+ str << "Clip:" << info.clip()
+ << "Flip:" << info.flip()
+ << "StepX:" << info.steps()[0]
+ << "StepY:" << info.steps()[1]
+ << "MinSizes:" << info.min_sizes().size()
+ << "MaxSizes:" << info.max_sizes().size()
+ << "ImgSizeX:" << info.img_size().x
+ << "ImgSizeY:" << info.img_size().y
+ << "Offset:" << info.offset()
+ << "Variances:" << info.variances().size();
+ str << "}";
+ return str.str();
+}
+
/** Formatted output of the KeyPoint type.
*
* @param[out] os Output stream
@@ -1780,6 +1804,29 @@ inline ::std::ostream &operator<<(::std::ostream &os, const std::vector<T> &args
return os;
}
+/** Formatted output of @ref PriorBoxLayerInfo.
+ *
+ * @param[out] os Output stream.
+ * @param[in] info Type to output.
+ *
+ * @return Modified output stream.
+ */
+inline ::std::ostream &operator<<(::std::ostream &os, const PriorBoxLayerInfo &info)
+{
+ os << "Clip:" << info.clip()
+ << "Flip:" << info.flip()
+ << "StepX:" << info.steps()[0]
+ << "StepY:" << info.steps()[1]
+ << "MinSizes:" << info.min_sizes()
+ << "MaxSizes:" << info.max_sizes()
+ << "ImgSizeX:" << info.img_size().x
+ << "ImgSizeY:" << info.img_size().y
+ << "Offset:" << info.offset()
+ << "Variances:" << info.variances();
+
+ return os;
+}
+
/** Formatted output of a vector of objects.
*
* @param[in] args Vector of objects to print