From 3ebef32816435516f68cefba689dba7216464154 Mon Sep 17 00:00:00 2001 From: Michele Di Giorgio Date: Wed, 21 Feb 2018 10:02:58 +0000 Subject: COMPMID-949: Optimizing CLDepthwiseConvolution3x3Kernel for FP16 Change-Id: I2af6544eab17004c5b3de56557cb2cc5efecc915 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/122181 Tested-by: Jenkins Reviewed-by: Michalis Spyrou --- src/core/CL/CLKernelLibrary.cpp | 6 +- src/core/CL/cl_kernels/depthwise_convolution.cl | 236 ++++++++++++++++++++- .../CLDepthwiseConvolutionLayer3x3Kernel.cpp | 26 ++- 3 files changed, 256 insertions(+), 12 deletions(-) (limited to 'src') diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp index d815ac1afc..3ff3acd697 100644 --- a/src/core/CL/CLKernelLibrary.cpp +++ b/src/core/CL/CLKernelLibrary.cpp @@ -192,8 +192,10 @@ const std::map CLKernelLibrary::_kernel_program_map = { "depthwise_convolution_3x3", "depthwise_convolution.cl" }, { "depthwise_convolution_3x3_f16", "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_convolution_3x3_stridex1_stridey1_bifrost_f16", "depthwise_convolution.cl" }, + { "depthwise_convolution_3x3_stridex2_stridey2_bifrost_f16", "depthwise_convolution.cl" }, + { "depthwise_convolution_3x3_stridex1_stridey1_bifrost_f32", "depthwise_convolution.cl" }, + { "depthwise_convolution_3x3_stridex2_stridey2_bifrost_f32", "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 f352138776..07e67f4f2c 100644 --- a/src/core/CL/cl_kernels/depthwise_convolution.cl +++ b/src/core/CL/cl_kernels/depthwise_convolution.cl @@ -218,6 +218,22 @@ __kernel void depthwise_convolution_3x3( acc.s1 = fma(src0.s3, weights_row0.s2, acc.s1); \ }) +#define CONVOLUTION1x3_BIFROST4X1_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); \ + acc.s2 = fma(src0.s2, weights_row0.s0, acc.s2); \ + acc.s2 = fma(src0.s3, weights_row0.s1, acc.s2); \ + acc.s2 = fma(src0.s4, weights_row0.s2, acc.s2); \ + acc.s3 = fma(src0.s3, weights_row0.s0, acc.s3); \ + acc.s3 = fma(src0.s4, weights_row0.s1, acc.s3); \ + acc.s3 = fma(src0.s5, weights_row0.s2, acc.s3); \ + }) + #define CONVOLUTION1x3_BIFROST2X1_STRIDE2(acc, src0, src1, weights_row0) \ ({ \ acc.s0 = fma(src0.s0, weights_row0.s0, acc.s0); \ @@ -228,6 +244,22 @@ __kernel void depthwise_convolution_3x3( acc.s1 = fma(src1.s0, weights_row0.s2, acc.s1); \ }) +#define CONVOLUTION1x3_BIFROST4X1_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(src0.s4, weights_row0.s2, acc.s1); \ + acc.s2 = fma(src0.s4, weights_row0.s0, acc.s2); \ + acc.s2 = fma(src0.s5, weights_row0.s1, acc.s2); \ + acc.s2 = fma(src0.s6, weights_row0.s2, acc.s2); \ + acc.s3 = fma(src0.s6, weights_row0.s0, acc.s3); \ + acc.s3 = fma(src0.s7, weights_row0.s1, acc.s3); \ + acc.s3 = fma(src1.s0, weights_row0.s2, acc.s3); \ + }) + /** 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 * @@ -260,7 +292,7 @@ __kernel void depthwise_convolution_3x3( * @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( +__kernel void depthwise_convolution_3x3_stridex1_stridey1_bifrost_f32( TENSOR3D_DECLARATION(src), TENSOR3D_DECLARATION(dst), TENSOR3D_DECLARATION(weights) @@ -287,13 +319,13 @@ __kernel void depthwise_convolution_3x3_stridex1_stridey1_bifrost( 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 + // Note: Since each work-item computes 4x2 elements, we need to load 6 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 + float4 src40 = vload4(0, (__global float *)(src_addr + 4 * src_stride_y)); // Row4 + float4 src50 = vload4(0, (__global float *)(src_addr + 5 * src_stride_y)); // Row5 CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels0, src00, weights_row0); CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels0, src10, weights_row1); @@ -357,7 +389,7 @@ __kernel void depthwise_convolution_3x3_stridex1_stridey1_bifrost( * @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( +__kernel void depthwise_convolution_3x3_stridex2_stridey2_bifrost_f32( TENSOR3D_DECLARATION(src), TENSOR3D_DECLARATION(dst), TENSOR3D_DECLARATION(weights) @@ -694,7 +726,7 @@ inline half4 convolution3x3_f16( * @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_ptr Pointer to the destination tensor. Supported data types: same as @p src_ptr * @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) @@ -702,7 +734,7 @@ inline half4 convolution3x3_f16( * @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_ptr Pointer to the weights tensor. Supported data types: same as @p src_ptr * @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) @@ -747,4 +779,194 @@ __kernel void depthwise_convolution_3x3_f16( vstore4(pixels, 0, (__global half *)dst.ptr); } #endif // defined(CONV_STRIDE_X) + +/** This OpenCL kernel is optimized for Bifrost architectures and computes the 16bit floating point 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: F16 + * @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: same as @p src_ptr + * @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: same as @p src_ptr + * @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: same as @p src_ptr + * @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_f16( + 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); + +#ifdef HAS_BIAS + Vector biases = CONVERT_TO_VECTOR_STRUCT_NO_STEP(biases); + + half bias = *((__global half *)(vector_offset(&biases, get_global_id(2)))); +#endif /* defined(HAS_BIAS) */ + + half4 pixels0 = 0.0f; + half4 pixels1 = 0.0f; + half4 pixels2 = 0.0f; + half4 pixels3 = 0.0f; + + __global uchar *weights_addr = (__global uchar *)weights.ptr; + __global uchar *src_addr = (__global uchar *)offset(&src, 0, 0); + + // Load the weights + half3 weights_row0 = vload3(0, (__global half *)(weights_addr + 0 * weights_stride_y)); + half3 weights_row1 = vload3(0, (__global half *)(weights_addr + 1 * weights_stride_y)); + half3 weights_row2 = vload3(0, (__global half *)(weights_addr + 2 * weights_stride_y)); + + // Note: Since each work-item computes 4x4 elements, we need to load 6 rows from the input tensor + half8 src00 = vload8(0, (__global half *)(src_addr + 0 * src_stride_y)); // Row0 + half8 src10 = vload8(0, (__global half *)(src_addr + 1 * src_stride_y)); // Row1 + half8 src20 = vload8(0, (__global half *)(src_addr + 2 * src_stride_y)); // Row2 + half8 src30 = vload8(0, (__global half *)(src_addr + 3 * src_stride_y)); // Row3 + half8 src40 = vload8(0, (__global half *)(src_addr + 4 * src_stride_y)); // Row4 + half8 src50 = vload8(0, (__global half *)(src_addr + 5 * src_stride_y)); // Row5 + + CONVOLUTION1x3_BIFROST4X1_STRIDE1(pixels0, src00, weights_row0); + CONVOLUTION1x3_BIFROST4X1_STRIDE1(pixels0, src10, weights_row1); + CONVOLUTION1x3_BIFROST4X1_STRIDE1(pixels0, src20, weights_row2); + CONVOLUTION1x3_BIFROST4X1_STRIDE1(pixels1, src10, weights_row0); + CONVOLUTION1x3_BIFROST4X1_STRIDE1(pixels1, src20, weights_row1); + CONVOLUTION1x3_BIFROST4X1_STRIDE1(pixels1, src30, weights_row2); + CONVOLUTION1x3_BIFROST4X1_STRIDE1(pixels2, src20, weights_row0); + CONVOLUTION1x3_BIFROST4X1_STRIDE1(pixels2, src30, weights_row1); + CONVOLUTION1x3_BIFROST4X1_STRIDE1(pixels2, src40, weights_row2); + CONVOLUTION1x3_BIFROST4X1_STRIDE1(pixels3, src30, weights_row0); + CONVOLUTION1x3_BIFROST4X1_STRIDE1(pixels3, src40, weights_row1); + CONVOLUTION1x3_BIFROST4X1_STRIDE1(pixels3, src50, weights_row2); + +#ifdef HAS_BIAS + pixels0 += (half4)bias; + pixels1 += (half4)bias; + pixels2 += (half4)bias; + pixels3 += (half4)bias; +#endif /* defined(HAS_BIAS) */ + + vstore4(pixels0, 0, (__global half *)(dst.ptr + 0 * dst_stride_y)); + vstore4(pixels1, 0, (__global half *)(dst.ptr + 1 * dst_stride_y)); + vstore4(pixels2, 0, (__global half *)(dst.ptr + 2 * dst_stride_y)); + vstore4(pixels3, 0, (__global half *)(dst.ptr + 3 * dst_stride_y)); +} + +/** This OpenCL kernel is optimized for Bifrost architectures and computes 16bit floating point 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: F16 + * @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: same as @p src_ptr + * @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: same as @p src_ptr + * @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: same as @p src_ptr + * @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_f16( + 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); + +#ifdef HAS_BIAS + Vector biases = CONVERT_TO_VECTOR_STRUCT_NO_STEP(biases); + + half bias = *((__global half *)(vector_offset(&biases, get_global_id(2)))); +#endif /* defined(HAS_BIAS) */ + + half4 pixels0 = 0.0f; + half4 pixels1 = 0.0f; + + __global uchar *weights_addr = (__global uchar *)weights.ptr; + __global uchar *src_addr = (__global uchar *)offset(&src, 0, 0); + + // Load the weights + half3 weights_row0 = vload3(0, (__global half *)(weights_addr + 0 * weights_stride_y)); + half3 weights_row1 = vload3(0, (__global half *)(weights_addr + 1 * weights_stride_y)); + half3 weights_row2 = vload3(0, (__global half *)(weights_addr + 2 * weights_stride_y)); + + // Note: Since each work-item computes 2x4 elements, we need to load 5 rows from the input tensor + half8 src00 = vload8(0, (__global half *)(src_addr + 0 * src_stride_y)); // Row0 + half2 src01 = vload2(4, (__global half *)(src_addr + 0 * src_stride_y)); // Row0 + half8 src10 = vload8(0, (__global half *)(src_addr + 1 * src_stride_y)); // Row1 + half2 src11 = vload2(4, (__global half *)(src_addr + 1 * src_stride_y)); // Row1 + half8 src20 = vload8(0, (__global half *)(src_addr + 2 * src_stride_y)); // Row2 + half2 src21 = vload2(4, (__global half *)(src_addr + 2 * src_stride_y)); // Row2 + half8 src30 = vload8(0, (__global half *)(src_addr + 3 * src_stride_y)); // Row3 + half2 src31 = vload2(4, (__global half *)(src_addr + 3 * src_stride_y)); // Row3 + half8 src40 = vload8(0, (__global half *)(src_addr + 4 * src_stride_y)); // Row4 + half2 src41 = vload2(4, (__global half *)(src_addr + 4 * src_stride_y)); // Row4 + + CONVOLUTION1x3_BIFROST4X1_STRIDE2(pixels0, src00, src01, weights_row0); + CONVOLUTION1x3_BIFROST4X1_STRIDE2(pixels0, src10, src11, weights_row1); + CONVOLUTION1x3_BIFROST4X1_STRIDE2(pixels0, src20, src21, weights_row2); + CONVOLUTION1x3_BIFROST4X1_STRIDE2(pixels1, src20, src21, weights_row0); + CONVOLUTION1x3_BIFROST4X1_STRIDE2(pixels1, src30, src31, weights_row1); + CONVOLUTION1x3_BIFROST4X1_STRIDE2(pixels1, src40, src41, weights_row2); + +#ifdef HAS_BIAS + pixels0 += (half4)bias; + pixels1 += (half4)bias; +#endif /* defined(HAS_BIAS) */ + + vstore4(pixels0, 0, (__global half *)(dst.ptr + 0 * dst_stride_y)); + vstore4(pixels1, 0, (__global half *)(dst.ptr + 1 * dst_stride_y)); +} #endif // defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) diff --git a/src/core/CL/kernels/CLDepthwiseConvolutionLayer3x3Kernel.cpp b/src/core/CL/kernels/CLDepthwiseConvolutionLayer3x3Kernel.cpp index d50e4d695e..9e5585cba4 100644 --- a/src/core/CL/kernels/CLDepthwiseConvolutionLayer3x3Kernel.cpp +++ b/src/core/CL/kernels/CLDepthwiseConvolutionLayer3x3Kernel.cpp @@ -147,6 +147,7 @@ void CLDepthwiseConvolutionLayer3x3Kernel::configure(const ICLTensor *input, con // Configure the local work size for Bifrost with a value obtained // via exhaustive autotuning for the MobileNets tensor shapes. const GPUTarget gpu_target = get_target(); + const bool is_bifrost = gpu_target_is_in(gpu_target, GPUTarget::G71, GPUTarget::G72); // Configure kernel window unsigned int num_elems_read_per_iteration_x = 0; @@ -178,12 +179,31 @@ void CLDepthwiseConvolutionLayer3x3Kernel::configure(const ICLTensor *input, con num_elems_read_per_iteration_x = 3 + (num_elems_written_per_iteration_x - 1) * _conv_stride_x; break; } + if(is_bifrost) + { + if(_conv_stride_x == 1 && _conv_stride_y == 1) + { + kernel_name = "depthwise_convolution_3x3_stridex1_stridey1_bifrost_f16"; + num_elems_read_per_iteration_x = 8; + num_elems_written_per_iteration_x = 4; + num_elems_read_per_iteration_y = 6; + num_elems_written_per_iteration_y = 4; + } + else if(_conv_stride_x == 2 && _conv_stride_y == 2) + { + kernel_name = "depthwise_convolution_3x3_stridex2_stridey2_bifrost_f16"; + num_elems_read_per_iteration_x = 10; + num_elems_written_per_iteration_x = 4; + num_elems_read_per_iteration_y = 5; + num_elems_written_per_iteration_y = 2; + } + } } - else if(input->info()->data_type() == DataType::F32 && gpu_target_is_in(gpu_target, GPUTarget::G71, GPUTarget::G72)) + else if(input->info()->data_type() == DataType::F32 && is_bifrost) { if(_conv_stride_x == 1 && _conv_stride_y == 1) { - kernel_name = "depthwise_convolution_3x3_stridex1_stridey1_bifrost"; + kernel_name = "depthwise_convolution_3x3_stridex1_stridey1_bifrost_f32"; num_elems_read_per_iteration_x = 4; num_elems_read_per_iteration_y = 6; num_elems_written_per_iteration_x = 2; @@ -191,7 +211,7 @@ void CLDepthwiseConvolutionLayer3x3Kernel::configure(const ICLTensor *input, con } else if(_conv_stride_x == 2 && _conv_stride_y == 2) { - kernel_name = "depthwise_convolution_3x3_stridex2_stridey2_bifrost"; + kernel_name = "depthwise_convolution_3x3_stridex2_stridey2_bifrost_f32"; num_elems_read_per_iteration_x = 6; num_elems_read_per_iteration_y = 5; num_elems_written_per_iteration_x = 2; -- cgit v1.2.1