COMPMID-2080 Create a new kernel for CLPad with CONSTANT mode

Change-Id: Ib67bacd40c13c7784ca9f5699a235f9d106baddb
Signed-off-by: Giorgio Arena <giorgio.arena@arm.com>
Reviewed-on: https://review.mlplatform.org/c/1739
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Giuseppe Rossini <giuseppe.rossini@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>

4 files changed, 285 insertions, 35 deletions

diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp
index 2b843e7785..b938a1813a 100644
--- a/src/core/CL/CLKernelLibrary.cpp
+++ b/src/core/CL/CLKernelLibrary.cpp
@@ -407,6 +407,7 @@ const std::map<std::string, std::string> CLKernelLibrary::_kernel_program_map =
     { "NV21_to_RGBA8888_bt709", "color_convert.cl" },
     { "NV21_to_YUV444_bt709", "color_convert.cl" },
     { "output_stage_quantized", "direct_convolution_1x1_3x3_5x5_quantized.cl" },
+    { "pad_layer", "pad_layer.cl" },
     { "permute", "permute.cl" },
     { "pixelwise_mul_complex", "pixelwise_mul_float.cl" },
     { "pixelwise_mul_float", "pixelwise_mul_float.cl" },
@@ -874,6 +875,10 @@ const std::map<std::string, std::string> CLKernelLibrary::_program_source_map =
 #include "./cl_kernels/optical_flow_pyramid_lk.clembed"
     },
     {
+        "pad_layer.cl",
+#include "./cl_kernels/pad_layer.clembed"
+    },
+    {
         "permute.cl",
 #include "./cl_kernels/permute.clembed"
     },
