Merge pull request #1878 from abergeron/gpuarray_doc

NEWS.txt

-.. _NEWS:
-
 =============
 Release Notes
 =============

doc/crei2013/advanced_theano.txt

@@ -15,7 +15,7 @@ Profiling
 
 Theano output:
 
-.. literalinclude:: logreg_profile.txt
+.. literalinclude:: logreg_profile.prof
 
 Compilation pipeline
 --------------------

doc/extending/inplace.txt

@@ -216,5 +216,3 @@ optimization can pre-check whether it will get rejected by using the
 which Ops can be performed inplace. You may then skip the optimization if it is
 incompatible with this check. Note however that this check does not cover all
 cases where an optimization may be rejected (it will not detect cycles).
-
-.. _optdb:

doc/extending/other_ops.txt

@@ -176,19 +176,19 @@ has more implemented distributions.
 The slowest is our wrapper on NumPy's random generator.
 
 We explain and provide advice on 3 possibles implementations of new
-distributions here::
+distributions here:
 
-1) Extend our wrapper around NumPy random functions.
+1. Extend our wrapper around NumPy random functions.
    See this `PR <https://github.com/Theano/Theano/pull/1607>`_ as an example.
 
-2) Extend MRG implementation by reusing existing Theano Op. Look into
+2. Extend MRG implementation by reusing existing Theano Op. Look into
    the ``theano/sandbox/rng_mrg.py`` file and grep for all code about
    binomial(). This distribution uses the output of the uniform
    distribution and converts it to a binomial distribution with
    existing Theano operations. The tests go in
    ``theano/sandbox/test_rng_mrg.py``
 
-3) Extend MRG implementation with a new Op that takes a uniform sample as
+3. Extend MRG implementation with a new Op that takes a uniform sample as
    input. Look in the ``theano/sandbox/{rng_mrg,multinomial}.py`` file
    and its test in ``theano/sandbox/test_multinomal.py``. This is
    recommended when current Theano ops aren't well suited to modify

doc/index.txt

@@ -13,9 +13,10 @@ arrays efficiently. Theano features:
 * **dynamic C code generation** -- Evaluate expressions faster.
 * **extensive unit-testing and self-verification** -- Detect and diagnose many types of mistake.
 
-Theano has been powering large-scale computationally intensive scientific investigations
-since 2007.  But it is also approachable enough to be used in the classroom
-(IFT6266 at the University of Montreal).
+Theano has been powering large-scale computationally intensive
+scientific investigations since 2007.  But it is also approachable
+enough to be used in the classroom (IFT6266 at the University of
+Montreal).
 
 News
 ====
@@ -59,22 +60,24 @@ directory, so that when you pull updates via Git, they will be
 automatically reflected the "installed" version. For more information about
 installation and configuration, see :ref:`installing Theano <install>`.
 
-Status
-======
+.. only:: html
 
-.. image:: https://secure.travis-ci.org/Theano/Theano.png?branch=master
-    :target: http://travis-ci.org/Theano/Theano/builds
+  Status
+  ======
 
-.. image:: https://pypip.in/v/Theano/badge.png
-    :target: https://crate.io/packages/Theano/
-    :alt: Latest PyPI version
+  .. image:: https://secure.travis-ci.org/Theano/Theano.png?branch=master
+      :target: http://travis-ci.org/Theano/Theano/builds
 
-.. image:: https://pypip.in/d/Theano/badge.png
-    :target: https://crate.io/packages/Theano/
-    :alt: Number of PyPI downloads
+  .. image:: https://pypip.in/v/Theano/badge.png
+      :target: https://crate.io/packages/Theano/
+      :alt: Latest PyPI version
 
-.. _available on PyPI: http://pypi.python.org/pypi/Theano
-.. _Related Projects: https://github.com/Theano/Theano/wiki/Related-projects
+  .. image:: https://pypip.in/d/Theano/badge.png
+      :target: https://crate.io/packages/Theano/
+      :alt: Number of PyPI downloads
+
+  .. _available on PyPI: http://pypi.python.org/pypi/Theano
+  .. _Related Projects: https://github.com/Theano/Theano/wiki/Related-projects
 
 Citing Theano
 ==============

doc/install.txt

@@ -69,12 +69,17 @@ The following libraries and software are optional:
         To be able to make picture of Theano computation graph.
 
     `NVIDIA CUDA drivers and SDK`_
-        Required for GPU code generation/execution. Only NVIDIA GPUs using
-        32-bit floating point numbers are currently supported.
+        Required for GPU code generation/execution on NVIDIA gpus
+
+    `libgpuarray`_
+        Required for GPU/CPU code generation on CUDA and OpenCL devices (see: :ref:`gpuarray`.)
+
+        :note: OpenCL support is still minimal for now.
 
 .. _LaTeX: http://www.latex-project.org/
 .. _dvipng: http://savannah.nongnu.org/projects/dvipng/
 .. _NVIDIA CUDA drivers and SDK: http://developer.nvidia.com/object/gpucomputing.html
+.. _libgpuarray: http://deeplearning.net/software/libgpuarray/installation.html
 
 Linux
 -----

doc/library/tensor/basic.txt

@@ -967,7 +967,7 @@ Reductions
      * a *list of ints* - computed along these axes
 
 .. function:: ptp(x, axis = None)
-    
+
     Range of values (maximum - minimum) along an axis.
     The name of the function comes from the acronym for peak to peak.
 
@@ -977,7 +977,7 @@ Reductions
                 flatten the array.
 
     :Returns: A new array holding the result.
- 
+
 
 Indexing
 ========
@@ -1544,8 +1544,9 @@ Linear Algebra
 
                  If an integer i, it is converted to an array containing
                  the last i dimensions of the first tensor and the first
-                 i dimensions of the second tensor (excluding the first 
-                 (batch) dimension):
+                 i dimensions of the second tensor (excluding the first
+                 (batch) dimension)::
+
                      axes = [range(a.ndim - i, b.ndim), range(1,i+1)]
 
                  If an array, its two elements must contain compatible axes
@@ -1555,17 +1556,17 @@ Linear Algebra
                  3rd axis of b must have the same shape; the same is true for
                  the 3rd axis of a and the 5th axis of b.
     :type axes: int or array-like of length 2
-    
+
     :returns: a tensor with shape equal to the concatenation of a's shape
               (less any dimensions that were summed over) and b's shape
               (less first dimension and any dimensions that were summed over).
     :rtype: tensor of tensordots
-    
-    A hybrid of batch_dot and tensordot, this function computes the 
-    tensordot product between the two tensors, by iterating over the 
+
+    A hybrid of batch_dot and tensordot, this function computes the
+    tensordot product between the two tensors, by iterating over the
     first dimension using scan to perform a sequence of tensordots.
-    
-    :note: See :func:`tensordot` and :func:`batched_dot` for 
+
+    :note: See :func:`tensordot` and :func:`batched_dot` for
         supplementary documentation.
 
 
@@ -1598,85 +1599,92 @@ Gradient / Differentiation
 
     :rtype: variable or list of variables (matching `wrt`)
     :returns: gradients of the cost with respect to each of the `wrt` terms
-    
+
 .. function:: subgraph_grad(wrt, end, start=None, cost=None, details=False)
 
-    With respect to `wrt`, computes gradients of cost and/or from existing 
-    `start` gradients, up to the `end` variables of a symbolic digraph. 
+    With respect to `wrt`, computes gradients of cost and/or from existing
+    `start` gradients, up to the `end` variables of a symbolic digraph.
     In other words, computes gradients for a subgraph of the
     symbolic theano function. Ignores all disconnected inputs.
-    
-    This can be useful when one needs to perform the gradient descent 
-    iteratively (e.g. one layer at a time in an MLP), or when a particular 
-    operation is not differentiable in theano (e.g. stochastic sampling 
-    from a multinomial). In the latter case, the gradient of the 
-    non-differentiable process could be approximated by user-defined 
-    formula, which could be calculated using the gradients of a cost 
-    with respect to samples (0s and 1s). These gradients are obtained 
-    by performing a subgraph_grad from the `cost` or previously known gradients 
-    (`start`) up to the outputs of the stochastic process (`end`). 
-    A dictionary mapping gradients obtained from the user-defined 
-    differentiation of the process, to variables, could then be fed into 
+
+    This can be useful when one needs to perform the gradient descent
+    iteratively (e.g. one layer at a time in an MLP), or when a particular
+    operation is not differentiable in theano (e.g. stochastic sampling
+    from a multinomial). In the latter case, the gradient of the
+    non-differentiable process could be approximated by user-defined
+    formula, which could be calculated using the gradients of a cost
+    with respect to samples (0s and 1s). These gradients are obtained
+    by performing a subgraph_grad from the `cost` or previously known gradients
+    (`start`) up to the outputs of the stochastic process (`end`).
+    A dictionary mapping gradients obtained from the user-defined
+    differentiation of the process, to variables, could then be fed into
     another subgraph_grad as `start` with any other `cost` (e.g. weight decay).
-    
+
     In an MLP, we could use subgraph_grad to iteratively backpropagate:
     >>> x, t = theano.tensor.fvector('x'), theano.tensor.fvector('t')
     >>> w1 = theano.shared(np.random.randn(3,4))
     >>> w2 = theano.shared(np.random.randn(4,2))
     >>> a1 = theano.tensor.tanh(theano.tensor.dot(x,w1))
     >>> a2 = theano.tensor.tanh(theano.tensor.dot(a1,w2))
-    >>> cost2 = theano.tensor.sqr(a2 - t).sum() 
+    >>> cost2 = theano.tensor.sqr(a2 - t).sum()
     >>> cost2 += theano.tensor.sqr(w2.sum())
     >>> cost1 = theano.tensor.sqr(w1.sum())
-    
+
     >>> params = [[w2],[w1]]
     >>> costs = [cost2,cost1]
     >>> grad_ends = [[a1], [x]]
-    
+
     >>> next_grad = None
     >>> param_grads = []
     >>> for i in xrange(2):
     >>>     param_grad, next_grad = theano.subgraph_grad(
-    >>>         wrt=params[i], end=grad_ends[i], 
+    >>>         wrt=params[i], end=grad_ends[i],
     >>>         start=next_grad, cost=costs[i]
     >>>     )
     >>>     next_grad = dict(zip(grad_ends[i], next_grad))
     >>>     param_grads.extend(param_grad)
-    
-    :type wrt : List of Variables.
-        Gradients are computed with respect to `wrt`.
-    
-    :type end : List of Variables.
-        Theano variables at which to end gradient descent
-        (they are considered constant in theano.grad). 
-        For convenience, the gradients with respect to these variables 
-        are also returned.
-    
-    :type start : Dictionary of Variables
-    :param start: If not None, a dictionary mapping variables to 
-            their gradients. This is useful when the gradient on some 
-            variables are known. These are used to compute the gradients
-            backwards up to the variables in `end` 
-            (they are used as known_grad in theano.grad).
-    
-    :type cost: Scalar (0-dimensional) Variable.
-    :param cost: 
-            Additional costs for which to compute the gradients.  
-            For example, these could be weight decay, an l1 constraint,
-            MSE, NLL, etc. May optionally be None if start is provided.
-            Warning : If the gradients of `cost` with respect to any 
-            of the `start` variables is already part of the `start` 
-            dictionary, then it may be counted twice with respect to `wrt` 
-            and `end`.
-    
-    :type details: bool.
-    :param details: When True, additionally returns the 
-        list of gradients from `start` and of `cost`, respectively, 
-        with respect to `wrt` (not `end`).
-    
+
+    :type wrt: list of variables
+    :param wrt:
+      Gradients are computed with respect to `wrt`.
+
+    :type end: list of variables
+    :param end:
+      Theano variables at which to end gradient descent (they are
+      considered constant in theano.grad).  For convenience, the
+      gradients with respect to these variables are also returned.
+
+    :type start: dictionary of variables
+    :param start:
+      If not None, a dictionary mapping variables to their
+      gradients. This is useful when the gradient on some variables
+      are known. These are used to compute the gradients backwards up
+      to the variables in `end` (they are used as known_grad in
+      theano.grad).
+
+    :type cost: scalar (0-dimensional) variable
+    :param cost:
+
+      Additional costs for which to compute the gradients.  For
+      example, these could be weight decay, an l1 constraint, MSE,
+      NLL, etc. May optionally be None if start is provided.
+
+      .. warning::
+
+        If the gradients of `cost` with respect to any of the `start`
+        variables is already part of the `start` dictionary, then it
+        may be counted twice with respect to `wrt` and `end`.
+
+    :type details: bool
+    :param details:
+
+      When True, additionally returns the list of gradients from
+      `start` and of `cost`, respectively, with respect to `wrt` (not
+      `end`).
+
     :rtype: Tuple of 2 or 4 Lists of Variables
-    
-    :return: Returns lists of gradients with respect to `wrt` and `end`, 
+
+    :return: Returns lists of gradients with respect to `wrt` and `end`,
             respectively.
 
 

doc/proposals/opencl.txt

-=======
-OpenCL
-=======
-
-Migrate the GPU code-generators to the PyCUDA style, and eventually to OpenCL.
-This means mainly to use a different kind of code-generation strategy.  The
-kernel itself is compiled, but the calling code remains in python or cython.  We
-would no longer generate entire C files this way, and no longer use the CLinker
-for GPU code.
-

doc/scripts/docgen.py

@@ -65,10 +65,11 @@ if __name__ == '__main__':
     options.update(dict([x, y or True] for x, y in
         getopt.getopt(sys.argv[1:],
                       'o:',
-                      ['epydoc', 'rst', 'help', 'nopdf'])[0]))
+                      ['epydoc', 'rst', 'help', 'nopdf', 'cache'])[0]))
     if options['--help']:
         print 'Usage: %s [OPTIONS]' % sys.argv[0]
         print '  -o <dir>: output the html files in the specified dir'
+        print '  --cache: use the doctree cache'
         print '  --rst: only compile the doc (requires sphinx)'
         print '  --nopdf: do not produce a PDF file from the doc, only HTML'
         print '  --epydoc: only compile the api documentation',
@@ -114,16 +115,22 @@ if __name__ == '__main__':
 
     if options['--all'] or options['--rst']:
         mkdir("doc")
-        import sphinx
         sys.path[0:0] = [os.path.join(throot, 'doc')]
-        sphinx.main(['', '-E', os.path.join(throot, 'doc'), '.'])
+        def call_sphinx(builder, workdir, extraopts=None):
+            import sphinx
+            if extraopts is None:
+                extraopts = []
+            if not options['--cache']:
+                extraopts.append('-E')
+            sphinx.main(['', '-b', builder] + extraopts +
+                        [os.path.join(throot, 'doc'), workdir])
+        call_sphinx('html', '.')
 
         if not options['--nopdf']:
             # Generate latex file in a temp directory
             import tempfile
             workdir = tempfile.mkdtemp()
-            sphinx.main(['', '-E', '-b', 'latex',
-                os.path.join(throot, 'doc'), workdir])
+            call_sphinx('latex', workdir)
             # Compile to PDF
             os.chdir(workdir)
             os.system('make')

