Huawei_Technology
/
mindspore
Not watched
1
0
Fork 0
Code
Releases 13 Wiki evaluate Activity
Issues 0 Pull Requests 0 Datasets Model Cloudbrain HPC
Browse Source

Abstract WRITEIMAGE and READIMAGE into the kernel

pull/15668/head
yeyunpeng2020 4 years ago
parent
6bd5dc8c18
commit
deeff1828d
45 changed files with 497 additions and 191 deletions
Split View
Diff Options
Show Stats Download Patch File Download Diff File
  1. +2
    -2
      mindspore/lite/src/runtime/gpu/opencl/opencl_runtime.cc
  2. +9
    -9
      mindspore/lite/src/runtime/kernel/opencl/cl/concat.cl
  3. +39
    -44
      mindspore/lite/src/runtime/kernel/opencl/cl/gather.cl
  4. +10
    -10
      mindspore/lite/src/runtime/kernel/opencl/cl/one_hot.cl
  5. +43
    -48
      mindspore/lite/src/runtime/kernel/opencl/cl/reshape.cl
  6. +9
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/activation.cc
  7. +9
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/argminmax.cc
  8. +9
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/arithmetic.cc
  9. +9
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/arithmetic_self.cc
  10. +7
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/batch_to_space_nd.cc
  11. +7
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/batchnorm.cc
  12. +28
    -2
      mindspore/lite/src/runtime/kernel/opencl/kernel/concat.cc
  13. +7
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/conv2d.cc
  14. +7
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/conv2d_transpose.cc
  15. +7
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/depthwise_conv2d.cc
  16. +7
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/fullconnection.cc
  17. +7
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/fusion_eltwise.cc
  18. +9
    -6
      mindspore/lite/src/runtime/kernel/opencl/kernel/gather.cc
  19. +8
    -2
      mindspore/lite/src/runtime/kernel/opencl/kernel/layer_norm.cc
  20. +7
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/matmul.cc
  21. +8
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/one_hot.cc
  22. +7
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/pad.cc
  23. +7
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/pooling2d.cc
  24. +7
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/power.cc
  25. +7
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/prelu.cc
  26. +7
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/reduce.cc
  27. +9
    -6
      mindspore/lite/src/runtime/kernel/opencl/kernel/reshape.cc
  28. +7
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/resize.cc
  29. +7
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/scale.cc
  30. +8
    -6
      mindspore/lite/src/runtime/kernel/opencl/kernel/softmax.cc
  31. +7
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/space_to_batch_nd.cc
  32. +7
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/space_to_depth.cc
  33. +7
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/sparse_to_dense.cc
  34. +7
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/split.cc
  35. +7
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/stack.cc
  36. +12
    -6
      mindspore/lite/src/runtime/kernel/opencl/kernel/strassen.cc
  37. +7
    -1
      mindspore/lite/src/runtime/kernel/opencl/kernel/strided_slice.cc
  38. +12
    -8
      mindspore/lite/src/runtime/kernel/opencl/kernel/transpose.cc
  39. +9
    -3
      mindspore/lite/src/runtime/kernel/opencl/kernel/winograd.cc
  40. +7
    -2
      mindspore/lite/src/runtime/kernel/opencl/utils.cc
  41. +2
    -1
      mindspore/lite/src/runtime/kernel/opencl/utils.h
  42. +1
    -1
      mindspore/lite/src/scheduler.cc
  43. +1
    -0
      mindspore/lite/test/models_gpu_fp16.cfg
  44. +42
    -0
      mindspore/lite/test/models_gpu_fp32.cfg
  45. +55
    -9
      mindspore/lite/test/run_benchmark_nets.sh

+ 2
- 2
mindspore/lite/src/runtime/gpu/opencl/opencl_runtime.cc View File

