diff options
| author | Michele Di Giorgio <michele.digiorgio@arm.com> | 2018-05-08 17:23:52 +0100 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:51:17 +0000 |
| commit | d24af8a3e8f7cbd38fd3142056241c0c9f63e46a (patch) | |
| tree | a6d90561bbc1b976708d120a576bf222059dd36b /src/core/CL/cl_kernels | |
| parent | 932b561159cd6a8c9230bbd0343790c85755846e (diff) | |
| download | ComputeLibrary-d24af8a3e8f7cbd38fd3142056241c0c9f63e46a.tar.gz | |
COMPMID-1125: Add support for FP16 in CLDepthwiseConvolution
Change-Id: I4838f5a8e4c33ed646cd05e0bb682fca635a29a3
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/130469
Tested-by: Jenkins <bsgcomp@arm.com>
Reviewed-by: Michalis Spyrou <michalis.spyrou@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels')
| -rw-r--r-- | src/core/CL/cl_kernels/depthwise_convolution.cl | 2 | ||||
| -rw-r--r-- | src/core/CL/cl_kernels/fill_border.cl | 6 |
2 files changed, 4 insertions, 4 deletions
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) |