COMPMID-1462 SSD support: Create CL PriorBox

Change-Id: I5bf5d751ec7c02d96c26a769f49d03ea23a248b7

4 files changed, 595 insertions, 0 deletions

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