Make GpuDnnConv support inplace operation.

theano/sandbox/cuda/cudnn_helper.h

@@ -103,11 +103,11 @@ cudnnConvolutionForward_v2(
   const cudnnTensorDescriptor_t destDesc,
   void *destData) {
   assert(*(float *)alpha == 1.0);
-  assert(*(float *)beta == 0.0);
+  assert(*(float *)beta == 1.0);
   return cudnnConvolutionForward(handle, srcDesc, srcData,
 				 filterDesc, filterData,
 				 convDesc, destDesc, destData,
-				 CUDNN_RESULT_NO_ACCUMULATE);
+				 CUDNN_RESULT_ACCUMULATE);
 }
 #define cudnnConvolutionForward cudnnConvolutionForward_v2
 
@@ -124,11 +124,11 @@ cudnnConvolutionBackwardFilter_v2(
   const cudnnFilterDescriptor_t gradDesc,
   void *gradData) {
   assert(*(float *)alpha == 1.0);
-  assert(*(float *)beta == 0.0);
+  assert(*(float *)beta == 1.0);
   return cudnnConvolutionBackwardFilter(handle, srcDesc, srcData,
 					diffDesc, diffData,
 					convDesc, gradDesc, gradData,
-					CUDNN_RESULT_NO_ACCUMULATE);
+					CUDNN_RESULT_ACCUMULATE);
 }
 
 #define cudnnConvolutionBackwardFilter cudnnConvolutionBackwardFilter_v2
@@ -146,7 +146,7 @@ cudnnConvolutionBackwardData_v2(
   const cudnnTensorDescriptor_t gradDesc,
   void *gradData) {
   assert(*(float *)alpha == 1.0);
-  assert(*(float *)beta == 0.0);
+  assert(*(float *)beta == 1.0);
   return cudnnConvolutionBackwardData(handle,
 				      (cudnnFilterDescriptor_t)filterDesc,
 				      filterData,
@@ -155,7 +155,7 @@ cudnnConvolutionBackwardData_v2(
 				      (cudnnConvolutionDescriptor_t)convDesc,
 				      (cudnnTensorDescriptor_t)gradDesc,
 				      gradData,
-				      CUDNN_RESULT_NO_ACCUMULATE);
+				      CUDNN_RESULT_ACCUMULATE);
 }
 
 #define cudnnConvolutionBackwardData cudnnConvolutionBackwardData_v2

theano/sandbox/cuda/dnn.py

@@ -2,7 +2,7 @@ import os
 import numpy
 
 import theano
-from theano import Apply, gof, tensor, config
+from theano import Apply, gof, tensor, config, Variable
 from theano.scalar import as_scalar, constant
 from theano.gradient import DisconnectedType
 from theano.gof import Optimizer, local_optimizer, COp
@@ -16,7 +16,8 @@ from theano.sandbox.cuda.type import CudaNdarrayType
 from theano.sandbox.cuda import GpuOp
 from theano.sandbox.cuda.basic_ops import (as_cuda_ndarray_variable,
                                            gpu_contiguous, HostFromGpu,
-                                           cp_on_negative_strides)
+                                           cp_on_negative_strides,
+                                           gpu_alloc)
 from theano.sandbox.cuda.blas import (GpuConv, GpuDownsampleFactorMax,
                                       GpuDownsampleFactorMaxGrad)
 from theano.sandbox.cuda.nnet import GpuSoftmax
@@ -344,9 +345,9 @@ _one = constant(numpy.asarray(1.0, dtype='float32'))
 def ensure_float(val, default, name):
     if val is None:
         return default.clone()
-    if not isinstnace(val, Variable):
+    if not isinstance(val, Variable):
         val = constant(val)
-    if not isisntance(val.type, theano.scalar.Scalar):
+    if not isinstance(val.type, theano.scalar.Scalar):
         raise TypeError("%s: expected a scalar value" % (name,))
     if not val.type.dtype == 'float32':
         raise TypeError("%s: type is not float32" % (name,))
@@ -361,9 +362,9 @@ class GpuDnnConv(DnnBase, COp):
     :param kernel:
     :param descr: the convolution descriptor
     """
-    __props__ = ('workmem',)
+    __props__ = ('workmem', 'inplace')
 
-    def __init__(self, workmem=None):
+    def __init__(self, workmem=None, inplace=False):
         """
         :param workmem: either 'none', 'small' or 'large'.  Default is
         the value of :attr:`config.dnn.conv.workmem`.
@@ -373,92 +374,105 @@ class GpuDnnConv(DnnBase, COp):
         if workmem is None:
             workmem = config.dnn.conv.workmem
         self.workmem = workmem
+        self.inplace = inplace
+        if self.inplace:
+            self.destroy_map = {0: [2]}
         assert self.workmem in ['none', 'small', 'large']
 
     def __setstate__(self, d):
         self.__dict__.update(d)
         if not hasattr(self, 'workmem'):
-            self.workmem = 'small'
+            self.workmem = 'none'
+        if not hasattr(self, 'inplace'):
+            self.inplace = False
 
     def get_op_params(self):
+        if self.inplace:
+            inpl_def = [('CONV_INPLACE', '1')]
+        else:
+            inpl_def = []
         if version() == -1:
-            return [('CONV_ALGO', "0")]
-        if self.workmem == 'none':
-            alg = 'CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_GEMM'
-        elif self.workmem == 'small':
-            alg = 'CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM'
-        elif self.workmem == 'large':
-            alg = 'CUDNN_CONVOLUTION_FWD_ALGO_GEMM'
-        return [('CONV_ALGO', alg)]
-
-    def make_node(self, img, kern, desc, alpha=None, beta=None):
+            alg_def = ('CONV_ALGO', "0")
+        else:
+            if self.workmem == 'none':
+                alg = 'CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_GEMM'
+            elif self.workmem == 'small':
+                alg = 'CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM'
+            elif self.workmem == 'large':
+                alg = 'CUDNN_CONVOLUTION_FWD_ALGO_GEMM'
+            alg_def = ('CONV_ALGO', alg)
+        return [alg_def] + inpl_def
+
+    def make_node(self, img, kern, output, desc, alpha=None):
         img = as_cuda_ndarray_variable(img)
         kern = as_cuda_ndarray_variable(kern)
+        output = as_cuda_ndarray_variable(output)
         if img.type.ndim != 4:
             raise TypeError('img must be 4D tensor')
         if kern.type.ndim != 4:
             raise TypeError('kern must be 4D tensor')
+        if output.type.ndim != 4:
+            raise TypeError('output must be a 4D tensor')
 
         if not isinstance(desc.type, CDataType) \
                 or desc.type.ctype != 'cudnnConvolutionDescriptor_t':
             raise TypeError('desc must be cudnnConvolutionDescriptor_t')
 
         alpha = ensure_float(alpha, _one, 'alpha')
-        beta = ensure_float(beta, _zero, 'beta')
 
-        broadcastable = (img.type.broadcastable[0],
-                         kern.type.broadcastable[0],
-                         False, False)
-        return Apply(self, [img, kern, desc, alpha, beta],
-                     [CudaNdarrayType(broadcastable)()])
+        return Apply(self, [img, kern, output, desc, alpha],
+                     [output.type()])
 
     def grad(self, inp, grads):
-        img, kerns, desc, alpha, beta = inp
+        img, kerns, output, desc, alpha = inp
         top, = grads
 
         top = cp_on_negative_strides(top)
 
-        d_img = GpuDnnConvGradI()(kerns, top, desc,
-                                  img.shape[2], img.shape[3])
-        d_kerns = GpuDnnConvGradW()(img, top, desc,
-                                    kerns.shape[2], kerns.shape[3])
+        d_img = GpuDnnConvGradI()(kerns, top, img.zeros_like(), desc)
+        d_kerns = GpuDnnConvGradW()(img, top, kerns.zeros_like(), desc)
 
-        return [d_img, d_kerns, DisconnectedType()(), DisconnectedType()(),
-                DisconnectedType()()]
+        return [d_img, d_kerns, output.zeros_like(),
+                DisconnectedType()(), DisconnectedType()()]
 
     def connection_pattern(self, node):
-        # not connected to desc, alpha, beta
-        return [[1], [1], [0], [0], [0]]
+        # not connected to desc, alpha
+        return [[1], [1], [1], [0], [0]]
 
