COMPMID-895 - Optimizing CLDepthwiseConvolution3x3Kernel

This patch brings the MACs utilisation up to 25 % when both stride_x and stride_y are equal to 1
Performance reported in the following confluence page:
https://confluence.arm.com/display/MLENG/Depthwise+convolution+3x3+FP32+performance%3A+ACL+18.02

Change-Id: Ida1b64be9a88805902a3d90194559b58eb1224a3
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/119068
Reviewed-by: Michalis Spyrou <michalis.spyrou@arm.com>
Tested-by: Jenkins <bsgcomp@arm.com>

5 files changed, 291 insertions, 68 deletions

diff --git a/examples/cl_sgemm.cpp b/examples/cl_sgemm.cpp
index fa57885450..966661b9b4 100644
--- a/examples/cl_sgemm.cpp
+++ b/examples/cl_sgemm.cpp
@@ -198,4 +198,4 @@ private:
 int main(int argc, char **argv)
 {
     return utils::run_example<CLSGEMMExample>(argc, argv);
-}
+}
\ No newline at end of file
diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp
index d7ae6e3127..c26d8d80a6 100644
--- a/src/core/CL/CLKernelLibrary.cpp
+++ b/src/core/CL/CLKernelLibrary.cpp
@@ -189,6 +189,8 @@ const std::map<std::string, std::string> CLKernelLibrary::_kernel_program_map =
     { "deconvolution_upsample", "deconvolution_layer.cl" },
     { "depthwise_convolution_3x3", "depthwise_convolution.cl" },
     { "depthwise_convolution_3x3_quantized", "depthwise_convolution_quantized.cl" },
+    { "depthwise_convolution_3x3_stridex1_stridey1_bifrost", "depthwise_convolution.cl" },
+    { "depthwise_convolution_3x3_stridex2_stridey2_bifrost", "depthwise_convolution.cl" },
     { "depthwise_im2col", "depthwise_convolution.cl" },
     { "depthwise_vector_to_tensor", "depthwise_convolution.cl" },
     { "depthwise_weights_reshape", "depthwise_convolution.cl" },
diff --git a/src/core/CL/cl_kernels/depthwise_convolution.cl b/src/core/CL/cl_kernels/depthwise_convolution.cl
index 89555a0cb6..ac94b693e3 100644
--- a/src/core/CL/cl_kernels/depthwise_convolution.cl
+++ b/src/core/CL/cl_kernels/depthwise_convolution.cl
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2017 ARM Limited.
+ * Copyright (c) 2017-2018 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -144,9 +144,9 @@ inline float2 convolution3x3(
     return pixels;
 }
 
-/** This function computes the horizontal integral of the image.
+/** This OpenCL kernel computes the depthwise convolution 3x3
  *
- * @param[in] src_ptr                               Pointer to the source image. Supported data types: U8
+ * @param[in] src_ptr                               Pointer to the source image. Supported data types: 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)
@@ -154,7 +154,7 @@ inline float2 convolution3x3(
  * @param[in] src_offset_first_element_in_bytes     The offset of the first element in the source image
  * @param[in] src_stride_z                          Stride of the source tensor in Z dimension (in bytes)
  * @param[in] src_step_z                            src_stride_z * number of elements along Y processed per workitem(in bytes)
- * @param[in] dst_ptr                               Pointer to the destination tensor. Supported data types: F16/F32
+ * @param[in] dst_ptr                               Pointer to the destination tensor. Supported data types: F32
  * @param[in] dst_stride_x                          Stride of the destination tensor 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 tensor in Y dimension (in bytes)
@@ -162,7 +162,7 @@ inline float2 convolution3x3(
  * @param[in] dst_stride_z                          Stride of the destination tensor in Z dimension (in bytes)
  * @param[in] dst_step_z                            dst_stride_z * number of elements along Y processed per workitem(in bytes)
  * @param[in] dst_offset_first_element_in_bytes     The offset of the first element in the destination tensor
- * @param[in] weights_ptr                           Pointer to the weights tensor. Supported data types: F16/F32
+ * @param[in] weights_ptr                           Pointer to the weights tensor. Supported data types: F32
  * @param[in] weights_stride_x                      Stride of the weights tensor in X dimension (in bytes)
  * @param[in] weights_step_x                        weights_stride_x * number of elements along X processed per workitem(in bytes)
  * @param[in] weights_stride_y                      Stride of the weights tensor in Y dimension (in bytes)
@@ -175,7 +175,6 @@ inline float2 convolution3x3(
  * @param[in] biases_step_x                         (Optional) biases_stride_x * number of elements along X processed per workitem(in bytes)
  * @param[in] biases_offset_first_element_in_bytes  (Optional) The offset of the first element in the biases vector
  */
