Add attributes for implementation selection in the gradient convolution

theano/sandbox/cuda/dnn_conv_base.c

@@ -3,14 +3,16 @@ cudnnTensorDescriptor_t APPLY_SPECIFIC(input);
 cudnnTensorDescriptor_t APPLY_SPECIFIC(output);
 cudnnFilterDescriptor_t APPLY_SPECIFIC(kerns);
 
-/* Keep track, from one execution to another, of the dimension of the inputs
-and the algorithm, if any, that was selected according to these dimensions
-and the amount of memory available at that time.
+/* 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.
 */
 int APPLY_SPECIFIC(previous_input_shape)[4];
 int APPLY_SPECIFIC(previous_kerns_shape)[4];
+int APPLY_SPECIFIC(previous_output_shape)[4];
 cudnnConvolutionFwdAlgo_t APPLY_SPECIFIC(previous_algo);
-
+cudnnConvolutionBwdFilterAlgo_t APPLY_SPECIFIC(previous_bwd_f_algo);
+cudnnConvolutionBwdDataAlgo_t APPLY_SPECIFIC(previous_bwd_d_algo);
 
 #section init_code_struct
 
@@ -34,15 +36,18 @@ if ((APPLY_SPECIFIC(err) = cudnnCreateFilterDescriptor(&APPLY_SPECIFIC(kerns)))
   FAIL;
 }
 
-APPLY_SPECIFIC(previous_input_shape)[0] = 0;
-APPLY_SPECIFIC(previous_input_shape)[1] = 0;
-APPLY_SPECIFIC(previous_input_shape)[2] = 0;
-APPLY_SPECIFIC(previous_input_shape)[3] = 0;
-APPLY_SPECIFIC(previous_kerns_shape)[0] = 0;
-APPLY_SPECIFIC(previous_kerns_shape)[1] = 0;
-APPLY_SPECIFIC(previous_kerns_shape)[2] = 0;
-APPLY_SPECIFIC(previous_kerns_shape)[3] = 0;
+for (int i = 0; i < 4; i++)
+{
+  APPLY_SPECIFIC(previous_input_shape)[i] = 0;
+  APPLY_SPECIFIC(previous_kerns_shape)[i] = 0;
+  APPLY_SPECIFIC(previous_output_shape)[i] = 0;
+}
+
+// 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;
 
 #section cleanup_code_struct