doc/tutorial/aliasing.txt

@@ -40,11 +40,11 @@ changes to values in that pool.
 
 * The default behaviour of a function is to return user-space values for
   outputs, and to expect user-space values for inputs.
-    
+
 The distinction between Theano-managed memory and user-managed memory can be
 broken down by some Theano functions (e.g. ``shared``, ``get_value`` and the
-constructors for ``In`` and ``Out``) by using a ``borrow=True`` flag. 
-This can make those methods faster (by avoiding copy operations) at the expense 
+constructors for ``In`` and ``Out``) by using a ``borrow=True`` flag.
+This can make those methods faster (by avoiding copy operations) at the expense
 of risking subtle bugs in the overall program (by aliasing memory).
 
 The rest of this section is aimed at helping you to understand when it is safe
@@ -91,7 +91,7 @@ and may occur only temporarily even if it occurs at all.
 It is not guaranteed to occur because if Theano is using a GPU device, then the
 ``borrow`` flag has no effect. It may occur only temporarily because
 if we call a Theano function that updates the value of *s_true* the aliasing
-relationship *may* or *may not* be broken (the function is allowed to 
+relationship *may* or *may not* be broken (the function is allowed to
 update the ``shared`` variable by modifying its buffer, which will preserve
 the aliasing, or by changing which buffer the variable points to, which
 will terminate the aliasing).
@@ -113,7 +113,7 @@ Borrowing when Accessing Value of Shared Variables
 Retrieving
 ----------
 
-A ``borrow`` argument can also be used to control how a ``shared`` variable's value is 
+A ``borrow`` argument can also be used to control how a ``shared`` variable's value is
 retrieved.
 
 
@@ -138,8 +138,8 @@ The reason that ``borrow=True`` might still make a copy is that the internal
 representation of a ``shared`` variable might not be what you expect.  When you
 create a ``shared`` variable by passing a NumPy array for example, then ``get_value()``
 must return a NumPy array too.  That's how Theano can make the GPU use
-transparent.  But when you are using a GPU (or in the future perhaps a remote machine), 
-then the numpy.ndarray is not the internal representation of your data. 
+transparent.  But when you are using a GPU (or in the future perhaps a remote machine),
+then the numpy.ndarray is not the internal representation of your data.
 If you really want Theano to return its internal representation *and never copy it*
 then you should use the ``return_internal_type=True`` argument to
 ``get_value``.  It will never cast the internal object (always return in
@@ -171,8 +171,8 @@ Assigning
 ---------
 
 ``Shared`` variables also have a ``set_value`` method that can accept an optional
-``borrow=True`` argument. The semantics are similar to those of creating a new 
-``shared`` variable - ``borrow=False`` is the default and ``borrow=True`` means 
+``borrow=True`` argument. The semantics are similar to those of creating a new
+``shared`` variable - ``borrow=False`` is the default and ``borrow=True`` means
 that Theano *may* reuse the buffer you provide as the internal storage for the variable.
 
 A standard pattern for manually updating the value of a ``shared`` variable is as
@@ -216,12 +216,13 @@ be costly.  Here are a few tips to ensure fast and efficient use of GPU memory a
 (Further information on the current implementation of the GPU version of ``set_value()`` can be found
 here: :ref:`libdoc_cuda_var`)
 
+.. _borrowfunction:
 
 Borrowing when Constructing Function Objects
 ============================================
 
 A ``borrow`` argument can also be provided to the ``In`` and ``Out`` objects
-that control how ``theano.function`` handles its argument[s] and return value[s]. 
+that control how ``theano.function`` handles its argument[s] and return value[s].
 
 .. If you modify this code, also change :
 .. theano/tests/test_tutorial.py:T_aliasing.test_aliasing_3
@@ -237,7 +238,7 @@ that control how ``theano.function`` handles its argument[s] and return value[s]
 Borrowing an input means that Theano will treat the argument you provide as if
 it were part of Theano's pool of temporaries.  Consequently, your input
 may be reused as a buffer (and overwritten!) during the computation of other variables in the
-course of evaluating that function (e.g. ``f``). 
+course of evaluating that function (e.g. ``f``).
 
 
 Borrowing an output means that Theano will not insist on allocating a fresh
@@ -258,13 +259,58 @@ combination of ``return_internal_type=True`` and ``borrow=True`` arguments to
 hints that give more flexibility to the compilation and optimization of the
 graph.
 
+For GPU graphs, this borrowing can have a major speed impact.  See the following code:
+
+.. code-block:: python
+
+  from theano import function, config, shared, sandbox, tensor, Out
+  import numpy
+  import time
+
+  vlen = 10 * 30 * 768  # 10 x # cores x # threads per core
+  iters = 1000
+
+  rng = numpy.random.RandomState(22)
+  x = shared(numpy.asarray(rng.rand(vlen), config.floatX))
+  f1 = function([], sandbox.cuda.basic_ops.gpu_from_host(tensor.exp(x)))
+  f2 = function([],
+                Out(sandbox.cuda.basic_ops.gpu_from_host(tensor.exp(x)),
+                    borrow=True))
+  t0 = time.time()
+  for i in xrange(iters):
+      r = f1()
+  t1 = time.time()
+  no_borrow = t1 - t0
+  t0 = time.time()
+  for i in xrange(iters):
+      r = f2()
+  t1 = time.time()
+  print 'Looping', iters, 'times took', no_borrow, 'seconds without borrow',
+  print 'and', t1 - t0, 'seconds with borrow.'
+  if numpy.any([isinstance(x.op, tensor.Elemwise) and
+                ('Gpu' not in type(x.op).__name__)
+                for x in f1.maker.fgraph.toposort()]):
+      print 'Used the cpu'
+  else:
+      print 'Used the gpu'
+
+Which produces this output:
+
+.. code-block:: text
+
+   $ THEANO_FLAGS=device=gpu0,floatX=float32 python test1.py
+   Using gpu device 0: GeForce GTX 275
+   Looping 1000 times took 0.368273973465 seconds without borrow and 0.0240728855133 seconds with borrow.
+   Used the gpu
+
 *Take home message:*
 
-When an input *x* to a function is not needed after the function returns and you
-would like to make it available to Theano as additional workspace, then consider
-marking it with ``In(x, borrow=True)``.  It may make the function faster and
-reduce its memory requirement.
-When a return value *y* is large (in terms of memory footprint), and you only need to read from it once, right
-away when it's returned, then consider marking it with an ``Out(y,
+When an input *x* to a function is not needed after the function
+returns and you would like to make it available to Theano as
+additional workspace, then consider marking it with ``In(x,
+borrow=True)``.  It may make the function faster and reduce its memory
+requirement.  When a return value *y* is large (in terms of memory
+footprint), and you only need to read from it once, right away when
+it's returned, then consider marking it with an ``Out(y,
 borrow=True)``.
 

doc/tutorial/modes.txt

@@ -137,7 +137,7 @@ Theano defines the following modes by name:
 - ``'DebugMode``: Verify the correctness of all optimizations, and compare C and Python 
     implementations. This mode can take much longer than the other modes, but can identify
     several kinds of problems.
-- ``'ProfileMode'``(deprecated): Same optimization as FAST_RUN, but print some profiling information.
+- ``'ProfileMode'`` (deprecated): Same optimization as FAST_RUN, but print some profiling information.
 
 The default mode is typically ``FAST_RUN``, but it can be controlled via
 the configuration variable :attr:`config.mode`,

doc/tutorial/profiling.txt

@@ -72,4 +72,4 @@ to run the example:
 
 The output:
 
-.. literalinclude:: profiling_example_out.txt
+.. literalinclude:: profiling_example_out.prof

doc/tutorial/using_gpu.txt

@@ -5,17 +5,22 @@
 Using the GPU
 =============
 
-For an introductory discussion of *Graphical Processing Units* (GPU) and their use for
-intensive parallel computation purposes, see `GPGPU <http://en.wikipedia.org/wiki/GPGPU>`_.
+For an introductory discussion of *Graphical Processing Units* (GPU)
+and their use for intensive parallel computation purposes, see `GPGPU
+<http://en.wikipedia.org/wiki/GPGPU>`_.
 
-One of Theano's design goals is to specify computations at an
-abstract level, so that the internal function compiler has a lot of flexibility
-about how to carry out those computations.  One of the ways we take advantage of
-this flexibility is in carrying out calculations on an Nvidia graphics card when
-the device present in the computer is CUDA-enabled.
+One of Theano's design goals is to specify computations at an abstract
+level, so that the internal function compiler has a lot of flexibility
+about how to carry out those computations.  One of the ways we take
+advantage of this flexibility is in carrying out calculations on a
+graphics card.
 
-Setting Up CUDA
-----------------
+There are two ways currently to use a gpu, one of which only supports NVIDIA cards (:ref:`cuda`) and the other, in development, that should support any OpenCL device as well as NVIDIA cards (:ref:`gpuarray`).
+
+.. _cuda:
+
+CUDA backend
+------------
 
 If you have not done so already, you will need to install Nvidia's
 GPU-programming toolchain (CUDA) and configure Theano to use it.
@@ -23,7 +28,7 @@ We provide installation instructions for :ref:`Linux <gpu_linux>`,
 :ref:`MacOS <gpu_macos>` and :ref:`Windows <gpu_windows>`.
 
 Testing Theano with GPU
------------------------
+~~~~~~~~~~~~~~~~~~~~~~~
 
 To see if your GPU is being used, cut and paste the following program into a
 file and run it.
@@ -60,10 +65,10 @@ The program just computes the ``exp()`` of a bunch of random numbers.
 Note that we use the ``shared`` function to
 make sure that the input *x* is stored on the graphics device.
 
-.. the following figures have been measured twice on BART3 on Aug 2nd 2012 with no other job running simultaneously 
+.. the following figures have been measured twice on BART3 on Aug 2nd 2012 with no other job running simultaneously
 
 If I run this program (in check1.py) with ``device=cpu``, my computer takes a little over 3 seconds,
-whereas on the GPU it takes just over 0.64 seconds. The GPU will not always produce the exact 
+whereas on the GPU it takes just over 0.64 seconds. The GPU will not always produce the exact
 same floating-point numbers as the CPU. As a benchmark, a loop that calls ``numpy.exp(x.get_value())`` takes about 46 seconds.
 
 .. code-block:: text
@@ -87,7 +92,7 @@ Note that GPU operations in Theano require for now ``floatX`` to be *float32* (s
 
 
 Returning a Handle to Device-Allocated Data
--------------------------------------------
+~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
 
 The speedup is not greater in the preceding example because the function is
 returning its result as a NumPy ndarray which has already been copied from the
@@ -139,144 +144,61 @@ The output from this program is
       1.62323296]
     Used the gpu
 
-Here we've shaved off about 50% of the run-time by simply not copying the
-resulting array back to the host.
-The object returned by each function call is now not a NumPy array but a
-"CudaNdarray" which can be converted to a NumPy ndarray by the normal
-NumPy casting mechanism.
-
-
-Running the GPU at Full Speed
-------------------------------
-
-To really get maximum performance in this simple example, we need to use an
-:class:`out<function.Out>` instance with the flag ``borrow=True`` to tell Theano not to copy
-the output it returns to us. This is because Theano pre-allocates memory for internal use 
-(like working buffers), and by default will never return a result that is aliased to one of
-its internal buffers: instead, it will copy the buffers associated to outputs into newly 
-allocated memory at each function call. This is to ensure that subsequent function calls will
-not overwrite previously computed outputs. Although this is normally what you want, our last
-example was so simple that it had the unwanted side-effect of really slowing things down.
-
-
-.. 
-    TODO:
-    The story here about copying and working buffers is misleading and potentially not correct
-    ... why exactly does borrow=True cut 75% of the runtime ???
-
-..  TODO: Answer by Olivier D: it sounds correct to me -- memory allocations must be slow.
-
-.. If you modify this code, also change :
-.. theano/tests/test_tutorial.py:T_using_gpu.test_using_gpu_3
-.. code-block:: python
-
-    from theano import function, config, shared, sandbox, Out
-    import theano.tensor as T
-    import numpy
-    import time
-
-    vlen = 10 * 30 * 768  # 10 x # cores x # threads per core
-    iters = 1000
-
-    rng = numpy.random.RandomState(22)
-    x = shared(numpy.asarray(rng.rand(vlen), config.floatX))
-    f = function([], 
-            Out(sandbox.cuda.basic_ops.gpu_from_host(T.exp(x)),
-                borrow=True))
-    print f.maker.fgraph.toposort()
-    t0 = time.time()
-    for i in xrange(iters):
-        r = f()
-    t1 = time.time()
-    print 'Looping %d times took' % iters, t1 - t0, 'seconds'
-    print 'Result is', r
-    print 'Numpy result is', numpy.asarray(r)
-    if numpy.any([isinstance(x.op, T.Elemwise) for x in f.maker.fgraph.toposort()]):
-        print 'Used the cpu'
-    else:
-        print 'Used the gpu'
-
-Running this version of the code takes just over 0.05 seconds, that is 60x faster than
-the CPU implementation!
-
-.. code-block:: text
-
-    With *flag* ``borrow=False``:
-
-    $ THEANO_FLAGS=mode=FAST_RUN,device=gpu,floatX=float32 python using_gpu_solution_1.py
-    Using gpu device 0: GeForce GTX 580
-    [GpuElemwise{exp,no_inplace}(<CudaNdarrayType(float32, vector)>)]
-    Looping 1000 times took 0.31614613533 seconds
-    Result is <CudaNdarray object at 0x77e9270>
-    Numpy result is [ 1.23178029  1.61879349  1.52278066 ...,  2.20771813  2.29967761
-      1.62323296]
-    Used the gpu
-
-    With *flag* ``borrow=True``:
-
-    $ THEANO_FLAGS=mode=FAST_RUN,device=gpu,floatX=float32 python using_gpu_solution_1.py
-    Using gpu device 0: GeForce GTX 580
-    [GpuElemwise{exp,no_inplace}(<CudaNdarrayType(float32, vector)>)]
-    Looping 1000 times took 0.0502779483795 seconds
-    Result is <CudaNdarray object at 0x83e5cb0>
-    Numpy result is [ 1.23178029  1.61879349  1.52278066 ...,  2.20771813  2.29967761
-      1.62323296]
-    Used the gpu
-
-
-This version of the code including the flag ``borrow=True`` is slightly less safe because if we had saved
-the *r* returned from one function call, we would have to take care and remember that its value might
-be over-written by a subsequent function call.  Although ``borrow=True`` makes a dramatic difference
-in this example, be careful!  The advantage of ``borrow=True`` is much weaker in larger graphs, and 
-there is a lot of potential for making a mistake by failing to account for the resulting memory aliasing.
+Here we've shaved off about 50% of the run-time by simply not copying
+the resulting array back to the host.  The object returned by each
+function call is now not a NumPy array but a "CudaNdarray" which can
+be converted to a NumPy ndarray by the normal NumPy casting mechanism
+using something like ``numpy.asarray()``.
 
+For even more speed you can play with the ``borrow`` flag.  See
+:ref:`borrowfunction`.
 
 What Can Be Accelerated on the GPU
-----------------------------------
+~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
 
 The performance characteristics will change as we continue to optimize our
 implementations, and vary from device to device, but to give a rough idea of
 what to expect right now:
 
-* Only computations 
+* Only computations
   with *float32* data-type can be accelerated. Better support for *float64* is expected in upcoming hardware but
-  *float64* computations are still relatively slow (Jan 2010).  
+  *float64* computations are still relatively slow (Jan 2010).
 * Matrix
   multiplication, convolution, and large element-wise operations can be
   accelerated a lot (5-50x) when arguments are large enough to keep 30
-  processors busy.  
+  processors busy.
 * Indexing,
   dimension-shuffling and  constant-time reshaping will be equally fast on GPU
   as on CPU.
-* Summation 
+* Summation
   over rows/columns of tensors can be a little slower on the GPU than on the CPU.
-* Copying 
+* Copying
   of large quantities of data to and from a device is relatively slow, and
   often cancels most of the advantage of one or two accelerated functions on
   that data.  Getting GPU performance largely hinges on making data transfer to
   the device pay off.
 
 Tips for Improving Performance on GPU
--------------------------------------
+~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
 
-* Consider 
+* Consider
   adding ``floatX=float32`` to your ``.theanorc`` file if you plan to do a lot of
   GPU work.
 * Use the Theano flag ``allow_gc=False``. See :ref:`gpu_async`
-* Prefer  
+* Prefer
   constructors like ``matrix``, ``vector`` and ``scalar`` to ``dmatrix``, ``dvector`` and
   ``dscalar`` because the former will give you *float32* variables when
   ``floatX=float32``.
-* Ensure 
+* Ensure
   that your output variables have a *float32* dtype and not *float64*.  The
   more *float32* variables are in your graph, the more work the GPU can do for
   you.
-* Minimize 
+* Minimize
   tranfers to the GPU device by using ``shared`` *float32* variables to store
   frequently-accessed data (see :func:`shared()<shared.shared>`).  When using
   the GPU, *float32* tensor ``shared`` variables are stored on the GPU by default to
   eliminate transfer time for GPU ops using those variables.
-* If you aren't happy with the performance you see, try building your functions with 
+* If you aren't happy with the performance you see, try building your functions with
   ``mode='ProfileMode'``. This should print some timing information at program
   termination. Is time being used sensibly?   If an op or Apply is
   taking more time than its share, then if you know something about GPU
@@ -296,7 +218,7 @@ Tips for Improving Performance on GPU
 .. _gpu_async:
 
 GPU Async capabilities
-----------------------
+~~~~~~~~~~~~~~~~~~~~~~
 
 Ever since Theano 0.6 we started to use the asynchronous capability of
 GPUs. This allows us to be faster but with the possibility that some
@@ -314,7 +236,7 @@ as it inserts synchronization points in the graph. Set the Theano flag
 usage.
 
 Changing the Value of Shared Variables
---------------------------------------
+~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
 
 To change the value of a ``shared`` variable, e.g. to provide new data to processes,
 use ``shared_variable.set_value(new_value)``. For a lot more detail about this,
@@ -322,12 +244,12 @@ see :ref:`aliasing`.
 
 
 Exercise
-========
+++++++++
 
 Consider again the logistic regression:
 
 .. code-block:: python
-    
+
     import numpy
     import theano
     import theano.tensor as T
@@ -386,9 +308,9 @@ Consider again the logistic regression:
     print "prediction on D"
     print predict(D[0])
 
-   
 
-Modify and execute this example to run on GPU with ``floatX=float32`` and 
+
+Modify and execute this example to run on GPU with ``floatX=float32`` and
 time it using the command line ``time python file.py``. (Of course, you may use some of your answer
 to the exercise in section :ref:`Configuration Settings and Compiling Mode<using_modes>`.)
 
@@ -407,17 +329,204 @@ What can be done to further increase the speed of the GPU version? Put your idea
    * There is a limit of one GPU per process.
    * Use the Theano flag ``device=gpu`` to require use of the GPU device.
    * Use ``device=gpu{0, 1, ...}`` to specify which GPU if you have more than one.
- 
+
    * Apply the Theano flag ``floatX=float32`` (through ``theano.config.floatX``) in your code.
    * ``Cast`` inputs before storing them into a ``shared`` variable.
    * Circumvent the automatic cast of *int32* with *float32* to *float64*:
-    
+
      * Insert manual cast in your code or use *[u]int{8,16}*.
      * Insert manual cast around the mean operator (this involves division by length, which is an *int64*).
      * Notice that a new casting mechanism is being developed.
 
 :download:`Solution<using_gpu_solution_1.py>`
 
+-------------------------------------------
+
+.. _gpuarray:
+
+GpuArray Backend
+----------------
+
+If you have not done so already, you will need to install libgpuarray
+as well as at least one computing toolkit.  Instructions for doing so
+are provided at `libgpuarray <http://deeplearning.net/software/libgpuarray/installation.html>`_.
+
+While all types of devices are supported if using OpenCL, for the
+remainder of this section, whatever compute device you are using will
+be referred to as GPU.
+
+.. warning::
+
+  While it is fully our intention to support OpenCL, as of May 2014
+  this support is still in its infancy.  A lot of very useful ops
+  still do not support it because they were ported from the old
+  backend with minimal change.
+
+Testing Theano with GPU
+~~~~~~~~~~~~~~~~~~~~~~~
+
+To see if your GPU is being used, cut and paste the following program
+into a file and run it.
+
+.. code-block:: python
+
+  from theano import function, config, shared, tensor, sandbox
+  import numpy
+  import time
+
+  vlen = 10 * 30 * 768  # 10 x #cores x # threads per core
+  iters = 1000
+
+  rng = numpy.random.RandomState(22)
+  x = shared(numpy.asarray(rng.rand(vlen), config.floatX))
+  f = function([], tensor.exp(x))
+  print f.maker.fgraph.toposort()
+  t0 = time.time()
+  for i in xrange(iters):
+      r = f()
+  t1 = time.time()
+  print 'Looping %d times took' % iters, t1 - t0, 'seconds'
+  print 'Result is', r
+  if numpy.any([isinstance(x.op, tensor.Elemwise) and
+                ('Gpu' not in type(x.op).__name__)
+                for x in f.maker.fgraph.toposort()]):
+      print 'Used the cpu'
+  else:
+      print 'Used the gpu'
+
+The program just compute ``exp()`` of a bunch of random numbers.  Note
+that we use the :func:`theano.shared` function to make sure that the
+input *x* is stored on the GPU.
+
+.. code-block:: text
+
+  $ THEANO_FLAGS=device=cpu python check1.py
+  [Elemwise{exp,no_inplace}(<TensorType(float64, vector)>)]
+  Looping 1000 times took 2.6071999073 seconds
+  Result is [ 1.23178032  1.61879341  1.52278065 ...,  2.20771815  2.29967753
+    1.62323285]
+  Used the cpu
+
+  $ THEANO_FLAGS=device=cuda0 python check1.py
+  Using device cuda0: GeForce GTX 275
+  [GpuElemwise{exp,no_inplace}(<GpuArray<float64>>), HostFromGpu(gpuarray)(GpuElemwise{exp,no_inplace}.0)]
+  Looping 1000 times took 2.28562092781 seconds
+  Result is [ 1.23178032  1.61879341  1.52278065 ...,  2.20771815  2.29967753
+    1.62323285]
+  Used the gpu
+
+Returning a Handle to Device-Allocated Data
+~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
+
+By default functions that execute on the GPU still return a standard
+numpy ndarray.  A transfer operation is inserted just before the
+results are returned to ensure a consistent interface with CPU code.
+This allows changing the deivce some code runs on by only replacing
+the value of the ``device`` flag without touching the code.
+
+If you don't mind a loss of flexibility, you can ask theano to return
+the GPU object directly.  The following code is modifed to do just that.
+
+.. code-block:: python
+  :emphasize-lines: 10,17
+
+  from theano import function, config, shared, tensor, sandbox
+  import numpy
+  import time
+
+  vlen = 10 * 30 * 768  # 10 x #cores x # threads per core
+  iters = 1000
+
+  rng = numpy.random.RandomState(22)
+  x = shared(numpy.asarray(rng.rand(vlen), config.floatX))
+  f = function([], sandbox.gpuarray.basic_ops.gpu_from_host(tensor.exp(x)))
+  print f.maker.fgraph.toposort()
+  t0 = time.time()
+  for i in xrange(iters):
+      r = f()
+  t1 = time.time()
+  print 'Looping %d times took' % iters, t1 - t0, 'seconds'
+  print 'Result is', numpy.asarray(r)
+  if numpy.any([isinstance(x.op, tensor.Elemwise) and
+                ('Gpu' not in type(x.op).__name__)
+                for x in f.maker.fgraph.toposort()]):
+      print 'Used the cpu'
+  else:
+      print 'Used the gpu'
+
+Here the :func:`theano.sandbox.gpuarray.basic.gpu_from_host` call
+means "copy input to the GPU".  However during the optimization phase,
+since the result will already be on th gpu, it will be removed.  It is
+used here to tell theano that we want the result on the GPU.
+
+The output is
+
+.. code-block:: text
+
+  $ THEANO_FLAGS=device=cuda0 python check2.py
+  Using device cuda0: GeForce GTX 275
+  [GpuElemwise{exp,no_inplace}(<GpuArray<float64>>)]
+  Looping 1000 times took 0.455810785294 seconds
+  Result is [ 1.23178032  1.61879341  1.52278065 ...,  2.20771815  2.29967753
+    1.62323285]
+  Used the gpu
+
+
+While the time per call appears to be much lower than the two previous
+invocations (and should indeed be lower, since we avoid a transfer)
+the massive speedup we obtained is in part due to asynchronous nature
+of execution on GPUs, meaning that the work isn't completed yet, just
+'launched'.  We'll talk about that later.
+
+The object returned is a GpuArray from pygpu.  It mostly acts as a
+numpy ndarray with some exceptions due to its data being on the GPU.
+You can copy it to the host and convert it to a regular ndarray by
+using usual numpy casting such as ``numpy.asarray()``.
+
+For even more speed, you can play with the ``borrow`` flag.  See
+:ref:`borrowfunction`.
+
+What Can be Accelerated on the GPU
+~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
+
+The performance characteristics will of course vary from device to
+device, and also as we refine our implementation.
+
+This backend supports all regular theano data types (float32, float64,
+int, ...) however GPU support varies and some units can't deal with
+double (float64) or small (less than 32 bits like int16) data types.
+You will get an error at compile time or runtime if this is the case.
+
+Complex support is untested and most likely completely broken.
+
+In general, large operations like matrix multiplication, or
+element-wise operations with large inputs, will be significatly
+faster.
+
+
+GPU Async Capabilities
+~~~~~~~~~~~~~~~~~~~~~~
+
+By default, all operations on the GPU are run asynchronously.  This
+means that they are only scheduled to run and the function returns.
+This is made somewhat transparently by the underlying libgpuarray.
+
+A forced synchronization point is introduced when doing memory
+transfers between device and host. Another is introduced when
+releasing active memory buffers on the GPU (active buffers are buffers
+that are still in use by a kernel).
+
+It is possible to force synchronization for a particular GpuArray by
+calling its ``sync()`` method.  This is useful to get accurate timings
+when doing benchmarks.
+
+The forced synchronization points interact with the garbage collection
+of the intermediate results.  To get the fastest speed possible, you
+should disable the garbage collector by using the theano flag
+``allow_gc=False``.  Be aware that this will increase memory usage
+sometimes significantly.
+
+
 -------------------------------------------
 
 
@@ -426,7 +535,7 @@ Software for Directly Programming a GPU
 
 Leaving aside Theano which is a meta-programmer, there are:
 
-* **CUDA**: GPU programming API by NVIDIA based on extension to C (CUDA C) 
+* **CUDA**: GPU programming API by NVIDIA based on extension to C (CUDA C)
 
   * Vendor-specific
 
@@ -438,17 +547,17 @@ Leaving aside Theano which is a meta-programmer, there are:
 
   * Fewer libraries, lesser spread.
 
-* **PyCUDA**: Python bindings to CUDA driver interface allow to access Nvidia's CUDA parallel 
+* **PyCUDA**: Python bindings to CUDA driver interface allow to access Nvidia's CUDA parallel
   computation API from Python
 
   * Convenience:
 
     Makes it easy to do GPU meta-programming from within Python.
- 
+
     Abstractions to compile low-level CUDA code from Python (``pycuda.driver.SourceModule``).
 
     GPU memory buffer (``pycuda.gpuarray.GPUArray``).
-  
+
     Helpful documentation.
 
   * Completeness: Binding to all of CUDA's driver API.
@@ -465,9 +574,9 @@ Leaving aside Theano which is a meta-programmer, there are:
 
     PyCUDA knows about dependencies (e.g. it won't detach from a context before all memory
     allocated in it is also freed).
-     
 
-  (This is adapted from PyCUDA's `documentation <http://documen.tician.de/pycuda/index.html>`_ 
+
+  (This is adapted from PyCUDA's `documentation <http://documen.tician.de/pycuda/index.html>`_
   and Andreas Kloeckner's `website <http://mathema.tician.de/software/pycuda>`_ on PyCUDA.)
 
 
@@ -488,7 +597,7 @@ The following resources will assist you in this learning process:
   * `NVIDIA's slides <http://www.sdsc.edu/us/training/assets/docs/NVIDIA-02-BasicsOfCUDA.pdf>`_
 
   * `Stein's (NYU) slides <http://www.cs.nyu.edu/manycores/cuda_many_cores.pdf>`_
-  
+
 * **CUDA API and CUDA C: Advanced**
 
   * `MIT IAP2009 CUDA <https://sites.google.com/site/cudaiap2009/home>`_
@@ -509,7 +618,7 @@ The following resources will assist you in this learning process:
 
   * `Kloeckner's slides <http://www.gputechconf.com/gtcnew/on-demand-gtc.php?sessionTopic=&searchByKeyword=kloeckner&submit=&select=+&sessionEvent=2&sessionYear=2010&sessionFormat=3>`_
 
-  * `Kloeckner' website <http://mathema.tician.de/software/pycuda>`_ 
+  * `Kloeckner' website <http://mathema.tician.de/software/pycuda>`_
 
 * **PYCUDA: Advanced**
 
@@ -528,7 +637,7 @@ you feel competent enough, you may try yourself on the corresponding exercises.
   import pycuda.autoinit
   import pycuda.driver as drv
   import numpy
-  
+
   from pycuda.compiler import SourceModule
   mod = SourceModule("""
   __global__ void multiply_them(float *dest, float *a, float *b)
@@ -539,10 +648,10 @@ you feel competent enough, you may try yourself on the corresponding exercises.
   """)
 
   multiply_them = mod.get_function("multiply_them")
-  
+
   a = numpy.random.randn(400).astype(numpy.float32)
   b = numpy.random.randn(400).astype(numpy.float32)
-  
+
   dest = numpy.zeros_like(a)
   multiply_them(
           drv.Out(dest), drv.In(a), drv.In(b),
@@ -553,7 +662,7 @@ you feel competent enough, you may try yourself on the corresponding exercises.
 
 
 Exercise
-========
+~~~~~~~~
 
 Run the preceding example.
 
@@ -604,7 +713,7 @@ Modify and execute to work for a matrix of shape (20, 10).
                 pycuda_fct(inputs[0][0], z[0], numpy.intc(inputs[0][0].size),
                            block=(512,1,1), grid=grid)
             return thunk
-    
+
 
 Use this code to test it:
 
@@ -616,8 +725,7 @@ Use this code to test it:
 
 
 Exercise
-========
-
+~~~~~~~~
 
 Run the preceding example.
 

theano/compile/builders.py

@@ -24,6 +24,7 @@ class OpFromGraph(gof.Op):
         - Add support for the GPU? Probably just need an opt to remove transfer
         - Add support to pickle this Op.
         - Add support/test with random generator
+
     :note:
         - We support shared variables in the inner graph. This is automatic and
           invisible to the user. They can be as input to the node or in the

theano/compile/ops.py

@@ -490,6 +490,7 @@ def register_rebroadcast_c_code(typ, code, version=()):
                  %(oname)s for the input and output C variable names
                  respectively.  %(axis)s for the axis that we need to
                  check. This code is put in a loop for all axis
+
     :param version: A number indicating the version of the code, for cache.
     """
     Rebroadcast.c_code_and_version[typ] = (code, version)
@@ -497,14 +498,18 @@ def register_rebroadcast_c_code(typ, code, version=()):
 
 class Rebroadcast(gof.Op):
     """
-    Change the input's broadcastable fields in
-    some predetermined way.
-    e.g.: Rebroadcast((0, True), (1, False))(x)
-          would make x broadcastable in axis 0
-          and not broadcastable in axis 1
-    See also the unbroadcast, addbroadcast and patternbroadcast functions.
-
-    ..note: work inplace and work for CudaNdarrayType
+    Change the input's broadcastable fields in some predetermined way.
+
+    :code:`Rebroadcast((0, True), (1, False))(x)` would make :code:`x`
+    broadcastable in axis 0 and not broadcastable in axis 1
+
+    .. seealso::
+
+      :func:`unbroadcast <theano.tensor.unbroadcast>`
+      :func:`addbroadcast <theano.tensor.addbroadcast>`
+      :func:`patternbroadcast <theano.tensor.patternbroadcast>`
+
+    ..note: works inplace and works for CudaNdarrayType
     """
     view_map = {0: [0]}
     # Mapping from Type to C code (and version) to use.

theano/gof/utils.py

@@ -95,15 +95,15 @@ def memoize(f):
 def deprecated(filename, msg=''):
     """Decorator which will print a warning message on the first call.
 
-    Use it like this:
+    Use it like this::
 
-    @deprecated('myfile', 'do something different...')
-    def fn_name(...)
-        ...
+      @deprecated('myfile', 'do something different...')
+      def fn_name(...)
+          ...
 
-    And it will print
+    And it will print::
 
-    WARNING myfile.fn_name deprecated. do something different...
+      WARNING myfile.fn_name deprecated. do something different...
 
     """
     def _deprecated(f):

theano/gradient.py

@@ -546,59 +546,65 @@ def grad(cost, wrt, consider_constant=None,
 
 def subgraph_grad(wrt, end, start=None, cost=None, details=False):
     '''
-    With respect to `wrt`, computes gradients of cost and/or from existing 
-    `start` gradients, up to the `end` variables of a symbolic digraph. 
-    In other words, computes gradients for a subgraph of the
-    symbolic theano function. Ignores all disconnected inputs.
+    With respect to `wrt`, computes gradients of cost and/or from
+    existing `start` gradients, up to the `end` variables of a
+    symbolic digraph.  In other words, computes gradients for a
+    subgraph of the symbolic theano function. Ignores all disconnected
+    inputs.
     
-    This can be useful when one needs to perform the gradient descent 
-    iteratively (e.g. one layer at a time in an MLP), or when a particular 
-    operation is not differentiable in theano (e.g. stochastic sampling 
-    from a multinomial). In the latter case, the gradient of the 
-    non-differentiable process could be approximated by user-defined 
-    formula, which could be calculated using the gradients of a cost 
-    with respect to samples (0s and 1s). These gradients are obtained 
-    by performing a subgraph_grad from the `cost` or previously known gradients 
-    (`start`) up to the outputs of the stochastic process (`end`). 
-    A dictionary mapping gradients obtained from the user-defined 
-    differentiation of the process, to variables, could then be fed into 
-    another subgraph_grad as `start` with any other `cost` (e.g. weight decay).
+    This can be useful when one needs to perform the gradient descent
+    iteratively (e.g. one layer at a time in an MLP), or when a
+    particular operation is not differentiable in theano
+    (e.g. stochastic sampling from a multinomial). In the latter case,
+    the gradient of the non-differentiable process could be
+    approximated by user-defined formula, which could be calculated
+    using the gradients of a cost with respect to samples (0s and
+    1s). These gradients are obtained by performing a subgraph_grad
+    from the `cost` or previously known gradients (`start`) up to the
+    outputs of the stochastic process (`end`).  A dictionary mapping
+    gradients obtained from the user-defined differentiation of the
+    process, to variables, could then be fed into another
+    subgraph_grad as `start` with any other `cost` (e.g. weight
+    decay).
     
-    :type wrt : List of Variables.
-        Gradients are computed with respect to `wrt`.
+    :type wrt: list of variables
+    :param wrt:
+      Gradients are computed with respect to `wrt`.
     
-    :type end : List of Variables.
-        Theano variables at which to end gradient descent
-        (they are considered constant in theano.grad). 
-        For convenience, the gradients with respect to these variables 
-        are also returned.
+    :type end: list of variables
+    :param end:
+      Theano variables at which to end gradient descent (they are
+      considered constant in theano.grad).  For convenience, the
+      gradients with respect to these variables are also returned.
     
-    :type start : Dictionary of Variables
-    :param start: If not None, a dictionary mapping variables to 
-            their gradients. This is useful when the gradient on some 
-            variables are known. These are used to compute the gradients
-            backwards up to the variables in `end` 
-            (they are used as known_grad in theano.grad).
+    :type start: dictionary of variables
+    :param start:
+      If not None, a dictionary mapping variables to their
+      gradients. This is useful when the gradient on some variables
+      are known. These are used to compute the gradients backwards up
+      to the variables in `end` (they are used as known_grad in
+      theano.grad).
     
-    :type cost: Scalar (0-dimensional) Variable.
+    :type cost: scalar (0-dimensional) variable
     :param cost: 
-            Additional costs for which to compute the gradients.  
-            For example, these could be weight decay, an l1 constraint,
-            MSE, NLL, etc. May optionally be None if start is provided.
-            Warning : If the gradients of `cost` with respect to any 
-            of the `start` variables is already part of the `start` 
-            dictionary, then it may be counted twice with respect to `wrt` 
-            and `end`.
+      Additional costs for which to compute the gradients.  For
+      example, these could be weight decay, an l1 constraint, MSE,
+      NLL, etc. May optionally be None if start is provided.  Warning
+      : If the gradients of `cost` with respect to any of the `start`
+      variables is already part of the `start` dictionary, then it may
+      be counted twice with respect to `wrt` and `end`.
     
-    :type details: bool.
-    :param details: When True, additionally returns the 
-        list of gradients from `start` and of `cost`, respectively, 
-        with respect to `wrt` (not `end`).
+    :type details: bool
+    :param details:
+      When True, additionally returns the list of gradients from
+      `start` and of `cost`, respectively, with respect to `wrt` (not
+      `end`).
     
     :rtype: Tuple of 2 or 4 Lists of Variables
     
     :return: Returns lists of gradients with respect to `wrt` and `end`, 
             respectively.
+
     '''
     assert ((cost is not None) or (start is not None))
     assert isinstance(end, list)

theano/sandbox/rng_mrg.py

@@ -999,8 +999,9 @@ def guess_n_streams(size, warn=True):
     """
     Return a guess at a good number of streams.
 
-    :param warn: If True, warn when a guess cannot be made (in which case
-    we return 60 * 256).
+    :param warn:
+      If True, warn when a guess cannot be made (in which case we
+      return 60 * 256).
     """
     # TODO: a smart way of choosing the number of streams, see #612.
     # Note that this code was moved out of `MRG_RandomStreams` so that it can
@@ -1134,20 +1135,25 @@ class MRG_RandomStreams(object):
         ndim may be a plain integer to supplement the missing
         information.
 
-        :param low: Lower bound of the interval on which values are sampled.
-        If the ``dtype`` arg is provided, ``low`` will be cast into dtype.
-        This bound is excluded.
+        :param low:
+          Lower bound of the interval on which values are sampled.  If
+          the ``dtype`` arg is provided, ``low`` will be cast into
+          dtype.  This bound is excluded.
 
-        :param high: Higher bound of the interval on which values are sampled.
-        If the ``dtype`` arg is provided, ``high`` will be cast into dtype.
-        This bound is excluded.
+        :param high:
+          Higher bound of the interval on which values are sampled.
+          If the ``dtype`` arg is provided, ``high`` will be cast into
+          dtype.  This bound is excluded.
 
-        :param size: Can be a list of integer or Theano variable
-                (ex: the shape of other Theano Variable)
+        :param size:
+          Can be a list of integer or Theano variable (ex: the shape
+          of other Theano Variable)
+
+        :param dtype:
+          The output data type. If dtype is not specified, it will be
+          inferred from the dtype of low and high, but will be at
+          least as precise as floatX.
 
-        :param dtype: The output data type. If dtype is not specified, it will
-        be inferred from the dtype of low and high, but will be at least as
-        precise as floatX.
         """
         low = as_tensor_variable(low)
         high = as_tensor_variable(high)
@@ -1274,14 +1280,18 @@ class MRG_RandomStreams(object):
     def normal(self, size, avg=0.0, std=1.0, ndim=None,
                dtype=None, nstreams=None):
         """
-        :param size: Can be a list of integers or Theano variables (ex: the
-        shape of another Theano Variable)
+        :param size:
+          Can be a list of integers or Theano variables (ex: the shape
+          of another Theano Variable)
+
+        :param dtype:
+          The output data type. If dtype is not specified, it will be
+          inferred from the dtype of low and high, but will be at
+          least as precise as floatX.
 
-        :param dtype: The output data type. If dtype is not specified, it will
-        be inferred from the dtype of low and high, but will be at least as
-        precise as floatX.
+        :param nstreams:
+          Number of streams.
 
-        :param nstreams: Number of streams.
         """
         # We need an even number of ]0,1[ samples. Then we split them
         # in two halves. First half becomes our U1's for Box-Muller,

theano/tensor/nnet/Conv3D.py

@@ -66,8 +66,10 @@ class Conv3D(theano.Op):
         b_ = T.as_tensor_variable(b)
         d_ = T.as_tensor_variable(d)
 
-        node = theano.Apply(self, inputs=[V_, W_,b_,d_], outputs = [ T.TensorType(V_.dtype, (V_.broadcastable[0],False,False,False, W_.broadcastable[0]))() ] )
+        bcast = (V_.broadcastable[0], False, False, False, W_.broadcastable[0])
 
+        node = theano.Apply(self, inputs=[V_, W_, b_, d_],
+                            outputs=[T.TensorType(V_.dtype, bcast)()])
 
         return node
 
@@ -118,8 +120,6 @@ class Conv3D(theano.Op):
         dCdW.name = 'Conv3D_dCdW(dCdH='+dCdH_name+',V='+V_name+',W='+W_name+')'
         dCdb.name = 'Conv3D_dCdb(dCdH='+dCdH_name+',V='+V_name+',W='+W_name+',b='+b_name+')'
 
-
-
         return [ dCdV, dCdW, dCdb, dCdd ]
 
     def perform(self, node, inputs, output_storage):
@@ -149,8 +149,7 @@ class Conv3D(theano.Op):
 
         rval = (batch_size,  output_height, output_width, output_dur, output_channels )
 
-
-        return [ rval ]
+        return [rval]
 
     def c_support_code(self):
         return blas_header_text()
@@ -174,7 +173,6 @@ class Conv3D(theano.Op):
 
         H = outputs[0]
 
-
         codeSource =  """
             ///////////// < code generated by Conv3D >
 
@@ -279,7 +277,6 @@ class Conv3D(theano.Op):
             const long long outputWidth = int( (vidWidth - filterWidth) / dc )+1;
             const long long outputDur = int( (vidDur - filterDur) / dt ) +1;
 
-
             npy_intp dims[5];
             dims[0] = batchSize;
             dims[4] = outputChannels;
@@ -287,8 +284,6 @@ class Conv3D(theano.Op):
             dims[2] = outputWidth;
             dims[3] = outputDur;
 
-
-
             if(!(%(H)s) || PyArray_DIMS(%(H)s)[0]!=dims[0] ||
             PyArray_DIMS(%(H)s)[1]!=dims[1] ||
             PyArray_DIMS(%(H)s)[2]!=dims[2] ||
@@ -303,10 +298,8 @@ class Conv3D(theano.Op):
             }
 { // extra scope so fail works
 
-
             #define ELEM_AT(x, i) * ( dtype_ ## x *) ( PyArray_BYTES(x) + (i) )
 
-
             const int ws0 = PyArray_STRIDES(%(W)s)[0];
             const int ws1 = PyArray_STRIDES(%(W)s)[1];
             const int ws2 = PyArray_STRIDES(%(W)s)[2];
@@ -319,22 +312,14 @@ class Conv3D(theano.Op):
             const int bs  = PyArray_STRIDES(%(b)s)[0];
             const int hs4 = PyArray_STRIDES(%(H)s)[4];
 
-
-
             // Compute H
             //H[i,j,x,y,t] = b_j + sum_k sum_l sum_m sum_z W[j,z,k,l,m] V[i,z, dr*r+k,dc*c+l,dt*t+m]
             //TODO: add special cases
             // ex: filterDur == 1 && batchSize == 1 && dt = 1  (for SFA)
             // ex: inputChannels == 1 """
 
-
-
-
-
-
-
-
-        #if the data types are not mixed, we can insert special case optimizations based on BLAS
+        # if the data types are not mixed, we can insert special case
+        # optimizations based on BLAS
         VV, WV, bv, dv = node.inputs
         HV = node.outputs[0]
         if (theano.config.blas.ldflags and
@@ -546,7 +531,7 @@ class Conv3D(theano.Op):
 
         return strutil.render_string(codeSource,locals())
 
-global conv3D
+
 conv3D = Conv3D()
 """
 3D "convolution" of multiple filters on a minibatch