-
 __kernel void depthwise_convolution_3x3(
     TENSOR3D_DECLARATION(src),
     TENSOR3D_DECLARATION(dst),
@@ -207,9 +206,214 @@ __kernel void depthwise_convolution_3x3(
 
     vstore2(pixels, 0, (__global float *)dst.ptr);
 }
-
 #endif //defined(CONV_STRIDE_X)
 
+#define CONVOLUTION1x3_BIFROST2X1_STRIDE1(acc, src0, weights_row0) \
+    ({                                                             \
+        acc.s0 = fma(src0.s0, weights_row0.s0, acc.s0);            \
+        acc.s0 = fma(src0.s1, weights_row0.s1, acc.s0);            \
+        acc.s0 = fma(src0.s2, weights_row0.s2, acc.s0);            \
+        acc.s1 = fma(src0.s1, weights_row0.s0, acc.s1);            \
+        acc.s1 = fma(src0.s2, weights_row0.s1, acc.s1);            \
+        acc.s1 = fma(src0.s3, weights_row0.s2, acc.s1);            \
+    })
+
+#define CONVOLUTION1x3_BIFROST2X1_STRIDE2(acc, src0, src1, weights_row0) \
+    ({                                                                   \
+        acc.s0 = fma(src0.s0, weights_row0.s0, acc.s0);                  \
+        acc.s0 = fma(src0.s1, weights_row0.s1, acc.s0);                  \
+        acc.s0 = fma(src0.s2, weights_row0.s2, acc.s0);                  \
+        acc.s1 = fma(src0.s2, weights_row0.s0, acc.s1);                  \
+        acc.s1 = fma(src0.s3, weights_row0.s1, acc.s1);                  \
+        acc.s1 = fma(src1.s0, weights_row0.s2, acc.s1);                  \
+    })
+
+/** This OpenCL kernel is optimized for Bifrost architectures and computes the depthwise convolution 3x3 when both
+ * stride_x and stride_y are equal to 1
+ *
+ * @param[in] src_ptr                               Pointer to the source image. Supported data types: 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_offset_first_element_in_bytes     The offset of the first element in the source image
+ * @param[in] src_stride_z                          Stride of the source tensor in Z dimension (in bytes)
+ * @param[in] src_step_z                            src_stride_z * number of elements along Y processed per workitem(in bytes)
+ * @param[in] dst_ptr                               Pointer to the destination tensor. Supported data types: F32
+ * @param[in] dst_stride_x                          Stride of the destination tensor 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 tensor 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 tensor in Z dimension (in bytes)
+ * @param[in] dst_step_z                            dst_stride_z * number of elements along Y processed per workitem(in bytes)
+ * @param[in] dst_offset_first_element_in_bytes     The offset of the first element in the destination tensor
+ * @param[in] weights_ptr                           Pointer to the weights tensor. Supported data types: F32
+ * @param[in] weights_stride_x                      Stride of the weights tensor in X dimension (in bytes)
+ * @param[in] weights_step_x                        weights_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in] weights_stride_y                      Stride of the weights tensor in Y dimension (in bytes)
+ * @param[in] weights_step_y                        weights_stride_y * number of elements along Y processed per workitem(in bytes)
+ * @param[in] weights_stride_z                      Stride of the weights tensor in Z dimension (in bytes)
+ * @param[in] weights_step_z                        weights_stride_z * number of elements along Y processed per workitem(in bytes)
+ * @param[in] weights_offset_first_element_in_bytes The offset of the first element in the biases vector
+ * @param[in] biases_ptr                            (Optional) Pointer to the biases vector. Supported data types: F32
+ * @param[in] biases_stride_x                       (Optional) Stride of the biases vector in X dimension (in bytes)
+ * @param[in] biases_step_x                         (Optional) biases_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in] biases_offset_first_element_in_bytes  (Optional) The offset of the first element in the biases vector
+ */
+__kernel void depthwise_convolution_3x3_stridex1_stridey1_bifrost(
+    TENSOR3D_DECLARATION(src),
+    TENSOR3D_DECLARATION(dst),
+    TENSOR3D_DECLARATION(weights)
+#if defined(HAS_BIAS)
+    ,
+    VECTOR_DECLARATION(biases)
+#endif //defined(HAS_BIAS)
+)
+{
+    Image    src     = CONVERT_TENSOR3D_TO_IMAGE_STRUCT(src);
+    Image    dst     = CONVERT_TENSOR3D_TO_IMAGE_STRUCT(dst);
+    Tensor3D weights = CONVERT_TO_TENSOR3D_STRUCT(weights);
+
+    float2 pixels0 = 0.0f;
+    float2 pixels1 = 0.0f;
+    float2 pixels2 = 0.0f;
+    float2 pixels3 = 0.0f;
+
+    __global uchar *weights_addr = (__global uchar *)weights.ptr;
+    __global uchar *src_addr     = (__global uchar *)offset(&src, 0, 0);
+
+    // Load the weights
+    float3 weights_row0 = vload3(0, (__global float *)(weights_addr + 0 * weights_stride_y));
+    float3 weights_row1 = vload3(0, (__global float *)(weights_addr + 1 * weights_stride_y));
+    float3 weights_row2 = vload3(0, (__global float *)(weights_addr + 2 * weights_stride_y));
+
+    // Note: Since each work-item computes 4x2 elements, we need to load 4 rows from the input tensor
+    float4 src00 = vload4(0, (__global float *)(src_addr + 0 * src_stride_y)); // Row0
+    float4 src10 = vload4(0, (__global float *)(src_addr + 1 * src_stride_y)); // Row1
+    float4 src20 = vload4(0, (__global float *)(src_addr + 2 * src_stride_y)); // Row2
+    float4 src30 = vload4(0, (__global float *)(src_addr + 3 * src_stride_y)); // Row3
+    float4 src40 = vload4(0, (__global float *)(src_addr + 4 * src_stride_y)); // Row3
+    float4 src50 = vload4(0, (__global float *)(src_addr + 5 * src_stride_y)); // Row3
+
+    CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels0, src00, weights_row0);
+    CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels0, src10, weights_row1);
+    CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels0, src20, weights_row2);
+    CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels1, src10, weights_row0);
+    CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels1, src20, weights_row1);
+    CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels1, src30, weights_row2);
+    CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels2, src20, weights_row0);
+    CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels2, src30, weights_row1);
+    CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels2, src40, weights_row2);
+    CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels3, src30, weights_row0);
+    CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels3, src40, weights_row1);
+    CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels3, src50, weights_row2);
+
+#ifdef HAS_BIAS
+    Vector biases = CONVERT_TO_VECTOR_STRUCT_NO_STEP(biases);
+
+    float bias = *((__global float *)(vector_offset(&biases, get_global_id(2))));
+
+    pixels0 += (float2)bias;
+    pixels1 += (float2)bias;
+    pixels2 += (float2)bias;
+    pixels3 += (float2)bias;
+#endif /* defined(HAS_BIAS) */
+
+    vstore2(pixels0, 0, (__global float *)(dst.ptr + 0 * dst_stride_y));
+    vstore2(pixels1, 0, (__global float *)(dst.ptr + 1 * dst_stride_y));
+    vstore2(pixels2, 0, (__global float *)(dst.ptr + 2 * dst_stride_y));
+    vstore2(pixels3, 0, (__global float *)(dst.ptr + 3 * dst_stride_y));
+}
+
+/** This OpenCL kernel is optimized for Bifrost architectures and computes the depthwise convolution 3x3 when both
+ * stride_x and stride_y are equal to 2
+ *
+ * @param[in] src_ptr                               Pointer to the source image. Supported data types: 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_offset_first_element_in_bytes     The offset of the first element in the source image
+ * @param[in] src_stride_z                          Stride of the source tensor in Z dimension (in bytes)
+ * @param[in] src_step_z                            src_stride_z * number of elements along Y processed per workitem(in bytes)
+ * @param[in] dst_ptr                               Pointer to the destination tensor. Supported data types: F32
+ * @param[in] dst_stride_x                          Stride of the destination tensor 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 tensor 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 tensor in Z dimension (in bytes)
+ * @param[in] dst_step_z                            dst_stride_z * number of elements along Y processed per workitem(in bytes)
+ * @param[in] dst_offset_first_element_in_bytes     The offset of the first element in the destination tensor
+ * @param[in] weights_ptr                           Pointer to the weights tensor. Supported data types: F32
+ * @param[in] weights_stride_x                      Stride of the weights tensor in X dimension (in bytes)
+ * @param[in] weights_step_x                        weights_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in] weights_stride_y                      Stride of the weights tensor in Y dimension (in bytes)
+ * @param[in] weights_step_y                        weights_stride_y * number of elements along Y processed per workitem(in bytes)
+ * @param[in] weights_stride_z                      Stride of the weights tensor in Z dimension (in bytes)
+ * @param[in] weights_step_z                        weights_stride_z * number of elements along Y processed per workitem(in bytes)
+ * @param[in] weights_offset_first_element_in_bytes The offset of the first element in the biases vector
+ * @param[in] biases_ptr                            (Optional) Pointer to the biases vector. Supported data types: F32
+ * @param[in] biases_stride_x                       (Optional) Stride of the biases vector in X dimension (in bytes)
+ * @param[in] biases_step_x                         (Optional) biases_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in] biases_offset_first_element_in_bytes  (Optional) The offset of the first element in the biases vector
+ */
+__kernel void depthwise_convolution_3x3_stridex2_stridey2_bifrost(
+    TENSOR3D_DECLARATION(src),
+    TENSOR3D_DECLARATION(dst),
+    TENSOR3D_DECLARATION(weights)
+#if defined(HAS_BIAS)
+    ,
+    VECTOR_DECLARATION(biases)
+#endif //defined(HAS_BIAS)
+)
+{
+    Image    src     = CONVERT_TENSOR3D_TO_IMAGE_STRUCT(src);
+    Image    dst     = CONVERT_TENSOR3D_TO_IMAGE_STRUCT(dst);
+    Tensor3D weights = CONVERT_TO_TENSOR3D_STRUCT(weights);
+
+    float2 pixels0 = 0.0f;
+    float2 pixels1 = 0.0f;
+
+    __global uchar *weights_addr = (__global uchar *)weights.ptr;
+    __global uchar *src_addr     = (__global uchar *)offset(&src, 0, 0);
+
+    // Load the weights
+    float3 weights_row0 = vload3(0, (__global float *)(weights_addr + 0 * weights_stride_y));
+    float3 weights_row1 = vload3(0, (__global float *)(weights_addr + 1 * weights_stride_y));
+    float3 weights_row2 = vload3(0, (__global float *)(weights_addr + 2 * weights_stride_y));
+
+    // Note: Since each work-item computes 4x2 elements, we need to load 5 rows from the input tensor
+    float4 src00 = vload4(0, (__global float *)(src_addr + 0 * src_stride_y)); // Row0
+    float2 src01 = vload2(2, (__global float *)(src_addr + 0 * src_stride_y)); // Row0
+    float4 src10 = vload4(0, (__global float *)(src_addr + 1 * src_stride_y)); // Row1
+    float2 src11 = vload2(2, (__global float *)(src_addr + 1 * src_stride_y)); // Row1
+    float4 src20 = vload4(0, (__global float *)(src_addr + 2 * src_stride_y)); // Row2
+    float2 src21 = vload2(2, (__global float *)(src_addr + 2 * src_stride_y)); // Row2
+    float4 src30 = vload4(0, (__global float *)(src_addr + 3 * src_stride_y)); // Row3
+    float2 src31 = vload2(2, (__global float *)(src_addr + 3 * src_stride_y)); // Row3
+    float4 src40 = vload4(0, (__global float *)(src_addr + 4 * src_stride_y)); // Row4
+    float2 src41 = vload2(2, (__global float *)(src_addr + 4 * src_stride_y)); // Row4
+
+    CONVOLUTION1x3_BIFROST2X1_STRIDE2(pixels0, src00, src01, weights_row0);
+    CONVOLUTION1x3_BIFROST2X1_STRIDE2(pixels0, src10, src11, weights_row1);
+    CONVOLUTION1x3_BIFROST2X1_STRIDE2(pixels0, src20, src21, weights_row2);
+    CONVOLUTION1x3_BIFROST2X1_STRIDE2(pixels1, src20, src21, weights_row0);
+    CONVOLUTION1x3_BIFROST2X1_STRIDE2(pixels1, src30, src31, weights_row1);
+    CONVOLUTION1x3_BIFROST2X1_STRIDE2(pixels1, src40, src41, weights_row2);
+
+#ifdef HAS_BIAS
+    Vector biases = CONVERT_TO_VECTOR_STRUCT_NO_STEP(biases);
+
+    float bias = *((__global float *)(vector_offset(&biases, get_global_id(2))));
+
+    pixels0 += (float2)bias;
+    pixels1 += (float2)bias;
+#endif /* defined(HAS_BIAS) */
+
+    vstore2(pixels0, 0, (__global float *)(dst.ptr + 0 * dst_stride_y));
+    vstore2(pixels1, 0, (__global float *)(dst.ptr + 1 * dst_stride_y));
+}
+
 #if defined(SRC_WIDTH) && defined(DATA_TYPE)
 /** This kernel reshapes each of the tensor's low three dimensions to single rows.
  *
@@ -288,7 +492,6 @@ __kernel void depthwise_weights_reshape(
  * @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 tensor
  */
