[Date Prev][Date Next] [Thread Prev][Thread Next] [Date Index] [Thread Index]

Re: Fixes for python-pyopencl and new upsteam release



Dnia 2010-10-21, czw o godzinie 20:07 +0100, Adam D. Barratt pisze:
> On Thu, 2010-10-21 at 19:36 +0200, Tomasz Rybak wrote:
> > Dnia 2010-10-18, pon o godzinie 21:13 +0100, Adam D. Barratt pisze:
> > > On Tue, 2010-10-12 at 11:23 +0200, Tomasz Rybak wrote:
> [...]
> > > > There has been 24 changes in upstream between version in Debian
> > > > and final 0.92. From those maybe 6 are not needed - documentation
> > > > and example fixes, fixing Python setup (which is disabled in package),
> > > > fixing Windows bug. If I would to backport fixes for errors, I would
> > > > need to include more than half of changes made by upstream.
> > > 
> > > It's a little difficult to tell from that description whether the
> > > changes would be acceptable.  Please could you provide a debdiff for the
> > > proposed upload?
> > 
> > Attached.
> > I can also attach full diff of source between version in Squeeze
> > and proposed version, if needed.
> 
> Thanks.  A debdiff between the source packages was what I intended so
> yes, please. :-)
> 
> Regards,
> 
> Adam
> 

Sorry.
pyopencl.debdiff contains debbiff
pyopencl.gitdiff contains diff from upstream git respository
between version in Squeeze and proposed one.

Regards.

-- 
Tomasz Rybak <bogomips@post.pl> GPG/PGP key ID: 2AD5 9860
Fingerprint A481 824E 7DD3 9C0E C40A  488E C654 FB33 2AD5 9860
http://member.acm.org/~tomaszrybak
diff -Nru pyopencl-0.92~beta+git20100709/aksetup_helper.py pyopencl-0.92/aksetup_helper.py
--- pyopencl-0.92~beta+git20100709/aksetup_helper.py	2010-07-10 18:03:27.000000000 +0200
+++ pyopencl-0.92/aksetup_helper.py	2010-10-21 19:10:19.000000000 +0200
@@ -2,8 +2,34 @@
 import distribute_setup
 distribute_setup.use_setuptools()
 
+import setuptools
 from setuptools import Extension
 
+if not hasattr(setuptools, "_distribute"):
+    print("-------------------------------------------------------------------------")
+    print("Setuptools conflict detected.")
+    print("-------------------------------------------------------------------------")
+    print("When I imported setuptools, I did not get the distribute version of")
+    print("setuptools, which is troubling--this package really wants to be used")
+    print("with distribute rather than the old setuptools package. More than likely,")
+    print("you have both distribute and setuptools installed, which is bad.")
+    print("")
+    print("See this page for more information:")
+    print("http://wiki.tiker.net/DistributeVsSetuptools";)
+    print("-------------------------------------------------------------------------")
+    print("I will continue after a short while, fingers crossed.")
+    print("-------------------------------------------------------------------------")
+
+    delay = 10
+
+    from time import sleep
+    import sys
+    while delay:
+        sys.stdout.write("Continuing in %d seconds...   \r" % delay)
+        sys.stdout.flush()
+        delay -= 1
+        sleep(1)
+
 def setup(*args, **kwargs):
     from setuptools import setup
     import traceback
@@ -14,16 +40,16 @@
     except SystemExit:
         raise
     except:
-        print "----------------------------------------------------------------------------"
-        print "Sorry, your build failed. Try rerunning configure.py with different options."
-        print "----------------------------------------------------------------------------"
+        print ("----------------------------------------------------------------------------")
+        print ("Sorry, your build failed. Try rerunning configure.py with different options.")
+        print ("----------------------------------------------------------------------------")
         raise
 
 
 
 
 class NumpyExtension(Extension):
-    # nicked from 
+    # nicked from
     # http://mail.python.org/pipermail/distutils-sig/2007-September/008253.html
     # solution by Michael Hoffmann
     def __init__(self, *args, **kwargs):
@@ -83,7 +109,7 @@
 
 # tools -----------------------------------------------------------------------
 def flatten(list):
-    """For an iterable of sub-iterables, generate each member of each 
+    """For an iterable of sub-iterables, generate each member of each
     sub-iterable in turn, i.e. a flattened version of that super-iterable.
 
     Example: Turn [[a,b,c],[d,e,f]] into [a,b,c,d,e,f].
@@ -110,20 +136,20 @@
 
     if (not schema.have_config() and not schema.have_global_config()
             and warn_about_no_config):
-        print "*************************************************************"
-        print "*** I have detected that you have not run configure.py."
-        print "*************************************************************"
-        print "*** Additionally, no global config files were found."
-        print "*** I will go ahead with the default configuration."
-        print "*** In all likelihood, this will not work out."
-        print "*** "
-        print "*** See README_SETUP.txt for more information."
-        print "*** "
-        print "*** If the build does fail, just re-run configure.py with the"
-        print "*** correct arguments, and then retry. Good luck!"
-        print "*************************************************************"
-        print "*** HIT Ctrl-C NOW IF THIS IS NOT WHAT YOU WANT"
-        print "*************************************************************"
+        print("*************************************************************")
+        print("*** I have detected that you have not run configure.py.")
+        print("*************************************************************")
+        print("*** Additionally, no global config files were found.")
+        print("*** I will go ahead with the default configuration.")
+        print("*** In all likelihood, this will not work out.")
+        print("*** ")
+        print("*** See README_SETUP.txt for more information.")
+        print("*** ")
+        print("*** If the build does fail, just re-run configure.py with the")
+        print("*** correct arguments, and then retry. Good luck!")
+        print("*************************************************************")
+        print("*** HIT Ctrl-C NOW IF THIS IS NOT WHAT YOU WANT")
+        print("*************************************************************")
 
         delay = 10
 
@@ -204,7 +230,7 @@
     return re.subn(r"\$\{([a-zA-Z0-9_]+)\}", my_repl, s)[0]
 
 def expand_value(v, options):
-    if isinstance(v, (str, unicode)):
+    if isinstance(v, str):
         return expand_str(v, options)
     elif isinstance(v, list):
         return [expand_value(i, options) for i in v]
@@ -246,15 +272,15 @@
         self.conf_dir = conf_dir
 
     def get_default_config(self):