diff --git a/src/core/CL/cl_kernels/pad_layer.cl b/src/core/CL/cl_kernels/pad_layer.cl
new file mode 100644
index 0000000000..ace2f0d3a0
--- /dev/null
+++ b/src/core/CL/cl_kernels/pad_layer.cl
@@ -0,0 +1,132 @@
+/*
+ * Copyright (c) 2019 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "helpers.h"
+
+#if defined(DATA_TYPE) && defined(VEC_SIZE) && defined(CONST_VAL)
+
+#define VEC_TYPE VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
+#define VEC_INT VEC_DATA_TYPE(int, VEC_SIZE)
+#define CONVERT_SELECT(x) CONVERT(x, VEC_DATA_TYPE(SELECT_DT, VEC_SIZE))
+
+#if VEC_SIZE == 1
+#define OFFSETS (int)0
+#elif VEC_SIZE == 2
+#define OFFSETS (int2)(0, 1)
+#elif VEC_SIZE == 4
+#define OFFSETS (int4)(0, 1, 2, 3)
+#elif VEC_SIZE == 8
+#define OFFSETS (int8)(0, 1, 2, 3, 4, 5, 6, 7)
+#elif VEC_SIZE == 16
+#define OFFSETS (int16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)
+#else // VEC_SIZE
+#error "Only 1, 2, 3, 4, 8 and 16 vector sizes allowed"
+#endif // VEC_SIZE
+
+/** Perform a pad operation
+ *
+ * @note Data type can be passed using the -DDATA_TYPE compile flag, e.g. -DDATA_TYPE=float
+ * @note Vector size must be passed using the -DVEC_SIZE compile flag, e.g. -DVEC_SIZE=4
+ * @note Constant value must be passed using the -DCONST_VAL compile flag, e.g. -DCONST_VAL=1.27
+ * @note Pad to add to the left must be passed using the -DPAD_LEFT compile flag, e.g. -DPAD_LEFT=5
+ * @note Input tensor's width must be passed using the -DSRC_WIDTH compile flag, e.g. -DSRC_WIDTH=224
+ * @note Data type to use for the select instruction must be passed using the -DSELECT_DT compile flag, e.g. -DSELECT_DT=float
+ * @note In case pad left is more than the vector size, the number of threads to skil alond the X axis must be passed using the
+ *       -DTHREADS_TO_SKIP_X compile flag, e.g. -DTHREADS_TO_SKIP_X=1. This is defined as (PAD_LEFT / VEC_SIZE)
+ * @note In pad also needs to be added to the top of the tensor, the following compile flags must be passed at compile time:
+ *       -# -DPAD_TOP: Pad to add to the top of the input tensor (e.g. -DPAD_TOP=3)
+ *       -# -DSRC_HEIGHT: Input tensor's height (e.g. -DSRC_HEIGHT=127)
+ * @note In pad also needs to be added to the depth of the tensor, the following compile flags must be passed at compile time:
+ *       -# -DPAD_NEAR: Pad to add before the first plane of the input tensor (e.g. -DPAD_NEAR=3)
+ *       -# -DSRC_DEPTH: Input tensor's depth (e.g. -DSRC_DEPTH=32)
+ *
+ * @param[in]  src_ptr                           Pointer to the source image. Supported data types: U8, S8, QASYMM8, U16, S16, U32, S32, F16, F32
+ * @param[in]  src_stride_x                      Stride of the source image in X dimension (in bytes)
+ * @param[in]  src_step_x                        src_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in]  src_stride_y                      Stride of the source image in Y dimension (in bytes)
+ * @param[in]  src_step_y                        src_stride_y * number of elements along Y processed per workitem(in bytes)
+ * @param[in]  src_stride_z                      Stride of the source image in Z dimension (in bytes)
+ * @param[in]  src_step_z                        src_stride_z * number of elements along Z processed per workitem(in bytes)
+ * @param[in]  src_offset_first_element_in_bytes The offset of the first element in the source image
+ * @param[out] dst_ptr                           Pointer to the destination image. Supported data types: same as @p src_ptr
+ * @param[in]  dst_stride_x                      Stride of the destination image in X dimension (in bytes)
+ * @param[in]  dst_step_x                        dst_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in]  dst_stride_y                      Stride of the destination image in Y dimension (in bytes)
+ * @param[in]  dst_step_y                        dst_stride_y * number of elements along Y processed per workitem(in bytes)
+ * @param[in]  dst_stride_z                      Stride of the destination image in Z dimension (in bytes)
+ * @param[in]  dst_step_z                        dst_stride_z * number of elements along Z processed per workitem(in bytes)
+ * @param[in]  dst_offset_first_element_in_bytes The offset of the first element in the destination image
+ */
+__kernel void pad_layer(TENSOR3D_DECLARATION(src),
+                        TENSOR3D_DECLARATION(dst))
+{
+    const int x = get_global_id(0);
+    const int y = get_global_id(1);
+    const int z = get_global_id(2);
+
+#if defined(PAD_NEAR)
+    if(z < PAD_NEAR || z >= (PAD_NEAR + SRC_DEPTH))
+    {
+        Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst);
+        VSTORE(VEC_SIZE)
+        ((VEC_TYPE)CONST_VAL, 0, (__global DATA_TYPE *)dst.ptr);
+    }
+    else
+    {
+#endif // defined(PAD_NEAR)
+
+        Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src);
+        Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst);
+
+#if defined(THREADS_TO_SKIP_X)
+        /* In case the pad left is greater than the vector size, and we are past the threads operating solely on pad values,
+         * the input pointer must be brought back along the X axis to start from the first non-pad values.
+         *
+         * E.g. with VEC_SIZE=2, PAD_LEFT=5, CONST_VAL=0 and 1D input |1 2 3 4 5 6|:
+         *  -# The first thread will compute the output values |0 0| since it detects (x_outs == (0, 1)) < PAD_LEFT
+         *  -# The second thread will compute the output values |0 0| since it detects (x_outs == (2, 3)) < PAD_LEFT
+         *  -# The third thread should compute |0 1|, however the input pointer is now ahead of ((x * VEC_SIZE) == 4) values, reading |4 5|
+         *  -# To detect this, we use ((PAD_LEFT / VEC_SIZE) == THREADS_TO_SKIP_X == 2) and check that it is >= to the current x
+         *  -# So, we bring the pointer back of THREADS_TO_SKIP_X threads, which means multiplying this constant by the input's step along the X axis
+         *  -# Now that the pointer is back of ((THREADS_TO_SKIP_X * src_step_x) == 4) values, it will read the desired values |0 1|
+         */
+        src.ptr -= select(0u, THREADS_TO_SKIP_X * src_step_x, x >= THREADS_TO_SKIP_X);
+#endif // defined(THREADS_TO_SKIP_X)
+#if defined(PAD_NEAR)
+        src.ptr -= PAD_NEAR * src_step_z;
+#endif // defined(PAD_NEAR)
+
+        VEC_TYPE src_vals = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)src.ptr);
+
+        VEC_INT xs_out = (VEC_INT)(x * VEC_SIZE) + OFFSETS;
+        VEC_INT cond   = xs_out < (VEC_INT)PAD_LEFT || xs_out >= (VEC_INT)(PAD_LEFT + SRC_WIDTH);
+#if defined(PAD_TOP)
+        cond |= (VEC_INT)y < (VEC_INT)PAD_TOP || (VEC_INT)y >= (VEC_INT)(PAD_TOP + SRC_HEIGHT);
+#endif // defined(PAD_TOP)
+        VSTORE(VEC_SIZE)
+        (select(src_vals, (VEC_TYPE)CONST_VAL, CONVERT_SELECT(cond)), 0, (__global DATA_TYPE *)dst.ptr);
+#if defined(PAD_NEAR)
+    }
+#endif // defined(PAD_NEAR)
+}
+#endif // defined(DATA_TYPE) && defined(VEC_SIZE) && defined(CONST_VAL)
diff --git a/src/core/CL/kernels/CLPadLayerKernel.cpp b/src/core/CL/kernels/CLPadLayerKernel.cpp
new file mode 100644
index 0000000000..d35164d309
--- /dev/null
+++ b/src/core/CL/kernels/CLPadLayerKernel.cpp
@@ -0,0 +1,143 @@
+/*
+ * Copyright (c) 2019 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/CL/kernels/CLPadLayerKernel.h"
+
+#include "arm_compute/core/CL/CLHelpers.h"
+
+namespace arm_compute
+{
+namespace
+{
+Status validate_arguments(const ITensorInfo *input, const ITensorInfo *output, const PaddingList &padding, PixelValue constant_value, PaddingMode mode)
+{
+    ARM_COMPUTE_UNUSED(input, output, constant_value);
+    ARM_COMPUTE_RETURN_ERROR_ON(padding.empty());
+    ARM_COMPUTE_RETURN_ERROR_ON_MSG(mode != PaddingMode::CONSTANT, "Only CONSTANT mode supported.");
+
+    return Status{};
+}
+
+std::pair<Status, Window> validate_and_configure_window(ITensorInfo *input, ITensorInfo *output, const PaddingList &padding, PixelValue constant_value, PaddingMode mode)
+{
+    ARM_COMPUTE_UNUSED(constant_value, mode);
+    const unsigned int num_elems_processed_per_iteration = std::min(16U, 32U / static_cast<unsigned int>(element_size_from_data_type(input->data_type())));
+
+    // Configure kernel window
+    Window win = calculate_max_window(*output, Steps(num_elems_processed_per_iteration));
+
+    const int input_start_x = -(padding.at(0).first % num_elems_processed_per_iteration);
+    const int input_start_y = padding.size() > 1 ? -padding.at(1).first : 0;
+
+    AccessWindowRectangle  input_access(input, input_start_x, input_start_y, num_elems_processed_per_iteration, 1);
+    AccessWindowHorizontal output_access(output, 0, num_elems_processed_per_iteration);
+
+    const bool window_changed = update_window_and_padding(win, input_access, output_access);
+    output_access.set_valid_region(win, ValidRegion(Coordinates(), output->tensor_shape()));
+
+    Status err = (window_changed) ? ARM_COMPUTE_CREATE_ERROR(ErrorCode::RUNTIME_ERROR, "Insufficient Padding!") : Status{};
+    return std::make_pair(err, win);
+}
+} // namespace
+
+CLPadLayerKernel::CLPadLayerKernel()
+    : _input(nullptr), _output(nullptr), _input_start_x(0), _input_start_y(0)
+{
+}
+
+void CLPadLayerKernel::configure(const ICLTensor *input, ICLTensor *output, const PaddingList &padding, PixelValue constant_value, PaddingMode mode)
+{
+    ARM_COMPUTE_ERROR_ON_NULLPTR(input, output);
+
+    // Output tensor auto initialisation if not yet initialized
+    auto_init_if_empty(*output->info(), *input->info()->clone());
+    ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(input->info(), output->info(), padding, constant_value, mode));
+
+    _input  = input;
+    _output = output;
+
+    // Set build options
+    const unsigned int num_elems_processed_per_iteration = std::min(16U, 32U / static_cast<unsigned int>(element_size_from_data_type(input->info()->data_type())));
+    _input_start_x                                       = -(padding.at(0).first % num_elems_processed_per_iteration);
+    _input_start_y                                       = padding.size() > 1 ? -padding.at(1).first : 0;
+
+    CLBuildOptions build_opts;
+    build_opts.add_option("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type()));
+    build_opts.add_option("-DVEC_SIZE=" + support::cpp11::to_string(num_elems_processed_per_iteration));
+    build_opts.add_option("-DCONST_VAL=" + string_from_pixel_value(constant_value, input->info()->data_type()));
+    build_opts.add_option("-DPAD_LEFT=" + support::cpp11::to_string(padding.at(0).first));
+    build_opts.add_option("-DSRC_WIDTH=" + support::cpp11::to_string(input->info()->dimension(0)));
+    build_opts.add_option("-DSELECT_DT=" + get_cl_select_type_from_data_type(input->info()->data_type()));
+    build_opts.add_option_if(padding.at(0).first > num_elems_processed_per_iteration, "-DTHREADS_TO_SKIP_X=" + support::cpp11::to_string(padding.at(0).first / num_elems_processed_per_iteration));
+
+    if(padding.size() > 1)
+    {
+        build_opts.add_option("-DPAD_TOP=" + support::cpp11::to_string(padding.at(1).first));
+        build_opts.add_option("-DSRC_HEIGHT=" + support::cpp11::to_string(input->info()->dimension(1)));
+
+        if(padding.size() > 2)
+        {
+            build_opts.add_option("-DPAD_NEAR=" + support::cpp11::to_string(padding.at(2).first));
+            build_opts.add_option("-DSRC_DEPTH=" + support::cpp11::to_string(input->info()->dimension(2)));
+        }
+    }
+
+    // Create kernel
+    _kernel = static_cast<cl::Kernel>(CLKernelLibrary::get().create_kernel("pad_layer", build_opts.options()));
+
+    // Configure kernel window
+    auto win_config = validate_and_configure_window(input->info(), output->info(), padding, constant_value, mode);
+    ARM_COMPUTE_ERROR_THROW_ON(win_config.first);
+    ICLKernel::configure_internal(win_config.second);
+}
+
+Status CLPadLayerKernel::validate(const ITensorInfo *input, const ITensorInfo *output, const PaddingList &padding, PixelValue constant_value, PaddingMode mode)
+{
+    ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(input, output, padding, constant_value, mode));
+    ARM_COMPUTE_RETURN_ON_ERROR(validate_and_configure_window(input->clone().get(), output->clone().get(), padding, constant_value, mode).first);
+
+    return Status{};
+}
+
+void CLPadLayerKernel::run(const Window &window, cl::CommandQueue &queue)
+{
+    ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this);
+    ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(ICLKernel::window(), window);
+
+    Window win_in = window;
+    win_in.adjust(Window::DimX, _input_start_x, true);
+    win_in.adjust(Window::DimY, _input_start_y, true);
+
+    Window slice_out = window.first_slice_window_3D();
+    Window slice_in  = win_in.first_slice_window_3D();
+    do
+    {
+        unsigned int idx = 0;
+        add_3D_tensor_argument(idx, _input, slice_in);
+        add_3D_tensor_argument(idx, _output, slice_out);
+
+        enqueue(queue, *this, slice_out);
+    }
+    while(window.slide_window_slice_3D(slice_out) && win_in.slide_window_slice_3D(slice_in));
+}
+} // namespace arm_compute
diff --git a/src/runtime/CL/functions/CLPadLayer.cpp b/src/runtime/CL/functions/CLPadLayer.cpp
index 99e312183a..dcd0a5b1ad 100644
--- a/src/runtime/CL/functions/CLPadLayer.cpp
+++ b/src/runtime/CL/functions/CLPadLayer.cpp
@@ -23,40 +23,14 @@
  */
 #include "arm_compute/runtime/CL/functions/CLPadLayer.h"
 