-    def infer_shape(self, node, shape):
-        b = shape[0][0]  # Number of inputs
-        h = shape[0][2]  # Height of input feature maps
-        w = shape[0][3]  # Width of input feature maps
-        nb = shape[1][0]  # Number of output feature maps
-        kh = shape[1][2]  # Height of each filter
-        kw = shape[1][3]  # Width of each filter
-        padh = 0
-        padw = 0
-        if (
-            not node.inputs[2].owner
-            or not isinstance(node.inputs[2].owner.op, GpuDnnConvDesc)
-        ):
-            raise theano.tensor.basic.ShareError("case not implemented and probably not needed")
-        desc = node.inputs[2].owner.op
-        sh, sw = desc.subsample
-        if desc.border_mode == 'full':
+    @staticmethod
+    def get_out_shape(ishape, kshape, border_mode, subsample):
+        """
+        This function computes the output shape for a convolution with
+        the specified parameters.  `ishape` and `kshape` can be symbolic
+        or scalar.
+        """
+        b = ishape[0]  # Number of inputs
+        h = ishape[2]  # Height of input feature maps
+        w = ishape[3]  # Width of input feature maps
+        nb = kshape[0]  # Number of output feature maps
+        kh = kshape[2]  # Height of each filter
+        kw = kshape[3]  # Width of each filter
+
+        sh, sw = subsample
+        if border_mode == 'full':
             padh = kh - 1
             padw = kw - 1
-        elif isinstance(desc.border_mode, tuple):
-            padh, padw = desc.border_mode
+        elif isinstance(border_mode, tuple):
+            padh, padw = border_mode
         else:
-            assert desc.border_mode == 'valid'
+            assert border_mode == 'valid'
+            padh = 0
+            padw = 0
 
-        return [(
+        return (
             b, nb,
             (h + 2*padh - kh)//sh + 1,
             (w + 2*padw - kw)//sw + 1
-        )]
+        )
+
+
+    def infer_shape(self, node, shape):
+        return [shape[2]]
 
 
 class GpuDnnConvGradW(DnnBase, COp):
@@ -470,62 +484,64 @@ class GpuDnnConvGradW(DnnBase, COp):
     :param descr: the convolution descriptor
 
     """
-    __props__ = ()
+    __props__ = ('inplace',)
 
-    def __init__(self):
+    def __init__(self, inplace=False):
         COp.__init__(self, ["dnn_base.c", "dnn_conv_base.c", "dnn_gw.c"],
                      "APPLY_SPECIFIC(conv_gw)")
+        self.inplace = inplace
+        if self.inplace:
+            self.destroy_map = {0: [2]}
+
+    def __setstate__(self, d):
+        self.__dict__.update(d)
+        if not hasattr(self, 'inplace'):
+            self.inplace = False
 
     def grad(self, inp, grads):
-        img, top, desc, h, w, alpha, beta = inp
+        img, top, output, desc, alpha = inp
         kerns, = grads
 
         kerns = gpu_contiguous(kerns)
 
-        d_img = GpuDnnConvGradI()(kerns, top, desc,
-                                  img.shape[2], img.shape[3])
-        d_top = GpuDnnConv()(img, kerns, desc)
+        d_img = GpuDnnConvGradI()(kerns, top, img.zeros_like(), desc)
+        d_top = GpuDnnConv()(img, kerns, top.zeros_like(), desc)
 
-        return (d_img, d_top, DisconnectedType()(), DisconnectedType()(),
-                DisconnectedType()(), DiconnnectedType()(),
-                DisconnectedType()())
+        return (d_img, d_top, output.zeros_like(),
+                DisconnectedType()(), DiconnnectedType()())
 
     def connection_pattern(self, node):
-        # not connected to desc, h, w, alpha, beta
-        return [[1], [1], [0], [0], [0], [0], [0]]
+        # not connected to desc, alpha
+        return [[1], [1], [1], [0], [0]]
 
-    def make_node(self, img, topgrad, desc, h, w, alpha=None, beta=None):
+    def get_op_params(self):
+        if self.inplace:
+            return [('CONV_INPLACE', '1')]
+        else:
+            return []
+
+    def make_node(self, img, topgrad, output, desc, alpha=None):
         img = as_cuda_ndarray_variable(img)
         topgrad = as_cuda_ndarray_variable(topgrad)
+        output = as_cuda_ndarray_variable(output)
         if img.type.ndim != 4:
             raise TypeError('img must be 4D tensor')
         if topgrad.type.ndim != 4:
             raise TypeError('topgrad must be 4D tensor')
+        if output.type.ndim != 4:
+            raise TypeError('output must be 4D tensor')
 
         if not isinstance(desc.type, CDataType) \
                 or desc.type.ctype != 'cudnnConvolutionDescriptor_t':
             raise TypeError('desc must be cudnnConvolutionDescriptor_t')
 
-        h = as_scalar(h)
-        w = as_scalar(w)
-
         alpha = ensure_float(alpha, _one, 'alpha')
-        beta = ensure_float(beta, _zero, 'beta')
 
-        broadcastable = [topgrad.type.broadcastable[1],
-                         img.type.broadcastable[1],
-                         False, False]
-
-        return Apply(self, [img, topgrad, desc, h, w, alpha, beta],
-                     [CudaNdarrayType(broadcastable)()])
+        return Apply(self, [img, topgrad, output, desc, alpha],
+                     [output.type()])
 
     def infer_shape(self, node, shape):
-        return [(
-            shape[1][1],
-            shape[0][1],
-            node.inputs[3],
-            node.inputs[4]
-        )]
+        return [shape[2]]
 
 
 class GpuDnnConvGradI(DnnBase, COp):
@@ -537,61 +553,58 @@ class GpuDnnConvGradI(DnnBase, COp):
     :param descr: the convolution descriptor
 
     """