-        return dict((opt.name, opt.default) 
+        return dict((opt.name, opt.default)
                 for opt in self.options)
-        
+
     def read_config_from_pyfile(self, filename):
         result = {}
         filevars = {}
-        execfile(filename, filevars)
+        exec(compile(open(filename, "r").read(), filename, "exec"), filevars)
 
-        for key, value in filevars.iteritems():
+        for key, value in filevars.items():
             if key in self.optdict:
                 result[key] = value
 
@@ -265,13 +291,13 @@
         filevars = {}
 
         try:
-            execfile(filename, filevars)
+            exec(compile(open(filename, "r").read(), filename, "exec"), filevars)
         except IOError:
             pass
 
         del filevars["__builtins__"]
 
-        for key, value in config.iteritems():
+        for key, value in config.items():
             if value is not None:
                 filevars[key] = value
 
@@ -296,7 +322,7 @@
         result = self.get_default_config()
 
         import os
-        
+
         confignames = []
         if self.global_conf_file is not None:
             confignames.append(self.global_conf_file)
@@ -328,16 +354,16 @@
         result = self.get_default_config_with_files()
         if os.access(cfile, os.R_OK):
             filevars = {}
-            execfile(cfile, filevars)
+            exec(compile(open(cfile, "r").read(), cfile, "exec"), filevars)
 
-            for key, value in filevars.iteritems():
+            for key, value in filevars.items():
                 if key in self.optdict:
                     result[key] = value
                 elif key == "__builtins__":
                     pass
                 else:
-                    raise KeyError, "invalid config key in %s: %s" % (
-                            cfile, key)
+                    raise KeyError("invalid config key in %s: %s" % (
+                            cfile, key))
 
         expand_options(result)
 
@@ -399,13 +425,13 @@
         return result
 
     def value_to_str(self, default):
-        return default 
+        return default
 
     def add_to_configparser(self, parser, default=None):
         default = default_or(default, self.default)
         default_str = self.value_to_str(default)
         parser.add_option(
-            "--" + self.as_option(), dest=self.name, 
+            "--" + self.as_option(), dest=self.name,
             default=default_str,
             metavar=self.metavar(), help=self.get_help(default))
 
@@ -417,7 +443,7 @@
         option = self.as_option()
 
         if not isinstance(self.default, bool):
-            raise ValueError, "Switch options must have a default"
+            raise ValueError("Switch options must have a default")
 
         if default is None:
             default = self.default
@@ -426,11 +452,11 @@
             action = "store_false"
         else:
             action = "store_true"
-            
+
         parser.add_option(
-            "--" + self.as_option(), 
-            dest=self.name, 
-            help=self.get_help(default), 
+            "--" + self.as_option(),
+            dest=self.name,
+            help=self.get_help(default),
             default=default,
             action=action)
 
@@ -458,7 +484,7 @@
 class IncludeDir(StringListOption):
     def __init__(self, lib_name, default=None, human_name=None, help=None):
         StringListOption.__init__(self, "%s_INC_DIR" % lib_name, default,
-                help=help or ("Include directories for %s" 
+                help=help or ("Include directories for %s"
                 % (human_name or humanize(lib_name))))
 
 class LibraryDir(StringListOption):
@@ -470,14 +496,14 @@
 class Libraries(StringListOption):
     def __init__(self, lib_name, default=None, human_name=None, help=None):
         StringListOption.__init__(self, "%s_LIBNAME" % lib_name, default,
-                help=help or ("Library names for %s (without lib or .so)" 
+                help=help or ("Library names for %s (without lib or .so)"
                 % (human_name or humanize(lib_name))))
 
 class BoostLibraries(Libraries):
     def __init__(self, lib_base_name):
-        Libraries.__init__(self, "BOOST_%s" % lib_base_name.upper(), 
+        Libraries.__init__(self, "BOOST_%s" % lib_base_name.upper(),
                 ["boost_%s-${BOOST_COMPILER}-mt" % lib_base_name],
-                help="Library names for Boost C++ %s library (without lib or .so)" 
+                help="Library names for Boost C++ %s library (without lib or .so)"
                     % humanize(lib_base_name))
 
 def set_up_shipped_boost_if_requested(conf):
@@ -491,22 +517,22 @@
 
     if conf["USE_SHIPPED_BOOST"]:
         if not exists("bpl-subset/bpl_subset/boost/version.hpp"):
-            print >>sys.stderr, "------------------------------------------------------------------------"
-            print >>sys.stderr, "The shipped Boost library was not found, but USE_SHIPPED_BOOST is True."
-            print >>sys.stderr, "(The files should be under bpl-subset/.)"
-            print >>sys.stderr, "------------------------------------------------------------------------"
-            print >>sys.stderr, "If you got this package from git, you probably want to do"
-            print >>sys.stderr, ""
-            print >>sys.stderr, " $ git submodule init"
-            print >>sys.stderr, " $ git submodule update"
-            print >>sys.stderr, ""
-            print >>sys.stderr, "to fetch what you are presently missing. If you got this from"
-            print >>sys.stderr, "a distributed package on the net, that package is broken and"
-            print >>sys.stderr, "should be fixed. For now, I will turn off 'USE_SHIPPED_BOOST'"
-            print >>sys.stderr, "to try and see if the build succeeds that way, but in the long"
-            print >>sys.stderr, "run you might want to either get the missing bits or turn"
-            print >>sys.stderr, "'USE_SHIPPED_BOOST' off."
-            print >>sys.stderr, "------------------------------------------------------------------------"
+            print("------------------------------------------------------------------------")
+            print("The shipped Boost library was not found, but USE_SHIPPED_BOOST is True.")
+            print("(The files should be under bpl-subset/.)")
+            print("------------------------------------------------------------------------")
+            print("If you got this package from git, you probably want to do")
+            print("")
+            print(" $ git submodule init")
+            print(" $ git submodule update")
+            print("")
+            print("to fetch what you are presently missing. If you got this from")
+            print("a distributed package on the net, that package is broken and")
+            print("should be fixed. For now, I will turn off 'USE_SHIPPED_BOOST'")
+            print("to try and see if the build succeeds that way, but in the long")
+            print("run you might want to either get the missing bits or turn")
+            print("'USE_SHIPPED_BOOST' off.")
+            print("------------------------------------------------------------------------")
             conf["USE_SHIPPED_BOOST"] = False
 
             delay = 10
@@ -541,7 +567,7 @@
             source_files += glob(
                     "bpl-subset/bpl_subset/libs/thread/src/pthread/*.cpp")
 
-        return (source_files, 
+        return (source_files,
                 {"BOOST_MULTI_INDEX_DISABLE_SERIALIZATION": 1}
                 )
     else:
@@ -552,7 +578,7 @@
     return [
         IncludeDir("BOOST", []),
         LibraryDir("BOOST", []),
-        Option("BOOST_COMPILER", default="gcc43", 
+        Option("BOOST_COMPILER", default="gcc43",
             help="The compiler with which Boost C++ was compiled, e.g. gcc43"),
         ]
 
@@ -568,29 +594,29 @@
     from setup import get_config_schema
     schema = get_config_schema()
     if schema.have_config():
-        print "************************************************************"
-        print "*** I have detected that you have already run configure."
-        print "*** I'm taking the configured values as defaults for this"
-        print "*** configure run. If you don't want this, delete the file"
-        print "*** %s." % schema.get_conf_file()
-        print "************************************************************"
+        print("************************************************************")
+        print("*** I have detected that you have already run configure.")
+        print("*** I'm taking the configured values as defaults for this")
+        print("*** configure run. If you don't want this, delete the file")
+        print("*** %s." % schema.get_conf_file())
+        print("************************************************************")
 
     import sys
 
     description = "generate a configuration file for this software package"
     parser = OptionParser(description=description)
     parser.add_option(
-	    "--python-exe", dest="python_exe", default=sys.executable,
-	    help="Which Python interpreter to use", metavar="PATH")
+            "--python-exe", dest="python_exe", default=sys.executable,
+            help="Which Python interpreter to use", metavar="PATH")
 
     parser.add_option("--prefix", default=None,
-	    help="Ignored")
+            help="Ignored")
     parser.add_option("--enable-shared", help="Ignored", action="store_false")
     parser.add_option("--disable-static", help="Ignored", action="store_false")
-    parser.add_option("--update-user", help="Update user config file (%s)" % schema.user_conf_file, 
+    parser.add_option("--update-user", help="Update user config file (%s)" % schema.user_conf_file,
             action="store_true")
-    parser.add_option("--update-global", 
-            help="Update global config file (%s)" % schema.global_conf_file, 
+    parser.add_option("--update-global",
+            help="Update global config file (%s)" % schema.global_conf_file,
             action="store_true")
 
     schema.add_to_configparser(parser, schema.read_config())
diff -Nru pyopencl-0.92~beta+git20100709/debian/changelog pyopencl-0.92/debian/changelog
--- pyopencl-0.92~beta+git20100709/debian/changelog	2010-07-12 22:52:35.000000000 +0200
+++ pyopencl-0.92/debian/changelog	2010-10-21 19:26:59.000000000 +0200
@@ -1,3 +1,23 @@
+pyopencl (0.92-1) UNRELEASED; urgency=low
+
+  * New upstream release
+  * Fixed Vcs-* URLs in debian/control; now they point to Alioth
+    instead of upstream git repository (Closes: #599875)
+  * Fixed Maintainer field in debian/copyrigth (Closes: #588873)
+  * Added missing licence for pyopencl/clrandom.py file (Closes: #599874)
+  * Fixed FTBFS on i386 (Closes: #599782)
+  * Forced build scripts not to use boost libraries included in the source
+  * Removed messages about missing Boost libraries; they are superfluous
+    as package is using Debian ones
+  * Update Standards-Version to 3.9.1; no changes necessary.
+  * Switch to debhelper compat level 8; no changes necessary.
+  * Changed dependencies:
+    * Depend on libnvidia-compiler instead of transitional libnvidia-compiler1
+    * Build-depend on nvidia-libopencl1 and khronos-opencl-headers
+      instead of nvidia-libopencl1-dev
+
+ -- Tomasz Rybak <bogomips@post.pl>  Mon, 18 Oct 2010 21:13:06 +0200
+
 pyopencl (0.92~beta+git20100709-1) unstable; urgency=low
 
   * Initial release
diff -Nru pyopencl-0.92~beta+git20100709/debian/compat pyopencl-0.92/debian/compat
--- pyopencl-0.92~beta+git20100709/debian/compat	2010-07-10 18:32:03.000000000 +0200
+++ pyopencl-0.92/debian/compat	2010-10-21 19:26:59.000000000 +0200
@@ -1 +1 @@
-7
+8
diff -Nru pyopencl-0.92~beta+git20100709/debian/control pyopencl-0.92/debian/control
--- pyopencl-0.92~beta+git20100709/debian/control	2010-07-10 18:32:03.000000000 +0200
+++ pyopencl-0.92/debian/control	2010-10-21 19:26:59.000000000 +0200
@@ -2,27 +2,29 @@
 Section: contrib/python
 Priority: optional
 Maintainer: Tomasz Rybak <bogomips@post.pl>
-Build-Depends: debhelper (>= 7.0.50~),
+Build-Depends: debhelper (>= 8),
  python-support,
  python-all-dev,
  python-setuptools,
- nvidia-libopencl1-dev | libopencl1-dev,
+ nvidia-libopencl1-dev | khronos-opencl-headers | libopencl1-dev | nvidia-current-dev,
+ nvidia-libopencl1,
  libboost-python-dev,
  mesa-common-dev,
  python-numpy,
  python-matplotlib,
  python-sphinx,
  python-pytools (>= 7)
-Standards-Version: 3.8.4
+Standards-Version: 3.9.1
 XS-Python-Version: >= 2.5
 Homepage: http://mathema.tician.de/software/pyopencl
-Vcs-Git: http://git.tiker.net/trees/pyopencl.git
-Vcs-Browser: http://git.tiker.net/?p=pyopencl.git;a=summary
+Vcs-Git: git://git.debian.org/git/collab-maint/python-pyopencl.git
+Vcs-Browser: http://git.debian.org/?p=collab-maint/python-pyopencl.git
 
 Package: python-pyopencl
 Architecture: amd64 i386
 Depends: ${shlibs:Depends}, ${misc:Depends}, ${python:Depends},
- libnvidia-compiler1,
+ libnvidia-compiler1 | libnvidia-compiler,
+ nvidia-libopencl1,
  nvidia-opencl-common,
  python-numpy,
  python-matplotlib,
diff -Nru pyopencl-0.92~beta+git20100709/debian/copyright pyopencl-0.92/debian/copyright
--- pyopencl-0.92~beta+git20100709/debian/copyright	2010-07-10 18:32:03.000000000 +0200
+++ pyopencl-0.92/debian/copyright	2010-10-21 19:26:59.000000000 +0200
@@ -1,6 +1,6 @@
 Format-Specification: http://svn.debian.org/wsvn/dep/web/deps/dep5.mdwn?op=file&rev=135
 Name: PyOpenCL
-Maintainer: Tomasz Rybak <bogomips@post.pl>
+Maintainer: Andreas Klöckner <lists@informa.tiker.net>
 Source: http://git.tiker.net/trees/pyopencl.git
  http://pypi.python.org/pypi/pyopencl/
  http://pypi.python.org/packages/source/p/pyopencl/
@@ -29,3 +29,26 @@
  FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
  OTHER DEALINGS IN THE SOFTWARE.
  
+Files: pyopencl/clrandom.py
+Copyright: Copyright (C) 1990, RSA Data Security, Inc.
+Licence:
+ Copyright (C) 1990, RSA Data Security, Inc. All rights reserved.
+ 
+ License to copy and use this software is granted provided that
+ it is identified as the "RSA Data Security, Inc. MD5 Message
+ Digest Algorithm" in all material mentioning or referencing this
+ software or this function.
+ 
+ License is also granted to make and use derivative works
+ provided that such works are identified as "derived from the RSA
+ Data Security, Inc. MD5 Message Digest Algorithm" in all
+ material mentioning or referencing the derived work.
+ 
+ RSA Data Security, Inc. makes no representations concerning
+ either the merchantability of this software or the suitability
+ of this software for any particular purpose.  It is provided "as
+ is" without express or implied warranty of any kind.
+ 
+ These notices must be retained in any copies of any part of this
+ documentation and/or software.
+
diff -Nru pyopencl-0.92~beta+git20100709/debian/patches/replace-setuptools.patch pyopencl-0.92/debian/patches/replace-setuptools.patch
--- pyopencl-0.92~beta+git20100709/debian/patches/replace-setuptools.patch	2010-07-10 18:32:03.000000000 +0200
+++ pyopencl-0.92/debian/patches/replace-setuptools.patch	2010-10-21 19:26:59.000000000 +0200
@@ -4,10 +4,10 @@
 Forwarded: not-needed
 Author: Tomasz Rybak <bogomips@post.pl>
 Last-Update: 2010-06-02
-Index: pyopencl-0,91.5+git20100531/MANIFEST.in
+Index: pyopencl-0.92~beta+git20100806/MANIFEST.in
 ===================================================================
---- pyopencl-0,91.5+git20100531.orig/MANIFEST.in	2010-06-02 17:30:59.000000000 +0200
-+++ pyopencl-0,91.5+git20100531/MANIFEST.in	2010-06-02 17:31:14.000000000 +0200
+--- pyopencl-0.92~beta+git20100806.orig/MANIFEST.in	2010-08-07 13:49:58.000000000 +0200
++++ pyopencl-0.92~beta+git20100806/MANIFEST.in	2010-08-07 15:13:27.000000000 +0200
 @@ -7,7 +7,6 @@
  include doc/*.py
  include doc/source/conf.py
@@ -16,15 +16,112 @@
  include configure.py
  include Makefile.in
  include aksetup_helper.py
-Index: pyopencl-0,91.5+git20100531/aksetup_helper.py
+Index: pyopencl-0.92~beta+git20100806/aksetup_helper.py
 ===================================================================
---- pyopencl-0,91.5+git20100531.orig/aksetup_helper.py	2010-06-02 17:30:55.000000000 +0200
-+++ pyopencl-0,91.5+git20100531/aksetup_helper.py	2010-06-02 17:31:06.000000000 +0200
-@@ -1,7 +1,3 @@
+--- pyopencl-0.92~beta+git20100806.orig/aksetup_helper.py	2010-08-07 13:49:58.000000000 +0200
++++ pyopencl-0.92~beta+git20100806/aksetup_helper.py	2010-08-07 15:13:52.000000000 +0200
+@@ -1,35 +1,6 @@
 -# dealings with ez_setup ------------------------------------------------------
 -import distribute_setup
 -distribute_setup.use_setuptools()
 -
+ import setuptools
  from setuptools import Extension
  
+-if not hasattr(setuptools, "_distribute"):
+-    print("-------------------------------------------------------------------------")
+-    print("Setuptools conflict detected.")
+-    print("-------------------------------------------------------------------------")
+-    print("When I imported setuptools, I did not get the distribute version of")
+-    print("setuptools, which is troubling--this package really wants to be used")
+-    print("with distribute rather than the old setuptools package. More than likely,")
+-    print("you have both distribute and setuptools installed, which is bad.")
+-    print("")
+-    print("See this page for more information:")
+-    print("http://wiki.tiker.net/DistributeVsSetuptools";)
+-    print("-------------------------------------------------------------------------")
+-    print("I will continue after a short while, fingers crossed.")
+-    print("-------------------------------------------------------------------------")
+-
+-    delay = 10
+-
+-    from time import sleep
+-    import sys
+-    while delay:
+-        sys.stdout.write("Continuing in %d seconds...   \r" % delay)
+-        sys.stdout.flush()
+-        delay -= 1
+-        sleep(1)
+-
  def setup(*args, **kwargs):
+     from setuptools import setup
+     import traceback
+@@ -134,33 +105,6 @@
+         from setup import get_config_schema
+         schema = get_config_schema()
+ 
+-    if (not schema.have_config() and not schema.have_global_config()
+-            and warn_about_no_config):
+-        print("*************************************************************")
+-        print("*** I have detected that you have not run configure.py.")
+-        print("*************************************************************")
+-        print("*** Additionally, no global config files were found.")
+-        print("*** I will go ahead with the default configuration.")
+-        print("*** In all likelihood, this will not work out.")
+-        print("*** ")
+-        print("*** See README_SETUP.txt for more information.")
+-        print("*** ")
+-        print("*** If the build does fail, just re-run configure.py with the")
+-        print("*** correct arguments, and then retry. Good luck!")
+-        print("*************************************************************")
+-        print("*** HIT Ctrl-C NOW IF THIS IS NOT WHAT YOU WANT")
+-        print("*************************************************************")
+-
+-        delay = 10
+-
+-        from time import sleep
+-        import sys
+-        while delay:
+-            sys.stdout.write("Continuing in %d seconds...   \r" % delay)
+-            sys.stdout.flush()
+-            delay -= 1
+-            sleep(1)
+-
+     return schema.read_config()
+ 
+ 
+@@ -517,34 +461,8 @@
+ 
+     if conf["USE_SHIPPED_BOOST"]:
+         if not exists("bpl-subset/bpl_subset/boost/version.hpp"):
+-            print("------------------------------------------------------------------------")
+-            print("The shipped Boost library was not found, but USE_SHIPPED_BOOST is True.")
+-            print("(The files should be under bpl-subset/.)")
+-            print("------------------------------------------------------------------------")
+-            print("If you got this package from git, you probably want to do")
+-            print("")
+-            print(" $ git submodule init")
+-            print(" $ git submodule update")
+-            print("")
+-            print("to fetch what you are presently missing. If you got this from")
+-            print("a distributed package on the net, that package is broken and")
+-            print("should be fixed. For now, I will turn off 'USE_SHIPPED_BOOST'")
+-            print("to try and see if the build succeeds that way, but in the long")
+-            print("run you might want to either get the missing bits or turn")
+-            print("'USE_SHIPPED_BOOST' off.")
+-            print("------------------------------------------------------------------------")
+             conf["USE_SHIPPED_BOOST"] = False
+ 
+-            delay = 10
+-
+-            from time import sleep
+-            import sys
+-            while delay:
+-                sys.stdout.write("Continuing in %d seconds...   \r" % delay)
+-                sys.stdout.flush()
+-                delay -= 1
+-                sleep(1)
+-
+     if conf["USE_SHIPPED_BOOST"]:
+         conf["BOOST_INC_DIR"] = ["bpl-subset/bpl_subset"]
+         conf["BOOST_LIB_DIR"] = []
diff -Nru pyopencl-0.92~beta+git20100709/debian/rules pyopencl-0.92/debian/rules
--- pyopencl-0.92~beta+git20100709/debian/rules	2010-07-10 18:32:03.000000000 +0200
+++ pyopencl-0.92/debian/rules	2010-10-21 19:26:59.000000000 +0200
@@ -8,11 +8,11 @@
 
 override_dh_auto_configure:
 	./configure.py --boost-python-libname=boost_python-py \
-		--cl-enable-gl
+		--cl-enable-gl --cl-inc-dir=/usr/include/nvidia-current
 
 override_dh_auto_build:
 	dh_auto_build --buildsystem=python_distutils
-	PYTHONPATH=../build/lib.$(DEB_BUILD_ARCH_OS)-$(DEB_BUILD_GNU_CPU)-$(firstword $(PYVERS))/ $(MAKE) -C doc html
+	PYTHONPATH=../$$(ls -d build/lib.*-*-$(firstword $(PYVERS)))/ $(MAKE) -C doc html
 
 override_dh_auto_install:
 	dh_auto_install --buildsystem=python_distutils --destdir=debian/$(PACKAGE_NAME)
@@ -28,14 +28,14 @@
 MODULE_NAME=pyopencl
 DEB_UPSTREAM_VERSION=$(shell dpkg-parsechangelog \
 	| sed -rne 's/^Version: ([^-]+).*/\1/p')
-GIT_REVISION=6fee76173b147a6c84364f8248649d52a9d6557f
+GIT_REVISION=dcd70ebc88337cc6474452b83ac02d4ae0c411a9
 GIT_URL=http://git.tiker.net/trees/pyopencl.git 
 
 get-orig-source:
 	rm -rf $(MODULE_NAME)-$(DEB_UPSTREAM_VERSION) $(MODULE_NAME)_$(DEB_UPSTREAM_VERSION).orig.tar.gz
 	git clone $(GIT_URL) $(MODULE_NAME)-$(DEB_UPSTREAM_VERSION)
 	cd $(MODULE_NAME)-$(DEB_UPSTREAM_VERSION) && git checkout $(GIT_REVISION)
-	rm -rf $(MODULE_NAME)-$(DEB_UPSTREAM_VERSION)/.git $(MODULE_NAME)-$(DEB_UPSTREAM_VERSION)/.gitignore
+	rm -rf $(MODULE_NAME)-$(DEB_UPSTREAM_VERSION)/.git $(MODULE_NAME)-$(DEB_UPSTREAM_VERSION)/.gitignore $(MODULE_NAME)-$(DEB_UPSTREAM_VERSION)/.gitmodules
 	tar czf $(MODULE_NAME)_$(DEB_UPSTREAM_VERSION).orig.tar.gz $(MODULE_NAME)-$(DEB_UPSTREAM_VERSION)
 	rm -rf $(MODULE_NAME)-$(DEB_UPSTREAM_VERSION)
 
diff -Nru pyopencl-0.92~beta+git20100709/distribute_setup.py pyopencl-0.92/distribute_setup.py
--- pyopencl-0.92~beta+git20100709/distribute_setup.py	2010-07-10 18:03:27.000000000 +0200
+++ pyopencl-0.92/distribute_setup.py	2010-10-21 19:10:19.000000000 +0200
@@ -46,19 +46,21 @@
             args = [quote(arg) for arg in args]
         return os.spawnl(os.P_WAIT, sys.executable, *args) == 0
 
-DEFAULT_VERSION = "0.6.4"
+DEFAULT_VERSION = "0.6.14"
 DEFAULT_URL = "http://pypi.python.org/packages/source/d/distribute/";
+SETUPTOOLS_FAKED_VERSION = "0.6c11"
+
 SETUPTOOLS_PKG_INFO = """\
 Metadata-Version: 1.0
 Name: setuptools
-Version: 0.6c9
+Version: %s
 Summary: xxxx
 Home-page: xxx
 Author: xxx
 Author-email: xxx
 License: xxx
 Description: xxx
-"""
+""" % SETUPTOOLS_FAKED_VERSION
 
 
 def _install(tarball):
@@ -79,12 +81,14 @@
 
         # installing
         log.warn('Installing Distribute')
-        assert _python_cmd('setup.py', 'install')
+        if not _python_cmd('setup.py', 'install'):
+            log.warn('Something went wrong during the installation.')
+            log.warn('See the error message above.')
     finally:
         os.chdir(old_wd)
 
 
-def _build_egg(tarball, to_dir):
+def _build_egg(egg, tarball, to_dir):
     # extracting the tarball
     tmpdir = tempfile.mkdtemp()
     log.warn('Extracting in %s', tmpdir)
@@ -104,27 +108,28 @@
         log.warn('Building a Distribute egg in %s', to_dir)
         _python_cmd('setup.py', '-q', 'bdist_egg', '--dist-dir', to_dir)
 
-        # returning the result
-        for file in os.listdir(to_dir):
-            if fnmatch.fnmatch(file, 'distribute-%s*.egg' % DEFAULT_VERSION):
-                return os.path.join(to_dir, file)
-
-        raise IOError('Could not build the egg.')
     finally:
         os.chdir(old_wd)
+    # returning the result
+    log.warn(egg)
+    if not os.path.exists(egg):
+        raise IOError('Could not build the egg.')
 
 
 def _do_download(version, download_base, to_dir, download_delay):
-    tarball = download_setuptools(version, download_base,
-                                  to_dir, download_delay)
-    egg = _build_egg(tarball, to_dir)
+    egg = os.path.join(to_dir, 'distribute-%s-py%d.%d.egg'
+                       % (version, sys.version_info[0], sys.version_info[1]))
+    if not os.path.exists(egg):
+        tarball = download_setuptools(version, download_base,
+                                      to_dir, download_delay)
+        _build_egg(egg, tarball, to_dir)
     sys.path.insert(0, egg)
     import setuptools
     setuptools.bootstrap_install_from = egg
 
 
 def use_setuptools(version=DEFAULT_VERSION, download_base=DEFAULT_URL,
-                   to_dir=os.curdir, download_delay=15):
+                   to_dir=os.curdir, download_delay=15, no_fake=True):
     # making sure we use the absolute path
     to_dir = os.path.abspath(to_dir)
     was_imported = 'pkg_resources' in sys.modules or \
@@ -133,21 +138,23 @@
         try:
             import pkg_resources
             if not hasattr(pkg_resources, '_distribute'):
-                fake_setuptools()
+                if not no_fake:
+                    _fake_setuptools()
                 raise ImportError
         except ImportError:
             return _do_download(version, download_base, to_dir, download_delay)
         try:
             pkg_resources.require("distribute>="+version)
             return
-        except pkg_resources.VersionConflict, e:
+        except pkg_resources.VersionConflict:
+            e = sys.exc_info()[1]
             if was_imported:
-                print >>sys.stderr, (
+                sys.stderr.write(
                 "The required version of distribute (>=%s) is not available,\n"
                 "and can't be installed while this script is running. Please\n"
                 "install a more recent version first, using\n"
                 "'easy_install -U distribute'."
-                "\n\n(Currently using %r)") % (version, e.args[0])
+                "\n\n(Currently using %r)\n" % (version, e.args[0]))
                 sys.exit(2)
             else:
                 del pkg_resources, sys.modules['pkg_resources']    # reload ok
@@ -157,7 +164,8 @@
             return _do_download(version, download_base, to_dir,
                                 download_delay)
     finally:
-        _create_fake_setuptools_pkg_info(to_dir)
+        if not no_fake:
+            _create_fake_setuptools_pkg_info(to_dir)
 
 def download_setuptools(version=DEFAULT_VERSION, download_base=DEFAULT_URL,
                         to_dir=os.curdir, delay=15):
@@ -171,7 +179,10 @@
     """
     # making sure we use the absolute path
     to_dir = os.path.abspath(to_dir)
-    import urllib2
+    try:
+        from urllib.request import urlopen
+    except ImportError:
+        from urllib2 import urlopen
     tgz_name = "distribute-%s.tar.gz" % version
     url = download_base + tgz_name
     saveto = os.path.join(to_dir, tgz_name)
@@ -179,7 +190,7 @@
     if not os.path.exists(saveto):  # Avoid repeated downloads
         try:
             log.warn("Downloading %s", url)
-            src = urllib2.urlopen(url)
+            src = urlopen(url)
             # Read/write all in one block, so we don't create a corrupt file
             # if the download is interrupted.
             data = src.read()
@@ -192,6 +203,29 @@
                 dst.close()
     return os.path.realpath(saveto)
 
+def _no_sandbox(function):
+    def __no_sandbox(*args, **kw):
+        try:
+            from setuptools.sandbox import DirectorySandbox
+            if not hasattr(DirectorySandbox, '_old'):
+                def violation(*args):
+                    pass
+                DirectorySandbox._old = DirectorySandbox._violation
+                DirectorySandbox._violation = violation
+                patched = True
+            else:
+                patched = False
+        except ImportError:
+            patched = False
+
+        try:
+            return function(*args, **kw)
+        finally:
+            if patched:
+                DirectorySandbox._violation = DirectorySandbox._old
+                del DirectorySandbox._old
+
+    return __no_sandbox
 
 def _patch_file(path, content):
     """Will backup the file then patch it"""
@@ -209,26 +243,17 @@
         f.close()
     return True
 
+_patch_file = _no_sandbox(_patch_file)
 
 def _same_content(path, content):
     return open(path).read() == content
 
-
 def _rename_path(path):
     new_name = path + '.OLD.%s' % time.time()
     log.warn('Renaming %s into %s', path, new_name)
-    try:
-        from setuptools.sandbox import DirectorySandbox
-        def _violation(*args):
-            pass
-        DirectorySandbox._violation = _violation
-    except ImportError:
-        pass
-
     os.rename(path, new_name)
     return new_name
 
-
 def _remove_flat_installation(placeholder):
     if not os.path.isdir(placeholder):
         log.warn('Unkown installation at %s', placeholder)
@@ -262,6 +287,7 @@
                      'Setuptools distribution', element)
     return True
 
+_remove_flat_installation = _no_sandbox(_remove_flat_installation)
 
 def _after_install(dist):
     log.warn('After install bootstrap.')
@@ -273,17 +299,20 @@
         log.warn('Could not find the install location')
         return
     pyver = '%s.%s' % (sys.version_info[0], sys.version_info[1])
-    setuptools_file = 'setuptools-0.6c9-py%s.egg-info' % pyver
+    setuptools_file = 'setuptools-%s-py%s.egg-info' % \
+            (SETUPTOOLS_FAKED_VERSION, pyver)
     pkg_info = os.path.join(placeholder, setuptools_file)
     if os.path.exists(pkg_info):
         log.warn('%s already exists', pkg_info)
         return
+
     log.warn('Creating %s', pkg_info)
     f = open(pkg_info, 'w')
     try:
         f.write(SETUPTOOLS_PKG_INFO)
     finally:
         f.close()
+
     pth_file = os.path.join(placeholder, 'setuptools.pth')
     log.warn('Creating %s', pth_file)
     f = open(pth_file, 'w')
@@ -292,6 +321,7 @@
     finally:
         f.close()
 
+_create_fake_setuptools_pkg_info = _no_sandbox(_create_fake_setuptools_pkg_info)
 
 def _patch_egg_dir(path):
     # let's check if it's already patched
@@ -311,10 +341,11 @@
         f.close()
     return True
 
+_patch_egg_dir = _no_sandbox(_patch_egg_dir)
 
 def _before_install():
     log.warn('Before install bootstrap.')
-    fake_setuptools()
+    _fake_setuptools()
 
 
 def _under_prefix(location):
@@ -330,12 +361,12 @@
                 if len(args) > index:
                     top_dir = args[index+1]
                     return location.startswith(top_dir)
-            elif option == '--user' and USER_SITE is not None:
-                return location.startswith(USER_SITE)
+        if arg == '--user' and USER_SITE is not None:
+            return location.startswith(USER_SITE)
     return True
 
 
-def fake_setuptools():
+def _fake_setuptools():
     log.warn('Scanning installed packages')
     try:
         import pkg_resources
@@ -344,7 +375,13 @@
         log.warn('Setuptools or Distribute does not seem to be installed.')
         return
     ws = pkg_resources.working_set
-    setuptools_dist = ws.find(pkg_resources.Requirement.parse('setuptools'))
+    try:
+        setuptools_dist = ws.find(pkg_resources.Requirement.parse('setuptools',
+                                  replacement=False))
+    except TypeError:
+        # old distribute API
+        setuptools_dist = ws.find(pkg_resources.Requirement.parse('setuptools'))
+
     if setuptools_dist is None:
         log.warn('No setuptools distribution found')
         return
@@ -384,6 +421,9 @@
 def _relaunch():
     log.warn('Relaunching...')
     # we have to relaunch the process
+    # pip marker to avoid a relaunch bug
+    if sys.argv[:3] == ['-c', 'install', '--single-version-externally-managed']:
+        sys.argv[0] = 'setup.py'
     args = [sys.executable] + sys.argv
     sys.exit(subprocess.call(args))
 
@@ -408,7 +448,7 @@
             # Extract directories with a safe mode.
             directories.append(tarinfo)
             tarinfo = copy.copy(tarinfo)
-            tarinfo.mode = 0700
+            tarinfo.mode = 448 # decimal for oct 0700
         self.extract(tarinfo, path)
 
     # Reverse sort directories.
@@ -427,7 +467,8 @@
             self.chown(tarinfo, dirpath)
             self.utime(tarinfo, dirpath)
             self.chmod(tarinfo, dirpath)
-        except ExtractError, e:
+        except ExtractError:
+            e = sys.exc_info()[1]
             if self.errorlevel > 1:
                 raise
             else:
diff -Nru pyopencl-0.92~beta+git20100709/doc/source/array.rst pyopencl-0.92/doc/source/array.rst
--- pyopencl-0.92~beta+git20100709/doc/source/array.rst	1970-01-01 01:00:00.000000000 +0100
+++ pyopencl-0.92/doc/source/array.rst	2010-10-21 19:10:19.000000000 +0200
@@ -0,0 +1,356 @@
+The :class:`Array` Class
+========================
+
+.. module:: pyopencl.array
+
+.. class:: DefaultAllocator(context, flags=pyopencl.mem_flags.READ_WRITE)
+
+    An Allocator that uses :class:`pyopencl.Buffer` with the given *flags*.
+
+    .. method:: __call__(self, size)
+
+.. class:: Array(context, shape, dtype, order="C", allocator=None, base=None, data=None, queue=None)
+
+    A :class:`numpy.ndarray` work-alike that stores its data and performs its
+    computations on the compute device.  *shape* and *dtype* work exactly as in
+    :mod:`numpy`.  Arithmetic methods in :class:`Array` support the
+    broadcasting of scalars. (e.g. `array+5`) If the
+
+    *allocator* is a callable that, upon being called with an argument of the number
+    of bytes to be allocated, returns an object that can be cast to an
+    :class:`int` representing the address of the newly allocated memory.
+    (See :class:`DefaultAllocator`.)
+
+    .. attribute :: data
+
+        The :class:`pyopencl.MemoryObject` instance created for the memory that backs
+        this :class:`Array`.
+
+    .. attribute :: shape
+
+        The tuple of lengths of each dimension in the array.
+
+    .. attribute :: dtype
+
+        The :class:`numpy.dtype` of the items in the GPU array.
+
+    .. attribute :: size
+
+        The number of meaningful entries in the array. Can also be computed by
+        multiplying up the numbers in :attr:`shape`.
+
+    .. attribute :: mem_size
+
+        The total number of entries, including padding, that are present in
+        the array.
+
+    .. attribute :: nbytes
+
+        The size of the entire array in bytes. Computed as :attr:`size` times
+        ``dtype.itemsize``.
+
+    .. method :: __len__()
+
+        Returns the size of the leading dimension of *self*.
+
+    .. method :: set(ary, queue=None, async=False)
+
+        Transfer the contents the :class:`numpy.ndarray` object *ary*
+        onto the device.
+
+        *ary* must have the same dtype and size (not necessarily shape) as *self*.
+
+
+    .. method :: get(queue=None, ary=None, async=False)
+
+        Transfer the contents of *self* into *ary* or a newly allocated
+        :mod:`numpy.ndarray`. If *ary* is given, it must have the right
+        size (not necessarily shape) and dtype.
+
+    .. method :: __str__()
+    .. method :: __repr__()
+
+    .. method :: mul_add(self, selffac, other, otherfac, queue=None):
+
+        Return `selffac*self + otherfac*other`.
+
+    .. method :: __add__(other)
+    .. method :: __sub__(other)
+    .. method :: __iadd__(other)
+    .. method :: __isub__(other)
+    .. method :: __neg__(other)
+    .. method :: __mul__(other)
+    .. method :: __div__(other)
+    .. method :: __rdiv__(other)
+    .. method :: __pow__(other)
+
+    .. method :: __abs__()
+
+        Return a :class:`Array` containing the absolute value of each
+        element of *self*.
+
+    .. UNDOC reverse()
+
+    .. method :: fill(scalar, queue=None)
+
+        Fill the array with *scalar*.
+
+    .. method :: astype(dtype, queue=None)
+
+        Return *self*, cast to *dtype*.
+
+Constructing :class:`Array` Instances
+----------------------------------------
+
+.. function:: to_device(context, queue, ary, allocator=None, async=False)
+
+    Return a :class:`Array` that is an exact copy of the :class:`numpy.ndarray`
+    instance *ary*.
+
+    See :class:`Array` for the meaning of *allocator*.
+
+.. function:: empty(context, shape, dtype, order="C", allocator=None, base=None, data=None, queue=None)
+
+    A synonym for the :class:`Array` constructor.
+
+.. function:: zeros(context, queue, shape, dtype, order="C", allocator=None)
+
+    Same as :func:`empty`, but the :class:`Array` is zero-initialized before
+    being returned.
+
+.. function:: empty_like(other_ary)
+
+    Make a new, uninitialized :class:`Array` having the same properties
+    as *other_ary*.
+
+.. function:: zeros_like(other_ary)
+
+    Make a new, zero-initialized :class:`Array` having the same properties
+    as *other_ary*.
+
+.. function:: arange(context, queue, start, stop, step, dtype=None)
+
+    Create a :class:`Array` filled with numbers spaced `step` apart,
+    starting from `start` and ending at `stop`.
+
+    For floating point arguments, the length of the result is
+    `ceil((stop - start)/step)`.  This rule may result in the last
+    element of the result being greater than `stop`.
+
+    *dtype*, if not specified, is taken as the largest common type
+    of *start*, *stop* and *step*.
+
+.. function:: take(a, indices, out=None, queue=None)
+
+    Return the :class:`Array` ``[a[indices[0]], ..., a[indices[n]]]``.
+    For the moment, *a* must be a type that can be bound to a texture.
+
+Conditionals
+^^^^^^^^^^^^
+
+.. function:: if_positive(criterion, then_, else_, out=None, queue=None)
+
+    Return an array like *then_*, which, for the element at index *i*,
+    contains *then_[i]* if *criterion[i]>0*, else *else_[i]*. (added in 0.94)
+
+.. function:: maximum(a, b, out=None, queue=None)
+
+    Return the elementwise maximum of *a* and *b*. (added in 0.94)
+
+.. function:: minimum(a, b, out=None, queue=None)
+
+    Return the elementwise minimum of *a* and *b*. (added in 0.94)
+
+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)
+.. TODO: atan2
+.. function:: atanh(array, queue=None)
+.. function:: atanpi(array, queue=None)
+.. TODO: 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
+-----------------------------------
+
+.. module:: pyopencl.clrandom
+
+.. function:: rand(context, queue, shape, dtype)
+
+    Return an array of `shape` filled with random values of `dtype`
+    in the range [0,1).
+
+Single-pass Custom Expression Evaluation
+----------------------------------------
+
+.. warning::
+
+    The following functionality is included in this documentation in the
+    hope that it may be useful, but its interface may change in future
+    revisions. Feedback is welcome.
+
+.. module:: pyopencl.elementwise
+
+Evaluating involved expressions on :class:`pyopencl.array.Array` instances 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.
+
+.. class:: ElementwiseKernel(context, arguments, operation, name="kernel", options=[])
+
+    Generate a kernel that takes a number of scalar or vector *arguments*
+    and performs the scalar *operation* on each entry of its arguments, if that
+    argument is a vector.
+
+    *arguments* is specified as a string formatted as a C argument list.
+    *operation* is specified as a C assignment statement, without a semicolon.
+    Vectors in *operation* should be indexed by the variable *i*.
+
+    *name* specifies the name as which the kernel is compiled, 
+    and *options* are passed unmodified to :meth:`pyopencl.Program.build`.
+
+    .. method:: __call__(*args)
+
+        Invoke the generated scalar kernel. The arguments may either be scalars or
+        :class:`GPUArray` instances.
+
+Here's a usage example::
+
+    import pyopencl as cl
+    import pyopencl.array as cl_array
+    import numpy
+
+    ctx = cl.create_some_context()
+    queue = cl.CommandQueue(ctx)
+
+    n = 10
+    a_gpu = cl_array.to_device(
+            ctx, queue, numpy.random.randn(n).astype(numpy.float32))
+    b_gpu = cl_array.to_device(
+            ctx, queue, numpy.random.randn(n).astype(numpy.float32))
+
+    from pyopencl.elementwise import ElementwiseKernel
+    lin_comb = ElementwiseKernel(ctx,
+            "float a, float *x, "
+            "float b, float *y, "
+            "float *z",
+            "z[i] = a*x[i] + b*y[i]",
+            "linear_combination")
+
+    c_gpu = cl_array.empty_like(a_gpu)
+    lin_comb(5, a_gpu, 6, b_gpu, c_gpu)
+
+    import numpy.linalg as la
+    assert la.norm((c_gpu - (5*a_gpu+6*b_gpu)).get()) < 1e-5
+
+(You can find this example as :file:`examples/demo_elementwise.py` in the PyOpenCL
+distribution.)
+
+Fast Fourier Transforms
+-----------------------
+
+Bogdan Opanchuk's `pyfft <http://pypi.python.org/pypi/pyfft>`_ package offers a
+variety of GPU-based FFT implementations.
+
diff -Nru pyopencl-0.92~beta+git20100709/doc/source/index.rst pyopencl-0.92/doc/source/index.rst
--- pyopencl-0.92~beta+git20100709/doc/source/index.rst	2010-07-10 18:03:27.000000000 +0200
+++ pyopencl-0.92/doc/source/index.rst	2010-10-21 19:10:19.000000000 +0200
@@ -70,6 +70,7 @@
     :maxdepth: 2
 
     runtime
+    array
     misc
 
 Note that this guide does not explain OpenCL programming and technology. Please 
diff -Nru pyopencl-0.92~beta+git20100709/doc/source/misc.rst pyopencl-0.92/doc/source/misc.rst
--- pyopencl-0.92~beta+git20100709/doc/source/misc.rst	2010-07-10 18:03:27.000000000 +0200
+++ pyopencl-0.92/doc/source/misc.rst	2010-10-21 19:10:19.000000000 +0200
@@ -82,9 +82,10 @@
   emphasize the importance of *loccal_size*.
 * Add :meth:`pyopencl.Kernel.set_scalar_arg_dtypes`.
 * Add support for the
-  `cl_nv_device_attribute_query <ghttp://www.khronos.org/registry/cl/extensions/khr/cl_nv_device_attribute_query.txt>`_
+  `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
 --------------
diff -Nru pyopencl-0.92~beta+git20100709/doc/source/runtime.rst pyopencl-0.92/doc/source/runtime.rst
--- pyopencl-0.92~beta+git20100709/doc/source/runtime.rst	2010-07-10 18:03:27.000000000 +0200
+++ pyopencl-0.92/doc/source/runtime.rst	2010-10-21 19:10:19.000000000 +0200
@@ -565,7 +565,7 @@
 
         See :class:`kernel_work_group_info` for values of *param*.
 
-    .. method:: set_arg(self, arg)
+    .. method:: set_arg(self, index, arg)
 
         *arg* may be
 
@@ -594,6 +594,8 @@
 
         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
diff -Nru pyopencl-0.92~beta+git20100709/examples/demo_mandelbrot.py pyopencl-0.92/examples/demo_mandelbrot.py
--- pyopencl-0.92~beta+git20100709/examples/demo_mandelbrot.py	2010-07-10 18:03:27.000000000 +0200
+++ pyopencl-0.92/examples/demo_mandelbrot.py	2010-10-21 19:10:19.000000000 +0200
@@ -27,41 +27,43 @@
 # Speed notes are listed in the same place
 
 # set width and height of window, more pixels take longer to calculate
-w = 400
-h = 400
+w = 256
+h = 256
 
 def calc_fractal_opencl(q, maxiter):
     ctx = cl.Context(cl.get_platforms()[0].get_devices())
     queue = cl.CommandQueue(ctx)
 
-    output = np.empty(q.shape, dtype=np.uint64)# resize(np.array(0,), q.shape)
+    output = np.empty(q.shape, dtype=np.uint16)
 
     mf = cl.mem_flags
     q_opencl = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=q)
     output_opencl = cl.Buffer(ctx, mf.WRITE_ONLY, output.nbytes)
 
     prg = cl.Program(ctx, """
+    #pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable
     __kernel void mandelbrot(__global float2 *q,
-                     __global long *output, long const maxiter)
+                     __global ushort *output, ushort const maxiter)
     {
         int gid = get_global_id(0);
         float nreal, real = 0;
         float imag = 0;
+
+        output[gid] = 0;
+
         for(int curiter = 0; curiter < maxiter; curiter++) {
-            nreal = real*real - imag*imag + q[gid][0];
-            imag = 2* real*imag + q[gid][1];
+            nreal = real*real - imag*imag + q[gid].x;
+            imag = 2* real*imag + q[gid].y;
             real = nreal;
 
-            if (real*real + imag*imag > 4.) {
+            if (real*real + imag*imag > 4.0f)
                  output[gid] = curiter;
-                 break;
-            }
         }
     }
     """).build()
 
-    prg.mandelbrot(queue, output.shape, None, q_opencl,
-            output_opencl, np.int32(maxiter))
+    prg.mandelbrot(queue, output.shape, (64,), q_opencl,
+            output_opencl, np.uint16(maxiter))
 
     cl.enqueue_read_buffer(queue, output_opencl, output).wait()
 
@@ -121,10 +123,10 @@
             self.root.mainloop()
 
 
-        def draw(self, x1, x2, y1, y2, maxiter=300):
+        def draw(self, x1, x2, y1, y2, maxiter=30):
             # draw the Mandelbrot set, from numpy example
-            xx = np.arange(x1, x2, (x2-x1)/w*2)
-            yy = np.arange(y2, y1, (y1-y2)/h*2) * 1j
+            xx = np.arange(x1, x2, (x2-x1)/w)
+            yy = np.arange(y2, y1, (y1-y2)/h) * 1j
             q = np.ravel(xx+yy[:, np.newaxis]).astype(np.complex64)
 
             start_main = time.time()
@@ -134,18 +136,20 @@
             secs = end_main - start_main
             print "Main took", secs
 
-            output = (output + (256*output) + (256**2)*output) * 8
-            # convert output to a string
-            self.mandel = output.tostring()
+            self.mandel = (output.reshape((h,w)) /
+                    float(output.max()) * 255.).astype(np.uint8)
 
         def create_image(self):
             """"
             create the image from the draw() string
             """
-            self.im = Image.new("RGB", (w/2, h/2))
             # you can experiment with these x and y ranges
             self.draw(-2.13, 0.77, -1.3, 1.3)
-            self.im.fromstring(self.mandel, "raw", "RGBX", 0, -1)
+            self.im = Image.fromarray(self.mandel)
+            self.im.putpalette(reduce(
+                lambda a,b: a+b, ((i,0,0) for i in range(255))
+            ))
+
 
         def create_label(self):
             # put the image on a label widget
diff -Nru pyopencl-0.92~beta+git20100709/examples/gl_interop_demo.py pyopencl-0.92/examples/gl_interop_demo.py
--- pyopencl-0.92~beta+git20100709/examples/gl_interop_demo.py	2010-07-10 18:03:27.000000000 +0200
+++ pyopencl-0.92/examples/gl_interop_demo.py	2010-10-21 19:10:19.000000000 +0200
@@ -1,75 +1,80 @@
-from OpenGL.GL import *
-from OpenGL.GLUT import *
-from OpenGL.raw.GL.VERSION.GL_1_5 import glBufferData as rawGlBufferData
-try:
-    from OpenGL.WGL import wglGetCurrentDisplay as GetCurrentDisplay, wglGetCurrentContext as GetCurrentContext
-except:
-    pass
-try:
-    from OpenGL.GLX import glXGetCurrentDisplay as GetCurrentDisplay, glXGetCurrentContext as GetCurrentContext
-except:
-    pass
-import pyopencl as cl
-
-
-n_vertices = 10000
-
-src = """
-
-__kernel void generate_sin(__global float2* a)
-{
-    int id = get_global_id(0);
-    int n = get_global_size(0);
-    float r = (float)id / (float)n;
-    float x = r * 16.0f * 3.1415f;
-    a[id].x = r * 2.0f - 1.0f;
-    a[id].y = native_sin(x);
-}
-
-"""
-
-def initialize():
-    plats = cl.get_platforms()
-    ctx_props = cl.context_properties
-    props = [(ctx_props.PLATFORM, plats[0]), (ctx_props.GL_CONTEXT_KHR,
-        GetCurrentContext()), (ctx_props.GLX_DISPLAY_KHR, GetCurrentDisplay())]
-    ctx = cl.Context(properties=props)
-    glClearColor(1, 1, 1, 1)
-    glColor(0, 0, 1)
-    vbo = glGenBuffers(1)
-    glBindBuffer(GL_ARRAY_BUFFER, vbo)
-    rawGlBufferData(GL_ARRAY_BUFFER, n_vertices * 2 * 4, None, GL_STATIC_DRAW)
-    glEnableClientState(GL_VERTEX_ARRAY)
-    glVertexPointer(2, GL_FLOAT, 0, None)
-    coords_dev = cl.GLBuffer(ctx, cl.mem_flags.READ_WRITE, int(vbo))
-    prog = cl.Program(ctx, src).build()
-    queue = cl.CommandQueue(ctx)
-    cl.enqueue_acquire_gl_objects(queue, [coords_dev])
-    prog.generate_sin(queue, (n_vertices,), None, coords_dev)
-    cl.enqueue_release_gl_objects(queue, [coords_dev])
-    queue.finish()
-    glFlush()
-
-def display():
-    glClear(GL_COLOR_BUFFER_BIT)
-    glDrawArrays(GL_LINE_STRIP, 0, n_vertices)
-    glFlush()
-
-def reshape(w, h):
-    glViewport(0, 0, w, h)
-    glMatrixMode(GL_PROJECTION)
-    glLoadIdentity()
-    glMatrixMode(GL_MODELVIEW)
-
-if __name__ == '__main__':
-    import sys
-    glutInit(sys.argv)
-    if len(sys.argv) > 1:
-        n_vertices = int(sys.argv[1])
-    glutInitWindowSize(800, 160)
-    glutInitWindowPosition(0, 0)
-    glutCreateWindow('OpenCL/OpenGL Interop Tutorial: Sin Generator')
-    glutDisplayFunc(display)
-    glutReshapeFunc(reshape)
-    initialize()
-    glutMainLoop()
+from OpenGL.GL import *
+from OpenGL.GLUT import *
+from OpenGL.raw.GL.VERSION.GL_1_5 import glBufferData as rawGlBufferData
+from OpenGL import platform, GLX, WGL
+import pyopencl as cl
+
+
+n_vertices = 10000
+
+src = """
+
+__kernel void generate_sin(__global float2* a)
+{
+    int id = get_global_id(0);
+    int n = get_global_size(0);
+    float r = (float)id / (float)n;
+    float x = r * 16.0f * 3.1415f;
+    a[id].x = r * 2.0f - 1.0f;
+    a[id].y = native_sin(x);
+}
+
+"""
+
+def initialize():
+    plats = cl.get_platforms()
+    ctx_props = cl.context_properties
+
+    props = [(ctx_props.PLATFORM, plats[0]), 
+            (ctx_props.GL_CONTEXT_KHR, platform.GetCurrentContext())]
+
+    import sys
+    if sys.platform == "linux2":
+        props.append(
+                (ctx_props.GLX_DISPLAY_KHR, 
+                    GLX.glXGetCurrentDisplay()))
+    elif sys.platform == "win32":
+        props.append(
+                (ctx_props.WGL_HDC_KHR, 
+                    WGL.wglGetCurrentDC()))
+    ctx = cl.Context(properties=props)
+
+    glClearColor(1, 1, 1, 1)
+    glColor(0, 0, 1)
+    vbo = glGenBuffers(1)
+    glBindBuffer(GL_ARRAY_BUFFER, vbo)
+    rawGlBufferData(GL_ARRAY_BUFFER, n_vertices * 2 * 4, None, GL_STATIC_DRAW)
+    glEnableClientState(GL_VERTEX_ARRAY)
+    glVertexPointer(2, GL_FLOAT, 0, None)
+    coords_dev = cl.GLBuffer(ctx, cl.mem_flags.READ_WRITE, int(vbo))
+    prog = cl.Program(ctx, src).build()
+    queue = cl.CommandQueue(ctx)
+    cl.enqueue_acquire_gl_objects(queue, [coords_dev])
+    prog.generate_sin(queue, (n_vertices,), None, coords_dev)
+    cl.enqueue_release_gl_objects(queue, [coords_dev])
+    queue.finish()
+    glFlush()
+
+def display():
+    glClear(GL_COLOR_BUFFER_BIT)
+    glDrawArrays(GL_LINE_STRIP, 0, n_vertices)
+    glFlush()
+
+def reshape(w, h):
+    glViewport(0, 0, w, h)
+    glMatrixMode(GL_PROJECTION)
+    glLoadIdentity()
+    glMatrixMode(GL_MODELVIEW)
+
+if __name__ == '__main__':
+    import sys
+    glutInit(sys.argv)
+    if len(sys.argv) > 1:
+        n_vertices = int(sys.argv[1])
+    glutInitWindowSize(800, 160)
+    glutInitWindowPosition(0, 0)
+    glutCreateWindow('OpenCL/OpenGL Interop Tutorial: Sin Generator')
+    glutDisplayFunc(display)
+    glutReshapeFunc(reshape)
+    initialize()
+    glutMainLoop()
diff -Nru pyopencl-0.92~beta+git20100709/.gitmodules pyopencl-0.92/.gitmodules
--- pyopencl-0.92~beta+git20100709/.gitmodules	2010-07-10 18:03:27.000000000 +0200
+++ pyopencl-0.92/.gitmodules	1970-01-01 01:00:00.000000000 +0100
@@ -1,3 +0,0 @@
-[submodule "bpl-subset"]
-	path = bpl-subset
-	url = http://git.tiker.net/trees/bpl-subset.git
diff -Nru pyopencl-0.92~beta+git20100709/MANIFEST.in pyopencl-0.92/MANIFEST.in
--- pyopencl-0.92~beta+git20100709/MANIFEST.in	2010-07-10 18:03:27.000000000 +0200
+++ pyopencl-0.92/MANIFEST.in	2010-10-21 19:10:19.000000000 +0200
@@ -13,4 +13,4 @@
 include aksetup_helper.py
 include README_SETUP.txt
 
-recursive-include bpl-subset
+recursive-include bpl-subset *.h *.hpp *.cpp *.html *.inl *.ipp *.pl *.txt
diff -Nru pyopencl-0.92~beta+git20100709/pyopencl/array.py pyopencl-0.92/pyopencl/array.py
--- pyopencl-0.92~beta+git20100709/pyopencl/array.py	2010-07-10 18:03:27.000000000 +0200
+++ pyopencl-0.92/pyopencl/array.py	2010-10-21 19:10:19.000000000 +0200
@@ -126,7 +126,7 @@
     work on an element-by-element basis, just like :class:`numpy.ndarray`.
     """
 
-    def __init__(self, context, shape, dtype, allocator=None,
+    def __init__(self, context, shape, dtype, order="C", allocator=None,
             base=None, data=None, queue=None):
         if allocator is None:
             allocator = DefaultAllocator(context)
@@ -147,6 +147,9 @@
 
         self.shape = shape
         self.dtype = numpy.dtype(dtype)
+        if order not in ["C", "F"]:
+            raise ValueError("order must be either 'C' or 'F'")
+        self.order = order
 
         self.mem_size = self.size = s
         self.nbytes = self.dtype.itemsize * self.size
@@ -180,7 +183,7 @@
 
     def get(self, queue=None, ary=None, async=False):
         if ary is None:
-            ary = numpy.empty(self.shape, self.dtype)
+            ary = numpy.empty(self.shape, self.dtype, order=self.order)
         else:
             if ary.size != self.size:
                 raise TypeError("'ary' has non-matching type")
@@ -344,12 +347,20 @@
         return result
 
     def __iadd__(self, other):
-        self._axpbyz(self, 1, self, 1, other)
-        return self
+        if isinstance(other, Array):
+            self._axpbyz(self, 1, self, 1, other)
+            return self
+        else:
+            self._axpbz(self, 1, self, other)
+            return self
 
     def __isub__(self, other):
-        self._axpbyz(self, 1, self, -1, other)
-        return self
+        if isinstance(other, Array):
+            self._axpbyz(self, 1, self, -1, other)
+            return self
+        else:
+            self._axpbz(self, 1, self, -other)
+            return self
 
     def __neg__(self):
         result = self._new_like_me()
@@ -470,30 +481,6 @@
         self._copy(result, self, queue=queue)
         return result
 
-    # slicing -----------------------------------------------------------------
-    def __getitem__(self, idx):
-        if idx == ():
-            return self
-
-        if len(self.shape) > 1:
-            raise NotImplementedError("multi-d slicing is not yet implemented")
-
-        if not isinstance(idx, slice):
-            raise ValueError("non-slice indexing not supported: %s" % (idx,))
-
-        l, = self.shape
-        start, stop, stride = idx.indices(l)
-
-        if stride != 1:
-            raise NotImplementedError("strided slicing is not yet implemented")
-
-        return Array(
-                shape=((stop-start)//stride,),
-                dtype=self.dtype,
-                allocator=self.allocator,
-                base=self,
-                data=int(self.data) + start*self.dtype.itemsize)
-
     # rich comparisons (or rather, lack thereof) ------------------------------
     def __eq__(self, other):
         raise NotImplementedError
@@ -517,7 +504,15 @@
 
 def to_device(context, queue, ary, allocator=None, async=False):
     """Converts a numpy array to a :class:`Array`."""
-    result = Array(context, ary.shape, ary.dtype, allocator,
+    if ary.flags.f_contiguous:
+        order = "F"
+    elif ary.flags.c_contiguous:
+        order = "C"
+    else:
+        raise ValueError("to_device only works on C- or Fortran-"
+                "contiguous arrays")
+
+    result = Array(context, ary.shape, ary.dtype, order, allocator,
             queue=queue)
     result.set(ary, async=async)
     return result
@@ -525,21 +520,22 @@
 
 empty = Array
 
-def zeros(context, queue, shape, dtype, allocator=None):
+def zeros(context, queue, shape, dtype, order="C", allocator=None):
     """Returns an array of the given shape and dtype filled with 0's."""
 
-    result = Array(context, shape, dtype, allocator, queue=queue)
+    result = Array(context, shape, dtype, 
+            order=order, allocator=allocator, queue=queue)
     result.fill(0)
     return result
 
 def empty_like(ary):
     result = Array(ary.context,
-            ary.shape, ary.dtype, ary.allocator, queue=ary.queue)
+            ary.shape, ary.dtype, allocator=ary.allocator, queue=ary.queue)
     return result
 
 def zeros_like(ary):
     result = Array(ary.context,
-            ary.shape, ary.dtype, ary.allocator, queue=ary.queue)
+            ary.shape, ary.dtype, allocator=ary.allocator, queue=ary.queue)
     result.fill(0)
     return result
 
@@ -641,7 +637,8 @@
 
 def take(a, indices, out=None, queue=None):
     if out is None:
-        out = Array(a.context, indices.shape, a.dtype, a.allocator,
+        out = Array(a.context, indices.shape, a.dtype, 
+                allocator=a.allocator,
                 queue=queue or a.queue)
 
     assert len(indices.shape) == 1
@@ -666,7 +663,8 @@
     vec_count = len(arrays)
 
     if out is None:
-        out = [Array(context, queue, indices.shape, a_dtype, a_allocator)
+        out = [Array(context, queue, indices.shape, a_dtype, 
+            allocator=a_allocator)
                 for i in range(vec_count)]
     else:
         if len(out) != len(arrays):
@@ -715,7 +713,8 @@
     vec_count = len(arrays)
 
     if out is None:
-        out = [Array(context, dest_shape, a_dtype, a_allocator, queue=queue)
+        out = [Array(context, dest_shape, a_dtype, 
+            allocator=a_allocator, queue=queue)
                 for i in range(vec_count)]
     else:
         if a_dtype != single_valued(o.dtype for o in out):
@@ -783,7 +782,7 @@
     vec_count = len(arrays)
 
     if out is None:
-        out = [Array(context, dest_shape, a_dtype, a_allocator, queue=queue)
+        out = [Array(context, dest_shape, a_dtype, allocator=a_allocator, queue=queue)
                 for i in range(vec_count)]
     else:
         if a_dtype != single_valued(o.dtype for o in out):
diff -Nru pyopencl-0.92~beta+git20100709/pyopencl/__init__.py pyopencl-0.92/pyopencl/__init__.py
--- pyopencl-0.92~beta+git20100709/pyopencl/__init__.py	2010-07-10 18:03:27.000000000 +0200
+++ pyopencl-0.92/pyopencl/__init__.py	2010-10-21 19:10:19.000000000 +0200
@@ -224,6 +224,10 @@
                 self.set_arg(i, arg)
         else:
             from struct import pack
+
+            if len(args) != len(arg_type_chars):
+                raise ValueError("length of argument type array and "
+                        "length of argument list do not agree")
             for i, (arg, arg_type_char) in enumerate(
                     zip(args, arg_type_chars)):
                 if arg_type_char:
diff -Nru pyopencl-0.92~beta+git20100709/pyopencl/version.py pyopencl-0.92/pyopencl/version.py
--- pyopencl-0.92~beta+git20100709/pyopencl/version.py	2010-07-10 18:03:27.000000000 +0200
+++ pyopencl-0.92/pyopencl/version.py	2010-10-21 19:10:19.000000000 +0200
@@ -1,5 +1,5 @@
 VERSION = (0, 92)
-VERSION_STATUS = "beta"
+VERSION_STATUS = ""
 VERSION_TEXT = ".".join(str(x) for x in VERSION) + VERSION_STATUS
 
 
diff -Nru pyopencl-0.92~beta+git20100709/src/wrapper/wrap_cl.hpp pyopencl-0.92/src/wrapper/wrap_cl.hpp
--- pyopencl-0.92~beta+git20100709/src/wrapper/wrap_cl.hpp	2010-07-10 18:03:27.000000000 +0200
+++ pyopencl-0.92/src/wrapper/wrap_cl.hpp	2010-10-21 19:10:19.000000000 +0200
@@ -20,6 +20,12 @@
 #include <CL/cl.h>
 // FIXME: Nvidia doesn't install cl_ext.h by default. Grr.
 // #include <CL/cl_ext.h>
+
+#ifdef _WIN32
+#define NOMINMAX
+#include <windows.h>
+#endif
+
 #ifdef HAVE_GL
 #include <GL/gl.h>
 #include <CL/cl_gl.h>
@@ -2575,7 +2581,8 @@
               std::vector<size_t> result;
               PYOPENCL_GET_VEC_INFO(KernelWorkGroup,
                   PYOPENCL_FIRST_ARG, param_name, result);
-              return py::list(result);
+
+              PYOPENCL_RETURN_VECTOR(size_t, result);
             }
           case CL_KERNEL_LOCAL_MEM_SIZE:
 #ifdef CL_VERSION_1_1
@@ -2875,8 +2882,8 @@
     std::vector<cl_context_properties> props
       = parse_context_properties(py_properties);
 
-    typedef CL_API_ENTRY cl_int CL_API_CALL
-      (*func_ptr_type)(const cl_context_properties * /* properties */,
+    typedef CL_API_ENTRY cl_int (CL_API_CALL
+      *func_ptr_type)(const cl_context_properties * /* properties */,
           cl_gl_context_info            /* param_name */,
           size_t                        /* param_value_size */,
           void *                        /* param_value */,
diff -Nru pyopencl-0.92~beta+git20100709/src/wrapper/wrap_constants.cpp pyopencl-0.92/src/wrapper/wrap_constants.cpp
--- pyopencl-0.92~beta+git20100709/src/wrapper/wrap_constants.cpp	2010-07-10 18:03:27.000000000 +0200
+++ pyopencl-0.92/src/wrapper/wrap_constants.cpp	2010-10-21 19:10:19.000000000 +0200
@@ -167,6 +167,7 @@
     ADD_ATTR(DEVICE_, QUEUE_PROPERTIES);
     ADD_ATTR(DEVICE_, NAME);
     ADD_ATTR(DEVICE_, VENDOR);
+    ADD_ATTR(, DRIVER_VERSION);
     ADD_ATTR(DEVICE_, VERSION);
     ADD_ATTR(DEVICE_, PROFILE);
     ADD_ATTR(DEVICE_, VERSION);
diff -Nru pyopencl-0.92~beta+git20100709/test/test_clmath.py pyopencl-0.92/test/test_clmath.py
--- pyopencl-0.92~beta+git20100709/test/test_clmath.py	1970-01-01 01:00:00.000000000 +0100
+++ pyopencl-0.92/test/test_clmath.py	2010-10-21 19:10:19.000000000 +0200
@@ -0,0 +1,182 @@
+from __future__ import division
+import math
+import numpy
+import pytools.test
+
+def have_cl():
+    try:
+        import pyopencl
+        return True
+    except:
+        return False
+
+if have_cl():
+    import pyopencl.array as cl_array
+    import pyopencl as cl
+    import pyopencl.clmath as clmath
+    from pyopencl.tools import pytest_generate_tests_for_pyopencl \
+            as pytest_generate_tests
+
+
+
+
+
+sizes = [10, 128, 1024, 1<<10, 1<<13]
+
+
+
+
+def has_double_support(dev):
+    for ext in dev.extensions.split(" "):
+        if ext == "cl_khr_fp64":
+            return True
+    return False
+
+
+
+
+numpy_func_names = {
+        "asin": "arcsin",
+        "acos": "arccos",
+        "atan": "arctan",
+        }
+
+
+
+
+def make_unary_function_test(name, (a, b)=(0, 1), threshold=0):
+    def test(ctx_getter):
+        context = ctx_getter()
+        queue = cl.CommandQueue(context)
+
+        gpu_func = getattr(clmath, name)
+        cpu_func = getattr(numpy, numpy_func_names.get(name, name))
+
+        if has_double_support(context.devices[0]):
+            dtypes = [numpy.float32, numpy.float64]
+        else:
+            dtypes = [numpy.float32]
+
+        for s in sizes:
+            for dtype in dtypes:
+                args = cl_array.arange(context, queue, a, b, (b-a)/s, 
+                        dtype=numpy.float32)
+                gpu_results = gpu_func(args).get()
+                cpu_results = cpu_func(args.get())
+
+                max_err = numpy.max(numpy.abs(cpu_results - gpu_results))
+                assert (max_err <= threshold).all(), \
+                        (max_err, name, dtype)
+
+    return pytools.test.mark_test.opencl(test)
+
+
+
+
+if have_cl():
+    test_ceil = make_unary_function_test("ceil", (-10, 10))
+    test_floor = make_unary_function_test("ceil", (-10, 10))
+    test_fabs = make_unary_function_test("fabs", (-10, 10))
+    test_exp = make_unary_function_test("exp", (-3, 3), 1e-5)
+    test_log = make_unary_function_test("log", (1e-5, 1), 1e-6)
+    test_log10 = make_unary_function_test("log10", (1e-5, 1), 5e-7)
+    test_sqrt = make_unary_function_test("sqrt", (1e-5, 1), 2e-7)
+
+    test_sin = make_unary_function_test("sin", (-10, 10), 1e-7)
+    test_cos = make_unary_function_test("cos", (-10, 10), 1e-7)
+    test_asin = make_unary_function_test("asin", (-0.9, 0.9), 5e-7)
+    test_acos = make_unary_function_test("acos", (-0.9, 0.9), 5e-7)
+    test_tan = make_unary_function_test("tan", 
+            (-math.pi/2 + 0.1, math.pi/2 - 0.1), 1e-5)
+    test_atan = make_unary_function_test("atan", (-10, 10), 2e-7)
+
+    test_sinh = make_unary_function_test("sinh", (-3, 3), 1e-6)
+    test_cosh = make_unary_function_test("cosh", (-3, 3), 1e-6)
+    test_tanh = make_unary_function_test("tanh", (-3, 3), 2e-6)
+
+
+
+
+@pytools.test.mark_test.opencl
+def test_fmod(ctx_getter):
+    context = ctx_getter()
+    queue = cl.CommandQueue(context)
+
+    for s in sizes:
+        a = cl_array.arange(context, queue, s, dtype=numpy.float32)/10
+        a2 = cl_array.arange(context, queue, s, dtype=numpy.float32)/45.2 + 0.1
+        b = clmath.fmod(a, a2)
+
+        a = a.get()
+        a2 = a2.get()
+        b = b.get()
+
+        for i in range(s):
+            assert math.fmod(a[i], a2[i]) == b[i]
+
+@pytools.test.mark_test.opencl
+def test_ldexp(ctx_getter):
+    context = ctx_getter()
+    queue = cl.CommandQueue(context)
+
+    for s in sizes:
+        a = cl_array.arange(context, queue, s, dtype=numpy.float32)
+        a2 = cl_array.arange(context, queue, s, dtype=numpy.float32)*1e-3
+        b = clmath.ldexp(a,a2)
+
+        a = a.get()
+        a2 = a2.get()
+        b = b.get()
+
+        for i in range(s):
+            assert math.ldexp(a[i], int(a2[i])) == b[i]
+
+@pytools.test.mark_test.opencl
+def test_modf(ctx_getter):
+    context = ctx_getter()
+    queue = cl.CommandQueue(context)
+
+    for s in sizes:
+        a = cl_array.arange(context, queue, s, dtype=numpy.float32)/10
+        fracpart, intpart = clmath.modf(a)
+
+        a = a.get()
+        intpart = intpart.get()
+        fracpart = fracpart.get()
+
+        for i in range(s):
+            fracpart_true, intpart_true = math.modf(a[i])
+
+            assert intpart_true == intpart[i]
+            assert abs(fracpart_true - fracpart[i]) < 1e-4
+
+@pytools.test.mark_test.opencl
+def test_frexp(ctx_getter):
+    context = ctx_getter()
+    queue = cl.CommandQueue(context)
+
+    for s in sizes:
+        a = cl_array.arange(context, queue, s, dtype=numpy.float32)/10
+        significands, exponents = clmath.frexp(a)
+
+        a = a.get()
+        significands = significands.get()
+        exponents = exponents.get()
+
+        for i in range(s):
+            sig_true, ex_true = math.frexp(a[i])
+
+            assert sig_true == significands[i]
+            assert ex_true == exponents[i]
+
+
+
+
+if __name__ == "__main__":
+    # make sure that import failures get reported, instead of skipping the tests.
+    import sys
+    if len(sys.argv) > 1:
+        exec sys.argv[1]
+    else:
+        from py.test.cmdline import main
+        main([__file__])
diff -Nru pyopencl-0.92~beta+git20100709/test/test_math.py pyopencl-0.92/test/test_math.py
--- pyopencl-0.92~beta+git20100709/test/test_math.py	2010-07-10 18:03:27.000000000 +0200
+++ pyopencl-0.92/test/test_math.py	1970-01-01 01:00:00.000000000 +0100
@@ -1,182 +0,0 @@
-from __future__ import division
-import math
-import numpy
-import pytools.test
-
-def have_cl():
-    try:
-        import pyopencl
-        return True
-    except:
-        return False
-
-if have_cl():
-    import pyopencl.array as cl_array
-    import pyopencl as cl
-    import pyopencl.clmath as clmath
-    from pyopencl.tools import pytest_generate_tests_for_pyopencl \
-            as pytest_generate_tests
-
-
-
-
-
-sizes = [10, 128, 1024, 1<<10, 1<<13]
-
-
-
-
-def has_double_support(dev):
-    for ext in dev.extensions.split(" "):
-        if ext == "cl_khr_fp64":
-            return True
-    return False
-
-
-
-
-numpy_func_names = {
-        "asin": "arcsin",
-        "acos": "arccos",
-        "atan": "arctan",
-        }
-
-
-
-
-def make_unary_function_test(name, (a, b)=(0, 1), threshold=0):
-    def test(ctx_getter):
-        context = ctx_getter()
-        queue = cl.CommandQueue(context)
-
-        gpu_func = getattr(clmath, name)
-        cpu_func = getattr(numpy, numpy_func_names.get(name, name))
-
-        if has_double_support(context.devices[0]):
-            dtypes = [numpy.float32, numpy.float64]
-        else:
-            dtypes = [numpy.float32]
-
-        for s in sizes:
-            for dtype in dtypes:
-                args = cl_array.arange(context, queue, a, b, (b-a)/s, 
-                        dtype=numpy.float32)
-                gpu_results = gpu_func(args).get()
-                cpu_results = cpu_func(args.get())
-
-                max_err = numpy.max(numpy.abs(cpu_results - gpu_results))
-                assert (max_err <= threshold).all(), \
-                        (max_err, name, dtype)
-
-    return pytools.test.mark_test.opencl(test)
-
-
-
-
-if have_cl():
-    test_ceil = make_unary_function_test("ceil", (-10, 10))
-    test_floor = make_unary_function_test("ceil", (-10, 10))
-    test_fabs = make_unary_function_test("fabs", (-10, 10))
-    test_exp = make_unary_function_test("exp", (-3, 3), 1e-5)
-    test_log = make_unary_function_test("log", (1e-5, 1), 5e-7)
-    test_log10 = make_unary_function_test("log10", (1e-5, 1), 3e-7)
-    test_sqrt = make_unary_function_test("sqrt", (1e-5, 1), 2e-7)
-
-    test_sin = make_unary_function_test("sin", (-10, 10), 1e-7)
-    test_cos = make_unary_function_test("cos", (-10, 10), 1e-7)
-    test_asin = make_unary_function_test("asin", (-0.9, 0.9), 5e-7)
-    test_acos = make_unary_function_test("acos", (-0.9, 0.9), 5e-7)
-    test_tan = make_unary_function_test("tan", 
-            (-math.pi/2 + 0.1, math.pi/2 - 0.1), 1e-5)
-    test_atan = make_unary_function_test("atan", (-10, 10), 2e-7)
-
-    test_sinh = make_unary_function_test("sinh", (-3, 3), 1e-6)
-    test_cosh = make_unary_function_test("cosh", (-3, 3), 1e-6)
-    test_tanh = make_unary_function_test("tanh", (-3, 3), 2e-6)
-
-
-
-
-@pytools.test.mark_test.opencl
-def test_fmod(ctx_getter):
-    context = ctx_getter()
-    queue = cl.CommandQueue(context)
-
-    for s in sizes:
-        a = cl_array.arange(context, queue, s, dtype=numpy.float32)/10
-        a2 = cl_array.arange(context, queue, s, dtype=numpy.float32)/45.2 + 0.1
-        b = clmath.fmod(a, a2)
-
-        a = a.get()
-        a2 = a2.get()
-        b = b.get()
-
-        for i in range(s):
-            assert math.fmod(a[i], a2[i]) == b[i]
-
-@pytools.test.mark_test.opencl
-def test_ldexp(ctx_getter):
-    context = ctx_getter()
-    queue = cl.CommandQueue(context)
-
-    for s in sizes:
-        a = cl_array.arange(context, queue, s, dtype=numpy.float32)
-        a2 = cl_array.arange(context, queue, s, dtype=numpy.float32)*1e-3
-        b = clmath.ldexp(a,a2)
-
-        a = a.get()
-        a2 = a2.get()
-        b = b.get()
-
-        for i in range(s):
-            assert math.ldexp(a[i], int(a2[i])) == b[i]
-
-@pytools.test.mark_test.opencl
-def test_modf(ctx_getter):
-    context = ctx_getter()
-    queue = cl.CommandQueue(context)
-
-    for s in sizes:
-        a = cl_array.arange(context, queue, s, dtype=numpy.float32)/10
-        fracpart, intpart = clmath.modf(a)
-
-        a = a.get()
-        intpart = intpart.get()
-        fracpart = fracpart.get()
-
-        for i in range(s):
-            fracpart_true, intpart_true = math.modf(a[i])
-
-            assert intpart_true == intpart[i]
-            assert abs(fracpart_true - fracpart[i]) < 1e-4
-
-@pytools.test.mark_test.opencl
-def test_frexp(ctx_getter):
-    context = ctx_getter()
-    queue = cl.CommandQueue(context)
-
-    for s in sizes:
-        a = cl_array.arange(context, queue, s, dtype=numpy.float32)/10
-        significands, exponents = clmath.frexp(a)
-
-        a = a.get()
-        significands = significands.get()
-        exponents = exponents.get()
-
-        for i in range(s):
-            sig_true, ex_true = math.frexp(a[i])
-
-            assert sig_true == significands[i]
-            assert ex_true == exponents[i]
-
-
-
-
-if __name__ == "__main__":
-    # make sure that import failures get reported, instead of skipping the tests.
-    import sys
-    if len(sys.argv) > 1:
-        exec sys.argv[1]
-    else:
-        from py.test.cmdline import main
-        main([__file__])
diff --git a/.gitignore b/.gitignore
index 8a52c24..2ed0d45 100644
--- a/.gitignore
+++ b/.gitignore
@@ -44,3 +44,4 @@ MANIFEST
 tmp
 temp*
 setuptools.pth
+distribute-*.tar.gz
diff --git a/MANIFEST.in b/MANIFEST.in
index fdfc08c..1426ac2 100644
--- a/MANIFEST.in
+++ b/MANIFEST.in
@@ -13,4 +13,4 @@ include Makefile.in
 include aksetup_helper.py
 include README_SETUP.txt
 
-recursive-include bpl-subset
+recursive-include bpl-subset *.h *.hpp *.cpp *.html *.inl *.ipp *.pl *.txt
diff --git a/aksetup_helper.py b/aksetup_helper.py
index 5637d68..252fd3f 100644
--- a/aksetup_helper.py
+++ b/aksetup_helper.py
@@ -2,8 +2,34 @@
 import distribute_setup
 distribute_setup.use_setuptools()
 
+import setuptools
 from setuptools import Extension
 
+if not hasattr(setuptools, "_distribute"):
+    print("-------------------------------------------------------------------------")
+    print("Setuptools conflict detected.")
+    print("-------------------------------------------------------------------------")
+    print("When I imported setuptools, I did not get the distribute version of")
+    print("setuptools, which is troubling--this package really wants to be used")
+    print("with distribute rather than the old setuptools package. More than likely,")
+    print("you have both distribute and setuptools installed, which is bad.")
+    print("")
+    print("See this page for more information:")
+    print("http://wiki.tiker.net/DistributeVsSetuptools";)
+    print("-------------------------------------------------------------------------")
+    print("I will continue after a short while, fingers crossed.")
+    print("-------------------------------------------------------------------------")
+
+    delay = 10
+
+    from time import sleep
+    import sys
+    while delay:
+        sys.stdout.write("Continuing in %d seconds...   \r" % delay)
+        sys.stdout.flush()
+        delay -= 1
+        sleep(1)
+
 def setup(*args, **kwargs):
     from setuptools import setup
     import traceback
@@ -14,16 +40,16 @@ def setup(*args, **kwargs):
     except SystemExit:
         raise
     except:
-        print "----------------------------------------------------------------------------"
-        print "Sorry, your build failed. Try rerunning configure.py with different options."
-        print "----------------------------------------------------------------------------"
+        print ("----------------------------------------------------------------------------")
+        print ("Sorry, your build failed. Try rerunning configure.py with different options.")
+        print ("----------------------------------------------------------------------------")
         raise
 
 
 
 
 class NumpyExtension(Extension):
-    # nicked from 
+    # nicked from
     # http://mail.python.org/pipermail/distutils-sig/2007-September/008253.html
     # solution by Michael Hoffmann
     def __init__(self, *args, **kwargs):
@@ -83,7 +109,7 @@ class HedgeExtension(PyUblasExtension):
 
 # tools -----------------------------------------------------------------------
 def flatten(list):
-    """For an iterable of sub-iterables, generate each member of each 
+    """For an iterable of sub-iterables, generate each member of each
     sub-iterable in turn, i.e. a flattened version of that super-iterable.
 
     Example: Turn [[a,b,c],[d,e,f]] into [a,b,c,d,e,f].
@@ -110,20 +136,20 @@ def get_config(schema=None, warn_about_no_config=True):
 
     if (not schema.have_config() and not schema.have_global_config()
             and warn_about_no_config):
-        print "*************************************************************"
-        print "*** I have detected that you have not run configure.py."
-        print "*************************************************************"
-        print "*** Additionally, no global config files were found."
-        print "*** I will go ahead with the default configuration."
-        print "*** In all likelihood, this will not work out."
-        print "*** "
-        print "*** See README_SETUP.txt for more information."
-        print "*** "
-        print "*** If the build does fail, just re-run configure.py with the"
-        print "*** correct arguments, and then retry. Good luck!"
-        print "*************************************************************"
-        print "*** HIT Ctrl-C NOW IF THIS IS NOT WHAT YOU WANT"
-        print "*************************************************************"
+        print("*************************************************************")
+        print("*** I have detected that you have not run configure.py.")
+        print("*************************************************************")
+        print("*** Additionally, no global config files were found.")
+        print("*** I will go ahead with the default configuration.")
+        print("*** In all likelihood, this will not work out.")
+        print("*** ")
+        print("*** See README_SETUP.txt for more information.")
+        print("*** ")
+        print("*** If the build does fail, just re-run configure.py with the")
+        print("*** correct arguments, and then retry. Good luck!")
+        print("*************************************************************")
+        print("*** HIT Ctrl-C NOW IF THIS IS NOT WHAT YOU WANT")
+        print("*************************************************************")
 
         delay = 10
 
@@ -204,7 +230,7 @@ def expand_str(s, options):
     return re.subn(r"\$\{([a-zA-Z0-9_]+)\}", my_repl, s)[0]
 
 def expand_value(v, options):
-    if isinstance(v, (str, unicode)):
+    if isinstance(v, str):
         return expand_str(v, options)
     elif isinstance(v, list):
         return [expand_value(i, options) for i in v]
@@ -246,15 +272,15 @@ class ConfigSchema:
         self.conf_dir = conf_dir
 
     def get_default_config(self):
-        return dict((opt.name, opt.default) 
+        return dict((opt.name, opt.default)
                 for opt in self.options)
-        
+
     def read_config_from_pyfile(self, filename):
         result = {}
         filevars = {}
-        execfile(filename, filevars)
+        exec(compile(open(filename, "r").read(), filename, "exec"), filevars)
 
-        for key, value in filevars.iteritems():
+        for key, value in filevars.items():
             if key in self.optdict:
                 result[key] = value
 
@@ -265,13 +291,13 @@ class ConfigSchema:
         filevars = {}
 
         try:
-            execfile(filename, filevars)
+            exec(compile(open(filename, "r").read(), filename, "exec"), filevars)
         except IOError:
             pass
 
         del filevars["__builtins__"]
 
-        for key, value in config.iteritems():
+        for key, value in config.items():
             if value is not None:
                 filevars[key] = value
 
@@ -296,7 +322,7 @@ class ConfigSchema:
         result = self.get_default_config()
 
         import os
-        
+
         confignames = []
         if self.global_conf_file is not None:
             confignames.append(self.global_conf_file)
@@ -328,16 +354,16 @@ class ConfigSchema:
         result = self.get_default_config_with_files()
         if os.access(cfile, os.R_OK):
             filevars = {}
-            execfile(cfile, filevars)
+            exec(compile(open(cfile, "r").read(), cfile, "exec"), filevars)
 
-            for key, value in filevars.iteritems():
+            for key, value in filevars.items():
                 if key in self.optdict:
                     result[key] = value
                 elif key == "__builtins__":
                     pass
                 else:
-                    raise KeyError, "invalid config key in %s: %s" % (
-                            cfile, key)
+                    raise KeyError("invalid config key in %s: %s" % (
+                            cfile, key))
 
         expand_options(result)
 
@@ -399,13 +425,13 @@ class Option(object):
         return result
 
     def value_to_str(self, default):
-        return default 
+        return default
 
     def add_to_configparser(self, parser, default=None):
         default = default_or(default, self.default)
         default_str = self.value_to_str(default)
         parser.add_option(
-            "--" + self.as_option(), dest=self.name, 
+            "--" + self.as_option(), dest=self.name,
             default=default_str,
             metavar=self.metavar(), help=self.get_help(default))
 
@@ -417,7 +443,7 @@ class Switch(Option):
         option = self.as_option()
 
         if not isinstance(self.default, bool):
-            raise ValueError, "Switch options must have a default"
+            raise ValueError("Switch options must have a default")
 
         if default is None:
             default = self.default
@@ -426,11 +452,11 @@ class Switch(Option):
             action = "store_false"
         else:
             action = "store_true"
-            
+
         parser.add_option(
-            "--" + self.as_option(), 
-            dest=self.name, 
-            help=self.get_help(default), 
+            "--" + self.as_option(),
+            dest=self.name,
+            help=self.get_help(default),
             default=default,
             action=action)
 
@@ -458,7 +484,7 @@ class StringListOption(Option):
 class IncludeDir(StringListOption):
     def __init__(self, lib_name, default=None, human_name=None, help=None):
         StringListOption.__init__(self, "%s_INC_DIR" % lib_name, default,
-                help=help or ("Include directories for %s" 
+                help=help or ("Include directories for %s"
                 % (human_name or humanize(lib_name))))
 
 class LibraryDir(StringListOption):
@@ -470,14 +496,14 @@ class LibraryDir(StringListOption):
 class Libraries(StringListOption):
     def __init__(self, lib_name, default=None, human_name=None, help=None):
         StringListOption.__init__(self, "%s_LIBNAME" % lib_name, default,
-                help=help or ("Library names for %s (without lib or .so)" 
+                help=help or ("Library names for %s (without lib or .so)"
                 % (human_name or humanize(lib_name))))
 
 class BoostLibraries(Libraries):
     def __init__(self, lib_base_name):
-        Libraries.__init__(self, "BOOST_%s" % lib_base_name.upper(), 
+        Libraries.__init__(self, "BOOST_%s" % lib_base_name.upper(),
                 ["boost_%s-${BOOST_COMPILER}-mt" % lib_base_name],
-                help="Library names for Boost C++ %s library (without lib or .so)" 
+                help="Library names for Boost C++ %s library (without lib or .so)"
                     % humanize(lib_base_name))
 
 def set_up_shipped_boost_if_requested(conf):
@@ -491,22 +517,22 @@ def set_up_shipped_boost_if_requested(conf):
 
     if conf["USE_SHIPPED_BOOST"]:
         if not exists("bpl-subset/bpl_subset/boost/version.hpp"):
-            print >>sys.stderr, "------------------------------------------------------------------------"
-            print >>sys.stderr, "The shipped Boost library was not found, but USE_SHIPPED_BOOST is True."
-            print >>sys.stderr, "(The files should be under bpl-subset/.)"
-            print >>sys.stderr, "------------------------------------------------------------------------"
-            print >>sys.stderr, "If you got this package from git, you probably want to do"
-            print >>sys.stderr, ""
-            print >>sys.stderr, " $ git submodule init"
-            print >>sys.stderr, " $ git submodule update"
-            print >>sys.stderr, ""
-            print >>sys.stderr, "to fetch what you are presently missing. If you got this from"
-            print >>sys.stderr, "a distributed package on the net, that package is broken and"
-            print >>sys.stderr, "should be fixed. For now, I will turn off 'USE_SHIPPED_BOOST'"
-            print >>sys.stderr, "to try and see if the build succeeds that way, but in the long"
-            print >>sys.stderr, "run you might want to either get the missing bits or turn"
-            print >>sys.stderr, "'USE_SHIPPED_BOOST' off."
-            print >>sys.stderr, "------------------------------------------------------------------------"
+            print("------------------------------------------------------------------------")
+            print("The shipped Boost library was not found, but USE_SHIPPED_BOOST is True.")
+            print("(The files should be under bpl-subset/.)")
+            print("------------------------------------------------------------------------")
+            print("If you got this package from git, you probably want to do")
+            print("")
+            print(" $ git submodule init")
+            print(" $ git submodule update")
+            print("")
+            print("to fetch what you are presently missing. If you got this from")
+            print("a distributed package on the net, that package is broken and")
+            print("should be fixed. For now, I will turn off 'USE_SHIPPED_BOOST'")
+            print("to try and see if the build succeeds that way, but in the long")
+            print("run you might want to either get the missing bits or turn")
+            print("'USE_SHIPPED_BOOST' off.")
+            print("------------------------------------------------------------------------")
             conf["USE_SHIPPED_BOOST"] = False
 
             delay = 10
@@ -541,7 +567,7 @@ def set_up_shipped_boost_if_requested(conf):
             source_files += glob(
                     "bpl-subset/bpl_subset/libs/thread/src/pthread/*.cpp")
 
-        return (source_files, 
+        return (source_files,
                 {"BOOST_MULTI_INDEX_DISABLE_SERIALIZATION": 1}
                 )
     else:
@@ -552,7 +578,7 @@ def make_boost_base_options():
     return [
         IncludeDir("BOOST", []),
         LibraryDir("BOOST", []),
-        Option("BOOST_COMPILER", default="gcc43", 
+        Option("BOOST_COMPILER", default="gcc43",
             help="The compiler with which Boost C++ was compiled, e.g. gcc43"),
         ]
 
@@ -568,29 +594,29 @@ def configure_frontend():
     from setup import get_config_schema
     schema = get_config_schema()
     if schema.have_config():
-        print "************************************************************"
-        print "*** I have detected that you have already run configure."
-        print "*** I'm taking the configured values as defaults for this"
-        print "*** configure run. If you don't want this, delete the file"
-        print "*** %s." % schema.get_conf_file()
-        print "************************************************************"
+        print("************************************************************")
+        print("*** I have detected that you have already run configure.")
+        print("*** I'm taking the configured values as defaults for this")
+        print("*** configure run. If you don't want this, delete the file")
+        print("*** %s." % schema.get_conf_file())
+        print("************************************************************")
 
     import sys
 
     description = "generate a configuration file for this software package"
     parser = OptionParser(description=description)
     parser.add_option(
-	    "--python-exe", dest="python_exe", default=sys.executable,
-	    help="Which Python interpreter to use", metavar="PATH")
+            "--python-exe", dest="python_exe", default=sys.executable,
+            help="Which Python interpreter to use", metavar="PATH")
 
     parser.add_option("--prefix", default=None,
-	    help="Ignored")
+            help="Ignored")
     parser.add_option("--enable-shared", help="Ignored", action="store_false")
     parser.add_option("--disable-static", help="Ignored", action="store_false")
-    parser.add_option("--update-user", help="Update user config file (%s)" % schema.user_conf_file, 
+    parser.add_option("--update-user", help="Update user config file (%s)" % schema.user_conf_file,
             action="store_true")
-    parser.add_option("--update-global", 
-            help="Update global config file (%s)" % schema.global_conf_file, 
+    parser.add_option("--update-global",
+            help="Update global config file (%s)" % schema.global_conf_file,
             action="store_true")
 
     schema.add_to_configparser(parser, schema.read_config())
diff --git a/distribute_setup.py b/distribute_setup.py
index 59424cc..3ea2e66 100644
--- a/distribute_setup.py
+++ b/distribute_setup.py
@@ -46,19 +46,21 @@ except ImportError:
             args = [quote(arg) for arg in args]
         return os.spawnl(os.P_WAIT, sys.executable, *args) == 0
 
-DEFAULT_VERSION = "0.6.4"
+DEFAULT_VERSION = "0.6.14"
 DEFAULT_URL = "http://pypi.python.org/packages/source/d/distribute/";
+SETUPTOOLS_FAKED_VERSION = "0.6c11"
+
 SETUPTOOLS_PKG_INFO = """\
 Metadata-Version: 1.0
 Name: setuptools
-Version: 0.6c9
+Version: %s
 Summary: xxxx
 Home-page: xxx
 Author: xxx
 Author-email: xxx
 License: xxx
 Description: xxx
-"""
+""" % SETUPTOOLS_FAKED_VERSION
 
 
 def _install(tarball):
@@ -79,12 +81,14 @@ def _install(tarball):
 
         # installing
         log.warn('Installing Distribute')
-        assert _python_cmd('setup.py', 'install')
+        if not _python_cmd('setup.py', 'install'):
+            log.warn('Something went wrong during the installation.')
+            log.warn('See the error message above.')
     finally:
         os.chdir(old_wd)
 
 
-def _build_egg(tarball, to_dir):
+def _build_egg(egg, tarball, to_dir):
     # extracting the tarball
     tmpdir = tempfile.mkdtemp()
     log.warn('Extracting in %s', tmpdir)
@@ -104,27 +108,28 @@ def _build_egg(tarball, to_dir):
         log.warn('Building a Distribute egg in %s', to_dir)
         _python_cmd('setup.py', '-q', 'bdist_egg', '--dist-dir', to_dir)
 
-        # returning the result
-        for file in os.listdir(to_dir):
-            if fnmatch.fnmatch(file, 'distribute-%s*.egg' % DEFAULT_VERSION):
-                return os.path.join(to_dir, file)
-
-        raise IOError('Could not build the egg.')
     finally:
         os.chdir(old_wd)
+    # returning the result
+    log.warn(egg)
+    if not os.path.exists(egg):
+        raise IOError('Could not build the egg.')
 
 
 def _do_download(version, download_base, to_dir, download_delay):
-    tarball = download_setuptools(version, download_base,
-                                  to_dir, download_delay)
-    egg = _build_egg(tarball, to_dir)
+    egg = os.path.join(to_dir, 'distribute-%s-py%d.%d.egg'
+                       % (version, sys.version_info[0], sys.version_info[1]))
+    if not os.path.exists(egg):
+        tarball = download_setuptools(version, download_base,
+                                      to_dir, download_delay)
+        _build_egg(egg, tarball, to_dir)
     sys.path.insert(0, egg)
     import setuptools
     setuptools.bootstrap_install_from = egg
 
 
 def use_setuptools(version=DEFAULT_VERSION, download_base=DEFAULT_URL,
-                   to_dir=os.curdir, download_delay=15):
+                   to_dir=os.curdir, download_delay=15, no_fake=True):
     # making sure we use the absolute path
     to_dir = os.path.abspath(to_dir)
     was_imported = 'pkg_resources' in sys.modules or \
@@ -133,21 +138,23 @@ def use_setuptools(version=DEFAULT_VERSION, download_base=DEFAULT_URL,
         try:
             import pkg_resources
             if not hasattr(pkg_resources, '_distribute'):
-                fake_setuptools()
+                if not no_fake:
+                    _fake_setuptools()
                 raise ImportError
         except ImportError:
             return _do_download(version, download_base, to_dir, download_delay)
         try:
             pkg_resources.require("distribute>="+version)
             return
-        except pkg_resources.VersionConflict, e:
+        except pkg_resources.VersionConflict:
+            e = sys.exc_info()[1]
             if was_imported:
-                print >>sys.stderr, (
+                sys.stderr.write(
                 "The required version of distribute (>=%s) is not available,\n"
                 "and can't be installed while this script is running. Please\n"
                 "install a more recent version first, using\n"
                 "'easy_install -U distribute'."
-                "\n\n(Currently using %r)") % (version, e.args[0])
+                "\n\n(Currently using %r)\n" % (version, e.args[0]))
                 sys.exit(2)
             else:
                 del pkg_resources, sys.modules['pkg_resources']    # reload ok
@@ -157,7 +164,8 @@ def use_setuptools(version=DEFAULT_VERSION, download_base=DEFAULT_URL,
             return _do_download(version, download_base, to_dir,
                                 download_delay)
     finally:
-        _create_fake_setuptools_pkg_info(to_dir)
+        if not no_fake:
+            _create_fake_setuptools_pkg_info(to_dir)
 
 def download_setuptools(version=DEFAULT_VERSION, download_base=DEFAULT_URL,
                         to_dir=os.curdir, delay=15):
@@ -171,7 +179,10 @@ def download_setuptools(version=DEFAULT_VERSION, download_base=DEFAULT_URL,
     """
     # making sure we use the absolute path
     to_dir = os.path.abspath(to_dir)
-    import urllib2
+    try:
+        from urllib.request import urlopen
+    except ImportError:
+        from urllib2 import urlopen
     tgz_name = "distribute-%s.tar.gz" % version
     url = download_base + tgz_name
     saveto = os.path.join(to_dir, tgz_name)
@@ -179,7 +190,7 @@ def download_setuptools(version=DEFAULT_VERSION, download_base=DEFAULT_URL,
     if not os.path.exists(saveto):  # Avoid repeated downloads
         try:
             log.warn("Downloading %s", url)
-            src = urllib2.urlopen(url)
+            src = urlopen(url)
             # Read/write all in one block, so we don't create a corrupt file
             # if the download is interrupted.
             data = src.read()
@@ -192,6 +203,29 @@ def download_setuptools(version=DEFAULT_VERSION, download_base=DEFAULT_URL,
                 dst.close()
     return os.path.realpath(saveto)
 
+def _no_sandbox(function):
+    def __no_sandbox(*args, **kw):
+        try:
+            from setuptools.sandbox import DirectorySandbox
+            if not hasattr(DirectorySandbox, '_old'):
+                def violation(*args):
+                    pass
+                DirectorySandbox._old = DirectorySandbox._violation
+                DirectorySandbox._violation = violation
+                patched = True
+            else:
+                patched = False
+        except ImportError:
+            patched = False
+
+        try:
+            return function(*args, **kw)
+        finally:
+            if patched:
+                DirectorySandbox._violation = DirectorySandbox._old
+                del DirectorySandbox._old
+
+    return __no_sandbox
 
 def _patch_file(path, content):
     """Will backup the file then patch it"""
@@ -209,26 +243,17 @@ def _patch_file(path, content):
         f.close()
     return True
 
+_patch_file = _no_sandbox(_patch_file)
 
 def _same_content(path, content):
     return open(path).read() == content
 
-
 def _rename_path(path):
     new_name = path + '.OLD.%s' % time.time()
     log.warn('Renaming %s into %s', path, new_name)
-    try:
-        from setuptools.sandbox import DirectorySandbox
-        def _violation(*args):
-            pass
-        DirectorySandbox._violation = _violation
-    except ImportError:
-        pass
-
     os.rename(path, new_name)
     return new_name
 
-
 def _remove_flat_installation(placeholder):
     if not os.path.isdir(placeholder):
         log.warn('Unkown installation at %s', placeholder)
@@ -262,6 +287,7 @@ def _remove_flat_installation(placeholder):
                      'Setuptools distribution', element)
     return True
 
+_remove_flat_installation = _no_sandbox(_remove_flat_installation)
 
 def _after_install(dist):
     log.warn('After install bootstrap.')
@@ -273,17 +299,20 @@ def _create_fake_setuptools_pkg_info(placeholder):
         log.warn('Could not find the install location')
         return
     pyver = '%s.%s' % (sys.version_info[0], sys.version_info[1])
-    setuptools_file = 'setuptools-0.6c9-py%s.egg-info' % pyver
+    setuptools_file = 'setuptools-%s-py%s.egg-info' % \
+            (SETUPTOOLS_FAKED_VERSION, pyver)
     pkg_info = os.path.join(placeholder, setuptools_file)
     if os.path.exists(pkg_info):
         log.warn('%s already exists', pkg_info)
         return
+
     log.warn('Creating %s', pkg_info)
     f = open(pkg_info, 'w')
     try:
         f.write(SETUPTOOLS_PKG_INFO)
     finally:
         f.close()
+
     pth_file = os.path.join(placeholder, 'setuptools.pth')
     log.warn('Creating %s', pth_file)
     f = open(pth_file, 'w')
@@ -292,6 +321,7 @@ def _create_fake_setuptools_pkg_info(placeholder):
     finally:
         f.close()
 
+_create_fake_setuptools_pkg_info = _no_sandbox(_create_fake_setuptools_pkg_info)
 
 def _patch_egg_dir(path):
     # let's check if it's already patched
@@ -311,10 +341,11 @@ def _patch_egg_dir(path):
         f.close()
     return True
 
+_patch_egg_dir = _no_sandbox(_patch_egg_dir)
 
 def _before_install():
     log.warn('Before install bootstrap.')
-    fake_setuptools()
+    _fake_setuptools()
 
 
 def _under_prefix(location):
@@ -330,12 +361,12 @@ def _under_prefix(location):
                 if len(args) > index:
                     top_dir = args[index+1]
                     return location.startswith(top_dir)
-            elif option == '--user' and USER_SITE is not None:
-                return location.startswith(USER_SITE)
+        if arg == '--user' and USER_SITE is not None:
+            return location.startswith(USER_SITE)
     return True
 
 
-def fake_setuptools():
+def _fake_setuptools():
     log.warn('Scanning installed packages')
     try:
         import pkg_resources
@@ -344,7 +375,13 @@ def fake_setuptools():
         log.warn('Setuptools or Distribute does not seem to be installed.')
         return
     ws = pkg_resources.working_set
-    setuptools_dist = ws.find(pkg_resources.Requirement.parse('setuptools'))
+    try:
+        setuptools_dist = ws.find(pkg_resources.Requirement.parse('setuptools',
+                                  replacement=False))
+    except TypeError:
+        # old distribute API
+        setuptools_dist = ws.find(pkg_resources.Requirement.parse('setuptools'))
+
     if setuptools_dist is None:
         log.warn('No setuptools distribution found')
         return
@@ -384,6 +421,9 @@ def fake_setuptools():
 def _relaunch():
     log.warn('Relaunching...')
     # we have to relaunch the process
+    # pip marker to avoid a relaunch bug
+    if sys.argv[:3] == ['-c', 'install', '--single-version-externally-managed']:
+        sys.argv[0] = 'setup.py'
     args = [sys.executable] + sys.argv
     sys.exit(subprocess.call(args))
 
@@ -408,7 +448,7 @@ def _extractall(self, path=".", members=None):
             # Extract directories with a safe mode.
             directories.append(tarinfo)
             tarinfo = copy.copy(tarinfo)
-            tarinfo.mode = 0700
+            tarinfo.mode = 448 # decimal for oct 0700
         self.extract(tarinfo, path)
 
     # Reverse sort directories.
@@ -427,7 +467,8 @@ def _extractall(self, path=".", members=None):
             self.chown(tarinfo, dirpath)
             self.utime(tarinfo, dirpath)
             self.chmod(tarinfo, dirpath)
-        except ExtractError, e:
+        except ExtractError:
+            e = sys.exc_info()[1]
             if self.errorlevel > 1:
                 raise
             else:
diff --git a/doc/source/array.rst b/doc/source/array.rst
new file mode 100644
index 0000000..93e5587
--- /dev/null
+++ b/doc/source/array.rst
@@ -0,0 +1,356 @@
+The :class:`Array` Class
+========================
+
+.. module:: pyopencl.array
+
+.. class:: DefaultAllocator(context, flags=pyopencl.mem_flags.READ_WRITE)
+
+    An Allocator that uses :class:`pyopencl.Buffer` with the given *flags*.
+
+    .. method:: __call__(self, size)
+
+.. class:: Array(context, shape, dtype, order="C", allocator=None, base=None, data=None, queue=None)
+
+    A :class:`numpy.ndarray` work-alike that stores its data and performs its
+    computations on the compute device.  *shape* and *dtype* work exactly as in
+    :mod:`numpy`.  Arithmetic methods in :class:`Array` support the
+    broadcasting of scalars. (e.g. `array+5`) If the
+
+    *allocator* is a callable that, upon being called with an argument of the number
+    of bytes to be allocated, returns an object that can be cast to an
+    :class:`int` representing the address of the newly allocated memory.
+    (See :class:`DefaultAllocator`.)
+
+    .. attribute :: data
+
+        The :class:`pyopencl.MemoryObject` instance created for the memory that backs
+        this :class:`Array`.
+
+    .. attribute :: shape
+
+        The tuple of lengths of each dimension in the array.
+
+    .. attribute :: dtype
+
+        The :class:`numpy.dtype` of the items in the GPU array.
+
+    .. attribute :: size
+
+        The number of meaningful entries in the array. Can also be computed by
+        multiplying up the numbers in :attr:`shape`.
+
+    .. attribute :: mem_size
+
+        The total number of entries, including padding, that are present in
+        the array.
+
+    .. attribute :: nbytes
+
+        The size of the entire array in bytes. Computed as :attr:`size` times
+        ``dtype.itemsize``.
+
+    .. method :: __len__()
+
+        Returns the size of the leading dimension of *self*.
+
+    .. method :: set(ary, queue=None, async=False)
+
+        Transfer the contents the :class:`numpy.ndarray` object *ary*
+        onto the device.
+
+        *ary* must have the same dtype and size (not necessarily shape) as *self*.
+
+
+    .. method :: get(queue=None, ary=None, async=False)
+
+        Transfer the contents of *self* into *ary* or a newly allocated
+        :mod:`numpy.ndarray`. If *ary* is given, it must have the right
+        size (not necessarily shape) and dtype.
+
+    .. method :: __str__()
+    .. method :: __repr__()
+
+    .. method :: mul_add(self, selffac, other, otherfac, queue=None):
+
+        Return `selffac*self + otherfac*other`.
+
+    .. method :: __add__(other)
+    .. method :: __sub__(other)
+    .. method :: __iadd__(other)
+    .. method :: __isub__(other)
+    .. method :: __neg__(other)
+    .. method :: __mul__(other)
+    .. method :: __div__(other)
+    .. method :: __rdiv__(other)
+    .. method :: __pow__(other)
+
+    .. method :: __abs__()
+
+        Return a :class:`Array` containing the absolute value of each
+        element of *self*.
+
+    .. UNDOC reverse()
+
+    .. method :: fill(scalar, queue=None)
+
+        Fill the array with *scalar*.
+
+    .. method :: astype(dtype, queue=None)
+
+        Return *self*, cast to *dtype*.
+
+Constructing :class:`Array` Instances
+----------------------------------------
+
+.. function:: to_device(context, queue, ary, allocator=None, async=False)
+
+    Return a :class:`Array` that is an exact copy of the :class:`numpy.ndarray`
+    instance *ary*.
+
+    See :class:`Array` for the meaning of *allocator*.
+
+.. function:: empty(context, shape, dtype, order="C", allocator=None, base=None, data=None, queue=None)
+
+    A synonym for the :class:`Array` constructor.
+
+.. function:: zeros(context, queue, shape, dtype, order="C", allocator=None)
+
+    Same as :func:`empty`, but the :class:`Array` is zero-initialized before
+    being returned.
+
+.. function:: empty_like(other_ary)
+
+    Make a new, uninitialized :class:`Array` having the same properties
+    as *other_ary*.
+
+.. function:: zeros_like(other_ary)
+
+    Make a new, zero-initialized :class:`Array` having the same properties
+    as *other_ary*.
+
+.. function:: arange(context, queue, start, stop, step, dtype=None)
+
+    Create a :class:`Array` filled with numbers spaced `step` apart,
+    starting from `start` and ending at `stop`.
+
+    For floating point arguments, the length of the result is
+    `ceil((stop - start)/step)`.  This rule may result in the last
+    element of the result being greater than `stop`.
+
+    *dtype*, if not specified, is taken as the largest common type
+    of *start*, *stop* and *step*.
+
+.. function:: take(a, indices, out=None, queue=None)
+
+    Return the :class:`Array` ``[a[indices[0]], ..., a[indices[n]]]``.
+    For the moment, *a* must be a type that can be bound to a texture.
+
+Conditionals
+^^^^^^^^^^^^
+
+.. function:: if_positive(criterion, then_, else_, out=None, queue=None)
+
+    Return an array like *then_*, which, for the element at index *i*,
+    contains *then_[i]* if *criterion[i]>0*, else *else_[i]*. (added in 0.94)
+
+.. function:: maximum(a, b, out=None, queue=None)
+
+    Return the elementwise maximum of *a* and *b*. (added in 0.94)
+
+.. function:: minimum(a, b, out=None, queue=None)
+
+    Return the elementwise minimum of *a* and *b*. (added in 0.94)
+
+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)
+.. TODO: atan2
+.. function:: atanh(array, queue=None)
+.. function:: atanpi(array, queue=None)
+.. TODO: 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
+-----------------------------------
+
+.. module:: pyopencl.clrandom
+
+.. function:: rand(context, queue, shape, dtype)
+
+    Return an array of `shape` filled with random values of `dtype`
+    in the range [0,1).
+
+Single-pass Custom Expression Evaluation
+----------------------------------------
+
+.. warning::
+
+    The following functionality is included in this documentation in the
+    hope that it may be useful, but its interface may change in future
+    revisions. Feedback is welcome.
+
+.. module:: pyopencl.elementwise
+
+Evaluating involved expressions on :class:`pyopencl.array.Array` instances 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.
+
+.. class:: ElementwiseKernel(context, arguments, operation, name="kernel", options=[])
+
+    Generate a kernel that takes a number of scalar or vector *arguments*
+    and performs the scalar *operation* on each entry of its arguments, if that
+    argument is a vector.
+
+    *arguments* is specified as a string formatted as a C argument list.
+    *operation* is specified as a C assignment statement, without a semicolon.
+    Vectors in *operation* should be indexed by the variable *i*.
+
+    *name* specifies the name as which the kernel is compiled, 
+    and *options* are passed unmodified to :meth:`pyopencl.Program.build`.
+
+    .. method:: __call__(*args)
+
+        Invoke the generated scalar kernel. The arguments may either be scalars or
+        :class:`GPUArray` instances.
+
+Here's a usage example::
+
+    import pyopencl as cl
+    import pyopencl.array as cl_array
+    import numpy
+
+    ctx = cl.create_some_context()
+    queue = cl.CommandQueue(ctx)
+
+    n = 10
+    a_gpu = cl_array.to_device(
+            ctx, queue, numpy.random.randn(n).astype(numpy.float32))
+    b_gpu = cl_array.to_device(
+            ctx, queue, numpy.random.randn(n).astype(numpy.float32))
+
+    from pyopencl.elementwise import ElementwiseKernel
+    lin_comb = ElementwiseKernel(ctx,
+            "float a, float *x, "
+            "float b, float *y, "
+            "float *z",
+            "z[i] = a*x[i] + b*y[i]",
+            "linear_combination")
+
+    c_gpu = cl_array.empty_like(a_gpu)
+    lin_comb(5, a_gpu, 6, b_gpu, c_gpu)
+
+    import numpy.linalg as la
+    assert la.norm((c_gpu - (5*a_gpu+6*b_gpu)).get()) < 1e-5
+
+(You can find this example as :file:`examples/demo_elementwise.py` in the PyOpenCL
+distribution.)
+
+Fast Fourier Transforms
+-----------------------
+
+Bogdan Opanchuk's `pyfft <http://pypi.python.org/pypi/pyfft>`_ package offers a
+variety of GPU-based FFT implementations.
+
diff --git a/doc/source/index.rst b/doc/source/index.rst
index 38456e9..c3bca32 100644
--- a/doc/source/index.rst
+++ b/doc/source/index.rst
@@ -70,6 +70,7 @@ Contents
     :maxdepth: 2
 
     runtime
+    array
     misc
 
 Note that this guide does not explain OpenCL programming and technology. Please 
diff --git a/doc/source/misc.rst b/doc/source/misc.rst
index 56ce060..bd92d09 100644
--- a/doc/source/misc.rst
+++ b/doc/source/misc.rst
@@ -82,9 +82,10 @@ Version 0.92
   emphasize the importance of *loccal_size*.
 * Add :meth:`pyopencl.Kernel.set_scalar_arg_dtypes`.
 * Add support for the
-  `cl_nv_device_attribute_query <ghttp://www.khronos.org/registry/cl/extensions/khr/cl_nv_device_attribute_query.txt>`_
+  `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
 --------------
diff --git a/doc/source/runtime.rst b/doc/source/runtime.rst
index c1c83d8..00c3ef2 100644
--- a/doc/source/runtime.rst
+++ b/doc/source/runtime.rst
@@ -565,7 +565,7 @@ Programs and Kernels
 
         See :class:`kernel_work_group_info` for values of *param*.
 
-    .. method:: set_arg(self, arg)
+    .. method:: set_arg(self, index, arg)
 
         *arg* may be
 
@@ -594,6 +594,8 @@ Programs and Kernels
 
         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
diff --git a/examples/demo_mandelbrot.py b/examples/demo_mandelbrot.py
index ddcf1b2..4919dbb 100644
--- a/examples/demo_mandelbrot.py
+++ b/examples/demo_mandelbrot.py
@@ -27,41 +27,43 @@ import pyopencl as cl
 # Speed notes are listed in the same place
 
 # set width and height of window, more pixels take longer to calculate
-w = 400
-h = 400
+w = 256
+h = 256
 
 def calc_fractal_opencl(q, maxiter):
     ctx = cl.Context(cl.get_platforms()[0].get_devices())
     queue = cl.CommandQueue(ctx)
 
-    output = np.empty(q.shape, dtype=np.uint64)# resize(np.array(0,), q.shape)
+    output = np.empty(q.shape, dtype=np.uint16)
 
     mf = cl.mem_flags
     q_opencl = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=q)
     output_opencl = cl.Buffer(ctx, mf.WRITE_ONLY, output.nbytes)
 
     prg = cl.Program(ctx, """
+    #pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable
     __kernel void mandelbrot(__global float2 *q,
-                     __global long *output, long const maxiter)
+                     __global ushort *output, ushort const maxiter)
     {
         int gid = get_global_id(0);
         float nreal, real = 0;
         float imag = 0;
+
+        output[gid] = 0;
+
         for(int curiter = 0; curiter < maxiter; curiter++) {
-            nreal = real*real - imag*imag + q[gid][0];
-            imag = 2* real*imag + q[gid][1];
+            nreal = real*real - imag*imag + q[gid].x;
+            imag = 2* real*imag + q[gid].y;
             real = nreal;
 
-            if (real*real + imag*imag > 4.) {
+            if (real*real + imag*imag > 4.0f)
                  output[gid] = curiter;
-                 break;
-            }
         }
     }
     """).build()
 
-    prg.mandelbrot(queue, output.shape, None, q_opencl,
-            output_opencl, np.int32(maxiter))
+    prg.mandelbrot(queue, output.shape, (64,), q_opencl,
+            output_opencl, np.uint16(maxiter))
 
     cl.enqueue_read_buffer(queue, output_opencl, output).wait()
 
@@ -121,10 +123,10 @@ if __name__ == '__main__':
             self.root.mainloop()
 
 
-        def draw(self, x1, x2, y1, y2, maxiter=300):
+        def draw(self, x1, x2, y1, y2, maxiter=30):
             # draw the Mandelbrot set, from numpy example
-            xx = np.arange(x1, x2, (x2-x1)/w*2)
-            yy = np.arange(y2, y1, (y1-y2)/h*2) * 1j
+            xx = np.arange(x1, x2, (x2-x1)/w)
+            yy = np.arange(y2, y1, (y1-y2)/h) * 1j
             q = np.ravel(xx+yy[:, np.newaxis]).astype(np.complex64)
 
             start_main = time.time()
@@ -134,18 +136,20 @@ if __name__ == '__main__':
             secs = end_main - start_main
             print "Main took", secs
 
-            output = (output + (256*output) + (256**2)*output) * 8
-            # convert output to a string
-            self.mandel = output.tostring()
+            self.mandel = (output.reshape((h,w)) /
+                    float(output.max()) * 255.).astype(np.uint8)
 
         def create_image(self):
             """"
             create the image from the draw() string
             """
-            self.im = Image.new("RGB", (w/2, h/2))
             # you can experiment with these x and y ranges
             self.draw(-2.13, 0.77, -1.3, 1.3)
-            self.im.fromstring(self.mandel, "raw", "RGBX", 0, -1)
+            self.im = Image.fromarray(self.mandel)
+            self.im.putpalette(reduce(
+                lambda a,b: a+b, ((i,0,0) for i in range(255))
+            ))
+
 
         def create_label(self):
             # put the image on a label widget
diff --git a/examples/gl_interop_demo.py b/examples/gl_interop_demo.py
index dd4bf3e..0246976 100644
--- a/examples/gl_interop_demo.py
+++ b/examples/gl_interop_demo.py
@@ -1,75 +1,80 @@
-from OpenGL.GL import *
-from OpenGL.GLUT import *
-from OpenGL.raw.GL.VERSION.GL_1_5 import glBufferData as rawGlBufferData
-try:
-    from OpenGL.WGL import wglGetCurrentDisplay as GetCurrentDisplay, wglGetCurrentContext as GetCurrentContext
-except:
-    pass
-try:
-    from OpenGL.GLX import glXGetCurrentDisplay as GetCurrentDisplay, glXGetCurrentContext as GetCurrentContext
-except:
-    pass
-import pyopencl as cl
-
-
-n_vertices = 10000
-
-src = """
-
-__kernel void generate_sin(__global float2* a)
-{
-    int id = get_global_id(0);
-    int n = get_global_size(0);
-    float r = (float)id / (float)n;
-    float x = r * 16.0f * 3.1415f;
-    a[id].x = r * 2.0f - 1.0f;
-    a[id].y = native_sin(x);
-}
-
-"""
-
-def initialize():
-    plats = cl.get_platforms()
-    ctx_props = cl.context_properties
-    props = [(ctx_props.PLATFORM, plats[0]), (ctx_props.GL_CONTEXT_KHR,
-        GetCurrentContext()), (ctx_props.GLX_DISPLAY_KHR, GetCurrentDisplay())]
-    ctx = cl.Context(properties=props)
-    glClearColor(1, 1, 1, 1)
-    glColor(0, 0, 1)
-    vbo = glGenBuffers(1)
-    glBindBuffer(GL_ARRAY_BUFFER, vbo)
-    rawGlBufferData(GL_ARRAY_BUFFER, n_vertices * 2 * 4, None, GL_STATIC_DRAW)
-    glEnableClientState(GL_VERTEX_ARRAY)
-    glVertexPointer(2, GL_FLOAT, 0, None)
-    coords_dev = cl.GLBuffer(ctx, cl.mem_flags.READ_WRITE, int(vbo))
-    prog = cl.Program(ctx, src).build()
-    queue = cl.CommandQueue(ctx)
-    cl.enqueue_acquire_gl_objects(queue, [coords_dev])
-    prog.generate_sin(queue, (n_vertices,), None, coords_dev)
-    cl.enqueue_release_gl_objects(queue, [coords_dev])
-    queue.finish()
-    glFlush()
-
-def display():
-    glClear(GL_COLOR_BUFFER_BIT)
-    glDrawArrays(GL_LINE_STRIP, 0, n_vertices)
-    glFlush()
-
-def reshape(w, h):
-    glViewport(0, 0, w, h)
-    glMatrixMode(GL_PROJECTION)
-    glLoadIdentity()
-    glMatrixMode(GL_MODELVIEW)
-
-if __name__ == '__main__':
-    import sys
-    glutInit(sys.argv)
-    if len(sys.argv) > 1:
-        n_vertices = int(sys.argv[1])
-    glutInitWindowSize(800, 160)
-    glutInitWindowPosition(0, 0)
-    glutCreateWindow('OpenCL/OpenGL Interop Tutorial: Sin Generator')
-    glutDisplayFunc(display)
-    glutReshapeFunc(reshape)
-    initialize()
-    glutMainLoop()
+from OpenGL.GL import *
+from OpenGL.GLUT import *
+from OpenGL.raw.GL.VERSION.GL_1_5 import glBufferData as rawGlBufferData
+from OpenGL import platform, GLX, WGL
+import pyopencl as cl
+
+
+n_vertices = 10000
+
+src = """
+
+__kernel void generate_sin(__global float2* a)
+{
+    int id = get_global_id(0);
+    int n = get_global_size(0);
+    float r = (float)id / (float)n;
+    float x = r * 16.0f * 3.1415f;
+    a[id].x = r * 2.0f - 1.0f;
+    a[id].y = native_sin(x);
+}
+
+"""
+
+def initialize():
+    plats = cl.get_platforms()
+    ctx_props = cl.context_properties
+
+    props = [(ctx_props.PLATFORM, plats[0]), 
+            (ctx_props.GL_CONTEXT_KHR, platform.GetCurrentContext())]
+
+    import sys
+    if sys.platform == "linux2":
+        props.append(
+                (ctx_props.GLX_DISPLAY_KHR, 
+                    GLX.glXGetCurrentDisplay()))
+    elif sys.platform == "win32":
+        props.append(
+                (ctx_props.WGL_HDC_KHR, 
+                    WGL.wglGetCurrentDC()))
+    ctx = cl.Context(properties=props)
+
+    glClearColor(1, 1, 1, 1)
+    glColor(0, 0, 1)
+    vbo = glGenBuffers(1)
+    glBindBuffer(GL_ARRAY_BUFFER, vbo)
+    rawGlBufferData(GL_ARRAY_BUFFER, n_vertices * 2 * 4, None, GL_STATIC_DRAW)
+    glEnableClientState(GL_VERTEX_ARRAY)
+    glVertexPointer(2, GL_FLOAT, 0, None)
+    coords_dev = cl.GLBuffer(ctx, cl.mem_flags.READ_WRITE, int(vbo))
+    prog = cl.Program(ctx, src).build()
+    queue = cl.CommandQueue(ctx)
+    cl.enqueue_acquire_gl_objects(queue, [coords_dev])
+    prog.generate_sin(queue, (n_vertices,), None, coords_dev)
+    cl.enqueue_release_gl_objects(queue, [coords_dev])
+    queue.finish()
+    glFlush()
+
+def display():
+    glClear(GL_COLOR_BUFFER_BIT)
+    glDrawArrays(GL_LINE_STRIP, 0, n_vertices)
+    glFlush()
+
+def reshape(w, h):
+    glViewport(0, 0, w, h)
+    glMatrixMode(GL_PROJECTION)
+    glLoadIdentity()
+    glMatrixMode(GL_MODELVIEW)
+
+if __name__ == '__main__':
+    import sys
+    glutInit(sys.argv)
+    if len(sys.argv) > 1:
+        n_vertices = int(sys.argv[1])
+    glutInitWindowSize(800, 160)
+    glutInitWindowPosition(0, 0)
+    glutCreateWindow('OpenCL/OpenGL Interop Tutorial: Sin Generator')
+    glutDisplayFunc(display)
+    glutReshapeFunc(reshape)
+    initialize()
+    glutMainLoop()
diff --git a/pyopencl/__init__.py b/pyopencl/__init__.py
index e9c4474..0108257 100644
--- a/pyopencl/__init__.py
+++ b/pyopencl/__init__.py
@@ -224,6 +224,10 @@ def _add_functionality():
                 self.set_arg(i, arg)
         else:
             from struct import pack
+
+            if len(args) != len(arg_type_chars):
+                raise ValueError("length of argument type array and "
+                        "length of argument list do not agree")
             for i, (arg, arg_type_char) in enumerate(
                     zip(args, arg_type_chars)):
                 if arg_type_char:
diff --git a/pyopencl/array.py b/pyopencl/array.py
index a84a8eb..a754305 100644
--- a/pyopencl/array.py
+++ b/pyopencl/array.py
@@ -126,7 +126,7 @@ class Array(object):
     work on an element-by-element basis, just like :class:`numpy.ndarray`.
     """
 
-    def __init__(self, context, shape, dtype, allocator=None,
+    def __init__(self, context, shape, dtype, order="C", allocator=None,
             base=None, data=None, queue=None):
         if allocator is None:
             allocator = DefaultAllocator(context)
@@ -147,6 +147,9 @@ class Array(object):
 
         self.shape = shape
         self.dtype = numpy.dtype(dtype)
+        if order not in ["C", "F"]:
+            raise ValueError("order must be either 'C' or 'F'")
+        self.order = order
 
         self.mem_size = self.size = s
         self.nbytes = self.dtype.itemsize * self.size
@@ -180,7 +183,7 @@ class Array(object):
 
     def get(self, queue=None, ary=None, async=False):
         if ary is None:
-            ary = numpy.empty(self.shape, self.dtype)
+            ary = numpy.empty(self.shape, self.dtype, order=self.order)
         else:
             if ary.size != self.size:
                 raise TypeError("'ary' has non-matching type")
@@ -344,12 +347,20 @@ class Array(object):
         return result
 
     def __iadd__(self, other):
-        self._axpbyz(self, 1, self, 1, other)
-        return self
+        if isinstance(other, Array):
+            self._axpbyz(self, 1, self, 1, other)
+            return self
+        else:
+            self._axpbz(self, 1, self, other)
+            return self
 
     def __isub__(self, other):
-        self._axpbyz(self, 1, self, -1, other)
-        return self
+        if isinstance(other, Array):
+            self._axpbyz(self, 1, self, -1, other)
+            return self
+        else:
+            self._axpbz(self, 1, self, -other)
+            return self
 
     def __neg__(self):
         result = self._new_like_me()
@@ -470,30 +481,6 @@ class Array(object):
         self._copy(result, self, queue=queue)
         return result
 
-    # slicing -----------------------------------------------------------------
-    def __getitem__(self, idx):
-        if idx == ():
-            return self
-
-        if len(self.shape) > 1:
-            raise NotImplementedError("multi-d slicing is not yet implemented")
-
-        if not isinstance(idx, slice):
-            raise ValueError("non-slice indexing not supported: %s" % (idx,))
-
-        l, = self.shape
-        start, stop, stride = idx.indices(l)
-
-        if stride != 1:
-            raise NotImplementedError("strided slicing is not yet implemented")
-
-        return Array(
-                shape=((stop-start)//stride,),
-                dtype=self.dtype,
-                allocator=self.allocator,
-                base=self,
-                data=int(self.data) + start*self.dtype.itemsize)
-
     # rich comparisons (or rather, lack thereof) ------------------------------
     def __eq__(self, other):
         raise NotImplementedError
@@ -517,7 +504,15 @@ class Array(object):
 
 def to_device(context, queue, ary, allocator=None, async=False):
     """Converts a numpy array to a :class:`Array`."""
-    result = Array(context, ary.shape, ary.dtype, allocator,
+    if ary.flags.f_contiguous:
+        order = "F"
+    elif ary.flags.c_contiguous:
+        order = "C"
+    else:
+        raise ValueError("to_device only works on C- or Fortran-"
+                "contiguous arrays")
+
+    result = Array(context, ary.shape, ary.dtype, order, allocator,
             queue=queue)
     result.set(ary, async=async)
     return result
@@ -525,21 +520,22 @@ def to_device(context, queue, ary, allocator=None, async=False):
 
 empty = Array
 
-def zeros(context, queue, shape, dtype, allocator=None):
+def zeros(context, queue, shape, dtype, order="C", allocator=None):
     """Returns an array of the given shape and dtype filled with 0's."""
 
-    result = Array(context, shape, dtype, allocator, queue=queue)
+    result = Array(context, shape, dtype, 
+            order=order, allocator=allocator, queue=queue)
     result.fill(0)
     return result
 
 def empty_like(ary):
     result = Array(ary.context,
-            ary.shape, ary.dtype, ary.allocator, queue=ary.queue)
+            ary.shape, ary.dtype, allocator=ary.allocator, queue=ary.queue)
     return result
 
 def zeros_like(ary):
     result = Array(ary.context,
-            ary.shape, ary.dtype, ary.allocator, queue=ary.queue)
+            ary.shape, ary.dtype, allocator=ary.allocator, queue=ary.queue)
     result.fill(0)
     return result
 
@@ -641,7 +637,8 @@ def _take(result, ary, indices):
 
 def take(a, indices, out=None, queue=None):
     if out is None:
-        out = Array(a.context, indices.shape, a.dtype, a.allocator,
+        out = Array(a.context, indices.shape, a.dtype, 
+                allocator=a.allocator,
                 queue=queue or a.queue)
 
     assert len(indices.shape) == 1
@@ -666,7 +663,8 @@ def multi_take(arrays, indices, out=None, queue=None):
     vec_count = len(arrays)
 
     if out is None:
-        out = [Array(context, queue, indices.shape, a_dtype, a_allocator)
+        out = [Array(context, queue, indices.shape, a_dtype, 
+            allocator=a_allocator)
                 for i in range(vec_count)]
     else:
         if len(out) != len(arrays):
@@ -715,7 +713,8 @@ def multi_take_put(arrays, dest_indices, src_indices, dest_shape=None,
     vec_count = len(arrays)
 
     if out is None:
-        out = [Array(context, dest_shape, a_dtype, a_allocator, queue=queue)
+        out = [Array(context, dest_shape, a_dtype, 
+            allocator=a_allocator, queue=queue)
                 for i in range(vec_count)]
     else:
         if a_dtype != single_valued(o.dtype for o in out):
@@ -783,7 +782,7 @@ def multi_put(arrays, dest_indices, dest_shape=None, out=None, queue=None):
     vec_count = len(arrays)
 
     if out is None:
-        out = [Array(context, dest_shape, a_dtype, a_allocator, queue=queue)
+        out = [Array(context, dest_shape, a_dtype, allocator=a_allocator, queue=queue)
                 for i in range(vec_count)]
     else:
         if a_dtype != single_valued(o.dtype for o in out):
diff --git a/pyopencl/version.py b/pyopencl/version.py
index 107a527..359274f 100644
--- a/pyopencl/version.py
+++ b/pyopencl/version.py
@@ -1,5 +1,5 @@
 VERSION = (0, 92)
-VERSION_STATUS = "beta"
+VERSION_STATUS = ""
 VERSION_TEXT = ".".join(str(x) for x in VERSION) + VERSION_STATUS
 
 
diff --git a/src/wrapper/wrap_cl.hpp b/src/wrapper/wrap_cl.hpp
index e001d90..233ea32 100644
--- a/src/wrapper/wrap_cl.hpp
+++ b/src/wrapper/wrap_cl.hpp
@@ -20,6 +20,12 @@
 #include <CL/cl.h>
 // FIXME: Nvidia doesn't install cl_ext.h by default. Grr.
 // #include <CL/cl_ext.h>
+
+#ifdef _WIN32
+#define NOMINMAX
+#include <windows.h>
+#endif
+
 #ifdef HAVE_GL
 #include <GL/gl.h>
 #include <CL/cl_gl.h>
@@ -2575,7 +2581,8 @@ namespace pyopencl
               std::vector<size_t> result;
               PYOPENCL_GET_VEC_INFO(KernelWorkGroup,
                   PYOPENCL_FIRST_ARG, param_name, result);
-              return py::list(result);
+
+              PYOPENCL_RETURN_VECTOR(size_t, result);
             }
           case CL_KERNEL_LOCAL_MEM_SIZE:
 #ifdef CL_VERSION_1_1
@@ -2875,8 +2882,8 @@ namespace pyopencl
     std::vector<cl_context_properties> props
       = parse_context_properties(py_properties);
 
-    typedef CL_API_ENTRY cl_int CL_API_CALL
-      (*func_ptr_type)(const cl_context_properties * /* properties */,
+    typedef CL_API_ENTRY cl_int (CL_API_CALL
+      *func_ptr_type)(const cl_context_properties * /* properties */,
           cl_gl_context_info            /* param_name */,
           size_t                        /* param_value_size */,
           void *                        /* param_value */,
diff --git a/src/wrapper/wrap_constants.cpp b/src/wrapper/wrap_constants.cpp
index 4f2548b..a81adc9 100644
--- a/src/wrapper/wrap_constants.cpp
+++ b/src/wrapper/wrap_constants.cpp
@@ -167,6 +167,7 @@ void pyopencl_expose_constants()
     ADD_ATTR(DEVICE_, QUEUE_PROPERTIES);
     ADD_ATTR(DEVICE_, NAME);
     ADD_ATTR(DEVICE_, VENDOR);
+    ADD_ATTR(, DRIVER_VERSION);
     ADD_ATTR(DEVICE_, VERSION);
     ADD_ATTR(DEVICE_, PROFILE);
     ADD_ATTR(DEVICE_, VERSION);
diff --git a/test/test_clmath.py b/test/test_clmath.py
new file mode 100644
index 0000000..58b6a09
--- /dev/null
+++ b/test/test_clmath.py
@@ -0,0 +1,182 @@
+from __future__ import division
+import math
+import numpy
+import pytools.test
+
+def have_cl():
+    try:
+        import pyopencl
+        return True
+    except:
+        return False
+
+if have_cl():
+    import pyopencl.array as cl_array
+    import pyopencl as cl
+    import pyopencl.clmath as clmath
+    from pyopencl.tools import pytest_generate_tests_for_pyopencl \
+            as pytest_generate_tests
+
+
+
+
+
+sizes = [10, 128, 1024, 1<<10, 1<<13]
+
+
+
+
+def has_double_support(dev):
+    for ext in dev.extensions.split(" "):
+        if ext == "cl_khr_fp64":
+            return True
+    return False
+
+
+
+
+numpy_func_names = {
+        "asin": "arcsin",
+        "acos": "arccos",
+        "atan": "arctan",
+        }
+
+
+
+
+def make_unary_function_test(name, (a, b)=(0, 1), threshold=0):
+    def test(ctx_getter):
+        context = ctx_getter()
+        queue = cl.CommandQueue(context)
+
+        gpu_func = getattr(clmath, name)
+        cpu_func = getattr(numpy, numpy_func_names.get(name, name))
+
+        if has_double_support(context.devices[0]):
+            dtypes = [numpy.float32, numpy.float64]
+        else:
+            dtypes = [numpy.float32]
+
+        for s in sizes:
+            for dtype in dtypes:
+                args = cl_array.arange(context, queue, a, b, (b-a)/s, 
+                        dtype=numpy.float32)
+                gpu_results = gpu_func(args).get()
+                cpu_results = cpu_func(args.get())
+
+                max_err = numpy.max(numpy.abs(cpu_results - gpu_results))
+                assert (max_err <= threshold).all(), \
+                        (max_err, name, dtype)
+
+    return pytools.test.mark_test.opencl(test)
+
+
+
+
+if have_cl():
+    test_ceil = make_unary_function_test("ceil", (-10, 10))
+    test_floor = make_unary_function_test("ceil", (-10, 10))
+    test_fabs = make_unary_function_test("fabs", (-10, 10))
+    test_exp = make_unary_function_test("exp", (-3, 3), 1e-5)
+    test_log = make_unary_function_test("log", (1e-5, 1), 1e-6)
+    test_log10 = make_unary_function_test("log10", (1e-5, 1), 5e-7)
+    test_sqrt = make_unary_function_test("sqrt", (1e-5, 1), 2e-7)
+
+    test_sin = make_unary_function_test("sin", (-10, 10), 1e-7)
+    test_cos = make_unary_function_test("cos", (-10, 10), 1e-7)
+    test_asin = make_unary_function_test("asin", (-0.9, 0.9), 5e-7)
+    test_acos = make_unary_function_test("acos", (-0.9, 0.9), 5e-7)
+    test_tan = make_unary_function_test("tan", 
+            (-math.pi/2 + 0.1, math.pi/2 - 0.1), 1e-5)
+    test_atan = make_unary_function_test("atan", (-10, 10), 2e-7)
+
+    test_sinh = make_unary_function_test("sinh", (-3, 3), 1e-6)
+    test_cosh = make_unary_function_test("cosh", (-3, 3), 1e-6)
+    test_tanh = make_unary_function_test("tanh", (-3, 3), 2e-6)
+
+
+
+
+@pytools.test.mark_test.opencl
+def test_fmod(ctx_getter):
+    context = ctx_getter()
+    queue = cl.CommandQueue(context)
+
+    for s in sizes:
+        a = cl_array.arange(context, queue, s, dtype=numpy.float32)/10
+        a2 = cl_array.arange(context, queue, s, dtype=numpy.float32)/45.2 + 0.1
+        b = clmath.fmod(a, a2)
+
+        a = a.get()
+        a2 = a2.get()
+        b = b.get()
+
+        for i in range(s):
+            assert math.fmod(a[i], a2[i]) == b[i]
+
+@pytools.test.mark_test.opencl
+def test_ldexp(ctx_getter):
+    context = ctx_getter()
+    queue = cl.CommandQueue(context)
+
+    for s in sizes:
+        a = cl_array.arange(context, queue, s, dtype=numpy.float32)
+        a2 = cl_array.arange(context, queue, s, dtype=numpy.float32)*1e-3
+        b = clmath.ldexp(a,a2)
+
+        a = a.get()
+        a2 = a2.get()
+        b = b.get()
+
+        for i in range(s):
+            assert math.ldexp(a[i], int(a2[i])) == b[i]
+
+@pytools.test.mark_test.opencl
+def test_modf(ctx_getter):
+    context = ctx_getter()
+    queue = cl.CommandQueue(context)
+
+    for s in sizes:
+        a = cl_array.arange(context, queue, s, dtype=numpy.float32)/10
+        fracpart, intpart = clmath.modf(a)
+
+        a = a.get()
+        intpart = intpart.get()
+        fracpart = fracpart.get()
+
+        for i in range(s):
+            fracpart_true, intpart_true = math.modf(a[i])
+
+            assert intpart_true == intpart[i]
+            assert abs(fracpart_true - fracpart[i]) < 1e-4
+
+@pytools.test.mark_test.opencl
+def test_frexp(ctx_getter):
+    context = ctx_getter()
+    queue = cl.CommandQueue(context)
+
+    for s in sizes:
+        a = cl_array.arange(context, queue, s, dtype=numpy.float32)/10
+        significands, exponents = clmath.frexp(a)
+
+        a = a.get()
+        significands = significands.get()
+        exponents = exponents.get()
+
+        for i in range(s):
+            sig_true, ex_true = math.frexp(a[i])
+
+            assert sig_true == significands[i]
+            assert ex_true == exponents[i]
+
+
+
+
+if __name__ == "__main__":
+    # make sure that import failures get reported, instead of skipping the tests.
+    import sys
+    if len(sys.argv) > 1:
+        exec sys.argv[1]
+    else:
+        from py.test.cmdline import main
+        main([__file__])
diff --git a/test/test_math.py b/test/test_math.py
deleted file mode 100644
index c7aabf4..0000000
--- a/test/test_math.py
+++ /dev/null
@@ -1,182 +0,0 @@
-from __future__ import division
-import math
-import numpy
-import pytools.test
-
-def have_cl():
-    try:
-        import pyopencl
-        return True
-    except:
-        return False
-
-if have_cl():
-    import pyopencl.array as cl_array
-    import pyopencl as cl
-    import pyopencl.clmath as clmath
-    from pyopencl.tools import pytest_generate_tests_for_pyopencl \
-            as pytest_generate_tests
-
-
-
-
-
-sizes = [10, 128, 1024, 1<<10, 1<<13]
-
-
-
-
-def has_double_support(dev):
-    for ext in dev.extensions.split(" "):
-        if ext == "cl_khr_fp64":
-            return True
-    return False
-
-
-
-
-numpy_func_names = {
-        "asin": "arcsin",
-        "acos": "arccos",
-        "atan": "arctan",
-        }
-
-
-
-
-def make_unary_function_test(name, (a, b)=(0, 1), threshold=0):
-    def test(ctx_getter):
-        context = ctx_getter()
-        queue = cl.CommandQueue(context)
-
-        gpu_func = getattr(clmath, name)
-        cpu_func = getattr(numpy, numpy_func_names.get(name, name))
-
-        if has_double_support(context.devices[0]):
-            dtypes = [numpy.float32, numpy.float64]
-        else:
-            dtypes = [numpy.float32]
-
-        for s in sizes:
-            for dtype in dtypes:
-                args = cl_array.arange(context, queue, a, b, (b-a)/s, 
-                        dtype=numpy.float32)
-                gpu_results = gpu_func(args).get()
-                cpu_results = cpu_func(args.get())
-
-                max_err = numpy.max(numpy.abs(cpu_results - gpu_results))
-                assert (max_err <= threshold).all(), \
-                        (max_err, name, dtype)
-
-    return pytools.test.mark_test.opencl(test)
-
-
-
-
-if have_cl():
-    test_ceil = make_unary_function_test("ceil", (-10, 10))
-    test_floor = make_unary_function_test("ceil", (-10, 10))
-    test_fabs = make_unary_function_test("fabs", (-10, 10))
-    test_exp = make_unary_function_test("exp", (-3, 3), 1e-5)
-    test_log = make_unary_function_test("log", (1e-5, 1), 5e-7)
-    test_log10 = make_unary_function_test("log10", (1e-5, 1), 3e-7)
-    test_sqrt = make_unary_function_test("sqrt", (1e-5, 1), 2e-7)
-
-    test_sin = make_unary_function_test("sin", (-10, 10), 1e-7)
-    test_cos = make_unary_function_test("cos", (-10, 10), 1e-7)
-    test_asin = make_unary_function_test("asin", (-0.9, 0.9), 5e-7)
-    test_acos = make_unary_function_test("acos", (-0.9, 0.9), 5e-7)
-    test_tan = make_unary_function_test("tan", 
-            (-math.pi/2 + 0.1, math.pi/2 - 0.1), 1e-5)
-    test_atan = make_unary_function_test("atan", (-10, 10), 2e-7)
-
-    test_sinh = make_unary_function_test("sinh", (-3, 3), 1e-6)
-    test_cosh = make_unary_function_test("cosh", (-3, 3), 1e-6)
-    test_tanh = make_unary_function_test("tanh", (-3, 3), 2e-6)
-
-
-
-
-@pytools.test.mark_test.opencl
-def test_fmod(ctx_getter):
-    context = ctx_getter()
-    queue = cl.CommandQueue(context)
-
-    for s in sizes:
-        a = cl_array.arange(context, queue, s, dtype=numpy.float32)/10
-        a2 = cl_array.arange(context, queue, s, dtype=numpy.float32)/45.2 + 0.1
-        b = clmath.fmod(a, a2)
-
-        a = a.get()
-        a2 = a2.get()
-        b = b.get()
-
-        for i in range(s):
-            assert math.fmod(a[i], a2[i]) == b[i]
-
-@pytools.test.mark_test.opencl
-def test_ldexp(ctx_getter):
-    context = ctx_getter()
-    queue = cl.CommandQueue(context)
-
-    for s in sizes:
-        a = cl_array.arange(context, queue, s, dtype=numpy.float32)
-        a2 = cl_array.arange(context, queue, s, dtype=numpy.float32)*1e-3
-        b = clmath.ldexp(a,a2)
-
-        a = a.get()
-        a2 = a2.get()
-        b = b.get()
-
-        for i in range(s):
-            assert math.ldexp(a[i], int(a2[i])) == b[i]
-
-@pytools.test.mark_test.opencl
-def test_modf(ctx_getter):
-    context = ctx_getter()
-    queue = cl.CommandQueue(context)
-
-    for s in sizes:
-        a = cl_array.arange(context, queue, s, dtype=numpy.float32)/10
-        fracpart, intpart = clmath.modf(a)
-
-        a = a.get()
-        intpart = intpart.get()
-        fracpart = fracpart.get()
-
-        for i in range(s):
-            fracpart_true, intpart_true = math.modf(a[i])
-
-            assert intpart_true == intpart[i]
-            assert abs(fracpart_true - fracpart[i]) < 1e-4
-
-@pytools.test.mark_test.opencl
-def test_frexp(ctx_getter):
-    context = ctx_getter()
-    queue = cl.CommandQueue(context)
-
-    for s in sizes:
-        a = cl_array.arange(context, queue, s, dtype=numpy.float32)/10
-        significands, exponents = clmath.frexp(a)
-
-        a = a.get()
-        significands = significands.get()
-        exponents = exponents.get()
-
-        for i in range(s):
-            sig_true, ex_true = math.frexp(a[i])
-
-            assert sig_true == significands[i]
-            assert ex_true == exponents[i]
-
-
-
-
-if __name__ == "__main__":
-    # make sure that import failures get reported, instead of skipping the tests.
-    import sys
-    if len(sys.argv) > 1:
-        exec sys.argv[1]
-    else:
-        from py.test.cmdline import main
-        main([__file__])

Attachment: signature.asc
Description: This is a digitally signed message part


Reply to: