Make dnn grad convolutions compatible with both V2 and V3

theano/sandbox/cuda/dnn.py

@@ -833,8 +833,9 @@ class GpuDnnConvGradW(DnnBase, COp):
         else:
             inplace_def = []
 
-        if version() == -1:
+        if version() == -1 or version() < (3000, 3000):
             alg_def = ('CONV_ALGO', '0')
+            alg_choose_def = ('CHOOSE_ALGO', '0')
         else:
             if self.workmem == 'none':
                 alg_def = ('CONV_ALGO', 'CUDNN_CONVOLUTION_BWD_FILTER_ALGO_0')
@@ -989,8 +990,9 @@ class GpuDnnConvGradI(DnnBase, COp):
         else:
             inplace_def = []
 
-        if version() == -1:
+        if version() == -1 or version() < (3000, 3000):
             alg_def = ('CONV_ALGO', '0')
+            alg_choose_def = ('CHOOSE_ALGO', '0')
         else:
             if self.workmem == 'none':
                 alg_def = ('CONV_ALGO', 'CUDNN_CONVOLUTION_BWD_DATA_ALGO_0')

theano/sandbox/cuda/dnn_conv_base.c

@@ -52,8 +52,10 @@ for (int i = 0; i < 5; i++)
 // Select default implementations for the case where the convolution
 // implementations should be selected based on the size of the data.
 APPLY_SPECIFIC(previous_algo) = CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM;
-APPLY_SPECIFIC(previous_bwd_f_algo) = CUDNN_CONVOLUTION_BWD_FILTER_ALGO_0;
-APPLY_SPECIFIC(previous_bwd_d_algo) = CUDNN_CONVOLUTION_BWD_DATA_ALGO_0;
+#if defined(CUDNN_VERSION) && CUDNN_VERSION >= 3000
+APPLY_SPECIFIC(previous_bwd_f_algo) = CUDNN_CONVOLUTION_BWD_FILTER_ALGO_1;
+APPLY_SPECIFIC(previous_bwd_d_algo) = CUDNN_CONVOLUTION_BWD_DATA_ALGO_1;
+#endif
 
 #section cleanup_code_struct
 

theano/sandbox/cuda/dnn_gi.c

@@ -31,6 +31,7 @@ APPLY_SPECIFIC(conv_gi)(CudaNdarray *kerns, CudaNdarray *output,
   if (c_set_tensorNd(*input, nb_dim, APPLY_SPECIFIC(input)) == -1)
     return 1;
 
+#if defined(CUDNN_VERSION) && CUDNN_VERSION >= 3000
   {
     size_t worksize;
     void *workspace;
@@ -180,6 +181,16 @@ APPLY_SPECIFIC(conv_gi)(CudaNdarray *kerns, CudaNdarray *output,
       (void *)&beta,
       APPLY_SPECIFIC(input), CudaNdarray_DEV_DATA(*input));
   }
+#else
+  err = cudnnConvolutionBackwardData(
+    _handle,
+    (void *)&alpha,
+    APPLY_SPECIFIC(kerns), CudaNdarray_DEV_DATA(kerns),
+    APPLY_SPECIFIC(output), CudaNdarray_DEV_DATA(output),
+    desc,
+    (void *)&beta,
+    APPLY_SPECIFIC(input), CudaNdarray_DEV_DATA(*input));
+#endif
 
   if (err != CUDNN_STATUS_SUCCESS) {
     PyErr_Format(PyExc_RuntimeError, "GpuDnnConvGradI: error doing operation: %s",

theano/sandbox/cuda/dnn_gw.c

@@ -31,6 +31,7 @@ APPLY_SPECIFIC(conv_gw)(CudaNdarray *input, CudaNdarray *output,
   if (c_set_filterNd(*kerns, nb_dim, APPLY_SPECIFIC(kerns)) == -1)
     return 1;
 
+#if defined(CUDNN_VERSION) && CUDNN_VERSION >= 3000
   {
     size_t worksize;
     void *workspace;
@@ -181,6 +182,16 @@ APPLY_SPECIFIC(conv_gw)(CudaNdarray *input, CudaNdarray *output,
       APPLY_SPECIFIC(kerns), CudaNdarray_DEV_DATA(*kerns));
 
   }
+#else
+  err = cudnnConvolutionBackwardFilter(
+    _handle,
+    (void *)&alpha,
+    APPLY_SPECIFIC(input), CudaNdarray_DEV_DATA(input),
+    APPLY_SPECIFIC(output), CudaNdarray_DEV_DATA(output),
+    desc,
+    (void *)&beta,
+    APPLY_SPECIFIC(kerns), CudaNdarray_DEV_DATA(*kerns));
+#endif
 
   if (err != CUDNN_STATUS_SUCCESS) {
     PyErr_Format(PyExc_RuntimeError, "GpuDnnConvGradW: error doing operation: %s",