-
 __kernel void depthwise_im2col(TENSOR3D_DECLARATION(src), TENSOR3D_DECLARATION(dst))
 {
     Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst);
diff --git a/src/core/CL/kernels/CLDepthwiseConvolutionLayer3x3Kernel.cpp b/src/core/CL/kernels/CLDepthwiseConvolutionLayer3x3Kernel.cpp
index a9167ee859..2a60f60723 100644
--- a/src/core/CL/kernels/CLDepthwiseConvolutionLayer3x3Kernel.cpp
+++ b/src/core/CL/kernels/CLDepthwiseConvolutionLayer3x3Kernel.cpp
@@ -98,76 +98,74 @@ void CLDepthwiseConvolutionLayer3x3Kernel::configure(const ICLTensor *input, con
     build_opts.add_option("-DCONV_STRIDE_X=" + support::cpp11::to_string(_conv_stride_x));
     build_opts.add_option_if(_biases != nullptr, "-DHAS_BIAS");
 
-    // Create kernel
-    std::string kernel_name = is_data_type_quantized_asymmetric(_input->info()->data_type()) ? "depthwise_convolution_3x3_quantized" : "depthwise_convolution_3x3";
-    _kernel                 = static_cast<cl::Kernel>(CLKernelLibrary::get().create_kernel(kernel_name, build_opts.options()));
+    // Configure the local work size for Bifrost with a value obtained
+    // via exhaustive autotuning for the MobileNets tensor shapes.
+    const GPUTarget gpu_target = get_arch_from_target(get_target());
 
-    // Set static arguments
-    if(is_data_type_quantized_asymmetric(_input->info()->data_type()))
-    {
-        float multiplier        = _input->info()->quantization_info().scale * _weights->info()->quantization_info().scale / _output->info()->quantization_info().scale;
-        int   output_multiplier = 0;
-        int   output_shift      = 0;
-        quantization::calculate_quantized_multiplier_less_than_one(multiplier, &output_multiplier, &output_shift);
+    // Configure kernel window
+    const unsigned int conv_pad_left   = std::max(conv_info.pad_left(), 1U);
+    const unsigned int conv_pad_top    = std::max(conv_info.pad_top(), 1U);
+    const unsigned int conv_pad_right  = std::max(conv_info.pad_right(), 1U);
+    const unsigned int conv_pad_bottom = std::max(conv_info.pad_bottom(), 1U);
 
-        unsigned int idx = 3 * num_arguments_per_3D_tensor() + ((_biases != nullptr) ? num_arguments_per_1D_tensor() : 0);
+    unsigned int num_elems_read_per_iteration_x    = 0;
+    unsigned int num_elems_read_per_iteration_y    = 0;
+    unsigned int num_elems_written_per_iteration_x = 0;
+    unsigned int num_elems_written_per_iteration_y = 0;
 
-        _kernel.setArg(idx++, -_input->info()->quantization_info().offset);
-        _kernel.setArg(idx++, -_weights->info()->quantization_info().offset);
-        _kernel.setArg(idx++, _output->info()->quantization_info().offset);
-        _kernel.setArg(idx++, output_multiplier);
-        _kernel.setArg(idx++, output_shift);
-    }
+    // Create kernel
+    std::string kernel_name;
 
-    // Configure the local work size for Bifrost with a value obtained
-    // via exhaustive autotuning for the MobileNets tensor shapes.
-    const GPUTarget gpu_target = get_arch_from_target(get_target());
-    if(gpu_target == GPUTarget::BIFROST)
+    if(input->info()->data_type() == DataType::F32 && gpu_target == GPUTarget::BIFROST)
     {
-        // Assume uniform padding and striding.
-        const size_t pad    = _conv_pad_left;
-        const size_t stride = _conv_stride_x;
-        const size_t width  = input->info()->dimension(0);
-        if(pad == 1)
+        if(_conv_stride_x == 1 && _conv_stride_y == 1)
         {
-            const size_t width_by_stride = width / stride;
-            if(width_by_stride == 28) // 56/2 or 28/1
-            {
-                _lws_hint = cl::NDRange(7, 4, 3);
-            }
-            else if(width_by_stride == 14) // 28/2 or 14/1
-            {
-                _lws_hint = cl::NDRange(7, 7, 4);
-            }
+            kernel_name                       = "depthwise_convolution_3x3_stridex1_stridey1_bifrost";
+            num_elems_read_per_iteration_x    = 4;
+            num_elems_read_per_iteration_y    = 6;
+            num_elems_written_per_iteration_x = 2;
+            num_elems_written_per_iteration_y = 4;
         }
-        else if(pad == 0)
+        else if(_conv_stride_x == 2 && _conv_stride_y == 2)
         {
-            if(width >= 56) // 56 or 112
-            {
-                _lws_hint = cl::NDRange(8, 5, 2);
-            }
-            else if(width >= 14) // 14 or 28
-            {
-                _lws_hint = cl::NDRange(1, 5, 2);
-            }
-            else // 7
-            {
-                _lws_hint = cl::NDRange(1, 1, 2);
-            }
+            kernel_name                       = "depthwise_convolution_3x3_stridex2_stridey2_bifrost";
+            num_elems_read_per_iteration_x    = 6;
+            num_elems_read_per_iteration_y    = 5;
+            num_elems_written_per_iteration_x = 2;
+            num_elems_written_per_iteration_y = 2;
         }
+        else
+        {
+            kernel_name                       = "depthwise_convolution_3x3";
+            num_elems_written_per_iteration_x = 8 / data_size_from_type(input->info()->data_type());
+            num_elems_written_per_iteration_y = 1;
+            num_elems_read_per_iteration_x    = 3 + (num_elems_written_per_iteration_x - 1) * _conv_stride_x;
+            num_elems_read_per_iteration_y    = 3;
+        }
+    }
+    else
+    {
+        kernel_name                       = is_data_type_quantized_asymmetric(_input->info()->data_type()) ? "depthwise_convolution_3x3_quantized" : "depthwise_convolution_3x3";
+        num_elems_written_per_iteration_x = 8 / data_size_from_type(input->info()->data_type());
+        num_elems_written_per_iteration_y = 1;
+        num_elems_read_per_iteration_x    = 3 + (num_elems_written_per_iteration_x - 1) * _conv_stride_x;
+        num_elems_read_per_iteration_y    = 3;
     }
 
-    // Configure kernel window
-    const unsigned int num_elems_processed_per_iteration = 8 / data_size_from_type(input->info()->data_type());
-    const unsigned int num_elems_written_per_iteration   = num_elems_processed_per_iteration;
-    const unsigned int num_elems_read_per_iteration      = 3 + (num_elems_processed_per_iteration - 1) * _conv_stride_x;
-    const unsigned int num_rows_read_per_iteration       = 3;
+    // Calculate right and bottom border
+    int input_width  = input->info()->dimension(0) + conv_pad_left + conv_pad_right;
+    int input_height = input->info()->dimension(1) + conv_pad_top + conv_pad_bottom;
+
+    // Add padding only if necessary or it would always result in a window_changed
+    input_width  = ceil_to_multiple(input_width, num_elems_read_per_iteration_x);
+    input_height = ceil_to_multiple(input_height, num_elems_read_per_iteration_y);
 
-    Window win = calculate_max_window(*output->info(), Steps(num_elems_processed_per_iteration));
+    // Create window and update padding
+    Window win = calculate_max_window(*output->info(), Steps(num_elems_written_per_iteration_x, num_elems_written_per_iteration_y));
 
-    AccessWindowRectangle  input_access(input->info(), -border_size().left, -border_size().top, num_elems_read_per_iteration, num_rows_read_per_iteration, _conv_stride_x, _conv_stride_y);
-    AccessWindowHorizontal output_access(output->info(), 0, num_elems_written_per_iteration);
-    AccessWindowStatic     weights_access(weights->info(), 0, 0, weights->info()->dimension(0), weights->info()->dimension(1));
+    AccessWindowStatic    input_access(input->info(), -conv_pad_left, -conv_pad_top, input_width, input_height);
+    AccessWindowStatic    weights_access(weights->info(), 0, 0, 3, 3);
+    AccessWindowRectangle output_access(output->info(), 0, 0, num_elems_written_per_iteration_x, num_elems_written_per_iteration_y);
 
     update_window_and_padding(win, input_access, weights_access, output_access);
 
@@ -175,8 +173,28 @@ void CLDepthwiseConvolutionLayer3x3Kernel::configure(const ICLTensor *input, con
 
     ICLKernel::configure(win);
 
+    _kernel = static_cast<cl::Kernel>(CLKernelLibrary::get().create_kernel(kernel_name, build_opts.options()));
+
+    // Set static arguments
+    if(is_data_type_quantized_asymmetric(_input->info()->data_type()))
+    {
+        float multiplier        = _input->info()->quantization_info().scale * _weights->info()->quantization_info().scale / _output->info()->quantization_info().scale;
+        int   output_multiplier = 0;
+        int   output_shift      = 0;
+        quantization::calculate_quantized_multiplier_less_than_one(multiplier, &output_multiplier, &output_shift);
+
+        unsigned int idx = 3 * num_arguments_per_3D_tensor() + ((_biases != nullptr) ? num_arguments_per_1D_tensor() : 0);
+
+        _kernel.setArg(idx++, -_input->info()->quantization_info().offset);
+        _kernel.setArg(idx++, -_weights->info()->quantization_info().offset);
+        _kernel.setArg(idx++, _output->info()->quantization_info().offset);
+        _kernel.setArg(idx++, output_multiplier);
+        _kernel.setArg(idx++, output_shift);
+    }
+
     // Set config_id for enabling LWS tuning
-    _config_id = "depthwise_convolution3x3_";
+    _config_id = kernel_name;
+    _config_id += "_";
     _config_id += lower_string(string_from_data_type(input->info()->data_type()));
     _config_id += "_";
     _config_id += support::cpp11::to_string(input->info()->dimension(0));
diff --git a/src/runtime/CL/CLTuner.cpp b/src/runtime/CL/CLTuner.cpp
index 917d56756a..cf5b5bce2d 100644
--- a/src/runtime/CL/CLTuner.cpp
+++ b/src/runtime/CL/CLTuner.cpp
@@ -204,4 +204,4 @@ cl_int Interceptor::operator()(cl_command_queue command_queue, cl_kernel kernel,
     _tuner.set_cl_kernel_event(tmp);
 
     return retval;
-}
\ No newline at end of file
+}