From 5af52746f7d2c4752c2eeab77e9d3f84d87ff185 Mon Sep 17 00:00:00 2001 From: Megvii Engine Team Date: Sun, 26 Sep 2021 19:51:34 +0800 Subject: [PATCH] fix(mgb): fix bug caused by conv filter size is too big GitOrigin-RevId: 7fe3a9fcf12e61e8d4a23d5052b953b907c3ef4a --- dnn/src/cuda/convolution/backward_data/chanwise.cpp | 6 ++++++ dnn/src/cuda/utils.cuh | 18 ++++++++++++------ 2 files changed, 18 insertions(+), 6 deletions(-) diff --git a/dnn/src/cuda/convolution/backward_data/chanwise.cpp b/dnn/src/cuda/convolution/backward_data/chanwise.cpp index d9bbc1a7..d156f605 100644 --- a/dnn/src/cuda/convolution/backward_data/chanwise.cpp +++ b/dnn/src/cuda/convolution/backward_data/chanwise.cpp @@ -20,6 +20,12 @@ using namespace convolution; bool ConvolutionBackwardDataImpl::AlgoChanwise::is_available( const SizeArgs& args) const { + auto kparam = chanwise::Param::from_fwd_args(args.as_fwd_args()); + auto&& device_prop = cuda::current_device_prop(); + if (device_prop.sharedMemPerBlock < + kparam.chl_mul * kparam.flt_h * kparam.flt_w * args.diff_layout->dtype.size()) { + return false; + } if (!args.grad_layout->is_contiguous() || !args.diff_layout->is_contiguous()) { return false; } diff --git a/dnn/src/cuda/utils.cuh b/dnn/src/cuda/utils.cuh index a4b93860..732f531a 100644 --- a/dnn/src/cuda/utils.cuh +++ b/dnn/src/cuda/utils.cuh @@ -24,12 +24,18 @@ #include "src/cuda/atomic_add.cuh" #include "src/cuda/cudnn_with_check.h" -#define cuda_check(_x) \ - do { \ - cudaError_t _err = (_x); \ - if (_err != cudaSuccess) { \ - ::megdnn::cuda::__throw_cuda_error__(_err, #_x); \ - } \ +#define cuda_check(_x) \ + do { \ + cudaError_t _err = (_x); \ + if (_err != cudaSuccess) { \ + std::string x = std::string(#_x); \ + char line[10]; \ + sprintf(line, "%d", __LINE__); \ + ::megdnn::cuda::__throw_cuda_error__( \ + _err, (x + " error file:" + std::string(__FILE__) + ":" + \ + std::string(line)) \ + .c_str()); \ + } \ } while (0) #define cublas_check(_x) \