Try to fix conv output checking and ultimate jenkins tests.

theano/gpuarray/c_code/dnn_conv_base.c

@@ -22,7 +22,7 @@ static int c_check_groups_for_conv(cudnnConvolutionDescriptor_t desc, int groups
   }
   return 1;
 #else
-  return groups;  
+  return groups;
 #endif
 }
 
@@ -43,7 +43,7 @@ if ((APPLY_SPECIFIC(err) = cudnnCreateTensorDescriptor(&APPLY_SPECIFIC(output)))
   FAIL;
 }
 if ((APPLY_SPECIFIC(err) = cudnnCreateFilterDescriptor(&APPLY_SPECIFIC(kerns))) != CUDNN_STATUS_SUCCESS) {
-  PyErr_Format(PyExc_MemoryError, "could not allocate filter descriptor: %s", 
+  PyErr_Format(PyExc_MemoryError, "could not allocate filter descriptor: %s",
 	       cudnnGetErrorString(APPLY_SPECIFIC(err)));
   FAIL;
 }
@@ -84,10 +84,10 @@ static cudnnStatus_t checkCudnnStatus(cudnnStatus_t err, const char* msg)
 }
 
 static int
-c_get_largest_free_block_size(PyGpuContextObject *c) 
+c_get_largest_free_block_size(PyGpuContextObject *c)
 {
   size_t free = 0;
-  
+
   int err2 = gpucontext_property(c->ctx, GA_CTX_PROP_LARGEST_MEMBLOCK, &free);
   if (err2 != GA_NO_ERROR) {
     PyErr_Format(PyExc_RuntimeError, "Error when trying to find the "
@@ -98,12 +98,62 @@ c_get_largest_free_block_size(PyGpuContextObject *c)
   return free;
 }
 
+/** Check if convolution output tensor has expected dimensions
+    depending on given inputs and number of groups.
+    return 0 if everything is ok, non-0 on error.
+**/
+static int dnn_check_convolution_output(cudnnConvolutionDescriptor_t convDesc,
+                                        cudnnTensorDescriptor_t inputDesc,
+                                        cudnnFilterDescriptor_t filterDesc,
+                                        size_t tensorNdim,
+                                        PyGpuArrayObject* output,
+                                        int groups) {
+    int expected_output_dims[5] = {0};
+    cudnnStatus_t err = cudnnGetConvolutionNdForwardOutputDim(convDesc, inputDesc, filterDesc,
+                                                              tensorNdim, expected_output_dims);
+    if (err != CUDNN_STATUS_SUCCESS) {
+      PyErr_Format(PyExc_RuntimeError, "error computing convolution output dim: %s",
+                   cudnnGetErrorString(err));
+      return 1;
+    }
+    if (tensorNdim == 4) {
+      if ((PyGpuArray_DIMS(output)[0] != expected_output_dims[0]) ||
+          (PyGpuArray_DIMS(output)[1] / groups != 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 %dx%dx%dx%d"
+                     " but received %ldx%ldx%ldx%ld",
+                     expected_output_dims[0], expected_output_dims[1] * groups,
+                     expected_output_dims[2], expected_output_dims[3],
+                     PyGpuArray_DIMS(output)[0], PyGpuArray_DIMS(output)[1],
+                     PyGpuArray_DIMS(output)[2], PyGpuArray_DIMS(output)[3]);
+        return 1;
+      }
+    } else if (tensorNdim == 5) {
+      if ((PyGpuArray_DIMS(output)[0] != expected_output_dims[0]) ||
+          (PyGpuArray_DIMS(output)[1] / groups != 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 %dx%dx%dx%dx%d"
+                     " but received %ldx%ldx%ldx%ldx%ld",
+                     expected_output_dims[0], expected_output_dims[1] * groups,
+                     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]);
+        return 1;
+      }
+    }
+    return 0;
+}
 
 static std::string shape(int* res, int size)
 {
     std::ostringstream s;
     if (size > 0) {
-      
+
       s << res[0];
       for (int i = 1; i < size; ++i)
         s <<',' << res[i];
@@ -140,7 +190,7 @@ static std::string shape(cudnnConvolutionDescriptor_t convDesc)
     int nDim;
     cudnnConvolutionMode_t mode;
     cudnnDataType_t        computeType;
-    
+
     int                                 padA[5];
     int                                 strideA[5];
     int                                 dilationA[5];
@@ -155,7 +205,7 @@ static std::string shape(cudnnConvolutionDescriptor_t convDesc)
                                          &computeType ),
         "error getting convolution description");
     if (PyErr_Occurred()) return "";
-    
+
     return (std::string("-mode ") +
             ((mode == CUDNN_CONVOLUTION) ? "conv" : "cross") +
             " -pad " +
@@ -185,43 +235,8 @@ static std::string dnn_conv_shape(cudnnTensorDescriptor_t inputDesc, PyGpuArrayO
     cudnnDataType_t  dType;
     std::ostringstream s;
     int expected_output_dims[5] = {0};
-    cudnnStatus_t err = cudnnGetConvolutionNdForwardOutputDim(convDesc, inputDesc, filterDesc,
-							      PyGpuArray_NDIM(filter), expected_output_dims);
-    if (err != CUDNN_STATUS_SUCCESS) {
-      PyErr_Format(PyExc_RuntimeError, "error computing convolution output dim: %s",
-                   cudnnGetErrorString(err));
-      return "";
-    }
-    if (PyGpuArray_NDIM(filter) == 4) {
-      if ((PyGpuArray_DIMS(output)[0] != expected_output_dims[0]) ||
-          (PyGpuArray_DIMS(output)[1] / groups  != 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 %dx%dx%dx%d"
-                     " but received gradient with shape %ldx%ldx%ldx%ld",
-                     expected_output_dims[0], expected_output_dims[1] / groups,
-                     expected_output_dims[2], expected_output_dims[3],
-                     PyGpuArray_DIMS(output)[0], PyGpuArray_DIMS(output)[1],
-                     PyGpuArray_DIMS(output)[2], PyGpuArray_DIMS(output)[3]);
-        return "";
-      }
-    } else if (PyGpuArray_NDIM(filter) == 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 %dx%dx%dx%dx%d"
-                     " 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]);
+    if (dnn_check_convolution_output(convDesc, inputDesc, filterDesc, PyGpuArray_NDIM(filter), output, groups) != 0)
         return "";
-      }
-    }
     std::string shapeInput = shape(inputDesc);
     std::string shapeFilter = shape(filterDesc, &dType);
     std::string shapeConvDesc = shape(convDesc);
@@ -229,8 +244,8 @@ static std::string dnn_conv_shape(cudnnTensorDescriptor_t inputDesc, PyGpuArrayO
         return "";
     s << "-g " << groups << " -dim " << shapeInput << " -filt " <<
       shapeFilter << " " << shapeConvDesc;
-    
-// there have to be entries for both aligned and not
+
+    // there have to be entries for both aligned and not.
     if (!all_aligned(dType, PyGpuArray_DEV_DATA(input), PyGpuArray_DEV_DATA(output), PyGpuArray_DEV_DATA(filter)))
     {
       s << " [unaligned]";
@@ -240,18 +255,18 @@ static std::string dnn_conv_shape(cudnnTensorDescriptor_t inputDesc, PyGpuArrayO
 
 static void dnn_conv_update_cache(const std::string& hash, const AlgoRec& rec)
 {
-  pthread_mutex_lock(&algoMutex);    
+  pthread_mutex_lock(&algoMutex);
   algoCache[hash] = rec;
   pthread_mutex_unlock(&algoMutex);
 }
 
 static const AlgoRec* dnn_conv_check_cache(const std::string& hash)
 {
-  pthread_mutex_lock(&algoMutex);    
+  pthread_mutex_lock(&algoMutex);
   const AlgoRec* ret = 0;
-  
+
   AlgoCache::iterator hit = algoCache.find(hash);
-  
+
   if (hit != algoCache.end())
     ret = &hit->second;
 

theano/gpuarray/c_code/dnn_gi.c

@@ -141,6 +141,11 @@ APPLY_SPECIFIC(conv_gi)(PyGpuArrayObject *kerns, PyGpuArrayObject *output,
     return 1;
   if (c_set_tensor_for_conv(*input, APPLY_SPECIFIC(input), groups) == -1)
     return 1;
+
+  if (0 != dnn_check_convolution_output(desc, APPLY_SPECIFIC(input), APPLY_SPECIFIC(kerns),
+                                        PyGpuArray_NDIM(kerns), output, groups))
+    return 1;
+
   size_t input_offset = PyGpuArray_STRIDE(*input, 0) / groups;
   size_t kern_offset = PyGpuArray_STRIDE(kerns, 0) * PyGpuArray_DIM(kerns, 0) / groups;
   size_t output_offset = PyGpuArray_STRIDE(output, 0) / groups;

theano/gpuarray/c_code/dnn_gw.c

@@ -128,6 +128,11 @@ APPLY_SPECIFIC(conv_gw)(PyGpuArrayObject *input, PyGpuArrayObject *output,
     return 1;
   if (c_set_filter(*kerns, APPLY_SPECIFIC(kerns), groups) == -1)
     return 1;
+
+  if (0 != dnn_check_convolution_output(desc, APPLY_SPECIFIC(input), APPLY_SPECIFIC(kerns),
+                                        PyGpuArray_NDIM(*kerns), output, groups))
+    return 1;
+
   size_t input_offset = PyGpuArray_STRIDE(input, 0) / groups;
   size_t kern_offset = PyGpuArray_STRIDE(*kerns, 0) * PyGpuArray_DIM(*kerns, 0) / groups;
   size_t output_offset = PyGpuArray_STRIDE(output, 0) / groups;