From d24af8a3e8f7cbd38fd3142056241c0c9f63e46a Mon Sep 17 00:00:00 2001 From: Michele Di Giorgio Date: Tue, 8 May 2018 17:23:52 +0100 Subject: COMPMID-1125: Add support for FP16 in CLDepthwiseConvolution Change-Id: I4838f5a8e4c33ed646cd05e0bb682fca635a29a3 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/130469 Tested-by: Jenkins Reviewed-by: Michalis Spyrou --- src/core/CL/cl_kernels/depthwise_convolution.cl | 2 +- src/core/CL/cl_kernels/fill_border.cl | 6 +++--- src/runtime/CL/functions/CLDepthwiseConvolutionLayer.cpp | 2 +- 3 files changed, 5 insertions(+), 5 deletions(-) (limited to 'src') diff --git a/src/core/CL/cl_kernels/depthwise_convolution.cl b/src/core/CL/cl_kernels/depthwise_convolution.cl index 21c28539ef..5f4247e5d3 100644 --- a/src/core/CL/cl_kernels/depthwise_convolution.cl +++ b/src/core/CL/cl_kernels/depthwise_convolution.cl @@ -500,7 +500,7 @@ __kernel void depthwise_weights_reshape( #if defined(HAS_BIAS) if(get_global_id(1) == 0) { - *((__global DATA_TYPE *)(output_ptr + SRC_WIDTH * get_global_size(1) * dst_stride_x)) = *((__global float *)(biases.ptr + get_global_id(2) * biases_stride_x)); + *((__global DATA_TYPE *)(output_ptr + SRC_WIDTH * get_global_size(1) * dst_stride_x)) = *((__global DATA_TYPE *)(biases.ptr + get_global_id(2) * biases_stride_x)); } #endif // defined(HAS_BIAS) } diff --git a/src/core/CL/cl_kernels/fill_border.cl b/src/core/CL/cl_kernels/fill_border.cl index fbd4f6ae0f..33a9495d66 100644 --- a/src/core/CL/cl_kernels/fill_border.cl +++ b/src/core/CL/cl_kernels/fill_border.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016, 2017 ARM Limited. + * Copyright (c) 2016-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -35,7 +35,7 @@ * @attention The border size for top, bottom, left, right needs to be passed at the compile time. * e.g. --DBORDER_SIZE_TOP=0 -DBORDER_SIZE_BOTTOM=2 -DBORDER_SIZE_LEFT=0 -DBORDER_SIZE_RIGHT=2 * - * @param[in,out] buf_ptr Pointer to the source image. Supported data types: U8, U16, S16, U32, S32, F32 + * @param[in,out] buf_ptr Pointer to the source image. Supported data types: U8/U16/S16/U32/S32/F16/F32 * @param[in] buf_stride_x Stride of the source image in X dimension (in bytes) * @param[in] buf_step_x buf_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] buf_stride_y Stride of the source image in Y dimension (in bytes) @@ -110,7 +110,7 @@ __kernel void fill_image_borders_replicate( * @attention The border size for top, bottom, left, right needs to be passed at the compile time. * e.g. --DBORDER_SIZE_TOP=0 -DBORDER_SIZE_BOTTOM=2 -DBORDER_SIZE_LEFT=0 -DBORDER_SIZE_RIGHT=2 * - * @param[out] buf_ptr Pointer to the source image. Supported data types: U8, U16, S16, U32, S32, F32 + * @param[out] buf_ptr Pointer to the source image. Supported data types: U8/U16/S16/U32/S32/F16/F32 * @param[in] buf_stride_x Stride of the source image in X dimension (in bytes) * @param[in] buf_step_x buf_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] buf_stride_y Stride of the source image in Y dimension (in bytes) diff --git a/src/runtime/CL/functions/CLDepthwiseConvolutionLayer.cpp b/src/runtime/CL/functions/CLDepthwiseConvolutionLayer.cpp index ea2f93b85d..88bb0c417d 100644 --- a/src/runtime/CL/functions/CLDepthwiseConvolutionLayer.cpp +++ b/src/runtime/CL/functions/CLDepthwiseConvolutionLayer.cpp @@ -80,7 +80,7 @@ CLDepthwiseConvolutionLayer::CLDepthwiseConvolutionLayer() void CLDepthwiseConvolutionLayer::configure(ICLTensor *input, const ICLTensor *weights, const ICLTensor *biases, ICLTensor *output, const PadStrideInfo &conv_info, unsigned int depth_multiplier) { - ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::QASYMM8, DataType::F32); + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::QASYMM8, DataType::F16, DataType::F32); ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input, weights); ARM_COMPUTE_ERROR_ON((input->info()->dimension(2) * depth_multiplier) != weights->info()->dimension(2)); -- cgit v1.2.1