From ff25226adbcae4e94cafff65f5fa405a8702bae9 Mon Sep 17 00:00:00 2001 From: Sven Date: Tue, 11 Jan 2022 11:45:29 +0800 Subject: [PATCH] [Internal] support prebuilt kernel into shared library (#260) Signed-off-by: xiang.zhang --- .../vx/internal/src/kernel/vsi_nn_kernel.c | 66 +++++++++++++++---- .../libnnext/ops/vx/conv1d_ovxlib_k1024.vx | 8 +-- .../libnnext/ops/vx/l2normalizescale_axis0.vx | 2 +- .../libnnext/ops/vx/vsi_nn_kernel_header.vx | 2 +- .../internal/src/libnnext/vsi_nn_vxkernel.c | 17 +---- 5 files changed, 61 insertions(+), 34 deletions(-) diff --git a/src/tim/vx/internal/src/kernel/vsi_nn_kernel.c b/src/tim/vx/internal/src/kernel/vsi_nn_kernel.c index b266a99..02526f5 100644 --- a/src/tim/vx/internal/src/kernel/vsi_nn_kernel.c +++ b/src/tim/vx/internal/src/kernel/vsi_nn_kernel.c @@ -79,10 +79,11 @@ static vx_program _create_program_from_code vsi_nn_kernel_t* kernel ); -static const void* _load_internal_executable +static const uint8_t* _load_internal_executable ( const char* source_name, - size_t* size + size_t* size, + vsi_nn_kernel_type_e type ); static char* _load_source_code_from_file @@ -216,21 +217,62 @@ static vsi_status _cpu_register return status; } /* _cpu_register() */ -static const void* _load_internal_executable +#if VSI_USE_VXC_BINARY +static const uint8_t* _load_bin ( const char* source_name, - size_t* size + size_t* size, + const vsi_nn_vx_bin_resource_item_type* source_map, + size_t source_map_size, + const char* tail + ) +{ + const uint8_t* source; + char source_path[VSI_NN_MAX_PATH]; + size_t n; + int i; + source = NULL; + n = snprintf( source_path, VSI_NN_MAX_PATH, "%s%s", source_name, tail ); + if( n == VSI_NN_MAX_PATH ) + { + VSILOGE("Kernel source path overflow %d/%d", n, VSI_NN_MAX_PATH); + *size = 0; + return NULL; + } + for( i = 0; i < (int)source_map_size; i++ ) + { + if( strncmp( source_map[i].name, source_path, VSI_NN_MAX_PATH ) == 0 ) + { + source = source_map[i].data; + *size = source_map[i].len; + break; + } + } + if( !source ) + { + *size = 0; + } + return source; +} /* _load_bin() */ +#endif + +static const uint8_t* _load_internal_executable + ( + const char* source_name, + size_t* size, + vsi_nn_kernel_type_e type ) { #if VSI_USE_VXC_BINARY - int i; - for( i = 0; i < vx_bin_resource_items_cnt; i++ ) + switch( type ) { - if( strncmp( vx_bin_resource_items[i].name, source_name, VSI_NN_MAX_PATH ) == 0 ) - { - *size = (size_t)vx_bin_resource_items[i].len; - return vx_bin_resource_items[i].data; - } + case VSI_NN_KERNEL_TYPE_EVIS: + return _load_bin( source_name, size, + vx_bin_resource_items_vx, vx_bin_resource_items_vx_cnt, "_vx" ); + break; + case VSI_NN_KERNEL_TYPE_CL: + return _load_bin( source_name, size, + vx_bin_resource_items_cl, vx_bin_resource_items_cl_cnt, "_cl" ); } #endif return NULL; @@ -393,7 +435,7 @@ static vx_program _create_program_from_executable memset( &program_info, 0, sizeof( kernel_program_info_t ) ); program_info.data = _load_internal_executable( - source_info->data[0], &program_info.size); + source_info->data[0], &program_info.size, kernel->type); program = vxCreateProgramWithBinary( graph->ctx->c, (const vx_uint8 *)program_info.data, program_info.size ); return program; diff --git a/src/tim/vx/internal/src/libnnext/ops/vx/conv1d_ovxlib_k1024.vx b/src/tim/vx/internal/src/libnnext/ops/vx/conv1d_ovxlib_k1024.vx index dc2497e..148f69f 100644 --- a/src/tim/vx/internal/src/libnnext/ops/vx/conv1d_ovxlib_k1024.vx +++ b/src/tim/vx/internal/src/libnnext/ops/vx/conv1d_ovxlib_k1024.vx @@ -104,10 +104,10 @@ __kernel void conv1d_U8U8I32toU8_K1024_LARGE( vxc_short8 w_zp = (short)weight_ZP; vxc_uchar16 input_val = 0, weight_val = 0; int temp = 0, i, j; - Tensor src_tensor = create_image_from_image2d(input, 1); - uchar *src_ptr_base = (uchar *)src_image.ptr; + Tensor src_tensor = create_tensor_from_image2d_array(input, 1); + uchar *src_ptr_base = (uchar *)src_tensor.ptr; uchar *src_ptr; - Tensor dst_tensor = create_image_from_image2d(output, 1); + Tensor dst_tensor = create_tensor_from_image2d_array(output, 1); uchar *dst_ptr = (uchar *)dst_tensor.ptr; temp = read_imagei(bias, coord.yz).x; @@ -116,7 +116,7 @@ __kernel void conv1d_U8U8I32toU8_K1024_LARGE( for (i = 0; i < input_height; i++) { - src_ptr = src_ptr_base + (coord.x + coord.z * src_image.stride_y); + src_ptr = src_ptr_base + (coord.x + coord.z * src_tensor.stride_y); for (j = 0; j < kernel_cnt_x16; j++) { VXC_ReadImage2DArray(weight_val, weight, coord_w, VXC_5BITOFFSET_XY(0, 0), \ diff --git a/src/tim/vx/internal/src/libnnext/ops/vx/l2normalizescale_axis0.vx b/src/tim/vx/internal/src/libnnext/ops/vx/l2normalizescale_axis0.vx index c358585..1221ed1 100644 --- a/src/tim/vx/internal/src/libnnext/ops/vx/l2normalizescale_axis0.vx +++ b/src/tim/vx/internal/src/libnnext/ops/vx/l2normalizescale_axis0.vx @@ -262,7 +262,7 @@ __kernel __attribute__((reqd_work_group_size(16, 1, 1))) \ { \ int lidx = get_local_id(0); \ int offset = get_global_id(0); \ - Image src_img = create_image_from_image2d(input, 1); + Image src_img = create_image_from_image2d(input, 1); \ uchar *src_ptr_base = (uchar *)src_img.ptr; \ uchar *src_ptr; \ vxc_uchar8 src0, src1; \ diff --git a/src/tim/vx/internal/src/libnnext/ops/vx/vsi_nn_kernel_header.vx b/src/tim/vx/internal/src/libnnext/ops/vx/vsi_nn_kernel_header.vx index f6ccacc..bd97b11 100644 --- a/src/tim/vx/internal/src/libnnext/ops/vx/vsi_nn_kernel_header.vx +++ b/src/tim/vx/internal/src/libnnext/ops/vx/vsi_nn_kernel_header.vx @@ -7,7 +7,7 @@ Description : ============================================================================ */ -#include "cl_viv_vx_ext.h" +#pragma OPENCL EXTENSION cl_viv_vx_extension : enable typedef struct Image { diff --git a/src/tim/vx/internal/src/libnnext/vsi_nn_vxkernel.c b/src/tim/vx/internal/src/libnnext/vsi_nn_vxkernel.c index 14558a9..cffc314 100644 --- a/src/tim/vx/internal/src/libnnext/vsi_nn_vxkernel.c +++ b/src/tim/vx/internal/src/libnnext/vsi_nn_vxkernel.c @@ -33,10 +33,6 @@ #include "libnnext/vsi_nn_vxkernel.h" #include "kernel/vsi_nn_kernel.h" #include "libnnext/vsi_nn_libnnext_resource.h" -#if VSI_USE_VXC_BINARY -/*this header can be only included once in all *.c files*/ -#include "libnnext/vx_bin/vxc_binaries.h" -#endif static char s_vx_resource_path[VSI_NN_MAX_PATH] = "VX"; @@ -454,7 +450,7 @@ vx_node vsi_nn_RegisterClientKernelAndCreateNode ) { /* - * Deprecated: use vsi_nn_RegisterClientKernelAndNewNode() insteatd. + * Deprecated: use vsi_nn_RegisterClientKernelAndNewNode() instead. */ vsi_nn_kernel_info_t kernel_info; char *resource_name[1] = {NULL}; @@ -543,17 +539,6 @@ const uint8_t * vsi_nn_VxBinResourceGetResource vx_size *len ) { -#if VSI_USE_VXC_BINARY - int i; - for (i = 0; i < vx_bin_resource_items_cnt; i++) - { - if (strncmp(vx_bin_resource_items[i].name, name, VSI_NN_MAX_PATH) == 0) - { - *len = vx_bin_resource_items[i].len; - return vx_bin_resource_items[i].data; - } - } -#endif return NULL; } /* vsi_nn_VxResourceGetBinResource() */