@@ -374,11 +374,11 @@ int OpenCLRuntime::BuildKernel(const cl::Kernel &kernel, const std::string &prog
if (fp16_enable_) {
build_option +=
" -DFP16_ENABLE=1 -DFLT=half -DFLT4=half4 -DFLT16=half16 -DAS_FLT4=as_half4 -DAS_UINT4=as_ushort4 -DUINT4=ushort4"
" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh -DTO_FLT=convert_half -DTO_FLT4=convert_half4";
" -DTO_FLT=convert_half -DTO_FLT4=convert_half4";
} else {
build_option +=
" -DFP16_ENABLE=0 -DFLT=float -DFLT4=float4 -DFLT16=float16 -DAS_FLT4=as_float4 -DAS_UINT4=as_uint4 -DUINT4=uint4"
" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef -DTO_FLT=convert_float -DTO_FLT4=convert_float4";
" -DTO_FLT=convert_float -DTO_FLT4=convert_float4";
}
build_option += " -DMAX_IMAGE2D_WIDTH=" + std::to_string(max_image2d_width_);
build_option =


+ 9
- 9
mindspore/lite/src/runtime/kernel/opencl/cl/concat.cl View File

@@ -11,7 +11,7 @@ __constant sampler_t smp_none = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE |
if (X >= output_shape.x * output_shape.y || Y >= output_shape.z || Z >= output_shape.w) { \
return; \
} \
FLT4 result;
TYPE4 result;

// axis = 1
#define DOConcat2inputaxis1_NHWC4 \
@@ -248,12 +248,12 @@ CONCAT2(2input, axis3, _NHWC4)
int Align_OutShape = output_shape.w; \
int index_output = (IN * output_shape.y + IH) * stride_w + IW * Align_OutShape * C4NUM;

int doconcat(__read_only image2d_t input, __global FLT *output, int Align_Shape, int4 input_shape, int IN, int IH,
int doconcat(__read_only image2d_t input, __global TYPE *output, int Align_Shape, int4 input_shape, int IN, int IH,
int Y, int index_output) {
int Remainder = input_shape.w % C4NUM;
for (int i = 0; i < Align_Shape; ++i) {
FLT4 result = READ_IMAGE(input, smp_none, (int2)((Y * Align_Shape + i), (IN * input_shape.y + IH)));
FLT result_temp[4] = {result.x, result.y, result.z, result.w};
TYPE4 result = READ_IMAGE(input, smp_none, (int2)((Y * Align_Shape + i), (IN * input_shape.y + IH)));
TYPE result_temp[4] = {result.x, result.y, result.z, result.w};
if ((i + 1) * C4NUM <= input_shape.w) {
for (int j = 0; j < C4NUM; ++j) {
output[index_output++] = result_temp[j];
@@ -268,7 +268,7 @@ int doconcat(__read_only image2d_t input, __global FLT *output, int Align_Shape,
}

__kernel void ConcatInput2UnAlign_NHWC4(__read_only image2d_t input0, __read_only image2d_t input1,
__global FLT *output, int4 input_shape0, int4 input_shape1, int stride_w,
__global TYPE *output, int4 input_shape0, int4 input_shape1, int stride_w,
int4 output_shape) {
CHECK_IDX_UNALIGN;
index_output = doconcat(input0, output, Align_Shape0, input_shape0, IN, IH, Y, index_output);
@@ -276,7 +276,7 @@ __kernel void ConcatInput2UnAlign_NHWC4(__read_only image2d_t input0, __read_onl
}

__kernel void ConcatInput3UnAlign_NHWC4(__read_only image2d_t input0, __read_only image2d_t input1,
__read_only image2d_t input2, __global FLT *output, int4 input_shape0,
__read_only image2d_t input2, __global TYPE *output, int4 input_shape0,
int4 input_shape1, int4 input_shape2, int stride_w, int4 output_shape) {
CHECK_IDX_UNALIGN;
int Align_Shape2 = UP_DIV(input_shape2.w, C4NUM);
@@ -287,7 +287,7 @@ __kernel void ConcatInput3UnAlign_NHWC4(__read_only image2d_t input0, __read_onl

__kernel void ConcatInput4UnAlign_NHWC4(__read_only image2d_t input0, __read_only image2d_t input1,
__read_only image2d_t input2, __read_only image2d_t input3,
__global FLT *output, int4 input_shape0, int4 input_shape1, int4 input_shape2,
__global TYPE *output, int4 input_shape0, int4 input_shape1, int4 input_shape2,
int4 input_shape3, int stride_w, int4 output_shape) {
CHECK_IDX_UNALIGN;
int Align_Shape2 = UP_DIV(input_shape2.w, C4NUM), Align_Shape3 = UP_DIV(input_shape3.w, C4NUM);
@@ -299,7 +299,7 @@ __kernel void ConcatInput4UnAlign_NHWC4(__read_only image2d_t input0, __read_onl

__kernel void ConcatInput5UnAlign_NHWC4(__read_only image2d_t input0, __read_only image2d_t input1,
__read_only image2d_t input2, __read_only image2d_t input3,
__read_only image2d_t input4, __global FLT *output, int4 input_shape0,
__read_only image2d_t input4, __global TYPE *output, int4 input_shape0,
int4 input_shape1, int4 input_shape2, int4 input_shape3, int4 input_shape4,
int stride_w, int4 output_shape) {
CHECK_IDX_UNALIGN;
@@ -315,7 +315,7 @@ __kernel void ConcatInput5UnAlign_NHWC4(__read_only image2d_t input0, __read_onl
__kernel void ConcatInput6UnAlign_NHWC4(__read_only image2d_t input0, __read_only image2d_t input1,
__read_only image2d_t input2, __read_only image2d_t input3,
__read_only image2d_t input4, __read_only image2d_t input5,
__global FLT *output, int4 input_shape0, int4 input_shape1, int4 input_shape2,
__global TYPE *output, int4 input_shape0, int4 input_shape1, int4 input_shape2,
int4 input_shape3, int4 input_shape4, int4 input_shape5, int stride_w,
int4 output_shape) {
CHECK_IDX_UNALIGN;


+ 39
- 44
mindspore/lite/src/runtime/kernel/opencl/cl/gather.cl View File

@@ -1,48 +1,43 @@
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
#define C4NUM 4
__constant sampler_t smp_zero = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;

#define GATHER(SUFFIX, READ_IMAGE, WRITE_IMAGE, TYPE) \
__kernel void gather_##SUFFIX(__write_only image2d_t dst_data, __read_only image2d_t src_data, \
__global int *indices, int4 src_size, int4 dst_size, int indices_num, int axis) { \
int X = get_global_id(0); \
int Y = get_global_id(1); \
int Z = get_global_id(2); \
if (X >= dst_size.x || Y >= dst_size.y * dst_size.w || Z >= dst_size.z || dst_size.y == 0) { \
return; \
} \
TYPE##4 res_data = (TYPE##4)(0.0f, 0.0f, 0.0f, 0.0f); \
int batch = Y / dst_size.y; \
int height = Y % dst_size.y; \
if (axis == 0) { \
res_data = READ_IMAGE(src_data, smp_zero, (int2)(X * src_size.z + Z, indices[batch] * src_size.y + height)); \
} else if (axis == 1) { \
res_data = READ_IMAGE(src_data, smp_zero, (int2)(X * src_size.z + Z, batch * src_size.y + indices[height])); \
} else if (axis == 2) { \
res_data = READ_IMAGE(src_data, smp_zero, (int2)(indices[X] * src_size.z + Z, batch * src_size.y + height)); \
} else if (axis == 3) { \
int offset[4] = {indices[Z * 4] / 4, indices[Z * 4 + 1] / 4, indices[Z * 4 + 2] / 4, indices[Z * 4 + 3] / 4}; \
TYPE tmp[4]; \
TYPE res_tmp[4]; \
for (int i = 0; i < indices_num; ++i) { \
TYPE##4 rd_data = (TYPE##4)(0.0f, 0.0f, 0.0f, 0.0f); \
rd_data = READ_IMAGE(src_data, smp_zero, (int2)(X * src_size.z + offset[i], batch * src_size.y + height)); \
if (i >= 1 && offset[i] != offset[i - 1]) { \
rd_data = READ_IMAGE(src_data, smp_zero, (int2)(X * src_size.z + offset[i], batch * src_size.y + height)); \
} \
tmp[0] = rd_data.x; \
tmp[1] = rd_data.y; \
tmp[2] = rd_data.z; \
tmp[3] = rd_data.w; \
res_tmp[i] = tmp[indices[Z * 4 + i] % 4]; \
} \
res_data.x = res_tmp[0]; \
res_data.y = res_tmp[1]; \
res_data.z = res_tmp[2]; \
res_data.w = res_tmp[3]; \
} \
WRITE_IMAGE(dst_data, (int2)(X * dst_size.z + Z, batch * dst_size.y + height), res_data); \
__kernel void gather(__write_only image2d_t dst_data, __read_only image2d_t src_data, __global int *indices,
int4 src_size, int4 dst_size, int indices_num, int axis) {
int X = get_global_id(0);
int Y = get_global_id(1);
int Z = get_global_id(2);
if (X >= dst_size.x || Y >= dst_size.y * dst_size.w || Z >= dst_size.z || dst_size.y == 0) {
return;
}
// GATHER(SUFFIX, READ_IMAGE, WRITE_IMAGE, TYPE)
GATHER(float, READ_IMAGE, WRITE_IMAGE, FLT);
GATHER(int, read_imagei, write_imagei, int);
TYPE4 res_data = (TYPE4)(0.0f, 0.0f, 0.0f, 0.0f);
int batch = Y / dst_size.y;
int height = Y % dst_size.y;
if (axis == 0) {
res_data = READ_IMAGE(src_data, smp_zero, (int2)(X * src_size.z + Z, indices[batch] * src_size.y + height));
} else if (axis == 1) {
res_data = READ_IMAGE(src_data, smp_zero, (int2)(X * src_size.z + Z, batch * src_size.y + indices[height]));
} else if (axis == 2) {
res_data = READ_IMAGE(src_data, smp_zero, (int2)(indices[X] * src_size.z + Z, batch * src_size.y + height));
} else if (axis == 3) {
int offset[4] = {indices[Z * 4] / 4, indices[Z * 4 + 1] / 4, indices[Z * 4 + 2] / 4, indices[Z * 4 + 3] / 4};
TYPE tmp[4];
TYPE res_tmp[4];
for (int i = 0; i < indices_num; ++i) {
TYPE4 rd_data = (TYPE4)(0.0f, 0.0f, 0.0f, 0.0f);
rd_data = READ_IMAGE(src_data, smp_zero, (int2)(X * src_size.z + offset[i], batch * src_size.y + height));
if (i >= 1 && offset[i] != offset[i - 1]) {
rd_data = READ_IMAGE(src_data, smp_zero, (int2)(X * src_size.z + offset[i], batch * src_size.y + height));
}
tmp[0] = rd_data.x;
tmp[1] = rd_data.y;
tmp[2] = rd_data.z;
tmp[3] = rd_data.w;
res_tmp[i] = tmp[indices[Z * 4 + i] % 4];
}
res_data.x = res_tmp[0];
res_data.y = res_tmp[1];
res_data.z = res_tmp[2];
res_data.w = res_tmp[3];
}
WRITE_IMAGE(dst_data, (int2)(X * dst_size.z + Z, batch * dst_size.y + height), res_data);
}

+ 10
- 10
mindspore/lite/src/runtime/kernel/opencl/cl/one_hot.cl View File

@@ -22,7 +22,7 @@ __kernel void OneHotAxis0(__read_only image2d_t src_data, __write_only image2d_t
int N = Z / out_shape.y;
int H = Z % out_shape.y;
int in_index = (H * out_shape.z + Y) * out_shape.w + X;
int4 indices = read_imagei(src_data, smp_zero, (int2)(in_index % in_image2d_shape.x, in_index / in_image2d_shape.x));
int4 indices = READ_IMAGE(src_data, smp_zero, (int2)(in_index % in_image2d_shape.x, in_index / in_image2d_shape.x));
int *indices_int = (int *)&indices;
for (int i = 0; i < C4NUM; i++) {
if (support_neg_index != 0 && indices_int[i] < 0) {
@@ -42,7 +42,7 @@ __kernel void OneHotAxis0(__read_only image2d_t src_data, __write_only image2d_t
if (4 * X + 3 < C) {
SET_ON_OR_OFF_VALUE(result.w, N, indices_int[3], on_value, off_value);
}
write_imagef(dst_data, (int2)(Y * out_shape.w + X, Z), result);
WRITE_IMAGE(dst_data, (int2)(Y * out_shape.w + X, Z), result);
}

__kernel void OneHotAxis1(__read_only image2d_t src_data, __write_only image2d_t dst_data, int2 in_image2d_shape,
@@ -54,7 +54,7 @@ __kernel void OneHotAxis1(__read_only image2d_t src_data, __write_only image2d_t
int N = Z / out_shape.y;
int H = Z % out_shape.y;
int in_index = (N * out_shape.z + Y) * out_shape.w + X;
int4 indices = read_imagei(src_data, smp_zero, (int2)(in_index % in_image2d_shape.x, in_index / in_image2d_shape.x));
int4 indices = READ_IMAGE(src_data, smp_zero, (int2)(in_index % in_image2d_shape.x, in_index / in_image2d_shape.x));
int *indices_int = (int *)&indices;
for (int i = 0; i < C4NUM; i++) {
if (support_neg_index != 0 && indices_int[i] < 0) {
@@ -74,7 +74,7 @@ __kernel void OneHotAxis1(__read_only image2d_t src_data, __write_only image2d_t
if (4 * X + 3 < C) {
SET_ON_OR_OFF_VALUE(result.w, H, indices_int[3], on_value, off_value);
}
write_imagef(dst_data, (int2)(Y * out_shape.w + X, Z), result);
WRITE_IMAGE(dst_data, (int2)(Y * out_shape.w + X, Z), result);
}

__kernel void OneHotAxis2(__read_only image2d_t src_data, __write_only image2d_t dst_data, int2 in_image2d_shape,
@@ -86,7 +86,7 @@ __kernel void OneHotAxis2(__read_only image2d_t src_data, __write_only image2d_t
int N = Z / out_shape.y;
int H = Z % out_shape.y;
int in_index = (N * out_shape.y + H) * out_shape.w + X;
int4 indices = read_imagei(src_data, smp_zero, (int2)(in_index % in_image2d_shape.x, in_index / in_image2d_shape.x));
int4 indices = READ_IMAGE(src_data, smp_zero, (int2)(in_index % in_image2d_shape.x, in_index / in_image2d_shape.x));
int *indices_int = (int *)&indices;
for (int i = 0; i < C4NUM; i++) {
if (support_neg_index != 0 && indices_int[i] < 0) {
@@ -106,7 +106,7 @@ __kernel void OneHotAxis2(__read_only image2d_t src_data, __write_only image2d_t
if (4 * X + 3 < C) {
SET_ON_OR_OFF_VALUE(result.w, Y, indices_int[3], on_value, off_value);
}
write_imagef(dst_data, (int2)(Y * out_shape.w + X, Z), result);
WRITE_IMAGE(dst_data, (int2)(Y * out_shape.w + X, Z), result);
}

__kernel void OneHotAxis3(__read_only image2d_t src_data, __write_only image2d_t dst_data, int2 in_image2d_shape,
@@ -121,7 +121,7 @@ __kernel void OneHotAxis3(__read_only image2d_t src_data, __write_only image2d_t
int in_index_c4 = (N * out_shape.y + H) * ci4_size + Y / 4;
int in_index_c4_remainder = Y % 4;
int4 indices =
read_imagei(src_data, smp_zero, (int2)(in_index_c4 % in_image2d_shape.x, in_index_c4 / in_image2d_shape.x));
READ_IMAGE(src_data, smp_zero, (int2)(in_index_c4 % in_image2d_shape.x, in_index_c4 / in_image2d_shape.x));
int *indices_int = (int *)&indices;
int index_one = indices_int[in_index_c4_remainder];
if (support_neg_index != 0 && index_one < 0) {
@@ -140,7 +140,7 @@ __kernel void OneHotAxis3(__read_only image2d_t src_data, __write_only image2d_t
if (4 * X + 3 < C) {
SET_ON_OR_OFF_VALUE(result.w, 4 * X + 3, index_one, on_value, off_value);
}
write_imagef(dst_data, (int2)(Y * out_shape.w + X, Z), result);
WRITE_IMAGE(dst_data, (int2)(Y * out_shape.w + X, Z), result);
}

__kernel void OneHot2DAxis3(__read_only image2d_t src_data, __write_only image2d_t dst_data, int2 in_image2d_shape,
@@ -150,7 +150,7 @@ __kernel void OneHot2DAxis3(__read_only image2d_t src_data, __write_only image2d
int Z = get_global_id(2); // H * N (out_shape.h is 1, so N == Z)
if (X >= out_shape.w || Y >= out_shape.z || Z >= out_shape.x * out_shape.y) return;
int in_index_c4_remainder = Z % 4;
int4 indices = read_imagei(src_data, smp_zero, (int2)(Z / C4NUM, 0));
int4 indices = READ_IMAGE(src_data, smp_zero, (int2)(Z / C4NUM, 0));
int *indices_int = (int *)&indices;
int index_one = indices_int[in_index_c4_remainder];
if (support_neg_index != 0 && index_one < 0) {
@@ -169,5 +169,5 @@ __kernel void OneHot2DAxis3(__read_only image2d_t src_data, __write_only image2d
if (4 * X + 3 < C) {
SET_ON_OR_OFF_VALUE(result.w, 4 * X + 3, index_one, on_value, off_value);
}
write_imagef(dst_data, (int2)(Y * out_shape.w + X, Z), result);
WRITE_IMAGE(dst_data, (int2)(Y * out_shape.w + X, Z), result);
}

+ 43
- 48
mindspore/lite/src/runtime/kernel/opencl/cl/reshape.cl View File

@@ -3,52 +3,47 @@
#define UP_DIV(x, y) (((x) + (y) - (1)) / (y))
__constant sampler_t smp_zero = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;

#define RESHAPE_NHWC4(SUFFIX, READ_IMAGE, WRITE_IMAGE, TYPE) \
__kernel void reshape_NHWC4_##SUFFIX(__read_only image2d_t src_data, __write_only image2d_t dst_data, int4 src_size, \
int4 dst_size) { \
int X = get_global_id(0); \
int Y = get_global_id(1); \
int CO4 = UP_DIV(dst_size.z, C4NUM); \
int CO4_rem = dst_size.z % C4NUM; \
if (X >= dst_size.x || Y > dst_size.y) { \
return; \
} \
int CI4 = UP_DIV(src_size.x, C4NUM); \
int CI4_rem = src_size.x % C4NUM; \
CI4_rem = (CI4_rem == 0) ? C4NUM : CI4_rem; \
int in_img_x = CI4 * src_size.y; \
TYPE##4 res = (TYPE##4)(0.0f); \
TYPE tmp[4]; \
TYPE res_tmp[4]; \
int gcnt = 0; \
if (CO4_rem == 0 && ((CI4_rem & 0x3) == 0)) { \
gcnt = X + dst_size.x * Y; \
res = READ_IMAGE(src_data, smp_zero, (int2)(gcnt % in_img_x, gcnt / in_img_x)); \
WRITE_IMAGE(dst_data, (int2)(X, Y), res); \
} else { \
int start = ((X / CO4 * dst_size.z + min(dst_size.z, (X % CO4) * C4NUM)) + dst_size.w * Y); \
gcnt = start / src_size.x * CI4 + (start % src_size.x) / C4NUM; \
start = start % src_size.x % C4NUM; \
for (int i = 0, n = 0, j = start; i < C4NUM; ++n, j = 0) { \
int X_src = (gcnt + n) % in_img_x; \
res = READ_IMAGE(src_data, smp_zero, (int2)(X_src, (gcnt + n) / in_img_x)); \
tmp[0] = res.x; \
tmp[1] = res.y; \
tmp[2] = res.z; \
tmp[3] = res.w; \
int k = (X_src % CI4) == (CI4 - 1) ? CI4_rem : C4NUM; \
for (; j < k && i < C4NUM; ++j, ++i) { \
res_tmp[i] = tmp[j]; \
} \
} \
res.x = res_tmp[0]; \
res.y = res_tmp[1]; \
res.z = res_tmp[2]; \
res.w = res_tmp[3]; \
WRITE_IMAGE(dst_data, (int2)(X, Y), res); \
} \
__kernel void reshape_NHWC4(__read_only image2d_t src_data, __write_only image2d_t dst_data, int4 src_size,
int4 dst_size) {
int X = get_global_id(0);
int Y = get_global_id(1);
int CO4 = UP_DIV(dst_size.z, C4NUM);
int CO4_rem = dst_size.z % C4NUM;
if (X >= dst_size.x || Y > dst_size.y) {
return;
}

// RESHAPE_NHWC4(SUFFIX, READ_IMAGE, WRITE_IMAGE, TYPE)
RESHAPE_NHWC4(float, READ_IMAGE, WRITE_IMAGE, FLT);
RESHAPE_NHWC4(int, read_imagei, write_imagei, int);
int CI4 = UP_DIV(src_size.x, C4NUM);
int CI4_rem = src_size.x % C4NUM;
CI4_rem = (CI4_rem == 0) ? C4NUM : CI4_rem;
int in_img_x = CI4 * src_size.y;
TYPE4 res = (TYPE4)(0.0f);
TYPE tmp[4];
TYPE res_tmp[4];
int gcnt = 0;
if (CO4_rem == 0 && ((CI4_rem & 0x3) == 0)) {
gcnt = X + dst_size.x * Y;
res = READ_IMAGE(src_data, smp_zero, (int2)(gcnt % in_img_x, gcnt / in_img_x));
WRITE_IMAGE(dst_data, (int2)(X, Y), res);
} else {
int start = ((X / CO4 * dst_size.z + min(dst_size.z, (X % CO4) * C4NUM)) + dst_size.w * Y);
gcnt = start / src_size.x * CI4 + (start % src_size.x) / C4NUM;
start = start % src_size.x % C4NUM;
for (int i = 0, n = 0, j = start; i < C4NUM; ++n, j = 0) {
int X_src = (gcnt + n) % in_img_x;
res = READ_IMAGE(src_data, smp_zero, (int2)(X_src, (gcnt + n) / in_img_x));
tmp[0] = res.x;
tmp[1] = res.y;
tmp[2] = res.z;
tmp[3] = res.w;
int k = (X_src % CI4) == (CI4 - 1) ? CI4_rem : C4NUM;
for (; j < k && i < C4NUM; ++j, ++i) {
res_tmp[i] = tmp[j];
}
}
res.x = res_tmp[0];
res.y = res_tmp[1];
res.z = res_tmp[2];
res.w = res_tmp[3];
WRITE_IMAGE(dst_data, (int2)(X, Y), res);
}
}

+ 9
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/activation.cc View File

@@ -73,7 +73,15 @@ int ActivationOpenCLKernel::Prepare() {
std::string program_name = "Activation";
ocl_runtime_->LoadSource(program_name, source);
std::string kernel_name = GetActTypeString(type_);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name);
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeInt32) {
build_options_ext = {" -DTYPE=int -DTYPE4=int4 -DWRITE_IMAGE=write_imagei -DREAD_IMAGE=read_imagei "};
} else if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DTYPE=float -DTYPE4=float4 -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DTYPE=half -DTYPE4=half4 -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
SetConstArgs();
SetGlobalLocal();
MS_LOG(DEBUG) << kernel_name << " init Done!";


+ 9
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/argminmax.cc View File

@@ -149,7 +149,15 @@ int ArgMinMaxOpenCLKernel::Prepare() {
std::string source = argminmax_source;
std::string program_name = "argminmax";
ocl_runtime_->LoadSource(program_name, source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name);
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeInt32) {
build_options_ext = {" -DTYPE=int -DTYPE4=int4 -DWRITE_IMAGE=write_imagei -DREAD_IMAGE=read_imagei "};
} else if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DTYPE=float -DTYPE4=float4 -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DTYPE=half -DTYPE4=half4 -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
#endif

auto *param = reinterpret_cast<ArgMinMaxParameter *>(this->op_parameter_);


+ 9
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/arithmetic.cc View File

@@ -191,7 +191,15 @@ int ArithmeticOpenCLKernel::Prepare() {
std::string program_name = "Arithmetic";
std::string source = arithmetic_source;
ocl_runtime_->LoadSource(program_name, source);
int error_code = ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name_);
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeInt32) {
build_options_ext = {" -DTYPE=int -DTYPE4=int4 -DWRITE_IMAGE=write_imagei -DREAD_IMAGE=read_imagei "};
} else if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DTYPE=float -DTYPE4=float4 -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DTYPE=half -DTYPE4=half4 -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
int error_code = ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name_, build_options_ext);
#endif
if (error_code != RET_OK) {
return error_code;


+ 9
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/arithmetic_self.cc View File

@@ -89,7 +89,15 @@ int ArithmeticSelfOpenCLKernel::Prepare() {
MS_LOG(DEBUG) << "execute kernel name : " << kernel_name;
std::string program_name = "ArithmeticSelf";
ocl_runtime_->LoadSource(program_name, arithmeticself_source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name);
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeInt32) {
build_options_ext = {" -DTYPE=int -DTYPE4=int4 -DWRITE_IMAGE=write_imagei -DREAD_IMAGE=read_imagei "};
} else if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DTYPE=float -DTYPE4=float4 -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DTYPE=half -DTYPE4=half4 -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
SetGlobalLocal();
SetConstArgs();
return RET_OK;


+ 7
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/batch_to_space_nd.cc View File

@@ -94,7 +94,13 @@ int BatchToSpaceNDOpenCLKernel::Prepare() {
std::string source = batch_to_space_nd_source;
std::string program_name = "batch_to_space_nd";
ocl_runtime_->LoadSource(program_name, source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name);
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
#endif

SetGlobalLocal();


+ 7
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/batchnorm.cc View File

@@ -168,7 +168,13 @@ int BatchNormOpenCLKernel::Prepare() {
std::string source = batchnorm_source;
std::string program_name = "Batch_normalization";
ocl_runtime_->LoadSource(program_name, source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name);
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
MS_LOG(DEBUG) << kernel_name << " Init Done!";
int ret = Initweight();
if (ret) {


+ 28
- 2
mindspore/lite/src/runtime/kernel/opencl/kernel/concat.cc View File

@@ -170,7 +170,24 @@ int ConcatOpenCLKernel::ConvertWeightToTensor() {
bool src_is_fp16 = in_tensor->data_type() == kNumberTypeFloat16;
PackNHWCToNHWC4(in_tensor->data_c(), weight.data(), src_is_fp16,
fp16_enable && in_tensor->data_type() != kNumberTypeInt32, in_shape);
size_t dtype = fp16_enable && in_tensor->data_type() != kNumberTypeInt32 ? CL_HALF_FLOAT : CL_FLOAT;
size_t dtype;
switch (in_tensor->data_type()) {
case kNumberTypeInt32: {
dtype = CL_SIGNED_INT32;
break;
}
case kNumberTypeFloat32: {
dtype = CL_FLOAT;
break;
}
case kNumberTypeFloat16: {
dtype = CL_HALF_FLOAT;
break;
}
default:
MS_LOG(ERROR) << "Unsupported data type is" << in_tensor->data_type();
return RET_ERROR;
}
ImageSize img_size{in_shape.width, in_shape.height, dtype};
auto weight_ptr_ = allocator->Malloc(img_size, weight.data());
weight_ptrs_.push_back(weight_ptr_);
@@ -207,7 +224,15 @@ int ConcatOpenCLKernel::Prepare() {
std::string source = concat_source;
std::string program_name = "Concat";
ocl_runtime_->LoadSource(program_name, source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, {});
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeInt32) {
build_options_ext = {" -DTYPE=int -DTYPE4=int4 -DWRITE_IMAGE=write_imagei -DREAD_IMAGE=read_imagei "};
} else if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DTYPE=float -DTYPE4=float4 -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DTYPE=half -DTYPE4=half4 -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
MS_LOG(DEBUG) << kernel_name << " Init Done!";
SetConstArgs();
SetGlobalLocal();
@@ -235,4 +260,5 @@ int ConcatOpenCLKernel::Run() {

REG_KERNEL(kGPU, kNumberTypeFloat32, PrimitiveType_Concat, OpenCLKernelCreator<ConcatOpenCLKernel>)
REG_KERNEL(kGPU, kNumberTypeFloat16, PrimitiveType_Concat, OpenCLKernelCreator<ConcatOpenCLKernel>)
REG_KERNEL(kGPU, kNumberTypeInt32, PrimitiveType_Concat, OpenCLKernelCreator<ConcatOpenCLKernel>)
} // namespace mindspore::kernel

+ 7
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/conv2d.cc View File

@@ -144,7 +144,13 @@ void Conv2DOpenCLKernel::BuildKernel() {
kernel_name << "_Img";
}
ocl_runtime_->LoadSource(program_name, GetActDefines() + conv2d_source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name.str());
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name.str(), build_options_ext);
}

void Conv2DOpenCLKernel::SetBlockSize() {


+ 7
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/conv2d_transpose.cc View File

@@ -65,7 +65,13 @@ int Conv2dTransposeOpenCLKernel::Prepare() {
std::string source = GetActDefines() + conv2d_transpose_source;
std::string program_name = "conv2d_transpose";
ocl_runtime_->LoadSource(program_name, source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name);
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
#endif
auto ret = InitWeights();
if (ret != RET_OK) {


+ 7
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/depthwise_conv2d.cc View File

@@ -85,7 +85,13 @@ int DepthwiseConv2dOpenCLKernel::Prepare() {
std::string program_name = "DepthwiseConv2d";
std::string source = depthwise_conv2d_source;
ocl_runtime_->LoadSource(program_name, source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name);
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
#endif
auto ret = InitWeights();
if (ret != RET_OK) {


+ 7
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/fullconnection.cc View File

@@ -105,7 +105,13 @@ int FullConnectionOpenCLKernel::Prepare() {
std::string source = fullconnection_source;
std::string program_name = "FullConnection";
ocl_runtime_->LoadSource(program_name, GetActDefines() + source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name);
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
#endif
auto ret = InitWeights();
if (ret != RET_OK) {


+ 7
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/fusion_eltwise.cc View File

@@ -147,7 +147,13 @@ int FusionEltwiseOpenCLKernel::Prepare() {
std::string program_name = "FusionEltwise\n" + source;
std::string kernel_name = "FusionEltwise";
ocl_runtime_->LoadSource(program_name, source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name);
std::vector<std::string> build_options_ext;
if (ocl_runtime_->GetFp16Enable()) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
} else {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
InitWeights();
SetGlobalLocal();
SetConstArgs();


+ 9
- 6
mindspore/lite/src/runtime/kernel/opencl/kernel/gather.cc View File

@@ -106,11 +106,6 @@ void GatherOpenCLKernel::SetGlobalLocal() {

int GatherOpenCLKernel::Prepare() {
std::string kernel_name = "gather";
if (desc_.data_type == kNumberTypeInt32) {
kernel_name += "_int";
} else {
kernel_name += "_float";
}
if (in_tensors_.at(0)->shape().size() == 1 && axis_ == 0) {
axis_ = 3;
}
@@ -119,7 +114,15 @@ int GatherOpenCLKernel::Prepare() {
#else
std::string program_name = "gather";
ocl_runtime_->LoadSource(program_name, gather_source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, {});
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeInt32) {
build_options_ext = {" -DTYPE=int -DTYPE4=int4 -DWRITE_IMAGE=write_imagei -DREAD_IMAGE=read_imagei "};
} else if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DTYPE=float -DTYPE4=float4 -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DTYPE=half -DTYPE4=half4 -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
#endif
if (in_tensors_.at(1)->IsConst()) {
intensor1_is_tensor = false;


+ 8
- 2
mindspore/lite/src/runtime/kernel/opencl/kernel/layer_norm.cc View File

@@ -171,9 +171,15 @@ int LayerNormOpenCLKernel::Prepare() {
std::string source = layer_norm_source;
std::string program_name = "LayerNormalization";
ocl_runtime_->LoadSource(program_name, source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name);
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
kernel_name_mean_var += "Axis" + std::to_string(normalized_axis_) + "NHWC4";
ocl_runtime_->BuildKernel(kernel_mean_var_, program_name, kernel_name_mean_var);
ocl_runtime_->BuildKernel(kernel_mean_var_, program_name, kernel_name_mean_var, build_options_ext);
MS_LOG(DEBUG) << kernel_name << " Init Done!";
SetConstArgs();
SetGlobalLocal();


+ 7
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/matmul.cc View File

@@ -92,7 +92,13 @@ int MatMulOpenCLKernel::Prepare() {
std::string source = matmul_source;
std::string program_name = "MatMul";
ocl_runtime_->LoadSource(program_name, source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name);
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);

#endif
SetConstArgs();


+ 8
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/one_hot.cc View File

@@ -51,10 +51,17 @@ int OneHotOpenCLKernel::Prepare() {
#ifdef PROGRAM_WITH_IL
kernel_ = ocl_runtime_->GetKernelFromBinary(kernel_name);
#else

std::string source = one_hot_source;
std::string program_name = "OneHot";
ocl_runtime_->LoadSource(program_name, source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, {});
std::vector<std::string> build_options_ext;
if (ocl_runtime_->GetFp16Enable()) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=write_imagei "};
} else {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagei "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
#endif
InitWeights();
SetConstArgs();


+ 7
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/pad.cc View File

@@ -74,7 +74,13 @@ int PadOpenCLKernel::Prepare() {
const std::string source = pad_source;
const std::string program_name = "Pad";
ocl_runtime_->LoadSource(program_name, source);
ocl_runtime_->BuildKernel(kernel_, program_name, "Pad");
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, "Pad", build_options_ext);
SetConstArgs();
return RET_OK;
}


+ 7
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/pooling2d.cc View File

@@ -82,7 +82,13 @@ int PoolingOpenCLKernel::Prepare() {
std::string source = pooling2d_source;
std::string program_name = "Pooling2d";
ocl_runtime_->LoadSource(program_name, source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name);
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
#endif
SetConstArgs();
SetGlobalLocal();


+ 7
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/power.cc View File

@@ -118,7 +118,13 @@ int PowerOpenCLKernel::Prepare() {
scale_ = param->scale_;
shift_ = param->shift_;
ocl_runtime_->LoadSource(program_name, source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name);
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
MS_LOG(DEBUG) << kernel_name << " Init Done!";
SetGlobalLocal();
SetConstArgs();


+ 7
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/prelu.cc View File

@@ -130,7 +130,13 @@ int PReluOpenCLKernel::Prepare() {
std::string program_name = "PRelu";
std::string kernel_name = "PRelu_" + std::string(weight_is_scalar ? "scalar" : "vector");
ocl_runtime_->LoadSource(program_name, source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name);
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
InitWeights();
MS_LOG(DEBUG) << program_name << " init Done!";
MS_LOG(DEBUG) << "kernel_name=: " << kernel_name << " init Done!";


+ 7
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/reduce.cc View File

@@ -186,7 +186,13 @@ int ReduceOpenCLKernel::Prepare() {
std::string source = reduce_source;
std::string program_name = "Reduce";
ocl_runtime_->LoadSource(program_name, source);
auto ret = ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name);
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
auto ret = ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
if (ret != RET_OK) {
return ret;
}


+ 9
- 6
mindspore/lite/src/runtime/kernel/opencl/kernel/reshape.cc View File

@@ -73,18 +73,21 @@ void ReshapeOpenCLKernel::SetGlobalLocal() {

int ReshapeOpenCLKernel::Prepare() {
std::string kernel_name = "reshape_NHWC4";
if (desc_.data_type == kNumberTypeInt32) {
kernel_name += "_int";
} else {
kernel_name += "_float";
}
#ifdef PROGRAM_WITH_IL
kernel_ = ocl_runtime_->GetKernelFromBinary(kernel_name);
#else
std::string source = reshape_source;
std::string program_name = "reshape";
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeInt32) {
build_options_ext = {" -DTYPE=int -DTYPE4=int4 -DWRITE_IMAGE=write_imagei -DREAD_IMAGE=read_imagei "};
} else if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DTYPE=float -DTYPE4=float4 -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DTYPE=half -DTYPE4=half4 -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->LoadSource(program_name, source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, {});
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
#endif

SetGlobalLocal();


+ 7
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/resize.cc View File

@@ -70,7 +70,13 @@ int ResizeOpenCLKernel::Prepare() {
std::string source = resize_source;
std::string program_name = "Resize";
ocl_runtime_->LoadSource(program_name, source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name);
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
#endif
SetConstArgs();
SetGlobalLocal();


+ 7
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/scale.cc View File

@@ -166,7 +166,13 @@ int ScaleOpenCLKernel::Prepare() {
std::string program_name = "Scale";
std::string source = GetActDefines() + scale_source;
ocl_runtime_->LoadSource(program_name, source);
error_code = ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name);
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
error_code = ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
#endif
if (error_code != RET_OK) {
return error_code;


+ 8
- 6
mindspore/lite/src/runtime/kernel/opencl/kernel/softmax.cc View File

@@ -84,13 +84,15 @@ int SoftmaxOpenCLKernel::Prepare() {
#else
std::string program_name = "Softmax";
ocl_runtime_->LoadSource(program_name, source);
std::vector<std::string> ext_build_opt;
if (out_tensors_[0]->data_type() == kNumberTypeFloat32) {
ext_build_opt.push_back("-DOUT_FLT4=convert_float4 -DWRITE_IMAGEOUT=write_imagef");
} else {
ext_build_opt.push_back("-DOUT_FLT4=convert_half4 -DWRITE_IMAGEOUT=write_imageh");
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {
" -DOUT_FLT4=convert_float4 -DWRITE_IMAGEOUT=write_imagef -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {
" -DOUT_FLT4=convert_half4 -DWRITE_IMAGEOUT=write_imageh -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, ext_build_opt);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
#endif
SetConstArgs();
SetGlobalLocal();


+ 7
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/space_to_batch_nd.cc View File

@@ -100,7 +100,13 @@ int SpaceToBatchNDOpenCLKernel::Prepare() {
std::string source = space_to_batch_nd_source;
std::string program_name = "space_to_batch_nd";
ocl_runtime_->LoadSource(program_name, source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name);
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
#endif

SetGlobalLocal();


+ 7
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/space_to_depth.cc View File

@@ -58,7 +58,13 @@ int SpaceToDepthOpenCLKernel::Prepare() {
std::string source = space_to_depth_source;
std::string program_name = "SpaceToDepth";
ocl_runtime_->LoadSource(program_name, source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name);
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
#endif
SetConstArgs();
SetGlobalLocal();


+ 7
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/sparse_to_dense.cc View File

@@ -148,7 +148,13 @@ int SparseToDenseOpenCLKernel::Prepare() {
std::string source = sparse_to_dense_source;
std::string program_name = "SparseToDense";
ocl_runtime_->LoadSource(program_name, source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name);
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);

if (in_tensors_.size() > 3) {
auto input_tensor3 = in_tensors_[3];


+ 7
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/split.cc View File

@@ -136,7 +136,13 @@ int SplitOpenCLKernel::Prepare() {
std::string source = split_source;
std::string program_name = "split";
ocl_runtime_->LoadSource(program_name, source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name);
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
MS_LOG(DEBUG) << kernel_name << " Init Done!";
SetConstArgs();
SetGlobalLocal();


+ 7
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/stack.cc View File

@@ -165,7 +165,13 @@ int StackOpenCLKernel::Prepare() {
std::string source = stack_source;
std::string program_name = "stack";
ocl_runtime_->LoadSource(program_name, source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name);
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
SetConstArgs();
SetGlobalLocal();



+ 12
- 6
mindspore/lite/src/runtime/kernel/opencl/kernel/strassen.cc View File

@@ -32,12 +32,18 @@ int StrassenOpenCLKernel::Prepare() {
std::string source = strassen_source;
std::string program_name = "MatMul";
ocl_runtime_->LoadSource(program_name, source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name);
ocl_runtime_->BuildKernel(kernel_IMG_add_sub_2, program_name, "MatMul_IMG_Add_Sub_2");
ocl_runtime_->BuildKernel(kernel_BUF_add_sub_2, program_name, "MatMul_BUF_Add_Sub_2");
ocl_runtime_->BuildKernel(kernel_back_result, program_name, "Strassen_Back_Result");
ocl_runtime_->BuildKernel(MatMul_StrassenBUFFilled, program_name, "MatMul_BUF_Filled");
ocl_runtime_->BuildKernel(MatMul_StrassenIMGFilled, program_name, "MatMul_IMG_Filled");
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
ocl_runtime_->BuildKernel(kernel_IMG_add_sub_2, program_name, "MatMul_IMG_Add_Sub_2", build_options_ext);
ocl_runtime_->BuildKernel(kernel_BUF_add_sub_2, program_name, "MatMul_BUF_Add_Sub_2", build_options_ext);
ocl_runtime_->BuildKernel(kernel_back_result, program_name, "Strassen_Back_Result", build_options_ext);
ocl_runtime_->BuildKernel(MatMul_StrassenBUFFilled, program_name, "MatMul_BUF_Filled", build_options_ext);
ocl_runtime_->BuildKernel(MatMul_StrassenIMGFilled, program_name, "MatMul_IMG_Filled", build_options_ext);
auto ret = InitWeights();
if (ret != RET_OK) {
return ret;


+ 7
- 1
mindspore/lite/src/runtime/kernel/opencl/kernel/strided_slice.cc View File

@@ -89,7 +89,13 @@ int StridedSliceOpenCLKernel::CheckSpecs() {
int StridedSliceOpenCLKernel::Prepare() {
std::string program_name = "strided_slice";
ocl_runtime_->LoadSource(program_name, strided_slice_source);
ocl_runtime_->BuildKernel(kernel_, program_name, "strided_slice");
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, "strided_slice", build_options_ext);
SetConstArgs();
SetGlobalLocal();
return RET_OK;


+ 12
- 8
mindspore/lite/src/runtime/kernel/opencl/kernel/transpose.cc View File

@@ -84,21 +84,19 @@ int TransposeOpenCLKernel::Prepare() {
perm_4d_[2] = 2;
perm_4d_[3] = 3;
}
std::string kernel_name = "transpose";

if (tensor_size_.N == 1 && perm_4d_[0] == 0 && perm_4d_[1] == 3 && perm_4d_[2] == 1 && perm_4d_[3] == 2) {
type_ = TransposeType::AXIS0312;
kernel_name += "_0312";
} else if (tensor_size_.N == 1 && perm_4d_[0] == 0 && perm_4d_[1] == 2 && perm_4d_[2] == 3 && perm_4d_[3] == 1) {
type_ = TransposeType::AXIS0231;
} else {
type_ = TransposeType::GENERAL;
}
std::string kernel_name = "transpose";
if (type_ == TransposeType::AXIS0312) {
kernel_name += "_0312";
} else if (type_ == TransposeType::AXIS0231) {
kernel_name += "_0231";
} else {
type_ = TransposeType::GENERAL;
kernel_name += "_general";
}

if (in_tensors_[0]->shape().size() == 4 &&
in_tensors_[0]->shape()[2] * UP_DIV(in_tensors_[0]->shape()[3], C4NUM) > ocl_runtime_->GetMaxImage2DWidth()) {
// just for input
@@ -112,7 +110,13 @@ int TransposeOpenCLKernel::Prepare() {
std::string source = transpose_source;
std::string program_name = "transpose";
ocl_runtime_->LoadSource(program_name, source);
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name);
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options_ext);
#endif
SetConstArgs();
SetGlobalLocal();


+ 9
- 3
mindspore/lite/src/runtime/kernel/opencl/kernel/winograd.cc View File

@@ -82,10 +82,16 @@ std::vector<float> GenerateWinogradFilter(void *src, TypeId dtype, size_t CO, si
void WinogradOpenCLKernel::BuildKernel() {
std::string program_name = "winograd";
ocl_runtime_->LoadSource(program_name, GetActDefines() + winograd_source);
ocl_runtime_->BuildKernel(kernel_4x4to36_, program_name, "Winograd4x4To36");
std::vector<std::string> build_options_ext;
if (desc_.data_type == kNumberTypeFloat32) {
build_options_ext = {" -DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef "};
} else if (desc_.data_type == kNumberTypeFloat16) {
build_options_ext = {" -DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh "};
}
ocl_runtime_->BuildKernel(kernel_4x4to36_, program_name, "Winograd4x4To36", build_options_ext);
ocl_runtime_->BuildKernel(kernel_, program_name,
filter_type_ == MemType::IMG ? "WinogradConv2D_Img" : "WinogradConv2D");
ocl_runtime_->BuildKernel(kernel_36to4x4_, program_name, "Winograd36To4x4");
filter_type_ == MemType::IMG ? "WinogradConv2D_Img" : "WinogradConv2D", build_options_ext);
ocl_runtime_->BuildKernel(kernel_36to4x4_, program_name, "Winograd36To4x4", build_options_ext);
}

void WinogradOpenCLKernel::InitFilter() {


+ 7
- 2
mindspore/lite/src/runtime/kernel/opencl/utils.cc View File

@@ -234,19 +234,24 @@ int GetBroadcastGpuAxis(int ndim, int ori_axis) {
return axis;
}

void PackNHWCToNHWC4(void *src, void *dst, bool src_is_fp16, bool dst_is_fp16, const GpuTensorInfo &tensor) {
void PackNHWCToNHWC4(void *src, void *dst, bool src_is_fp16, bool dst_is_fp16, const GpuTensorInfo &tensor,
int data_type) {
MS_ASSERT(src);
MS_ASSERT(dst);
auto src_fp16 = reinterpret_cast<float16_t *>(src);
auto src_fp32 = reinterpret_cast<float32_t *>(src);
auto src_int32 = reinterpret_cast<int32_t *>(src);
auto dst_fp16 = reinterpret_cast<float16_t *>(dst);
auto dst_fp32 = reinterpret_cast<float32_t *>(dst);
auto dst_int32 = reinterpret_cast<int32_t *>(dst);
for (int n = 0, src_idx = 0; n < tensor.N; n++) {
for (int h = 0; h < tensor.H; ++h) {
for (int w = 0; w < tensor.W; ++w) {
for (int c = 0; c < tensor.C; ++c, ++src_idx) {
int dst_idx = ((n * tensor.H + h) * tensor.W + w) * tensor.Slice * C4NUM + c;
if (dst_is_fp16) {
if (data_type == kNumberTypeInt32) {
dst_int32[dst_idx] = src_int32[src_idx];
} else if (dst_is_fp16) {
dst_fp16[dst_idx] = src_is_fp16 ? src_fp16[src_idx] : static_cast<float16_t>(src_fp32[src_idx]);
} else {
dst_fp32[dst_idx] = src_is_fp16 ? static_cast<float32_t>(src_fp16[src_idx]) : src_fp32[src_idx];


+ 2
- 1
mindspore/lite/src/runtime/kernel/opencl/utils.h View File

@@ -58,7 +58,8 @@ int WriteToBin(const std::string &file_path, void *data, size_t size);

int GetBroadcastGpuAxis(int ndim, int ori_axis);

void PackNHWCToNHWC4(void *src, void *dst, bool src_is_fp16, bool dst_is_fp16, const GpuTensorInfo &tensor);
void PackNHWCToNHWC4(void *src, void *dst, bool src_is_fp16, bool dst_is_fp16, const GpuTensorInfo &tensor,
int data_type = kNumberTypeFloat32);

int CheckParamLikeTensor(const std::string &kernel_name, const std::string &tensor_name, lite::Tensor *tensor,
TypeId expect_data_type, const std::vector<int> &expect_shape);


+ 1
- 1
mindspore/lite/src/scheduler.cc View File

@@ -394,7 +394,7 @@ kernel::LiteKernel *Scheduler::FindGpuKernel(const std::vector<Tensor *> &in_ten
if (context_->IsGpuEnabled()) {
// support more data type like int32
kernel::KernelKey gpu_desc{kGPU, desc.data_type, desc.type};
if (context_->IsGpuFloat16Enabled()) {
if (desc.data_type == kNumberTypeFloat32 && context_->IsGpuFloat16Enabled()) {
gpu_desc.data_type = kNumberTypeFloat16;
}



+ 1
- 0
mindspore/lite/test/models_gpu_fp16.cfg View File

@@ -9,3 +9,4 @@ mtk_model_emotions_0727_nosoftmax.tflite
landmark
PoseNet_dla_17_x512_tmp
plat_isface
ml_location_lane_counter.onnx;5.5

+ 42
- 0
mindspore/lite/test/models_gpu_fp32.cfg View File

@@ -57,3 +57,45 @@ mtk_model_face_dress.pb;0.5;1;1,128,128,3
hiai_model_normalize_object_scene_ps_20200519.pb;0.5;1;1,224,224,3
hiai_label_and_video.pb;0.5;1;1,224,224,3
tinyyolov2-8.onnx;0.5;1;1,416,416,3
mtk_detect-mbv2-shortcut-400-400-simplified.onnx
emotion-ferplus-8.onnx
rcnn-ilsvrc13-9.onnx
shufflenet-v2-10.onnx
squeezenet1.1-7.onnx
ml_table_detection_fp32_tmp.onnx
ml_table_segment.onnx
googlenet-9.onnx
inception-v1-9.onnx
shufflenet-9.onnx
ml_face_3d.onnx
gts_version-RFB-320_simplified.onnx
mnist-8.onnx
ml_video_edit_judge.onnx
ml_video_edit_vignet.onnx
hdc_mobilenet_1w_class.onnx
ml_video_edit_imitate_filter.onnx
ml_edu_kit_hand_detection.onnx
ml_edu_kit_hand_key_position.onnx
mtk_detect-deeper-halfdeeper-mbv1-shortcut-400-400_nopostprocess_simplified_onnx.onnx
mtk_detect-mbv1-shortcut-400-400_nopostprocess_simplified_onnx.onnx
mtk_detect-deeper-halfdeeper-mbv1-lastearlySSD-shortcut-400-400_nopostprocess_simplified_onnx.onnx
ml_2012_ocr_detection_tmp.onnx
ml_video_edit_enhance_update_tmp.onnx
bloom_hongmo_detection_tmp.onnx
Q_face_recognition.onnx
Q888_iris_detect.onnx
ml_ocr_bank_card_detection_inception_tmp
ml_ocr_detect_20200305
Q_iMaxDN_RGB_385_p_RGB_RGB_pb2tflite.tflite
Q_iMaxSR_RGB_385_p_pb2tflite.tflite
mtk_age_gender.pb
mtk_model_ckpt.pb
Q_inception-249970-672-11-16.pb
Q_crnn_screen_slim400w_more_20w.pb
hiai_ssd_mobilenetv2_object.pb
hiai_humanDetection.pb
mtk_face_features_v1.pb
Q_crnn_ori_75w_slim_norm.pb
Q_crnn_ori_v2_405001_notrans_nopre.pb
bolt_segment.pb
ml_location_lane_counter.onnx;2

+ 55
- 9
mindspore/lite/test/run_benchmark_nets.sh View File

@@ -2123,15 +2123,17 @@ function Run_gpu() {
input_files=$input_files${data_path}'input/'$model_name'.ms.bin_'$i','
done
fi
if [[ ${accuracy_limit} == "" ]]; then
accuracy_limit="0.5"
fi
echo ${model_name} >> "${run_gpu_log_file}"
echo 'cd /data/local/tmp/benchmark_test' > adb_run_cmd.txt
if [[ $input_shapes == "" ]]; then
echo 'export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/data/local/tmp/benchmark_test;./benchmark --device=GPU --modelFile='${model_name}'.ms --inDataFile='${input_files}' --benchmarkDataFile='${output_file} >> "${run_gpu_log_file}"
echo 'export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/data/local/tmp/benchmark_test;./benchmark --device=GPU --modelFile='${model_name}'.ms --inDataFile='${input_files}' --benchmarkDataFile='${output_file} >> adb_run_cmd.txt
echo 'export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/data/local/tmp/benchmark_test;./benchmark --device=GPU --accuracyThreshold='${accuracy_limit}' --modelFile='${model_name}'.ms --inDataFile='${input_files}' --benchmarkDataFile='${output_file} >> "${run_gpu_log_file}"
echo 'export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/data/local/tmp/benchmark_test;./benchmark --device=GPU --accuracyThreshold='${accuracy_limit}' --modelFile='${model_name}'.ms --inDataFile='${input_files}' --benchmarkDataFile='${output_file} >> adb_run_cmd.txt
else
echo 'export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/data/local/tmp/benchmark_test;./benchmark --inputShapes='${input_shapes}' --accuracyThreshold='${accuracy_limit}' --device=GPU --modelFile='${model_name}'.ms --inDataFile='${input_files}' --benchmarkDataFile='${output_file} >> "${run_gpu_log_file}"
echo 'export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/data/local/tmp/benchmark_test;./benchmark --inputShapes='${input_shapes}' --accuracyThreshold='${accuracy_limit}' --device=GPU --modelFile='${model_name}'.ms --inDataFile='${input_files}' --benchmarkDataFile='${output_file} >> adb_run_cmd.txt

fi

adb -s ${device_id} shell < adb_run_cmd.txt >> "${run_gpu_log_file}"
@@ -2148,17 +2150,39 @@ function Run_gpu() {
if [[ $model_name == \#* ]]; then
continue
fi
model_name=`echo ${line} | awk -F ';' '{print $1}'`
accuracy_limit=`echo ${line} | awk -F ';' '{print $2}'`
input_num=`echo ${line} | awk -F ';' '{print $3}'`
input_shapes=`echo ${line} | awk -F ';' '{print $4}'`
input_files=""
data_path="/data/local/tmp/input_output/"
output_file=${data_path}'output/'${model_name}'.ms.out'
if [[ ${input_num} == "" || ${input_num} == 1 ]]; then
input_files=/data/local/tmp/input_output/input/${model_name}.ms.bin
else
for i in $(seq 1 $input_num)
do
input_files=$input_files${data_path}'input/'$model_name'.ms.bin_'$i','
done
fi
if [[ ${accuracy_limit} == "" ]]; then
accuracy_limit="5"
fi
echo ${model_name} >> "${run_gpu_log_file}"
echo 'cd /data/local/tmp/benchmark_test' > adb_run_cmd.txt
echo 'export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/data/local/tmp/benchmark_test;./benchmark --device=GPU --modelFile='${model_name}'.ms --inDataFile=/data/local/tmp/input_output/input/'${model_name}'.ms.bin --benchmarkDataFile=/data/local/tmp/input_output/output/'${model_name}'.ms.out --enableFp16=true --accuracyThreshold=5' >> "${run_gpu_log_file}"
echo 'export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/data/local/tmp/benchmark_test;./benchmark --device=GPU --modelFile='${model_name}'.ms --inDataFile=/data/local/tmp/input_output/input/'${model_name}'.ms.bin --benchmarkDataFile=/data/local/tmp/input_output/output/'${model_name}'.ms.out --enableFp16=true --accuracyThreshold=5' >> adb_run_cmd.txt
if [[ $input_shapes == "" ]]; then
echo 'export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/data/local/tmp/benchmark_test;./benchmark --device=GPU --enableFp16=true --accuracyThreshold='${accuracy_limit}' --modelFile='${model_name}'.ms --inDataFile='${input_files}' --benchmarkDataFile='${output_file} >> "${run_gpu_log_file}"
echo 'export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/data/local/tmp/benchmark_test;./benchmark --device=GPU --enableFp16=true --accuracyThreshold='${accuracy_limit}' --modelFile='${model_name}'.ms --inDataFile='${input_files}' --benchmarkDataFile='${output_file} >> adb_run_cmd.txt
else
echo 'export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/data/local/tmp/benchmark_test;./benchmark --device=GPU --enableFp16=true --inputShapes='${input_shapes}' --accuracyThreshold='${accuracy_limit}' --modelFile='${model_name}'.ms --inDataFile='${input_files}' --benchmarkDataFile='${output_file} >> "${run_gpu_log_file}"
echo 'export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/data/local/tmp/benchmark_test;./benchmark --device=GPU --enableFp16=true --inputShapes='${input_shapes}' --accuracyThreshold='${accuracy_limit}' --modelFile='${model_name}'.ms --inDataFile='${input_files}' --benchmarkDataFile='${output_file} >> adb_run_cmd.txt
fi
adb -s ${device_id} shell < adb_run_cmd.txt >> "${run_gpu_log_file}"
if [ $? = 0 ]; then
run_result='arm64_gpu_fp16: '${model_name}' pass'; echo ${run_result} >> ${run_benchmark_result_file}
else
run_result='arm64_gpu_fp16: '${model_name}' failed'; echo ${run_result} >> ${run_benchmark_result_file}; return 1
fi
#sleep 1
done < ${models_gpu_fp16_config}

# Run GPU weightquant converted models:
@@ -2167,17 +2191,39 @@ function Run_gpu() {
if [[ $model_name == \#* ]]; then
continue
fi
model_name=`echo ${line} | awk -F ';' '{print $1}'`
accuracy_limit=`echo ${line} | awk -F ';' '{print $2}'`
input_num=`echo ${line} | awk -F ';' '{print $3}'`
input_shapes=`echo ${line} | awk -F ';' '{print $4}'`
input_files=""
data_path="/data/local/tmp/input_output/"
output_file=${data_path}'output/'${model_name}'.ms.out'
if [[ ${input_num} == "" || ${input_num} == 1 ]]; then
input_files=/data/local/tmp/input_output/input/${model_name}.ms.bin
else
for i in $(seq 1 $input_num)
do
input_files=$input_files${data_path}'input/'$model_name'.ms.bin_'$i','
done
fi
if [[ ${accuracy_limit} == "" ]]; then
accuracy_limit="5"
fi
echo ${model_name} >> "${run_gpu_log_file}"
echo 'cd /data/local/tmp/benchmark_test' > adb_run_cmd.txt
echo 'export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/data/local/tmp/benchmark_test;./benchmark --device=GPU --modelFile='${model_name}'_weightquant.ms --inDataFile=/data/local/tmp/input_output/input/'${model_name}'.ms.bin --benchmarkDataFile=/data/local/tmp/input_output/output/'${model_name}'.ms.out --enableFp16=true --accuracyThreshold=5' >> "${run_gpu_log_file}"
echo 'export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/data/local/tmp/benchmark_test;./benchmark --device=GPU --modelFile='${model_name}'_weightquant.ms --inDataFile=/data/local/tmp/input_output/input/'${model_name}'.ms.bin --benchmarkDataFile=/data/local/tmp/input_output/output/'${model_name}'.ms.out --enableFp16=true --accuracyThreshold=5' >> adb_run_cmd.txt
if [[ $input_shapes == "" ]]; then
echo 'export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/data/local/tmp/benchmark_test;./benchmark --device=GPU --enableFp16=true --accuracyThreshold='${accuracy_limit}' --modelFile='${model_name}'.ms --inDataFile='${input_files}' --benchmarkDataFile='${output_file} >> "${run_gpu_log_file}"
echo 'export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/data/local/tmp/benchmark_test;./benchmark --device=GPU --enableFp16=true --accuracyThreshold='${accuracy_limit}' --modelFile='${model_name}'.ms --inDataFile='${input_files}' --benchmarkDataFile='${output_file} >> adb_run_cmd.txt
else
echo 'export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/data/local/tmp/benchmark_test;./benchmark --device=GPU --enableFp16=true --inputShapes='${input_shapes}' --accuracyThreshold='${accuracy_limit}' --modelFile='${model_name}'.ms --inDataFile='${input_files}' --benchmarkDataFile='${output_file} >> "${run_gpu_log_file}"
echo 'export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/data/local/tmp/benchmark_test;./benchmark --device=GPU --enableFp16=true --inputShapes='${input_shapes}' --accuracyThreshold='${accuracy_limit}' --modelFile='${model_name}'.ms --inDataFile='${input_files}' --benchmarkDataFile='${output_file} >> adb_run_cmd.txt
fi
adb -s ${device_id} shell < adb_run_cmd.txt >> "${run_gpu_log_file}"
if [ $? = 0 ]; then
run_result='arm64_gpu_weightquant: '${model_name}' pass'; echo ${run_result} >> ${run_benchmark_result_file}
else
run_result='arm64_gpu_weightquant: '${model_name}' failed'; echo ${run_result} >> ${run_benchmark_result_file}; return 1
fi
#sleep 1
done < ${models_gpu_weightquant_config}
}



Write Preview
Loading…
Cancel
Save