Merge pull request #6368 from notoraptor/fix-cudnn-conv-timed-algos-with-beta

theano/gpuarray/c_code/dnn_fwd.c

@@ -234,14 +234,24 @@ APPLY_SPECIFIC(conv_fwd)(PyGpuArrayObject *input, PyGpuArrayObject *kerns,
         if (input->ga.typecode == GA_HALF)
           c_set_math_type_for_conv(desc, CUDNN_TENSOR_OP_MATH);
 
+        /* cudnnFindConvolutionForwardAlgorithmEx() may write to output.
+           We don't want that if output is used in computation (ie. if beta != 0). */
+        PyGpuArrayObject* o = *output;
+        if (beta != 0) {
+            o = pygpu_empty(PyGpuArray_NDIM(*output), PyGpuArray_DIMS(*output), (*output)->ga.typecode, GA_C_ORDER, c, Py_None);
+        }
+
         // We don't sync the buffer as we don't care about the values.
         err = cudnnFindConvolutionForwardAlgorithmEx(
           params->handle, APPLY_SPECIFIC(input), PyGpuArray_DEV_DATA(input),
           APPLY_SPECIFIC(kerns), PyGpuArray_DEV_DATA(kerns),
-          desc, APPLY_SPECIFIC(output), PyGpuArray_DEV_DATA(*output),
+          desc, APPLY_SPECIFIC(output), PyGpuArray_DEV_DATA(o),
           1, &count, &choice, *(void **)tmpmem,
           maxfree);
         gpudata_release(tmpmem);
+        if (beta != 0) {
+            Py_XDECREF(o);
+        }
 
         if (err != CUDNN_STATUS_SUCCESS) {
           PyErr_Format(PyExc_RuntimeError,

theano/gpuarray/c_code/dnn_gi.c

@@ -204,12 +204,22 @@ APPLY_SPECIFIC(conv_gi)(PyGpuArrayObject *kerns, PyGpuArrayObject *output,
           return -1;
         }
 
+        /* cudnnFindConvolutionBackwardDataAlgorithmEx() may write to output (input).
+           We don't want that if output is used in computation (ie. if beta != 0). */
+        PyGpuArrayObject* ip = *input;
+        if (beta != 0) {
+            ip = pygpu_empty(PyGpuArray_NDIM(*input), PyGpuArray_DIMS(*input), (*input)->ga.typecode, GA_C_ORDER, c, Py_None);
+        }
+
         err = cudnnFindConvolutionBackwardDataAlgorithmEx(
           params->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(ip),
           1, &count, &choice, *(void **)tmpmem, maxfree);
         gpudata_release(tmpmem);
+        if (beta != 0) {
+            Py_XDECREF(ip);
+        }
 
         if (err != CUDNN_STATUS_SUCCESS) {
           PyErr_Format(PyExc_RuntimeError, "error selecting convolution algo: %s",

theano/gpuarray/c_code/dnn_gw.c

@@ -191,12 +191,22 @@ APPLY_SPECIFIC(conv_gw)(PyGpuArrayObject *input, PyGpuArrayObject *output,
           return -1;
         }
 
+        /* cudnnFindConvolutionBackwardFilterAlgorithmEx() may write to kernels output (kerns).
+           We don't want that if output is used in computation (ie. if beta != 0). */
+        PyGpuArrayObject* k = *kerns;
+        if (beta != 0) {
+            k = pygpu_empty(PyGpuArray_NDIM(*kerns), PyGpuArray_DIMS(*kerns), (*kerns)->ga.typecode, GA_C_ORDER, c, Py_None);
+        }
+
         err = cudnnFindConvolutionBackwardFilterAlgorithmEx(
           params->handle, APPLY_SPECIFIC(input), PyGpuArray_DEV_DATA(input),
           APPLY_SPECIFIC(output), PyGpuArray_DEV_DATA(output), desc,
-          APPLY_SPECIFIC(kerns), PyGpuArray_DEV_DATA(*kerns),
+          APPLY_SPECIFIC(kerns), PyGpuArray_DEV_DATA(k),
           1, &count, &choice, *(void **)tmpmem, maxfree);
         gpudata_release(tmpmem);
+        if (beta != 0) {
+            Py_XDECREF(k);
+        }
 
         if (err != CUDNN_STATUS_SUCCESS) {
           PyErr_Format(PyExc_RuntimeError,