derivative.cl File Reference

#include "helpers.h"

Go to the source code of this file.

Functions
__kernel void	derivative (__global uchar *src_ptr, uint src_stride_x, uint src_step_x, uint src_stride_y, uint src_step_y, uint src_offset_first_element_in_bytes, __global uchar *dst_gx_ptr, uint dst_gx_stride_x, uint dst_gx_step_x, uint dst_gx_stride_y, uint dst_gx_step_y, uint dst_gx_offset_first_element_in_bytes, __global uchar *dst_gy_ptr, uint dst_gy_stride_x, uint dst_gy_step_x, uint dst_gy_stride_y, uint dst_gy_step_y, uint dst_gy_offset_first_element_in_bytes)
	This OpenCL kernel that computes the first-order derivative. More...

Function Documentation

__kernel void derivative	(	__global uchar *	src_ptr,
		uint	src_stride_x,
		uint	src_step_x,
		uint	src_stride_y,
		uint	src_step_y,
		uint	src_offset_first_element_in_bytes,
		__global uchar *	dst_gx_ptr,
		uint	dst_gx_stride_x,
		uint	dst_gx_step_x,
		uint	dst_gx_stride_y,
		uint	dst_gx_step_y,
		uint	dst_gx_offset_first_element_in_bytes,
		__global uchar *	dst_gy_ptr,
		uint	dst_gy_stride_x,
		uint	dst_gy_step_x,
		uint	dst_gy_stride_y,
		uint	dst_gy_step_y,
		uint	dst_gy_offset_first_element_in_bytes
	)

This OpenCL kernel that computes the first-order derivative.

Attention

To enable computation of the X gradient -DGRAD_X must be passed at compile time, while computation of the Y gradient is performed when -DGRAD_Y is used. You can use both when computation of both gradients is required.

Parameters

[in]	src_ptr	Pointer to the source image. Supported data types: U8
[in]	src_stride_x	Stride of the source image in X dimension (in bytes)
[in]	src_step_x	src_stride_x * number of elements along X processed per workitem(in bytes)
[in]	src_stride_y	Stride of the source image in Y dimension (in bytes)
[in]	src_step_y	src_stride_y * number of elements along Y processed per workitem(in bytes)
[in]	src_offset_first_element_in_bytes	The offset of the first element in the source image
[out]	dst_gx_ptr	Pointer to the destination image. Supported data types: S16
[in]	dst_gx_stride_x	Stride of the destination image in X dimension (in bytes)
[in]	dst_gx_step_x	dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
[in]	dst_gx_stride_y	Stride of the destination image in Y dimension (in bytes)
[in]	dst_gx_step_y	dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
[in]	dst_gx_offset_first_element_in_bytes	The offset of the first element in the destination image
[out]	dst_gy_ptr	Pointer to the destination image. Supported data types: S16
[in]	dst_gy_stride_x	Stride of the destination image in X dimension (in bytes)
[in]	dst_gy_step_x	dst_gy_stride_x * number of elements along X processed per workitem(in bytes)
[in]	dst_gy_stride_y	Stride of the destination image in Y dimension (in bytes)
[in]	dst_gy_step_y	dst_gy_stride_y * number of elements along Y processed per workitem(in bytes)
[in]	dst_gy_offset_first_element_in_bytes	The offset of the first element in the destination image

Definition at line 50 of file derivative.cl.

References CONVERT_TO_IMAGE_STRUCT, offset(), Image::ptr, and arm_compute::test::validation::src.

Referenced by arm_compute::test::validation::DATA_TEST_CASE(), and DerivativeValidationFixture< TensorType, AccessorType, FunctionType, T, U >::setup().

{
    Image src = CONVERT_TO_IMAGE_STRUCT(src);
#ifdef GRAD_X
    Image dst_gx = CONVERT_TO_IMAGE_STRUCT(dst_gx);
#endif /* GRAD_X */
#ifdef GRAD_Y
    Image dst_gy = CONVERT_TO_IMAGE_STRUCT(dst_gy);
#endif /* GRAD_Y */

#ifdef GRAD_X
    short16 l_data = convert_short16(vload16(0, offset(&src, -1, 0)));
    short16 r_data = convert_short16(vload16(0, offset(&src, 1, 0)));
    vstore16(r_data - l_data, 0, ((__global short *)dst_gx.ptr));
#endif /* GRAD_X */
#ifdef GRAD_Y
    short16 t_data = convert_short16(vload16(0, offset(&src, 0, -1)));
    short16 b_data = convert_short16(vload16(0, offset(&src, 0, 1)));
    vstore16(b_data - t_data, 0, ((__global short *)dst_gy.ptr));
#endif /* GRAD_Y */
}