fix(mgb): fix bug caused by conv filter size is too big

GitOrigin-RevId: 7fe3a9fcf1

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;
     }

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)                                       \