From 0a887922c73bbe7c5d42b1eb3ae55730f0d9a139 Mon Sep 17 00:00:00 2001 From: Michalis Spyrou Date: Mon, 11 Jun 2018 16:30:23 +0100 Subject: COMPMID-1222 Implementing CLArithmeticDivision - FP32 / FP16 Change-Id: I2e3f725ef5ed1454755086b9640ab84a81f4d40e Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/135170 Reviewed-by: Anthony Barbier Tested-by: Jenkins --- src/core/CL/cl_kernels/arithmetic_op.cl | 56 +++++++++++++++++++++++++++++++-- 1 file changed, 54 insertions(+), 2 deletions(-) (limited to 'src/core/CL/cl_kernels/arithmetic_op.cl') diff --git a/src/core/CL/cl_kernels/arithmetic_op.cl b/src/core/CL/cl_kernels/arithmetic_op.cl index 12963473c5..8bd28230b7 100644 --- a/src/core/CL/cl_kernels/arithmetic_op.cl +++ b/src/core/CL/cl_kernels/arithmetic_op.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016, 2017 ARM Limited. + * Copyright (c) 2016-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -35,6 +35,8 @@ #define SUB(x, y) (x) - (y) #endif /* SATURATE */ +#define DIV(x, y) (x) / (y) + /** This function adds two tensors. * * @attention The input and output data_types need to be passed at compile time using -DDATA_TYPE_IN1, -DDATA_TYPE_IN2 and -DDATA_TYPE_OUT: @@ -86,7 +88,7 @@ __kernel void arithmetic_add( vstore16(ADD(in_a, in_b), 0, (__global DATA_TYPE_OUT *)out.ptr); } -/** This function subtracts one tensors from another. +/** This function subtracts one tensor from another. * * @attention The input and output data_types need to be passed at compile time using -DDATA_TYPE_IN1, -DDATA_TYPE_IN2 and -DDATA_TYPE_OUT: * e.g. -DDATA_TYPE_IN1=uchar -DDATA_TYPE_IN2=uchar -DDATA_TYPE_OUT=short @@ -136,3 +138,53 @@ __kernel void arithmetic_sub( // Calculate and store result vstore16(SUB(in_a, in_b), 0, (__global DATA_TYPE_OUT *)out.ptr); } + +/** This function divides one tensor from another. + * + * @attention The input and output data_types need to be passed at compile time using -DDATA_TYPE_IN1, -DDATA_TYPE_IN2 and -DDATA_TYPE_OUT: + * e.g. -DDATA_TYPE_IN1=float -DDATA_TYPE_IN2=float -DDATA_TYPE_OUT=float + * + * @param[in] in1_ptr Pointer to the source tensor. Supported data types: F16/F32 + * @param[in] in1_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in] in1_step_x in1_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] in1_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] in1_step_y in1_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] in1_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] in1_step_z in1_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] in1_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[in] in2_ptr Pointer to the source tensor. Supported data types: Same as @p in1_ptr + * @param[in] in2_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in] in2_step_x in2_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] in2_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] in2_step_y in2_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] in2_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] in2_step_z in2_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] in2_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[out] out_ptr Pointer to the destination tensor. Supported data types: Same as @p in1_ptr + * @param[in] out_stride_x Stride of the destination tensor in X dimension (in bytes) + * @param[in] out_step_x out_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] out_stride_y Stride of the destination tensor in Y dimension (in bytes) + * @param[in] out_step_y out_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] out_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] out_step_z out_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] out_offset_first_element_in_bytes The offset of the first element in the destination tensor + */ +__kernel void arithmetic_div( + TENSOR3D_DECLARATION(in1), + TENSOR3D_DECLARATION(in2), + TENSOR3D_DECLARATION(out)) +{ + // Get pixels pointer + Tensor3D in1 = CONVERT_TO_TENSOR3D_STRUCT(in1); + Tensor3D in2 = CONVERT_TO_TENSOR3D_STRUCT(in2); + Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(out); + + // Load values + VEC_DATA_TYPE(DATA_TYPE_OUT, 16) + in_a = CONVERT(vload16(0, (__global DATA_TYPE_IN1 *)in1.ptr), VEC_DATA_TYPE(DATA_TYPE_OUT, 16)); + VEC_DATA_TYPE(DATA_TYPE_OUT, 16) + in_b = CONVERT(vload16(0, (__global DATA_TYPE_IN2 *)in2.ptr), VEC_DATA_TYPE(DATA_TYPE_OUT, 16)); + + // Calculate and store result + vstore16(DIV(in_a, in_b), 0, (__global DATA_TYPE_OUT *)out.ptr); +} -- cgit v1.2.1