Support V2 in GpuDnnConv

theano/sandbox/cuda/dnn.py

@@ -549,6 +549,18 @@ class GpuDnnConv(DnnBase, COp):
         self.inplace = inplace
         if self.inplace:
             self.destroy_map = {0: [2]}
+
+        # In CuDNN version older than V3, the FFT implementation and the
+        # option to time the different implementations to get the fastest
+        # are both unavailable.
+        if version() < (3000, 3000):
+            if self.workmem == 'fft':
+                raise RuntimeError("CuDNN's FFT convolution is only available "
+                                   "starting at CuDNN v3")
+            elif self.workmem == 'time':
+                raise RuntimeError("CuDNN's convolution timing option is only "
+                                   "available starting at CuDNN v3")
+
         assert self.workmem in ['none', 'small', 'large', 'fft', 'time',
                                 'guess']
 
@@ -564,11 +576,12 @@ class GpuDnnConv(DnnBase, COp):
             inpl_def = [('CONV_INPLACE', '1')]
         else:
             inpl_def = []
+
+        choose_alg = '0'
+        choose_alg_time = '0'
         if version() == -1:
-            alg_def = ('CONV_ALGO', "0")
+            alg = "0"
         else:
-            choose_alg = '0'
-            choose_alg_time = '0'
             if self.workmem == 'none':
                 alg = 'CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_GEMM'
             elif self.workmem == 'small':
@@ -589,9 +602,9 @@ class GpuDnnConv(DnnBase, COp):
                 choose_alg = '1'
                 choose_alg_time = '1'
 
-            alg_def = ('CONV_ALGO', alg)
-            alg_choose_def = ('CHOOSE_ALGO', choose_alg)
-            alg_choose_time_def = ('CHOOSE_ALGO_TIME', choose_alg_time)
+        alg_def = ('CONV_ALGO', alg)
+        alg_choose_def = ('CHOOSE_ALGO', choose_alg)
+        alg_choose_time_def = ('CHOOSE_ALGO_TIME', choose_alg_time)
 
         return [alg_def, alg_choose_def, alg_choose_time_def] + inpl_def
 

theano/sandbox/cuda/dnn_conv_base.c

@@ -6,13 +6,19 @@ cudnnFilterDescriptor_t APPLY_SPECIFIC(kerns);
 /* Keep track, from one execution to another, of the dimension of the data
 and the algorithms, if any, that were selected according to these dimensions
 and according to the amount of memory available at that time.
+
+Note : Implementation selection for backward convolution only exists starting
+at V3.
 */
 int APPLY_SPECIFIC(previous_input_shape)[5];
 int APPLY_SPECIFIC(previous_kerns_shape)[5];
 int APPLY_SPECIFIC(previous_output_shape)[5];
 cudnnConvolutionFwdAlgo_t APPLY_SPECIFIC(previous_algo);
+
+#if defined(CUDNN_VERSION) && CUDNN_VERSION >= 3000
 cudnnConvolutionBwdFilterAlgo_t APPLY_SPECIFIC(previous_bwd_f_algo);
 cudnnConvolutionBwdDataAlgo_t APPLY_SPECIFIC(previous_bwd_d_algo);
+#endif
 
 #section init_code_struct
 

theano/sandbox/cuda/dnn_fwd.c

@@ -73,6 +73,7 @@ APPLY_SPECIFIC(conv_fwd)(CudaNdarray *input, CudaNdarray *kerns,
         // CuDNN time every implementation and choose the best one.
         if (CHOOSE_ALGO_TIME)
         {
+#if defined(CUDNN_VERSION) && CUDNN_VERSION >= 3000
           // Time the different implementations to choose the best one
           int requestedCount = 1;
           int count;
@@ -93,6 +94,7 @@ APPLY_SPECIFIC(conv_fwd)(CudaNdarray *input, CudaNdarray *kerns,
           }
 
           chosen_algo = choosen_algo_perf.algo;
+#endif
         }
         else
         {
@@ -138,12 +140,14 @@ APPLY_SPECIFIC(conv_fwd)(CudaNdarray *input, CudaNdarray *kerns,
       chosen_algo = CONV_ALGO;
     }
 
-    // The FFT implementation does not support strides, 1x1 filters or
-    // inputs with a spatial dimension larger than 1024.
+    // The FFT implementation (only in V3 and onward) does not support strides,
+    // 1x1 filters or inputs with a spatial dimension larger than 1024.
     // If the chosen implementation is FFT, validate that it can be used
     // on the current data and default on a safe implementation if it
     // can't.
-    // Following code is 2d-specific, but it is fine as ftt is define only for 2d-filters
+    // Following code is 2d-specific, but it is fine as ftt is defined only for
+    // 2d-filters
+#if defined(CUDNN_VERSION) && CUDNN_VERSION >= 3000
     if (chosen_algo == CUDNN_CONVOLUTION_FWD_ALGO_FFT && nb_dim == 4)
     {
 
@@ -178,6 +182,7 @@ APPLY_SPECIFIC(conv_fwd)(CudaNdarray *input, CudaNdarray *kerns,
         chosen_algo = CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_GEMM;
       }
     }
+#endif
 
     err = cudnnGetConvolutionForwardWorkspaceSize(_handle,
                                                   APPLY_SPECIFIC(input),