DnnConv grad ops in gpuarray backend check shape consistency.

theano/gpuarray/dnn_gi.c

@@ -77,6 +77,48 @@ APPLY_SPECIFIC(conv_gi)(PyGpuArrayObject *kerns, PyGpuArrayObject *output,
 
   cuda_enter(c->ctx);
 
+  int expected_output_dims[5] = {0};
+  err = cudnnGetConvolutionNdForwardOutputDim(desc, APPLY_SPECIFIC(input), APPLY_SPECIFIC(kerns),
+                                              PyGpuArray_NDIM(im), expected_output_dims);
+  if (err != CUDNN_STATUS_SUCCESS) {
+    PyErr_Format(PyExc_RuntimeError, "error computing convolution output dim: %s",
+                 cudnnGetErrorString(err));
+    cuda_exit(c->ctx);
+    return 1;
+  }
+  if (PyGpuArray_NDIM(im) == 4) {
+    if ((PyGpuArray_DIMS(output)[0] != expected_output_dims[0]) ||
+        (PyGpuArray_DIMS(output)[1] != expected_output_dims[1]) ||
+        (PyGpuArray_DIMS(output)[2] != expected_output_dims[2]) ||
+        (PyGpuArray_DIMS(output)[3] != expected_output_dims[3])) {
+      PyErr_Format(PyExc_ValueError, "impossible convolution output dim: expected %ldx%ldx%ldx%ld"
+                                     " but received gradient with shape %ldx%ldx%ldx%ld",
+                   expected_output_dims[0], expected_output_dims[1],
+                   expected_output_dims[2], expected_output_dims[3],
+                   PyGpuArray_DIMS(output)[0], PyGpuArray_DIMS(output)[1],
+                   PyGpuArray_DIMS(output)[2], PyGpuArray_DIMS(output)[3]);
+      cuda_exit(c->ctx);
+      return 1;
+    }
+  } else if (PyGpuArray_NDIM(im) == 5) {
+    if ((PyGpuArray_DIMS(output)[0] != expected_output_dims[0]) ||
+        (PyGpuArray_DIMS(output)[1] != expected_output_dims[1]) ||
+        (PyGpuArray_DIMS(output)[2] != expected_output_dims[2]) ||
+        (PyGpuArray_DIMS(output)[3] != expected_output_dims[3]) ||
+        (PyGpuArray_DIMS(output)[4] != expected_output_dims[4])) {
+      PyErr_Format(PyExc_ValueError, "impossible convolution output dim: expected %ldx%ldx%ldx%ldx%ld"
+                                     " but received gradient with shape %ldx%ldx%ldx%ldx%ld",
+                   expected_output_dims[0], expected_output_dims[1],
+                   expected_output_dims[2], expected_output_dims[3],
+                   expected_output_dims[4],
+                   PyGpuArray_DIMS(output)[0], PyGpuArray_DIMS(output)[1],
+                   PyGpuArray_DIMS(output)[2], PyGpuArray_DIMS(output)[3],
+                   PyGpuArray_DIMS(output)[4]);
+      cuda_exit(c->ctx);
+      return 1;
+    }
+  }
+
 #ifdef CHOOSE_ALGO
 #ifndef CHOOSE_ONCE
   reuse_algo = 1;

theano/gpuarray/dnn_gw.c

@@ -77,6 +77,48 @@ APPLY_SPECIFIC(conv_gw)(PyGpuArrayObject *input, PyGpuArrayObject *output,
 
   cuda_enter(c->ctx);
 
+  int expected_output_dims[5] = {0};
+  err = cudnnGetConvolutionNdForwardOutputDim(desc, APPLY_SPECIFIC(input), APPLY_SPECIFIC(kerns),
+                                              PyGpuArray_NDIM(input), expected_output_dims);
+  if (err != CUDNN_STATUS_SUCCESS) {
+    PyErr_Format(PyExc_RuntimeError, "error computing convolution output dim: %s",
+                 cudnnGetErrorString(err));
+    cuda_exit(c->ctx);
+    return 1;
+  }
+  if (PyGpuArray_NDIM(input) == 4) {
+    if ((PyGpuArray_DIMS(output)[0] != expected_output_dims[0]) ||
+        (PyGpuArray_DIMS(output)[1] != expected_output_dims[1]) ||
+        (PyGpuArray_DIMS(output)[2] != expected_output_dims[2]) ||
+        (PyGpuArray_DIMS(output)[3] != expected_output_dims[3])) {
+      PyErr_Format(PyExc_ValueError, "impossible convolution output dim: expected %ldx%ldx%dx%ld"
+                                     " but received gradient with shape %ldx%ldx%dx%ld",
+                   expected_output_dims[0], expected_output_dims[1],
+                   expected_output_dims[2], expected_output_dims[3],
+                   PyGpuArray_DIMS(output)[0], PyGpuArray_DIMS(output)[1],
+                   PyGpuArray_DIMS(output)[2], PyGpuArray_DIMS(output)[3]);
+      cuda_exit(c->ctx);
+      return 1;
+    }
+  } else if (PyGpuArray_NDIM(input) == 5) {
+    if ((PyGpuArray_DIMS(output)[0] != expected_output_dims[0]) ||
+        (PyGpuArray_DIMS(output)[1] != expected_output_dims[1]) ||
+        (PyGpuArray_DIMS(output)[2] != expected_output_dims[2]) ||
+        (PyGpuArray_DIMS(output)[3] != expected_output_dims[3]) ||
+        (PyGpuArray_DIMS(output)[4] != expected_output_dims[4])) {
+      PyErr_Format(PyExc_ValueError, "impossible convolution output dim: expected %ldx%ldx%ldx%ldx%ld"
+                                     " but received gradient with shape %ldx%ldx%ldx%ldx%ld",
+                   expected_output_dims[0], expected_output_dims[1],
+                   expected_output_dims[2], expected_output_dims[3],
+                   expected_output_dims[4],
+                   PyGpuArray_DIMS(output)[0], PyGpuArray_DIMS(output)[1],
+                   PyGpuArray_DIMS(output)[2], PyGpuArray_DIMS(output)[3],
+                   PyGpuArray_DIMS(output)[4]);
+      cuda_exit(c->ctx);
+      return 1;
+    }
+  }
+
 #ifdef CHOOSE_ALGO
 #ifndef CHOOSE_ONCE
   reuse_algo = 1;

theano/gpuarray/tests/test_dnn.py

@@ -12,6 +12,7 @@ import theano.tensor as T
 import theano.tests.unittest_tools as utt
 from theano.tensor.signal.pool import pool_2d, pool_3d
 from theano.tensor.signal.pool import Pool, MaxPoolGrad, AveragePoolGrad
+from theano.tensor.nnet.abstract_conv import get_conv_output_shape
 
 from .. import dnn
 from ..basic_ops import GpuAllocEmpty
@@ -628,56 +629,50 @@ class TestDnnInferShapes(utt.InferShapeTester):
                         [(1, 1, 1), (2, 2, 2)],
                         'none')
 
-    def _test_conv_gradw(self, img, kerns, out, img_val, kern_vals, border_mode, conv_mode, subsample):
+    def _test_conv_gradw(self, img, topgrad, kerns, img_shape, kerns_shape, border_mode, conv_mode, subsample):
         if not dnn.dnn_available(test_ctx_name):
             raise SkipTest(dnn.dnn_available.msg)
 
+        topgrad_shape = get_conv_output_shape(img_shape, kerns_shape,
+                                              border_mode, subsample)
+
         img_val = numpy.asarray(
-            img_val,
+            numpy.random.rand(*img_shape),
             dtype=theano.config.floatX
         )
-        kern_vals = numpy.asarray(
-            kern_vals,
+        topgrad_vals = numpy.asarray(
+            numpy.random.rand(*topgrad_shape),
             dtype=theano.config.floatX
         )
 
-        temp_img = img.dimshuffle(1, 0, 2, 3)
-        temp_kerns = kerns
-        if conv_mode == 'conv':
-            temp_kerns = temp_kerns[:, :, ::-1, ::-1]
-        temp_kerns = temp_kerns.dimshuffle(1, 0, 2, 3)
-        shape = (
-            kern_vals.shape[1], img_val.shape[1],
-            img_val.shape[2] - kern_vals.shape[2] + 1,
-            img_val.shape[3] - kern_vals.shape[3] + 1
-        )
-        out_vals = numpy.zeros(shape, dtype=theano.config.floatX)
+        kerns_vals = numpy.zeros(kerns_shape, dtype=theano.config.floatX)
+        kerns_shape = theano.shared(numpy.asarray(kerns_shape))
         desc = dnn.GpuDnnConvDesc(
             border_mode=border_mode,
             subsample=subsample,
             conv_mode=conv_mode,
             precision=set_precision(theano.config.floatX)
-        )(out.shape)
+        )(kerns_shape)
         conv_grad_w = dnn.GpuDnnConvGradW()(
-            temp_img,
-            temp_kerns,
-            out,
+            img,
+            topgrad,
+            kerns,
             desc,
         )
         self._compile_and_check(
-            [temp_img, temp_kerns, out],
+            [img, topgrad, kerns],
             [conv_grad_w],
-            [img_val, kern_vals, out_vals],
+            [img_val, topgrad_vals, kerns_vals],
             dnn.GpuDnnConvGradW
         )
 
     @parameterized.expand(product(border_modes, conv_modes), utt.custom_name_func)
     def test_conv_gradw(self, border_mode, conv_mode):
         self._test_conv_gradw(T.tensor4('img'),
+                              T.tensor4('topgrad'),
                               T.tensor4('kerns'),
-                              T.tensor4('out'),
-                              numpy.random.rand(2, 5, 6, 8),
-                              numpy.random.rand(2, 1, 5, 6),
+                              (5, 2, 6, 13),
+                              (1, 2, 3, 7),
                               border_mode,
                               conv_mode,
                               (1, 1))