-    __props__ = ()
+    __props__ = ('inplace',)
 
-    def __init__(self):
+    def __init__(self, inplace=False):
         COp.__init__(self, ["dnn_base.c", "dnn_conv_base.c", "dnn_gi.c"],
                      "APPLY_SPECIFIC(conv_gi)")
+        self.inplace = inplace
+        if self.inplace:
+            self.destroy_map = {0: [2]}
 
     def grad(self, inp, grads):
-        kerns, top, desc, h, w, alpha, beta = inp
+        kerns, top, output, desc, alpha = inp
         img, = grads
 
         img = cp_on_negative_strides(img)
 
-        d_kerns = GpuDnnConvGradW()(img, top, desc,
-                                    kerns.shape[2], kerns.shape[3])
-        d_top = GpuDnnConv()(img, kerns, desc)
-        return (d_kerns, d_top, DisconnectedType()(), DisconnectedType()(),
-                DisconnectedType()(), DisconnectedType()(),
-                DisconnectedType()())
+        d_kerns = GpuDnnConvGradW()(img, top, kerns.zeros_like(), desc)
+        d_top = GpuDnnConv()(img, kerns, top.zeros_like(), desc)
+        return (d_kerns, d_top, output.zeros_like(),
+                DisconnectedType()(), DisconnectedType()())
 
     def connection_pattern(self, node):
-        # not connected to desc, h, w, alpha, beta
-        return [[1], [1], [0], [0], [0], [0], [0]]
+        # not connected to desc, alpha
+        return [[1], [1], [1], [0], [0]]
 
-    def make_node(self, kern, topgrad, desc, h, w, alpha=None, beta=None):
+    def get_op_params(self):
+        if self.inplace:
+            return [('CONV_INPLACE', '1')]
+        else:
+            return []
+
+    def make_node(self, kern, topgrad, output, desc, alpha=None):
         kern = as_cuda_ndarray_variable(kern)
         topgrad = as_cuda_ndarray_variable(topgrad)
+        output = as_cuda_ndarray_variable(output)
         if kern.type.ndim != 4:
             raise TypeError('kern must be 4D tensor')
         if topgrad.type.ndim != 4:
             raise TypeError('topgrad must be 4D tensor')
+        if output.type.ndim != 4:
+            raise TypeError('output must be 4D tensor')
 
         if not isinstance(desc.type, CDataType) \
                 or desc.type.ctype != 'cudnnConvolutionDescriptor_t':
             raise TypeError('desc must be cudnnConvolutionDescriptor_t')
 
-        h = as_scalar(h)
-        w = as_scalar(w)
-
         alpha = ensure_float(alpha, _one, 'alpha')
-        beta = ensure_float(beta, _zero, 'beta')
-
-        broadcastable = [topgrad.type.broadcastable[0],
-                         kern.type.broadcastable[1],
-                         False, False]
 
-        return Apply(self, [kern, topgrad, desc, h, w, alpha, beta],
-                     [CudaNdarrayType(broadcastable)()])
+        return Apply(self, [kern, topgrad, output, desc, alpha],
+                     [output.type()])
 
     def infer_shape(self, node, shape):
