Skip to content

Add a Box3x3 tiling kernel and tiling library #12

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
ccaolii wants to merge 12 commits into
base: openvx_1.2
Choose a base branch
from
Open

Add a Box3x3 tiling kernel and tiling library #12

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
12 commits
Select commit Hold shift + click to select a range
d56899d
Add a Box3x3 tiling kernel and tiling library
Jiaqing-Cao Jul 17, 2019
e8106a2
Add Neon intrinsics in Box3x3 tiling kernel
Jiaqing-Cao Jul 19, 2019
93c8294
Add Phase, And, Or, Xor, and Not Tiling kernels
Jiaqing-Cao Aug 13, 2019
fcba88c
Add Multiply, Threshold and ColorConvert Tiling kernels
Jiaqing-Cao Aug 20, 2019
76449ff
Add NonLinear, Magnitude and OpenCL Warp and some other kernels
Jiaqing-Cao Aug 22, 2019
356154d
Add Convolve U8 OpenCL kernel
Jiaqing-Cao Aug 27, 2019
d94b210
Add Erode3x3, Dilate3x3 and Median3x3 Tiling kernels, and add validat…
Jiaqing-Cao Aug 28, 2019
26a7c93
Add add, sub, min, max, sobel3x3 tiling kernels
Jiaqing-Cao Sep 2, 2019
ef9b639
Add ConvertDepth, WeightedAverage, WarpAffine and WarpPerspective Til…
Jiaqing-Cao Sep 10, 2019
2298399
Add AbsDiff and IntergralImage Tiling kernel
Jiaqing-Cao Sep 17, 2019
b09036d
Add Remap and Convolve Tiling kernel
Jiaqing-Cao Sep 23, 2019
580734c
Add Hog, FastCorners, LBP, LUT, Scale, HalfScaleGaussian, ChannelComb…
Jiaqing-Cao Oct 8, 2019
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
9 changes: 8 additions & 1 deletion examples/vx_tiling_ext.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -220,7 +220,8 @@ static vx_status VX_CALLBACK vxAlphaOutputValidator(vx_node node, vx_uint32 inde
return status;
}


