Merge pull request #5614 from lamblin/fix_5558

theano/gpuarray/dnn.py

@@ -733,6 +733,13 @@ class GpuDnnConvGradW(DnnBase):
                           'with certain cuDNN algorithms depending on the compute capability of your GPU '
                           'if subsample is not (1, 1, 1). If you encounter problems, consider '
                           'setting the theano flag "dnn.conv.algo_bwd_filter" to "none".')
+        if (self.algo not in ('none', 'deterministic', 'fft', 'small') and
+                beta is not None and theano.tensor.extract_constant(beta) != 1):
+            warnings.warn('cuDNN backward filter operation for convolutions may produce bad results '
+                          'with certain cuDNN algorithms depending on the compute capability of your GPU '
+                          'if beta != 1. If you encounter problems, consider '
+                          'setting the theano flag "dnn.conv.algo_bwd_filter" to '
+                          '"none", "deterministic", "fft", or "small".')
         ctx_name = infer_context_name(img, topgrad, output)
         img = as_gpuarray_variable(img, ctx_name)
         topgrad = as_gpuarray_variable(topgrad, ctx_name)

theano/gpuarray/dnn_gi.c

@@ -167,7 +167,7 @@ APPLY_SPECIFIC(conv_gi)(PyGpuArrayObject *kerns, PyGpuArrayObject *output,
     err = cudnnFindConvolutionBackwardDataAlgorithmEx(
       _handle, APPLY_SPECIFIC(kerns), PyGpuArray_DEV_DATA(kerns),
       APPLY_SPECIFIC(output), PyGpuArray_DEV_DATA(output), desc,
-      APPLY_SPECIFIC(input), PyGpuArray_DEV_DATA(input),
+      APPLY_SPECIFIC(input), PyGpuArray_DEV_DATA(*input),
       1, &count, &choice, *(void **)tmpmem, free);
     gpudata_release(tmpmem);
 

theano/gpuarray/dnn_gw.c

@@ -34,7 +34,7 @@ APPLY_SPECIFIC(conv_gw)(PyGpuArrayObject *input, PyGpuArrayObject *output,
 
   if (PyGpuArray_DIMS(input)[1] != PyGpuArray_DIMS(km)[1]) {
     PyErr_SetString(PyExc_ValueError,
-		    "GpuDnnConv images and kernel must have the same stack size");
+                    "GpuDnnConv images and kernel must have the same stack size");
     return 1;
   }
 
@@ -165,8 +165,10 @@ APPLY_SPECIFIC(conv_gw)(PyGpuArrayObject *input, PyGpuArrayObject *output,
     }
 
     err = cudnnFindConvolutionBackwardFilterAlgorithmEx(
-      _handle, APPLY_SPECIFIC(input), APPLY_SPECIFIC(output), desc,
-      APPLY_SPECIFIC(kerns), 1, &count, &choice, *(void **)tmpmem, free);
+      _handle, APPLY_SPECIFIC(input), PyGpuArray_DEV_DATA(input),
+      APPLY_SPECIFIC(output), PyGpuArray_DEV_DATA(output), desc,
+      APPLY_SPECIFIC(kerns), PyGpuArray_DEV_DATA(*kerns),
+      1, &count, &choice, *(void **)tmpmem, free);
     gpudata_release(tmpmem);
 
     if (err != CUDNN_STATUS_SUCCESS) {