Browse Source

!15914 fix opencl tanh bug & add ci model

From: @yeyunpeng2020
Reviewed-by: @zhanghaibo5,@ddwsky
Signed-off-by: @zhanghaibo5
pull/15914/MERGE
mindspore-ci-bot Gitee 5 years ago
parent
commit
150066d00c
3 changed files with 64 additions and 8 deletions
  1. +2
    -6
      mindspore/lite/src/runtime/kernel/opencl/cl/fullconnection.cl
  2. +2
    -2
      mindspore/lite/src/runtime/kernel/opencl/cl/softmax.cl
  3. +60
    -0
      mindspore/lite/test/models_gpu_fp32.cfg

+ 2
- 6
mindspore/lite/src/runtime/kernel/opencl/cl/fullconnection.cl View File

@@ -33,9 +33,7 @@ __kernel void FullConnection(__read_only image2d_t input, __write_only image2d_t
} else if (act_type == ActivationType_RELU6) {
result = clamp(result, (FLT4)(0.0f), (FLT4)(6.0f));
} else if (act_type == ActivationType_TANH) {
FLT4 exp0 = exp(result);
FLT4 exp1 = exp(-result);
result = (exp0 - exp1) / (exp0 + exp1);
result = tanh(result);
}
WRITE_IMAGE(output, (int2)(gidx, gidz), result);
}
@@ -75,9 +73,7 @@ __kernel void FullConnectionWeightVar(__read_only image2d_t input, __write_only
} else if (act_type == ActivationType_RELU6) {
result = clamp(result, (FLT4)(0.0f), (FLT4)(6.0f));
} else if (act_type == ActivationType_TANH) {
FLT4 exp0 = exp(result);
FLT4 exp1 = exp(-result);
result = (exp0 - exp1) / (exp0 + exp1);
result = tanh(result);
}
WRITE_IMAGE(output, (int2)(gidx, gidz), result);
}


+ 2
- 2
mindspore/lite/src/runtime/kernel/opencl/cl/softmax.cl View File

@@ -36,14 +36,14 @@ __kernel void SoftmaxAxis3_NHWC4(__read_only image2d_t input, __write_only image
sum += dot(exp(t - input_max_f4), (float4)(1.f));
}
float4 t = convert_float4(READ_IMAGE(input, smp_zero, (int2)(Y * C4 + C4 - 1, n * H + X)));
sum += dot(exp(min(t - input_max_f4, 0)), mask);
sum += dot(exp(min(t - input_max_f4, (float4)(0.f))), mask);
for (int d = 0; d < C4 - 1; ++d) {
float4 result = convert_float4(READ_IMAGE(input, smp_zero, (int2)(Y * C4 + d, n * H + X)));
result = exp(result - input_max_f4) / sum;
WRITE_IMAGEOUT(output, (int2)(Y * C4 + d, n * H + X), OUT_FLT4(result));
}
float4 result = convert_float4(READ_IMAGE(input, smp_zero, (int2)(Y * C4 + C4 - 1, n * H + X)));
result = exp(min(result - input_max_f4, 0)) / sum;
result = exp(min(result - input_max_f4, (float4)(0.f))) / sum;
result = result * mask;
WRITE_IMAGEOUT(output, (int2)(Y * C4 + C4 - 1, n * H + X), OUT_FLT4(result));
}


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

@@ -199,3 +199,63 @@ mtk_AADB_HADB_MBV2_model_f16.tflite
mtk_AADB_HADB_MBV3_model_f16.tflite
mtk_model_emotions_0725_fp16.tflite
mtk_face_features_v1_fp16.tflite
Q888_age_gender_orderd.tflite
emotion
gender_res_large_deploy
glasses
hat
isface
ml_bank_detect_0312_tmp
ml_face_div_parsing
ml_hardware_eyeclose
Mnet6_0312_extract_pay
pose_3d
hiai_face_RFB-Epoch-170-no-transpose
tracking
detect-deeper-halfdeeper-mbv1-shortcut-400-400_nopostprocess_simplified
hiai_face_detect_rfb
hiai_face_isface
hiai_face_landmark
hiai_face_pose_tuku
ml_hand_detection
ml_ocr_sfz_detect_0325_tmp
ml_hardware_liveness
ml_liveness_detect_landmark_tmp
ml_face_contour
2012_ATLANTA_1class_20190621_v4.x_nomean
ml_ocr_sfz_add_final_0325
ml_hardware_pose
ml_bank_recog
2012_ATLANTA_10class_20190131_v4.0
mnet
recognition
ml_face_landmark
model_hebing_3branch
hiai_cv_focusShootOCRModel_07
hiai_cv_focusShootOCRModel_03
hiai_cv_focusShootOCRModel_01
hiai_face_hat1
hiai_cv_focusShootOCRModel_04
hiai_cv_focusShootOCRModel_06
hiai_cpu_face_hat
hiai_video_seg
hiai_semantic_seg
hiai_human_seg
hiai_face_recognition_1
hiai_cpu_face_detect
hiai_cpu_face_attr
hiai_face_attr1
retinaface
deconvs_model
ml_location_scene_division
ml_tabel_recog
6c_seg_nomean_20200610
ml_video_edit_img_segment
ml_video_edit_video_segment_gauss_adaptis_part1
ml_video_edit_Mnet
ml_video_edit_detect
ml_video_edit_MnetN367_extract_1010_pay
ml_video_edit_person_divison_pic
ml_video_edit_reid
ml_video_edit_v10_best_model_nomean_20200723
inception-v2-9.onnx

Loading…
Cancel
Save