//Move this struct into "include/VX/vx_khr_tiling.h"
#if 0
/*! [publish_support] */
typedef struct _vx_tiling_kernel_t {
/*! kernel name */
Expand All @@ -246,6 +247,7 @@ typedef struct _vx_tiling_kernel_t {
/*! border information. */
vx_border_t border;
} vx_tiling_kernel_t;
#endif

static vx_tiling_kernel_t tiling_kernels[] = {
{"org.khronos.openvx.tiling_gaussian_3x3",
Expand All @@ -255,6 +257,7 @@ static vx_tiling_kernel_t tiling_kernels[] = {
2,
{{VX_INPUT, VX_TYPE_IMAGE, VX_PARAMETER_STATE_REQUIRED},
{VX_OUTPUT, VX_TYPE_IMAGE, VX_PARAMETER_STATE_REQUIRED}},
NULL,
vxFilterInputValidator,
vxFilterOutputValidator,
{1, 1},
Expand All @@ -269,6 +272,7 @@ static vx_tiling_kernel_t tiling_kernels[] = {
{{VX_INPUT, VX_TYPE_IMAGE, VX_PARAMETER_STATE_REQUIRED},
{VX_INPUT, VX_TYPE_SCALAR, VX_PARAMETER_STATE_REQUIRED},
{VX_OUTPUT, VX_TYPE_IMAGE, VX_PARAMETER_STATE_REQUIRED}},
NULL,
vxAlphaInputValidator,
vxAlphaOutputValidator,
{1, 1},
Expand All @@ -282,6 +286,7 @@ static vx_tiling_kernel_t tiling_kernels[] = {
2,
{{VX_INPUT, VX_TYPE_IMAGE, VX_PARAMETER_STATE_REQUIRED},
{VX_OUTPUT, VX_TYPE_IMAGE, VX_PARAMETER_STATE_REQUIRED}},
NULL,
vxFilterInputValidator,
vxFilterOutputValidator,
{1, 1},
Expand All @@ -296,6 +301,7 @@ static vx_tiling_kernel_t tiling_kernels[] = {
{{VX_INPUT, VX_TYPE_IMAGE, VX_PARAMETER_STATE_REQUIRED},
{VX_INPUT, VX_TYPE_IMAGE, VX_PARAMETER_STATE_REQUIRED},
{VX_OUTPUT, VX_TYPE_IMAGE, VX_PARAMETER_STATE_REQUIRED}},
NULL,
vxAddInputValidator,
vxAddOutputValidator,
{1, 1},
Expand All @@ -319,6 +325,7 @@ VX_API_ENTRY vx_status VX_API_CALL vxPublishKernels(vx_context context)
tiling_kernels[k].flexible_function,
tiling_kernels[k].fast_function,
tiling_kernels[k].num_params,
tiling_kernels[k].validate,
tiling_kernels[k].input_validator,
tiling_kernels[k].output_validator);
if (kernel)
Expand Down
3 changes: 3 additions & 0 deletions kernels/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,4 +19,7 @@
add_subdirectory( c_model )
add_subdirectory( debug )
add_subdirectory( extras )
if (OPENVX_USE_TILING)
add_subdirectory( tiling )
endif (OPENVX_USE_TILING)

10 changes: 10 additions & 0 deletions kernels/opencl/vx_and.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,10 @@

__kernel void vx_and(int asx, int asy, __global uchar *a,
int bsx, int bsy, __global uchar *b,
int csx, int csy, __global uchar *c)
{
int x = get_global_id(0);
int y = get_global_id(1);

c[y * csy + x * csx] = a[y * asy + x * asx] & b[y * bsy + x * bsx];
}
87 changes: 87 additions & 0 deletions kernels/opencl/vx_box3x3.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,87 @@

//Define 3 types of border
#define VX_ID_KHRONOS 0x000
#define VX_ENUM_BORDER 0x0C
#define VX_ENUM_BASE(vendor, id) (((vendor) << 20) | (id << 12))

#define VX_BORDER_UNDEFINED VX_ENUM_BASE(VX_ID_KHRONOS, VX_ENUM_BORDER) + 0x0
#define VX_BORDER_CONSTANT VX_ENUM_BASE(VX_ID_KHRONOS, VX_ENUM_BORDER) + 0x1
#define VX_BORDER_REPLICATE VX_ENUM_BASE(VX_ID_KHRONOS, VX_ENUM_BORDER) + 0x2

#define BOX3x3 sum += (uint)src[x_top * ssx + y_top * ssy]; \
sum += (uint)src[x * ssx + y_top * ssy]; \
sum += (uint)src[x_bot * ssx + y_top * ssy]; \
sum += (uint)src[x_top * ssx + y * ssy]; \
sum += (uint)src[x * ssx + y * ssy]; \
sum += (uint)src[x_bot * ssx + y * ssy]; \
sum += (uint)src[x_top * ssx + y_bot * ssy]; \
sum += (uint)src[x * ssx + y_bot * ssy]; \
sum += (uint)src[x_bot * ssx + y_bot * ssy]; \
sum = sum / 9; \
dst[x * dsx + y * dsy] = (uchar)sum; \


__kernel void vx_box3x3(int ssx, int ssy, __global uchar *src,
int bordermode, uchar const_vaule,
int dsx, int dsy, __global uchar *dst)
{
const int x = get_global_id(0);
const int y = get_global_id(1);
const size_t high_x = get_global_size(0);
const size_t high_y = get_global_size(1);
uint sum = 0;

int y_top = y - 1;
int y_bot = y + 1;
int x_top = x - 1;
int x_bot = x + 1;

int ky, kx;
uint dest_index = 0;

if (bordermode == VX_BORDER_CONSTANT)
{
uchar pixel[9];
// Calculate border
if (y == 0 || x == 0 || x == high_x - 1 || y == high_y - 1)
{
for (ky = -1; ky <= 1; ++ky)
{
int yy = y + ky;
int ccase_y = yy < 0 || yy >= high_y;

for (kx = -1; kx <= 1; ++kx, ++dest_index)
{
int xx = x + kx;
int ccase = ccase_y || xx < 0 || xx >= high_x;

if (!ccase)
pixel[dest_index] = src[xx * ssx + yy * ssy];
else
pixel[dest_index] = const_vaule;
}
}

sum = pixel[0] + pixel[1] + pixel[2] + pixel[3] + pixel[4] + pixel[5] + pixel[6] + pixel[7] + pixel[8];

sum = sum / 9;
dst[x * dsx + y * dsy] = (uchar)sum;
}
else
{
BOX3x3;
}
}
else
{
if (bordermode == VX_BORDER_REPLICATE)
{
y_top = y_top < 0 ? 0 : y - 1;
y_bot = y_bot >= high_y ? high_y - 1 : y + 1;
x_top = x_top < 0 ? 0 : x - 1;
x_bot = x_bot >= high_x ? high_x - 1 : x + 1;
}

BOX3x3;
}
}
93 changes: 93 additions & 0 deletions kernels/opencl/vx_convolve.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,93 @@

//Define 3 types of border
#define VX_ID_KHRONOS 0x000
#define VX_ENUM_BORDER 0x0C
#define VX_ENUM_BASE(vendor, id) (((vendor) << 20) | (id << 12))

#define VX_BORDER_UNDEFINED VX_ENUM_BASE(VX_ID_KHRONOS, VX_ENUM_BORDER) + 0x0
#define VX_BORDER_CONSTANT VX_ENUM_BASE(VX_ID_KHRONOS, VX_ENUM_BORDER) + 0x1
#define VX_BORDER_REPLICATE VX_ENUM_BASE(VX_ID_KHRONOS, VX_ENUM_BORDER) + 0x2

#define C_MAX_CONVOLUTION_DIM (15)
#define UINT8_MAX 255

#define Convolve \
uchar slice[C_MAX_CONVOLUTION_DIM * C_MAX_CONVOLUTION_DIM] = { 0 }; \
uint center_x = x, center_y = y; \
int width = high_x, height = high_y; \
int ky, kx; \
uint dest_index = 0; \
\
if( bordermode == VX_BORDER_REPLICATE || bordermode == VX_BORDER_UNDEFINED ) \
{ \
for (ky = -(int)conv_radius_y; ky <= (int)conv_radius_y; ++ky) \
{ \
int yy = (int)(center_y + ky); \
yy = yy < 0 ? 0 : yy >= height ? height - 1 : yy; \
\
for (kx = -(int)conv_radius_x; kx <= (int)conv_radius_x; ++kx, ++dest_index) \
{ \
int xx = (int)(center_x + kx); \
xx = xx < 0 ? 0 : xx >= width ? width - 1 : xx; \
slice[dest_index] = src[xx * ssx + yy * ssy]; \
} \
} \
} \
else if( bordermode == VX_BORDER_CONSTANT ) \
{ \
for (ky = -(int)conv_radius_y; ky <= (int)conv_radius_y; ++ky) \
{ \
int yy = (int)(center_y + ky); \
int ccase_y = yy < 0 || yy >= height; \
\
for (kx = -(int)conv_radius_x; kx <= (int)conv_radius_x; ++kx, ++dest_index) \
{ \
int xx = (int)(center_x + kx); \
int ccase = ccase_y || xx < 0 || xx >= width; \
if( !ccase ) \
slice[dest_index] = src[xx * ssx + yy * ssy]; \
else \
slice[dest_index] = (uchar)const_vaule; \
} \
} \
} \
\
for (int i = 0; i < (int)(conv_width * conv_height); ++i) \
sum += conv_mat[conv_width * conv_height - 1 - i] * slice[i]; \
\
value = sum / (int)scale; \
\
if (value < 0) dst[x * dsx + y * dsy] = 0; \
else if (value > UINT8_MAX) dst[x * dsx + y * dsy] = UINT8_MAX; \
else dst[x * dsx + y * dsy] = value;

__kernel void vx_Convolve(int ssx, int ssy, __global uchar *src,
int bordermode, uchar const_vaule,
uint conv_width, uint conv_height,
uint scale, __global short *conv_mat,
int dsx, int dsy, __global uchar *dst)
{
const int x = get_global_id(0);
const int y = get_global_id(1);

int low_x = 0, low_y = 0;
int high_x = get_global_size(0);
int high_y = get_global_size(1);
int sum = 0;
int value = 0;

int conv_radius_x, conv_radius_y;
conv_radius_x = (int)conv_width / 2;
conv_radius_y = (int)conv_height / 2;

if (bordermode == VX_BORDER_UNDEFINED)
{
low_x = conv_radius_x;
high_x = ((high_x >= (uint)conv_radius_x) ? high_x - conv_radius_x : 0);
low_y = conv_radius_y;
high_y = ((high_y >= (uint)conv_radius_y) ? high_y - conv_radius_y : 0);
}

Convolve;

}
93 changes: 93 additions & 0 deletions kernels/opencl/vx_dilate3x3.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,93 @@
//Define 3 types of border
#define VX_ID_KHRONOS 0x000
#define VX_ENUM_BORDER 0x0C
#define VX_ENUM_BASE(vendor, id) (((vendor) << 20) | (id << 12))

#define VX_BORDER_UNDEFINED VX_ENUM_BASE(VX_ID_KHRONOS, VX_ENUM_BORDER) + 0x0
#define VX_BORDER_CONSTANT VX_ENUM_BASE(VX_ID_KHRONOS, VX_ENUM_BORDER) + 0x1
#define VX_BORDER_REPLICATE VX_ENUM_BASE(VX_ID_KHRONOS, VX_ENUM_BORDER) + 0x2

uchar max_op(uchar a, uchar b)
{
return a > b ? a : b;
}

#define DILATE3x3 pixels[0] = src[x_top * ssx + y_top * ssy]; \
pixels[1] = src[x * ssx + y_top * ssy]; \
pixels[2] = src[x_bot * ssx + y_top * ssy]; \
pixels[3] = src[x_top * ssx + y * ssy]; \
pixels[4] = src[x * ssx + y * ssy]; \
pixels[5] = src[x_bot * ssx + y * ssy]; \
pixels[6] = src[x_top * ssx + y_bot * ssy]; \
pixels[7] = src[x * ssx + y_bot * ssy]; \
pixels[8] = src[x_bot * ssx + y_bot * ssy]; \
max_value = pixels[0]; \
for (i = 1; i < 9; i++) \
max_value = max_op(max_value, pixels[i]); \
dst[x * dsx + y * dsy] = max_value; \

__kernel void vx_dilate3x3(int ssx, int ssy, __global uchar *src,
int bordermode, uchar const_vaule,
int dsx, int dsy, __global uchar *dst)
{
const int x = get_global_id(0);
const int y = get_global_id(1);
const size_t high_x = get_global_size(0);
const size_t high_y = get_global_size(1);
uint sum = 0;

int y_top = y - 1;
int y_bot = y + 1;
int x_top = x - 1;
int x_bot = x + 1;

int ky, kx, i;
uint dest_index = 0;
uchar pixels[9], max_value;

if (bordermode == VX_BORDER_CONSTANT)
{
// Calculate border
if (y == 0 || x == 0 || x == high_x - 1 || y == high_y - 1)
{
for (ky = -1; ky <= 1; ++ky)
{
int yy = y + ky;
int ccase_y = yy < 0 || yy >= high_y;

for (kx = -1; kx <= 1; ++kx, ++dest_index)
{
int xx = x + kx;
int ccase = ccase_y || xx < 0 || xx >= high_x;

if (!ccase)
pixels[dest_index] = src[xx * ssx + yy * ssy];
else
pixels[dest_index] = const_vaule;
}
}

max_value = pixels[0];
for (i = 1; i < 9; i++)
max_value = max_op(max_value, pixels[i]);

dst[x * dsx + y * dsy] = max_value;
}
else
{
DILATE3x3;
}
}
else
{
if (bordermode == VX_BORDER_REPLICATE)
{
y_top = y_top < 0 ? 0 : y - 1;
y_bot = y_bot >= high_y ? high_y - 1 : y + 1;
x_top = x_top < 0 ? 0 : x - 1;
x_bot = x_bot >= high_x ? high_x - 1 : x + 1;
}

DILATE3x3;
}
}
Loading