-        return [(
-            shape[1][0],
-            shape[0][1],
-            node.inputs[3],
-            node.inputs[4]
-        )]
+        return [shape[2]]
 
 
 def dnn_conv(img, kerns, border_mode='valid', subsample=(1, 1),
@@ -620,32 +633,31 @@ def dnn_conv(img, kerns, border_mode='valid', subsample=(1, 1),
     :param workmem: Specify the amount of working memory allowed.
       More memory is usually faster.  One of 'none', 'small' or
       'large'.  (default is None which takes its value from
-      config.dnn.conv.workmem)
+      :attr:`config.dnn.conv.workmem`)
 
 
     :warning: The cuDNN library only works with GPU that have a compute
       capability of 3.0 or higer.  This means that older GPU will not
       work with this Op.
-    :note: The working memory of the op is influenced by
-      :attr:`config.dnn.conv.workmem`.
     """
     fgraph = getattr(img, 'fgraph', None) or getattr(kerns, 'fgraph', None)
     if (border_mode == 'valid' and subsample == (1,1) and
         direction_hint == 'bprop weights'):
         # Special case: We are asked to use GpuDnnConvGradW. We need to set
         # up a suitable 'fake' convolution to compute the gradient for.
-        img = gpu_contiguous(img.dimshuffle(1, 0, 2, 3))
+        img = cp_on_negative_strides(img.dimshuffle(1, 0, 2, 3))
         if conv_mode == 'conv':
             # We need to flip manually. These 'kerns' are not the kernels
             # that would be flipped by conv_mode='conv' in GpuDnnConvGradW.
             kerns = kerns[:, :, ::-1, ::-1]
         kerns = gpu_contiguous(kerns.dimshuffle(1, 0, 2, 3))
-        shape = theano.tensor.stack(kerns.shape[1], img.shape[1],
-                                    img.shape[2] - kerns.shape[2] + 1,
-                                    img.shape[3] - kerns.shape[3] + 1)
+        shape2 = shape_i(img, 2, fgraph) - shape_i(kerns, 2, fgraph) + 1
+        shape3 = shape_i(img, 3, fgraph) - shape_i(kerns, 3, fgraph) + 1
+        out = gpu_alloc(_zero.clone(), shape_i(kerns, 1, fgraph),
+                        shape_i(img, 1, fgraph), shape2, shape3)
         desc = GpuDnnConvDesc(border_mode='valid', subsample=(1, 1),
-                              conv_mode='cross')(img.shape, shape)
-        conv = GpuDnnConvGradW()(img, kerns, desc, shape[2], shape[3])
+                              conv_mode='cross')(img.shape, out.shape)
+        conv = GpuDnnConvGradW()(img, kerns, out, desc)
         return as_cuda_ndarray_variable(conv.dimshuffle(1, 0, 2, 3))
 
     elif (border_mode == 'full' and subsample == (1, 1) and
@@ -653,17 +665,16 @@ def dnn_conv(img, kerns, border_mode='valid', subsample=(1, 1),
         # Special case: We can be faster by using GpuDnnConvGradI to compute
         # the full convolution as the backward pass of a valid convolution.
         # We just need to set up a suitable 'fake' valid convolution.
-        img = gpu_contiguous(img)
+        img = cp_on_negative_strides(img)
         kerns = gpu_contiguous(kerns.dimshuffle(1, 0, 2, 3))
         conv_mode = 'cross' if conv_mode == 'conv' else 'conv'
         shape2 = shape_i(img, 2, fgraph) + shape_i(kerns, 2, fgraph) - 1
         shape3 = shape_i(img, 3, fgraph) + shape_i(kerns, 3, fgraph) - 1
-        shape = theano.tensor.stack(shape_i(img, 0, fgraph),
-                                    shape_i(kerns, 1, fgraph),
-                                    shape2, shape3)
+        out = gpu_alloc(_zero.clone(), shape_i(img, 0, fgraph),
+                        shape_i(kerns, 1, fgraph), shape2, shape3)
         desc = GpuDnnConvDesc(border_mode='valid', subsample=(1, 1),
-                              conv_mode=conv_mode)(shape, kerns.shape)
-        return GpuDnnConvGradI()(kerns, img, desc, shape2, shape3)
+                              conv_mode=conv_mode)(out.shape, kerns.shape)
+        return GpuDnnConvGradI()(kerns, img, out, desc)
 
     # Standard case: We use GpuDnnConv with suitable padding.
     # cp_on_negative_strides will return a gpu_contiguous copy
@@ -678,7 +689,12 @@ def dnn_conv(img, kerns, border_mode='valid', subsample=(1, 1),
         # algorithm.
         if workmem is None or workmem == 'small':
             workmem = 'none'
-    return GpuDnnConv(workmem=workmem)(img, kerns, desc)
+    out_shp = GpuDnnConv.get_out_shape(img.shape, kerns.shape, border_mode,
+                                       subsample)
+    out = gpu_alloc(_zero.clone(),
+                    out_shp[0], out_shp[1],
+                    out_shp[2], out_shp[3])
+    return GpuDnnConv(workmem=workmem)(img, kerns, out, desc)
 
 
 class GpuDnnPoolDesc(GpuOp):
@@ -1455,6 +1471,27 @@ if True:
                     rval, node.outputs[0].type.broadcastable)
             return [rval]
 
+    @register_opt('cudnn')
+    @local_optimizer([GpuDnnConv], inplace=True)
+    def local_dnn_conv_inplace(node):
+        if type(node.op) != GpuDnnConv or node.op.inplace == True:
+            return
+        return [GpuDnnConv(workmem=node.op.workmem, inplace=True)(*node.inputs)]
+
+    @register_opt('cudnn')
+    @local_optimizer([GpuDnnConvGradW], inplace=True)
+    def local_dnn_convgw_inplace(node):
+        if type(node.op) != GpuDnnConvGradW or node.op.inplace == True:
+            return
+        return [GpuDnnConvGradW(inplace=True)(*node.inputs)]
+
+    @register_opt('cudnn')
+    @local_optimizer([GpuDnnConvGradI], inplace=True)
+    def local_dnn_convgi_inplace(node):
+        if type(node.op) != GpuDnnConvGradI or node.op.inplace == True:
+            return
+        return [GpuDnnConvGradI(inplace=True)(*node.inputs)]
+
     @register_opt('cudnn')
     @local_optimizer([GpuDownsampleFactorMax])
     def local_pool_dnn(node):

theano/sandbox/cuda/dnn_fwd.c

@@ -2,9 +2,8 @@
 
 int
 APPLY_SPECIFIC(conv_fwd)(CudaNdarray *input, CudaNdarray *kerns,
-                         cudnnConvolutionDescriptor_t desc,
-                         float alpha, float beta,
-                         CudaNdarray **output) {
+                         CudaNdarray *om, cudnnConvolutionDescriptor_t desc,
+                         float alpha, CudaNdarray **output) {
   cudnnStatus_t err = CUDNN_STATUS_SUCCESS;
 
   if (c_set_tensor4d(input, APPLY_SPECIFIC(input)) == -1)
@@ -12,23 +11,16 @@ APPLY_SPECIFIC(conv_fwd)(CudaNdarray *input, CudaNdarray *kerns,
   if (c_set_filter(kerns, APPLY_SPECIFIC(kerns)) == -1)
     return 1;
 
-  {
-    int out_dims[4];
-    err = cudnnGetConvolution2dForwardOutputDim(
-      desc,
-      APPLY_SPECIFIC(input),
-      APPLY_SPECIFIC(kerns),
-      &out_dims[0], &out_dims[1], &out_dims[2], &out_dims[3]);
-    if (err != CUDNN_STATUS_SUCCESS) {
-      PyErr_Format(PyExc_RuntimeError,
-		   "GpuDnnConv: error while computing the output shape: %s",
-		   cudnnGetErrorString(err));
-      return 1;
-    }
-    if (CudaNdarray_prep_output(output, 4, out_dims) != 0) {
-      return 1;
-    }
-  }
+#ifdef CONV_INPLACE
+  Py_XDECREF(*output);
+  *output = om;
+  Py_INCREF(*output);
+#else
+  if (CudaNdarray_prep_output(output, 4, CudaNdarray_HOST_DIMS(om)) != 0)
+    return 1;
+  if (CudaNdarray_CopyFromCudaNdarray(*output, om))
+    return 1;
+#endif
 
   if (c_set_tensor4d(*output, APPLY_SPECIFIC(output)) == -1)
     return 1;
@@ -55,6 +47,8 @@ APPLY_SPECIFIC(conv_fwd)(CudaNdarray *input, CudaNdarray *kerns,
     if (workspace == NULL && worksize != 0)
       return 1;
 
+    const float beta = 1;
+
     err = cudnnConvolutionForward(
       _handle,
       (void *)&alpha,

theano/sandbox/cuda/dnn_gi.c

@@ -2,9 +2,8 @@
 
 int
 APPLY_SPECIFIC(conv_gi)(CudaNdarray *kerns, CudaNdarray *output,
-			cudnnConvolutionDescriptor_t desc,
-			int h, int w, float alpha, float beta,
-			CudaNdarray **input) {
+                        CudaNdarray *im, cudnnConvolutionDescriptor_t desc,
+                        float alpha, CudaNdarray **input) {
   cudnnStatus_t err = CUDNN_STATUS_SUCCESS;
 
   if (c_set_tensor4d(output, APPLY_SPECIFIC(output)) == -1)
@@ -12,33 +11,33 @@ APPLY_SPECIFIC(conv_gi)(CudaNdarray *kerns, CudaNdarray *output,
   if (c_set_filter(kerns, APPLY_SPECIFIC(kerns)) == -1)
     return 1;
 
-  {
-    int out_dims[4];
-    out_dims[0] = CudaNdarray_HOST_DIMS(output)[0];
-    out_dims[1] = CudaNdarray_HOST_DIMS(kerns)[1];
-    out_dims[2] = h;
-    out_dims[3] = w;
-    if (CudaNdarray_prep_output(input, 4, out_dims) != 0) {
-      return 1;
-    }
-  }
+#ifdef CONV_INPLACE
+  Py_XDECREF(*input);
+  *input = im;
+  Py_INCREF(*input);
+#else
+  if (CudaNdarray_prep_output(input, 4, CudaNdarray_HOST_DIMS(im)) != 0)
+    return 1;
+  if (CudaNdarray_CopyFromCudaNdarray(*input, im))
+    return 1;
+#endif
 
   if (c_set_tensor4d(*input, APPLY_SPECIFIC(input)) == -1)
     return 1;
 
-  {
-    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));
-  }
+  const float beta = 1;
+
+  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));
   if (err != CUDNN_STATUS_SUCCESS) {
     PyErr_Format(PyExc_RuntimeError, "GpuDnnConvGradI: error doing operation: %s",
-		 cudnnGetErrorString(err));
+                 cudnnGetErrorString(err));
     return 1;
   }
   return 0;

theano/sandbox/cuda/dnn_gw.c

@@ -2,9 +2,8 @@
 
 int 
 APPLY_SPECIFIC(conv_gw)(CudaNdarray *input, CudaNdarray *output,
-			cudnnConvolutionDescriptor_t desc,
-			int h, int w, float alpha, float beta,
-			CudaNdarray **kerns) {
+                        CudaNdarray *km, cudnnConvolutionDescriptor_t desc,
+                        float alpha, CudaNdarray **kerns) {
   cudnnStatus_t err = CUDNN_STATUS_SUCCESS;
 
   if (c_set_tensor4d(input, APPLY_SPECIFIC(input)) == -1)
@@ -12,33 +11,33 @@ APPLY_SPECIFIC(conv_gw)(CudaNdarray *input, CudaNdarray *output,
   if (c_set_tensor4d(output, APPLY_SPECIFIC(output)) == -1)
     return 1;
 
-  {
-    int out_dims[4];
-    out_dims[0] = CudaNdarray_HOST_DIMS(output)[1];
-    out_dims[1] = CudaNdarray_HOST_DIMS(input)[1];
-    out_dims[2] = h;
-    out_dims[3] = w;
-    if (CudaNdarray_prep_output(kerns, 4, out_dims) != 0) {
-      return 1;
-    }
-  }
+#ifdef CONV_INPLACE
+  Py_XDECREF(*kerns);
+  *kerns = km;
+  Py_INCREF(*kerns);
+#else
+  if (CudaNdarray_prep_output(kerns, 4, CudaNdarray_HOST_DIMS(km)) != 0)
+    return 1;
+  if (CudaNdarray_CopyFromCudaNdarray(*kerns, km))
+    return 1;
+#endif
 
   if (c_set_filter(*kerns, APPLY_SPECIFIC(kerns)) == -1)
     return 1;
 
-  {
-    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));
-  }
+  const float beta = 1;
+
+  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));
   if (err != CUDNN_STATUS_SUCCESS) {
     PyErr_Format(PyExc_RuntimeError, "GpuDnnConvGradW: error doing operation: %s",
-		 cudnnGetErrorString(err));
+                 cudnnGetErrorString(err));
     return 1;
   }
   return 0;