[Internal] support prebuilt kernel into shared library (#260)

Signed-off-by: xiang.zhang <xiang.zhang@verisilicon.com>

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;
+    switch( type )
-    for( i = 0; i < vx_bin_resource_items_cnt; i++ )
     {
-        if( strncmp( vx_bin_resource_items[i].name, source_name, VSI_NN_MAX_PATH ) == 0 )
+        case VSI_NN_KERNEL_TYPE_EVIS:
-        {
+            return _load_bin( source_name, size,
-            *size = (size_t)vx_bin_resource_items[i].len;
+                vx_bin_resource_items_vx, vx_bin_resource_items_vx_cnt, "_vx" );
-            return vx_bin_resource_items[i].data;
+            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;

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);
+    Tensor src_tensor = create_tensor_from_image2d_array(input, 1);
-    uchar *src_ptr_base = (uchar *)src_image.ptr;
+    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), \

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; \

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
 {

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() */