-#include "arm_compute/core/CL/ICLTensor.h"
-#include "arm_compute/core/Types.h"
 #include "arm_compute/core/utils/misc/ShapeCalculator.h"
-#include "support/ToolchainSupport.h"
 
 namespace arm_compute
 {
 CLPadLayer::CLPadLayer()
-    : _copy_kernel(), _mode(), _padding(), _memset_kernel(), _num_dimensions(0), _slice_functions(), _concat_functions(), _slice_results(), _concat_results()
+    : _pad_kernel(), _copy_kernel(), _mode(), _padding(), _num_dimensions(0), _slice_functions(), _concat_functions(), _slice_results(), _concat_results()
 {
 }
-
-void CLPadLayer::configure_constant_mode(ICLTensor *input, ICLTensor *output, const PaddingList &padding, const PixelValue constant_value)
-{
-    // Set the pages of the output to the constant_value.
-    _memset_kernel.configure(output, constant_value);
-
-    // Fill out padding list with zeroes.
-    PaddingList padding_extended = padding;
-    for(size_t i = padding.size(); i < TensorShape::num_max_dimensions; i++)
-    {
-        padding_extended.emplace_back(PaddingInfo{ 0, 0 });
-    }
-
-    // Create a window within the output tensor where the input will be copied.
-    Window copy_window = Window();
-    for(uint32_t i = 0; i < output->info()->num_dimensions(); ++i)
-    {
-        copy_window.set(i, Window::Dimension(padding_extended[i].first, padding_extended[i].first + input->info()->dimension(i), 1));
-    }
-    // Copy the input to the output, leaving the padding filled with the constant_value.
-    _copy_kernel.configure(input, output, PaddingList(), &copy_window);
-}
-
 void CLPadLayer::configure_reflect_symmetric_mode(ICLTensor *input, ICLTensor *output)
 {
     int64_t last_padding_dimension = _padding.size() - 1;
@@ -163,14 +137,12 @@ void CLPadLayer::configure_reflect_symmetric_mode(ICLTensor *input, ICLTensor *o
 
 void CLPadLayer::configure(ICLTensor *input, ICLTensor *output, const PaddingList &padding, PixelValue constant_value, PaddingMode mode)
 {
-    ARM_COMPUTE_ERROR_THROW_ON(validate(input->info(), output->info(), padding, constant_value, mode));
-
     _padding = padding;
     _mode    = mode;
 
     TensorShape padded_shape = misc::shape_calculator::compute_padded_shape(input->info()->tensor_shape(), _padding);
-
     auto_init_if_empty(*output->info(), input->info()->clone()->set_tensor_shape(padded_shape));
+    ARM_COMPUTE_ERROR_THROW_ON(validate(input->info(), output->info(), padding, constant_value, mode));
 
     // Find the last dimension requiring padding so that it is known when to write to output and whether any padding is applied.
     int64_t last_padding_dimension = _padding.size() - 1;
@@ -188,7 +160,7 @@ void CLPadLayer::configure(ICLTensor *input, ICLTensor *output, const PaddingLis
         {
             case PaddingMode::CONSTANT:
             {
-                configure_constant_mode(input, output, padding, constant_value);
+                _pad_kernel.configure(input, output, padding, constant_value, mode);
                 break;
             }
             case PaddingMode::REFLECT:
@@ -231,18 +203,17 @@ Status CLPadLayer::validate(const ITensorInfo *input, const ITensorInfo *output,
         ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DIMENSIONS(output->tensor_shape(), padded_shape);
         ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(output, input);
         ARM_COMPUTE_RETURN_ON_ERROR(CLCopyKernel::validate(input, output, PaddingList(), &copy_window));
-        ARM_COMPUTE_RETURN_ON_ERROR(CLMemsetKernel::validate(output, constant_value));
     }
     else
     {
         ARM_COMPUTE_RETURN_ON_ERROR(CLCopyKernel::validate(input, &input->clone()->set_tensor_shape(padded_shape), PaddingList(), &copy_window));
-        ARM_COMPUTE_RETURN_ON_ERROR(CLMemsetKernel::validate(&input->clone()->set_tensor_shape(padded_shape), constant_value));
     }
 
     switch(mode)
     {
         case PaddingMode::CONSTANT:
         {
+            ARM_COMPUTE_RETURN_ON_ERROR(CLPadLayerKernel::validate(input, output, padding, constant_value, mode));
             break;
         }
         case PaddingMode::REFLECT:
@@ -279,8 +250,7 @@ void CLPadLayer::run()
         {
             case PaddingMode::CONSTANT:
             {
-                CLScheduler::get().enqueue(_memset_kernel, false);
-                CLScheduler::get().enqueue(_copy_kernel, true);
+                CLScheduler::get().enqueue(_pad_kernel, false);
                 break;
             }
             case PaddingMode::REFLECT: