From e06e9c7133620b354ee7ce51af3cd2961d1f2216 Mon Sep 17 00:00:00 2001 From: wandongdong Date: Thu, 27 Aug 2020 17:58:38 +0800 Subject: [PATCH] add print info for GPU subgraph --- mindspore/lite/src/lite_session.cc | 1 + .../kernel/opencl/cl/depthwise_conv2d.cl | 44 ++++++++++--------- .../kernel/opencl/kernel/conv2d_transpose.cc | 4 +- .../kernel/opencl/kernel/convolution.cc | 28 ++++++------ .../kernel/opencl/kernel/depthwise_conv2d.cc | 28 +++++++----- .../runtime/kernel/opencl/kernel/matmul.cc | 4 +- .../runtime/kernel/opencl/kernel/to_format.cc | 6 ++- .../lite/src/runtime/opencl/opencl_runtime.cc | 19 ++++---- mindspore/lite/src/scheduler.cc | 14 +++++- mindspore/lite/src/scheduler.h | 3 +- .../kernel/opencl/depthwise_conv2d_tests.cc | 13 +++--- 11 files changed, 94 insertions(+), 70 deletions(-) diff --git a/mindspore/lite/src/lite_session.cc b/mindspore/lite/src/lite_session.cc index 7c09737fad..a046988323 100644 --- a/mindspore/lite/src/lite_session.cc +++ b/mindspore/lite/src/lite_session.cc @@ -273,6 +273,7 @@ int LiteSession::Init(Context *context) { if (context_->device_ctx_.type == DT_GPU) { auto opencl_runtime = lite::opencl::OpenCLRuntime::GetInstance(); opencl_runtime->Init(); + MS_LOG(INFO) << "Init OpenCL runtime."; } #endif executor = new Executor(); diff --git a/mindspore/lite/src/runtime/kernel/opencl/cl/depthwise_conv2d.cl b/mindspore/lite/src/runtime/kernel/opencl/cl/depthwise_conv2d.cl index 17afaae4a9..0c1b83d444 100644 --- a/mindspore/lite/src/runtime/kernel/opencl/cl/depthwise_conv2d.cl +++ b/mindspore/lite/src/runtime/kernel/opencl/cl/depthwise_conv2d.cl @@ -1,8 +1,9 @@ #pragma OPENCL EXTENSION cl_khr_fp16 : enable __constant sampler_t smp_zero = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST; __kernel void DepthwiseConv2d_IMG_NC4HW4(__read_only image2d_t src_data, __global FLT4 *filter, __global FLT4 *bias, - float relu_clip, __write_only image2d_t dst_data, int2 kernel_size, - int2 stride, int2 padding, int2 dilation, int4 src_size, int4 dst_size) { + __write_only image2d_t dst_data, int2 kernel_size, + int2 stride, int2 padding, int2 dilation, int4 src_size, int4 dst_size, + float relu_clip_min, float relu_clip_max) { int X = get_global_id(0); int Y = get_global_id(1); int Z = get_global_id(2); @@ -27,13 +28,14 @@ __kernel void DepthwiseConv2d_IMG_NC4HW4(__read_only image2d_t src_data, __globa } FLT4 bias_p = bias[Z]; FLT4 res = TO_FLT4(r) + bias_p; - res = clamp(res, (FLT)(0.0f), (FLT)(relu_clip)); + res = clamp(res, (FLT)(relu_clip_min), (FLT)(relu_clip_max)); WRITE_IMAGE(dst_data, (int2)(X, (Z * dst_size.y + Y)), res); } __kernel void DepthwiseConv2d_IMG_NHWC4(__read_only image2d_t src_data, __global FLT4 *filter, __global FLT4 *bias, - float relu_clip, __write_only image2d_t dst_data, int2 kernel_size, - int2 stride, int2 padding, int2 dilation, int4 src_size, int4 dst_size) { + __write_only image2d_t dst_data, int2 kernel_size, + int2 stride, int2 padding, int2 dilation, int4 src_size, int4 dst_size, + float relu_clip_min, float relu_clip_max) { int X = get_global_id(0); int Y = get_global_id(1); int Z = get_global_id(2); @@ -58,13 +60,14 @@ __kernel void DepthwiseConv2d_IMG_NHWC4(__read_only image2d_t src_data, __global } FLT4 bias_p = bias[Z]; FLT4 res = TO_FLT4(r) + bias_p; - res = clamp(res, (FLT)(0.0f), (FLT)(relu_clip)); + res = clamp(res, (FLT)(relu_clip_min), (FLT)(relu_clip_max)); WRITE_IMAGE(dst_data, (int2)(X * dst_size.z + Z, Y), res); } __kernel void DepthwiseConv2d_IMG_NHWC4_1x1(__read_only image2d_t src_data, __global FLT4 *filter, __global FLT4 *bias, - float relu_clip, __write_only image2d_t dst_data, int2 kernel_size, - int2 stride, int2 padding, int2 dilation, int4 src_size, int4 dst_size) { + __write_only image2d_t dst_data, int2 kernel_size, int2 stride, + int2 padding, int2 dilation, int4 src_size, int4 dst_size, + float relu_clip_min, float relu_clip_max) { int X = get_global_id(0); int Y = get_global_id(1); int Z = get_global_id(2); @@ -81,7 +84,6 @@ __kernel void DepthwiseConv2d_IMG_NHWC4_1x1(__read_only image2d_t src_data, __gl bool outside_x = x_c < 0 || x_c >= src_size.x; if (!outside_x && !outside_y) { FLT4 flt_p = filter[fx_c]; - // FLT4 src_p =src_data[((y_c * src_size.x + x_c) * src_size.z + Z)]; FLT4 src_p = READ_IMAGE(src_data, smp_zero, (int2)(Z, (y_c * src_size.x + x_c) * src_size.z)); r += TO_FLT4(src_p * flt_p); } @@ -89,13 +91,13 @@ __kernel void DepthwiseConv2d_IMG_NHWC4_1x1(__read_only image2d_t src_data, __gl } FLT4 bias_p = bias[Z]; FLT4 res = TO_FLT4(r) + bias_p; - res = clamp(res, (FLT)(0.0f), (FLT)(relu_clip)); - // dst_data[((Y * dst_size.x + X) * dst_size.z + Z)] = res; + res = clamp(res, (FLT)(relu_clip_min), (FLT)(relu_clip_max)); WRITE_IMAGE(dst_data, (int2)(Z, (Y * dst_size.x + X) * dst_size.z), res); } __kernel void DepthwiseConv2d_BUF_NC4HW4(__global FLT4 *src_data, __global FLT4 *filter, __global FLT4 *bias, - float relu_clip, __global FLT4 *dst_data, int2 kernel_size, int2 stride, - int2 padding, int2 dilation, int4 src_size, int4 dst_size) { + __global FLT4 *dst_data, int2 kernel_size, int2 stride, + int2 padding, int2 dilation, int4 src_size, int4 dst_size, + float relu_clip_min, float relu_clip_max) { int X = get_global_id(0); int Y = get_global_id(1); int Z = get_global_id(2); @@ -120,13 +122,14 @@ __kernel void DepthwiseConv2d_BUF_NC4HW4(__global FLT4 *src_data, __global FLT4 } FLT4 bias_p = bias[Z]; FLT4 res = TO_FLT4(r) + bias_p; - res = clamp(res, (FLT)(0.0f), (FLT)(relu_clip)); + res = clamp(res, (FLT)(relu_clip_min), (FLT)(relu_clip_max)); dst_data[(((Z)*dst_size.y + (Y)) * dst_size.x + (X))] = res; } __kernel void DepthwiseConv2d_BUF_NHWC4(__global FLT4 *src_data, __global FLT4 *filter, __global FLT4 *bias, - float relu_clip, __global FLT4 *dst_data, int2 kernel_size, int2 stride, - int2 padding, int2 dilation, int4 src_size, int4 dst_size) { + __global FLT4 *dst_data, int2 kernel_size, int2 stride, + int2 padding, int2 dilation, int4 src_size, int4 dst_size, + float relu_clip_min, float relu_clip_max) { int X = get_global_id(0); int Y = get_global_id(1); int Z = get_global_id(2); @@ -151,13 +154,14 @@ __kernel void DepthwiseConv2d_BUF_NHWC4(__global FLT4 *src_data, __global FLT4 * } FLT4 bias_p = bias[Z]; FLT4 res = TO_FLT4(r) + bias_p; - res = clamp(res, (FLT)(0.0f), (FLT)(relu_clip)); + res = clamp(res, (FLT)(relu_clip_min), (FLT)(relu_clip_max)); dst_data[((Y * dst_size.x + X) * dst_size.z + Z)] = res; } __kernel void DepthwiseConv2d_BUF_NHWC4_1x1(__global FLT4 *src_data, __global FLT4 *filter, __global FLT4 *bias, - float relu_clip, __global FLT4 *dst_data, int2 kernel_size, int2 stride, - int2 padding, int2 dilation, int4 src_size, int4 dst_size) { + __global FLT4 *dst_data, int2 kernel_size, int2 stride, + int2 padding, int2 dilation, int4 src_size, int4 dst_size, + float relu_clip_min, float relu_clip_max) { int X = get_global_id(0); int Y = get_global_id(1); int Z = get_global_id(2); @@ -181,6 +185,6 @@ __kernel void DepthwiseConv2d_BUF_NHWC4_1x1(__global FLT4 *src_data, __global FL } FLT4 bias_p = bias[Z]; FLT4 res = TO_FLT4(r) + bias_p; - res = clamp(res, (FLT)(0.0f), (FLT)(relu_clip)); + res = clamp(res, (FLT)(relu_clip_min), (FLT)(relu_clip_max)); dst_data[((Y * dst_size.x + X) * dst_size.z + Z)] = res; } diff --git a/mindspore/lite/src/runtime/kernel/opencl/kernel/conv2d_transpose.cc b/mindspore/lite/src/runtime/kernel/opencl/kernel/conv2d_transpose.cc index c316350ac9..609ebc0dcc 100644 --- a/mindspore/lite/src/runtime/kernel/opencl/kernel/conv2d_transpose.cc +++ b/mindspore/lite/src/runtime/kernel/opencl/kernel/conv2d_transpose.cc @@ -186,8 +186,8 @@ int Conv2dTransposeOpenCLKernel::Run() { cl_int4 dst_size = {oh, ow, UP_DIV(co, C4NUM), 1}; int arg_cnt = 0; ocl_runtime->SetKernelArg(kernel_, arg_cnt++, in_tensors_[0]->Data()); - ocl_runtime->SetKernelArg(kernel_, arg_cnt++, padWeight_); - ocl_runtime->SetKernelArg(kernel_, arg_cnt++, bias_); + ocl_runtime->SetKernelArg(kernel_, arg_cnt++, padWeight_, lite::opencl::MemType::BUF); + ocl_runtime->SetKernelArg(kernel_, arg_cnt++, bias_, lite::opencl::MemType::BUF); ocl_runtime->SetKernelArg(kernel_, arg_cnt++, out_tensors_[0]->Data()); ocl_runtime->SetKernelArg(kernel_, arg_cnt++, kernel_size); ocl_runtime->SetKernelArg(kernel_, arg_cnt++, stride); diff --git a/mindspore/lite/src/runtime/kernel/opencl/kernel/convolution.cc b/mindspore/lite/src/runtime/kernel/opencl/kernel/convolution.cc index e75e15f676..8200549b6d 100644 --- a/mindspore/lite/src/runtime/kernel/opencl/kernel/convolution.cc +++ b/mindspore/lite/src/runtime/kernel/opencl/kernel/convolution.cc @@ -254,34 +254,34 @@ int ConvolutionOpenCLKernel::Run() { arg_cn = 0; cl_int4 _4x4to36_in_shape = {1, IH, IW, CI_SLICES}; cl_int4 _4x4to36_out_shape = {1, 36, TILES_XY, CI_SLICES}; - ocl_runtime->SetKernelArg(kernel_4x4to36, arg_cn++, in_tensors_[0]->Data()); - ocl_runtime->SetKernelArg(kernel_4x4to36, arg_cn++, winograd_mem0_); + ocl_runtime->SetKernelArg(kernel_4x4to36, arg_cn++, in_tensors_[0]->Data(), lite::opencl::MemType::IMG); + ocl_runtime->SetKernelArg(kernel_4x4to36, arg_cn++, winograd_mem0_, lite::opencl::MemType::IMG); ocl_runtime->SetKernelArg(kernel_4x4to36, arg_cn++, _4x4to36_in_shape); ocl_runtime->SetKernelArg(kernel_4x4to36, arg_cn++, _4x4to36_out_shape); arg_cn = 0; cl_int4 conv_in_shape = {1, 36, TILES_XY, CI_SLICES}; cl_int4 conv_out_shape = {1, 36, TILES_XY, CO_SLICES}; - ocl_runtime->SetKernelArg(kernel_conv, arg_cn++, winograd_mem0_); - ocl_runtime->SetKernelArg(kernel_conv, arg_cn++, winograd_mem1_); - ocl_runtime->SetKernelArg(kernel_conv, arg_cn++, packed_weight_); + ocl_runtime->SetKernelArg(kernel_conv, arg_cn++, winograd_mem0_, lite::opencl::MemType::IMG); + ocl_runtime->SetKernelArg(kernel_conv, arg_cn++, winograd_mem1_, lite::opencl::MemType::IMG); + ocl_runtime->SetKernelArg(kernel_conv, arg_cn++, packed_weight_, lite::opencl::MemType::BUF); ocl_runtime->SetKernelArg(kernel_conv, arg_cn++, conv_in_shape); ocl_runtime->SetKernelArg(kernel_conv, arg_cn++, conv_out_shape); arg_cn = 0; cl_int4 _36to4x4_in_shape = {1, 16, TILES_XY, CO_SLICES}; cl_int4 _36to4x4_out_shape = {1, OH, OW, CO_SLICES}; - ocl_runtime->SetKernelArg(kernel_36to4x4, arg_cn++, winograd_mem1_); - ocl_runtime->SetKernelArg(kernel_36to4x4, arg_cn++, out_tensors_[0]->Data()); - ocl_runtime->SetKernelArg(kernel_36to4x4, arg_cn++, packed_bias_); + ocl_runtime->SetKernelArg(kernel_36to4x4, arg_cn++, winograd_mem1_, lite::opencl::MemType::IMG); + ocl_runtime->SetKernelArg(kernel_36to4x4, arg_cn++, out_tensors_[0]->Data(), lite::opencl::MemType::IMG); + ocl_runtime->SetKernelArg(kernel_36to4x4, arg_cn++, packed_bias_, lite::opencl::MemType::BUF); ocl_runtime->SetKernelArg(kernel_36to4x4, arg_cn++, _36to4x4_in_shape); ocl_runtime->SetKernelArg(kernel_36to4x4, arg_cn++, _36to4x4_out_shape); } else { arg_cn = 0; - ocl_runtime->SetKernelArg(kernel_conv, arg_cn++, in_tensors_[0]->Data()); - ocl_runtime->SetKernelArg(kernel_conv, arg_cn++, out_tensors_[0]->Data()); - ocl_runtime->SetKernelArg(kernel_conv, arg_cn++, packed_weight_); - ocl_runtime->SetKernelArg(kernel_conv, arg_cn++, packed_bias_); + ocl_runtime->SetKernelArg(kernel_conv, arg_cn++, in_tensors_[0]->Data(), lite::opencl::MemType::IMG); + ocl_runtime->SetKernelArg(kernel_conv, arg_cn++, out_tensors_[0]->Data(), lite::opencl::MemType::IMG); + ocl_runtime->SetKernelArg(kernel_conv, arg_cn++, packed_weight_, lite::opencl::MemType::BUF); + ocl_runtime->SetKernelArg(kernel_conv, arg_cn++, packed_bias_, lite::opencl::MemType::BUF); } if (use_winograd_) { @@ -594,9 +594,9 @@ std::string ConvolutionOpenCLKernel::CodeGenWinograd36To4x4() { auto param = reinterpret_cast(op_parameter_); if (param->act_type_ == ActType_Relu) { - code += " acc = max(acc, (float4)(0.0f));\n"; + code += " acc = max(acc, (FLT4)(0.0f));\n"; } else if (param->act_type_ == ActType_Relu6) { - code += " acc = clamp(acc, (float4)(0.0f), (float4)(6.0f));\n"; + code += " acc = clamp(acc, (FLT4)(0.0f), (FLT4)(6.0f));\n"; } code += diff --git a/mindspore/lite/src/runtime/kernel/opencl/kernel/depthwise_conv2d.cc b/mindspore/lite/src/runtime/kernel/opencl/kernel/depthwise_conv2d.cc index 06e833c692..1ed841409c 100644 --- a/mindspore/lite/src/runtime/kernel/opencl/kernel/depthwise_conv2d.cc +++ b/mindspore/lite/src/runtime/kernel/opencl/kernel/depthwise_conv2d.cc @@ -15,6 +15,7 @@ */ #include "src/runtime/kernel/opencl/kernel/depthwise_conv2d.h" +#include #include #include #include @@ -180,7 +181,8 @@ int DepthwiseConv2dOpenCLKernel::Run() { std::vector local; GetLocalSize(0, global, &local); - float relu_clip1 = 6.0; + std::map> relu_clips{ + {ActType_No, {FLT_MIN, FLT_MAX}}, {ActType_Relu, {0.0, FLT_MAX}}, {ActType_Relu6, {0, 6.0}}}; cl_int2 kernel_size = {parameter->kernel_h_, parameter->kernel_w_}; cl_int2 stride = {parameter->stride_h_, parameter->stride_w_}; cl_int2 padding = {-parameter->pad_u_, -parameter->pad_l_}; @@ -189,17 +191,19 @@ int DepthwiseConv2dOpenCLKernel::Run() { cl_int4 dst_size = {(cl_int)out_tensors_[0]->Width(), (cl_int)out_tensors_[0]->Height(), (cl_int)CO4, (cl_int)out_tensors_[0]->Batch()}; - ocl_runtime->SetKernelArg(kernel_, 1, packed_weight_); - ocl_runtime->SetKernelArg(kernel_, 2, bias_data_); - ocl_runtime->SetKernelArg(kernel_, 3, relu_clip1); - ocl_runtime->SetKernelArg(kernel_, 5, kernel_size); - ocl_runtime->SetKernelArg(kernel_, 6, stride); - ocl_runtime->SetKernelArg(kernel_, 7, padding); - ocl_runtime->SetKernelArg(kernel_, 8, dilation); - ocl_runtime->SetKernelArg(kernel_, 9, src_size); - ocl_runtime->SetKernelArg(kernel_, 10, dst_size); - ocl_runtime->SetKernelArg(kernel_, 0, in_tensors_[0]->Data()); - ocl_runtime->SetKernelArg(kernel_, 4, out_tensors_[0]->Data()); + int arg_cnt = 0; + ocl_runtime->SetKernelArg(kernel_, arg_cnt++, in_tensors_[0]->Data()); + ocl_runtime->SetKernelArg(kernel_, arg_cnt++, packed_weight_, lite::opencl::MemType::BUF); + ocl_runtime->SetKernelArg(kernel_, arg_cnt++, bias_data_, lite::opencl::MemType::BUF); + ocl_runtime->SetKernelArg(kernel_, arg_cnt++, out_tensors_[0]->Data()); + ocl_runtime->SetKernelArg(kernel_, arg_cnt++, kernel_size); + ocl_runtime->SetKernelArg(kernel_, arg_cnt++, stride); + ocl_runtime->SetKernelArg(kernel_, arg_cnt++, padding); + ocl_runtime->SetKernelArg(kernel_, arg_cnt++, dilation); + ocl_runtime->SetKernelArg(kernel_, arg_cnt++, src_size); + ocl_runtime->SetKernelArg(kernel_, arg_cnt++, dst_size); + ocl_runtime->SetKernelArg(kernel_, arg_cnt++, relu_clips[parameter->act_type_].first); + ocl_runtime->SetKernelArg(kernel_, arg_cnt++, relu_clips[parameter->act_type_].second); ocl_runtime->RunKernel(kernel_, global, local, nullptr); return RET_OK; } diff --git a/mindspore/lite/src/runtime/kernel/opencl/kernel/matmul.cc b/mindspore/lite/src/runtime/kernel/opencl/kernel/matmul.cc index c379156e4c..0356963ad3 100644 --- a/mindspore/lite/src/runtime/kernel/opencl/kernel/matmul.cc +++ b/mindspore/lite/src/runtime/kernel/opencl/kernel/matmul.cc @@ -163,8 +163,8 @@ int MatMulOpenCLKernel::Run() { std::vector global = {UP_ROUND(sizeCO.s[1], local[0]), 4}; int arg_count = 0; ocl_runtime->SetKernelArg(kernel_, arg_count++, in_tensors_[0]->Data()); - ocl_runtime->SetKernelArg(kernel_, arg_count++, padWeight_); - ocl_runtime->SetKernelArg(kernel_, arg_count++, bias_); + ocl_runtime->SetKernelArg(kernel_, arg_count++, padWeight_, lite::opencl::MemType::BUF); + ocl_runtime->SetKernelArg(kernel_, arg_count++, bias_, lite::opencl::MemType::BUF); ocl_runtime->SetKernelArg(kernel_, arg_count++, out_tensors_[0]->Data()); ocl_runtime->SetKernelArg(kernel_, arg_count++, sizeCI); ocl_runtime->SetKernelArg(kernel_, arg_count++, sizeCO); diff --git a/mindspore/lite/src/runtime/kernel/opencl/kernel/to_format.cc b/mindspore/lite/src/runtime/kernel/opencl/kernel/to_format.cc index 7ac206df78..aeb3209fc1 100644 --- a/mindspore/lite/src/runtime/kernel/opencl/kernel/to_format.cc +++ b/mindspore/lite/src/runtime/kernel/opencl/kernel/to_format.cc @@ -147,8 +147,10 @@ int ToFormatOpenCLKernel::Run() { cl_int4 shape{(cl_int)nhwc_shape_[0], (cl_int)nhwc_shape_[1], (cl_int)nhwc_shape_[2], (cl_int)nhwc_shape_[3]}; cl_int4 gsize{(cl_int)global[0], (cl_int)global[1], (cl_int)global[2], 1}; - ocl_runtime->SetKernelArg(kernel_, 0, in_tensors_[0]->Data()); - ocl_runtime->SetKernelArg(kernel_, 1, out_tensors_[0]->Data()); + auto src_mem_type = (out_mem_type_ == OpenCLMemType::IMG) ? lite::opencl::MemType::BUF : lite::opencl::MemType::IMG; + auto dst_mem_type = (out_mem_type_ == OpenCLMemType::IMG) ? lite::opencl::MemType::IMG : lite::opencl::MemType::BUF; + ocl_runtime->SetKernelArg(kernel_, 0, in_tensors_[0]->Data(), src_mem_type); + ocl_runtime->SetKernelArg(kernel_, 1, out_tensors_[0]->Data(), dst_mem_type); ocl_runtime->SetKernelArg(kernel_, 2, gsize); ocl_runtime->SetKernelArg(kernel_, 3, shape); ocl_runtime->RunKernel(kernel_, global, local, nullptr); diff --git a/mindspore/lite/src/runtime/opencl/opencl_runtime.cc b/mindspore/lite/src/runtime/opencl/opencl_runtime.cc index 8d214d5a98..0f681c6ce0 100644 --- a/mindspore/lite/src/runtime/opencl/opencl_runtime.cc +++ b/mindspore/lite/src/runtime/opencl/opencl_runtime.cc @@ -275,20 +275,23 @@ int OpenCLRuntime::BuildKernel(cl::Kernel &kernel, const std::string &program_na // fp16 enable, kernel will use half and read_imageh and write_imageh. build_options_str = "-DFLT=half -DFLT4=half4 -DFLT16=half16 " - "-DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh -DTO_FLT4=convert_half4"; + "-DWRITE_IMAGE=write_imageh -DREAD_IMAGE=read_imageh -DTO_FLT4=convert_half4 "; } else { // fp16 not enable, kernel will use float and read_imagef and write_imagef. build_options_str = "-DFLT=float -DFLT4=float4 -DFLT16=float16 " - "-DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef -DTO_FLT4=convert_float4"; + "-DWRITE_IMAGE=write_imagef -DREAD_IMAGE=read_imagef -DTO_FLT4=convert_float4 "; } - build_options_str = std::accumulate( - build_options.begin(), build_options.end(), build_options_str, - [](const std::string &options, const std::string &option) -> std::string { return options + " " + option; }); + auto build_options_ext = std::accumulate( + build_options.begin(), build_options.end(), std::string(""), + [](const std::string &options, const std::string &option) -> std::string { + auto res = options + " " + option; + return res; + }); build_options_str += default_build_opts_; // program identifier = program_name + build_options - std::string build_program_key = program_name + build_options_str; + std::string build_program_key = program_name + build_options_str + build_options_ext; auto build_program_it = program_map_.find(build_program_key); cl::Program program; @@ -409,9 +412,7 @@ GpuInfo OpenCLRuntime::ParseGpuInfo(std::string device_name, std::string device_ bool OpenCLRuntime::LoadSource(const std::string &program_name, const std::string &source) { auto it_source = g_opencl_program_map.find(program_name); - if (it_source != g_opencl_program_map.end()) { - it_source->second = source; - } else { + if (it_source == g_opencl_program_map.end()) { g_opencl_program_map.emplace(program_name, source); } return true; diff --git a/mindspore/lite/src/scheduler.cc b/mindspore/lite/src/scheduler.cc index 7b42ee7316..250ab9b610 100644 --- a/mindspore/lite/src/scheduler.cc +++ b/mindspore/lite/src/scheduler.cc @@ -16,6 +16,8 @@ #include "src/scheduler.h" #include +#include +#include #include "include/errorcode.h" #include "src/kernel_registry.h" #include "src/common/graph_util.h" @@ -140,7 +142,7 @@ int Scheduler::InitOp2Kernel(const lite::Model *model, std::vectorat(size_t(outIndexes->GetAs(j)))); } auto *primitive = model->GetOp(cNode->name()->str()); - auto *kernel = this->ScheduleNode(inputs, outputs, primitive); + auto *kernel = this->ScheduleNode(inputs, outputs, primitive, cNode); if (nullptr == kernel) { MS_LOG(ERROR) << "ScheduleNode return nullptr, name: " << cNode->name()->str() << ", type: " << schema::EnumNamePrimitiveType(cNode->primitive()->value_type()); @@ -176,6 +178,7 @@ void Scheduler::ConstructSubgraphs(std::vector *kernels) { } std::vector subgraph_kernels; + size_t sub_cnt{0}; for (auto temp_kernels : sub_kernels_list) { kernel::KERNEL_ARCH arch = temp_kernels.front()->desc().arch; if (arch == kernel::KERNEL_ARCH::kCPU) { @@ -194,7 +197,10 @@ void Scheduler::ConstructSubgraphs(std::vector *kernels) { } else { auto subgraph_kernel = CreateSubKernel(temp_kernels, arch); subgraph_kernels.emplace_back(subgraph_kernel); + std::string arch_name = (arch == kernel::KERNEL_ARCH::kGPU) ? "GPU" : "NPU"; + MS_LOG(INFO) << arch_name << " subgraph id" << sub_cnt << " created."; } + ++sub_cnt; } kernels->clear(); kernels->insert(kernels->begin(), subgraph_kernels.begin(), subgraph_kernels.end()); @@ -223,7 +229,7 @@ kernel::LiteKernel *Scheduler::CreateSubKernel(const std::vector &in_tensors, const std::vector &out_tensors, - const mindspore::lite::PrimitiveC *primitive) { + const mindspore::lite::PrimitiveC *primitive, const schema::CNode *cnode) { MS_ASSERT(nullptr != primitive); auto data_type = in_tensors.front()->data_type(); kernel::KernelKey desc{kernel::KERNEL_ARCH::kCPU, data_type, static_cast(primitive->Type())}; @@ -233,6 +239,10 @@ kernel::LiteKernel *Scheduler::ScheduleNode(const std::vector if (nullptr != kernel) { kernel->set_desc(desc); return kernel; + } else { + MS_LOG(ERROR) << "Not supported GPU Op " + << schema::EnumNamePrimitiveType(static_cast(primitive->Type())) << " " + << (cnode->name()->str()); } } diff --git a/mindspore/lite/src/scheduler.h b/mindspore/lite/src/scheduler.h index 0815d9c6be..aa39383c07 100644 --- a/mindspore/lite/src/scheduler.h +++ b/mindspore/lite/src/scheduler.h @@ -35,7 +35,8 @@ class Scheduler { protected: kernel::LiteKernel *ScheduleNode(const std::vector &in_tensors, const std::vector &out_tensors, - const mindspore::lite::PrimitiveC *primitive); + const mindspore::lite::PrimitiveC *primitive, + const schema::CNode *cnode); private: int InitOp2Kernel(const lite::Model *model, std::vector *tensors, diff --git a/mindspore/lite/test/ut/src/runtime/kernel/opencl/depthwise_conv2d_tests.cc b/mindspore/lite/test/ut/src/runtime/kernel/opencl/depthwise_conv2d_tests.cc index 693f6b97c1..7ea580d672 100644 --- a/mindspore/lite/test/ut/src/runtime/kernel/opencl/depthwise_conv2d_tests.cc +++ b/mindspore/lite/test/ut/src/runtime/kernel/opencl/depthwise_conv2d_tests.cc @@ -207,7 +207,7 @@ TEST_F(TestConvolutionDwOpenCL, NoPadNC4HW4Fp32) { float gnd_data[] = {3.3848767, 1.4446403, 1.8428744, 1.3194335, 2.5873442, 2.1384869, 2.04022, 1.1872686, 2.2294958, 1.6570128, 2.465089, 1.4294086, 2.7941442, 1.7871612, 2.188921, 1.0601988}; - DepthWiseTestMain(conv_param.get(), input_data, weight_data, gnd_data, schema::Format_NC4HW4); + DepthWiseTestMain(conv_param.release(), input_data, weight_data, gnd_data, schema::Format_NC4HW4); } TEST_F(TestConvolutionDwOpenCL, PadNC4HW4Fp32) { @@ -279,7 +279,7 @@ TEST_F(TestConvolutionDwOpenCL, PadNC4HW4Fp32) { 0.8749627, 0.8953936, 0.5093431, 1.5496738, 0.54936385, 0.7683113, 1.165742, 1.3682933, 1.0517888, 0.59817517, 0.75649744, 1.2075498, 0.38804203}; - DepthWiseTestMain(conv_param.get(), input_data, weight_data, gnd_data, schema::Format_NC4HW4); + DepthWiseTestMain(conv_param.release(), input_data, weight_data, gnd_data, schema::Format_NC4HW4); } TEST_F(TestConvolutionDwOpenCL, NoPadNHWC4Fp32) { @@ -324,7 +324,8 @@ TEST_F(TestConvolutionDwOpenCL, NoPadNHWC4Fp32) { float gnd_data[] = {3.3848767, 1.4446403, 1.8428744, 1.3194335, 2.5873442, 2.1384869, 2.04022, 1.1872686, 2.2294958, 1.6570128, 2.465089, 1.4294086, 2.7941442, 1.7871612, 2.188921, 1.0601988}; - DepthWiseTestMain(conv_param.get(), input_data, weight_data, gnd_data, schema::Format_NHWC4); + DepthWiseTestMain(conv_param.release(), input_data, weight_data, gnd_data, schema::Format_NHWC4); + // delete conv_param; } TEST_F(TestConvolutionDwOpenCL, PadNHWC4Fp32) { @@ -396,7 +397,7 @@ TEST_F(TestConvolutionDwOpenCL, PadNHWC4Fp32) { 0.8749627, 0.8953936, 0.5093431, 1.5496738, 0.54936385, 0.7683113, 1.165742, 1.3682933, 1.0517888, 0.59817517, 0.75649744, 1.2075498, 0.38804203}; - DepthWiseTestMain(conv_param.get(), input_data, weight_data, gnd_data, schema::Format_NHWC4); + DepthWiseTestMain(conv_param.release(), input_data, weight_data, gnd_data, schema::Format_NHWC4); } TEST_F(TestConvolutionDwOpenCL, NoPadNHWC4Fp16) { @@ -443,7 +444,7 @@ TEST_F(TestConvolutionDwOpenCL, NoPadNHWC4Fp16) { 2.2294958, 1.6570128, 2.465089, 1.4294086, 2.7941442, 1.7871612, 2.188921, 1.0601988}; lite::opencl::OpenCLRuntime::GetInstance()->SetFp16Enable(true); - DepthWiseTestMain(conv_param.get(), input_data, weight_data, gnd_data, schema::Format_NHWC4, + DepthWiseTestMain(conv_param.release(), input_data, weight_data, gnd_data, schema::Format_NHWC4, kNumberTypeFloat16, true, 1e-2); } @@ -517,7 +518,7 @@ TEST_F(TestConvolutionDwOpenCL, PadNHWC4Fp16) { 1.0517888, 0.59817517, 0.75649744, 1.2075498, 0.38804203}; lite::opencl::OpenCLRuntime::GetInstance()->SetFp16Enable(true); - DepthWiseTestMain(conv_param.get(), input_data, weight_data, gnd_data, schema::Format_NHWC4, + DepthWiseTestMain(conv_param.release(), input_data, weight_data, gnd_data, schema::Format_NHWC4, kNumberTypeFloat16, true, 1e-2); }