[packages/python-pyopencl] - updated to 2015.1 - added doc patch - build also python3 module

qboosh qboosh at pld-linux.org
Sun Mar 15 15:21:00 CET 2015


commit f58472434a95bd1a0323c0d4a821ec7d0c45a95c
Author: Jakub Bogusz <qboosh at pld-linux.org>
Date:   Sun Mar 15 15:22:15 2015 +0100

    - updated to 2015.1
    - added doc patch
    - build also python3 module

 python-pyopencl-doc.patch | 2523 +++++++++++++++++++++++++++++++++++++++++++++
 python-pyopencl.spec      |  189 +++-
 2 files changed, 2687 insertions(+), 25 deletions(-)
---
diff --git a/python-pyopencl.spec b/python-pyopencl.spec
index 4620d22..2cca772 100644
--- a/python-pyopencl.spec
+++ b/python-pyopencl.spec
@@ -1,26 +1,52 @@
-Summary:	Python wrapper for OpenCL
-Summary(pl.UTF-8):	Pythonowy interfejs do OpenCL
+#
+# Conditional build:
+%bcond_without	python2	# CPython 2.x module
+%bcond_without	python3	# CPython 3.x module
+%bcond_without	doc	# Sphinx documentation
+#
+%if %{without python2}
+%undefine	with_doc}
+%endif
+Summary:	Python 2 wrapper for OpenCL
+Summary(pl.UTF-8):	Interfejs Pythona 2 do OpenCL
 Name:		python-pyopencl
-Version:	2012.1
+Version:	2015.1
 Release:	1
 License:	MIT
-Group:		Development/Languages/Python
-Source0:	http://pypi.python.org/packages/source/p/pyopencl/pyopencl-%{version}.tar.gz
-# Source0-md5:	e6ae70d89d086af4636e4fbb2c806368
+Group:		Libraries/Python
+#Source0Download: https://pypi.python.org/pypi/pyopencl
+Source0:	https://pypi.python.org/packages/source/p/pyopencl/pyopencl-%{version}.tar.gz
+# Source0-md5:	c7b9dd0bb113ad852dae6fdadd417899
+Patch0:		%{name}-doc.patch
 URL:		http://mathema.tician.de/software/pyopencl
 BuildRequires:	OpenCL-devel >= 1.1
-BuildRequires:	boost-devel
 BuildRequires:	libstdc++-devel
-BuildRequires:	python-devel
+BuildRequires:	rpmbuild(macros) >= 1.616
+%if %{with python2}
+BuildRequires:	boost-python-devel
+BuildRequires:	python-appdirs >= 1.4.0
+BuildRequires:	python-devel >= 1:2.4
 BuildRequires:	python-decorator >= 3.2.0
 BuildRequires:	python-distribute
 BuildRequires:	python-numpy-devel
-BuildRequires:	python-pytools >= 2011.2
-BuildRequires:	rpmbuild(macros) >= 1.219
-BuildRequires:	sphinx-pdg
+BuildRequires:	python-pytest >= 2
+BuildRequires:	python-pytools >= 2014.2
+%{?with_doc:BuildRequires:	sphinx-pdg}
+%endif
+%if %{with python3}
+BuildRequires:	boost-python3-devel
+BuildRequires:	python3-appdirs >= 1.4.0
+BuildRequires:	python3-devel >= 1:3.2
+BuildRequires:	python3-decorator >= 3.2.0
+BuildRequires:	python3-distribute
+BuildRequires:	python3-numpy-devel
+BuildRequires:	python3-pytest >= 2
+BuildRequires:	python3-pytools >= 2014.2
+%endif
 Requires:	OpenCL >= 1.1
+Requires:	python-appdirs >= 1.4.0
 Requires:	python-decorator >= 3.2.0
-Requires:	python-pytools >= 2011.2
+Requires:	python-pytools >= 2014.2
 Requires:	python-numpy
 Suggests:	python-Mako >= 0.3.6
 BuildRoot:	%{tmpdir}/%{name}-%{version}-root-%(id -u -n)
@@ -36,47 +62,160 @@ zrównoleglonych jednostek obliczeniowych. Próbuje zaoferować
 możliwości obliczeniowe w tym samym stylu, co siostrzany projekt
 PyCUDA.
 
+%package -n python3-pyopencl
+Summary:	Python 3 wrapper for OpenCL
+Summary(pl.UTF-8):	Interfejs Pythona 3 do OpenCL
+Group:		Libraries/Python
+Requires:	OpenCL >= 1.1
+Requires:	python3-appdirs >= 1.4.0
+Requires:	python3-decorator >= 3.2.0
+Requires:	python3-pytools >= 2014.2
+Requires:	python3-numpy
+Suggests:	python3-Mako >= 0.3.6
+
+%description -n python3-pyopencl
+PyOpenCL lets you access GPUs and other massively parallel compute
+devices from Python. It tries to offer computing goodness in the
+spirit of its sister project PyCUDA.
+
+%description -n python3-pyopencl -l pl.UTF-8
+PyOpenCL pozwala na dostęp z poziomu Pythona do GPU i innych znacznie
+zrównoleglonych jednostek obliczeniowych. Próbuje zaoferować
+możliwości obliczeniowe w tym samym stylu, co siostrzany projekt
+PyCUDA.
+
+%package apidocs
+Summary:	Documentation for PyOpenCL module
+Summary(pl.UTF-8):	Dokumentacja modułu PyOpenCL
+Group:		Documentation
+
+%description apidocs
+Documentation for PyOpenCL module.
+
+%description apidocs -l pl.UTF-8
+Dokumentacja modułu PyOpenCL.
+
+%package examples
+Summary:	Examples for PyOpenCL module
+Summary(pl.UTF-8):	Przykłady do modułu PyOpenCL
+Group:		Documentation
+
+%description examples
+Examples for PyOpenCL module.
+
+%description examples -l pl.UTF-8
+Przykłady do modułu PyOpenCL.
+
 %prep
 %setup -q -n pyopencl-%{version}
+%patch0 -p1
 
 %build
-./configure.py \
-	--python-exe=%{__python} \
-	--boost-inc-dir=%{_includedir} \
-	--boost-lib-dir=%{_libdir} \
-	--cxxflags="%(echo %{rpmcxxflags} | sed -e 's/ \+/,/g')" \
-	--no-cl-enable-device-fission \
-	--no-use-shipped-boost
 # NOTES:
-# --ldflags doesn't accept commas
 # --cl-enable-gl not supported by Mesa 9.0 (missing clEnqueueReleaseGLObjects symbol)
 # device-fission requires glGetExtensionFunctionAddress (missing in Mesa 9.0)
+%define	configopts \\\
+	CXXFLAGS="%{rpmcxxflags}" \\\
+	LDFLAGS="%{rpmldflags}" \\\
+	--boost-inc-dir=%{_includedir} \\\
+	--boost-lib-dir=%{_libdir} \\\
+	--no-cl-enable-device-fission \\\
+	--no-use-shipped-boost \\\
+	%{nil}
 
-%{__python} setup.py build
+%if %{with python2}
+install -d build-2
+./configure.py \
+	%{configopts} \
+	--python-exe=%{__python} \
+	--boost-python-libname=boost_python
+
+%{__python} setup.py build --build-base build-2
 
+%if %{with doc}
 %{__make} -C doc html \
-	PYTHONPATH=$(echo $(pwd)/build/lib.*)
+	PYTHONPATH="$(echo $(pwd)/build-2/lib.*):$(pwd)"
+%endif
+%{__mv} siteconf.py siteconf-2.py
+
+%endif
+
+%if %{with python3}
+install -d build-3
+./configure.py \
+	%{configopts} \
+	--python-exe=%{__python3} \
+	--boost-python-libname=boost_python3
+
+%{__python3} setup.py build --build-base build-3
+
+%{__mv} siteconf.py siteconf-3.py
+%endif
 
 %install
 rm -rf $RPM_BUILD_ROOT
 
-%{__python} setup.py install \
+%if %{with python2}
+cp -af siteconf-2.py siteconf.py
+%{__python} setup.py \
+	build --build-base build-2 \
+	install --skip-build \
 	--optimize=2 \
 	--root=$RPM_BUILD_ROOT
 
 %py_postclean
+%endif
+
+%if %{with python3}
+cp -af siteconf-3.py siteconf.py
+%{__python3} setup.py \
+	build --build-base build-3 \
+	install --skip-build \
+	--optimize=2 \
+	--root=$RPM_BUILD_ROOT
+%endif
+
+install -d $RPM_BUILD_ROOT%{_examplesdir}/%{name}-%{version}
+cp -a examples/* $RPM_BUILD_ROOT%{_examplesdir}/%{name}-%{version}
 
 %clean
 rm -rf $RPM_BUILD_ROOT
 
+%if %{with python2}
 %files
 %defattr(644,root,root,755)
-%doc README doc/build/html/* examples
+%doc README.rst
 %dir %{py_sitedir}/pyopencl
 %attr(755,root,root) %{py_sitedir}/pyopencl/_cl.so
 %attr(755,root,root) %{py_sitedir}/pyopencl/_pvt_struct.so
 %{py_sitedir}/pyopencl/*.py[co]
 %{py_sitedir}/pyopencl/characterize
+%{py_sitedir}/pyopencl/cl
 %{py_sitedir}/pyopencl/compyte
 %{py_sitedir}/pyopencl-%{version}-py*.egg-info
-%{_includedir}/pyopencl
+%endif
+
+%if %{with python3}
+%files -n python3-pyopencl
+%defattr(644,root,root,755)
+%doc README.rst
+%dir %{py3_sitedir}/pyopencl
+%attr(755,root,root) %{py3_sitedir}/pyopencl/_cl.cpython-*.so
+%attr(755,root,root) %{py3_sitedir}/pyopencl/_pvt_struct.cpython-*.so
+%{py3_sitedir}/pyopencl/*.py
+%{py3_sitedir}/pyopencl/__pycache__
+%{py3_sitedir}/pyopencl/characterize
+%{py3_sitedir}/pyopencl/cl
+%{py3_sitedir}/pyopencl/compyte
+%{py3_sitedir}/pyopencl-%{version}-py*.egg-info
+%endif
+
+%if %{with doc}
+%files apidocs
+%defattr(644,root,root,755)
+%doc doc/_build/html/*
+%endif
+
+%files examples
+%defattr(644,root,root,755)
+%{_examplesdir}/%{name}-%{version}
diff --git a/python-pyopencl-doc.patch b/python-pyopencl-doc.patch
new file mode 100644
index 0000000..12ba778
--- /dev/null
+++ b/python-pyopencl-doc.patch
@@ -0,0 +1,2523 @@
+diff -Nur pyopencl-2015.1/doc/algorithm.rst pyopencl/doc/algorithm.rst
+--- pyopencl-2015.1/doc/algorithm.rst	1970-01-01 01:00:00.000000000 +0100
++++ pyopencl/doc/algorithm.rst	2015-03-05 20:44:55.746519776 +0100
+@@ -0,0 +1,274 @@
++Parallel Algorithms
++===================
++
++.. include:: subst.rst
++
++Element-wise expression evalution ("map")
++-----------------------------------------
++
++.. module:: pyopencl.elementwise
++
++Evaluating involved expressions on :class:`pyopencl.array.Array` instances by
++using overloaded operators can be somewhat inefficient, because a new temporary
++is created for each intermediate result. The functionality in the module
++:mod:`pyopencl.elementwise` contains tools to help generate kernels that
++evaluate multi-stage expressions on one or several operands in a single pass.
++
++.. autoclass:: ElementwiseKernel(context, arguments, operation, name="kernel", preamble="", options=[])
++
++    .. method:: __call__(*args, wait_for=None)
++
++        Invoke the generated scalar kernel. The arguments may either be scalars or
++        :class:`GPUArray` instances.
++
++        |std-enqueue-blurb|
++
++Here's a usage example::
++
++.. literalinclude:: ../examples/demo_elementwise.py
++
++(You can find this example as
++:download:`examples/demo_elementwise.py <../examples/demo_elementwise.py>`
++in the PyOpenCL distribution.)
++
++.. _custom-reductions:
++
++Sums and counts ("reduce")
++--------------------------
++
++.. module:: pyopencl.reduction
++
++.. class:: ReductionKernel(ctx, dtype_out, neutral, reduce_expr, map_expr=None, arguments=None, name="reduce_kernel", options=[], preamble="")
++
++    Generate a kernel that takes a number of scalar or vector *arguments*
++    (at least one vector argument), performs the *map_expr* on each entry of
++    the vector argument and then the *reduce_expr* on the outcome of that.
++    *neutral* serves as an initial value. *preamble* offers the possibility
++    to add preprocessor directives and other code (such as helper functions)
++    to be added before the actual reduction kernel code.
++
++    Vectors in *map_expr* should be indexed by the variable *i*. *reduce_expr*
++    uses the formal values "a" and "b" to indicate two operands of a binary
++    reduction operation. If you do not specify a *map_expr*, ``in[i]`` is
++    automatically assumed and treated as the only one input argument.
++
++    *dtype_out* specifies the :class:`numpy.dtype` in which the reduction is
++    performed and in which the result is returned. *neutral* is specified as
++    float or integer formatted as string. *reduce_expr* and *map_expr* are
++    specified as string formatted operations and *arguments* is specified as a
++    string formatted as a C argument list. *name* specifies the name as which
++    the kernel is compiled. *options* are passed unmodified to
++    :meth:`pyopencl.Program.build`. *preamble* specifies a string of code that
++    is inserted before the actual kernels.
++
++    .. method:: __call__(*args, queue=None, wait_for=None, return_event=False, out=None)
++
++        |explain-waitfor|
++
++        With *out* the resulting single-entry :class:`pyopencl.array.Array` can
++        be specified. Because offsets are supported one can store results
++        anywhere (e.g. ``out=a[3]``).
++
++        :return: the resulting scalar as a single-entry :class:`pyopencl.array.Array`
++            if *return_event* is *False*, otherwise a tuple ``(scalar_array, event)``.
++
++        .. note::
++
++            The returned :class:`pyopencl.Event` corresponds only to part of the
++            execution of the reduction. It is not suitable for profiling.
++
++    .. versionadded:: 2011.1
++
++    .. versionchanged:: 2014.2
++
++        Added *out* parameter.
++
++Here's a usage example::
++
++    a = pyopencl.array.arange(queue, 400, dtype=numpy.float32)
++    b = pyopencl.array.arange(queue, 400, dtype=numpy.float32)
++
++    krnl = ReductionKernel(ctx, numpy.float32, neutral="0",
++            reduce_expr="a+b", map_expr="x[i]*y[i]",
++            arguments="__global float *x, __global float *y")
++
++    my_dot_prod = krnl(a, b).get()
++
++.. _custom-scan:
++
++Prefix Sums ("scan")
++--------------------
++
++.. module:: pyopencl.scan
++
++.. |scan_extra_args| replace:: a list of tuples *(name, value)* specifying
++    extra arguments to pass to the scan procedure. For version 2013.1,
++    *value* must be a of a :mod:`numpy` sized scalar type. As of version 2013.2,
++    *value* may also be a :class:`pyopencl.array.Array`.
++.. |preamble| replace:: A snippet of C that is inserted into the compiled kernel
++    before the actual kernel function. May be used for, e.g. type definitions
++    or include statements.
++
++A prefix sum is a running sum of an array, as provided by
++e.g. :mod:`numpy.cumsum`::
++
++    >>> import numpy as np
++    >>> a = [1,1,1,1,1,2,2,2,2,2]
++    >>> np.cumsum(a)
++    array([ 1,  2,  3,  4,  5,  7,  9, 11, 13, 15])
++
++This is a very simple example of what a scan can do. It turns out that scans
++are significantly more versatile. They are a basic building block of many
++non-trivial parallel algorithms. Many of the operations enabled by scans seem
++difficult to parallelize because of loop-carried dependencies.
++
++.. seealso::
++
++    `Prefix sums and their applications <http://citeseerx.ist.psu.edu/viewdoc/summary?doi=10.1.1.128.6230>`_, by Guy Blelloch.
++        This article gives an overview of some surprising applications of scans.
++
++    :ref:`predefined-scans`
++        These operations built into PyOpenCL are realized using :class:`GenericScanKernel`.
++
++Usage Example
++^^^^^^^^^^^^^
++
++This example illustrates the implementation of a simplified version of
++:func:`pyopencl.algorithm.copy_if`,
++which copies integers from an array into the (variable-size) output if they are
++greater than 300::
++
++    knl = GenericScanKernel(
++            ctx, np.int32,
++            arguments="__global int *ary, __global int *out",
++            input_expr="(ary[i] > 300) ? 1 : 0",
++            scan_expr="a+b", neutral="0",
++            output_statement="""
++                if (prev_item != item) out[item-1] = ary[i];
++                """)
++
++    out = a.copy()
++    knl(a, out)
++
++    a_host = a.get()
++    out_host = a_host[a_host > 300]
++
++    assert (out_host == out.get()[:len(out_host)]).all()
++
++The value being scanned over is a number of flags indicating whether each array
++element is greater than 300. These flags are computed by *input_expr*. The
++prefix sum over this array gives a running count of array items greater than
++300. The *output_statement* the compares `prev_item` (the previous item's scan
++result, i.e. index) to `item` (the current item's scan result, i.e.
++index). If they differ, i.e. if the predicate was satisfied at this
++position, then the item is stored in the output at the computed index.
++
++This example does not make use of the following advanced features also available
++in PyOpenCL:
++
++* Segmented scans
++
++* Access to the previous item in *input_expr* (e.g. for comparisons)
++  See the `implementation <https://github.com/pyopencl/pyopencl/blob/master/pyopencl/scan.py#L1353>`_ of :func:`unique` for an example.
++
++Making Custom Scan Kernels
++^^^^^^^^^^^^^^^^^^^^^^^^^^
++
++.. versionadded:: 2013.1
++
++.. autoclass:: GenericScanKernel
++
++    .. method:: __call__(*args, allocator=None, queue=None, size=None, wait_for=None)
++
++        *queue* and *allocator* default to the ones provided on the first
++        :class:`pyopencl.array.Array` in *args*. *size* may specify the
++        length of the scan to be carried out. If not given, this length
++        is inferred from the first array argument passed.
++
++        |std-enqueue-blurb|
++
++        .. note::
++
++            The returned :class:`pyopencl.Event` corresponds only to part of the
++            execution of the scan. It is not suitable for profiling.
++
++Debugging aids
++~~~~~~~~~~~~~~
++
++.. class:: GenericDebugScanKernel
++
++    Performs the same function and has the same interface as
++    :class:`GenericScanKernel`, but uses a dead-simple, sequential scan.  Works
++    best on CPU platforms, and helps isolate bugs in scans by removing the
++    potential for issues originating in parallel execution.
++
++.. _predefined-scans:
++
++Simple / Legacy Interface
++^^^^^^^^^^^^^^^^^^^^^^^^^
++
++.. class:: ExclusiveScanKernel(ctx, dtype, scan_expr, neutral, name_prefix="scan", options=[], preamble="", devices=None)
++
++    Generates a kernel that can compute a `prefix sum <https://secure.wikimedia.org/wikipedia/en/wiki/Prefix_sum>`_
++    using any associative operation given as *scan_expr*.
++    *scan_expr* uses the formal values "a" and "b" to indicate two operands of
++    an associative binary operation. *neutral* is the neutral element
++    of *scan_expr*, obeying *scan_expr(a, neutral) == a*.
++
++    *dtype* specifies the type of the arrays being operated on.
++    *name_prefix* is used for kernel names to ensure recognizability
++    in profiles and logs. *options* is a list of compiler options to use
++    when building. *preamble* specifies a string of code that is
++    inserted before the actual kernels. *devices* may be used to restrict
++    the set of devices on which the kernel is meant to run. (defaults
++    to all devices in the context *ctx*.
++
++    .. method:: __call__(self, input_ary, output_ary=None, allocator=None, queue=None)
++
++.. class:: InclusiveScanKernel(dtype, scan_expr, neutral=None, name_prefix="scan", options=[], preamble="", devices=None)
++
++    Works like :class:`ExclusiveScanKernel`.
++
++    .. versionchanged:: 2013.1
++        *neutral* is now always required.
++
++For the array `[1,2,3]`, inclusive scan results in `[1,3,6]`, and exclusive
++scan results in `[0,1,3]`.
++
++Here's a usage example::
++
++    knl = InclusiveScanKernel(context, np.int32, "a+b")
++
++    n = 2**20-2**18+5
++    host_data = np.random.randint(0, 10, n).astype(np.int32)
++    dev_data = cl_array.to_device(queue, host_data)
++
++    knl(dev_data)
++    assert (dev_data.get() == np.cumsum(host_data, axis=0)).all()
++
++Predicated copies ("partition", "unique", ...)
++----------------------------------------------
++
++.. module:: pyopencl.algorithm
++
++.. autofunction:: copy_if
++
++.. autofunction:: remove_if
++
++.. autofunction:: partition
++
++.. autofunction:: unique
++
++Sorting (radix sort)
++--------------------
++
++.. autoclass:: RadixSort
++
++    .. automethod:: __call__
++
++Building many variable-size lists
++---------------------------------
++
++.. autoclass:: ListOfListsBuilder
++
++    .. automethod:: __call__
+diff -Nur pyopencl-2015.1/doc/array.rst pyopencl/doc/array.rst
+--- pyopencl-2015.1/doc/array.rst	1970-01-01 01:00:00.000000000 +0100
++++ pyopencl/doc/array.rst	2015-03-05 20:44:55.746519776 +0100
+@@ -0,0 +1,293 @@
++Multi-dimensional arrays
++========================
++
++.. module:: pyopencl.array
++
++The functionality in this module provides something of a work-alike for
++:mod:`numpy` arrays, but with all operations executed on the CL compute device.
++
++Data Types
++----------
++
++PyOpenCL provides some amount of integration between the :mod:`numpy`
++type system, as represented by :class:`numpy.dtype`, and the types
++available in OpenCL. All the simple scalar types map straightforwardly
++to their CL counterparts.
++
++.. _vector-types:
++
++Vector Types
++^^^^^^^^^^^^
++
++.. class :: vec
++
++    All of OpenCL's supported vector types, such as `float3` and `long4` are
++    available as :mod:`numpy` data types within this class. These
++    :class:`numpy.dtype` instances have field names of `x`, `y`, `z`, and `w`
++    just like their OpenCL counterparts. They will work both for parameter passing
++    to kernels as well as for passing data back and forth between kernels and
++    Python code. For each type, a `make_type` function is also provided (e.g.
++    `make_float3(x,y,z)`).
++
++    If you want to construct a pre-initialized vector type you have three new
++    functions to choose from:
++
++    * `zeros_type()`
++    * `ones_type()`
++    * `filled_type(fill_value)`
++
++    .. versionadded:: 2014.1
++
++    .. versionchanged:: 2014.1
++        The `make_type` functions have a default value (0) for each component.
++        Relying on the default values has been deprecated. Either specify all
++        components or use one of th new flavors mentioned above for constructing
++        a vector.
++
++Custom data types
++^^^^^^^^^^^^^^^^^
++
++If you would like to use your own (struct/union/whatever) data types in array
++operations where you supply operation source code, define those types in the
++*preamble* passed to :class:`pyopencl.elementwise.ElementwiseKernel`,
++:class:`pyopencl.reduction.ReductionKernel` (or similar), and let PyOpenCL know
++about them using this function:
++
++.. currentmodule:: pyopencl.tools
++
++.. autofunction:: get_or_register_dtype
++
++.. exception:: TypeNameNotKnown
++
++    .. versionadded:: 2013.1
++
++.. function:: register_dtype(dtype, name)
++
++    .. versionchanged:: 2013.1
++        This function has been deprecated. It is recommended that you develop
++        against the new interface, :func:`get_or_register_dtype`.
++
++.. function:: dtype_to_ctype(dtype)
++
++    Returns a C name registered for *dtype*.
++
++    .. versionadded: 2013.1
++
++This function helps with producing C/OpenCL declarations for structured
++:class:`numpy.dtype` instances:
++
++.. autofunction:: match_dtype_to_c_struct
++
++A more complete example of how to use custom structured types can be
++found in :file:`examples/demo-struct-reduce.py` in the PyOpenCL
++distribution.
++
++.. currentmodule:: pyopencl.array
++
++Complex Numbers
++^^^^^^^^^^^^^^^
++
++PyOpenCL's :class:`Array` type supports complex numbers out of the box, by
++simply using the corresponding :mod:`numpy` types.
++
++If you would like to use this support in your own kernels, here's how to
++proceed: Since OpenCL 1.2 (and earlier) do not specify native complex number
++support, PyOpenCL works around that deficiency. By saying::
++
++    #include <pyopencl-complex.h>
++
++in your kernel, you get complex types `cfloat_t` and `cdouble_t`, along with
++functions defined on them such as `cfloat_mul(a, b)` or `cdouble_log(z)`.
++Elementwise kernels automatically include the header if your kernel has
++complex input or output.
++See the `source file
++<https://github.com/pyopencl/pyopencl/blob/master/src/cl/pyopencl-complex.h>`_
++for a precise list of what's available.
++
++If you need double precision support, please::
++
++    #define PYOPENCL_DEFINE_CDOUBLE
++
++before including the header, as DP support apparently cannot be reliably
++autodetected.
++
++Under the hood, the complex types are simply `float2` and `double2`.
++
++.. warning::
++    Note that, at the OpenCL source code level, addition (real + complex) and
++    multiplication (complex*complex) are defined for e.g. `float2`, but yield
++    wrong results, so that you need to use the corresponding functions.
++    (The :mod:`Array` type implements complex arithmetic as you remember it,
++    without any idiotic quirks like this.)
++
++.. versionadded:: 2012.1
++
++The :class:`Array` Class
++------------------------
++
++.. autoclass:: Array
++
++.. autoexception:: ArrayHasOffsetError
++
++Constructing :class:`Array` Instances
++^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
++
++.. autofunction:: to_device
++.. function:: empty(queue, shape, dtype, order="C", allocator=None, data=None)
++
++    A synonym for the :class:`Array` constructor.
++
++.. autofunction:: zeros
++.. autofunction:: empty_like
++.. autofunction:: zeros_like
++.. autofunction:: arange
++.. autofunction:: take
++.. autofunction:: concatenate
++
++Conditionals
++^^^^^^^^^^^^
++
++.. autofunction:: if_positive
++.. autofunction:: maximum
++.. autofunction:: minimum
++
++.. _reductions:
++
++Reductions
++^^^^^^^^^^
++
++.. autofunction:: sum
++.. autofunction:: dot
++.. autofunction:: vdot
++.. autofunction:: subset_dot
++.. autofunction:: max
++.. autofunction:: min
++.. autofunction:: subset_max
++.. autofunction:: subset_min
++
++See also :ref:`custom-reductions`.
++
++Elementwise Functions on :class:`Arrray` Instances
++--------------------------------------------------
++
++.. module:: pyopencl.clmath
++
++The :mod:`pyopencl.clmath` module contains exposes array versions of the C
++functions available in the OpenCL standard. (See table 6.8 in the spec.)
++
++.. function:: acos(array, queue=None)
++.. function:: acosh(array, queue=None)
++.. function:: acospi(array, queue=None)
++
++.. function:: asin(array, queue=None)
++.. function:: asinh(array, queue=None)
++.. function:: asinpi(array, queue=None)
++
++.. function:: atan(array, queue=None)
++.. autofunction:: atan2
++.. function:: atanh(array, queue=None)
++.. function:: atanpi(array, queue=None)
++.. autofunction:: atan2pi
++
++.. function:: cbrt(array, queue=None)
++.. function:: ceil(array, queue=None)
++.. TODO: copysign
++
++.. function:: cos(array, queue=None)
++.. function:: cosh(array, queue=None)
++.. function:: cospi(array, queue=None)
++
++.. function:: erfc(array, queue=None)
++.. function:: erf(array, queue=None)
++.. function:: exp(array, queue=None)
++.. function:: exp2(array, queue=None)
++.. function:: exp10(array, queue=None)
++.. function:: expm1(array, queue=None)
++
++.. function:: fabs(array, queue=None)
++.. TODO: fdim
++.. function:: floor(array, queue=None)
++.. TODO: fma
++.. TODO: fmax
++.. TODO: fmin
++
++.. function:: fmod(arg, mod, queue=None)
++
++    Return the floating point remainder of the division `arg/mod`,
++    for each element in `arg` and `mod`.
++
++.. TODO: fract
++
++
++.. function:: frexp(arg, queue=None)
++
++    Return a tuple `(significands, exponents)` such that
++    `arg == significand * 2**exponent`.
++
++.. TODO: hypot
++
++.. function:: ilogb(array, queue=None)
++.. function:: ldexp(significand, exponent, queue=None)
++
++    Return a new array of floating point values composed from the
++    entries of `significand` and `exponent`, paired together as
++    `result = significand * 2**exponent`.
++
++
++.. function:: lgamma(array, queue=None)
++.. TODO: lgamma_r
++
++.. function:: log(array, queue=None)
++.. function:: log2(array, queue=None)
++.. function:: log10(array, queue=None)
++.. function:: log1p(array, queue=None)
++.. function:: logb(array, queue=None)
++
++.. TODO: mad
++.. TODO: maxmag
++.. TODO: minmag
++
++
++.. function:: modf(arg, queue=None)
++
++    Return a tuple `(fracpart, intpart)` of arrays containing the
++    integer and fractional parts of `arg`.
++
++.. function:: nan(array, queue=None)
++
++.. TODO: nextafter
++.. TODO: remainder
++.. TODO: remquo
++
++.. function:: rint(array, queue=None)
++.. TODO: rootn
++.. function:: round(array, queue=None)
++
++.. function:: sin(array, queue=None)
++.. TODO: sincos
++.. function:: sinh(array, queue=None)
++.. function:: sinpi(array, queue=None)
++
++.. function:: sqrt(array, queue=None)
++
++.. function:: tan(array, queue=None)
++.. function:: tanh(array, queue=None)
++.. function:: tanpi(array, queue=None)
++.. function:: tgamma(array, queue=None)
++.. function:: trunc(array, queue=None)
++
++Generating Arrays of Random Numbers
++-----------------------------------
++
++.. automodule:: pyopencl.clrandom
++
++    .. autoclass:: RanluxGenerator
++
++        .. automethod:: fill_uniform
++        .. automethod:: uniform
++        .. automethod:: fill_normal
++        .. automethod:: normal
++        .. automethod:: synchronize
++
++    .. autofunction:: rand
++    .. autofunction:: fill_rand
+diff -Nur pyopencl-2015.1/doc/.gitignore pyopencl/doc/.gitignore
+--- pyopencl-2015.1/doc/.gitignore	1970-01-01 01:00:00.000000000 +0100
++++ pyopencl/doc/.gitignore	2015-03-05 20:44:55.746519776 +0100
+@@ -0,0 +1 @@
++constants.inc
+diff -Nur pyopencl-2015.1/doc/howto.rst pyopencl/doc/howto.rst
+--- pyopencl-2015.1/doc/howto.rst	1970-01-01 01:00:00.000000000 +0100
++++ pyopencl/doc/howto.rst	2015-03-05 20:44:55.746519776 +0100
+@@ -0,0 +1,103 @@
++How-tos
++=======
++
++How to use struct types with PyOpenCL
++-------------------------------------
++
++We import and initialize PyOpenCL as usual:
++
++.. doctest::
++    :options: +ELLIPSIS
++
++    >>> import numpy as np
++    >>> import pyopencl as cl
++    >>> import pyopencl.tools
++    >>> import pyopencl.array
++
++    >>> ctx = cl.create_some_context(interactive=False)
++    >>> queue = cl.CommandQueue(ctx)
++
++Then, suppose we would like to declare a struct consisting of an integer and a
++floating point number. We first create a :class:`numpy.dtype` along these
++lines:
++
++.. doctest::
++
++    >>> my_struct = np.dtype([("field1", np.int32), ("field2", np.float32)])
++    >>> print my_struct
++    [('field1', '<i4'), ('field2', '<f4')]
++
++.. note::
++
++    Not all :mod:`numpy` dtypes are supported yet. For example strings (and
++    generally things that have a shape of their own) are not supported.
++
++Since OpenCL C may have a different opinion for :mod:`numpy` on how the struct
++should be laid out, for example because of `alignment
++<https://en.wikipedia.org/wiki/Data_structure_alignment>`_. So as a first step, we
++match our dtype against CL's version:
++
++.. doctest::
++
++    >>> my_struct, my_struct_c_decl = cl.tools.match_dtype_to_c_struct(
++    ...    ctx.devices[0], "my_struct", my_struct)
++    >>> print my_struct_c_decl
++    typedef struct {
++      int field1;
++      float field2;
++    } my_struct;
++    <BLANKLINE>
++    <BLANKLINE>
++
++We then tell PyOpenCL about our new type.
++
++.. doctest::
++
++    >>> my_struct = cl.tools.get_or_register_dtype("my_struct", my_struct)
++
++Next, we can create some data of that type on the host and transfer it to
++the device:
++
++.. doctest::
++
++    >>> ary_host = np.empty(20, my_struct)
++    >>> ary_host["field1"].fill(217)
++    >>> ary_host["field2"].fill(1000)
++    >>> ary_host[13]["field2"] = 12
++    >>> print ary_host
++    [(217, 1000.0) (217, 1000.0) (217, 1000.0) (217, 1000.0) (217, 1000.0)
++     (217, 1000.0) (217, 1000.0) (217, 1000.0) (217, 1000.0) (217, 1000.0)
++     (217, 1000.0) (217, 1000.0) (217, 1000.0) (217, 12.0) (217, 1000.0)
++     (217, 1000.0) (217, 1000.0) (217, 1000.0) (217, 1000.0) (217, 1000.0)]
++
++    >>> ary = cl.array.to_device(queue, ary_host)
++
++We can then operate on the array with our own kernels:
++
++.. doctest::
++
++    >>> prg = cl.Program(ctx, my_struct_c_decl + """
++    ...     __kernel void set_to_1(__global my_struct *a)
++    ...     {
++    ...         a[get_global_id(0)].field1 = 1;
++    ...     }
++    ...     """).build()
++
++    >>> evt = prg.set_to_1(queue, ary.shape, None, ary.data)
++    >>> print ary
++    [(1, 1000.0) (1, 1000.0) (1, 1000.0) (1, 1000.0) (1, 1000.0) (1, 1000.0)
++     (1, 1000.0) (1, 1000.0) (1, 1000.0) (1, 1000.0) (1, 1000.0) (1, 1000.0)
++     (1, 1000.0) (1, 12.0) (1, 1000.0) (1, 1000.0) (1, 1000.0) (1, 1000.0)
++     (1, 1000.0) (1, 1000.0)]
++
++as well as with PyOpenCL's built-in operations:
++
++    >>> from pyopencl.elementwise import ElementwiseKernel
++    >>> elwise = ElementwiseKernel(ctx, "my_struct *a", "a[i].field1 = 2;",
++    ...    preamble=my_struct_c_decl)
++    >>> evt = elwise(ary)
++    >>> print ary
++    [(2, 1000.0) (2, 1000.0) (2, 1000.0) (2, 1000.0) (2, 1000.0) (2, 1000.0)
++     (2, 1000.0) (2, 1000.0) (2, 1000.0) (2, 1000.0) (2, 1000.0) (2, 1000.0)
++     (2, 1000.0) (2, 12.0) (2, 1000.0) (2, 1000.0) (2, 1000.0) (2, 1000.0)
++     (2, 1000.0) (2, 1000.0)]
+diff -Nur pyopencl-2015.1/doc/index.rst pyopencl/doc/index.rst
+--- pyopencl-2015.1/doc/index.rst	1970-01-01 01:00:00.000000000 +0100
++++ pyopencl/doc/index.rst	2015-03-05 20:44:55.746519776 +0100
+@@ -0,0 +1,102 @@
++Welcome to PyOpenCL's documentation!
++====================================
++
++PyOpenCL gives you easy, Pythonic access to the `OpenCL
++<http://www.khronos.org/opencl/>`_ parallel computation API.
++What makes PyOpenCL special?
++
++* Object cleanup tied to lifetime of objects. This idiom,
++  often called
++  `RAII <http://en.wikipedia.org/wiki/Resource_Acquisition_Is_Initialization>`_
++  in C++, makes it much easier to write correct, leak- and
++  crash-free code.
++
++* Completeness. PyOpenCL puts the full power of OpenCL's API at your
++  disposal, if you wish. Every obscure `get_info()` query and
++  all CL calls are accessible.
++
++* Automatic Error Checking. All errors are automatically translated
++  into Python exceptions.
++
++* Speed. PyOpenCL's base layer is written in C++, so all the niceties above
++  are virtually free.
++
++* Helpful Documentation. You're looking at it. ;)
++
++* Liberal license. PyOpenCL is open-source under the
++  :ref:`MIT license <license>`
++  and free for commercial, academic, and private use.
++
++Here's an example, to give you an impression:
++
++.. literalinclude:: ../examples/demo.py
++
++(You can find this example as
++:download:`examples/demo.py <../examples/demo.py>` in the PyOpenCL
++source distribution.)
++
++Tutorials
++=========
++
++* Gaston Hillar's `two-part article series
++  <http://www.drdobbs.com/open-source/easy-opencl-with-python/240162614>`_
++  in Dr. Dobb's Journal provides a friendly introduction to PyOpenCL.
++* `Simon McIntosh-Smith <http://www.cs.bris.ac.uk/~simonm/>`_
++  and `Tom Deakin <http://www.tomdeakin.com/>`_'s course
++  `Hands-on OpenCL <http://handsonopencl.github.io/>`_ contains
++  both `lecture slides <https://github.com/HandsOnOpenCL/Lecture-Slides/releases>`_
++  and `excercises (with solutions) <https://github.com/HandsOnOpenCL/Exercises-Solutions>`_
++  (The course covers PyOpenCL as well as OpenCL's C and C++ APIs.)
++* PyOpenCL course at `PASI <http://bu.edu/pasi>`_: Parts
++  `1 <https://www.youtube.com/watch?v=X9mflbX1NL8>`_
++  `2 <https://www.youtube.com/watch?v=MqvfCE_bKOg>`_
++  `3 <https://www.youtube.com/watch?v=TAvKmV7CuUw>`_
++  `4 <https://www.youtube.com/watch?v=SsuJ0LvZW1Q>`_
++  (YouTube, 2011)
++* PyOpenCL course at `DTU GPULab <http://gpulab.imm.dtu.dk/>`_ and
++  `Simula <http://simula.no/>`_ (2011):
++  `Lecture 1 <http://tiker.net/pub/simula-pyopencl-lec1.pdf>`_
++  `Lecture 2 <http://tiker.net/pub/simula-pyopencl-lec2.pdf>`_
++  `Problem set 1 <http://tiker.net/pub/simula-pyopencl-probset1.pdf>`_
++  `Problem set 2 <http://tiker.net/pub/simula-pyopencl-probset2.pdf>`_
++* Ian Johnson's `PyOpenCL tutorial <http://enja.org/2011/02/22/adventures-in-pyopencl-part-1-getting-started-with-python/>`_.
++
++Software that works with or enhances PyOpenCL
++=============================================
++
++* Bogdan Opanchuk's `reikna <http://pypi.python.org/pypi/reikna>`_ offers a
++  variety of GPU-based algorithms (FFT, random number generation, matrix
++  multiplication) designed to work with :class:`pyopencl.array.Array` objects.
++
++* Gregor Thalhammer's `gpyfft <https://github.com/geggo/gpyfft>`_ provides a
++  Python wrapper for the OpenCL FFT library clFFT from AMD.
++
++If you know of a piece of software you feel that should be on this list, please
++let me know, or, even better, send a patch!
++
++Contents
++========
++
++.. toctree::
++    :maxdepth: 2
++
++    runtime
++    array
++    algorithm
++    howto
++    tools
++    misc
++
++Note that this guide does not explain OpenCL programming and technology. Please
++refer to the official `Khronos OpenCL documentation <http://khronos.org/opencl>`_
++for that.
++
++PyOpenCL also has its own `web site <http://mathema.tician.de/software/pyopencl>`_,
++where you can find updates, new versions, documentation, and support.
++
++Indices and tables
++==================
++
++* :ref:`genindex`
++* :ref:`modindex`
++* :ref:`search`
+diff -Nur pyopencl-2015.1/doc/misc.rst pyopencl/doc/misc.rst
+--- pyopencl-2015.1/doc/misc.rst	1970-01-01 01:00:00.000000000 +0100
++++ pyopencl/doc/misc.rst	2015-03-05 20:44:55.749853109 +0100
+@@ -0,0 +1,510 @@
++Installation
++============
++
++Installation information is maintained collaboratively on the
++`PyOpenCL Wiki <http://wiki.tiker.net/PyOpenCL/Installation>`_.
++
++Acknowledgments
++===============
++
++Too many to list. Please see the
++`commit log <https://github.com/pyopencl/pyopencl/commits/master>`_
++for detailed acknowledgments.
++
++Tips
++====
++
++Syntax highlighting
++-------------------
++
++You can obtain Vim syntax highlighting for OpenCL C inlined in Python by
++checking `this file
++<https://github.com/pyopencl/pyopencl/blob/master/contrib/pyopencl.vim>`_.
++
++Note that the triple-quoted strings containing the source must start with
++`"""//CL// ..."""`.
++
++IPython integration
++-------------------
++
++PyOpenCL comes with IPython integration, which lets you seamlessly integrate
++PyOpenCL kernels into your IPython notebooks. Simply load the PyOpenCL 
++IPython extension using::
++
++    %load_ext pyopencl.ipython_ext
++
++and then use the ``%%cl_kernel`` 'cell-magic' command. See `this notebook
++<http://nbviewer.ipython.org/urls/raw.githubusercontent.com/pyopencl/pyopencl/master/examples/ipython-demo.ipynb>`_
++(which ships with PyOpenCL) for a demonstration.
++
++You can pass build options to be used for building the program executable by using the ``-o`` flag on the first line of the cell (next to the ``%%cl_kernel`` directive). For example: `%%cl_kernel -o "-cl-fast-relaxed-math"``.
++
++There are also line magics: ``cl_load_edit_kernel`` which will load a file into the next cell (adding ``cl_kernel`` to the first line) and ``cl_kernel_from_file`` which will compile kernels from a file (as if you copy-and-pasted the contents of the file to a cell with ``cl_kernel``). Boths of these magics take options ``-f`` to specify the file and optionally ``-o`` for build options.
++
++.. versionadded:: 2014.1
++
++Guidelines
++==========
++
++.. _api-compatibility:
++
++API Stability
++-------------
++
++I consider PyOpenCL's API "stable".  That doesn't mean it can't
++change. But if it does, your code will generally continue to run. It
++may however start spewing warnings about things you need to change to
++stay compatible with future versions.
++
++Deprecation warnings will be around for a whole year, as identified by the
++first number in the release name.  (the "2014" in "2014.1") I.e. a function
++that was deprecated in 2014.n will generally be removed in 2015.n (or perhaps
++later). Further, the stability promise applies for any code that's part of a
++released version. It doesn't apply to undocumented bits of the API, and it
++doesn't apply to unreleased code downloaded from git.
++
++.. _versus-c:
++
++Relation with OpenCL's C Bindings
++---------------------------------
++
++We've tried to follow these guidelines when binding the OpenCL's
++C interface to Python:
++
++* Remove the `cl_`, `CL_` and `cl` prefix from data types, macros and
++  function names.
++* Follow :pep:`8`, i.e.
++
++  * Make function names lowercase.
++  * If a data type or function name is composed of more than one word,
++    separate the words with a single underscore.
++
++* `get_info` functions become attributes.
++* Object creation is done by constructors, to the extent possible.
++  (i.e. minimize use of "factory functions")
++
++* If an operation involves two or more "complex" objects (like e.g. a
++  kernel enqueue involves a kernel and a queue), refuse the temptation
++  to guess which one should get a method for the operation.
++  Instead, simply leave that command to be a function.
++
++.. _interoperability:
++
++Interoperability with other OpenCL software
++-------------------------------------------
++
++Just about every object in :mod:`pyopencl` supports the following
++interface (here shown as an example for :class:`pyopencl.MemoryObject`,
++from which :class:`pyopencl.Buffer` and :class:`pyopencl.Image` inherit):
++
++* :meth:`pyopencl.MemoryObject.from_int_ptr`
++* :attr:`pyopencl.MemoryObject.int_ptr`
++
++This allows retrieving the C-level pointer to an OpenCL object as a Python
++integer, which may then be passed to other C libraries whose interfaces expose
++OpenCL objects. It also allows turning C-level OpenCL objects obtained from
++other software to be turned into the corresponding :mod:`pyopencl` objects.
++
++.. versionadded:: 2013.2
++
++User-visible Changes
++====================
++
++Version 2014.1
++--------------
++.. note::
++
++    This version is currently under development. You can get snapshots from
++    PyOpenCL's `git repository <https://github.com/pyopencl/pyopencl>`_
++
++* :ref:`ipython-integration`
++* Bug fixes
++
++Version 2013.2
++--------------
++
++* Add :meth:`pyopencl.array.Array.map_to_host`.
++* Support *strides* on :func:`pyopencl.enqueue_map_buffer` and
++  :func:`pyopencl.enqueue_map_image`.
++* :class:`pyopencl.ImageFormat` was made comparable and hashable.
++* :mod:`pyopencl.reduction` supports slicing (contributed by Alex Nitz)
++* Added :ref:`interoperability`
++* Bug fixes
++
++Version 2013.1
++--------------
++
++* Vastly improved :ref:`custom-scan`.
++* Add :func:`pyopencl.tools.match_dtype_to_c_struct`,
++  for better integration of the CL and :mod:`numpy` type systems.
++* More/improved Bessel functions.
++  See `the source <https://github.com/pyopencl/pyopencl/tree/master/src/cl>`_.
++* Add :envvar:`PYOPENCL_NO_CACHE` environment variable to aid debugging.
++  (e.g. with AMD's CPU implementation, see
++  `their programming guide <http://developer.amd.com/sdks/AMDAPPSDK/assets/AMD_Accelerated_Parallel_Processing_OpenCL_Programming_Guide.pdf>`_)
++* Deprecated :func:`pyopencl.tools.register_dtype` in favor of
++  :func:`pyopencl.tools.get_or_register_dtype`.
++* Clean up the :class:`pyopencl.array.Array` constructor interface.
++* Deprecate :class:`pyopencl.array.DefaultAllocator`.
++* Deprecate :class:`pyopencl.tools.CLAllocator`.
++* Introduce :class:`pyopencl.tools.DeferredAllocator`, :class:`pyopencl.tools.ImmediateAllocator`.
++* Allow arrays whose beginning does not coincide with the beginning of their
++  :attr:`pyopencl.array.Array.data` :class:`pyopencl.Buffer`.
++  See :attr:`pyopencl.array.Array.base_data` and :attr:`pyopencl.array.Array.offset`.
++  Note that not all functions in PyOpenCL support such arrays just yet. These
++  will fail with :exc:`pyopencl.array.ArrayHasOffsetError`.
++* Add :meth:`pyopencl.array.Array.__getitem__` and :meth:`pyopencl.array.Array.__setitem__`,
++  supporting generic slicing.
++
++  It is *possible* to create non-contiguous arrays using this functionality.
++  Most operations (elementwise etc.) will not work on such arrays.
++
++  Note also that some operations (specifically, reductions and scans) on sliced
++  arrays that start past the beginning of the original array will fail for now.
++  This will be fixed in a future release.
++
++* :class:`pyopencl.CommandQueue` may be used as a context manager (in a ``with`` statement)
++* Add :func:`pyopencl.clmath.atan2`, :func:`pyopencl.clmath.atan2pi`.
++* Add :func:`pyopencl.array.concatenate`.
++* Add :meth:`pyopencl.Kernel.capture_call`.
++
++.. note::
++
++    The addition of :meth:`pyopencl.array.Array.__getitem__` has an unintended
++    consequence due to `numpy bug 3375
++    <https://github.com/numpy/numpy/issues/3375>`_.  For instance, this
++    expression::
++
++        numpy.float32(5) * some_pyopencl_array
++
++    may take a very long time to execute. This is because :mod:`numpy` first
++    builds an object array of (compute-device) scalars (!) before it decides that
++    that's probably not such a bright idea and finally calls
++    :meth:`pyopencl.array.Array.__rmul__`.
++
++    Note that only left arithmetic operations of :class:`pyopencl.array.Array`
++    by :mod:`numpy` scalars are affected. Python's number types (:class:`float` etc.)
++    are unaffected, as are right multiplications.
++
++    If a program that used to run fast suddenly runs extremely slowly, it is
++    likely that this bug is to blame.
++
++    Here's what you can do:
++
++    * Use Python scalars instead of :mod:`numpy` scalars.
++    * Switch to right multiplications if possible.
++    * Use a patched :mod:`numpy`. See the bug report linked above for a pull
++      request with a fix.
++    * Switch to a fixed version of :mod:`numpy` when available.
++
++Version 2012.1
++--------------
++
++* Support for complex numbers.
++* Support for Bessel functions. (experimental)
++* Numerous fixes.
++
++Version 2011.2
++--------------
++
++* Add :func:`pyopencl.enqueue_migrate_mem_object`.
++* Add :func:`pyopencl.image_from_array`.
++* IMPORTANT BUGFIX: Kernel caching was broken for all the 2011.1.x releases, with
++  severe consequences on the execution time of :class:`pyopencl.array.Array`
++  operations.
++  Henrik Andresen at a `PyOpenCL workshop at DTU <http://gpulab.imm.dtu.dk/courses.html>`_
++  first noticed the strange timings.
++* All comparable PyOpenCL objects are now also hashable.
++* Add :func:`pyopencl.tools.context_dependent_memoize` to the documented
++  functionality.
++* Base :mod:`pyopencl.clrandom` on `RANLUXCL <https://bitbucket.org/ivarun/ranluxcl>`_,
++  add functionality.
++* Add :class:`pyopencl.NannyEvent` objects.
++* Add :mod:`pyopencl.characterize`.
++* Ensure compatibility with OS X Lion.
++* Add :func:`pyopencl.tools.register_dtype` to enable scan/reduction on struct types.
++* :func:`pyopencl.enqueue_migrate_mem_object` was renamed
++  :func:`pyopencl.enqueue_migrate_mem_object_ext`.
++  :func:`pyopencl.enqueue_migrate_mem_object` now refers to the OpenCL 1.2 function
++  of this name, if available.
++* :func:`pyopencl.create_sub_devices` was renamed
++  :func:`pyopencl.create_sub_devices_ext`.
++  :func:`pyopencl.create_sub_devices` now refers to the OpenCL 1.2 function
++  of this name, if available.
++* Alpha support for OpenCL 1.2.
++
++Version 2011.1.2
++----------------
++
++* More bug fixes.
++
++Version 2011.1.1
++----------------
++
++* Fixes for Python 3 compatibility. (with work by Christoph Gohlke)
++
++Version 2011.1
++--------------
++
++* All *is_blocking* parameters now default to *True* to avoid
++  crashy-by-default behavior. (suggested by Jan Meinke)
++  In particular, this change affects
++  :func:`pyopencl.enqueue_read_buffer`,
++  :func:`pyopencl.enqueue_write_buffer`,
++  :func:`pyopencl.enqueue_read_buffer_rect`,
++  :func:`pyopencl.enqueue_write_buffer_rect`,
++  :func:`pyopencl.enqueue_read_image`,
++  :func:`pyopencl.enqueue_write_image`,
++  :func:`pyopencl.enqueue_map_buffer`,
++  :func:`pyopencl.enqueue_map_image`.
++* Add :mod:`pyopencl.reduction`.
++* Add :ref:`reductions`.
++* Add :mod:`pyopencl.scan`.
++* Add :meth:`pyopencl.MemoryObject.get_host_array`.
++* Deprecate context arguments of
++  :func:`pyopencl.array.to_device`,
++  :func:`pyopencl.array.zeros`,
++  :func:`pyopencl.array.arange`.
++* Make construction of :class:`pyopencl.array.Array` more flexible (*cqa* argument.)
++* Add :ref:`memory-pools`.
++* Add vector types, see :class:`pyopencl.array.vec`.
++* Add :attr:`pyopencl.array.Array.strides`, :attr:`pyopencl.array.Array.flags`.
++  Allow the creation of arrys in C and Fortran order.
++* Add :func:`pyopencl.enqueue_copy`. Deprecate all other transfer functions.
++* Add support for numerous extensions, among them device fission.
++* Add a compiler cache.
++* Add the 'g_times_l' keyword arg to kernel execution.
++
++Version 0.92
++------------
++
++* Add support for OpenCL 1.1.
++* Add support for the
++  `cl_khr_gl_sharing <ghttp://www.khronos.org/registry/cl/extensions/khr/cl_khr_gl_sharing.txt>`_
++  extension, leading to working GL interoperability.
++* Add :meth:`pyopencl.Kernel.set_args`.
++* The call signature of :meth:`pyopencl.Kernel.__call__` changed to
++  emphasize the importance of *local_size*.
++* Add :meth:`pyopencl.Kernel.set_scalar_arg_dtypes`.
++* Add support for the
++  `cl_nv_device_attribute_query <http://www.khronos.org/registry/cl/extensions/khr/cl_nv_device_attribute_query.txt>`_
++  extension.
++* Add :meth:`pyopencl.array.Array` and related functionality.
++* Make build not depend on Boost C++.
++
++Version 0.91.5
++--------------
++
++* Add :attr:`pyopencl.ImageFormat.channel_count`,
++  :attr:`pyopencl.ImageFormat.dtype_size`,
++  :attr:`pyopencl.ImageFormat.itemsize`.
++* Add missing :func:`pyopencl.enqueue_copy_buffer`.
++* Add :func:`pyopencl.create_some_context`.
++* Add :func:`pyopencl.enqueue_barrier`, which was previously missing.
++
++Version 0.91.4
++--------------
++
++A bugfix release. No user-visible changes.
++
++Version 0.91.3
++--------------
++
++* All parameters named *host_buffer* were renamed *hostbuf* for consistency
++  with the :class:`pyopencl.Buffer` constructor introduced in 0.91.
++  Compatibility code is in place.
++* The :class:`pyopencl.Image` constructor does not need a *shape* parameter if the
++  given *hostbuf* has *hostbuf.shape*.
++* The :class:`pyopencl.Context` constructor can now be called without parameters.
++
++Version 0.91.2
++--------------
++
++* :meth:`pyopencl.Program.build` now captures build logs and adds them
++  to the exception text.
++* Deprecate :func:`pyopencl.create_context_from_type` in favor of second
++  form of :class:`pyopencl.Context` constructor
++* Introduce :class:`pyopencl.LocalMemory`.
++* Document kernel invocation and :meth:`pyopencl.Kernel.set_arg`.
++
++Version 0.91.1
++--------------
++
++* Fixed a number of bugs, notably involving :class:`pyopencl.Sampler`.
++* :class:`pyopencl.Device`, :class:`pyopencl.Platform`,
++  :class:`pyopencl.Context` now have nicer string representations.
++* Add :attr:`Image.shape`. (suggested by David Garcia)
++
++Version 0.91
++------------
++
++* Add :ref:`gl-interop`.
++* Add a test suite.
++* Fix numerous `get_info` bugs. (reports by David Garcia and the test suite)
++* Add :meth:`pyopencl.ImageFormat.__repr__`.
++* Add :meth:`pyopencl.addressing_mode.to_string` and colleagues.
++* The `pitch` arguments to
++  :func:`pyopencl.create_image_2d`,
++  :func:`pyopencl.create_image_3d`,
++  :func:`pyopencl.enqueue_read_image`, and
++  :func:`pyopencl.enqueue_write_image`
++  are now defaulted to zero. The argument order of `enqueue_{read,write}_image`
++  has changed for this reason.
++* Deprecate
++  :func:`pyopencl.create_image_2d`,
++  :func:`pyopencl.create_image_3d`
++  in favor of the :class:`pyopencl.Image` constructor.
++* Deprecate
++  :func:`pyopencl.create_program_with_source`,
++  :func:`pyopencl.create_program_with_binary`
++  in favor of the :class:`pyopencl.Program` constructor.
++* Deprecate
++  :func:`pyopencl.create_buffer`,
++  :func:`pyopencl.create_host_buffer`
++  in favor of the :class:`pyopencl.Buffer` constructor.
++* :meth:`pyopencl.MemoryObject.get_image_info` now actually exists.
++* Add :attr:`pyopencl.MemoryObject.image.info`.
++* Fix API tracing.
++* Add constructor arguments to :class:`pyopencl.ImageFormat`.  (suggested by David Garcia)
++
++Version 0.90.4
++--------------
++
++* Add build fixes for Windows and OS X.
++
++Version 0.90.3
++--------------
++
++* Fix a GNU-ism in the C++ code of the wrapper.
++
++Version 0.90.2
++--------------
++
++* Fix :meth:`pyopencl.Platform.get_info`.
++* Fix passing properties to :class:`pyopencl.CommandQueue`.
++  Also fix related documentation.
++
++Version 0.90.1
++--------------
++
++* Fix building on the Mac.
++
++Version 0.90
++------------
++
++* Initial release.
++
++.. _license:
++
++License
++=======
++
++PyOpenCL is licensed to you under the MIT/X Consortium license:
++
++Copyright (c) 2009-13 Andreas Klöckner and Contributors.
++
++Permission is hereby granted, free of charge, to any person
++obtaining a copy of this software and associated documentation
++files (the "Software"), to deal in the Software without
++restriction, including without limitation the rights to use,
++copy, modify, merge, publish, distribute, sublicense, and/or sell
++copies of the Software, and to permit persons to whom the
++Software is furnished to do so, subject to the following
++conditions:
++
++The above copyright notice and this permission notice shall be
++included in all copies or substantial portions of the Software.
++
++THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
++EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES
++OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
++NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT
++HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY,
++WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
++FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
++OTHER DEALINGS IN THE SOFTWARE.
++
++PyOpenCL includes derivatives of parts of the `Thrust
++<https://code.google.com/p/thrust/>`_ computing package (in particular the scan
++implementation). These parts are licensed as follows:
++
++    Copyright 2008-2011 NVIDIA Corporation
++
++    Licensed under the Apache License, Version 2.0 (the "License");
++    you may not use this file except in compliance with the License.
++    You may obtain a copy of the License at
++
++        <http://www.apache.org/licenses/LICENSE-2.0>
++
++    Unless required by applicable law or agreed to in writing, software
++    distributed under the License is distributed on an "AS IS" BASIS,
++    WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
++    See the License for the specific language governing permissions and
++    limitations under the License.
++
++.. note::
++
++    If you use Apache-licensed parts, be aware that these may be incompatible
++    with software licensed exclusively under GPL2.  (Most software is licensed
++    as GPL2 or later, in which case this is not an issue.)
++
++PyOpenCL includes the RANLUXCL random number generator:
++
++    Copyright (c) 2011 Ivar Ursin Nikolaisen
++
++    Permission is hereby granted, free of charge, to any person obtaining a copy of this
++    software and associated documentation files (the "Software"), to deal in the Software
++    without restriction, including without limitation the rights to use, copy, modify,
++    merge, publish, distribute, sublicense, and/or sell copies of the Software, and to
++    permit persons to whom the Software is furnished to do so, subject to the following
++    conditions:
++
++    The above copyright notice and this permission notice shall be included in all copies
++    or substantial portions of the Software.
++
++    THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IMPLIED,
++    INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS FOR A
++    PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT
++    HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF
++    CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE
++    OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
++
++Frequently Asked Questions
++==========================
++
++The FAQ is maintained collaboratively on the
++`Wiki FAQ page <http://wiki.tiker.net/PyOpenCL/FrequentlyAskedQuestions>`_.
++
++Citing PyOpenCL
++===============
++
++We are not asking you to gratuitously cite PyOpenCL in work that is otherwise
++unrelated to software. That said, if you do discuss some of the development
++aspects of your code and would like to highlight a few of the ideas behind
++PyOpenCL, feel free to cite `this article
++<http://dx.doi.org/10.1016/j.parco.2011.09.001>`_:
++
++    Andreas Klöckner, Nicolas Pinto, Yunsup Lee, Bryan Catanzaro, Paul Ivanov,
++    Ahmed Fasih, PyCUDA and PyOpenCL: A scripting-based approach to GPU
++    run-time code generation, Parallel Computing, Volume 38, Issue 3, March
++    2012, Pages 157-174.
++
++Here's a Bibtex entry for your convenience::
++
++    @article{kloeckner_pycuda_2012,
++       author = {{Kl{\"o}ckner}, Andreas
++            and {Pinto}, Nicolas
++            and {Lee}, Yunsup
++            and {Catanzaro}, B.
++            and {Ivanov}, Paul
++            and {Fasih}, Ahmed },
++       title = "{PyCUDA and PyOpenCL: A Scripting-Based Approach to GPU Run-Time Code Generation}",
++       journal = "Parallel Computing",
++       volume = "38",
++       number = "3",
++       pages = "157--174",
++       year = "2012",
++       issn = "0167-8191",
++       doi = "10.1016/j.parco.2011.09.001",
++    }
++
+diff -Nur pyopencl-2015.1/doc/runtime.rst pyopencl/doc/runtime.rst
+--- pyopencl-2015.1/doc/runtime.rst	1970-01-01 01:00:00.000000000 +0100
++++ pyopencl/doc/runtime.rst	2015-03-05 20:44:55.749853109 +0100
+@@ -0,0 +1,975 @@
++.. _reference-doc:
++
++.. include:: subst.rst
++
++OpenCL Runtime
++==============
++
++Version Queries
++---------------
++
++.. module:: pyopencl
++.. moduleauthor:: Andreas Kloeckner <inform at tiker.net>
++
++.. data:: VERSION
++
++    Gives the numeric version of PyOpenCL as a variable-length tuple
++    of integers. Enables easy version checks such as
++    *VERSION >= (0, 93)*.
++
++.. data:: VERSION_STATUS
++
++    A text string such as `"rc4"` or `"beta"` qualifying the status
++    of the release.
++
++.. data:: VERSION_TEXT
++
++    The full release name (such as `"0.93rc4"`) in string form.
++
++.. function:: get_cl_header_version()
++
++    Return a variable-length tuple of integers representing the
++    version of the OpenCL header against which PyOpenCL was
++    compiled.
++
++    .. versionadded:: 0.92
++
++.. _errors:
++
++Error Reporting
++---------------
++
++.. class:: Error
++
++    Base class for all PyOpenCL exceptions.
++
++.. class:: MemoryError
++
++.. class:: LogicError
++
++.. class:: RuntimeError
++
++Constants
++---------
++
++.. include:: constants.inc
++
++Platforms, Devices and Contexts
++-------------------------------
++
++.. function:: get_platforms()
++
++    Return a list of :class:`Platform` instances.
++
++.. class:: Platform
++
++    .. attribute:: info
++
++        Lower case versions of the :class:`platform_info` constants
++        may be used as attributes on instances of this class
++        to directly query info attributes.
++
++    .. method:: get_info(param)
++
++        See :class:`platform_info` for values of *param*.
++
++    .. method:: get_devices(device_type=device_type.ALL)
++
++        Return a list of devices matching *device_type*.
++        See :class:`device_type` for values of *device_type*.
++
++        .. versionchanged:: 2013.2
++
++            This used to raise an exception if no matching
++            devices were found. Now, it will simply return
++            an empty list.
++
++    .. automethod:: from_int_ptr
++    .. autoattribute:: int_ptr
++
++    |comparable|
++
++.. class:: Device
++
++    .. attribute:: info
++
++        Lower case versions of the :class:`device_info` constants
++        may be used as attributes on instances of this class
++        to directly query info attributes.
++
++    .. method:: get_info(param)
++
++        See :class:`device_info` for values of *param*.
++
++    .. automethod:: from_int_ptr
++    .. autoattribute:: int_ptr
++
++    Two instances of this class may be compared using *=="* and *"!="*.
++
++.. class:: Context(devices=None, properties=None, dev_type=None)
++
++    Create a new context. *properties* is a list of key-value
++    tuples, where each key must be one of :class:`context_properties`.
++    At most one of *devices* and *dev_type* may be not `None`, where
++    *devices* is a list of :class:`Device` instances, and
++    *dev_type* is one of the :class:`device_type` constants.
++    If neither is specified, a context with a *dev_type* of
++    :attr:`device_type.DEFAULT` is created.
++
++    .. note::
++
++        Calling the constructor with no arguments will fail for recent
++        CL drivers that support the OpenCL ICD. If you want similar,
++        just-give-me-a-context-already behavior, we recommend
++        :func:`create_some_context`. See, e.g. this
++        `explanation by AMD <http://developer.amd.com/support/KnowledgeBase/Lists/KnowledgeBase/DispForm.aspx?ID=71>`_.
++
++    .. note::
++
++        For
++        :attr:`context_properties.CL_GL_CONTEXT_KHR`,
++        :attr:`context_properties.CL_EGL_DISPLAY_KHR`,
++        :attr:`context_properties.CL_GLX_DISPLAY_KHR`,
++        :attr:`context_properties.CL_WGL_HDC_KHR`, and
++        :attr:`context_properties.CL_CGL_SHAREGROUP_KHR`
++        :attr:`context_properties.CL_CGL_SHAREGROUP_APPLE`
++        the value in the key-value pair is a PyOpenGL context or display
++        instance.
++
++    .. versionchanged:: 0.91.2
++        Constructor arguments *dev_type* added.
++
++    .. attribute:: info
++
++        Lower case versions of the :class:`context_info` constants
++        may be used as attributes on instances of this class
++        to directly query info attributes.
++
++    .. method:: get_info(param)
++
++        See :class:`context_info` for values of *param*.
++
++    .. method:: create_sub_devices(properties)
++
++        *properties* is an array of one (or more) of the forms::
++
++            [ dpp.EQUALLY, 8]
++            [ dpp.BY_COUNTS, 5, 7, 9, dpp.PARTITION_BY_COUNTS_LIST_END]
++            [ dpp.BY_NAMES, 5, 7, 9, dpp.PARTITION_BY_NAMES_LIST_END]
++            [ dpp.BY_AFFINITY_DOMAIN, dad.L1_CACHE]
++
++        where `dpp` represents :class:`device_partition_property`
++        and `dad` represent :class:`device_affinity_domain`.
++
++        `PROPERTIES_LIST_END_EXT` is added automatically.
++
++        Only available with CL 1.2.
++
++        .. versionadded:: 2011.2
++
++    .. method:: create_sub_devices_ext(properties)
++
++        *properties* is an array of one (or more) of the forms::
++
++            [ dppe.EQUALLY, 8]
++            [ dppe.BY_COUNTS, 5, 7, 9, dppe.PARTITION_BY_COUNTS_LIST_END]
++            [ dppe.BY_NAMES, 5, 7, 9, dppe.PARTITION_BY_NAMES_LIST_END]
++            [ dppe.BY_AFFINITY_DOMAIN, ad.L1_CACHE]
++
++        where `dppe` represents :class:`device_partition_property_ext`
++        and `ad` represent :class:`affinity_domain_ext`.
++
++        `PROPERTIES_LIST_END_EXT` is added automatically.
++
++        Only available with the `cl_ext_device_fission` extension.
++
++        .. versionadded:: 2011.1
++
++    .. automethod:: from_int_ptr
++    .. autoattribute:: int_ptr
++
++    |comparable|
++
++.. function:: create_some_context(interactive=True)
++
++    Create a :class:`Context` 'somehow'.
++
++    If multiple choices for platform and/or device exist, *interactive*
++    is True, and *sys.stdin.isatty()* is also True,
++    then the user is queried about which device should be chosen.
++    Otherwise, a device is chosen in an implementation-defined manner.
++
++Command Queues and Events
++-------------------------
++
++.. class:: CommandQueue(context, device=None, properties=None)
++
++    Create a new command queue. *properties* is a bit field
++    consisting of :class:`command_queue_properties` values.
++
++    if *device* is None, one of the devices in *context* is chosen
++    in an implementation-defined manner.
++
++    A :class:`CommandQueue` may be used as a context manager, like this::
++
++        with cl.CommandQueue(self.cl_context) as queue:
++            enqueue_stuff(queue, ...)
++
++    :meth:`finish` is automatically called at the end of the context.
++
++    .. versionadded:: 2013.1
++
++        Context manager capability.
++
++    .. attribute:: info
++
++        Lower case versions of the :class:`command_queue_info` constants
++        may be used as attributes on instances of this class
++        to directly query info attributes.
++
++    .. method:: get_info(param)
++
++        See :class:`command_queue_info` for values of *param*.
++
++    .. method:: set_property(prop, enable)
++
++        See :class:`command_queue_properties` for possible values of *prop*.
++        *enable* is a :class:`bool`.
++
++        Unavailable in OpenCL 1.1 and newer.
++
++    .. method:: flush()
++    .. method:: finish()
++
++    .. automethod:: from_int_ptr
++    .. autoattribute:: int_ptr
++
++    |comparable|
++
++.. class:: Event
++
++    .. attribute:: info
++
++        Lower case versions of the :class:`event_info` constants
++        may be used as attributes on instances of this class
++        to directly query info attributes.
++
++    .. attribute:: profile.info
++
++        Lower case versions of the :class:`profiling_info` constants
++        may be used as attributes on the attribute `profile` of this
++        class to directly query profiling info.
++
++        For example, you may use *evt.profile.end* instead of
++        *evt.get_profiling_info(pyopencl.profiling_info.END)*.
++
++    .. method:: get_info(param)
++
++        See :class:`event_info` for values of *param*.
++
++    .. method:: get_profiling_info(param)
++
++        See :class:`profiling_info` for values of *param*.
++        See :attr:`profile` for an easier way of obtaining
++        the same information.
++
++    .. method:: wait()
++
++    .. automethod:: from_int_ptr
++    .. autoattribute:: int_ptr
++
++    |comparable|
++
++.. function:: wait_for_events(events)
++
++.. function:: enqueue_barrier(queue, wait_for=None)
++
++    Enqueues a barrier operation. which ensures that all queued commands in
++    command_queue have finished execution. This command is a synchronization
++    point.
++
++    .. versionadded:: 0.91.5
++    .. versionchanged:: 2011.2
++        Takes *wait_for* and returns an :class:`Event`
++
++.. function:: enqueue_marker(queue, wait_for=None)
++
++    Returns an :class:`Event`.
++
++    .. versionchanged:: 2011.2
++        Takes *wait_for*.
++
++.. class:: UserEvent(context)
++
++    A subclass of :class:`Event`. Only available with OpenCL 1.1 and newer.
++
++    .. versionadded:: 0.92
++
++    .. method:: set_status(status)
++
++        See :class:`command_execution_status` for possible values of *status*.
++
++.. class:: NannyEvent
++
++    Transfers between host and device return events of this type. They hold
++    a reference to the host-side buffer and wait for the transfer to complete
++    when they are freed. Therefore, they can safely release the reference to
++    the object they're guarding upon destruction.
++
++    A subclass of :class:`Event`.
++
++    .. versionadded:: 2011.2
++
++    .. method:: get_ward()
++
++    .. method:: wait()
++
++        In addition to performing the same wait as :meth:`Event.wait()`, this
++        method also releases the reference to the guarded object.
++
++Memory
++------
++
++.. class:: MemoryObject
++
++    .. attribute:: info
++
++        Lower case versions of the :class:`mem_info` constants
++        may be used as attributes on instances of this class
++        to directly query info attributes.
++
++    .. attribute:: hostbuf
++
++    .. method:: get_info(param)
++
++        See :class:`mem_info` for values of *param*.
++
++    .. method:: release()
++
++    .. method:: get_host_array(shape, dtype, order="C")
++
++        Return the memory object's associated host memory
++        area as a :class:`numpy.ndarray` of the given *shape*,
++        *dtype* and *order*.
++
++    .. automethod:: from_int_ptr
++    .. autoattribute:: int_ptr
++
++    |comparable|
++
++.. function:: enqueue_migrate_mem_objects(queue, mem_objects, flags=0, wait_for=None)
++
++    :param flags: from :class:`mem_migration_flags`
++
++    .. versionadded:: 2011.2
++
++    Only available with CL 1.2.
++
++.. function:: enqueue_migrate_mem_object_ext(queue, mem_objects, flags=0, wait_for=None)
++
++    :param flags: from :class:`migrate_mem_object_flags_ext`
++
++    .. versionadded:: 2011.2
++
++    Only available with the `cl_ext_migrate_memobject`
++    extension.
++
++Buffers
++^^^^^^^
++
++.. class:: Buffer(context, flags, size=0, hostbuf=None)
++
++    Create a :class:`Buffer`.
++    See :class:`mem_flags` for values of *flags*.
++    If *hostbuf* is specified, *size* defaults to the size of
++    the specified buffer if it is passed as zero.
++
++    :class:`Buffer` inherits from :class:`MemoryObject`.
++
++    .. note::
++
++        Python also defines a type of `buffer object
++        <https://docs.python.org/3.4/c-api/buffer.html>`_,
++        and PyOpenCL interacts with those, too, as the host-side
++        target of :func:`enqueue_copy`. Make sure to always be
++        clear on whether a :class:`Buffer` or a Python buffer
++        object is needed.
++
++    Note that actual memory allocation in OpenCL may be deferred.
++    Buffers are attached to a :class:`Context` and are only
++    moved to a device once the buffer is used on that device.
++    That is also the point when out-of-memory errors will occur.
++    If you'd like to be sure that there's enough memory for
++    your allocation, either use :func:`enqueue_migrate_mem_objects`
++    (if available) or simply perform a small transfer to the
++    buffer. See also :class:`pyopencl.tools.ImmediateAllocator`.
++
++    .. method:: get_sub_region(origin, size, flags=0)
++
++        Only available in OpenCL 1.1 and newer.
++
++    .. method:: __getitem__(slc)
++
++        *slc* is a :class:`slice` object indicating from which byte index range
++        a sub-buffer is to be created. The *flags* argument of
++        :meth:`get_sub_region` is set to the same flags with which *self* was
++        created.
++
++
++.. function:: enqueue_fill_buffer(queue, mem, pattern, offset, size, wait_for=None)
++
++    :arg pattern: a buffer object (likely a :class:`numpy.ndarray`)
++
++    |std-enqueue-blurb|
++
++    Only available with CL 1.2.
++
++    .. versionadded:: 2011.2
++
++Image Formats
++^^^^^^^^^^^^^
++
++.. class:: ImageFormat([channel_order, channel_type])
++
++
++    .. attribute:: channel_order
++
++        See :class:`channel_order` for possible values.
++
++    .. attribute:: channel_data_type
++
++        See :class:`channel_type` for possible values.
++
++    .. attribute:: channel_count
++
++        .. versionadded:: 0.91.5
++
++    .. attribute:: dtype_size
++
++        .. versionadded:: 0.91.5
++
++    .. attribute:: itemsize
++
++        .. versionadded:: 0.91.5
++
++    .. method:: __repr__
++
++        Returns a :class:`str` representation of the image format.
++
++        .. versionadded:: 0.91
++
++    |comparable|
++
++    .. versionchanged:: 0.91
++
++        Constructor arguments added.
++
++    .. versionchanged:: 2013.2
++
++        :class:`ImageFormat` was made comparable and hashable
++
++.. function:: get_supported_image_formats(context, flags, image_type)
++
++    See :class:`mem_flags` for possible values of *flags*
++    and :class:`mem_object_type` for possible values of *image_type*.
++
++Images
++^^^^^^
++.. class:: Image(context, flags, format, shape=None, pitches=None, hostbuf=None, is_array=False, buffer=None):
++
++    See :class:`mem_flags` for values of *flags*.
++    *shape* is a 2- or 3-tuple. *format* is an instance of :class:`ImageFormat`.
++    *pitches* is a 1-tuple for 2D images and a 2-tuple for 3D images, indicating
++    the distance in bytes from one scan line to the next, and from one 2D image
++    slice to the next.
++
++    If *hostbuf* is given and *shape* is `None`, then *hostbuf.shape* is
++    used as the *shape* parameter.
++
++    :class:`Image` inherits from :class:`MemoryObject`.
++
++    .. note::
++
++        If you want to load images from :mod:`numpy.ndarray` instances or read images
++        back into them, be aware that OpenCL images expect the *x* dimension to vary
++        fastest, whereas in the default (C) order of :mod:`numpy` arrays, the last index
++        varies fastest. If your array is arranged in the wrong order in memory,
++        there are two possible fixes for this:
++
++        * Convert the array to Fortran (column-major) order using :func:`numpy.asarray`.
++
++        * Pass *ary.T.copy()* to the image creation function.
++
++    .. versionadded:: 0.91
++
++    .. versionchanged:: 2011.2
++        Added *is_array* and *buffer*, which are only available on CL 1.2 and newer.
++
++    .. attribute:: info
++
++        Lower case versions of the :class:`mem_info`
++        and :class:`image_info` constants
++        may be used as attributes on instances of this class
++        to directly query info attributes.
++
++    .. attribute:: shape
++
++        Return the value of the *shape* constructor argument as a :class:`tuple`.
++
++    .. method:: get_image_info(param)
++
++        See :class:`image_info` for values of *param*.
++
++    .. method:: release()
++
++    |comparable|
++
++.. function:: image_from_array(ctx, ary, num_channels=None, mode="r", norm_int=False)
++
++    Build a 2D or 3D :class:`Image` from the :class:`numpy.ndarray` *ary*.  If
++    *num_channels* is greater than one, the last dimension of *ary* must be
++    identical to *num_channels*. *ary* must be in C order. If *num_channels* is
++    not given, it defaults to 1 for scalar types and the number of entries
++    for :ref:`vector-types`.
++
++    The :class:`ImageFormat` is chosen as the first *num_channels* components
++    of "RGBA".
++
++    :param mode: "r" or "w" for read/write
++
++    .. note::
++
++        When reading from the image object, the indices passed to `read_imagef` are
++        in the reverse order from what they would be when accessing *ary* from
++        Python.
++
++    If *norm_int* is `True`, then the integer values are normalized to a floating
++    point scale of 0..1 when read.
++
++    .. versionadded:: 2011.2
++
++.. function:: enqueue_fill_image(queue, mem, color, origin, region, wait_for=None)
++
++    :arg color: a buffer object (likely a :class:`numpy.ndarray`)
++
++    |std-enqueue-blurb|
++
++    Only available with CL 1.2.
++
++    .. versionadded:: 2011.2
++
++Transfers
++^^^^^^^^^
++
++.. autofunction:: enqueue_copy(queue, dest, src, **kwargs)
++
++Mapping Memory into Host Address Space
++^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
++
++.. class:: MemoryMap
++
++    .. method:: release(queue=None, wait_for=None)
++
++.. function:: enqueue_map_buffer(queue, buf, flags, offset, shape, dtype, order="C", strides=None, wait_for=None, is_blocking=True)
++
++    |explain-waitfor|
++    *shape*, *dtype*, and *order* have the same meaning
++    as in :func:`numpy.empty`.
++    See :class:`map_flags` for possible values of *flags*.
++    *strides*, if given, overrides *order*.
++
++    :return: a tuple *(array, event)*. *array* is a
++        :class:`numpy.ndarray` representing the host side
++        of the map. Its *.base* member contains a
++        :class:`MemoryMap`.
++
++    .. versionchanged:: 2011.1
++        *is_blocking* now defaults to True.
++
++    .. versionchanged:: 2013.1
++        *order* now defaults to "C".
++
++    .. versionchanged:: 2013.2
++        Added *strides* argument.
++
++.. function:: enqueue_map_image(queue, buf, flags, origin, region, shape, dtype, order="C", strides=None, wait_for=None, is_blocking=True)
++
++    |explain-waitfor|
++    *shape*, *dtype*, and *order* have the same meaning
++    as in :func:`numpy.empty`.
++    See :class:`map_flags` for possible values of *flags*.
++    *strides*, if given, overrides *order*.
++
++    :return: a tuple *(array, event)*. *array* is a
++        :class:`numpy.ndarray` representing the host side
++        of the map. Its *.base* member contains a
++        :class:`MemoryMap`.
++
++    .. versionchanged:: 2011.1
++        *is_blocking* now defaults to True.
++
++    .. versionchanged:: 2013.1
++        *order* now defaults to "C".
++
++    .. versionchanged:: 2013.2
++        Added *strides* argument.
++
++Samplers
++^^^^^^^^
++
++.. class:: Sampler(context, normalized_coords, addressing_mode, filter_mode)
++
++    *normalized_coords* is a :class:`bool` indicating whether
++    to use coordinates between 0 and 1 (*True*) or the texture's
++    natural pixel size (*False*).
++    See :class:`addressing_mode` and :class:`filter_mode` for possible
++    argument values.
++
++    .. attribute:: info
++
++        Lower case versions of the :class:`sampler_info` constants
++        may be used as attributes on instances of this class
++        to directly query info attributes.
++
++    .. method:: get_info(param)
++
++        See :class:`sampler_info` for values of *param*.
++
++    .. automethod:: from_int_ptr
++    .. autoattribute:: int_ptr
++
++    |comparable|
++
++Programs and Kernels
++--------------------
++
++.. class:: Program(context, src)
++           Program(context, devices, binaries)
++
++    *binaries* must contain one binary for each entry in *devices*.
++
++    .. attribute:: info
++
++        Lower case versions of the :class:`program_info` constants
++        may be used as attributes on instances of this class
++        to directly query info attributes.
++
++    .. method:: get_info(param)
++
++        See :class:`program_info` for values of *param*.
++
++    .. method:: get_build_info(device, param)
++
++        See :class:`program_build_info` for values of *param*.
++
++    .. method:: build(options=[], devices=None)
++
++        *options* is a string of compiler flags.
++        Returns *self*.
++
++        By default, built binaries are cached in an on-disk cache called
++        :file:`pyopencl-compiler-cache-vN-uidNAME-pyVERSION` in the directory
++        returned by :func:`tempfile.gettempdir`.  By setting the environment
++        variable :envvar:`PYOPENCL_NO_CACHE` to any non-empty value, this
++        caching is suppressed.  Any options found in the environment variable
++        :envvar:`PYOPENCL_BUILD_OPTIONS` will be appended to *options*.
++
++        .. versionchanged:: 2011.1
++            *options* may now also be a :class:`list` of :class:`str`.
++
++        .. versionchanged:: 2013.1
++            Added :envvar:`PYOPENCL_NO_CACHE`.
++            Added :envvar:`PYOPENCL_BUILD_OPTIONS`.
++
++    .. method:: compile(self, options=[], devices=None, headers=[])
++
++        :param headers: a list of tuples *(name, program)*.
++
++        Only available with CL 1.2.
++
++        .. versionadded:: 2011.2
++
++    .. attribute:: kernel_name
++
++        :class:`Kernel` objects can be produced from a built
++        (see :meth:`build`) program simply by attribute lookup.
++
++        .. note::
++
++            The :class:`program_info` attributes live
++            in the same name space and take precedence over
++            :class:`Kernel` names.
++
++    .. method:: all_kernels()
++
++        Returns a list of all :class:`Kernel` objects in the :class:`Program`.
++
++    .. automethod:: from_int_ptr
++    .. autoattribute:: int_ptr
++
++    |comparable|
++
++.. function:: create_program_with_built_in_kernels(context, devices, kernel_names)
++
++    Only available with CL 1.2.
++
++    .. versionadded:: 2011.2
++
++.. function:: link_program(context, programs, options=[], devices=None)
++
++    Only available with CL 1.2.
++
++    .. versionadded:: 2011.2
++
++.. function:: unload_platform_compiler(platform)
++
++    Only available with CL 1.2.
++
++    .. versionadded:: 2011.2
++
++.. class:: Kernel(program, name)
++
++    .. attribute:: info
++
++        Lower case versions of the :class:`kernel_info` constants
++        may be used as attributes on instances of this class
++        to directly query info attributes.
++
++    .. method:: get_info(param)
++
++        See :class:`kernel_info` for values of *param*.
++
++    .. method:: get_work_group_info(param, device)
++
++        See :class:`kernel_work_group_info` for values of *param*.
++
++    .. method:: get_arg_info(arg_index, param)
++
++        See :class:`kernel_arg_info` for values of *param*.
++
++        Only available in OpenCL 1.2 and newer.
++
++    .. method:: set_arg(self, index, arg)
++
++        *arg* may be
++
++        * `None`: This may be passed for `__global` memory references
++          to pass a NULL pointer to the kernel.
++        * Anything that satisfies the Python buffer interface,
++          in particular :class:`numpy.ndarray`, :class:`str`,
++          or :mod:`numpy`'s sized scalars, such as :class:`numpy.int32`
++          or :class:`numpy.float64`.
++
++          .. note::
++
++              Note that Python's own :class:`int` or :class:`float`
++              objects will not work out of the box. See
++              :meth:`Kernel.set_scalar_arg_dtypes` for a way to make
++              them work. Alternatively, the standard library module
++              :mod:`struct` can be used to convert Python's native
++              number types to binary data in a :class:`str`.
++
++        * An instance of :class:`MemoryObject`. (e.g. :class:`Buffer`,
++          :class:`Image`, etc.)
++        * An instance of :class:`LocalMemory`.
++        * An instance of :class:`Sampler`.
++
++    .. method:: set_args(self, *args)
++
++        Invoke :meth:`set_arg` on each element of *args* in turn.
++
++        .. versionadded:: 0.92
++
++    .. method:: set_scalar_arg_dtypes(arg_dtypes)
++
++        Inform the wrapper about the sized types of scalar
++        :class:`Kernel` arguments. For each argument,
++        *arg_dtypes* contains an entry. For non-scalars,
++        this must be *None*. For scalars, it must be an
++        object acceptable to the :class:`numpy.dtype`
++        constructor, indicating that the corresponding
++        scalar argument is of that type.
++
++        After invoking this function with the proper information,
++        most suitable number types will automatically be
++        cast to the right type for kernel invocation.
++
++        .. note ::
++
++            The information set by this rountine is attached to a single kernel
++            instance. A new kernel instance is created every time you use
++            `program.kernel` attribute access. The following will therefore not
++            work::
++
++               prg = cl.Program(...).build()
++               prg.kernel.set_scalar_arg_dtypes(...)
++               prg.kernel(queue, n_globals, None, args)
++
++
++    .. method:: __call__(queue, global_size, local_size, *args, global_offset=None, wait_for=None, g_times_l=False)
++
++        Use :func:`enqueue_nd_range_kernel` to enqueue a kernel execution, after using
++        :meth:`set_args` to set each argument in turn. See the documentation for
++        :meth:`set_arg` to see what argument types are allowed.
++        |std-enqueue-blurb|
++
++        *None* may be passed for local_size.
++
++        If *g_times_l* is specified, the global size will be multiplied by the
++        local size. (which makes the behavior more like Nvidia CUDA) In this case,
++        *global_size* and *local_size* also do not have to have the same number
++        of dimensions.
++
++        .. note::
++
++            :meth:`__call__` is *not* thread-safe. It sets the arguments using :meth:`set_args`
++            and then runs :func:`enqueue_nd_range_kernel`. Another thread could race it
++            in doing the same things, with undefined outcome. This issue is inherited
++            from the C-level OpenCL API. The recommended solution is to make a kernel
++            (i.e. access `prg.kernel_name`, which corresponds to making a new kernel)
++            for every thread that may enqueue calls to the kernel.
++
++            A solution involving implicit locks was discussed and decided against on the
++            mailing list in `October 2012
++            <http://lists.tiker.net/pipermail/pyopencl/2012-October/001311.html>`_.
++
++        .. versionchanged:: 0.92
++            *local_size* was promoted to third positional argument from being a
++            keyword argument. The old keyword argument usage will continue to
++            be accepted with a warning throughout the 0.92 release cycle.
++            This is a backward-compatible change (just barely!) because
++            *local_size* as third positional argument can only be a
++            :class:`tuple` or *None*.  :class:`tuple` instances are never valid
++            :class:`Kernel` arguments, and *None* is valid as an argument, but
++            its treatment in the wrapper had a bug (now fixed) that prevented
++            it from working.
++
++        .. versionchanged:: 2011.1
++            Added the *g_times_l* keyword arg.
++
++    .. method:: capture_call(filename, queue, global_size, local_size, *args, global_offset=None, wait_for=None, g_times_l=False)
++
++        This method supports the exact same interface as :meth:`__call__`, but
++        instead of invoking the kernel, it writes a self-contained PyOpenCL program
++        to *filename* that reproduces this invocation. Data and kernel source code
++        will be packaged up in *filename*'s source code.
++
++        This is mainly intended as a debugging aid. For example, it can be used
++        to automate the task of creating a small, self-contained test case for
++        an observed problem. It can also help separate a misbehaving kernel from
++        a potentially large or time-consuming outer code.
++
++        To use, simply change::
++
++            evt = my_kernel(queue, gsize, lsize, arg1, arg2, ...)
++
++        to::
++
++            evt = my_kernel.capture_call("bug.py", queue, gsize, lsize, arg1, arg2, ...)
++
++        .. versionadded:: 2013.1
++
++    .. automethod:: from_int_ptr
++    .. autoattribute:: int_ptr
++
++    |comparable|
++
++.. class:: LocalMemory(size)
++
++    A helper class to pass `__local` memory arguments to kernels.
++
++    .. versionadded:: 0.91.2
++
++    .. attribute:: size
++
++        The size of local buffer in bytes to be provided.
++
++.. function:: enqueue_nd_range_kernel(queue, kernel, global_work_size, local_work_size, global_work_offset=None, wait_for=None, g_times_l=False)
++
++    |std-enqueue-blurb|
++
++    If *g_times_l* is specified, the global size will be multiplied by the
++    local size. (which makes the behavior more like Nvidia CUDA) In this case,
++    *global_size* and *local_size* also do not have to have the same number
++    of dimensions.
++
++    .. versionchanged:: 2011.1
++        Added the *g_times_l* keyword arg.
++
++
++.. function:: enqueue_task(queue, kernel, wait_for=None)
++
++    |std-enqueue-blurb|
++
++.. _gl-interop:
++
++GL Interoperability
++-------------------
++
++Functionality in this section is only available when PyOpenCL is compiled
++with GL support. See :func:`have_gl`.
++
++.. versionadded:: 0.91
++
++.. function:: have_gl()
++
++    Return *True* if PyOpenCL was compiled with OpenGL interoperability, otherwise *False*.
++
++.. function:: get_gl_sharing_context_properties()
++
++    Return a :class:`list` of :class:`context_properties` that will
++    allow a newly created context to share the currently active GL
++    context.
++
++.. function:: get_apple_cgl_share_group()
++
++    Get share group handle for current CGL context.
++
++    Apple OS X only.
++
++    .. versionadded:: 2011.1
++
++.. class:: GLBuffer(context, flags, bufobj)
++
++    :class:`GLBuffer` inherits from :class:`MemoryObject`.
++
++    .. attribute:: gl_object
++
++.. class:: GLRenderBuffer(context, flags, bufobj)
++
++    :class:`GLRenderBuffer` inherits from :class:`MemoryObject`.
++
++    .. attribute:: gl_object
++
++.. class:: GLTexture(context, flags, texture_target, miplevel, texture, dims)
++
++    *dims* is either 2 or 3.
++    :class:`GLTexture` inherits from :class:`Image`.
++
++    .. attribute:: gl_object
++
++    .. method:: get_gl_texture_info(param)
++
++        See :class:`gl_texture_info` for values of *param*.  Only available when PyOpenCL is compiled with GL support. See :func:`have_gl`.
++
++.. function:: enqueue_acquire_gl_objects(queue, mem_objects, wait_for=None)
++
++    *mem_objects* is a list of :class:`MemoryObject` instances.
++    |std-enqueue-blurb|
++
++.. function:: enqueue_release_gl_objects(queue, mem_objects, wait_for=None)
++
++    *mem_objects* is a list of :class:`MemoryObject` instances. |std-enqueue-blurb|
++
++.. function:: get_gl_context_info_khr(properties, param_name, platform=None)
++
++    Get information on which CL device corresponds to a given
++    GL/EGL/WGL/CGL device.
++
++    See the :class:`Context` constructor for the meaning of
++    *properties* and :class:`gl_context_info` for *param_name*.
++
++
++    .. versionchanged:: 2011.2
++        Accepts the *platform* argument.  Using *platform* equal to None is
++        deprecated as of PyOpenCL 2011.2.
+diff -Nur pyopencl-2015.1/doc/_static/akdoc.css pyopencl/doc/_static/akdoc.css
+--- pyopencl-2015.1/doc/_static/akdoc.css	1970-01-01 01:00:00.000000000 +0100
++++ pyopencl/doc/_static/akdoc.css	2015-03-05 20:44:55.746519776 +0100
+@@ -0,0 +1,59 @@
++pre {
++  line-height: 110%;
++}
++
++.footer {
++  background-color: #eee;
++}
++
++body > div.container {
++  margin-top:10px;
++}
++
++dd {
++  margin-left: 40px;
++}
++
++tt.descname {
++  font-size: 100%;
++}
++
++code {
++  color: rgb(51,51,51);
++}
++
++h1 {
++  padding-bottom:7px;
++  border-bottom: 1px solid #ccc;
++}
++
++h2 {
++  padding-bottom:5px;
++  border-bottom: 1px solid #ccc;
++}
++
++h3 {
++  padding-bottom:5px;
++  border-bottom: 1px solid #ccc;
++}
++
++.rubric {
++  font-size: 120%;
++  padding-bottom:1px;
++  border-bottom: 1px solid #ccc;
++}
++
++.headerlink {
++  padding-left: 1ex;
++  padding-right: 1ex;
++}
++
++a.headerlink:hover {
++  text-decoration: none;
++}
++
++blockquote p {
++  font-size: 100%;
++  font-weight: normal;
++  line-height: normal;
++};
+diff -Nur pyopencl-2015.1/doc/subst.rst pyopencl/doc/subst.rst
+--- pyopencl-2015.1/doc/subst.rst	1970-01-01 01:00:00.000000000 +0100
++++ pyopencl/doc/subst.rst	2015-03-05 20:44:55.749853109 +0100
+@@ -0,0 +1,13 @@
++.. |comparable| replace:: Instances of this class are hashable, and two
++    instances of this class may be compared using *"=="* and *"!="*.
++    (Hashability was added in version 2011.2.)
++
++.. |buf-iface| replace:: must implement the Python buffer interface.
++    (e.g. by being an :class:`numpy.ndarray`)
++.. |explain-waitfor| replace:: *wait_for*
++    may either be *None* or a list of :class:`pyopencl.Event` instances for
++    whose completion this command waits before starting exeuction.
++.. |std-enqueue-blurb| replace:: Returns a new :class:`pyopencl.Event`. |explain-waitfor|
++
++.. |copy-depr| replace:: **Note:** This function is deprecated as of PyOpenCL 2011.1.
++        Use :func:`enqueue_copy` instead.
+diff -Nur pyopencl-2015.1/doc/_templates/layout.html pyopencl/doc/_templates/layout.html
+--- pyopencl-2015.1/doc/_templates/layout.html	1970-01-01 01:00:00.000000000 +0100
++++ pyopencl/doc/_templates/layout.html	2015-03-05 20:44:55.746519776 +0100
+@@ -0,0 +1,2 @@
++{% extends "!layout.html" %}
++{% set bootswatch_css_custom = ['_static/akdoc.css']%}
+diff -Nur pyopencl-2015.1/doc/tools.rst pyopencl/doc/tools.rst
+--- pyopencl-2015.1/doc/tools.rst	1970-01-01 01:00:00.000000000 +0100
++++ pyopencl/doc/tools.rst	2015-03-05 20:44:55.749853109 +0100
+@@ -0,0 +1,140 @@
++Built-in Utilities
++==================
++
++.. module:: pyopencl.tools
++
++.. _memory-pools:
++
++Memory Pools
++------------
++
++The constructor :func:`pyopencl.Buffer` can consume a fairly large amount of
++processing time if it is invoked very frequently. For example, code based on
++:class:`pyopencl.array.Array` can easily run into this issue because a
++fresh memory area is allocated for each intermediate result. Memory pools are a
++remedy for this problem based on the observation that often many of the block
++allocations are of the same sizes as previously used ones.
++
++Then, instead of fully returning the memory to the system and incurring the 
++associated reallocation overhead, the pool holds on to the memory and uses it
++to satisfy future allocations of similarly-sized blocks. The pool reacts
++appropriately to out-of-memory conditions as long as all memory allocations
++are made through it. Allocations performed from outside of the pool may run
++into spurious out-of-memory conditions due to the pool owning much or all of
++the available memory.
++
++Using :class:`pyopencl.array.Array` instances with a :class:`MemoryPool` is
++not complicated::
++
++    mem_pool = cl_tools.MemoryPool(cl_tools.ImmediateAllocator(queue))
++    a_dev = cl_array.arange(queue, 2000, dtype=np.float32, allocator=mem_pool)
++
++.. class:: PooledBuffer
++
++    An object representing a :class:`MemoryPool`-based allocation of
++    device memory.  Once this object is deleted, its associated device
++    memory is returned to the pool. This supports the same interface
++    as :class:`pyopencl.Buffer`.
++
++.. class:: DeferredAllocator(context, mem_flags=pyopencl.mem_flags.READ_WRITE)
++
++    *mem_flags* takes its values from :class:`pyopencl.mem_flags` and corresponds
++    to the *flags* argument of :class:`pyopencl.Buffer`. DeferredAllocator
++    has the same semantics as regular OpenCL buffer allocation, i.e. it may
++    promise memory to be available that may (in any call to a buffer-using
++    CL function) turn out to not exist later on. (Allocations in CL are
++    bound to contexts, not devices, and memory availability depends on which
++    device the buffer is used with.)
++
++    .. versionchanged::
++        In version 2013.1, :class:`CLAllocator` was deprecated and replaced
++        by :class:`DeferredAllocator`.
++
++    .. method:: __call__(size)
++
++        Allocate a :class:`pyopencl.Buffer` of the given *size*.
++
++.. class:: ImmediateAllocator(queue, mem_flags=pyopencl.mem_flags.READ_WRITE)
++
++    *mem_flags* takes its values from :class:`pyopencl.mem_flags` and corresponds
++    to the *flags* argument of :class:`pyopencl.Buffer`. DeferredAllocator
++    has the same semantics as regular OpenCL buffer allocation, i.e. it may
++    promise memory to be available that later on (in any call to a buffer-using
++    CL function).
++
++    .. versionadded:: 2013.1
++
++    .. method:: __call__(size)
++
++        Allocate a :class:`pyopencl.Buffer` of the given *size*.
++
++.. class:: MemoryPool(allocator)
++
++    A memory pool for OpenCL device memory. *allocator* must be an instance of
++    one of the above classes, and should be an :class:`ImmediateAllocator`.
++    The memory pool assumes that allocation failures are reported
++    by the allocator immediately, and not in the OpenCL-typical
++    deferred manner.
++
++    .. attribute:: held_blocks
++
++        The number of unused blocks being held by this pool.
++
++    .. attribute:: active_blocks
++
++        The number of blocks in active use that have been allocated
++        through this pool.
++
++    .. method:: allocate(size)
++
++        Return a :class:`PooledBuffer` of the given *size*.
++
++    .. method:: __call__(size)
++
++        Synoynm for :meth:`allocate` to match :class:`CLAllocator` interface.
++
++        .. versionadded: 2011.2
++
++    .. method:: free_held
++
++        Free all unused memory that the pool is currently holding.
++
++    .. method:: stop_holding
++
++        Instruct the memory to start immediately freeing memory returned
++        to it, instead of holding it for future allocations.
++        Implicitly calls :meth:`free_held`.
++        This is useful as a cleanup action when a memory pool falls out
++        of use.
++
++CL-Object-dependent Caching
++---------------------------
++
++.. autofunction:: first_arg_dependent_memoize
++.. autofunction:: clear_first_arg_caches
++
++Testing
++-------
++
++.. function:: pytest_generate_tests_for_pyopencl(metafunc)
++
++    Using the line::
++
++        from pyopencl.tools import pytest_generate_tests_for_pyopencl \
++                as pytest_generate_tests
++
++    in your `pytest <http://pytest.org>`_ test scripts allows you to use the
++    arguments *ctx_factory*, *device*, or *platform* in your test functions,
++    and they will automatically be run for each OpenCL device/platform in the
++    system, as appropriate.
++
++    The following two environment variables are also supported to control
++    device/platform choice::
++
++        PYOPENCL_TEST=0:0,1;intel=i5,i7
++
++Device Characterization
++-----------------------
++
++.. automodule:: pyopencl.characterize
++    :members:
+diff -Nur pyopencl-2015.1/doc/upload-docs.sh pyopencl/doc/upload-docs.sh
+--- pyopencl-2015.1/doc/upload-docs.sh	1970-01-01 01:00:00.000000000 +0100
++++ pyopencl/doc/upload-docs.sh	2015-03-05 20:44:55.749853109 +0100
+@@ -0,0 +1,3 @@
++#! /bin/sh
++
++rsync --progress --verbose --archive --delete _build/html/* doc-upload:doc/pyopencl
================================================================

---- gitweb:

http://git.pld-linux.org/gitweb.cgi/packages/python-pyopencl.git/commitdiff/f58472434a95bd1a0323c0d4a821ec7d0c45a95c




More information about the pld-cvs-commit mailing list