p83571920
/
mindspore
forked from mindspore-Ecosystem/mindspore
1
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

fix opencl winograd fp16 bug

This commit is contained in:
wangdongxu 2021-01-29 10:53:27 +08:00
parent e2907c1280
commit e3387e65e4
2 changed files with 8 additions and 4 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
mindspore/lite/src/runtime/kernel/opencl/cl/winograd.cl
View File

@ -52,6 +52,10 @@ __kernel void Winograd4x4To36(__read_only image2d_t input, // height=N*H
for (int x = 0; x < 6; x++) {
acc += BtD_row[x] * Bt[y * 6 + x];
}
#if FP16_ENABLE
acc = min(acc, HALF_MAX);
acc = max(acc, -HALF_MAX);
#endif
WRITE_IMAGE(output, (int2)(tile_hw, y_idx + y), acc);
}
}

8
mindspore/lite/src/runtime/opencl/opencl_runtime.cc
View File

@ -367,12 +367,12 @@ int OpenCLRuntime::BuildKernel(cl::Kernel &kernel, const std::string &program_na
std::string build_option = default_build_option_;
if (fp16_enable_) {
build_option +=
" -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";
" -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";
} else {
build_option +=
" -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";
" -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";
}
build_option =
std::accumulate(build_options_ext.begin(), build_options_ext.end(), build_option,