VERSION, cmd/geth: bumped version 1.4.3

[release/1.4.3] miner: fixed pending state by not shutting down update loop
(cherry picked from commit a824c3f02f)
2026-03-01 22:46:48 +01:00 · 2016-05-10 13:46:03 +02:00 · 2016-05-10 13:46:03 +02:00 · 2016-05-10 13:46:02 +02:00 · 2016-05-09 22:06:27 +02:00 · 2016-05-09 22:04:40 +02:00
2329 changed files with 163301 additions and 473006 deletions
--- a/.github/CONTRIBUTING.md
+++ b/.github/CONTRIBUTING.md
@@ -8,7 +8,7 @@ and help.
 ## Contributing

 If you'd like to contribute to go-ethereum please fork, fix, commit and
-send a pull request. Commits which do not comply with the coding standards
+send a pull request. Commits who do not comply with the coding standards
 are ignored (use gofmt!). If you send pull requests make absolute sure that you
 commit on the `develop` branch and that you do not merge to master.
 Commits that are directly based on master are simply ignored.
--- a/.gitignore
+++ b/.gitignore
@@ -13,7 +13,8 @@
 .ethtest
 */**/*tx_database*
 */**/*dapps*
-build/_vendor/pkg
+Godeps/_workspace/pkg
+Godeps/_workspace/bin

 #*
 .#*
@@ -22,11 +23,17 @@ build/_vendor/pkg
 .project
 .settings

+deploy/osx/Mist.app
+deploy/osx/Mist\ Installer.dmg
+cmd/mist/assets/ext/ethereum.js/
+
 # used by the Makefile
 /build/_workspace/
 /build/bin/
-/geth*.zip

 # travis
 profile.tmp
 profile.cov
+
+# vagrant
+.vagrant
--- a/.gitmodules
+++ b/.gitmodules
@@ -0,0 +1,3 @@
+[submodule "cmd/mist/assets/ext/ethereum.js"]
+	path = cmd/mist/assets/ext/ethereum.js
+	url = https://github.com/ethereum/web3.js
--- a/.mailmap
+++ b/.mailmap
@@ -17,7 +17,6 @@ Taylor Gerring <taylor.gerring@gmail.com> <taylor.gerring@ethereum.org>
 Bas van Kervel <bas@ethdev.com>
 Bas van Kervel <bas@ethdev.com> <basvankervel@ziggo.nl>
 Bas van Kervel <bas@ethdev.com> <basvankervel@gmail.com>
-Bas van Kervel <bas@ethdev.com> <bas-vk@users.noreply.github.com>

 Sven Ehlert <sven@ethdev.com>

@@ -63,30 +62,4 @@ Joseph Chow <ethereum@outlook.com> ethers <TODO>

 Enrique Fynn <enriquefynn@gmail.com>

-Vincent G <caktux@gmail.com>
-
-RJ Catalano <rj@erisindustries.com>
-
-Nchinda Nchinda <nchinda2@gmail.com>
-
-Aron Fischer <homotopycolimit@users.noreply.github.com>
-
-Vlad Gluhovsky <gluk256@users.noreply.github.com>
-
-Ville Sundell <github@solarius.fi>
-
-Elliot Shepherd <elliot@identitii.com>
-
-Yohann Léon <sybiload@gmail.com>
-
-Gregg Dourgarian <greggd@tempworks.com>
-
-Casey Detrio <cdetrio@gmail.com>
-
-Jens Agerberg <github@agerberg.me>
-
-Nick Johnson <arachnid@notdot.net>
-
-Henning Diedrich <hd@eonblast.com>
-Henning Diedrich <hd@eonblast.com> Drake Burroughs <wildfyre@hotmail.com>
-
+Vincent G <caktux@gmail.com>
--- a/.travis.yml
+++ b/.travis.yml
@@ -1,100 +1,29 @@
 language: go
-go_import_path: github.com/ethereum/go-ethereum
-sudo: false
-matrix:
-  include:
-    - os: linux
-      dist: trusty
-      go: 1.5.4
-      env:
-        - GO15VENDOREXPERIMENT=1
-    - os: linux
-      dist: trusty
-      go: 1.6.2
-    - os: linux
-      dist: trusty
-      go: 1.7
-    - os: osx
-      go: 1.7
-
-    # This builder does the Ubuntu PPA and Linux Azure uploads
-    - os: linux
-      dist: trusty
-      sudo: required
-      go: 1.7
-      env:
-        - ubuntu-ppa
-        - azure-linux
-      addons:
-        apt:
-          packages:
-            - devscripts
-            - debhelper
-            - dput
-            - gcc-multilib
-      script:
-        # Build for the primary platforms that Trusty can manage
-        - go run build/ci.go debsrc -signer "Go Ethereum Linux Builder <geth-ci@ethereum.org>" -upload ppa:ethereum/ethereum
-        - go run build/ci.go install
-        - go run build/ci.go archive -type tar -signer LINUX_SIGNING_KEY -upload gethstore/builds
-        - go run build/ci.go install -arch 386
-        - go run build/ci.go archive -arch 386 -type tar -signer LINUX_SIGNING_KEY -upload gethstore/builds
-
-        # Switch over GCC to cross compilation (breaks 386, hence why do it here only)
-        - sudo -E apt-get -yq --no-install-suggests --no-install-recommends --force-yes install gcc-arm-linux-gnueabi libc6-dev-armel-cross gcc-arm-linux-gnueabihf libc6-dev-armhf-cross gcc-aarch64-linux-gnu libc6-dev-arm64-cross
-        - sudo ln -s /usr/include/asm-generic /usr/include/asm
-
-        - GOARM=5 CC=arm-linux-gnueabi-gcc go run build/ci.go install -arch arm
-        - GOARM=5 go run build/ci.go archive -arch arm -type tar -signer LINUX_SIGNING_KEY -upload gethstore/builds
-        - GOARM=6 CC=arm-linux-gnueabi-gcc go run build/ci.go install -arch arm
-        - GOARM=6 go run build/ci.go archive -arch arm -type tar -signer LINUX_SIGNING_KEY -upload gethstore/builds
-        - GOARM=7 CC=arm-linux-gnueabihf-gcc go run build/ci.go install -arch arm
-        - GOARM=7 go run build/ci.go archive -arch arm -type tar -signer LINUX_SIGNING_KEY -upload gethstore/builds
-        - CC=aarch64-linux-gnu-gcc go run build/ci.go install -arch arm64
-        - go run build/ci.go archive -arch arm64 -type tar -signer LINUX_SIGNING_KEY -upload gethstore/builds
-
-    # This builder does the OSX Azure, Android Maven and Azure and iOS CocoaPods and Azure uploads
-    - os: osx
-      go: 1.7
-      env:
-        - azure-osx
-        - mobile
-      script:
-        - go run build/ci.go install
-        - go run build/ci.go archive -type tar -signer OSX_SIGNING_KEY -upload gethstore/builds
-
-        # Build the Android archive and upload it to Maven Central and Azure
-        - brew update
-        - brew install android-sdk maven gpg
-        - alias gpg="gpg2"
-
-        - export ANDROID_HOME=/usr/local/opt/android-sdk
-        - echo "y" | android update sdk --no-ui --filter `android list sdk | grep "SDK Platform Android" | grep -E 'API 15|API 19|API 24' | awk '{print $1}' | cut -d '-' -f 1 | tr '\n' ','`
-
-        - go run build/ci.go aar -signer ANDROID_SIGNING_KEY -deploy https://oss.sonatype.org -upload gethstore/builds
-
-        # Build the iOS framework and upload it to CocoaPods and Azure
-        - gem uninstall cocoapods -a
-        - gem install cocoapods --pre
-
-        - mv ~/.cocoapods/repos/master ~/.cocoapods/repos/master.bak
-        - sed -i '.bak' 's/repo.join/!repo.join/g' $(dirname `gem which cocoapods`)/cocoapods/sources_manager.rb
-        - if [ "$TRAVIS_PULL_REQUEST" = "false" ]; then git clone --depth=1 https://github.com/CocoaPods/Specs.git ~/.cocoapods/repos/master && pod setup --verbose; fi
-
-        - xctool -version
-        - xcrun simctl list
-
-        - go run build/ci.go xcode -signer IOS_SIGNING_KEY -deploy trunk -upload gethstore/builds
-
+go:
+  - 1.4.2
 install:
+  # - go get code.google.com/p/go.tools/cmd/goimports
+  # - go get github.com/golang/lint/golint
+  # - go get golang.org/x/tools/cmd/vet 
  - go get golang.org/x/tools/cmd/cover
+before_script:
+  # - gofmt -l -w .
+  # - goimports -l -w .
+  # - golint .
+  # - go vet ./...
+  # - go test -race ./...
 script:
-  - go run build/ci.go install
-  - go run build/ci.go test -coverage -vet
-
+  - make travis-test-with-coverage
+after_success:
+  - bash <(curl -s https://codecov.io/bash)
+env:
+  global:
+    - secure: "U2U1AmkU4NJBgKR/uUAebQY87cNL0+1JHjnLOmmXwxYYyj5ralWb1aSuSH3qSXiT93qLBmtaUkuv9fberHVqrbAeVlztVdUsKAq7JMQH+M99iFkC9UiRMqHmtjWJ0ok4COD1sRYixxi21wb/JrMe3M1iL4QJVS61iltjHhVdM64="
+sudo: false
 notifications:
  webhooks:
    urls:
      - https://webhooks.gitter.im/e/e09ccdce1048c5e03445
-    on_success: change
+    on_success: change 
    on_failure: always
+    on_start: false 
--- a/29
+++ b/29
@@ -1,63 +1,36 @@
 # This is the official list of go-ethereum authors for copyright purposes.

-Ales Katona <ales@coinbase.com>
 Alex Leverington <alex@ethdev.com>
 Alexandre Van de Sande <alex.vandesande@ethdev.com>
-Aron Fischer <homotopycolimit@users.noreply.github.com>
 Bas van Kervel <bas@ethdev.com>
-Benjamin Brent <benjamin@benjaminbrent.com>
-Casey Detrio <cdetrio@gmail.com>
 Christoph Jentzsch <jentzsch.software@gmail.com>
 Daniel A. Nagy <nagy.da@gmail.com>
-Elliot Shepherd <elliot@identitii.com>
+Drake Burroughs <wildfyre@hotmail.com>
 Enrique Fynn <enriquefynn@gmail.com>
 Ethan Buchman <ethan@coinculture.info>
 Fabian Vogelsteller <fabian@frozeman.de>
-Fabio Berger <fabioberger1991@gmail.com>
 Felix Lange <fjl@twurst.com>
-Gregg Dourgarian <greggd@tempworks.com>
 Gustav Simonsson <gustav.simonsson@gmail.com>
-Hao Bryan Cheng <haobcheng@gmail.com>
-Henning Diedrich <hd@eonblast.com>
 Isidoro Ghezzi <isidoro.ghezzi@icloud.com>
 Jae Kwon <jkwon.work@gmail.com>
 Jason Carver <jacarver@linkedin.com>
 Jeff R. Allen <jra@nella.org>
 Jeffrey Wilcke <jeffrey@ethereum.org>
-Jens Agerberg <github@agerberg.me>
-Jonathan Brown <jbrown@bluedroplet.com>
 Joseph Chow <ethereum@outlook.com>
-Justin Clark-Casey <justincc@justincc.org>
-Kenji Siu <kenji@isuntv.com>
 Kobi Gurkan <kobigurk@gmail.com>
 Lefteris Karapetsas <lefteris@refu.co>
 Leif Jurvetson <leijurv@gmail.com>
 Maran Hidskes <maran.hidskes@gmail.com>
 Marek Kotewicz <marek.kotewicz@gmail.com>
-Martin Holst Swende <martin@swende.se>
-Matthew Di Ferrante <mattdf@users.noreply.github.com>
 Matthew Wampler-Doty <matthew.wampler.doty@gmail.com>
-Nchinda Nchinda <nchinda2@gmail.com>
 Nick Dodson <silentcicero@outlook.com>
-Nick Johnson <arachnid@notdot.net>
-Paulo L F Casaretto <pcasaretto@gmail.com>
 Peter Pratscher <pratscher@gmail.com>
 Péter Szilágyi <peterke@gmail.com>
-RJ Catalano <rj@erisindustries.com>
 Ramesh Nair <ram@hiddentao.com>
 Ricardo Catalinas Jiménez <r@untroubled.be>
 Rémy Roy <remyroy@remyroy.com>
-Stein Dekker <dekker.stein@gmail.com>
-Steven Roose <stevenroose@gmail.com>
 Taylor Gerring <taylor.gerring@gmail.com>
-Thomas Bocek <tom@tomp2p.net>
-Tosh Camille <tochecamille@gmail.com>
 Viktor Trón <viktor.tron@gmail.com>
-Ville Sundell <github@solarius.fi>
 Vincent G <caktux@gmail.com>
 Vitalik Buterin <v@buterin.com>
-Vlad Gluhovsky <gluk256@users.noreply.github.com>
-Yohann Léon <sybiload@gmail.com>
-Yoichi Hirai <i@yoichihirai.com>
 Zsolt Felföldi <zsfelfoldi@gmail.com>
-ΞTHΞЯSPHΞЯΞ <{viktor.tron,nagydani,zsfelfoldi}@gmail.com>
--- a/14
+++ b/14
@@ -1,14 +0,0 @@
-FROM alpine:3.3
-
-ADD . /go-ethereum
-RUN \
-  apk add --update git go make gcc musl-dev         && \
-  (cd go-ethereum && make geth)                     && \
-  cp go-ethereum/build/bin/geth /geth               && \
-  apk del git go make gcc musl-dev                  && \
-  rm -rf /go-ethereum && rm -rf /var/cache/apk/*
-
-EXPOSE 8545
-EXPOSE 30303
-
-ENTRYPOINT ["/geth"]
--- a/Godeps/Godeps.json
+++ b/Godeps/Godeps.json
@@ -0,0 +1,324 @@
+{
+	"ImportPath": "github.com/ethereum/go-ethereum",
+	"GoVersion": "go1.5.2",
+	"Packages": [
+		"./..."
+	],
+	"Deps": [
+		{
+			"ImportPath": "github.com/Gustav-Simonsson/go-opencl/cl",
+			"Rev": "593e01cfc4f3353585015321e01951d4a907d3ef"
+		},
+		{
+			"ImportPath": "github.com/cespare/cp",
+			"Rev": "165db2f241fd235aec29ba6d9b1ccd5f1c14637c"
+		},
+		{
+			"ImportPath": "github.com/codegangsta/cli",
+			"Comment": "1.2.0-215-g0ab42fd",
+			"Rev": "0ab42fd482c27cf2c95e7794ad3bb2082c2ab2d7"
+		},
+		{
+			"ImportPath": "github.com/davecgh/go-spew/spew",
+			"Rev": "5215b55f46b2b919f50a1df0eaa5886afe4e3b3d"
+		},
+		{
+			"ImportPath": "github.com/ethereum/ethash",
+			"Comment": "v23.1-245-g25b32de",
+			"Rev": "25b32de0c0271065c28c3719c2bfe86959d72f0c"
+		},
+		{
+			"ImportPath": "github.com/fatih/color",
+			"Comment": "v0.1-12-g9aae6aa",
+			"Rev": "9aae6aaa22315390f03959adca2c4d395b02fcef"
+		},
+		{
+			"ImportPath": "github.com/gizak/termui",
+			"Rev": "08a5d3f67b7d9ec87830ea39c48e570a1f18531f"
+		},
+		{
+			"ImportPath": "github.com/golang/snappy",
+			"Rev": "799c780093d646c1b79d30894e22512c319fa137"
+		},
+		{
+			"ImportPath": "github.com/hashicorp/golang-lru",
+			"Rev": "a0d98a5f288019575c6d1f4bb1573fef2d1fcdc4"
+		},
+		{
+			"ImportPath": "github.com/hashicorp/golang-lru/simplelru",
+			"Rev": "a0d98a5f288019575c6d1f4bb1573fef2d1fcdc4"
+		},
+		{
+			"ImportPath": "github.com/huin/goupnp",
+			"Rev": "46bde78b11f3f021f2a511df138be9e2fc7506e8"
+		},
+		{
+			"ImportPath": "github.com/huin/goupnp/dcps/internetgateway1",
+			"Rev": "46bde78b11f3f021f2a511df138be9e2fc7506e8"
+		},
+		{
+			"ImportPath": "github.com/huin/goupnp/dcps/internetgateway2",
+			"Rev": "46bde78b11f3f021f2a511df138be9e2fc7506e8"
+		},
+		{
+			"ImportPath": "github.com/huin/goupnp/httpu",
+			"Rev": "46bde78b11f3f021f2a511df138be9e2fc7506e8"
+		},
+		{
+			"ImportPath": "github.com/huin/goupnp/scpd",
+			"Rev": "46bde78b11f3f021f2a511df138be9e2fc7506e8"
+		},
+		{
+			"ImportPath": "github.com/huin/goupnp/soap",
+			"Rev": "46bde78b11f3f021f2a511df138be9e2fc7506e8"
+		},
+		{
+			"ImportPath": "github.com/huin/goupnp/ssdp",
+			"Rev": "46bde78b11f3f021f2a511df138be9e2fc7506e8"
+		},
+		{
+			"ImportPath": "github.com/jackpal/gateway",
+			"Rev": "192609c58b8985e645cbe82ddcb28a4362ca0fdc"
+		},
+		{
+			"ImportPath": "github.com/jackpal/go-nat-pmp",
+			"Rev": "46523a463303c6ede3ddfe45bde1c7ed52ebaacd"
+		},
+		{
+			"ImportPath": "github.com/mattn/go-colorable",
+			"Rev": "9fdad7c47650b7d2e1da50644c1f4ba7f172f252"
+		},
+		{
+			"ImportPath": "github.com/mattn/go-isatty",
+			"Rev": "56b76bdf51f7708750eac80fa38b952bb9f32639"
+		},
+		{
+			"ImportPath": "github.com/mattn/go-runewidth",
+			"Comment": "travisish-44-ge882a96",
+			"Rev": "e882a96ec18dd43fa283187b66af74497c9101c0"
+		},
+		{
+			"ImportPath": "github.com/microsoft/go-winio",
+			"Comment": "v0.2.0",
+			"Rev": "9e2895e5f6c3f16473b91d37fae6e89990a4520c"
+		},
+		{
+			"ImportPath": "github.com/nsf/termbox-go",
+			"Rev": "362329b0aa6447eadd52edd8d660ec1dff470295"
+		},
+		{
+			"ImportPath": "github.com/pborman/uuid",
+			"Comment": "v1.0-6-g0f1a469",
+			"Rev": "0f1a46960a86dcdf5dd30d3e6568a497a997909f"
+		},
+		{
+			"ImportPath": "github.com/peterh/liner",
+			"Rev": "ad1edfd30321d8f006ccf05f1e0524adeb943060"
+		},
+		{
+			"ImportPath": "github.com/rcrowley/go-metrics",
+			"Rev": "51425a2415d21afadfd55cd93432c0bc69e9598d"
+		},
+		{
+			"ImportPath": "github.com/rjeczalik/notify",
+			"Rev": "5dd6205716539662f8f14ab513552b41eab69d5d"
+		},
+		{
+			"ImportPath": "github.com/robertkrimen/otto",
+			"Rev": "53221230c215611a90762720c9042ac782ef74ee"
+		},
+		{
+			"ImportPath": "github.com/robertkrimen/otto/ast",
+			"Rev": "53221230c215611a90762720c9042ac782ef74ee"
+		},
+		{
+			"ImportPath": "github.com/robertkrimen/otto/dbg",
+			"Rev": "53221230c215611a90762720c9042ac782ef74ee"
+		},
+		{
+			"ImportPath": "github.com/robertkrimen/otto/file",
+			"Rev": "53221230c215611a90762720c9042ac782ef74ee"
+		},
+		{
+			"ImportPath": "github.com/robertkrimen/otto/parser",
+			"Rev": "53221230c215611a90762720c9042ac782ef74ee"
+		},
+		{
+			"ImportPath": "github.com/robertkrimen/otto/registry",
+			"Rev": "53221230c215611a90762720c9042ac782ef74ee"
+		},
+		{
+			"ImportPath": "github.com/robertkrimen/otto/token",
+			"Rev": "53221230c215611a90762720c9042ac782ef74ee"
+		},
+		{
+			"ImportPath": "github.com/rs/cors",
+			"Rev": "5950cf11d77f8a61b432a25dd4d444b4ced01379"
+		},
+		{
+			"ImportPath": "github.com/syndtr/goleveldb/leveldb",
+			"Rev": "917f41c560270110ceb73c5b38be2a9127387071"
+		},
+		{
+			"ImportPath": "github.com/syndtr/goleveldb/leveldb/cache",
+			"Rev": "917f41c560270110ceb73c5b38be2a9127387071"
+		},
+		{
+			"ImportPath": "github.com/syndtr/goleveldb/leveldb/comparer",
+			"Rev": "917f41c560270110ceb73c5b38be2a9127387071"
+		},
+		{
+			"ImportPath": "github.com/syndtr/goleveldb/leveldb/errors",
+			"Rev": "917f41c560270110ceb73c5b38be2a9127387071"
+		},
+		{
+			"ImportPath": "github.com/syndtr/goleveldb/leveldb/filter",
+			"Rev": "917f41c560270110ceb73c5b38be2a9127387071"
+		},
+		{
+			"ImportPath": "github.com/syndtr/goleveldb/leveldb/iterator",
+			"Rev": "917f41c560270110ceb73c5b38be2a9127387071"
+		},
+		{
+			"ImportPath": "github.com/syndtr/goleveldb/leveldb/journal",
+			"Rev": "917f41c560270110ceb73c5b38be2a9127387071"
+		},
+		{
+			"ImportPath": "github.com/syndtr/goleveldb/leveldb/memdb",
+			"Rev": "917f41c560270110ceb73c5b38be2a9127387071"
+		},
+		{
+			"ImportPath": "github.com/syndtr/goleveldb/leveldb/opt",
+			"Rev": "917f41c560270110ceb73c5b38be2a9127387071"
+		},
+		{
+			"ImportPath": "github.com/syndtr/goleveldb/leveldb/storage",
+			"Rev": "917f41c560270110ceb73c5b38be2a9127387071"
+		},
+		{
+			"ImportPath": "github.com/syndtr/goleveldb/leveldb/table",
+			"Rev": "917f41c560270110ceb73c5b38be2a9127387071"
+		},
+		{
+			"ImportPath": "github.com/syndtr/goleveldb/leveldb/util",
+			"Rev": "917f41c560270110ceb73c5b38be2a9127387071"
+		},
+		{
+			"ImportPath": "golang.org/x/crypto/pbkdf2",
+			"Rev": "1f22c0103821b9390939b6776727195525381532"
+		},
+		{
+			"ImportPath": "golang.org/x/crypto/ripemd160",
+			"Rev": "1f22c0103821b9390939b6776727195525381532"
+		},
+		{
+			"ImportPath": "golang.org/x/crypto/scrypt",
+			"Rev": "1f22c0103821b9390939b6776727195525381532"
+		},
+		{
+			"ImportPath": "golang.org/x/net/context",
+			"Rev": "8968c61983e8f51a91b8c0ef25bf739278c89634"
+		},
+		{
+			"ImportPath": "golang.org/x/net/html",
+			"Rev": "8968c61983e8f51a91b8c0ef25bf739278c89634"
+		},
+		{
+			"ImportPath": "golang.org/x/net/html/atom",
+			"Rev": "8968c61983e8f51a91b8c0ef25bf739278c89634"
+		},
+		{
+			"ImportPath": "golang.org/x/net/html/charset",
+			"Rev": "8968c61983e8f51a91b8c0ef25bf739278c89634"
+		},
+		{
+			"ImportPath": "golang.org/x/net/websocket",
+			"Rev": "8968c61983e8f51a91b8c0ef25bf739278c89634"
+		},
+		{
+			"ImportPath": "golang.org/x/sys/unix",
+			"Rev": "50c6bc5e4292a1d4e65c6e9be5f53be28bcbe28e"
+		},
+		{
+			"ImportPath": "golang.org/x/text/encoding",
+			"Rev": "09761194ac5034a97b2bfad4f5b896b0ac350b3e"
+		},
+		{
+			"ImportPath": "golang.org/x/text/encoding/charmap",
+			"Rev": "09761194ac5034a97b2bfad4f5b896b0ac350b3e"
+		},
+		{
+			"ImportPath": "golang.org/x/text/encoding/htmlindex",
+			"Rev": "09761194ac5034a97b2bfad4f5b896b0ac350b3e"
+		},
+		{
+			"ImportPath": "golang.org/x/text/encoding/internal",
+			"Rev": "09761194ac5034a97b2bfad4f5b896b0ac350b3e"
+		},
+		{
+			"ImportPath": "golang.org/x/text/encoding/internal/identifier",
+			"Rev": "09761194ac5034a97b2bfad4f5b896b0ac350b3e"
+		},
+		{
+			"ImportPath": "golang.org/x/text/encoding/japanese",
+			"Rev": "09761194ac5034a97b2bfad4f5b896b0ac350b3e"
+		},
+		{
+			"ImportPath": "golang.org/x/text/encoding/korean",
+			"Rev": "09761194ac5034a97b2bfad4f5b896b0ac350b3e"
+		},
+		{
+			"ImportPath": "golang.org/x/text/encoding/simplifiedchinese",
+			"Rev": "09761194ac5034a97b2bfad4f5b896b0ac350b3e"
+		},
+		{
+			"ImportPath": "golang.org/x/text/encoding/traditionalchinese",
+			"Rev": "09761194ac5034a97b2bfad4f5b896b0ac350b3e"
+		},
+		{
+			"ImportPath": "golang.org/x/text/encoding/unicode",
+			"Rev": "09761194ac5034a97b2bfad4f5b896b0ac350b3e"
+		},
+		{
+			"ImportPath": "golang.org/x/text/internal/tag",
+			"Rev": "09761194ac5034a97b2bfad4f5b896b0ac350b3e"
+		},
+		{
+			"ImportPath": "golang.org/x/text/internal/utf8internal",
+			"Rev": "09761194ac5034a97b2bfad4f5b896b0ac350b3e"
+		},
+		{
+			"ImportPath": "golang.org/x/text/language",
+			"Rev": "09761194ac5034a97b2bfad4f5b896b0ac350b3e"
+		},
+		{
+			"ImportPath": "golang.org/x/text/runes",
+			"Rev": "09761194ac5034a97b2bfad4f5b896b0ac350b3e"
+		},
+		{
+			"ImportPath": "golang.org/x/text/transform",
+			"Rev": "09761194ac5034a97b2bfad4f5b896b0ac350b3e"
+		},
+		{
+			"ImportPath": "golang.org/x/tools/go/ast/astutil",
+			"Rev": "758728c4b28cfbac299730969ef8f655c4761283"
+		},
+		{
+			"ImportPath": "golang.org/x/tools/imports",
+			"Rev": "758728c4b28cfbac299730969ef8f655c4761283"
+		},
+		{
+			"ImportPath": "gopkg.in/check.v1",
+			"Rev": "4f90aeace3a26ad7021961c297b22c42160c7b25"
+		},
+		{
+			"ImportPath": "gopkg.in/fatih/set.v0",
+			"Comment": "v0.1.0-3-g27c4092",
+			"Rev": "27c40922c40b43fe04554d8223a402af3ea333f3"
+		},
+		{
+			"ImportPath": "gopkg.in/karalabe/cookiejar.v2/collections/prque",
+			"Rev": "8dcd6a7f4951f6ff3ee9cbb919a06d8925822e57"
+		}
+	]
+}
--- a/Godeps/Readme
+++ b/Godeps/Readme
@@ -0,0 +1,5 @@
+This directory tree is generated automatically by godep.
+
+Please do not edit.
+
+See https://github.com/tools/godep for more information.
--- a/Godeps/_workspace/.gitignore
+++ b/Godeps/_workspace/.gitignore
@@ -0,0 +1,2 @@
+/pkg
+/bin
--- a/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/cl.go
+++ b/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/cl.go
@@ -0,0 +1,26 @@
+/*
+Package cl provides a binding to the OpenCL api. It's mostly a low-level
+wrapper that avoids adding functionality while still making the interface
+a little more friendly and easy to use.
+
+Resource life-cycle management:
+
+For any CL object that gets created (buffer, queue, kernel, etc..) you should
+call object.Release() when finished with it to free the CL resources. This
+explicitely calls the related clXXXRelease method for the type. However,
+as a fallback there is a finalizer set for every resource item that takes
+care of it (eventually) if Release isn't called. In this way you can have
+better control over the life cycle of resources while having a fall back
+to avoid leaks. This is similar to how file handles and such are handled
+in the Go standard packages.
+*/
+package cl
+
+// #include "headers/1.2/opencl.h"
+// #cgo CFLAGS: -Iheaders/1.2
+// #cgo darwin LDFLAGS: -framework OpenCL
+// #cgo linux LDFLAGS: -lOpenCL
+import "C"
+import "errors"
+
+var ErrUnsupported = errors.New("cl: unsupported")
--- a/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/context.go
+++ b/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/context.go
@@ -0,0 +1,161 @@
+package cl
+
+// #include <stdlib.h>
+// #ifdef __APPLE__
+// #include "OpenCL/opencl.h"
+// #else
+// #include "cl.h"
+// #endif
+import "C"
+
+import (
+	"runtime"
+	"unsafe"
+)
+
+const maxImageFormats = 256
+
+type Context struct {
+	clContext C.cl_context
+	devices   []*Device
+}
+
+type MemObject struct {
+	clMem C.cl_mem
+	size  int
+}
+
+func releaseContext(c *Context) {
+	if c.clContext != nil {
+		C.clReleaseContext(c.clContext)
+		c.clContext = nil
+	}
+}
+
+func releaseMemObject(b *MemObject) {
+	if b.clMem != nil {
+		C.clReleaseMemObject(b.clMem)
+		b.clMem = nil
+	}
+}
+
+func newMemObject(mo C.cl_mem, size int) *MemObject {
+	memObject := &MemObject{clMem: mo, size: size}
+	runtime.SetFinalizer(memObject, releaseMemObject)
+	return memObject
+}
+
+func (b *MemObject) Release() {
+	releaseMemObject(b)
+}
+
+// TODO: properties
+func CreateContext(devices []*Device) (*Context, error) {
+	deviceIds := buildDeviceIdList(devices)
+	var err C.cl_int
+	clContext := C.clCreateContext(nil, C.cl_uint(len(devices)), &deviceIds[0], nil, nil, &err)
+	if err != C.CL_SUCCESS {
+		return nil, toError(err)
+	}
+	if clContext == nil {
+		return nil, ErrUnknown
+	}
+	context := &Context{clContext: clContext, devices: devices}
+	runtime.SetFinalizer(context, releaseContext)
+	return context, nil
+}
+
+func (ctx *Context) GetSupportedImageFormats(flags MemFlag, imageType MemObjectType) ([]ImageFormat, error) {
+	var formats [maxImageFormats]C.cl_image_format
+	var nFormats C.cl_uint
+	if err := C.clGetSupportedImageFormats(ctx.clContext, C.cl_mem_flags(flags), C.cl_mem_object_type(imageType), maxImageFormats, &formats[0], &nFormats); err != C.CL_SUCCESS {
+		return nil, toError(err)
+	}
+	fmts := make([]ImageFormat, nFormats)
+	for i, f := range formats[:nFormats] {
+		fmts[i] = ImageFormat{
+			ChannelOrder:    ChannelOrder(f.image_channel_order),
+			ChannelDataType: ChannelDataType(f.image_channel_data_type),
+		}
+	}
+	return fmts, nil
+}
+
+func (ctx *Context) CreateCommandQueue(device *Device, properties CommandQueueProperty) (*CommandQueue, error) {
+	var err C.cl_int
+	clQueue := C.clCreateCommandQueue(ctx.clContext, device.id, C.cl_command_queue_properties(properties), &err)
+	if err != C.CL_SUCCESS {
+		return nil, toError(err)
+	}
+	if clQueue == nil {
+		return nil, ErrUnknown
+	}
+	commandQueue := &CommandQueue{clQueue: clQueue, device: device}
+	runtime.SetFinalizer(commandQueue, releaseCommandQueue)
+	return commandQueue, nil
+}
+
+func (ctx *Context) CreateProgramWithSource(sources []string) (*Program, error) {
+	cSources := make([]*C.char, len(sources))
+	for i, s := range sources {
+		cs := C.CString(s)
+		cSources[i] = cs
+		defer C.free(unsafe.Pointer(cs))
+	}
+	var err C.cl_int
+	clProgram := C.clCreateProgramWithSource(ctx.clContext, C.cl_uint(len(sources)), &cSources[0], nil, &err)
+	if err != C.CL_SUCCESS {
+		return nil, toError(err)
+	}
+	if clProgram == nil {
+		return nil, ErrUnknown
+	}
+	program := &Program{clProgram: clProgram, devices: ctx.devices}
+	runtime.SetFinalizer(program, releaseProgram)
+	return program, nil
+}
+
+func (ctx *Context) CreateBufferUnsafe(flags MemFlag, size int, dataPtr unsafe.Pointer) (*MemObject, error) {
+	var err C.cl_int
+	clBuffer := C.clCreateBuffer(ctx.clContext, C.cl_mem_flags(flags), C.size_t(size), dataPtr, &err)
+	if err != C.CL_SUCCESS {
+		return nil, toError(err)
+	}
+	if clBuffer == nil {
+		return nil, ErrUnknown
+	}
+	return newMemObject(clBuffer, size), nil
+}
+
+func (ctx *Context) CreateEmptyBuffer(flags MemFlag, size int) (*MemObject, error) {
+	return ctx.CreateBufferUnsafe(flags, size, nil)
+}
+
+func (ctx *Context) CreateEmptyBufferFloat32(flags MemFlag, size int) (*MemObject, error) {
+	return ctx.CreateBufferUnsafe(flags, 4*size, nil)
+}
+
+func (ctx *Context) CreateBuffer(flags MemFlag, data []byte) (*MemObject, error) {
+	return ctx.CreateBufferUnsafe(flags, len(data), unsafe.Pointer(&data[0]))
+}
+
+//float64
+func (ctx *Context) CreateBufferFloat32(flags MemFlag, data []float32) (*MemObject, error) {
+	return ctx.CreateBufferUnsafe(flags, 4*len(data), unsafe.Pointer(&data[0]))
+}
+
+func (ctx *Context) CreateUserEvent() (*Event, error) {
+	var err C.cl_int
+	clEvent := C.clCreateUserEvent(ctx.clContext, &err)
+	if err != C.CL_SUCCESS {
+		return nil, toError(err)
+	}
+	return newEvent(clEvent), nil
+}
+
+func (ctx *Context) Release() {
+	releaseContext(ctx)
+}
+
+// http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateSubBuffer.html
+// func (memObject *MemObject) CreateSubBuffer(flags MemFlag, bufferCreateType BufferCreateType, )
--- a/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/device.go
+++ b/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/device.go
@@ -0,0 +1,510 @@
+package cl
+
+// #ifdef __APPLE__
+// #include "OpenCL/opencl.h"
+// #else
+// #include "cl.h"
+// #include "cl_ext.h"
+// #endif
+import "C"
+
+import (
+	"strings"
+	"unsafe"
+)
+
+const maxDeviceCount = 64
+
+type DeviceType uint
+
+const (
+	DeviceTypeCPU         DeviceType = C.CL_DEVICE_TYPE_CPU
+	DeviceTypeGPU         DeviceType = C.CL_DEVICE_TYPE_GPU
+	DeviceTypeAccelerator DeviceType = C.CL_DEVICE_TYPE_ACCELERATOR
+	DeviceTypeDefault     DeviceType = C.CL_DEVICE_TYPE_DEFAULT
+	DeviceTypeAll         DeviceType = C.CL_DEVICE_TYPE_ALL
+)
+
+type FPConfig int
+
+const (
+	FPConfigDenorm         FPConfig = C.CL_FP_DENORM           // denorms are supported
+	FPConfigInfNaN         FPConfig = C.CL_FP_INF_NAN          // INF and NaNs are supported
+	FPConfigRoundToNearest FPConfig = C.CL_FP_ROUND_TO_NEAREST // round to nearest even rounding mode supported
+	FPConfigRoundToZero    FPConfig = C.CL_FP_ROUND_TO_ZERO    // round to zero rounding mode supported
+	FPConfigRoundToInf     FPConfig = C.CL_FP_ROUND_TO_INF     // round to positive and negative infinity rounding modes supported
+	FPConfigFMA            FPConfig = C.CL_FP_FMA              // IEEE754-2008 fused multiply-add is supported
+	FPConfigSoftFloat      FPConfig = C.CL_FP_SOFT_FLOAT       // Basic floating-point operations (such as addition, subtraction, multiplication) are implemented in software
+)
+
+var fpConfigNameMap = map[FPConfig]string{
+	FPConfigDenorm:         "Denorm",
+	FPConfigInfNaN:         "InfNaN",
+	FPConfigRoundToNearest: "RoundToNearest",
+	FPConfigRoundToZero:    "RoundToZero",
+	FPConfigRoundToInf:     "RoundToInf",
+	FPConfigFMA:            "FMA",
+	FPConfigSoftFloat:      "SoftFloat",
+}
+
+func (c FPConfig) String() string {
+	var parts []string
+	for bit, name := range fpConfigNameMap {
+		if c&bit != 0 {
+			parts = append(parts, name)
+		}
+	}
+	if parts == nil {
+		return ""
+	}
+	return strings.Join(parts, "|")
+}
+
+func (dt DeviceType) String() string {
+	var parts []string
+	if dt&DeviceTypeCPU != 0 {
+		parts = append(parts, "CPU")
+	}
+	if dt&DeviceTypeGPU != 0 {
+		parts = append(parts, "GPU")
+	}
+	if dt&DeviceTypeAccelerator != 0 {
+		parts = append(parts, "Accelerator")
+	}
+	if dt&DeviceTypeDefault != 0 {
+		parts = append(parts, "Default")
+	}
+	if parts == nil {
+		parts = append(parts, "None")
+	}
+	return strings.Join(parts, "|")
+}
+
+type Device struct {
+	id C.cl_device_id
+}
+
+func buildDeviceIdList(devices []*Device) []C.cl_device_id {
+	deviceIds := make([]C.cl_device_id, len(devices))
+	for i, d := range devices {
+		deviceIds[i] = d.id
+	}
+	return deviceIds
+}
+
+// Obtain the list of devices available on a platform. 'platform' refers
+// to the platform returned by GetPlatforms or can be nil. If platform
+// is nil, the behavior is implementation-defined.
+func GetDevices(platform *Platform, deviceType DeviceType) ([]*Device, error) {
+	var deviceIds [maxDeviceCount]C.cl_device_id
+	var numDevices C.cl_uint
+	var platformId C.cl_platform_id
+	if platform != nil {
+		platformId = platform.id
+	}
+	if err := C.clGetDeviceIDs(platformId, C.cl_device_type(deviceType), C.cl_uint(maxDeviceCount), &deviceIds[0], &numDevices); err != C.CL_SUCCESS {
+		return nil, toError(err)
+	}
+	if numDevices > maxDeviceCount {
+		numDevices = maxDeviceCount
+	}
+	devices := make([]*Device, numDevices)
+	for i := 0; i < int(numDevices); i++ {
+		devices[i] = &Device{id: deviceIds[i]}
+	}
+	return devices, nil
+}
+
+func (d *Device) nullableId() C.cl_device_id {
+	if d == nil {
+		return nil
+	}
+	return d.id
+}
+
+func (d *Device) GetInfoString(param C.cl_device_info, panicOnError bool) (string, error) {
+	var strC [1024]C.char
+	var strN C.size_t
+	if err := C.clGetDeviceInfo(d.id, param, 1024, unsafe.Pointer(&strC), &strN); err != C.CL_SUCCESS {
+		if panicOnError {
+			panic("Should never fail")
+		}
+		return "", toError(err)
+	}
+
+	// OpenCL strings are NUL-terminated, and the terminator is included in strN
+	// Go strings aren't NUL-terminated, so subtract 1 from the length
+	return C.GoStringN((*C.char)(unsafe.Pointer(&strC)), C.int(strN-1)), nil
+}
+
+func (d *Device) getInfoUint(param C.cl_device_info, panicOnError bool) (uint, error) {
+	var val C.cl_uint
+	if err := C.clGetDeviceInfo(d.id, param, C.size_t(unsafe.Sizeof(val)), unsafe.Pointer(&val), nil); err != C.CL_SUCCESS {
+		if panicOnError {
+			panic("Should never fail")
+		}
+		return 0, toError(err)
+	}
+	return uint(val), nil
+}
+
+func (d *Device) getInfoSize(param C.cl_device_info, panicOnError bool) (int, error) {
+	var val C.size_t
+	if err := C.clGetDeviceInfo(d.id, param, C.size_t(unsafe.Sizeof(val)), unsafe.Pointer(&val), nil); err != C.CL_SUCCESS {
+		if panicOnError {
+			panic("Should never fail")
+		}
+		return 0, toError(err)
+	}
+	return int(val), nil
+}
+
+func (d *Device) getInfoUlong(param C.cl_device_info, panicOnError bool) (int64, error) {
+	var val C.cl_ulong
+	if err := C.clGetDeviceInfo(d.id, param, C.size_t(unsafe.Sizeof(val)), unsafe.Pointer(&val), nil); err != C.CL_SUCCESS {
+		if panicOnError {
+			panic("Should never fail")
+		}
+		return 0, toError(err)
+	}
+	return int64(val), nil
+}
+
+func (d *Device) getInfoBool(param C.cl_device_info, panicOnError bool) (bool, error) {
+	var val C.cl_bool
+	if err := C.clGetDeviceInfo(d.id, param, C.size_t(unsafe.Sizeof(val)), unsafe.Pointer(&val), nil); err != C.CL_SUCCESS {
+		if panicOnError {
+			panic("Should never fail")
+		}
+		return false, toError(err)
+	}
+	return val == C.CL_TRUE, nil
+}
+
+func (d *Device) Name() string {
+	str, _ := d.GetInfoString(C.CL_DEVICE_NAME, true)
+	return str
+}
+
+func (d *Device) Vendor() string {
+	str, _ := d.GetInfoString(C.CL_DEVICE_VENDOR, true)
+	return str
+}
+
+func (d *Device) Extensions() string {
+	str, _ := d.GetInfoString(C.CL_DEVICE_EXTENSIONS, true)
+	return str
+}
+
+func (d *Device) OpenCLCVersion() string {
+	str, _ := d.GetInfoString(C.CL_DEVICE_OPENCL_C_VERSION, true)
+	return str
+}
+
+func (d *Device) Profile() string {
+	str, _ := d.GetInfoString(C.CL_DEVICE_PROFILE, true)
+	return str
+}
+
+func (d *Device) Version() string {
+	str, _ := d.GetInfoString(C.CL_DEVICE_VERSION, true)
+	return str
+}
+
+func (d *Device) DriverVersion() string {
+	str, _ := d.GetInfoString(C.CL_DRIVER_VERSION, true)
+	return str
+}
+
+// The default compute device address space size specified as an
+// unsigned integer value in bits. Currently supported values are 32 or 64 bits.
+func (d *Device) AddressBits() int {
+	val, _ := d.getInfoUint(C.CL_DEVICE_ADDRESS_BITS, true)
+	return int(val)
+}
+
+// Size of global memory cache line in bytes.
+func (d *Device) GlobalMemCachelineSize() int {
+	val, _ := d.getInfoUint(C.CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE, true)
+	return int(val)
+}
+
+// Maximum configured clock frequency of the device in MHz.
+func (d *Device) MaxClockFrequency() int {
+	val, _ := d.getInfoUint(C.CL_DEVICE_MAX_CLOCK_FREQUENCY, true)
+	return int(val)
+}
+
+// The number of parallel compute units on the OpenCL device.
+// A work-group executes on a single compute unit. The minimum value is 1.
+func (d *Device) MaxComputeUnits() int {
+	val, _ := d.getInfoUint(C.CL_DEVICE_MAX_COMPUTE_UNITS, true)
+	return int(val)
+}
+
+// Max number of arguments declared with the __constant qualifier in a kernel.
+// The minimum value is 8 for devices that are not of type CL_DEVICE_TYPE_CUSTOM.
+func (d *Device) MaxConstantArgs() int {
+	val, _ := d.getInfoUint(C.CL_DEVICE_MAX_CONSTANT_ARGS, true)
+	return int(val)
+}
+
+// Max number of simultaneous image objects that can be read by a kernel.
+// The minimum value is 128 if CL_DEVICE_IMAGE_SUPPORT is CL_TRUE.
+func (d *Device) MaxReadImageArgs() int {
+	val, _ := d.getInfoUint(C.CL_DEVICE_MAX_READ_IMAGE_ARGS, true)
+	return int(val)
+}
+
+// Maximum number of samplers that can be used in a kernel. The minimum
+// value is 16 if CL_DEVICE_IMAGE_SUPPORT is CL_TRUE. (Also see sampler_t.)
+func (d *Device) MaxSamplers() int {
+	val, _ := d.getInfoUint(C.CL_DEVICE_MAX_SAMPLERS, true)
+	return int(val)
+}
+
+// Maximum dimensions that specify the global and local work-item IDs used
+// by the data parallel execution model. (Refer to clEnqueueNDRangeKernel).
+// The minimum value is 3 for devices that are not of type CL_DEVICE_TYPE_CUSTOM.
+func (d *Device) MaxWorkItemDimensions() int {
+	val, _ := d.getInfoUint(C.CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, true)
+	return int(val)
+}
+
+// Max number of simultaneous image objects that can be written to by a
+// kernel. The minimum value is 8 if CL_DEVICE_IMAGE_SUPPORT is CL_TRUE.
+func (d *Device) MaxWriteImageArgs() int {
+	val, _ := d.getInfoUint(C.CL_DEVICE_MAX_WRITE_IMAGE_ARGS, true)
+	return int(val)
+}
+
+// The minimum value is the size (in bits) of the largest OpenCL built-in
+// data type supported by the device (long16 in FULL profile, long16 or
+// int16 in EMBEDDED profile) for devices that are not of type CL_DEVICE_TYPE_CUSTOM.
+func (d *Device) MemBaseAddrAlign() int {
+	val, _ := d.getInfoUint(C.CL_DEVICE_MEM_BASE_ADDR_ALIGN, true)
+	return int(val)
+}
+
+func (d *Device) NativeVectorWidthChar() int {
+	val, _ := d.getInfoUint(C.CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR, true)
+	return int(val)
+}
+
+func (d *Device) NativeVectorWidthShort() int {
+	val, _ := d.getInfoUint(C.CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT, true)
+	return int(val)
+}
+
+func (d *Device) NativeVectorWidthInt() int {
+	val, _ := d.getInfoUint(C.CL_DEVICE_NATIVE_VECTOR_WIDTH_INT, true)
+	return int(val)
+}
+
+func (d *Device) NativeVectorWidthLong() int {
+	val, _ := d.getInfoUint(C.CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG, true)
+	return int(val)
+}
+
+func (d *Device) NativeVectorWidthFloat() int {
+	val, _ := d.getInfoUint(C.CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT, true)
+	return int(val)
+}
+
+func (d *Device) NativeVectorWidthDouble() int {
+	val, _ := d.getInfoUint(C.CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE, true)
+	return int(val)
+}
+
+func (d *Device) NativeVectorWidthHalf() int {
+	val, _ := d.getInfoUint(C.CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF, true)
+	return int(val)
+}
+
+// Max height of 2D image in pixels. The minimum value is 8192
+// if CL_DEVICE_IMAGE_SUPPORT is CL_TRUE.
+func (d *Device) Image2DMaxHeight() int {
+	val, _ := d.getInfoSize(C.CL_DEVICE_IMAGE2D_MAX_HEIGHT, true)
+	return int(val)
+}
+
+// Max width of 2D image or 1D image not created from a buffer object in
+// pixels. The minimum value is 8192 if CL_DEVICE_IMAGE_SUPPORT is CL_TRUE.
+func (d *Device) Image2DMaxWidth() int {
+	val, _ := d.getInfoSize(C.CL_DEVICE_IMAGE2D_MAX_WIDTH, true)
+	return int(val)
+}
+
+// Max depth of 3D image in pixels. The minimum value is 2048 if CL_DEVICE_IMAGE_SUPPORT is CL_TRUE.
+func (d *Device) Image3DMaxDepth() int {
+	val, _ := d.getInfoSize(C.CL_DEVICE_IMAGE3D_MAX_DEPTH, true)
+	return int(val)
+}
+
+// Max height of 3D image in pixels. The minimum value is 2048 if CL_DEVICE_IMAGE_SUPPORT is CL_TRUE.
+func (d *Device) Image3DMaxHeight() int {
+	val, _ := d.getInfoSize(C.CL_DEVICE_IMAGE3D_MAX_HEIGHT, true)
+	return int(val)
+}
+
+// Max width of 3D image in pixels. The minimum value is 2048 if CL_DEVICE_IMAGE_SUPPORT is CL_TRUE.
+func (d *Device) Image3DMaxWidth() int {
+	val, _ := d.getInfoSize(C.CL_DEVICE_IMAGE3D_MAX_WIDTH, true)
+	return int(val)
+}
+
+// Max size in bytes of the arguments that can be passed to a kernel. The
+// minimum value is 1024 for devices that are not of type CL_DEVICE_TYPE_CUSTOM.
+// For this minimum value, only a maximum of 128 arguments can be passed to a kernel.
+func (d *Device) MaxParameterSize() int {
+	val, _ := d.getInfoSize(C.CL_DEVICE_MAX_PARAMETER_SIZE, true)
+	return int(val)
+}
+
+// Maximum number of work-items in a work-group executing a kernel on a
+// single compute unit, using the data parallel execution model. (Refer
+// to clEnqueueNDRangeKernel). The minimum value is 1.
+func (d *Device) MaxWorkGroupSize() int {
+	val, _ := d.getInfoSize(C.CL_DEVICE_MAX_WORK_GROUP_SIZE, true)
+	return int(val)
+}
+
+// Describes the resolution of device timer. This is measured in nanoseconds.
+func (d *Device) ProfilingTimerResolution() int {
+	val, _ := d.getInfoSize(C.CL_DEVICE_PROFILING_TIMER_RESOLUTION, true)
+	return int(val)
+}
+
+// Size of local memory arena in bytes. The minimum value is 32 KB for
+// devices that are not of type CL_DEVICE_TYPE_CUSTOM.
+func (d *Device) LocalMemSize() int64 {
+	val, _ := d.getInfoUlong(C.CL_DEVICE_LOCAL_MEM_SIZE, true)
+	return val
+}
+
+// Max size in bytes of a constant buffer allocation. The minimum value is
+// 64 KB for devices that are not of type CL_DEVICE_TYPE_CUSTOM.
+func (d *Device) MaxConstantBufferSize() int64 {
+	val, _ := d.getInfoUlong(C.CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, true)
+	return val
+}
+
+// Max size of memory object allocation in bytes. The minimum value is max
+// (1/4th of CL_DEVICE_GLOBAL_MEM_SIZE, 128*1024*1024) for devices that are
+// not of type CL_DEVICE_TYPE_CUSTOM.
+func (d *Device) MaxMemAllocSize() int64 {
+	val, _ := d.getInfoUlong(C.CL_DEVICE_MAX_MEM_ALLOC_SIZE, true)
+	return val
+}
+
+// Size of global device memory in bytes.
+func (d *Device) GlobalMemSize() int64 {
+	val, _ := d.getInfoUlong(C.CL_DEVICE_GLOBAL_MEM_SIZE, true)
+	return val
+}
+
+func (d *Device) Available() bool {
+	val, _ := d.getInfoBool(C.CL_DEVICE_AVAILABLE, true)
+	return val
+}
+
+func (d *Device) CompilerAvailable() bool {
+	val, _ := d.getInfoBool(C.CL_DEVICE_COMPILER_AVAILABLE, true)
+	return val
+}
+
+func (d *Device) EndianLittle() bool {
+	val, _ := d.getInfoBool(C.CL_DEVICE_ENDIAN_LITTLE, true)
+	return val
+}
+
+// Is CL_TRUE if the device implements error correction for all
+// accesses to compute device memory (global and constant). Is
+// CL_FALSE if the device does not implement such error correction.
+func (d *Device) ErrorCorrectionSupport() bool {
+	val, _ := d.getInfoBool(C.CL_DEVICE_ERROR_CORRECTION_SUPPORT, true)
+	return val
+}
+
+func (d *Device) HostUnifiedMemory() bool {
+	val, _ := d.getInfoBool(C.CL_DEVICE_HOST_UNIFIED_MEMORY, true)
+	return val
+}
+
+func (d *Device) ImageSupport() bool {
+	val, _ := d.getInfoBool(C.CL_DEVICE_IMAGE_SUPPORT, true)
+	return val
+}
+
+func (d *Device) Type() DeviceType {
+	var deviceType C.cl_device_type
+	if err := C.clGetDeviceInfo(d.id, C.CL_DEVICE_TYPE, C.size_t(unsafe.Sizeof(deviceType)), unsafe.Pointer(&deviceType), nil); err != C.CL_SUCCESS {
+		panic("Failed to get device type")
+	}
+	return DeviceType(deviceType)
+}
+
+// Describes double precision floating-point capability of the OpenCL device
+func (d *Device) DoubleFPConfig() FPConfig {
+	var fpConfig C.cl_device_fp_config
+	if err := C.clGetDeviceInfo(d.id, C.CL_DEVICE_DOUBLE_FP_CONFIG, C.size_t(unsafe.Sizeof(fpConfig)), unsafe.Pointer(&fpConfig), nil); err != C.CL_SUCCESS {
+		panic("Failed to get double FP config")
+	}
+	return FPConfig(fpConfig)
+}
+
+// Describes the OPTIONAL half precision floating-point capability of the OpenCL device
+func (d *Device) HalfFPConfig() FPConfig {
+	var fpConfig C.cl_device_fp_config
+	err := C.clGetDeviceInfo(d.id, C.CL_DEVICE_HALF_FP_CONFIG, C.size_t(unsafe.Sizeof(fpConfig)), unsafe.Pointer(&fpConfig), nil)
+	if err != C.CL_SUCCESS {
+		return FPConfig(0)
+	}
+	return FPConfig(fpConfig)
+}
+
+// Type of local memory supported. This can be set to CL_LOCAL implying dedicated
+// local memory storage such as SRAM, or CL_GLOBAL. For custom devices, CL_NONE
+// can also be returned indicating no local memory support.
+func (d *Device) LocalMemType() LocalMemType {
+	var memType C.cl_device_local_mem_type
+	if err := C.clGetDeviceInfo(d.id, C.CL_DEVICE_LOCAL_MEM_TYPE, C.size_t(unsafe.Sizeof(memType)), unsafe.Pointer(&memType), nil); err != C.CL_SUCCESS {
+		return LocalMemType(C.CL_NONE)
+	}
+	return LocalMemType(memType)
+}
+
+// Describes the execution capabilities of the device. The mandated minimum capability is CL_EXEC_KERNEL.
+func (d *Device) ExecutionCapabilities() ExecCapability {
+	var execCap C.cl_device_exec_capabilities
+	if err := C.clGetDeviceInfo(d.id, C.CL_DEVICE_EXECUTION_CAPABILITIES, C.size_t(unsafe.Sizeof(execCap)), unsafe.Pointer(&execCap), nil); err != C.CL_SUCCESS {
+		panic("Failed to get execution capabilities")
+	}
+	return ExecCapability(execCap)
+}
+
+func (d *Device) GlobalMemCacheType() MemCacheType {
+	var memType C.cl_device_mem_cache_type
+	if err := C.clGetDeviceInfo(d.id, C.CL_DEVICE_GLOBAL_MEM_CACHE_TYPE, C.size_t(unsafe.Sizeof(memType)), unsafe.Pointer(&memType), nil); err != C.CL_SUCCESS {
+		return MemCacheType(C.CL_NONE)
+	}
+	return MemCacheType(memType)
+}
+
+// Maximum number of work-items that can be specified in each dimension of the work-group to clEnqueueNDRangeKernel.
+//
+// Returns n size_t entries, where n is the value returned by the query for CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS.
+//
+// The minimum value is (1, 1, 1) for devices that are not of type CL_DEVICE_TYPE_CUSTOM.
+func (d *Device) MaxWorkItemSizes() []int {
+	dims := d.MaxWorkItemDimensions()
+	sizes := make([]C.size_t, dims)
+	if err := C.clGetDeviceInfo(d.id, C.CL_DEVICE_MAX_WORK_ITEM_SIZES, C.size_t(int(unsafe.Sizeof(sizes[0]))*dims), unsafe.Pointer(&sizes[0]), nil); err != C.CL_SUCCESS {
+		panic("Failed to get max work item sizes")
+	}
+	intSizes := make([]int, dims)
+	for i, s := range sizes {
+		intSizes[i] = int(s)
+	}
+	return intSizes
+}
--- a/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/device12.go
+++ b/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/device12.go
@@ -0,0 +1,51 @@
+// +build cl12
+
+package cl
+
+// #ifdef __APPLE__
+// #include "OpenCL/opencl.h"
+// #else
+// #include "cl.h"
+// #endif
+import "C"
+import "unsafe"
+
+const FPConfigCorrectlyRoundedDivideSqrt FPConfig = C.CL_FP_CORRECTLY_ROUNDED_DIVIDE_SQRT
+
+func init() {
+	fpConfigNameMap[FPConfigCorrectlyRoundedDivideSqrt] = "CorrectlyRoundedDivideSqrt"
+}
+
+func (d *Device) BuiltInKernels() string {
+	str, _ := d.getInfoString(C.CL_DEVICE_BUILT_IN_KERNELS, true)
+	return str
+}
+
+// Is CL_FALSE if the implementation does not have a linker available. Is CL_TRUE if the linker is available. This can be CL_FALSE for the embedded platform profile only. This must be CL_TRUE if CL_DEVICE_COMPILER_AVAILABLE is CL_TRUE
+func (d *Device) LinkerAvailable() bool {
+	val, _ := d.getInfoBool(C.CL_DEVICE_LINKER_AVAILABLE, true)
+	return val
+}
+
+func (d *Device) ParentDevice() *Device {
+	var deviceId C.cl_device_id
+	if err := C.clGetDeviceInfo(d.id, C.CL_DEVICE_PARENT_DEVICE, C.size_t(unsafe.Sizeof(deviceId)), unsafe.Pointer(&deviceId), nil); err != C.CL_SUCCESS {
+		panic("ParentDevice failed")
+	}
+	if deviceId == nil {
+		return nil
+	}
+	return &Device{id: deviceId}
+}
+
+// Max number of pixels for a 1D image created from a buffer object. The minimum value is 65536 if CL_DEVICE_IMAGE_SUPPORT is CL_TRUE.
+func (d *Device) ImageMaxBufferSize() int {
+	val, _ := d.getInfoSize(C.CL_DEVICE_IMAGE_MAX_BUFFER_SIZE, true)
+	return int(val)
+}
+
+// Max number of images in a 1D or 2D image array. The minimum value is 2048 if CL_DEVICE_IMAGE_SUPPORT is CL_TRUE
+func (d *Device) ImageMaxArraySize() int {
+	val, _ := d.getInfoSize(C.CL_DEVICE_IMAGE_MAX_ARRAY_SIZE, true)
+	return int(val)
+}
--- a/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/headers/1.2/cl.h
+++ b/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/headers/1.2/cl.h
--- a/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/headers/1.2/cl_ext.h
+++ b/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/headers/1.2/cl_ext.h
@@ -0,0 +1,315 @@
+/*******************************************************************************
+ * Copyright (c) 2008-2013 The Khronos Group Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and/or associated documentation files (the
+ * "Materials"), to deal in the Materials without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Materials, and to
+ * permit persons to whom the Materials are furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Materials.
+ *
+ * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+ * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+ ******************************************************************************/
+
+/* $Revision: 11928 $ on $Date: 2010-07-13 09:04:56 -0700 (Tue, 13 Jul 2010) $ */
+
+/* cl_ext.h contains OpenCL extensions which don't have external */
+/* (OpenGL, D3D) dependencies.                                   */
+
+#ifndef __CL_EXT_H
+#define __CL_EXT_H
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+#ifdef __APPLE__
+#include <AvailabilityMacros.h>
+#endif
+
+#include <cl.h>
+
+/* cl_khr_fp16 extension - no extension #define since it has no functions  */
+#define CL_DEVICE_HALF_FP_CONFIG                    0x1033
+
+/* Memory object destruction
+ *
+ * Apple extension for use to manage externally allocated buffers used with cl_mem objects with CL_MEM_USE_HOST_PTR
+ *
+ * Registers a user callback function that will be called when the memory object is deleted and its resources 
+ * freed. Each call to clSetMemObjectCallbackFn registers the specified user callback function on a callback 
+ * stack associated with memobj. The registered user callback functions are called in the reverse order in 
+ * which they were registered. The user callback functions are called and then the memory object is deleted 
+ * and its resources freed. This provides a mechanism for the application (and libraries) using memobj to be 
+ * notified when the memory referenced by host_ptr, specified when the memory object is created and used as 
+ * the storage bits for the memory object, can be reused or freed.
+ *
+ * The application may not call CL api's with the cl_mem object passed to the pfn_notify.
+ *
+ * Please check for the "cl_APPLE_SetMemObjectDestructor" extension using clGetDeviceInfo(CL_DEVICE_EXTENSIONS)
+ * before using.
+ */
+#define cl_APPLE_SetMemObjectDestructor 1
+cl_int  CL_API_ENTRY clSetMemObjectDestructorAPPLE(  cl_mem /* memobj */, 
+                                        void (* /*pfn_notify*/)( cl_mem /* memobj */, void* /*user_data*/), 
+                                        void * /*user_data */ )             CL_EXT_SUFFIX__VERSION_1_0;  
+
+
+/* Context Logging Functions
+ *
+ * The next three convenience functions are intended to be used as the pfn_notify parameter to clCreateContext().
+ * Please check for the "cl_APPLE_ContextLoggingFunctions" extension using clGetDeviceInfo(CL_DEVICE_EXTENSIONS)
+ * before using.
+ *
+ * clLogMessagesToSystemLog fowards on all log messages to the Apple System Logger 
+ */
+#define cl_APPLE_ContextLoggingFunctions 1
+extern void CL_API_ENTRY clLogMessagesToSystemLogAPPLE(  const char * /* errstr */, 
+                                            const void * /* private_info */, 
+                                            size_t       /* cb */, 
+                                            void *       /* user_data */ )  CL_EXT_SUFFIX__VERSION_1_0;
+
+/* clLogMessagesToStdout sends all log messages to the file descriptor stdout */
+extern void CL_API_ENTRY clLogMessagesToStdoutAPPLE(   const char * /* errstr */, 
+                                          const void * /* private_info */, 
+                                          size_t       /* cb */, 
+                                          void *       /* user_data */ )    CL_EXT_SUFFIX__VERSION_1_0;
+
+/* clLogMessagesToStderr sends all log messages to the file descriptor stderr */
+extern void CL_API_ENTRY clLogMessagesToStderrAPPLE(   const char * /* errstr */, 
+                                          const void * /* private_info */, 
+                                          size_t       /* cb */, 
+                                          void *       /* user_data */ )    CL_EXT_SUFFIX__VERSION_1_0;
+
+
+/************************ 
+* cl_khr_icd extension *                                                  
+************************/
+#define cl_khr_icd 1
+
+/* cl_platform_info                                                        */
+#define CL_PLATFORM_ICD_SUFFIX_KHR                  0x0920
+
+/* Additional Error Codes                                                  */
+#define CL_PLATFORM_NOT_FOUND_KHR                   -1001
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clIcdGetPlatformIDsKHR(cl_uint          /* num_entries */,
+                       cl_platform_id * /* platforms */,
+                       cl_uint *        /* num_platforms */);
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clIcdGetPlatformIDsKHR_fn)(
+    cl_uint          /* num_entries */,
+    cl_platform_id * /* platforms */,
+    cl_uint *        /* num_platforms */);
+
+
+/* Extension: cl_khr_image2D_buffer
+ *
+ * This extension allows a 2D image to be created from a cl_mem buffer without a copy.
+ * The type associated with a 2D image created from a buffer in an OpenCL program is image2d_t.
+ * Both the sampler and sampler-less read_image built-in functions are supported for 2D images
+ * and 2D images created from a buffer.  Similarly, the write_image built-ins are also supported
+ * for 2D images created from a buffer.
+ *
+ * When the 2D image from buffer is created, the client must specify the width,
+ * height, image format (i.e. channel order and channel data type) and optionally the row pitch
+ *
+ * The pitch specified must be a multiple of CL_DEVICE_IMAGE_PITCH_ALIGNMENT pixels.
+ * The base address of the buffer must be aligned to CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT pixels.
+ */
+    
+/*************************************
+ * cl_khr_initalize_memory extension *
+ *************************************/
+    
+#define CL_CONTEXT_MEMORY_INITIALIZE_KHR            0x200E
+    
+    
+/**************************************
+ * cl_khr_terminate_context extension *
+ **************************************/
+    
+#define CL_DEVICE_TERMINATE_CAPABILITY_KHR          0x200F
+#define CL_CONTEXT_TERMINATE_KHR                    0x2010
+
+#define cl_khr_terminate_context 1
+extern CL_API_ENTRY cl_int CL_API_CALL clTerminateContextKHR(cl_context /* context */) CL_EXT_SUFFIX__VERSION_1_2;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clTerminateContextKHR_fn)(cl_context /* context */) CL_EXT_SUFFIX__VERSION_1_2;
+    
+    
+/*
+ * Extension: cl_khr_spir
+ *
+ * This extension adds support to create an OpenCL program object from a 
+ * Standard Portable Intermediate Representation (SPIR) instance
+ */
+
+#define CL_DEVICE_SPIR_VERSIONS                     0x40E0
+#define CL_PROGRAM_BINARY_TYPE_INTERMEDIATE         0x40E1
+
+
+/******************************************
+* cl_nv_device_attribute_query extension *
+******************************************/
+/* cl_nv_device_attribute_query extension - no extension #define since it has no functions */
+#define CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV       0x4000
+#define CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV       0x4001
+#define CL_DEVICE_REGISTERS_PER_BLOCK_NV            0x4002
+#define CL_DEVICE_WARP_SIZE_NV                      0x4003
+#define CL_DEVICE_GPU_OVERLAP_NV                    0x4004
+#define CL_DEVICE_KERNEL_EXEC_TIMEOUT_NV            0x4005
+#define CL_DEVICE_INTEGRATED_MEMORY_NV              0x4006
+
+/*********************************
+* cl_amd_device_attribute_query *
+*********************************/
+#define CL_DEVICE_PROFILING_TIMER_OFFSET_AMD        0x4036
+
+/*********************************
+* cl_arm_printf extension
+*********************************/
+#define CL_PRINTF_CALLBACK_ARM                      0x40B0
+#define CL_PRINTF_BUFFERSIZE_ARM                    0x40B1
+
+#ifdef CL_VERSION_1_1
+   /***********************************
+    * cl_ext_device_fission extension *
+    ***********************************/
+    #define cl_ext_device_fission   1
+    
+    extern CL_API_ENTRY cl_int CL_API_CALL
+    clReleaseDeviceEXT( cl_device_id /*device*/ ) CL_EXT_SUFFIX__VERSION_1_1; 
+    
+    typedef CL_API_ENTRY cl_int 
+    (CL_API_CALL *clReleaseDeviceEXT_fn)( cl_device_id /*device*/ ) CL_EXT_SUFFIX__VERSION_1_1;
+
+    extern CL_API_ENTRY cl_int CL_API_CALL
+    clRetainDeviceEXT( cl_device_id /*device*/ ) CL_EXT_SUFFIX__VERSION_1_1; 
+    
+    typedef CL_API_ENTRY cl_int 
+    (CL_API_CALL *clRetainDeviceEXT_fn)( cl_device_id /*device*/ ) CL_EXT_SUFFIX__VERSION_1_1;
+
+    typedef cl_ulong  cl_device_partition_property_ext;
+    extern CL_API_ENTRY cl_int CL_API_CALL
+    clCreateSubDevicesEXT(  cl_device_id /*in_device*/,
+                            const cl_device_partition_property_ext * /* properties */,
+                            cl_uint /*num_entries*/,
+                            cl_device_id * /*out_devices*/,
+                            cl_uint * /*num_devices*/ ) CL_EXT_SUFFIX__VERSION_1_1;
+
+    typedef CL_API_ENTRY cl_int 
+    ( CL_API_CALL * clCreateSubDevicesEXT_fn)(  cl_device_id /*in_device*/,
+                                                const cl_device_partition_property_ext * /* properties */,
+                                                cl_uint /*num_entries*/,
+                                                cl_device_id * /*out_devices*/,
+                                                cl_uint * /*num_devices*/ ) CL_EXT_SUFFIX__VERSION_1_1;
+
+    /* cl_device_partition_property_ext */
+    #define CL_DEVICE_PARTITION_EQUALLY_EXT             0x4050
+    #define CL_DEVICE_PARTITION_BY_COUNTS_EXT           0x4051
+    #define CL_DEVICE_PARTITION_BY_NAMES_EXT            0x4052
+    #define CL_DEVICE_PARTITION_BY_AFFINITY_DOMAIN_EXT  0x4053
+    
+    /* clDeviceGetInfo selectors */
+    #define CL_DEVICE_PARENT_DEVICE_EXT                 0x4054
+    #define CL_DEVICE_PARTITION_TYPES_EXT               0x4055
+    #define CL_DEVICE_AFFINITY_DOMAINS_EXT              0x4056
+    #define CL_DEVICE_REFERENCE_COUNT_EXT               0x4057
+    #define CL_DEVICE_PARTITION_STYLE_EXT               0x4058
+    
+    /* error codes */
+    #define CL_DEVICE_PARTITION_FAILED_EXT              -1057
+    #define CL_INVALID_PARTITION_COUNT_EXT              -1058
+    #define CL_INVALID_PARTITION_NAME_EXT               -1059
+    
+    /* CL_AFFINITY_DOMAINs */
+    #define CL_AFFINITY_DOMAIN_L1_CACHE_EXT             0x1
+    #define CL_AFFINITY_DOMAIN_L2_CACHE_EXT             0x2
+    #define CL_AFFINITY_DOMAIN_L3_CACHE_EXT             0x3
+    #define CL_AFFINITY_DOMAIN_L4_CACHE_EXT             0x4
+    #define CL_AFFINITY_DOMAIN_NUMA_EXT                 0x10
+    #define CL_AFFINITY_DOMAIN_NEXT_FISSIONABLE_EXT     0x100
+    
+    /* cl_device_partition_property_ext list terminators */
+    #define CL_PROPERTIES_LIST_END_EXT                  ((cl_device_partition_property_ext) 0)
+    #define CL_PARTITION_BY_COUNTS_LIST_END_EXT         ((cl_device_partition_property_ext) 0)
+    #define CL_PARTITION_BY_NAMES_LIST_END_EXT          ((cl_device_partition_property_ext) 0 - 1)
+
+/*********************************
+* cl_qcom_ext_host_ptr extension
+*********************************/
+
+#define CL_MEM_EXT_HOST_PTR_QCOM                  (1 << 29)
+
+#define CL_DEVICE_EXT_MEM_PADDING_IN_BYTES_QCOM   0x40A0      
+#define CL_DEVICE_PAGE_SIZE_QCOM                  0x40A1
+#define CL_IMAGE_ROW_ALIGNMENT_QCOM               0x40A2
+#define CL_IMAGE_SLICE_ALIGNMENT_QCOM             0x40A3
+#define CL_MEM_HOST_UNCACHED_QCOM                 0x40A4
+#define CL_MEM_HOST_WRITEBACK_QCOM                0x40A5
+#define CL_MEM_HOST_WRITETHROUGH_QCOM             0x40A6
+#define CL_MEM_HOST_WRITE_COMBINING_QCOM          0x40A7
+
+typedef cl_uint                                   cl_image_pitch_info_qcom;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clGetDeviceImageInfoQCOM(cl_device_id             device,
+                         size_t                   image_width,
+                         size_t                   image_height,
+                         const cl_image_format   *image_format,
+                         cl_image_pitch_info_qcom param_name,
+                         size_t                   param_value_size,
+                         void                    *param_value,
+                         size_t                  *param_value_size_ret);
+
+typedef struct _cl_mem_ext_host_ptr
+{
+    /* Type of external memory allocation. */
+    /* Legal values will be defined in layered extensions. */
+    cl_uint  allocation_type;
+            
+	/* Host cache policy for this external memory allocation. */
+    cl_uint  host_cache_policy;
+
+} cl_mem_ext_host_ptr;
+
+/*********************************
+* cl_qcom_ion_host_ptr extension
+*********************************/
+
+#define CL_MEM_ION_HOST_PTR_QCOM                  0x40A8
+
+typedef struct _cl_mem_ion_host_ptr
+{
+    /* Type of external memory allocation. */
+    /* Must be CL_MEM_ION_HOST_PTR_QCOM for ION allocations. */
+    cl_mem_ext_host_ptr  ext_host_ptr;
+
+    /* ION file descriptor */
+    int                  ion_filedesc;
+            
+    /* Host pointer to the ION allocated memory */
+    void*                ion_hostptr;
+
+} cl_mem_ion_host_ptr;
+
+#endif /* CL_VERSION_1_1 */
+
+#ifdef __cplusplus
+}
+#endif
+
+
+#endif /* __CL_EXT_H */
--- a/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/headers/1.2/cl_gl.h
+++ b/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/headers/1.2/cl_gl.h
@@ -0,0 +1,158 @@
+/**********************************************************************************
+ * Copyright (c) 2008 - 2012 The Khronos Group Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and/or associated documentation files (the
+ * "Materials"), to deal in the Materials without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Materials, and to
+ * permit persons to whom the Materials are furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Materials.
+ *
+ * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+ * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+ **********************************************************************************/
+
+#ifndef __OPENCL_CL_GL_H
+#define __OPENCL_CL_GL_H
+
+#include <cl.h>
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+typedef cl_uint     cl_gl_object_type;
+typedef cl_uint     cl_gl_texture_info;
+typedef cl_uint     cl_gl_platform_info;
+typedef struct __GLsync *cl_GLsync;
+
+/* cl_gl_object_type = 0x2000 - 0x200F enum values are currently taken           */
+#define CL_GL_OBJECT_BUFFER                     0x2000
+#define CL_GL_OBJECT_TEXTURE2D                  0x2001
+#define CL_GL_OBJECT_TEXTURE3D                  0x2002
+#define CL_GL_OBJECT_RENDERBUFFER               0x2003
+#define CL_GL_OBJECT_TEXTURE2D_ARRAY            0x200E
+#define CL_GL_OBJECT_TEXTURE1D                  0x200F
+#define CL_GL_OBJECT_TEXTURE1D_ARRAY            0x2010
+#define CL_GL_OBJECT_TEXTURE_BUFFER             0x2011
+
+/* cl_gl_texture_info           */
+#define CL_GL_TEXTURE_TARGET                    0x2004
+#define CL_GL_MIPMAP_LEVEL                      0x2005
+#define CL_GL_NUM_SAMPLES                       0x2012
+
+
+extern CL_API_ENTRY cl_mem CL_API_CALL
+clCreateFromGLBuffer(cl_context     /* context */,
+                     cl_mem_flags   /* flags */,
+                     cl_GLuint      /* bufobj */,
+                     int *          /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_mem CL_API_CALL
+clCreateFromGLTexture(cl_context      /* context */,
+                      cl_mem_flags    /* flags */,
+                      cl_GLenum       /* target */,
+                      cl_GLint        /* miplevel */,
+                      cl_GLuint       /* texture */,
+                      cl_int *        /* errcode_ret */) CL_API_SUFFIX__VERSION_1_2;
+    
+extern CL_API_ENTRY cl_mem CL_API_CALL
+clCreateFromGLRenderbuffer(cl_context   /* context */,
+                           cl_mem_flags /* flags */,
+                           cl_GLuint    /* renderbuffer */,
+                           cl_int *     /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clGetGLObjectInfo(cl_mem                /* memobj */,
+                  cl_gl_object_type *   /* gl_object_type */,
+                  cl_GLuint *           /* gl_object_name */) CL_API_SUFFIX__VERSION_1_0;
+                  
+extern CL_API_ENTRY cl_int CL_API_CALL
+clGetGLTextureInfo(cl_mem               /* memobj */,
+                   cl_gl_texture_info   /* param_name */,
+                   size_t               /* param_value_size */,
+                   void *               /* param_value */,
+                   size_t *             /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clEnqueueAcquireGLObjects(cl_command_queue      /* command_queue */,
+                          cl_uint               /* num_objects */,
+                          const cl_mem *        /* mem_objects */,
+                          cl_uint               /* num_events_in_wait_list */,
+                          const cl_event *      /* event_wait_list */,
+                          cl_event *            /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clEnqueueReleaseGLObjects(cl_command_queue      /* command_queue */,
+                          cl_uint               /* num_objects */,
+                          const cl_mem *        /* mem_objects */,
+                          cl_uint               /* num_events_in_wait_list */,
+                          const cl_event *      /* event_wait_list */,
+                          cl_event *            /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+
+/* Deprecated OpenCL 1.1 APIs */
+extern CL_API_ENTRY CL_EXT_PREFIX__VERSION_1_1_DEPRECATED cl_mem CL_API_CALL
+clCreateFromGLTexture2D(cl_context      /* context */,
+                        cl_mem_flags    /* flags */,
+                        cl_GLenum       /* target */,
+                        cl_GLint        /* miplevel */,
+                        cl_GLuint       /* texture */,
+                        cl_int *        /* errcode_ret */) CL_EXT_SUFFIX__VERSION_1_1_DEPRECATED;
+    
+extern CL_API_ENTRY CL_EXT_PREFIX__VERSION_1_1_DEPRECATED cl_mem CL_API_CALL
+clCreateFromGLTexture3D(cl_context      /* context */,
+                        cl_mem_flags    /* flags */,
+                        cl_GLenum       /* target */,
+                        cl_GLint        /* miplevel */,
+                        cl_GLuint       /* texture */,
+                        cl_int *        /* errcode_ret */) CL_EXT_SUFFIX__VERSION_1_1_DEPRECATED;
+    
+/* cl_khr_gl_sharing extension  */
+    
+#define cl_khr_gl_sharing 1
+    
+typedef cl_uint     cl_gl_context_info;
+    
+/* Additional Error Codes  */
+#define CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR  -1000
+    
+/* cl_gl_context_info  */
+#define CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR    0x2006
+#define CL_DEVICES_FOR_GL_CONTEXT_KHR           0x2007
+    
+/* Additional cl_context_properties  */
+#define CL_GL_CONTEXT_KHR                       0x2008
+#define CL_EGL_DISPLAY_KHR                      0x2009
+#define CL_GLX_DISPLAY_KHR                      0x200A
+#define CL_WGL_HDC_KHR                          0x200B
+#define CL_CGL_SHAREGROUP_KHR                   0x200C
+    
+extern CL_API_ENTRY cl_int CL_API_CALL
+clGetGLContextInfoKHR(const cl_context_properties * /* properties */,
+                      cl_gl_context_info            /* param_name */,
+                      size_t                        /* param_value_size */,
+                      void *                        /* param_value */,
+                      size_t *                      /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+    
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clGetGLContextInfoKHR_fn)(
+    const cl_context_properties * properties,
+    cl_gl_context_info            param_name,
+    size_t                        param_value_size,
+    void *                        param_value,
+    size_t *                      param_value_size_ret);
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif  /* __OPENCL_CL_GL_H */
--- a/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/headers/1.2/cl_gl_ext.h
+++ b/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/headers/1.2/cl_gl_ext.h
@@ -0,0 +1,65 @@
+/**********************************************************************************
+ * Copyright (c) 2008-2012 The Khronos Group Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and/or associated documentation files (the
+ * "Materials"), to deal in the Materials without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Materials, and to
+ * permit persons to whom the Materials are furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Materials.
+ *
+ * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+ * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+ **********************************************************************************/
+
+/* $Revision: 11708 $ on $Date: 2010-06-13 23:36:24 -0700 (Sun, 13 Jun 2010) $ */
+
+/* cl_gl_ext.h contains vendor (non-KHR) OpenCL extensions which have           */
+/* OpenGL dependencies.                                                         */
+
+#ifndef __OPENCL_CL_GL_EXT_H
+#define __OPENCL_CL_GL_EXT_H
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+#include <cl_gl.h>
+
+/*
+ * For each extension, follow this template
+ *  cl_VEN_extname extension  */
+/* #define cl_VEN_extname 1
+ * ... define new types, if any
+ * ... define new tokens, if any
+ * ... define new APIs, if any
+ *
+ *  If you need GLtypes here, mirror them with a cl_GLtype, rather than including a GL header
+ *  This allows us to avoid having to decide whether to include GL headers or GLES here.
+ */
+
+/* 
+ *  cl_khr_gl_event  extension
+ *  See section 9.9 in the OpenCL 1.1 spec for more information
+ */
+#define CL_COMMAND_GL_FENCE_SYNC_OBJECT_KHR     0x200D
+
+extern CL_API_ENTRY cl_event CL_API_CALL
+clCreateEventFromGLsyncKHR(cl_context           /* context */,
+                           cl_GLsync            /* cl_GLsync */,
+                           cl_int *             /* errcode_ret */) CL_EXT_SUFFIX__VERSION_1_1;
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif	/* __OPENCL_CL_GL_EXT_H  */
--- a/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/headers/1.2/cl_platform.h
+++ b/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/headers/1.2/cl_platform.h
--- a/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/headers/1.2/opencl.h
+++ b/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/headers/1.2/opencl.h
@@ -0,0 +1,43 @@
+/*******************************************************************************
+ * Copyright (c) 2008-2012 The Khronos Group Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and/or associated documentation files (the
+ * "Materials"), to deal in the Materials without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Materials, and to
+ * permit persons to whom the Materials are furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Materials.
+ *
+ * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+ * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+ ******************************************************************************/
+
+/* $Revision: 11708 $ on $Date: 2010-06-13 23:36:24 -0700 (Sun, 13 Jun 2010) $ */
+
+#ifndef __OPENCL_H
+#define __OPENCL_H
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+#include <cl.h>
+#include <cl_gl.h>
+#include <cl_gl_ext.h>
+#include <cl_ext.h>
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif  /* __OPENCL_H   */
+
--- a/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/image.go
+++ b/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/image.go
@@ -0,0 +1,83 @@
+// +build cl12
+
+package cl
+
+// #ifdef __APPLE__
+// #include "OpenCL/opencl.h"
+// #else
+// #include "cl.h"
+// #endif
+import "C"
+import (
+	"image"
+	"unsafe"
+)
+
+func (ctx *Context) CreateImage(flags MemFlag, imageFormat ImageFormat, imageDesc ImageDescription, data []byte) (*MemObject, error) {
+	format := imageFormat.toCl()
+	desc := imageDesc.toCl()
+	var dataPtr unsafe.Pointer
+	if data != nil {
+		dataPtr = unsafe.Pointer(&data[0])
+	}
+	var err C.cl_int
+	clBuffer := C.clCreateImage(ctx.clContext, C.cl_mem_flags(flags), &format, &desc, dataPtr, &err)
+	if err != C.CL_SUCCESS {
+		return nil, toError(err)
+	}
+	if clBuffer == nil {
+		return nil, ErrUnknown
+	}
+	return newMemObject(clBuffer, len(data)), nil
+}
+
+func (ctx *Context) CreateImageSimple(flags MemFlag, width, height int, channelOrder ChannelOrder, channelDataType ChannelDataType, data []byte) (*MemObject, error) {
+	format := ImageFormat{channelOrder, channelDataType}
+	desc := ImageDescription{
+		Type:   MemObjectTypeImage2D,
+		Width:  width,
+		Height: height,
+	}
+	return ctx.CreateImage(flags, format, desc, data)
+}
+
+func (ctx *Context) CreateImageFromImage(flags MemFlag, img image.Image) (*MemObject, error) {
+	switch m := img.(type) {
+	case *image.Gray:
+		format := ImageFormat{ChannelOrderIntensity, ChannelDataTypeUNormInt8}
+		desc := ImageDescription{
+			Type:     MemObjectTypeImage2D,
+			Width:    m.Bounds().Dx(),
+			Height:   m.Bounds().Dy(),
+			RowPitch: m.Stride,
+		}
+		return ctx.CreateImage(flags, format, desc, m.Pix)
+	case *image.RGBA:
+		format := ImageFormat{ChannelOrderRGBA, ChannelDataTypeUNormInt8}
+		desc := ImageDescription{
+			Type:     MemObjectTypeImage2D,
+			Width:    m.Bounds().Dx(),
+			Height:   m.Bounds().Dy(),
+			RowPitch: m.Stride,
+		}
+		return ctx.CreateImage(flags, format, desc, m.Pix)
+	}
+
+	b := img.Bounds()
+	w := b.Dx()
+	h := b.Dy()
+	data := make([]byte, w*h*4)
+	dataOffset := 0
+	for y := 0; y < h; y++ {
+		for x := 0; x < w; x++ {
+			c := img.At(x+b.Min.X, y+b.Min.Y)
+			r, g, b, a := c.RGBA()
+			data[dataOffset] = uint8(r >> 8)
+			data[dataOffset+1] = uint8(g >> 8)
+			data[dataOffset+2] = uint8(b >> 8)
+			data[dataOffset+3] = uint8(a >> 8)
+			dataOffset += 4
+		}
+	}
+	return ctx.CreateImageSimple(flags, w, h, ChannelOrderRGBA, ChannelDataTypeUNormInt8, data)
+}
--- a/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/kernel.go
+++ b/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/kernel.go
@@ -0,0 +1,127 @@
+package cl
+
+// #ifdef __APPLE__
+// #include "OpenCL/opencl.h"
+// #else
+// #include "cl.h"
+// #endif
+import "C"
+
+import (
+	"fmt"
+	"unsafe"
+)
+
+type ErrUnsupportedArgumentType struct {
+	Index int
+	Value interface{}
+}
+
+func (e ErrUnsupportedArgumentType) Error() string {
+	return fmt.Sprintf("cl: unsupported argument type for index %d: %+v", e.Index, e.Value)
+}
+
+type Kernel struct {
+	clKernel C.cl_kernel
+	name     string
+}
+
+type LocalBuffer int
+
+func releaseKernel(k *Kernel) {
+	if k.clKernel != nil {
+		C.clReleaseKernel(k.clKernel)
+		k.clKernel = nil
+	}
+}
+
+func (k *Kernel) Release() {
+	releaseKernel(k)
+}
+
+func (k *Kernel) SetArgs(args ...interface{}) error {
+	for index, arg := range args {
+		if err := k.SetArg(index, arg); err != nil {
+			return err
+		}
+	}
+	return nil
+}
+
+func (k *Kernel) SetArg(index int, arg interface{}) error {
+	switch val := arg.(type) {
+	case uint8:
+		return k.SetArgUint8(index, val)
+	case int8:
+		return k.SetArgInt8(index, val)
+	case uint32:
+		return k.SetArgUint32(index, val)
+	case uint64:
+		return k.SetArgUint64(index, val)
+	case int32:
+		return k.SetArgInt32(index, val)
+	case float32:
+		return k.SetArgFloat32(index, val)
+	case *MemObject:
+		return k.SetArgBuffer(index, val)
+	case LocalBuffer:
+		return k.SetArgLocal(index, int(val))
+	default:
+		return ErrUnsupportedArgumentType{Index: index, Value: arg}
+	}
+}
+
+func (k *Kernel) SetArgBuffer(index int, buffer *MemObject) error {
+	return k.SetArgUnsafe(index, int(unsafe.Sizeof(buffer.clMem)), unsafe.Pointer(&buffer.clMem))
+}
+
+func (k *Kernel) SetArgFloat32(index int, val float32) error {
+	return k.SetArgUnsafe(index, int(unsafe.Sizeof(val)), unsafe.Pointer(&val))
+}
+
+func (k *Kernel) SetArgInt8(index int, val int8) error {
+	return k.SetArgUnsafe(index, int(unsafe.Sizeof(val)), unsafe.Pointer(&val))
+}
+
+func (k *Kernel) SetArgUint8(index int, val uint8) error {
+	return k.SetArgUnsafe(index, int(unsafe.Sizeof(val)), unsafe.Pointer(&val))
+}
+
+func (k *Kernel) SetArgInt32(index int, val int32) error {
+	return k.SetArgUnsafe(index, int(unsafe.Sizeof(val)), unsafe.Pointer(&val))
+}
+
+func (k *Kernel) SetArgUint32(index int, val uint32) error {
+	return k.SetArgUnsafe(index, int(unsafe.Sizeof(val)), unsafe.Pointer(&val))
+}
+
+func (k *Kernel) SetArgUint64(index int, val uint64) error {
+	return k.SetArgUnsafe(index, int(unsafe.Sizeof(val)), unsafe.Pointer(&val))
+}
+
+func (k *Kernel) SetArgLocal(index int, size int) error {
+	return k.SetArgUnsafe(index, size, nil)
+}
+
+func (k *Kernel) SetArgUnsafe(index, argSize int, arg unsafe.Pointer) error {
+	//fmt.Println("FUNKY: ", index, argSize)
+	return toError(C.clSetKernelArg(k.clKernel, C.cl_uint(index), C.size_t(argSize), arg))
+}
+
+func (k *Kernel) PreferredWorkGroupSizeMultiple(device *Device) (int, error) {
+	var size C.size_t
+	err := C.clGetKernelWorkGroupInfo(k.clKernel, device.nullableId(), C.CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, C.size_t(unsafe.Sizeof(size)), unsafe.Pointer(&size), nil)
+	return int(size), toError(err)
+}
+
+func (k *Kernel) WorkGroupSize(device *Device) (int, error) {
+	var size C.size_t
+	err := C.clGetKernelWorkGroupInfo(k.clKernel, device.nullableId(), C.CL_KERNEL_WORK_GROUP_SIZE, C.size_t(unsafe.Sizeof(size)), unsafe.Pointer(&size), nil)
+	return int(size), toError(err)
+}
+
+func (k *Kernel) NumArgs() (int, error) {
+	var num C.cl_uint
+	err := C.clGetKernelInfo(k.clKernel, C.CL_KERNEL_NUM_ARGS, C.size_t(unsafe.Sizeof(num)), unsafe.Pointer(&num), nil)
+	return int(num), toError(err)
+}
--- a/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/kernel10.go
+++ b/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/kernel10.go
@@ -0,0 +1,7 @@
+// +build !cl12
+
+package cl
+
+func (k *Kernel) ArgName(index int) (string, error) {
+	return "", ErrUnsupported
+}
--- a/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/kernel12.go
+++ b/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/kernel12.go
@@ -0,0 +1,20 @@
+// +build cl12
+
+package cl
+
+// #ifdef __APPLE__
+// #include "OpenCL/opencl.h"
+// #else
+// #include "cl.h"
+// #endif
+import "C"
+import "unsafe"
+
+func (k *Kernel) ArgName(index int) (string, error) {
+	var strC [1024]byte
+	var strN C.size_t
+	if err := C.clGetKernelArgInfo(k.clKernel, C.cl_uint(index), C.CL_KERNEL_ARG_NAME, 1024, unsafe.Pointer(&strC[0]), &strN); err != C.CL_SUCCESS {
+		return "", toError(err)
+	}
+	return string(strC[:strN]), nil
+}
--- a/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/platform.go
+++ b/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/platform.go
@@ -0,0 +1,83 @@
+package cl
+
+// #ifdef __APPLE__
+// #include "OpenCL/opencl.h"
+// #else
+// #include "cl.h"
+// #endif
+import "C"
+
+import "unsafe"
+
+const maxPlatforms = 32
+
+type Platform struct {
+	id C.cl_platform_id
+}
+
+// Obtain the list of platforms available.
+func GetPlatforms() ([]*Platform, error) {
+	var platformIds [maxPlatforms]C.cl_platform_id
+	var nPlatforms C.cl_uint
+	if err := C.clGetPlatformIDs(C.cl_uint(maxPlatforms), &platformIds[0], &nPlatforms); err != C.CL_SUCCESS {
+		return nil, toError(err)
+	}
+	platforms := make([]*Platform, nPlatforms)
+	for i := 0; i < int(nPlatforms); i++ {
+		platforms[i] = &Platform{id: platformIds[i]}
+	}
+	return platforms, nil
+}
+
+func (p *Platform) GetDevices(deviceType DeviceType) ([]*Device, error) {
+	return GetDevices(p, deviceType)
+}
+
+func (p *Platform) getInfoString(param C.cl_platform_info) (string, error) {
+	var strC [2048]byte
+	var strN C.size_t
+	if err := C.clGetPlatformInfo(p.id, param, 2048, unsafe.Pointer(&strC[0]), &strN); err != C.CL_SUCCESS {
+		return "", toError(err)
+	}
+	return string(strC[:(strN - 1)]), nil
+}
+
+func (p *Platform) Name() string {
+	if str, err := p.getInfoString(C.CL_PLATFORM_NAME); err != nil {
+		panic("Platform.Name() should never fail")
+	} else {
+		return str
+	}
+}
+
+func (p *Platform) Vendor() string {
+	if str, err := p.getInfoString(C.CL_PLATFORM_VENDOR); err != nil {
+		panic("Platform.Vendor() should never fail")
+	} else {
+		return str
+	}
+}
+
+func (p *Platform) Profile() string {
+	if str, err := p.getInfoString(C.CL_PLATFORM_PROFILE); err != nil {
+		panic("Platform.Profile() should never fail")
+	} else {
+		return str
+	}
+}
+
+func (p *Platform) Version() string {
+	if str, err := p.getInfoString(C.CL_PLATFORM_VERSION); err != nil {
+		panic("Platform.Version() should never fail")
+	} else {
+		return str
+	}
+}
+
+func (p *Platform) Extensions() string {
+	if str, err := p.getInfoString(C.CL_PLATFORM_EXTENSIONS); err != nil {
+		panic("Platform.Extensions() should never fail")
+	} else {
+		return str
+	}
+}
--- a/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/program.go
+++ b/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/program.go
@@ -0,0 +1,105 @@
+package cl
+
+// #include <stdlib.h>
+// #ifdef __APPLE__
+// #include "OpenCL/opencl.h"
+// #else
+// #include "cl.h"
+// #endif
+import "C"
+
+import (
+	"fmt"
+	"runtime"
+	"unsafe"
+)
+
+type BuildError struct {
+	Message string
+	Device  *Device
+}
+
+func (e BuildError) Error() string {
+	if e.Device != nil {
+		return fmt.Sprintf("cl: build error on %q: %s", e.Device.Name(), e.Message)
+	} else {
+		return fmt.Sprintf("cl: build error: %s", e.Message)
+	}
+}
+
+type Program struct {
+	clProgram C.cl_program
+	devices   []*Device
+}
+
+func releaseProgram(p *Program) {
+	if p.clProgram != nil {
+		C.clReleaseProgram(p.clProgram)
+		p.clProgram = nil
+	}
+}
+
+func (p *Program) Release() {
+	releaseProgram(p)
+}
+
+func (p *Program) BuildProgram(devices []*Device, options string) error {
+	var cOptions *C.char
+	if options != "" {
+		cOptions = C.CString(options)
+		defer C.free(unsafe.Pointer(cOptions))
+	}
+	var deviceList []C.cl_device_id
+	var deviceListPtr *C.cl_device_id
+	numDevices := C.cl_uint(len(devices))
+	if devices != nil && len(devices) > 0 {
+		deviceList = buildDeviceIdList(devices)
+		deviceListPtr = &deviceList[0]
+	}
+	if err := C.clBuildProgram(p.clProgram, numDevices, deviceListPtr, cOptions, nil, nil); err != C.CL_SUCCESS {
+		buffer := make([]byte, 4096)
+		var bLen C.size_t
+		var err C.cl_int
+
+		for _, dev := range p.devices {
+			for i := 2; i >= 0; i-- {
+				err = C.clGetProgramBuildInfo(p.clProgram, dev.id, C.CL_PROGRAM_BUILD_LOG, C.size_t(len(buffer)), unsafe.Pointer(&buffer[0]), &bLen)
+				if err == C.CL_INVALID_VALUE && i > 0 && bLen < 1024*1024 {
+					// INVALID_VALUE probably means our buffer isn't large enough
+					buffer = make([]byte, bLen)
+				} else {
+					break
+				}
+			}
+			if err != C.CL_SUCCESS {
+				return toError(err)
+			}
+
+			if bLen > 1 {
+				return BuildError{
+					Device:  dev,
+					Message: string(buffer[:bLen-1]),
+				}
+			}
+		}
+
+		return BuildError{
+			Device:  nil,
+			Message: "build failed and produced no log entries",
+		}
+	}
+	return nil
+}
+
+func (p *Program) CreateKernel(name string) (*Kernel, error) {
+	cName := C.CString(name)
+	defer C.free(unsafe.Pointer(cName))
+	var err C.cl_int
+	clKernel := C.clCreateKernel(p.clProgram, cName, &err)
+	if err != C.CL_SUCCESS {
+		return nil, toError(err)
+	}
+	kernel := &Kernel{clKernel: clKernel, name: name}
+	runtime.SetFinalizer(kernel, releaseKernel)
+	return kernel, nil
+}
--- a/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/queue.go
+++ b/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/queue.go
@@ -0,0 +1,193 @@
+package cl
+
+// #ifdef __APPLE__
+// #include "OpenCL/opencl.h"
+// #else
+// #include "cl.h"
+// #endif
+import "C"
+
+import "unsafe"
+
+type CommandQueueProperty int
+
+const (
+	CommandQueueOutOfOrderExecModeEnable CommandQueueProperty = C.CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE
+	CommandQueueProfilingEnable          CommandQueueProperty = C.CL_QUEUE_PROFILING_ENABLE
+)
+
+type CommandQueue struct {
+	clQueue C.cl_command_queue
+	device  *Device
+}
+
+func releaseCommandQueue(q *CommandQueue) {
+	if q.clQueue != nil {
+		C.clReleaseCommandQueue(q.clQueue)
+		q.clQueue = nil
+	}
+}
+
+// Call clReleaseCommandQueue on the CommandQueue. Using the CommandQueue after Release will cause a panick.
+func (q *CommandQueue) Release() {
+	releaseCommandQueue(q)
+}
+
+// Blocks until all previously queued OpenCL commands in a command-queue are issued to the associated device and have completed.
+func (q *CommandQueue) Finish() error {
+	return toError(C.clFinish(q.clQueue))
+}
+
+// Issues all previously queued OpenCL commands in a command-queue to the device associated with the command-queue.
+func (q *CommandQueue) Flush() error {
+	return toError(C.clFlush(q.clQueue))
+}
+
+// Enqueues a command to map a region of the buffer object given by buffer into the host address space and returns a pointer to this mapped region.
+func (q *CommandQueue) EnqueueMapBuffer(buffer *MemObject, blocking bool, flags MapFlag, offset, size int, eventWaitList []*Event) (*MappedMemObject, *Event, error) {
+	var event C.cl_event
+	var err C.cl_int
+	ptr := C.clEnqueueMapBuffer(q.clQueue, buffer.clMem, clBool(blocking), flags.toCl(), C.size_t(offset), C.size_t(size), C.cl_uint(len(eventWaitList)), eventListPtr(eventWaitList), &event, &err)
+	if err != C.CL_SUCCESS {
+		return nil, nil, toError(err)
+	}
+	ev := newEvent(event)
+	if ptr == nil {
+		return nil, ev, ErrUnknown
+	}
+	return &MappedMemObject{ptr: ptr, size: size}, ev, nil
+}
+
+// Enqueues a command to map a region of an image object into the host address space and returns a pointer to this mapped region.
+func (q *CommandQueue) EnqueueMapImage(buffer *MemObject, blocking bool, flags MapFlag, origin, region [3]int, eventWaitList []*Event) (*MappedMemObject, *Event, error) {
+	cOrigin := sizeT3(origin)
+	cRegion := sizeT3(region)
+	var event C.cl_event
+	var err C.cl_int
+	var rowPitch, slicePitch C.size_t
+	ptr := C.clEnqueueMapImage(q.clQueue, buffer.clMem, clBool(blocking), flags.toCl(), &cOrigin[0], &cRegion[0], &rowPitch, &slicePitch, C.cl_uint(len(eventWaitList)), eventListPtr(eventWaitList), &event, &err)
+	if err != C.CL_SUCCESS {
+		return nil, nil, toError(err)
+	}
+	ev := newEvent(event)
+	if ptr == nil {
+		return nil, ev, ErrUnknown
+	}
+	size := 0 // TODO: could calculate this
+	return &MappedMemObject{ptr: ptr, size: size, rowPitch: int(rowPitch), slicePitch: int(slicePitch)}, ev, nil
+}
+
+// Enqueues a command to unmap a previously mapped region of a memory object.
+func (q *CommandQueue) EnqueueUnmapMemObject(buffer *MemObject, mappedObj *MappedMemObject, eventWaitList []*Event) (*Event, error) {
+	var event C.cl_event
+	if err := C.clEnqueueUnmapMemObject(q.clQueue, buffer.clMem, mappedObj.ptr, C.cl_uint(len(eventWaitList)), eventListPtr(eventWaitList), &event); err != C.CL_SUCCESS {
+		return nil, toError(err)
+	}
+	return newEvent(event), nil
+}
+
+// Enqueues a command to copy a buffer object to another buffer object.
+func (q *CommandQueue) EnqueueCopyBuffer(srcBuffer, dstBuffer *MemObject, srcOffset, dstOffset, byteCount int, eventWaitList []*Event) (*Event, error) {
+	var event C.cl_event
+	err := toError(C.clEnqueueCopyBuffer(q.clQueue, srcBuffer.clMem, dstBuffer.clMem, C.size_t(srcOffset), C.size_t(dstOffset), C.size_t(byteCount), C.cl_uint(len(eventWaitList)), eventListPtr(eventWaitList), &event))
+	return newEvent(event), err
+}
+
+// Enqueue commands to write to a buffer object from host memory.
+func (q *CommandQueue) EnqueueWriteBuffer(buffer *MemObject, blocking bool, offset, dataSize int, dataPtr unsafe.Pointer, eventWaitList []*Event) (*Event, error) {
+	var event C.cl_event
+	err := toError(C.clEnqueueWriteBuffer(q.clQueue, buffer.clMem, clBool(blocking), C.size_t(offset), C.size_t(dataSize), dataPtr, C.cl_uint(len(eventWaitList)), eventListPtr(eventWaitList), &event))
+	return newEvent(event), err
+}
+
+func (q *CommandQueue) EnqueueWriteBufferFloat32(buffer *MemObject, blocking bool, offset int, data []float32, eventWaitList []*Event) (*Event, error) {
+	dataPtr := unsafe.Pointer(&data[0])
+	dataSize := int(unsafe.Sizeof(data[0])) * len(data)
+	return q.EnqueueWriteBuffer(buffer, blocking, offset, dataSize, dataPtr, eventWaitList)
+}
+
+// Enqueue commands to read from a buffer object to host memory.
+func (q *CommandQueue) EnqueueReadBuffer(buffer *MemObject, blocking bool, offset, dataSize int, dataPtr unsafe.Pointer, eventWaitList []*Event) (*Event, error) {
+	var event C.cl_event
+	err := toError(C.clEnqueueReadBuffer(q.clQueue, buffer.clMem, clBool(blocking), C.size_t(offset), C.size_t(dataSize), dataPtr, C.cl_uint(len(eventWaitList)), eventListPtr(eventWaitList), &event))
+	return newEvent(event), err
+}
+
+func (q *CommandQueue) EnqueueReadBufferFloat32(buffer *MemObject, blocking bool, offset int, data []float32, eventWaitList []*Event) (*Event, error) {
+	dataPtr := unsafe.Pointer(&data[0])
+	dataSize := int(unsafe.Sizeof(data[0])) * len(data)
+	return q.EnqueueReadBuffer(buffer, blocking, offset, dataSize, dataPtr, eventWaitList)
+}
+
+// Enqueues a command to execute a kernel on a device.
+func (q *CommandQueue) EnqueueNDRangeKernel(kernel *Kernel, globalWorkOffset, globalWorkSize, localWorkSize []int, eventWaitList []*Event) (*Event, error) {
+	workDim := len(globalWorkSize)
+	var globalWorkOffsetList []C.size_t
+	var globalWorkOffsetPtr *C.size_t
+	if globalWorkOffset != nil {
+		globalWorkOffsetList = make([]C.size_t, len(globalWorkOffset))
+		for i, off := range globalWorkOffset {
+			globalWorkOffsetList[i] = C.size_t(off)
+		}
+		globalWorkOffsetPtr = &globalWorkOffsetList[0]
+	}
+	var globalWorkSizeList []C.size_t
+	var globalWorkSizePtr *C.size_t
+	if globalWorkSize != nil {
+		globalWorkSizeList = make([]C.size_t, len(globalWorkSize))
+		for i, off := range globalWorkSize {
+			globalWorkSizeList[i] = C.size_t(off)
+		}
+		globalWorkSizePtr = &globalWorkSizeList[0]
+	}
+	var localWorkSizeList []C.size_t
+	var localWorkSizePtr *C.size_t
+	if localWorkSize != nil {
+		localWorkSizeList = make([]C.size_t, len(localWorkSize))
+		for i, off := range localWorkSize {
+			localWorkSizeList[i] = C.size_t(off)
+		}
+		localWorkSizePtr = &localWorkSizeList[0]
+	}
+	var event C.cl_event
+	err := toError(C.clEnqueueNDRangeKernel(q.clQueue, kernel.clKernel, C.cl_uint(workDim), globalWorkOffsetPtr, globalWorkSizePtr, localWorkSizePtr, C.cl_uint(len(eventWaitList)), eventListPtr(eventWaitList), &event))
+	return newEvent(event), err
+}
+
+// Enqueues a command to read from a 2D or 3D image object to host memory.
+func (q *CommandQueue) EnqueueReadImage(image *MemObject, blocking bool, origin, region [3]int, rowPitch, slicePitch int, data []byte, eventWaitList []*Event) (*Event, error) {
+	cOrigin := sizeT3(origin)
+	cRegion := sizeT3(region)
+	var event C.cl_event
+	err := toError(C.clEnqueueReadImage(q.clQueue, image.clMem, clBool(blocking), &cOrigin[0], &cRegion[0], C.size_t(rowPitch), C.size_t(slicePitch), unsafe.Pointer(&data[0]), C.cl_uint(len(eventWaitList)), eventListPtr(eventWaitList), &event))
+	return newEvent(event), err
+}
+
+// Enqueues a command to write from a 2D or 3D image object to host memory.
+func (q *CommandQueue) EnqueueWriteImage(image *MemObject, blocking bool, origin, region [3]int, rowPitch, slicePitch int, data []byte, eventWaitList []*Event) (*Event, error) {
+	cOrigin := sizeT3(origin)
+	cRegion := sizeT3(region)
+	var event C.cl_event
+	err := toError(C.clEnqueueWriteImage(q.clQueue, image.clMem, clBool(blocking), &cOrigin[0], &cRegion[0], C.size_t(rowPitch), C.size_t(slicePitch), unsafe.Pointer(&data[0]), C.cl_uint(len(eventWaitList)), eventListPtr(eventWaitList), &event))
+	return newEvent(event), err
+}
+
+func (q *CommandQueue) EnqueueFillBuffer(buffer *MemObject, pattern unsafe.Pointer, patternSize, offset, size int, eventWaitList []*Event) (*Event, error) {
+	var event C.cl_event
+	err := toError(C.clEnqueueFillBuffer(q.clQueue, buffer.clMem, pattern, C.size_t(patternSize), C.size_t(offset), C.size_t(size), C.cl_uint(len(eventWaitList)), eventListPtr(eventWaitList), &event))
+	return newEvent(event), err
+}
+
+// A synchronization point that enqueues a barrier operation.
+func (q *CommandQueue) EnqueueBarrierWithWaitList(eventWaitList []*Event) (*Event, error) {
+	var event C.cl_event
+	err := toError(C.clEnqueueBarrierWithWaitList(q.clQueue, C.cl_uint(len(eventWaitList)), eventListPtr(eventWaitList), &event))
+	return newEvent(event), err
+}
+
+// Enqueues a marker command which waits for either a list of events to complete, or all previously enqueued commands to complete.
+func (q *CommandQueue) EnqueueMarkerWithWaitList(eventWaitList []*Event) (*Event, error) {
+	var event C.cl_event
+	err := toError(C.clEnqueueMarkerWithWaitList(q.clQueue, C.cl_uint(len(eventWaitList)), eventListPtr(eventWaitList), &event))
+	return newEvent(event), err
+}
--- a/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/types.go
+++ b/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/types.go
@@ -0,0 +1,487 @@
+package cl
+
+// #ifdef __APPLE__
+// #include "OpenCL/opencl.h"
+// #else
+// #include "cl.h"
+// #endif
+import "C"
+
+import (
+	"errors"
+	"fmt"
+	"reflect"
+	"runtime"
+	"strings"
+	"unsafe"
+)
+
+var (
+	ErrUnknown = errors.New("cl: unknown error") // Generally an unexpected result from an OpenCL function (e.g. CL_SUCCESS but null pointer)
+)
+
+type ErrOther int
+
+func (e ErrOther) Error() string {
+	return fmt.Sprintf("cl: error %d", int(e))
+}
+
+var (
+	ErrDeviceNotFound                     = errors.New("cl: Device Not Found")
+	ErrDeviceNotAvailable                 = errors.New("cl: Device Not Available")
+	ErrCompilerNotAvailable               = errors.New("cl: Compiler Not Available")
+	ErrMemObjectAllocationFailure         = errors.New("cl: Mem Object Allocation Failure")
+	ErrOutOfResources                     = errors.New("cl: Out Of Resources")
+	ErrOutOfHostMemory                    = errors.New("cl: Out Of Host Memory")
+	ErrProfilingInfoNotAvailable          = errors.New("cl: Profiling Info Not Available")
+	ErrMemCopyOverlap                     = errors.New("cl: Mem Copy Overlap")
+	ErrImageFormatMismatch                = errors.New("cl: Image Format Mismatch")
+	ErrImageFormatNotSupported            = errors.New("cl: Image Format Not Supported")
+	ErrBuildProgramFailure                = errors.New("cl: Build Program Failure")
+	ErrMapFailure                         = errors.New("cl: Map Failure")
+	ErrMisalignedSubBufferOffset          = errors.New("cl: Misaligned Sub Buffer Offset")
+	ErrExecStatusErrorForEventsInWaitList = errors.New("cl: Exec Status Error For Events In Wait List")
+	ErrCompileProgramFailure              = errors.New("cl: Compile Program Failure")
+	ErrLinkerNotAvailable                 = errors.New("cl: Linker Not Available")
+	ErrLinkProgramFailure                 = errors.New("cl: Link Program Failure")
+	ErrDevicePartitionFailed              = errors.New("cl: Device Partition Failed")
+	ErrKernelArgInfoNotAvailable          = errors.New("cl: Kernel Arg Info Not Available")
+	ErrInvalidValue                       = errors.New("cl: Invalid Value")
+	ErrInvalidDeviceType                  = errors.New("cl: Invalid Device Type")
+	ErrInvalidPlatform                    = errors.New("cl: Invalid Platform")
+	ErrInvalidDevice                      = errors.New("cl: Invalid Device")
+	ErrInvalidContext                     = errors.New("cl: Invalid Context")
+	ErrInvalidQueueProperties             = errors.New("cl: Invalid Queue Properties")
+	ErrInvalidCommandQueue                = errors.New("cl: Invalid Command Queue")
+	ErrInvalidHostPtr                     = errors.New("cl: Invalid Host Ptr")
+	ErrInvalidMemObject                   = errors.New("cl: Invalid Mem Object")
+	ErrInvalidImageFormatDescriptor       = errors.New("cl: Invalid Image Format Descriptor")
+	ErrInvalidImageSize                   = errors.New("cl: Invalid Image Size")
+	ErrInvalidSampler                     = errors.New("cl: Invalid Sampler")
+	ErrInvalidBinary                      = errors.New("cl: Invalid Binary")
+	ErrInvalidBuildOptions                = errors.New("cl: Invalid Build Options")
+	ErrInvalidProgram                     = errors.New("cl: Invalid Program")
+	ErrInvalidProgramExecutable           = errors.New("cl: Invalid Program Executable")
+	ErrInvalidKernelName                  = errors.New("cl: Invalid Kernel Name")
+	ErrInvalidKernelDefinition            = errors.New("cl: Invalid Kernel Definition")
+	ErrInvalidKernel                      = errors.New("cl: Invalid Kernel")
+	ErrInvalidArgIndex                    = errors.New("cl: Invalid Arg Index")
+	ErrInvalidArgValue                    = errors.New("cl: Invalid Arg Value")
+	ErrInvalidArgSize                     = errors.New("cl: Invalid Arg Size")
+	ErrInvalidKernelArgs                  = errors.New("cl: Invalid Kernel Args")
+	ErrInvalidWorkDimension               = errors.New("cl: Invalid Work Dimension")
+	ErrInvalidWorkGroupSize               = errors.New("cl: Invalid Work Group Size")
+	ErrInvalidWorkItemSize                = errors.New("cl: Invalid Work Item Size")
+	ErrInvalidGlobalOffset                = errors.New("cl: Invalid Global Offset")
+	ErrInvalidEventWaitList               = errors.New("cl: Invalid Event Wait List")
+	ErrInvalidEvent                       = errors.New("cl: Invalid Event")
+	ErrInvalidOperation                   = errors.New("cl: Invalid Operation")
+	ErrInvalidGlObject                    = errors.New("cl: Invalid Gl Object")
+	ErrInvalidBufferSize                  = errors.New("cl: Invalid Buffer Size")
+	ErrInvalidMipLevel                    = errors.New("cl: Invalid Mip Level")
+	ErrInvalidGlobalWorkSize              = errors.New("cl: Invalid Global Work Size")
+	ErrInvalidProperty                    = errors.New("cl: Invalid Property")
+	ErrInvalidImageDescriptor             = errors.New("cl: Invalid Image Descriptor")
+	ErrInvalidCompilerOptions             = errors.New("cl: Invalid Compiler Options")
+	ErrInvalidLinkerOptions               = errors.New("cl: Invalid Linker Options")
+	ErrInvalidDevicePartitionCount        = errors.New("cl: Invalid Device Partition Count")
+)
+var errorMap = map[C.cl_int]error{
+	C.CL_SUCCESS:                                   nil,
+	C.CL_DEVICE_NOT_FOUND:                          ErrDeviceNotFound,
+	C.CL_DEVICE_NOT_AVAILABLE:                      ErrDeviceNotAvailable,
+	C.CL_COMPILER_NOT_AVAILABLE:                    ErrCompilerNotAvailable,
+	C.CL_MEM_OBJECT_ALLOCATION_FAILURE:             ErrMemObjectAllocationFailure,
+	C.CL_OUT_OF_RESOURCES:                          ErrOutOfResources,
+	C.CL_OUT_OF_HOST_MEMORY:                        ErrOutOfHostMemory,
+	C.CL_PROFILING_INFO_NOT_AVAILABLE:              ErrProfilingInfoNotAvailable,
+	C.CL_MEM_COPY_OVERLAP:                          ErrMemCopyOverlap,
+	C.CL_IMAGE_FORMAT_MISMATCH:                     ErrImageFormatMismatch,
+	C.CL_IMAGE_FORMAT_NOT_SUPPORTED:                ErrImageFormatNotSupported,
+	C.CL_BUILD_PROGRAM_FAILURE:                     ErrBuildProgramFailure,
+	C.CL_MAP_FAILURE:                               ErrMapFailure,
+	C.CL_MISALIGNED_SUB_BUFFER_OFFSET:              ErrMisalignedSubBufferOffset,
+	C.CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST: ErrExecStatusErrorForEventsInWaitList,
+	C.CL_INVALID_VALUE:                             ErrInvalidValue,
+	C.CL_INVALID_DEVICE_TYPE:                       ErrInvalidDeviceType,
+	C.CL_INVALID_PLATFORM:                          ErrInvalidPlatform,
+	C.CL_INVALID_DEVICE:                            ErrInvalidDevice,
+	C.CL_INVALID_CONTEXT:                           ErrInvalidContext,
+	C.CL_INVALID_QUEUE_PROPERTIES:                  ErrInvalidQueueProperties,
+	C.CL_INVALID_COMMAND_QUEUE:                     ErrInvalidCommandQueue,
+	C.CL_INVALID_HOST_PTR:                          ErrInvalidHostPtr,
+	C.CL_INVALID_MEM_OBJECT:                        ErrInvalidMemObject,
+	C.CL_INVALID_IMAGE_FORMAT_DESCRIPTOR:           ErrInvalidImageFormatDescriptor,
+	C.CL_INVALID_IMAGE_SIZE:                        ErrInvalidImageSize,
+	C.CL_INVALID_SAMPLER:                           ErrInvalidSampler,
+	C.CL_INVALID_BINARY:                            ErrInvalidBinary,
+	C.CL_INVALID_BUILD_OPTIONS:                     ErrInvalidBuildOptions,
+	C.CL_INVALID_PROGRAM:                           ErrInvalidProgram,
+	C.CL_INVALID_PROGRAM_EXECUTABLE:                ErrInvalidProgramExecutable,
+	C.CL_INVALID_KERNEL_NAME:                       ErrInvalidKernelName,
+	C.CL_INVALID_KERNEL_DEFINITION:                 ErrInvalidKernelDefinition,
+	C.CL_INVALID_KERNEL:                            ErrInvalidKernel,
+	C.CL_INVALID_ARG_INDEX:                         ErrInvalidArgIndex,
+	C.CL_INVALID_ARG_VALUE:                         ErrInvalidArgValue,
+	C.CL_INVALID_ARG_SIZE:                          ErrInvalidArgSize,
+	C.CL_INVALID_KERNEL_ARGS:                       ErrInvalidKernelArgs,
+	C.CL_INVALID_WORK_DIMENSION:                    ErrInvalidWorkDimension,
+	C.CL_INVALID_WORK_GROUP_SIZE:                   ErrInvalidWorkGroupSize,
+	C.CL_INVALID_WORK_ITEM_SIZE:                    ErrInvalidWorkItemSize,
+	C.CL_INVALID_GLOBAL_OFFSET:                     ErrInvalidGlobalOffset,
+	C.CL_INVALID_EVENT_WAIT_LIST:                   ErrInvalidEventWaitList,
+	C.CL_INVALID_EVENT:                             ErrInvalidEvent,
+	C.CL_INVALID_OPERATION:                         ErrInvalidOperation,
+	C.CL_INVALID_GL_OBJECT:                         ErrInvalidGlObject,
+	C.CL_INVALID_BUFFER_SIZE:                       ErrInvalidBufferSize,
+	C.CL_INVALID_MIP_LEVEL:                         ErrInvalidMipLevel,
+	C.CL_INVALID_GLOBAL_WORK_SIZE:                  ErrInvalidGlobalWorkSize,
+	C.CL_INVALID_PROPERTY:                          ErrInvalidProperty,
+}
+
+func toError(code C.cl_int) error {
+	if err, ok := errorMap[code]; ok {
+		return err
+	}
+	return ErrOther(code)
+}
+
+type LocalMemType int
+
+const (
+	LocalMemTypeNone   LocalMemType = C.CL_NONE
+	LocalMemTypeGlobal LocalMemType = C.CL_GLOBAL
+	LocalMemTypeLocal  LocalMemType = C.CL_LOCAL
+)
+
+var localMemTypeMap = map[LocalMemType]string{
+	LocalMemTypeNone:   "None",
+	LocalMemTypeGlobal: "Global",
+	LocalMemTypeLocal:  "Local",
+}
+
+func (t LocalMemType) String() string {
+	name := localMemTypeMap[t]
+	if name == "" {
+		name = "Unknown"
+	}
+	return name
+}
+
+type ExecCapability int
+
+const (
+	ExecCapabilityKernel       ExecCapability = C.CL_EXEC_KERNEL        // The OpenCL device can execute OpenCL kernels.
+	ExecCapabilityNativeKernel ExecCapability = C.CL_EXEC_NATIVE_KERNEL // The OpenCL device can execute native kernels.
+)
+
+func (ec ExecCapability) String() string {
+	var parts []string
+	if ec&ExecCapabilityKernel != 0 {
+		parts = append(parts, "Kernel")
+	}
+	if ec&ExecCapabilityNativeKernel != 0 {
+		parts = append(parts, "NativeKernel")
+	}
+	if parts == nil {
+		return ""
+	}
+	return strings.Join(parts, "|")
+}
+
+type MemCacheType int
+
+const (
+	MemCacheTypeNone           MemCacheType = C.CL_NONE
+	MemCacheTypeReadOnlyCache  MemCacheType = C.CL_READ_ONLY_CACHE
+	MemCacheTypeReadWriteCache MemCacheType = C.CL_READ_WRITE_CACHE
+)
+
+func (ct MemCacheType) String() string {
+	switch ct {
+	case MemCacheTypeNone:
+		return "None"
+	case MemCacheTypeReadOnlyCache:
+		return "ReadOnly"
+	case MemCacheTypeReadWriteCache:
+		return "ReadWrite"
+	}
+	return fmt.Sprintf("Unknown(%x)", int(ct))
+}
+
+type MemFlag int
+
+const (
+	MemReadWrite    MemFlag = C.CL_MEM_READ_WRITE
+	MemWriteOnly    MemFlag = C.CL_MEM_WRITE_ONLY
+	MemReadOnly     MemFlag = C.CL_MEM_READ_ONLY
+	MemUseHostPtr   MemFlag = C.CL_MEM_USE_HOST_PTR
+	MemAllocHostPtr MemFlag = C.CL_MEM_ALLOC_HOST_PTR
+	MemCopyHostPtr  MemFlag = C.CL_MEM_COPY_HOST_PTR
+
+	MemWriteOnlyHost MemFlag = C.CL_MEM_HOST_WRITE_ONLY
+	MemReadOnlyHost  MemFlag = C.CL_MEM_HOST_READ_ONLY
+	MemNoAccessHost  MemFlag = C.CL_MEM_HOST_NO_ACCESS
+)
+
+type MemObjectType int
+
+const (
+	MemObjectTypeBuffer  MemObjectType = C.CL_MEM_OBJECT_BUFFER
+	MemObjectTypeImage2D MemObjectType = C.CL_MEM_OBJECT_IMAGE2D
+	MemObjectTypeImage3D MemObjectType = C.CL_MEM_OBJECT_IMAGE3D
+)
+
+type MapFlag int
+
+const (
+	// This flag specifies that the region being mapped in the memory object is being mapped for reading.
+	MapFlagRead                  MapFlag = C.CL_MAP_READ
+	MapFlagWrite                 MapFlag = C.CL_MAP_WRITE
+	MapFlagWriteInvalidateRegion MapFlag = C.CL_MAP_WRITE_INVALIDATE_REGION
+)
+
+func (mf MapFlag) toCl() C.cl_map_flags {
+	return C.cl_map_flags(mf)
+}
+
+type ChannelOrder int
+
+const (
+	ChannelOrderR         ChannelOrder = C.CL_R
+	ChannelOrderA         ChannelOrder = C.CL_A
+	ChannelOrderRG        ChannelOrder = C.CL_RG
+	ChannelOrderRA        ChannelOrder = C.CL_RA
+	ChannelOrderRGB       ChannelOrder = C.CL_RGB
+	ChannelOrderRGBA      ChannelOrder = C.CL_RGBA
+	ChannelOrderBGRA      ChannelOrder = C.CL_BGRA
+	ChannelOrderARGB      ChannelOrder = C.CL_ARGB
+	ChannelOrderIntensity ChannelOrder = C.CL_INTENSITY
+	ChannelOrderLuminance ChannelOrder = C.CL_LUMINANCE
+	ChannelOrderRx        ChannelOrder = C.CL_Rx
+	ChannelOrderRGx       ChannelOrder = C.CL_RGx
+	ChannelOrderRGBx      ChannelOrder = C.CL_RGBx
+)
+
+var channelOrderNameMap = map[ChannelOrder]string{
+	ChannelOrderR:         "R",
+	ChannelOrderA:         "A",
+	ChannelOrderRG:        "RG",
+	ChannelOrderRA:        "RA",
+	ChannelOrderRGB:       "RGB",
+	ChannelOrderRGBA:      "RGBA",
+	ChannelOrderBGRA:      "BGRA",
+	ChannelOrderARGB:      "ARGB",
+	ChannelOrderIntensity: "Intensity",
+	ChannelOrderLuminance: "Luminance",
+	ChannelOrderRx:        "Rx",
+	ChannelOrderRGx:       "RGx",
+	ChannelOrderRGBx:      "RGBx",
+}
+
+func (co ChannelOrder) String() string {
+	name := channelOrderNameMap[co]
+	if name == "" {
+		name = fmt.Sprintf("Unknown(%x)", int(co))
+	}
+	return name
+}
+
+type ChannelDataType int
+
+const (
+	ChannelDataTypeSNormInt8      ChannelDataType = C.CL_SNORM_INT8
+	ChannelDataTypeSNormInt16     ChannelDataType = C.CL_SNORM_INT16
+	ChannelDataTypeUNormInt8      ChannelDataType = C.CL_UNORM_INT8
+	ChannelDataTypeUNormInt16     ChannelDataType = C.CL_UNORM_INT16
+	ChannelDataTypeUNormShort565  ChannelDataType = C.CL_UNORM_SHORT_565
+	ChannelDataTypeUNormShort555  ChannelDataType = C.CL_UNORM_SHORT_555
+	ChannelDataTypeUNormInt101010 ChannelDataType = C.CL_UNORM_INT_101010
+	ChannelDataTypeSignedInt8     ChannelDataType = C.CL_SIGNED_INT8
+	ChannelDataTypeSignedInt16    ChannelDataType = C.CL_SIGNED_INT16
+	ChannelDataTypeSignedInt32    ChannelDataType = C.CL_SIGNED_INT32
+	ChannelDataTypeUnsignedInt8   ChannelDataType = C.CL_UNSIGNED_INT8
+	ChannelDataTypeUnsignedInt16  ChannelDataType = C.CL_UNSIGNED_INT16
+	ChannelDataTypeUnsignedInt32  ChannelDataType = C.CL_UNSIGNED_INT32
+	ChannelDataTypeHalfFloat      ChannelDataType = C.CL_HALF_FLOAT
+	ChannelDataTypeFloat          ChannelDataType = C.CL_FLOAT
+)
+
+var channelDataTypeNameMap = map[ChannelDataType]string{
+	ChannelDataTypeSNormInt8:      "SNormInt8",
+	ChannelDataTypeSNormInt16:     "SNormInt16",
+	ChannelDataTypeUNormInt8:      "UNormInt8",
+	ChannelDataTypeUNormInt16:     "UNormInt16",
+	ChannelDataTypeUNormShort565:  "UNormShort565",
+	ChannelDataTypeUNormShort555:  "UNormShort555",
+	ChannelDataTypeUNormInt101010: "UNormInt101010",
+	ChannelDataTypeSignedInt8:     "SignedInt8",
+	ChannelDataTypeSignedInt16:    "SignedInt16",
+	ChannelDataTypeSignedInt32:    "SignedInt32",
+	ChannelDataTypeUnsignedInt8:   "UnsignedInt8",
+	ChannelDataTypeUnsignedInt16:  "UnsignedInt16",
+	ChannelDataTypeUnsignedInt32:  "UnsignedInt32",
+	ChannelDataTypeHalfFloat:      "HalfFloat",
+	ChannelDataTypeFloat:          "Float",
+}
+
+func (ct ChannelDataType) String() string {
+	name := channelDataTypeNameMap[ct]
+	if name == "" {
+		name = fmt.Sprintf("Unknown(%x)", int(ct))
+	}
+	return name
+}
+
+type ImageFormat struct {
+	ChannelOrder    ChannelOrder
+	ChannelDataType ChannelDataType
+}
+
+func (f ImageFormat) toCl() C.cl_image_format {
+	var format C.cl_image_format
+	format.image_channel_order = C.cl_channel_order(f.ChannelOrder)
+	format.image_channel_data_type = C.cl_channel_type(f.ChannelDataType)
+	return format
+}
+
+type ProfilingInfo int
+
+const (
+	// A 64-bit value that describes the current device time counter in
+	// nanoseconds when the command identified by event is enqueued in
+	// a command-queue by the host.
+	ProfilingInfoCommandQueued ProfilingInfo = C.CL_PROFILING_COMMAND_QUEUED
+	// A 64-bit value that describes the current device time counter in
+	// nanoseconds when the command identified by event that has been
+	// enqueued is submitted by the host to the device associated with the command-queue.
+	ProfilingInfoCommandSubmit ProfilingInfo = C.CL_PROFILING_COMMAND_SUBMIT
+	// A 64-bit value that describes the current device time counter in
+	// nanoseconds when the command identified by event starts execution on the device.
+	ProfilingInfoCommandStart ProfilingInfo = C.CL_PROFILING_COMMAND_START
+	// A 64-bit value that describes the current device time counter in
+	// nanoseconds when the command identified by event has finished
+	// execution on the device.
+	ProfilingInfoCommandEnd ProfilingInfo = C.CL_PROFILING_COMMAND_END
+)
+
+type CommmandExecStatus int
+
+const (
+	CommmandExecStatusComplete  CommmandExecStatus = C.CL_COMPLETE
+	CommmandExecStatusRunning   CommmandExecStatus = C.CL_RUNNING
+	CommmandExecStatusSubmitted CommmandExecStatus = C.CL_SUBMITTED
+	CommmandExecStatusQueued    CommmandExecStatus = C.CL_QUEUED
+)
+
+type Event struct {
+	clEvent C.cl_event
+}
+
+func releaseEvent(ev *Event) {
+	if ev.clEvent != nil {
+		C.clReleaseEvent(ev.clEvent)
+		ev.clEvent = nil
+	}
+}
+
+func (e *Event) Release() {
+	releaseEvent(e)
+}
+
+func (e *Event) GetEventProfilingInfo(paramName ProfilingInfo) (int64, error) {
+	var paramValue C.cl_ulong
+	if err := C.clGetEventProfilingInfo(e.clEvent, C.cl_profiling_info(paramName), C.size_t(unsafe.Sizeof(paramValue)), unsafe.Pointer(&paramValue), nil); err != C.CL_SUCCESS {
+		return 0, toError(err)
+	}
+	return int64(paramValue), nil
+}
+
+// Sets the execution status of a user event object.
+//
+// `status` specifies the new execution status to be set and
+// can be CL_COMPLETE or a negative integer value to indicate
+// an error. A negative integer value causes all enqueued commands
+// that wait on this user event to be terminated. clSetUserEventStatus
+// can only be called once to change the execution status of event.
+func (e *Event) SetUserEventStatus(status int) error {
+	return toError(C.clSetUserEventStatus(e.clEvent, C.cl_int(status)))
+}
+
+// Waits on the host thread for commands identified by event objects in
+// events to complete. A command is considered complete if its execution
+// status is CL_COMPLETE or a negative value. The events specified in
+// event_list act as synchronization points.
+//
+// If the cl_khr_gl_event extension is enabled, event objects can also be
+// used to reflect the status of an OpenGL sync object. The sync object
+// in turn refers to a fence command executing in an OpenGL command
+// stream. This provides another method of coordinating sharing of buffers
+// and images between OpenGL and OpenCL.
+func WaitForEvents(events []*Event) error {
+	return toError(C.clWaitForEvents(C.cl_uint(len(events)), eventListPtr(events)))
+}
+
+func newEvent(clEvent C.cl_event) *Event {
+	ev := &Event{clEvent: clEvent}
+	runtime.SetFinalizer(ev, releaseEvent)
+	return ev
+}
+
+func eventListPtr(el []*Event) *C.cl_event {
+	if el == nil {
+		return nil
+	}
+	elist := make([]C.cl_event, len(el))
+	for i, e := range el {
+		elist[i] = e.clEvent
+	}
+	return (*C.cl_event)(&elist[0])
+}
+
+func clBool(b bool) C.cl_bool {
+	if b {
+		return C.CL_TRUE
+	}
+	return C.CL_FALSE
+}
+
+func sizeT3(i3 [3]int) [3]C.size_t {
+	var val [3]C.size_t
+	val[0] = C.size_t(i3[0])
+	val[1] = C.size_t(i3[1])
+	val[2] = C.size_t(i3[2])
+	return val
+}
+
+type MappedMemObject struct {
+	ptr        unsafe.Pointer
+	size       int
+	rowPitch   int
+	slicePitch int
+}
+
+func (mb *MappedMemObject) ByteSlice() []byte {
+	var byteSlice []byte
+	sliceHeader := (*reflect.SliceHeader)(unsafe.Pointer(&byteSlice))
+	sliceHeader.Cap = mb.size
+	sliceHeader.Len = mb.size
+	sliceHeader.Data = uintptr(mb.ptr)
+	return byteSlice
+}
+
+func (mb *MappedMemObject) Ptr() unsafe.Pointer {
+	return mb.ptr
+}
+
+func (mb *MappedMemObject) Size() int {
+	return mb.size
+}
+
+func (mb *MappedMemObject) RowPitch() int {
+	return mb.rowPitch
+}
+
+func (mb *MappedMemObject) SlicePitch() int {
+	return mb.slicePitch
+}
--- a/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/types12.go
+++ b/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/types12.go
@@ -0,0 +1,71 @@
+// +build cl12
+
+package cl
+
+// #ifdef __APPLE__
+// #include "OpenCL/opencl.h"
+// #else
+// #include "cl.h"
+// #endif
+import "C"
+
+const (
+	ChannelDataTypeUNormInt24  ChannelDataType = C.CL_UNORM_INT24
+	ChannelOrderDepth          ChannelOrder    = C.CL_DEPTH
+	ChannelOrderDepthStencil   ChannelOrder    = C.CL_DEPTH_STENCIL
+	MemHostNoAccess            MemFlag         = C.CL_MEM_HOST_NO_ACCESS  // OpenCL 1.2
+	MemHostReadOnly            MemFlag         = C.CL_MEM_HOST_READ_ONLY  // OpenCL 1.2
+	MemHostWriteOnly           MemFlag         = C.CL_MEM_HOST_WRITE_ONLY // OpenCL 1.2
+	MemObjectTypeImage1D       MemObjectType   = C.CL_MEM_OBJECT_IMAGE1D
+	MemObjectTypeImage1DArray  MemObjectType   = C.CL_MEM_OBJECT_IMAGE1D_ARRAY
+	MemObjectTypeImage1DBuffer MemObjectType   = C.CL_MEM_OBJECT_IMAGE1D_BUFFER
+	MemObjectTypeImage2DArray  MemObjectType   = C.CL_MEM_OBJECT_IMAGE2D_ARRAY
+	// This flag specifies that the region being mapped in the memory object is being mapped for writing.
+	//
+	// The contents of the region being mapped are to be discarded. This is typically the case when the
+	// region being mapped is overwritten by the host. This flag allows the implementation to no longer
+	// guarantee that the pointer returned by clEnqueueMapBuffer or clEnqueueMapImage contains the
+	// latest bits in the region being mapped which can be a significant performance enhancement.
+	MapFlagWriteInvalidateRegion MapFlag = C.CL_MAP_WRITE_INVALIDATE_REGION
+)
+
+func init() {
+	errorMap[C.CL_COMPILE_PROGRAM_FAILURE] = ErrCompileProgramFailure
+	errorMap[C.CL_DEVICE_PARTITION_FAILED] = ErrDevicePartitionFailed
+	errorMap[C.CL_INVALID_COMPILER_OPTIONS] = ErrInvalidCompilerOptions
+	errorMap[C.CL_INVALID_DEVICE_PARTITION_COUNT] = ErrInvalidDevicePartitionCount
+	errorMap[C.CL_INVALID_IMAGE_DESCRIPTOR] = ErrInvalidImageDescriptor
+	errorMap[C.CL_INVALID_LINKER_OPTIONS] = ErrInvalidLinkerOptions
+	errorMap[C.CL_KERNEL_ARG_INFO_NOT_AVAILABLE] = ErrKernelArgInfoNotAvailable
+	errorMap[C.CL_LINK_PROGRAM_FAILURE] = ErrLinkProgramFailure
+	errorMap[C.CL_LINKER_NOT_AVAILABLE] = ErrLinkerNotAvailable
+	channelOrderNameMap[ChannelOrderDepth] = "Depth"
+	channelOrderNameMap[ChannelOrderDepthStencil] = "DepthStencil"
+	channelDataTypeNameMap[ChannelDataTypeUNormInt24] = "UNormInt24"
+}
+
+type ImageDescription struct {
+	Type                            MemObjectType
+	Width, Height, Depth            int
+	ArraySize, RowPitch, SlicePitch int
+	NumMipLevels, NumSamples        int
+	Buffer                          *MemObject
+}
+
+func (d ImageDescription) toCl() C.cl_image_desc {
+	var desc C.cl_image_desc
+	desc.image_type = C.cl_mem_object_type(d.Type)
+	desc.image_width = C.size_t(d.Width)
+	desc.image_height = C.size_t(d.Height)
+	desc.image_depth = C.size_t(d.Depth)
+	desc.image_array_size = C.size_t(d.ArraySize)
+	desc.image_row_pitch = C.size_t(d.RowPitch)
+	desc.image_slice_pitch = C.size_t(d.SlicePitch)
+	desc.num_mip_levels = C.cl_uint(d.NumMipLevels)
+	desc.num_samples = C.cl_uint(d.NumSamples)
+	desc.buffer = nil
+	if d.Buffer != nil {
+		desc.buffer = d.Buffer.clMem
+	}
+	return desc
+}
--- a/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/types_darwin.go
+++ b/Godeps/_workspace/src/github.com/Gustav-Simonsson/go-opencl/cl/types_darwin.go
@@ -0,0 +1,45 @@
+package cl
+
+// #ifdef __APPLE__
+// #include "OpenCL/opencl.h"
+// #else
+// #include "cl.h"
+// #endif
+import "C"
+
+// Extension: cl_APPLE_fixed_alpha_channel_orders
+//
+// These selectors may be passed to clCreateImage2D() in the cl_image_format.image_channel_order field.
+// They are like CL_BGRA and CL_ARGB except that the alpha channel to be ignored.  On calls to read_imagef,
+// the alpha will be 0xff (1.0f) if the sample falls in the image and 0 if it does not fall in the image.
+// On calls to write_imagef, the alpha value is ignored and 0xff (1.0f) is written. These formats are
+// currently only available for the CL_UNORM_INT8 cl_channel_type. They are intended to support legacy
+// image formats.
+const (
+	ChannelOrder1RGBApple ChannelOrder = C.CL_1RGB_APPLE // Introduced in MacOS X.7.
+	ChannelOrderBGR1Apple ChannelOrder = C.CL_BGR1_APPLE // Introduced in MacOS X.7.
+)
+
+// Extension: cl_APPLE_biased_fixed_point_image_formats
+//
+// This selector may be passed to clCreateImage2D() in the cl_image_format.image_channel_data_type field.
+// It defines a biased signed 1.14 fixed point storage format, with range [-1, 3). The conversion from
+// float to this fixed point format is defined as follows:
+//
+//      ushort float_to_sfixed14( float x ){
+//          int i = convert_int_sat_rte( x * 0x1.0p14f );         // scale [-1, 3.0) to [-16384, 3*16384), round to nearest integer
+//          i = add_sat( i, 0x4000 );                             // apply bias, to convert to [0, 65535) range
+//          return convert_ushort_sat(i);                         // clamp to destination size
+//      }
+//
+// The inverse conversion is the reverse process. The formats are currently only available on the CPU with
+// the CL_RGBA channel layout.
+const (
+	ChannelDataTypeSFixed14Apple ChannelDataType = C.CL_SFIXED14_APPLE // Introduced in MacOS X.7.
+)
+
+func init() {
+	channelOrderNameMap[ChannelOrder1RGBApple] = "1RGBApple"
+	channelOrderNameMap[ChannelOrderBGR1Apple] = "RGB1Apple"
+	channelDataTypeNameMap[ChannelDataTypeSFixed14Apple] = "SFixed14Apple"
+}
--- a/Godeps/_workspace/src/github.com/cespare/cp/LICENSE.txt
+++ b/Godeps/_workspace/src/github.com/cespare/cp/LICENSE.txt
--- a/Godeps/_workspace/src/github.com/cespare/cp/README.md
+++ b/Godeps/_workspace/src/github.com/cespare/cp/README.md
--- a/Godeps/_workspace/src/github.com/cespare/cp/cp.go
+++ b/Godeps/_workspace/src/github.com/cespare/cp/cp.go
--- a/Godeps/_workspace/src/github.com/codegangsta/cli/.travis.yml
+++ b/Godeps/_workspace/src/github.com/codegangsta/cli/.travis.yml
@@ -0,0 +1,19 @@
+language: go
+sudo: false
+
+go:
+- 1.0.3
+- 1.1.2
+- 1.2.2
+- 1.3.3
+- 1.4.2
+- 1.5.1
+- tip
+
+matrix:
+  allow_failures:
+    - go: tip
+
+script:
+- go vet ./...
+- go test -v ./...
--- a/Godeps/_workspace/src/github.com/codegangsta/cli/LICENSE
+++ b/Godeps/_workspace/src/github.com/codegangsta/cli/LICENSE
@@ -0,0 +1,21 @@
+Copyright (C) 2013 Jeremy Saenz
+All Rights Reserved.
+
+MIT LICENSE
+
+Permission is hereby granted, free of charge, to any person obtaining a copy of
+this software and associated documentation files (the "Software"), to deal in
+the Software without restriction, including without limitation the rights to
+use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of
+the Software, and to permit persons to whom the Software is furnished to do so,
+subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in all
+copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS
+FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR
+COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER
+IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN
+CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
--- a/Godeps/_workspace/src/github.com/codegangsta/cli/README.md
+++ b/Godeps/_workspace/src/github.com/codegangsta/cli/README.md
@@ -0,0 +1,352 @@
+[![Coverage](http://gocover.io/_badge/github.com/codegangsta/cli?0)](http://gocover.io/github.com/codegangsta/cli)
+[![Build Status](https://travis-ci.org/codegangsta/cli.svg?branch=master)](https://travis-ci.org/codegangsta/cli)
+[![GoDoc](https://godoc.org/github.com/codegangsta/cli?status.svg)](https://godoc.org/github.com/codegangsta/cli)
+
+# cli.go
+
+`cli.go` is simple, fast, and fun package for building command line apps in Go. The goal is to enable developers to write fast and distributable command line applications in an expressive way.
+
+## Overview
+
+Command line apps are usually so tiny that there is absolutely no reason why your code should *not* be self-documenting. Things like generating help text and parsing command flags/options should not hinder productivity when writing a command line app.
+
+**This is where `cli.go` comes into play.** `cli.go` makes command line programming fun, organized, and expressive!
+
+## Installation
+
+Make sure you have a working Go environment (go 1.1+ is *required*). [See the install instructions](http://golang.org/doc/install.html).
+
+To install `cli.go`, simply run:
+```
+$ go get github.com/codegangsta/cli
+```
+
+Make sure your `PATH` includes to the `$GOPATH/bin` directory so your commands can be easily used:
+```
+export PATH=$PATH:$GOPATH/bin
+```
+
+## Getting Started
+
+One of the philosophies behind `cli.go` is that an API should be playful and full of discovery. So a `cli.go` app can be as little as one line of code in `main()`. 
+
+``` go
+package main
+
+import (
+  "os"
+  "github.com/codegangsta/cli"
+)
+
+func main() {
+  cli.NewApp().Run(os.Args)
+}
+```
+
+This app will run and show help text, but is not very useful. Let's give an action to execute and some help documentation:
+
+``` go
+package main
+
+import (
+  "os"
+  "github.com/codegangsta/cli"
+)
+
+func main() {
+  app := cli.NewApp()
+  app.Name = "boom"
+  app.Usage = "make an explosive entrance"
+  app.Action = func(c *cli.Context) {
+    println("boom! I say!")
+  }
+  
+  app.Run(os.Args)
+}
+```
+
+Running this already gives you a ton of functionality, plus support for things like subcommands and flags, which are covered below.
+
+## Example
+
+Being a programmer can be a lonely job. Thankfully by the power of automation that is not the case! Let's create a greeter app to fend off our demons of loneliness!
+
+Start by creating a directory named `greet`, and within it, add a file, `greet.go` with the following code in it:
+
+``` go
+package main
+
+import (
+  "os"
+  "github.com/codegangsta/cli"
+)
+
+func main() {
+  app := cli.NewApp()
+  app.Name = "greet"
+  app.Usage = "fight the loneliness!"
+  app.Action = func(c *cli.Context) {
+    println("Hello friend!")
+  }
+
+  app.Run(os.Args)
+}
+```
+
+Install our command to the `$GOPATH/bin` directory:
+
+```
+$ go install
+```
+
+Finally run our new command:
+
+```
+$ greet
+Hello friend!
+```
+
+`cli.go` also generates neat help text:
+
+```
+$ greet help
+NAME:
+    greet - fight the loneliness!
+
+USAGE:
+    greet [global options] command [command options] [arguments...]
+
+VERSION:
+    0.0.0
+
+COMMANDS:
+    help, h  Shows a list of commands or help for one command
+
+GLOBAL OPTIONS
+    --version	Shows version information
+```
+
+### Arguments
+
+You can lookup arguments by calling the `Args` function on `cli.Context`.
+
+``` go
+...
+app.Action = func(c *cli.Context) {
+  println("Hello", c.Args()[0])
+}
+...
+```
+
+### Flags
+
+Setting and querying flags is simple.
+
+``` go
+...
+app.Flags = []cli.Flag {
+  cli.StringFlag{
+    Name: "lang",
+    Value: "english",
+    Usage: "language for the greeting",
+  },
+}
+app.Action = func(c *cli.Context) {
+  name := "someone"
+  if len(c.Args()) > 0 {
+    name = c.Args()[0]
+  }
+  if c.String("lang") == "spanish" {
+    println("Hola", name)
+  } else {
+    println("Hello", name)
+  }
+}
+...
+```
+
+You can also set a destination variable for a flag, to which the content will be scanned.
+
+``` go
+...
+var language string
+app.Flags = []cli.Flag {
+  cli.StringFlag{
+    Name:        "lang",
+    Value:       "english",
+    Usage:       "language for the greeting",
+    Destination: &language,
+  },
+}
+app.Action = func(c *cli.Context) {
+  name := "someone"
+  if len(c.Args()) > 0 {
+    name = c.Args()[0]
+  }
+  if language == "spanish" {
+    println("Hola", name)
+  } else {
+    println("Hello", name)
+  }
+}
+...
+```
+
+See full list of flags at http://godoc.org/github.com/codegangsta/cli
+
+#### Alternate Names
+
+You can set alternate (or short) names for flags by providing a comma-delimited list for the `Name`. e.g.
+
+``` go
+app.Flags = []cli.Flag {
+  cli.StringFlag{
+    Name: "lang, l",
+    Value: "english",
+    Usage: "language for the greeting",
+  },
+}
+```
+
+That flag can then be set with `--lang spanish` or `-l spanish`. Note that giving two different forms of the same flag in the same command invocation is an error.
+
+#### Values from the Environment
+
+You can also have the default value set from the environment via `EnvVar`.  e.g.
+
+``` go
+app.Flags = []cli.Flag {
+  cli.StringFlag{
+    Name: "lang, l",
+    Value: "english",
+    Usage: "language for the greeting",
+    EnvVar: "APP_LANG",
+  },
+}
+```
+
+The `EnvVar` may also be given as a comma-delimited "cascade", where the first environment variable that resolves is used as the default.
+
+``` go
+app.Flags = []cli.Flag {
+  cli.StringFlag{
+    Name: "lang, l",
+    Value: "english",
+    Usage: "language for the greeting",
+    EnvVar: "LEGACY_COMPAT_LANG,APP_LANG,LANG",
+  },
+}
+```
+
+### Subcommands
+
+Subcommands can be defined for a more git-like command line app.
+
+```go
+...
+app.Commands = []cli.Command{
+  {
+    Name:      "add",
+    Aliases:     []string{"a"},
+    Usage:     "add a task to the list",
+    Action: func(c *cli.Context) {
+      println("added task: ", c.Args().First())
+    },
+  },
+  {
+    Name:      "complete",
+    Aliases:     []string{"c"},
+    Usage:     "complete a task on the list",
+    Action: func(c *cli.Context) {
+      println("completed task: ", c.Args().First())
+    },
+  },
+  {
+    Name:      "template",
+    Aliases:     []string{"r"},
+    Usage:     "options for task templates",
+    Subcommands: []cli.Command{
+      {
+        Name:  "add",
+        Usage: "add a new template",
+        Action: func(c *cli.Context) {
+            println("new task template: ", c.Args().First())
+        },
+      },
+      {
+        Name:  "remove",
+        Usage: "remove an existing template",
+        Action: func(c *cli.Context) {
+          println("removed task template: ", c.Args().First())
+        },
+      },
+    },
+  },
+}
+...
+```
+
+### Bash Completion
+
+You can enable completion commands by setting the `EnableBashCompletion`
+flag on the `App` object.  By default, this setting will only auto-complete to
+show an app's subcommands, but you can write your own completion methods for
+the App or its subcommands.
+
+```go
+...
+var tasks = []string{"cook", "clean", "laundry", "eat", "sleep", "code"}
+app := cli.NewApp()
+app.EnableBashCompletion = true
+app.Commands = []cli.Command{
+  {
+    Name:  "complete",
+    Aliases: []string{"c"},
+    Usage: "complete a task on the list",
+    Action: func(c *cli.Context) {
+       println("completed task: ", c.Args().First())
+    },
+    BashComplete: func(c *cli.Context) {
+      // This will complete if no args are passed
+      if len(c.Args()) > 0 {
+        return
+      }
+      for _, t := range tasks {
+        fmt.Println(t)
+      }
+    },
+  }
+}
+...
+```
+
+#### To Enable
+
+Source the `autocomplete/bash_autocomplete` file in your `.bashrc` file while
+setting the `PROG` variable to the name of your program:
+
+`PROG=myprogram source /.../cli/autocomplete/bash_autocomplete`
+
+#### To Distribute
+
+Copy `autocomplete/bash_autocomplete` into `/etc/bash_completion.d/` and rename
+it to the name of the program you wish to add autocomplete support for (or
+automatically install it there if you are distributing a package). Don't forget
+to source the file to make it active in the current shell.
+
+```
+sudo cp src/bash_autocomplete /etc/bash_completion.d/<myprogram>
+source /etc/bash_completion.d/<myprogram>
+```
+
+Alternatively, you can just document that users should source the generic
+`autocomplete/bash_autocomplete` in their bash configuration with `$PROG` set
+to the name of their program (as above).
+
+## Contribution Guidelines
+
+Feel free to put up a pull request to fix a bug or maybe add a feature. I will give it a code review and make sure that it does not break backwards compatibility. If I or any other collaborators agree that it is in line with the vision of the project, we will work with you to get the code into a mergeable state and merge it into the master branch.
+
+If you have contributed something significant to the project, I will most likely add you as a collaborator. As a collaborator you are given the ability to merge others pull requests. It is very important that new code does not break existing code, so be careful about what code you do choose to merge. If you have any questions feel free to link @codegangsta to the issue in question and we can review it together.
+
+If you feel like you have contributed to the project but have not yet been added as a collaborator, I probably forgot to add you. Hit @codegangsta up over email and we will get it figured out.
--- a/Godeps/_workspace/src/github.com/codegangsta/cli/app.go
+++ b/Godeps/_workspace/src/github.com/codegangsta/cli/app.go
@@ -5,28 +5,10 @@ import (
 	"io"
 	"io/ioutil"
 	"os"
-	"path/filepath"
-	"reflect"
-	"sort"
-	"strings"
+	"path"
 	"time"
 )

-var (
-	changeLogURL                    = "https://github.com/urfave/cli/blob/master/CHANGELOG.md"
-	appActionDeprecationURL         = fmt.Sprintf("%s#deprecated-cli-app-action-signature", changeLogURL)
-	runAndExitOnErrorDeprecationURL = fmt.Sprintf("%s#deprecated-cli-app-runandexitonerror", changeLogURL)
-
-	contactSysadmin = "This is an error in the application.  Please contact the distributor of this application if this is not you."
-
-	errNonFuncAction = NewExitError("ERROR invalid Action type.  "+
-		fmt.Sprintf("Must be a func of type `cli.ActionFunc`.  %s", contactSysadmin)+
-		fmt.Sprintf("See %s", appActionDeprecationURL), 2)
-	errInvalidActionSignature = NewExitError("ERROR invalid Action signature.  "+
-		fmt.Sprintf("Must be `cli.ActionFunc`.  %s", contactSysadmin)+
-		fmt.Sprintf("See %s", appActionDeprecationURL), 2)
-)
-
 // App is the main structure of a cli application. It is recommended that
 // an app be created with the cli.NewApp() function
 type App struct {
@@ -50,27 +32,24 @@ type App struct {
 	EnableBashCompletion bool
 	// Boolean to hide built-in help command
 	HideHelp bool
-	// Boolean to hide built-in version flag and the VERSION section of help
+	// Boolean to hide built-in version flag
 	HideVersion bool
-	// Populate on app startup, only gettable through method Categories()
-	categories CommandCategories
 	// An action to execute when the bash-completion flag is set
-	BashComplete BashCompleteFunc
+	BashComplete func(context *Context)
 	// An action to execute before any subcommands are run, but after the context is ready
 	// If a non-nil error is returned, no subcommands are run
-	Before BeforeFunc
+	Before func(context *Context) error
 	// An action to execute after any subcommands are run, but after the subcommand has finished
 	// It is run even if Action() panics
-	After AfterFunc
+	After func(context *Context) error
 	// The action to execute when no subcommands are specified
-	// Expects a `cli.ActionFunc` but will accept the *deprecated* signature of `func(*cli.Context) {}`
-	// *Note*: support for the deprecated `Action` signature will be removed in a future version
-	Action interface{}
-
+	Action func(context *Context)
 	// Execute this function if the proper command cannot be found
-	CommandNotFound CommandNotFoundFunc
-	// Execute this function if an usage error occurs
-	OnUsageError OnUsageErrorFunc
+	CommandNotFound func(context *Context, command string)
+	// Execute this function, if an usage error occurs. This is useful for displaying customized usage error messages.
+	// This function is able to replace the original error messages.
+	// If this function is not set, the "Incorrect usage" is displayed and the execution is interrupted.
+	OnUsageError func(context *Context, err error, isSubcommand bool) error
 	// Compilation date
 	Compiled time.Time
 	// List of all authors who contributed
@@ -83,12 +62,6 @@ type App struct {
 	Email string
 	// Writer writer to write output to
 	Writer io.Writer
-	// ErrWriter writes error output
-	ErrWriter io.Writer
-	// Other custom info
-	Metadata map[string]interface{}
-
-	didSetup bool
 }

 // Tries to find out when this binary was compiled.
@@ -101,12 +74,11 @@ func compileTime() time.Time {
 	return info.ModTime()
 }

-// NewApp creates a new cli Application with some reasonable defaults for Name,
-// Usage, Version and Action.
+// Creates a new cli Application with some reasonable defaults for Name, Usage, Version and Action.
 func NewApp() *App {
 	return &App{
-		Name:         filepath.Base(os.Args[0]),
-		HelpName:     filepath.Base(os.Args[0]),
+		Name:         path.Base(os.Args[0]),
+		HelpName:     path.Base(os.Args[0]),
 		Usage:        "A new cli application",
 		UsageText:    "",
 		Version:      "0.0.0",
@@ -117,16 +89,8 @@ func NewApp() *App {
 	}
 }

-// Setup runs initialization code to ensure all data structures are ready for
-// `Run` or inspection prior to `Run`.  It is internally called by `Run`, but
-// will return early if setup has already happened.
-func (a *App) Setup() {
-	if a.didSetup {
-		return
-	}
-
-	a.didSetup = true
-
+// Entry point to the cli app. Parses the arguments slice and routes to the proper flag/args combination
+func (a *App) Run(arguments []string) (err error) {
 	if a.Author != "" || a.Email != "" {
 		a.Authors = append(a.Authors, Author{Name: a.Author, Email: a.Email})
 	}
@@ -140,6 +104,7 @@ func (a *App) Setup() {
 	}
 	a.Commands = newCmds

+	// append help to commands
 	if a.Command(helpCommand.Name) == nil && !a.HideHelp {
 		a.Commands = append(a.Commands, helpCommand)
 		if (HelpFlag != BoolFlag{}) {
@@ -147,6 +112,7 @@ func (a *App) Setup() {
 		}
 	}

+	//append version/help flags
 	if a.EnableBashCompletion {
 		a.appendFlag(BashCompletionFlag)
 	}
@@ -155,18 +121,6 @@ func (a *App) Setup() {
 		a.appendFlag(VersionFlag)
 	}

-	a.categories = CommandCategories{}
-	for _, command := range a.Commands {
-		a.categories = a.categories.AddCommand(command.Category, command)
-	}
-	sort.Sort(a.categories)
-}
-
-// Run is the entry point to the cli app. Parses the arguments slice and routes
-// to the proper flag/args combination
-func (a *App) Run(arguments []string) (err error) {
-	a.Setup()
-
 	// parse flags
 	set := flagSet(a.Name, a.Flags)
 	set.SetOutput(ioutil.Discard)
@@ -186,12 +140,12 @@ func (a *App) Run(arguments []string) (err error) {
 	if err != nil {
 		if a.OnUsageError != nil {
 			err := a.OnUsageError(context, err, false)
-			HandleExitCoder(err)
+			return err
+		} else {
+			fmt.Fprintf(a.Writer, "%s\n\n", "Incorrect Usage.")
+			ShowAppHelp(context)
 			return err
 		}
-		fmt.Fprintf(a.Writer, "%s\n\n", "Incorrect Usage.")
-		ShowAppHelp(context)
-		return err
 	}

 	if !a.HideHelp && checkHelp(context) {
@@ -217,12 +171,10 @@ func (a *App) Run(arguments []string) (err error) {
 	}

 	if a.Before != nil {
-		beforeErr := a.Before(context)
-		if beforeErr != nil {
-			fmt.Fprintf(a.Writer, "%v\n\n", beforeErr)
+		err = a.Before(context)
+		if err != nil {
+			fmt.Fprintf(a.Writer, "%v\n\n", err)
 			ShowAppHelp(context)
-			HandleExitCoder(beforeErr)
-			err = beforeErr
 			return err
 		}
 	}
@@ -237,26 +189,19 @@ func (a *App) Run(arguments []string) (err error) {
 	}

 	// Run default Action
-	err = HandleAction(a.Action, context)
-
-	HandleExitCoder(err)
-	return err
+	a.Action(context)
+	return nil
 }

-// RunAndExitOnError calls .Run() and exits non-zero if an error was returned
-//
-// Deprecated: instead you should return an error that fulfills cli.ExitCoder
-// to cli.App.Run. This will cause the application to exit with the given eror
-// code in the cli.ExitCoder
+// Another entry point to the cli app, takes care of passing arguments and error handling
 func (a *App) RunAndExitOnError() {
 	if err := a.Run(os.Args); err != nil {
-		fmt.Fprintln(a.errWriter(), err)
-		OsExiter(1)
+		fmt.Fprintln(os.Stderr, err)
+		os.Exit(1)
 	}
 }

-// RunAsSubcommand invokes the subcommand given the context, parses ctx.Args() to
-// generate command-specific flags
+// Invokes the subcommand given the context, parses ctx.Args() to generate command-specific flags
 func (a *App) RunAsSubcommand(ctx *Context) (err error) {
 	// append help to commands
 	if len(a.Commands) > 0 {
@@ -307,12 +252,12 @@ func (a *App) RunAsSubcommand(ctx *Context) (err error) {
 	if err != nil {
 		if a.OnUsageError != nil {
 			err = a.OnUsageError(context, err, true)
-			HandleExitCoder(err)
+			return err
+		} else {
+			fmt.Fprintf(a.Writer, "%s\n\n", "Incorrect Usage.")
+			ShowSubcommandHelp(context)
 			return err
 		}
-		fmt.Fprintf(a.Writer, "%s\n\n", "Incorrect Usage.")
-		ShowSubcommandHelp(context)
-		return err
 	}

 	if len(a.Commands) > 0 {
@@ -329,7 +274,6 @@ func (a *App) RunAsSubcommand(ctx *Context) (err error) {
 		defer func() {
 			afterErr := a.After(context)
 			if afterErr != nil {
-				HandleExitCoder(err)
 				if err != nil {
 					err = NewMultiError(err, afterErr)
 				} else {
@@ -340,10 +284,8 @@ func (a *App) RunAsSubcommand(ctx *Context) (err error) {
 	}

 	if a.Before != nil {
-		beforeErr := a.Before(context)
-		if beforeErr != nil {
-			HandleExitCoder(beforeErr)
-			err = beforeErr
+		err := a.Before(context)
+		if err != nil {
 			return err
 		}
 	}
@@ -358,13 +300,12 @@ func (a *App) RunAsSubcommand(ctx *Context) (err error) {
 	}

 	// Run default Action
-	err = HandleAction(a.Action, context)
+	a.Action(context)

-	HandleExitCoder(err)
-	return err
+	return nil
 }

-// Command returns the named command on App. Returns nil if the command does not exist
+// Returns the named command on App. Returns nil if the command does not exist
 func (a *App) Command(name string) *Command {
 	for _, c := range a.Commands {
 		if c.HasName(name) {
@@ -375,46 +316,6 @@ func (a *App) Command(name string) *Command {
 	return nil
 }

-// Categories returns a slice containing all the categories with the commands they contain
-func (a *App) Categories() CommandCategories {
-	return a.categories
-}
-
-// VisibleCategories returns a slice of categories and commands that are
-// Hidden=false
-func (a *App) VisibleCategories() []*CommandCategory {
-	ret := []*CommandCategory{}
-	for _, category := range a.categories {
-		if visible := func() *CommandCategory {
-			for _, command := range category.Commands {
-				if !command.Hidden {
-					return category
-				}
-			}
-			return nil
-		}(); visible != nil {
-			ret = append(ret, visible)
-		}
-	}
-	return ret
-}
-
-// VisibleCommands returns a slice of the Commands with Hidden=false
-func (a *App) VisibleCommands() []Command {
-	ret := []Command{}
-	for _, command := range a.Commands {
-		if !command.Hidden {
-			ret = append(ret, command)
-		}
-	}
-	return ret
-}
-
-// VisibleFlags returns a slice of the Flags with Hidden=false
-func (a *App) VisibleFlags() []Flag {
-	return visibleFlags(a.Flags)
-}
-
 func (a *App) hasFlag(flag Flag) bool {
 	for _, f := range a.Flags {
 		if flag == f {
@@ -425,16 +326,6 @@ func (a *App) hasFlag(flag Flag) bool {
 	return false
 }

-func (a *App) errWriter() io.Writer {
-
-	// When the app ErrWriter is nil use the package level one.
-	if a.ErrWriter == nil {
-		return ErrWriter
-	}
-
-	return a.ErrWriter
-}
-
 func (a *App) appendFlag(flag Flag) {
 	if !a.hasFlag(flag) {
 		a.Flags = append(a.Flags, flag)
@@ -456,42 +347,3 @@ func (a Author) String() string {

 	return fmt.Sprintf("%v %v", a.Name, e)
 }
-
-// HandleAction uses ✧✧✧reflection✧✧✧ to figure out if the given Action is an
-// ActionFunc, a func with the legacy signature for Action, or some other
-// invalid thing.  If it's an ActionFunc or a func with the legacy signature for
-// Action, the func is run!
-func HandleAction(action interface{}, context *Context) (err error) {
-	defer func() {
-		if r := recover(); r != nil {
-			// Try to detect a known reflection error from *this scope*, rather than
-			// swallowing all panics that may happen when calling an Action func.
-			s := fmt.Sprintf("%v", r)
-			if strings.HasPrefix(s, "reflect: ") && strings.Contains(s, "too many input arguments") {
-				err = NewExitError(fmt.Sprintf("ERROR unknown Action error: %v.  See %s", r, appActionDeprecationURL), 2)
-			} else {
-				panic(r)
-			}
-		}
-	}()
-
-	if reflect.TypeOf(action).Kind() != reflect.Func {
-		return errNonFuncAction
-	}
-
-	vals := reflect.ValueOf(action).Call([]reflect.Value{reflect.ValueOf(context)})
-
-	if len(vals) == 0 {
-		return nil
-	}
-
-	if len(vals) > 1 {
-		return errInvalidActionSignature
-	}
-
-	if retErr, ok := vals[0].Interface().(error); vals[0].IsValid() && ok {
-		return retErr
-	}
-
-	return err
-}
--- a/Godeps/_workspace/src/github.com/codegangsta/cli/appveyor.yml
+++ b/Godeps/_workspace/src/github.com/codegangsta/cli/appveyor.yml
@@ -0,0 +1,16 @@
+version: "{build}"
+
+os: Windows Server 2012 R2
+
+install:
+  - go version
+  - go env
+
+build_script:
+  - cd %APPVEYOR_BUILD_FOLDER%
+  - go vet ./...
+  - go test -v ./...
+
+test: off
+
+deploy: off
--- a/Godeps/_workspace/src/github.com/codegangsta/cli/autocomplete/bash_autocomplete
+++ b/Godeps/_workspace/src/github.com/codegangsta/cli/autocomplete/bash_autocomplete
@@ -0,0 +1,14 @@
+#! /bin/bash
+
+: ${PROG:=$(basename ${BASH_SOURCE})}
+
+_cli_bash_autocomplete() {
+     local cur opts base
+     COMPREPLY=()
+     cur="${COMP_WORDS[COMP_CWORD]}"
+     opts=$( ${COMP_WORDS[@]:0:$COMP_CWORD} --generate-bash-completion )
+     COMPREPLY=( $(compgen -W "${opts}" -- ${cur}) )
+     return 0
+ }
+  
+ complete -F _cli_bash_autocomplete $PROG
--- a/Godeps/_workspace/src/github.com/codegangsta/cli/autocomplete/zsh_autocomplete
+++ b/Godeps/_workspace/src/github.com/codegangsta/cli/autocomplete/zsh_autocomplete
@@ -0,0 +1,5 @@
+autoload -U compinit && compinit
+autoload -U bashcompinit && bashcompinit
+
+script_dir=$(dirname $0)
+source ${script_dir}/bash_autocomplete
--- a/Godeps/_workspace/src/github.com/codegangsta/cli/cli.go
+++ b/Godeps/_workspace/src/github.com/codegangsta/cli/cli.go
@@ -10,10 +10,31 @@
 //     app := cli.NewApp()
 //     app.Name = "greet"
 //     app.Usage = "say a greeting"
-//     app.Action = func(c *cli.Context) error {
+//     app.Action = func(c *cli.Context) {
 //       println("Greetings")
 //     }
 //
 //     app.Run(os.Args)
 //   }
 package cli
+
+import (
+	"strings"
+)
+
+type MultiError struct {
+	Errors []error
+}
+
+func NewMultiError(err ...error) MultiError {
+	return MultiError{Errors: err}
+}
+
+func (m MultiError) Error() string {
+	errs := make([]string, len(m.Errors))
+	for i, err := range m.Errors {
+		errs[i] = err.Error()
+	}
+
+	return strings.Join(errs, "\n")
+}
--- a/Godeps/_workspace/src/github.com/codegangsta/cli/command.go
+++ b/Godeps/_workspace/src/github.com/codegangsta/cli/command.go
@@ -3,7 +3,6 @@ package cli
 import (
 	"fmt"
 	"io/ioutil"
-	"sort"
 	"strings"
 )

@@ -23,40 +22,35 @@ type Command struct {
 	Description string
 	// A short description of the arguments of this command
 	ArgsUsage string
-	// The category the command is part of
-	Category string
 	// The function to call when checking for bash command completions
-	BashComplete BashCompleteFunc
+	BashComplete func(context *Context)
 	// An action to execute before any sub-subcommands are run, but after the context is ready
 	// If a non-nil error is returned, no sub-subcommands are run
-	Before BeforeFunc
+	Before func(context *Context) error
 	// An action to execute after any subcommands are run, but after the subcommand has finished
 	// It is run even if Action() panics
-	After AfterFunc
+	After func(context *Context) error
 	// The function to call when this command is invoked
-	Action interface{}
-	// TODO: replace `Action: interface{}` with `Action: ActionFunc` once some kind
-	// of deprecation period has passed, maybe?
-
-	// Execute this function if a usage error occurs.
-	OnUsageError OnUsageErrorFunc
+	Action func(context *Context)
+	// Execute this function, if an usage error occurs. This is useful for displaying customized usage error messages.
+	// This function is able to replace the original error messages.
+	// If this function is not set, the "Incorrect usage" is displayed and the execution is interrupted.
+	OnUsageError func(context *Context, err error) error
 	// List of child commands
-	Subcommands Commands
+	Subcommands []Command
 	// List of flags to parse
 	Flags []Flag
 	// Treat all flags as normal arguments if true
 	SkipFlagParsing bool
 	// Boolean to hide built-in help command
 	HideHelp bool
-	// Boolean to hide this command from help or completion
-	Hidden bool

 	// Full name of command for help, defaults to full command name, including parent commands.
 	HelpName        string
 	commandNamePath []string
 }

-// FullName returns the full name of the command.
+// Returns the full name of the command.
 // For subcommands this ensures that parent commands are part of the command path
 func (c Command) FullName() string {
 	if c.commandNamePath == nil {
@@ -65,10 +59,7 @@ func (c Command) FullName() string {
 	return strings.Join(c.commandNamePath, " ")
 }

-// Commands is a slice of Command
-type Commands []Command
-
-// Run invokes the command given the context, parses ctx.Args() to generate command-specific flags
+// Invokes the command given the context, parses ctx.Args() to generate command-specific flags
 func (c Command) Run(ctx *Context) (err error) {
 	if len(c.Subcommands) > 0 {
 		return c.startApp(ctx)
@@ -129,14 +120,14 @@ func (c Command) Run(ctx *Context) (err error) {

 	if err != nil {
 		if c.OnUsageError != nil {
-			err := c.OnUsageError(ctx, err, false)
-			HandleExitCoder(err)
+			err := c.OnUsageError(ctx, err)
+			return err
+		} else {
+			fmt.Fprintln(ctx.App.Writer, "Incorrect Usage.")
+			fmt.Fprintln(ctx.App.Writer)
+			ShowCommandHelp(ctx, c.Name)
 			return err
 		}
-		fmt.Fprintln(ctx.App.Writer, "Incorrect Usage.")
-		fmt.Fprintln(ctx.App.Writer)
-		ShowCommandHelp(ctx, c.Name)
-		return err
 	}

 	nerr := normalizeFlags(c.Flags, set)
@@ -146,7 +137,6 @@ func (c Command) Run(ctx *Context) (err error) {
 		ShowCommandHelp(ctx, c.Name)
 		return nerr
 	}
-
 	context := NewContext(ctx.App, set, ctx)

 	if checkCommandCompletions(context, c.Name) {
@@ -161,7 +151,6 @@ func (c Command) Run(ctx *Context) (err error) {
 		defer func() {
 			afterErr := c.After(context)
 			if afterErr != nil {
-				HandleExitCoder(err)
 				if err != nil {
 					err = NewMultiError(err, afterErr)
 				} else {
@@ -172,26 +161,20 @@ func (c Command) Run(ctx *Context) (err error) {
 	}

 	if c.Before != nil {
-		err = c.Before(context)
+		err := c.Before(context)
 		if err != nil {
 			fmt.Fprintln(ctx.App.Writer, err)
 			fmt.Fprintln(ctx.App.Writer)
 			ShowCommandHelp(ctx, c.Name)
-			HandleExitCoder(err)
 			return err
 		}
 	}

 	context.Command = c
-	err = HandleAction(c.Action, context)
-
-	if err != nil {
-		HandleExitCoder(err)
-	}
-	return err
+	c.Action(context)
+	return nil
 }

-// Names returns the names including short names and aliases.
 func (c Command) Names() []string {
 	names := []string{c.Name}

@@ -202,7 +185,7 @@ func (c Command) Names() []string {
 	return append(names, c.Aliases...)
 }

-// HasName returns true if Command.Name or Command.ShortName matches given name
+// Returns true if Command.Name or Command.ShortName matches given name
 func (c Command) HasName(name string) bool {
 	for _, n := range c.Names() {
 		if n == name {
@@ -214,7 +197,7 @@ func (c Command) HasName(name string) bool {

 func (c Command) startApp(ctx *Context) error {
 	app := NewApp()
-	app.Metadata = ctx.App.Metadata
+
 	// set the name and usage
 	app.Name = fmt.Sprintf("%s %s", ctx.App.Name, c.Name)
 	if c.HelpName == "" {
@@ -244,13 +227,6 @@ func (c Command) startApp(ctx *Context) error {
 	app.Email = ctx.App.Email
 	app.Writer = ctx.App.Writer

-	app.categories = CommandCategories{}
-	for _, command := range c.Subcommands {
-		app.categories = app.categories.AddCommand(command.Category, command)
-	}
-
-	sort.Sort(app.categories)
-
 	// bash completion
 	app.EnableBashCompletion = ctx.App.EnableBashCompletion
 	if c.BashComplete != nil {
@@ -272,8 +248,3 @@ func (c Command) startApp(ctx *Context) error {

 	return app.RunAsSubcommand(ctx)
 }
-
-// VisibleFlags returns a slice of the Flags with Hidden=false
-func (c Command) VisibleFlags() []Flag {
-	return visibleFlags(c.Flags)
-}
--- a/Godeps/_workspace/src/github.com/codegangsta/cli/context.go
+++ b/Godeps/_workspace/src/github.com/codegangsta/cli/context.go
@@ -21,83 +21,57 @@ type Context struct {
 	parentContext  *Context
 }

-// NewContext creates a new context. For use in when invoking an App or Command action.
+// Creates a new context. For use in when invoking an App or Command action.
 func NewContext(app *App, set *flag.FlagSet, parentCtx *Context) *Context {
 	return &Context{App: app, flagSet: set, parentContext: parentCtx}
 }

-// Int looks up the value of a local int flag, returns 0 if no int flag exists
+// Looks up the value of a local int flag, returns 0 if no int flag exists
 func (c *Context) Int(name string) int {
 	return lookupInt(name, c.flagSet)
 }

-// Int64 looks up the value of a local int flag, returns 0 if no int flag exists
-func (c *Context) Int64(name string) int64 {
-	return lookupInt64(name, c.flagSet)
-}
-
-// Uint looks up the value of a local int flag, returns 0 if no int flag exists
-func (c *Context) Uint(name string) uint {
-	return lookupUint(name, c.flagSet)
-}
-
-// Uint64 looks up the value of a local int flag, returns 0 if no int flag exists
-func (c *Context) Uint64(name string) uint64 {
-	return lookupUint64(name, c.flagSet)
-}
-
-// Duration looks up the value of a local time.Duration flag, returns 0 if no
-// time.Duration flag exists
+// Looks up the value of a local time.Duration flag, returns 0 if no time.Duration flag exists
 func (c *Context) Duration(name string) time.Duration {
 	return lookupDuration(name, c.flagSet)
 }

-// Float64 looks up the value of a local float64 flag, returns 0 if no float64
-// flag exists
+// Looks up the value of a local float64 flag, returns 0 if no float64 flag exists
 func (c *Context) Float64(name string) float64 {
 	return lookupFloat64(name, c.flagSet)
 }

-// Bool looks up the value of a local bool flag, returns false if no bool flag exists
+// Looks up the value of a local bool flag, returns false if no bool flag exists
 func (c *Context) Bool(name string) bool {
 	return lookupBool(name, c.flagSet)
 }

-// BoolT looks up the value of a local boolT flag, returns false if no bool flag exists
+// Looks up the value of a local boolT flag, returns false if no bool flag exists
 func (c *Context) BoolT(name string) bool {
 	return lookupBoolT(name, c.flagSet)
 }

-// String looks up the value of a local string flag, returns "" if no string flag exists
+// Looks up the value of a local string flag, returns "" if no string flag exists
 func (c *Context) String(name string) string {
 	return lookupString(name, c.flagSet)
 }

-// StringSlice looks up the value of a local string slice flag, returns nil if no
-// string slice flag exists
+// Looks up the value of a local string slice flag, returns nil if no string slice flag exists
 func (c *Context) StringSlice(name string) []string {
 	return lookupStringSlice(name, c.flagSet)
 }

-// IntSlice looks up the value of a local int slice flag, returns nil if no int
-// slice flag exists
+// Looks up the value of a local int slice flag, returns nil if no int slice flag exists
 func (c *Context) IntSlice(name string) []int {
 	return lookupIntSlice(name, c.flagSet)
 }

-// Int64Slice looks up the value of a local int slice flag, returns nil if no int
-// slice flag exists
-func (c *Context) Int64Slice(name string) []int64 {
-	return lookupInt64Slice(name, c.flagSet)
-}
-
-// Generic looks up the value of a local generic flag, returns nil if no generic
-// flag exists
+// Looks up the value of a local generic flag, returns nil if no generic flag exists
 func (c *Context) Generic(name string) interface{} {
 	return lookupGeneric(name, c.flagSet)
 }

-// GlobalInt looks up the value of a global int flag, returns 0 if no int flag exists
+// Looks up the value of a global int flag, returns 0 if no int flag exists
 func (c *Context) GlobalInt(name string) int {
 	if fs := lookupGlobalFlagSet(name, c); fs != nil {
 		return lookupInt(name, fs)
@@ -105,41 +79,7 @@ func (c *Context) GlobalInt(name string) int {
 	return 0
 }

-// GlobalInt64 looks up the value of a global int flag, returns 0 if no int flag exists
-func (c *Context) GlobalInt64(name string) int64 {
-	if fs := lookupGlobalFlagSet(name, c); fs != nil {
-		return lookupInt64(name, fs)
-	}
-	return 0
-}
-
-// GlobalUint looks up the value of a global int flag, returns 0 if no int flag exists
-func (c *Context) GlobalUint(name string) uint {
-	if fs := lookupGlobalFlagSet(name, c); fs != nil {
-		return lookupUint(name, fs)
-	}
-	return 0
-}
-
-// GlobalUint64 looks up the value of a global int flag, returns 0 if no int flag exists
-func (c *Context) GlobalUint64(name string) uint64 {
-	if fs := lookupGlobalFlagSet(name, c); fs != nil {
-		return lookupUint64(name, fs)
-	}
-	return 0
-}
-
-// GlobalFloat64 looks up the value of a global float64 flag, returns float64(0)
-// if no float64 flag exists
-func (c *Context) GlobalFloat64(name string) float64 {
-	if fs := lookupGlobalFlagSet(name, c); fs != nil {
-		return lookupFloat64(name, fs)
-	}
-	return float64(0)
-}
-
-// GlobalDuration looks up the value of a global time.Duration flag, returns 0
-// if no time.Duration flag exists
+// Looks up the value of a global time.Duration flag, returns 0 if no time.Duration flag exists
 func (c *Context) GlobalDuration(name string) time.Duration {
 	if fs := lookupGlobalFlagSet(name, c); fs != nil {
 		return lookupDuration(name, fs)
@@ -147,8 +87,7 @@ func (c *Context) GlobalDuration(name string) time.Duration {
 	return 0
 }

-// GlobalBool looks up the value of a global bool flag, returns false if no bool
-// flag exists
+// Looks up the value of a global bool flag, returns false if no bool flag exists
 func (c *Context) GlobalBool(name string) bool {
 	if fs := lookupGlobalFlagSet(name, c); fs != nil {
 		return lookupBool(name, fs)
@@ -156,17 +95,7 @@ func (c *Context) GlobalBool(name string) bool {
 	return false
 }

-// GlobalBoolT looks up the value of a global bool flag, returns true if no bool
-// flag exists
-func (c *Context) GlobalBoolT(name string) bool {
-	if fs := lookupGlobalFlagSet(name, c); fs != nil {
-		return lookupBoolT(name, fs)
-	}
-	return false
-}
-
-// GlobalString looks up the value of a global string flag, returns "" if no
-// string flag exists
+// Looks up the value of a global string flag, returns "" if no string flag exists
 func (c *Context) GlobalString(name string) string {
 	if fs := lookupGlobalFlagSet(name, c); fs != nil {
 		return lookupString(name, fs)
@@ -174,8 +103,7 @@ func (c *Context) GlobalString(name string) string {
 	return ""
 }

-// GlobalStringSlice looks up the value of a global string slice flag, returns
-// nil if no string slice flag exists
+// Looks up the value of a global string slice flag, returns nil if no string slice flag exists
 func (c *Context) GlobalStringSlice(name string) []string {
 	if fs := lookupGlobalFlagSet(name, c); fs != nil {
 		return lookupStringSlice(name, fs)
@@ -183,8 +111,7 @@ func (c *Context) GlobalStringSlice(name string) []string {
 	return nil
 }

-// GlobalIntSlice looks up the value of a global int slice flag, returns nil if
-// no int slice flag exists
+// Looks up the value of a global int slice flag, returns nil if no int slice flag exists
 func (c *Context) GlobalIntSlice(name string) []int {
 	if fs := lookupGlobalFlagSet(name, c); fs != nil {
 		return lookupIntSlice(name, fs)
@@ -192,17 +119,7 @@ func (c *Context) GlobalIntSlice(name string) []int {
 	return nil
 }

-// GlobalInt64Slice looks up the value of a global int slice flag, returns nil if
-// no int slice flag exists
-func (c *Context) GlobalInt64Slice(name string) []int64 {
-	if fs := lookupGlobalFlagSet(name, c); fs != nil {
-		return lookupInt64Slice(name, fs)
-	}
-	return nil
-}
-
-// GlobalGeneric looks up the value of a global generic flag, returns nil if no
-// generic flag exists
+// Looks up the value of a global generic flag, returns nil if no generic flag exists
 func (c *Context) GlobalGeneric(name string) interface{} {
 	if fs := lookupGlobalFlagSet(name, c); fs != nil {
 		return lookupGeneric(name, fs)
@@ -210,22 +127,12 @@ func (c *Context) GlobalGeneric(name string) interface{} {
 	return nil
 }

-// NumFlags returns the number of flags set
+// Returns the number of flags set
 func (c *Context) NumFlags() int {
 	return c.flagSet.NFlag()
 }

-// Set sets a context flag to a value.
-func (c *Context) Set(name, value string) error {
-	return c.flagSet.Set(name, value)
-}
-
-// GlobalSet sets a context flag to a value on the global flagset
-func (c *Context) GlobalSet(name, value string) error {
-	return globalContext(c).flagSet.Set(name, value)
-}
-
-// IsSet determines if the flag was actually set
+// Determines if the flag was actually set
 func (c *Context) IsSet(name string) bool {
 	if c.setFlags == nil {
 		c.setFlags = make(map[string]bool)
@@ -236,7 +143,7 @@ func (c *Context) IsSet(name string) bool {
 	return c.setFlags[name] == true
 }

-// GlobalIsSet determines if the global flag was actually set
+// Determines if the global flag was actually set
 func (c *Context) GlobalIsSet(name string) bool {
 	if c.globalSetFlags == nil {
 		c.globalSetFlags = make(map[string]bool)
@@ -253,7 +160,7 @@ func (c *Context) GlobalIsSet(name string) bool {
 	return c.globalSetFlags[name]
 }

-// FlagNames returns a slice of flag names used in this context.
+// Returns a slice of flag names used in this context.
 func (c *Context) FlagNames() (names []string) {
 	for _, flag := range c.Command.Flags {
 		name := strings.Split(flag.GetName(), ",")[0]
@@ -265,7 +172,7 @@ func (c *Context) FlagNames() (names []string) {
 	return
 }

-// GlobalFlagNames returns a slice of global flag names used by the app.
+// Returns a slice of global flag names used by the app.
 func (c *Context) GlobalFlagNames() (names []string) {
 	for _, flag := range c.App.Flags {
 		name := strings.Split(flag.GetName(), ",")[0]
@@ -277,26 +184,20 @@ func (c *Context) GlobalFlagNames() (names []string) {
 	return
 }

-// Parent returns the parent context, if any
+// Returns the parent context, if any
 func (c *Context) Parent() *Context {
 	return c.parentContext
 }

-// Args contains apps console arguments
 type Args []string

-// Args returns the command line arguments associated with the context.
+// Returns the command line arguments associated with the context.
 func (c *Context) Args() Args {
 	args := Args(c.flagSet.Args())
 	return args
 }

-// NArg returns the number of the command line arguments.
-func (c *Context) NArg() int {
-	return len(c.Args())
-}
-
-// Get returns the nth argument, or else a blank string
+// Returns the nth argument, or else a blank string
 func (a Args) Get(n int) string {
 	if len(a) > n {
 		return a[n]
@@ -304,12 +205,12 @@ func (a Args) Get(n int) string {
 	return ""
 }

-// First returns the first argument, or else a blank string
+// Returns the first argument, or else a blank string
 func (a Args) First() string {
 	return a.Get(0)
 }

-// Tail returns the rest of the arguments (not the first one)
+// Return the rest of the arguments (not the first one)
 // or else an empty string slice
 func (a Args) Tail() []string {
 	if len(a) >= 2 {
@@ -318,12 +219,12 @@ func (a Args) Tail() []string {
 	return []string{}
 }

-// Present checks if there are any arguments present
+// Checks if there are any arguments present
 func (a Args) Present() bool {
 	return len(a) != 0
 }

-// Swap swaps arguments at the given indexes
+// Swaps arguments at the given indexes
 func (a Args) Swap(from, to int) error {
 	if from >= len(a) || to >= len(a) {
 		return errors.New("index out of range")
@@ -332,19 +233,6 @@ func (a Args) Swap(from, to int) error {
 	return nil
 }

-func globalContext(ctx *Context) *Context {
-	if ctx == nil {
-		return nil
-	}
-
-	for {
-		if ctx.parentContext == nil {
-			return ctx
-		}
-		ctx = ctx.parentContext
-	}
-}
-
 func lookupGlobalFlagSet(name string, ctx *Context) *flag.FlagSet {
 	if ctx.parentContext != nil {
 		ctx = ctx.parentContext
@@ -360,46 +248,7 @@ func lookupGlobalFlagSet(name string, ctx *Context) *flag.FlagSet {
 func lookupInt(name string, set *flag.FlagSet) int {
 	f := set.Lookup(name)
 	if f != nil {
-		val, err := strconv.ParseInt(f.Value.String(), 0, 64)
-		if err != nil {
-			return 0
-		}
-		return int(val)
-	}
-
-	return 0
-}
-
-func lookupInt64(name string, set *flag.FlagSet) int64 {
-	f := set.Lookup(name)
-	if f != nil {
-		val, err := strconv.ParseInt(f.Value.String(), 0, 64)
-		if err != nil {
-			return 0
-		}
-		return val
-	}
-
-	return 0
-}
-
-func lookupUint(name string, set *flag.FlagSet) uint {
-	f := set.Lookup(name)
-	if f != nil {
-		val, err := strconv.ParseUint(f.Value.String(), 0, 64)
-		if err != nil {
-			return 0
-		}
-		return uint(val)
-	}
-
-	return 0
-}
-
-func lookupUint64(name string, set *flag.FlagSet) uint64 {
-	f := set.Lookup(name)
-	if f != nil {
-		val, err := strconv.ParseUint(f.Value.String(), 0, 64)
+		val, err := strconv.Atoi(f.Value.String())
 		if err != nil {
 			return 0
 		}
@@ -463,16 +312,6 @@ func lookupIntSlice(name string, set *flag.FlagSet) []int {
 	return nil
 }

-func lookupInt64Slice(name string, set *flag.FlagSet) []int64 {
-	f := set.Lookup(name)
-	if f != nil {
-		return (f.Value.(*Int64Slice)).Value()
-
-	}
-
-	return nil
-}
-
 func lookupGeneric(name string, set *flag.FlagSet) interface{} {
 	f := set.Lookup(name)
 	if f != nil {
--- a/Godeps/_workspace/src/github.com/codegangsta/cli/flag.go
+++ b/Godeps/_workspace/src/github.com/codegangsta/cli/flag.go
@@ -4,28 +4,24 @@ import (
 	"flag"
 	"fmt"
 	"os"
-	"reflect"
 	"runtime"
 	"strconv"
 	"strings"
 	"time"
 )

-const defaultPlaceholder = "value"
-
-// BashCompletionFlag enables bash-completion for all commands and subcommands
+// This flag enables bash-completion for all commands and subcommands
 var BashCompletionFlag = BoolFlag{
-	Name:   "generate-bash-completion",
-	Hidden: true,
+	Name: "generate-bash-completion",
 }

-// VersionFlag prints the version for the application
+// This flag prints the version for the application
 var VersionFlag = BoolFlag{
 	Name:  "version, v",
 	Usage: "print the version",
 }

-// HelpFlag prints the help for all commands and subcommands
+// This flag prints the help for all commands and subcommands
 // Set to the zero value (BoolFlag{}) to disable flag -- keeps subcommand
 // unless HideHelp is set to true)
 var HelpFlag = BoolFlag{
@@ -33,10 +29,6 @@ var HelpFlag = BoolFlag{
 	Usage: "show help",
 }

-// FlagStringer converts a flag definition to a string. This is used by help
-// to display a flag.
-var FlagStringer FlagStringFunc = stringifyFlag
-
 // Flag is a common interface related to parsing flags in cli.
 // For more advanced flag parsing techniques, it is recommended that
 // this interface be implemented.
@@ -76,14 +68,24 @@ type GenericFlag struct {
 	Value  Generic
 	Usage  string
 	EnvVar string
-	Hidden bool
 }

 // String returns the string representation of the generic flag to display the
 // help text to the user (uses the String() method of the generic flag to show
 // the value)
 func (f GenericFlag) String() string {
-	return FlagStringer(f)
+	return withEnvHint(f.EnvVar, fmt.Sprintf("%s %v\t%v", prefixedNames(f.Name), f.FormatValueHelp(), f.Usage))
+}
+
+func (f GenericFlag) FormatValueHelp() string {
+	if f.Value == nil {
+		return ""
+	}
+	s := f.Value.String()
+	if len(s) == 0 {
+		return ""
+	}
+	return fmt.Sprintf("\"%s\"", s)
 }

 // Apply takes the flagset and calls Set on the generic flag with the value
@@ -105,7 +107,6 @@ func (f GenericFlag) Apply(set *flag.FlagSet) {
 	})
 }

-// GetName returns the name of a flag.
 func (f GenericFlag) GetName() string {
 	return f.Name
 }
@@ -129,19 +130,20 @@ func (f *StringSlice) Value() []string {
 	return *f
 }

-// StringSliceFlag is a string flag that can be specified multiple times on the
+// StringSlice is a string flag that can be specified multiple times on the
 // command-line
 type StringSliceFlag struct {
 	Name   string
 	Value  *StringSlice
 	Usage  string
 	EnvVar string
-	Hidden bool
 }

 // String returns the usage
 func (f StringSliceFlag) String() string {
-	return FlagStringer(f)
+	firstName := strings.Trim(strings.Split(f.Name, ",")[0], " ")
+	pref := prefixFor(firstName)
+	return withEnvHint(f.EnvVar, fmt.Sprintf("%s [%v]\t%v", prefixedNames(f.Name), pref+firstName+" option "+pref+firstName+" option", f.Usage))
 }

 // Apply populates the flag given the flag set and environment
@@ -169,12 +171,11 @@ func (f StringSliceFlag) Apply(set *flag.FlagSet) {
 	})
 }

-// GetName returns the name of a flag.
 func (f StringSliceFlag) GetName() string {
 	return f.Name
 }

-// IntSlice is an opaque type for []int to satisfy flag.Value
+// StringSlice is an opaque type for []int to satisfy flag.Value
 type IntSlice []int

 // Set parses the value into an integer and appends it to the list of values
@@ -182,14 +183,15 @@ func (f *IntSlice) Set(value string) error {
 	tmp, err := strconv.Atoi(value)
 	if err != nil {
 		return err
+	} else {
+		*f = append(*f, tmp)
 	}
-	*f = append(*f, tmp)
 	return nil
 }

 // String returns a readable representation of this value (for usage defaults)
 func (f *IntSlice) String() string {
-	return fmt.Sprintf("%#v", *f)
+	return fmt.Sprintf("%d", *f)
 }

 // Value returns the slice of ints set by this flag
@@ -204,12 +206,13 @@ type IntSliceFlag struct {
 	Value  *IntSlice
 	Usage  string
 	EnvVar string
-	Hidden bool
 }

 // String returns the usage
 func (f IntSliceFlag) String() string {
-	return FlagStringer(f)
+	firstName := strings.Trim(strings.Split(f.Name, ",")[0], " ")
+	pref := prefixFor(firstName)
+	return withEnvHint(f.EnvVar, fmt.Sprintf("%s [%v]\t%v", prefixedNames(f.Name), pref+firstName+" option "+pref+firstName+" option", f.Usage))
 }

 // Apply populates the flag given the flag set and environment
@@ -223,7 +226,7 @@ func (f IntSliceFlag) Apply(set *flag.FlagSet) {
 					s = strings.TrimSpace(s)
 					err := newVal.Set(s)
 					if err != nil {
-						fmt.Fprintf(ErrWriter, err.Error())
+						fmt.Fprintf(os.Stderr, err.Error())
 					}
 				}
 				f.Value = newVal
@@ -240,94 +243,21 @@ func (f IntSliceFlag) Apply(set *flag.FlagSet) {
 	})
 }

-// GetName returns the name of the flag.
 func (f IntSliceFlag) GetName() string {
 	return f.Name
 }

-// Int64Slice is an opaque type for []int to satisfy flag.Value
-type Int64Slice []int64
-
-// Set parses the value into an integer and appends it to the list of values
-func (f *Int64Slice) Set(value string) error {
-	tmp, err := strconv.ParseInt(value, 10, 64)
-	if err != nil {
-		return err
-	}
-	*f = append(*f, tmp)
-	return nil
-}
-
-// String returns a readable representation of this value (for usage defaults)
-func (f *Int64Slice) String() string {
-	return fmt.Sprintf("%#v", *f)
-}
-
-// Value returns the slice of ints set by this flag
-func (f *Int64Slice) Value() []int64 {
-	return *f
-}
-
-// Int64SliceFlag is an int flag that can be specified multiple times on the
-// command-line
-type Int64SliceFlag struct {
-	Name   string
-	Value  *Int64Slice
-	Usage  string
-	EnvVar string
-	Hidden bool
-}
-
-// String returns the usage
-func (f Int64SliceFlag) String() string {
-	return FlagStringer(f)
-}
-
-// Apply populates the flag given the flag set and environment
-func (f Int64SliceFlag) Apply(set *flag.FlagSet) {
-	if f.EnvVar != "" {
-		for _, envVar := range strings.Split(f.EnvVar, ",") {
-			envVar = strings.TrimSpace(envVar)
-			if envVal := os.Getenv(envVar); envVal != "" {
-				newVal := &Int64Slice{}
-				for _, s := range strings.Split(envVal, ",") {
-					s = strings.TrimSpace(s)
-					err := newVal.Set(s)
-					if err != nil {
-						fmt.Fprintf(ErrWriter, err.Error())
-					}
-				}
-				f.Value = newVal
-				break
-			}
-		}
-	}
-
-	eachName(f.Name, func(name string) {
-		if f.Value == nil {
-			f.Value = &Int64Slice{}
-		}
-		set.Var(f.Value, name, f.Usage)
-	})
-}
-
-// GetName returns the name of the flag.
-func (f Int64SliceFlag) GetName() string {
-	return f.Name
-}
-
 // BoolFlag is a switch that defaults to false
 type BoolFlag struct {
 	Name        string
 	Usage       string
 	EnvVar      string
 	Destination *bool
-	Hidden      bool
 }

 // String returns a readable representation of this value (for usage defaults)
 func (f BoolFlag) String() string {
-	return FlagStringer(f)
+	return withEnvHint(f.EnvVar, fmt.Sprintf("%s\t%v", prefixedNames(f.Name), f.Usage))
 }

 // Apply populates the flag given the flag set and environment
@@ -355,7 +285,6 @@ func (f BoolFlag) Apply(set *flag.FlagSet) {
 	})
 }

-// GetName returns the name of the flag.
 func (f BoolFlag) GetName() string {
 	return f.Name
 }
@@ -367,12 +296,11 @@ type BoolTFlag struct {
 	Usage       string
 	EnvVar      string
 	Destination *bool
-	Hidden      bool
 }

 // String returns a readable representation of this value (for usage defaults)
 func (f BoolTFlag) String() string {
-	return FlagStringer(f)
+	return withEnvHint(f.EnvVar, fmt.Sprintf("%s\t%v", prefixedNames(f.Name), f.Usage))
 }

 // Apply populates the flag given the flag set and environment
@@ -400,7 +328,6 @@ func (f BoolTFlag) Apply(set *flag.FlagSet) {
 	})
 }

-// GetName returns the name of the flag.
 func (f BoolTFlag) GetName() string {
 	return f.Name
 }
@@ -412,12 +339,19 @@ type StringFlag struct {
 	Usage       string
 	EnvVar      string
 	Destination *string
-	Hidden      bool
 }

 // String returns the usage
 func (f StringFlag) String() string {
-	return FlagStringer(f)
+	return withEnvHint(f.EnvVar, fmt.Sprintf("%s %v\t%v", prefixedNames(f.Name), f.FormatValueHelp(), f.Usage))
+}
+
+func (f StringFlag) FormatValueHelp() string {
+	s := f.Value
+	if len(s) == 0 {
+		return ""
+	}
+	return fmt.Sprintf("\"%s\"", s)
 }

 // Apply populates the flag given the flag set and environment
@@ -441,24 +375,23 @@ func (f StringFlag) Apply(set *flag.FlagSet) {
 	})
 }

-// GetName returns the name of the flag.
 func (f StringFlag) GetName() string {
 	return f.Name
 }

 // IntFlag is a flag that takes an integer
+// Errors if the value provided cannot be parsed
 type IntFlag struct {
 	Name        string
 	Value       int
 	Usage       string
 	EnvVar      string
 	Destination *int
-	Hidden      bool
 }

 // String returns the usage
 func (f IntFlag) String() string {
-	return FlagStringer(f)
+	return withEnvHint(f.EnvVar, fmt.Sprintf("%s \"%v\"\t%v", prefixedNames(f.Name), f.Value, f.Usage))
 }

 // Apply populates the flag given the flag set and environment
@@ -485,143 +418,10 @@ func (f IntFlag) Apply(set *flag.FlagSet) {
 	})
 }

-// GetName returns the name of the flag.
 func (f IntFlag) GetName() string {
 	return f.Name
 }

-// Int64Flag is a flag that takes a 64-bit integer
-type Int64Flag struct {
-	Name        string
-	Value       int64
-	Usage       string
-	EnvVar      string
-	Destination *int64
-	Hidden      bool
-}
-
-// String returns the usage
-func (f Int64Flag) String() string {
-	return FlagStringer(f)
-}
-
-// Apply populates the flag given the flag set and environment
-func (f Int64Flag) Apply(set *flag.FlagSet) {
-	if f.EnvVar != "" {
-		for _, envVar := range strings.Split(f.EnvVar, ",") {
-			envVar = strings.TrimSpace(envVar)
-			if envVal := os.Getenv(envVar); envVal != "" {
-				envValInt, err := strconv.ParseInt(envVal, 0, 64)
-				if err == nil {
-					f.Value = envValInt
-					break
-				}
-			}
-		}
-	}
-
-	eachName(f.Name, func(name string) {
-		if f.Destination != nil {
-			set.Int64Var(f.Destination, name, f.Value, f.Usage)
-			return
-		}
-		set.Int64(name, f.Value, f.Usage)
-	})
-}
-
-// GetName returns the name of the flag.
-func (f Int64Flag) GetName() string {
-	return f.Name
-}
-
-// UintFlag is a flag that takes an unsigned integer
-type UintFlag struct {
-	Name        string
-	Value       uint
-	Usage       string
-	EnvVar      string
-	Destination *uint
-	Hidden      bool
-}
-
-// String returns the usage
-func (f UintFlag) String() string {
-	return FlagStringer(f)
-}
-
-// Apply populates the flag given the flag set and environment
-func (f UintFlag) Apply(set *flag.FlagSet) {
-	if f.EnvVar != "" {
-		for _, envVar := range strings.Split(f.EnvVar, ",") {
-			envVar = strings.TrimSpace(envVar)
-			if envVal := os.Getenv(envVar); envVal != "" {
-				envValInt, err := strconv.ParseUint(envVal, 0, 64)
-				if err == nil {
-					f.Value = uint(envValInt)
-					break
-				}
-			}
-		}
-	}
-
-	eachName(f.Name, func(name string) {
-		if f.Destination != nil {
-			set.UintVar(f.Destination, name, f.Value, f.Usage)
-			return
-		}
-		set.Uint(name, f.Value, f.Usage)
-	})
-}
-
-// GetName returns the name of the flag.
-func (f UintFlag) GetName() string {
-	return f.Name
-}
-
-// Uint64Flag is a flag that takes an unsigned 64-bit integer
-type Uint64Flag struct {
-	Name        string
-	Value       uint64
-	Usage       string
-	EnvVar      string
-	Destination *uint64
-	Hidden      bool
-}
-
-// String returns the usage
-func (f Uint64Flag) String() string {
-	return FlagStringer(f)
-}
-
-// Apply populates the flag given the flag set and environment
-func (f Uint64Flag) Apply(set *flag.FlagSet) {
-	if f.EnvVar != "" {
-		for _, envVar := range strings.Split(f.EnvVar, ",") {
-			envVar = strings.TrimSpace(envVar)
-			if envVal := os.Getenv(envVar); envVal != "" {
-				envValInt, err := strconv.ParseUint(envVal, 0, 64)
-				if err == nil {
-					f.Value = uint64(envValInt)
-					break
-				}
-			}
-		}
-	}
-
-	eachName(f.Name, func(name string) {
-		if f.Destination != nil {
-			set.Uint64Var(f.Destination, name, f.Value, f.Usage)
-			return
-		}
-		set.Uint64(name, f.Value, f.Usage)
-	})
-}
-
-// GetName returns the name of the flag.
-func (f Uint64Flag) GetName() string {
-	return f.Name
-}
-
 // DurationFlag is a flag that takes a duration specified in Go's duration
 // format: https://golang.org/pkg/time/#ParseDuration
 type DurationFlag struct {
@@ -630,12 +430,11 @@ type DurationFlag struct {
 	Usage       string
 	EnvVar      string
 	Destination *time.Duration
-	Hidden      bool
 }

 // String returns a readable representation of this value (for usage defaults)
 func (f DurationFlag) String() string {
-	return FlagStringer(f)
+	return withEnvHint(f.EnvVar, fmt.Sprintf("%s \"%v\"\t%v", prefixedNames(f.Name), f.Value, f.Usage))
 }

 // Apply populates the flag given the flag set and environment
@@ -662,24 +461,23 @@ func (f DurationFlag) Apply(set *flag.FlagSet) {
 	})
 }

-// GetName returns the name of the flag.
 func (f DurationFlag) GetName() string {
 	return f.Name
 }

 // Float64Flag is a flag that takes an float value
+// Errors if the value provided cannot be parsed
 type Float64Flag struct {
 	Name        string
 	Value       float64
 	Usage       string
 	EnvVar      string
 	Destination *float64
-	Hidden      bool
 }

 // String returns the usage
 func (f Float64Flag) String() string {
-	return FlagStringer(f)
+	return withEnvHint(f.EnvVar, fmt.Sprintf("%s \"%v\"\t%v", prefixedNames(f.Name), f.Value, f.Usage))
 }

 // Apply populates the flag given the flag set and environment
@@ -705,21 +503,10 @@ func (f Float64Flag) Apply(set *flag.FlagSet) {
 	})
 }

-// GetName returns the name of the flag.
 func (f Float64Flag) GetName() string {
 	return f.Name
 }

-func visibleFlags(fl []Flag) []Flag {
-	visible := []Flag{}
-	for _, flag := range fl {
-		if !flagValue(flag).FieldByName("Hidden").Bool() {
-			visible = append(visible, flag)
-		}
-	}
-	return visible
-}
-
 func prefixFor(name string) (prefix string) {
 	if len(name) == 1 {
 		prefix = "-"
@@ -730,37 +517,16 @@ func prefixFor(name string) (prefix string) {
 	return
 }

-// Returns the placeholder, if any, and the unquoted usage string.
-func unquoteUsage(usage string) (string, string) {
-	for i := 0; i < len(usage); i++ {
-		if usage[i] == '`' {
-			for j := i + 1; j < len(usage); j++ {
-				if usage[j] == '`' {
-					name := usage[i+1 : j]
-					usage = usage[:i] + name + usage[j+1:]
-					return name, usage
-				}
-			}
-			break
-		}
-	}
-	return "", usage
-}
-
-func prefixedNames(fullName, placeholder string) string {
-	var prefixed string
+func prefixedNames(fullName string) (prefixed string) {
 	parts := strings.Split(fullName, ",")
 	for i, name := range parts {
 		name = strings.Trim(name, " ")
 		prefixed += prefixFor(name) + name
-		if placeholder != "" {
-			prefixed += " " + placeholder
-		}
 		if i < len(parts)-1 {
 			prefixed += ", "
 		}
 	}
-	return prefixed
+	return
 }

 func withEnvHint(envVar, str string) string {
@@ -778,105 +544,3 @@ func withEnvHint(envVar, str string) string {
 	}
 	return str + envText
 }
-
-func flagValue(f Flag) reflect.Value {
-	fv := reflect.ValueOf(f)
-	for fv.Kind() == reflect.Ptr {
-		fv = reflect.Indirect(fv)
-	}
-	return fv
-}
-
-func stringifyFlag(f Flag) string {
-	fv := flagValue(f)
-
-	switch f.(type) {
-	case IntSliceFlag:
-		return withEnvHint(fv.FieldByName("EnvVar").String(),
-			stringifyIntSliceFlag(f.(IntSliceFlag)))
-	case Int64SliceFlag:
-		return withEnvHint(fv.FieldByName("EnvVar").String(),
-			stringifyInt64SliceFlag(f.(Int64SliceFlag)))
-	case StringSliceFlag:
-		return withEnvHint(fv.FieldByName("EnvVar").String(),
-			stringifyStringSliceFlag(f.(StringSliceFlag)))
-	}
-
-	placeholder, usage := unquoteUsage(fv.FieldByName("Usage").String())
-
-	needsPlaceholder := false
-	defaultValueString := ""
-	val := fv.FieldByName("Value")
-
-	if val.IsValid() {
-		needsPlaceholder = true
-		defaultValueString = fmt.Sprintf(" (default: %v)", val.Interface())
-
-		if val.Kind() == reflect.String && val.String() != "" {
-			defaultValueString = fmt.Sprintf(" (default: %q)", val.String())
-		}
-	}
-
-	if defaultValueString == " (default: )" {
-		defaultValueString = ""
-	}
-
-	if needsPlaceholder && placeholder == "" {
-		placeholder = defaultPlaceholder
-	}
-
-	usageWithDefault := strings.TrimSpace(fmt.Sprintf("%s%s", usage, defaultValueString))
-
-	return withEnvHint(fv.FieldByName("EnvVar").String(),
-		fmt.Sprintf("%s\t%s", prefixedNames(fv.FieldByName("Name").String(), placeholder), usageWithDefault))
-}
-
-func stringifyIntSliceFlag(f IntSliceFlag) string {
-	defaultVals := []string{}
-	if f.Value != nil && len(f.Value.Value()) > 0 {
-		for _, i := range f.Value.Value() {
-			defaultVals = append(defaultVals, fmt.Sprintf("%d", i))
-		}
-	}
-
-	return stringifySliceFlag(f.Usage, f.Name, defaultVals)
-}
-
-func stringifyInt64SliceFlag(f Int64SliceFlag) string {
-	defaultVals := []string{}
-	if f.Value != nil && len(f.Value.Value()) > 0 {
-		for _, i := range f.Value.Value() {
-			defaultVals = append(defaultVals, fmt.Sprintf("%d", i))
-		}
-	}
-
-	return stringifySliceFlag(f.Usage, f.Name, defaultVals)
-}
-
-func stringifyStringSliceFlag(f StringSliceFlag) string {
-	defaultVals := []string{}
-	if f.Value != nil && len(f.Value.Value()) > 0 {
-		for _, s := range f.Value.Value() {
-			if len(s) > 0 {
-				defaultVals = append(defaultVals, fmt.Sprintf("%q", s))
-			}
-		}
-	}
-
-	return stringifySliceFlag(f.Usage, f.Name, defaultVals)
-}
-
-func stringifySliceFlag(usage, name string, defaultVals []string) string {
-	placeholder, usage := unquoteUsage(usage)
-	if placeholder == "" {
-		placeholder = defaultPlaceholder
-	}
-
-	defaultVal := ""
-	if len(defaultVals) > 0 {
-		defaultVal = fmt.Sprintf(" (default: %s)", strings.Join(defaultVals, ", "))
-	}
-
-	usageWithDefault := strings.TrimSpace(fmt.Sprintf("%s%s", usage, defaultVal))
-	return fmt.Sprintf("%s\t%s", prefixedNames(name, placeholder), usageWithDefault)
-}
--- a/Godeps/_workspace/src/github.com/codegangsta/cli/help.go
+++ b/Godeps/_workspace/src/github.com/codegangsta/cli/help.go
@@ -3,74 +3,68 @@ package cli
 import (
 	"fmt"
 	"io"
-	"os"
 	"strings"
 	"text/tabwriter"
 	"text/template"
 )

-// AppHelpTemplate is the text template for the Default help topic.
+// The text template for the Default help topic.
 // cli.go uses text/template to render templates. You can
 // render custom help text by setting this variable.
 var AppHelpTemplate = `NAME:
   {{.Name}} - {{.Usage}}

 USAGE:
-   {{if .UsageText}}{{.UsageText}}{{else}}{{.HelpName}} {{if .VisibleFlags}}[global options]{{end}}{{if .Commands}} command [command options]{{end}} {{if .ArgsUsage}}{{.ArgsUsage}}{{else}}[arguments...]{{end}}{{end}}
-   {{if .Version}}{{if not .HideVersion}}
+   {{if .UsageText}}{{.UsageText}}{{else}}{{.HelpName}} {{if .Flags}}[global options]{{end}}{{if .Commands}} command [command options]{{end}} {{if .ArgsUsage}}{{.ArgsUsage}}{{else}}[arguments...]{{end}}{{end}}
+   {{if .Version}}
 VERSION:
   {{.Version}}
-   {{end}}{{end}}{{if len .Authors}}
+   {{end}}{{if len .Authors}}
 AUTHOR(S):
-   {{range .Authors}}{{.}}{{end}}
-   {{end}}{{if .VisibleCommands}}
-COMMANDS:{{range .VisibleCategories}}{{if .Name}}
-   {{.Name}}:{{end}}{{range .VisibleCommands}}
-     {{join .Names ", "}}{{"\t"}}{{.Usage}}{{end}}
-{{end}}{{end}}{{if .VisibleFlags}}
+   {{range .Authors}}{{ . }}{{end}}
+   {{end}}{{if .Commands}}
+COMMANDS:
+   {{range .Commands}}{{join .Names ", "}}{{ "\t" }}{{.Usage}}
+   {{end}}{{end}}{{if .Flags}}
 GLOBAL OPTIONS:
-   {{range .VisibleFlags}}{{.}}
-   {{end}}{{end}}{{if .Copyright}}
+   {{range .Flags}}{{.}}
+   {{end}}{{end}}{{if .Copyright }}
 COPYRIGHT:
   {{.Copyright}}
   {{end}}
 `

-// CommandHelpTemplate is the text template for the command help topic.
+// The text template for the command help topic.
 // cli.go uses text/template to render templates. You can
 // render custom help text by setting this variable.
 var CommandHelpTemplate = `NAME:
   {{.HelpName}} - {{.Usage}}

 USAGE:
-   {{.HelpName}}{{if .VisibleFlags}} [command options]{{end}} {{if .ArgsUsage}}{{.ArgsUsage}}{{else}}[arguments...]{{end}}{{if .Category}}
-
-CATEGORY:
-   {{.Category}}{{end}}{{if .Description}}
+   {{.HelpName}}{{if .Flags}} [command options]{{end}} {{if .ArgsUsage}}{{.ArgsUsage}}{{else}}[arguments...]{{end}}{{if .Description}}

 DESCRIPTION:
-   {{.Description}}{{end}}{{if .VisibleFlags}}
+   {{.Description}}{{end}}{{if .Flags}}

 OPTIONS:
-   {{range .VisibleFlags}}{{.}}
-   {{end}}{{end}}
+   {{range .Flags}}{{.}}
+   {{end}}{{ end }}
 `

-// SubcommandHelpTemplate is the text template for the subcommand help topic.
+// The text template for the subcommand help topic.
 // cli.go uses text/template to render templates. You can
 // render custom help text by setting this variable.
 var SubcommandHelpTemplate = `NAME:
   {{.HelpName}} - {{.Usage}}

 USAGE:
-   {{.HelpName}} command{{if .VisibleFlags}} [command options]{{end}} {{if .ArgsUsage}}{{.ArgsUsage}}{{else}}[arguments...]{{end}}
+   {{.HelpName}} command{{if .Flags}} [command options]{{end}} {{if .ArgsUsage}}{{.ArgsUsage}}{{else}}[arguments...]{{end}}

-COMMANDS:{{range .VisibleCategories}}{{if .Name}}
-   {{.Name}}:{{end}}{{range .VisibleCommands}}
-     {{join .Names ", "}}{{"\t"}}{{.Usage}}{{end}}
-{{end}}{{if .VisibleFlags}}
+COMMANDS:
+   {{range .Commands}}{{join .Names ", "}}{{ "\t" }}{{.Usage}}
+   {{end}}{{if .Flags}}
 OPTIONS:
-   {{range .VisibleFlags}}{{.}}
+   {{range .Flags}}{{.}}
   {{end}}{{end}}
 `

@@ -79,14 +73,13 @@ var helpCommand = Command{
 	Aliases:   []string{"h"},
 	Usage:     "Shows a list of commands or help for one command",
 	ArgsUsage: "[command]",
-	Action: func(c *Context) error {
+	Action: func(c *Context) {
 		args := c.Args()
 		if args.Present() {
-			return ShowCommandHelp(c, args.First())
+			ShowCommandHelp(c, args.First())
+		} else {
+			ShowAppHelp(c)
 		}
-
-		ShowAppHelp(c)
-		return nil
 	},
 }

@@ -95,74 +88,65 @@ var helpSubcommand = Command{
 	Aliases:   []string{"h"},
 	Usage:     "Shows a list of commands or help for one command",
 	ArgsUsage: "[command]",
-	Action: func(c *Context) error {
+	Action: func(c *Context) {
 		args := c.Args()
 		if args.Present() {
-			return ShowCommandHelp(c, args.First())
+			ShowCommandHelp(c, args.First())
+		} else {
+			ShowSubcommandHelp(c)
 		}
-
-		return ShowSubcommandHelp(c)
 	},
 }

 // Prints help for the App or Command
 type helpPrinter func(w io.Writer, templ string, data interface{})

-// HelpPrinter is a function that writes the help output. If not set a default
-// is used. The function signature is:
-// func(w io.Writer, templ string, data interface{})
 var HelpPrinter helpPrinter = printHelp

-// VersionPrinter prints the version for the App
+// Prints version for the App
 var VersionPrinter = printVersion

-// ShowAppHelp is an action that displays the help.
-func ShowAppHelp(c *Context) error {
+func ShowAppHelp(c *Context) {
 	HelpPrinter(c.App.Writer, AppHelpTemplate, c.App)
-	return nil
 }

-// DefaultAppComplete prints the list of subcommands as the default app completion method
+// Prints the list of subcommands as the default app completion method
 func DefaultAppComplete(c *Context) {
 	for _, command := range c.App.Commands {
-		if command.Hidden {
-			continue
-		}
 		for _, name := range command.Names() {
 			fmt.Fprintln(c.App.Writer, name)
 		}
 	}
 }

-// ShowCommandHelp prints help for the given command
-func ShowCommandHelp(ctx *Context, command string) error {
+// Prints help for the given command
+func ShowCommandHelp(ctx *Context, command string) {
 	// show the subcommand help for a command with subcommands
 	if command == "" {
 		HelpPrinter(ctx.App.Writer, SubcommandHelpTemplate, ctx.App)
-		return nil
+		return
 	}

 	for _, c := range ctx.App.Commands {
 		if c.HasName(command) {
 			HelpPrinter(ctx.App.Writer, CommandHelpTemplate, c)
-			return nil
+			return
 		}
 	}

-	if ctx.App.CommandNotFound == nil {
-		return NewExitError(fmt.Sprintf("No help topic for '%v'", command), 3)
+	if ctx.App.CommandNotFound != nil {
+		ctx.App.CommandNotFound(ctx, command)
+	} else {
+		fmt.Fprintf(ctx.App.Writer, "No help topic for '%v'\n", command)
 	}
-
-	ctx.App.CommandNotFound(ctx, command)
-	return nil
 }

-// ShowSubcommandHelp prints help for the given subcommand
-func ShowSubcommandHelp(c *Context) error {
-	return ShowCommandHelp(c, c.Command.Name)
+// Prints help for the given subcommand
+func ShowSubcommandHelp(c *Context) {
+	ShowCommandHelp(c, c.Command.Name)
 }

-// ShowVersion prints the version number of the App
+// Prints the version number of the App
 func ShowVersion(c *Context) {
 	VersionPrinter(c)
 }
@@ -171,7 +155,7 @@ func printVersion(c *Context) {
 	fmt.Fprintf(c.App.Writer, "%v version %v\n", c.App.Name, c.App.Version)
 }

-// ShowCompletions prints the lists of commands within a given context
+// Prints the lists of commands within a given context
 func ShowCompletions(c *Context) {
 	a := c.App
 	if a != nil && a.BashComplete != nil {
@@ -179,7 +163,7 @@ func ShowCompletions(c *Context) {
 	}
 }

-// ShowCommandCompletions prints the custom completions for a given command
+// Prints the custom completions for a given command
 func ShowCommandCompletions(ctx *Context, command string) {
 	c := ctx.App.Command(command)
 	if c != nil && c.BashComplete != nil {
@@ -192,15 +176,12 @@ func printHelp(out io.Writer, templ string, data interface{}) {
 		"join": strings.Join,
 	}

-	w := tabwriter.NewWriter(out, 1, 8, 2, ' ', 0)
+	w := tabwriter.NewWriter(out, 0, 8, 1, '\t', 0)
 	t := template.Must(template.New("help").Funcs(funcMap).Parse(templ))
 	err := t.Execute(w, data)
 	if err != nil {
 		// If the writer is closed, t.Execute will fail, and there's nothing
-		// we can do to recover.
-		if os.Getenv("CLI_TEMPLATE_ERROR_DEBUG") != "" {
-			fmt.Fprintf(ErrWriter, "CLI TEMPLATE ERROR: %#v\n", err)
-		}
+		// we can do to recover. We could send this to os.Stderr if we need.
 		return
 	}
 	w.Flush()
--- a/Godeps/_workspace/src/github.com/davecgh/go-spew/LICENSE
+++ b/Godeps/_workspace/src/github.com/davecgh/go-spew/LICENSE
@@ -1,6 +1,4 @@
-ISC License
-
-Copyright (c) 2012-2016 Dave Collins <dave@davec.name>
+Copyright (c) 2012-2013 Dave Collins <dave@davec.name>

 Permission to use, copy, modify, and distribute this software for any
 purpose with or without fee is hereby granted, provided that the above
--- a/Godeps/_workspace/src/github.com/davecgh/go-spew/spew/bypass.go
+++ b/Godeps/_workspace/src/github.com/davecgh/go-spew/spew/bypass.go
@@ -1,4 +1,4 @@
-// Copyright (c) 2015-2016 Dave Collins <dave@davec.name>
+// Copyright (c) 2015 Dave Collins <dave@davec.name>
 //
 // Permission to use, copy, modify, and distribute this software for any
 // purpose with or without fee is hereby granted, provided that the above
@@ -13,10 +13,9 @@
 // OR IN CONNECTION WITH THE USE OR PERFORMANCE OF THIS SOFTWARE.

 // NOTE: Due to the following build constraints, this file will only be compiled
-// when the code is not running on Google App Engine, compiled by GopherJS, and
-// "-tags safe" is not added to the go build command line.  The "disableunsafe"
-// tag is deprecated and thus should not be used.
-// +build !js,!appengine,!safe,!disableunsafe
+// when the code is not running on Google App Engine and "-tags disableunsafe"
+// is not added to the go build command line.
+// +build !appengine,!disableunsafe

 package spew

--- a/Godeps/_workspace/src/github.com/davecgh/go-spew/spew/bypasssafe.go
+++ b/Godeps/_workspace/src/github.com/davecgh/go-spew/spew/bypasssafe.go
@@ -1,4 +1,4 @@
-// Copyright (c) 2015-2016 Dave Collins <dave@davec.name>
+// Copyright (c) 2015 Dave Collins <dave@davec.name>
 //
 // Permission to use, copy, modify, and distribute this software for any
 // purpose with or without fee is hereby granted, provided that the above
@@ -13,10 +13,9 @@
 // OR IN CONNECTION WITH THE USE OR PERFORMANCE OF THIS SOFTWARE.

 // NOTE: Due to the following build constraints, this file will only be compiled
-// when the code is running on Google App Engine, compiled by GopherJS, or
-// "-tags safe" is added to the go build command line.  The "disableunsafe"
-// tag is deprecated and thus should not be used.
-// +build js appengine safe disableunsafe
+// when either the code is running on Google App Engine or "-tags disableunsafe"
+// is added to the go build command line.
+// +build appengine disableunsafe

 package spew

--- a/Godeps/_workspace/src/github.com/davecgh/go-spew/spew/common.go
+++ b/Godeps/_workspace/src/github.com/davecgh/go-spew/spew/common.go
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2013-2016 Dave Collins <dave@davec.name>
+ * Copyright (c) 2013 Dave Collins <dave@davec.name>
 *
 * Permission to use, copy, modify, and distribute this software for any
 * purpose with or without fee is hereby granted, provided that the above
--- a/Godeps/_workspace/src/github.com/davecgh/go-spew/spew/config.go
+++ b/Godeps/_workspace/src/github.com/davecgh/go-spew/spew/config.go
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2013-2016 Dave Collins <dave@davec.name>
+ * Copyright (c) 2013 Dave Collins <dave@davec.name>
 *
 * Permission to use, copy, modify, and distribute this software for any
 * purpose with or without fee is hereby granted, provided that the above
@@ -64,18 +64,9 @@ type ConfigState struct {
 	// inside these interface methods.  As a result, this option relies on
 	// access to the unsafe package, so it will not have any effect when
 	// running in environments without access to the unsafe package such as
-	// Google App Engine or with the "safe" build tag specified.
+	// Google App Engine or with the "disableunsafe" build tag specified.
 	DisablePointerMethods bool

-	// DisablePointerAddresses specifies whether to disable the printing of
-	// pointer addresses. This is useful when diffing data structures in tests.
-	DisablePointerAddresses bool
-
-	// DisableCapacities specifies whether to disable the printing of capacities
-	// for arrays, slices, maps and channels. This is useful when diffing
-	// data structures in tests.
-	DisableCapacities bool
-
 	// ContinueOnMethod specifies whether or not recursion should continue once
 	// a custom error or Stringer interface is invoked.  The default, false,
 	// means it will print the results of invoking the custom error or Stringer
--- a/Godeps/_workspace/src/github.com/davecgh/go-spew/spew/doc.go
+++ b/Godeps/_workspace/src/github.com/davecgh/go-spew/spew/doc.go
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2013-2016 Dave Collins <dave@davec.name>
+ * Copyright (c) 2013 Dave Collins <dave@davec.name>
 *
 * Permission to use, copy, modify, and distribute this software for any
 * purpose with or without fee is hereby granted, provided that the above
@@ -91,15 +91,6 @@ The following configuration options are available:
 		which only accept pointer receivers from non-pointer variables.
 		Pointer method invocation is enabled by default.

-	* DisablePointerAddresses
-		DisablePointerAddresses specifies whether to disable the printing of
-		pointer addresses. This is useful when diffing data structures in tests.
-
-	* DisableCapacities
-		DisableCapacities specifies whether to disable the printing of
-		capacities for arrays, slices, maps and channels. This is useful when
-		diffing data structures in tests.
-
 	* ContinueOnMethod
 		Enables recursion into types after invoking error and Stringer interface
 		methods. Recursion after method invocation is disabled by default.
--- a/Godeps/_workspace/src/github.com/davecgh/go-spew/spew/dump.go
+++ b/Godeps/_workspace/src/github.com/davecgh/go-spew/spew/dump.go
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2013-2016 Dave Collins <dave@davec.name>
+ * Copyright (c) 2013 Dave Collins <dave@davec.name>
 *
 * Permission to use, copy, modify, and distribute this software for any
 * purpose with or without fee is hereby granted, provided that the above
@@ -129,7 +129,7 @@ func (d *dumpState) dumpPtr(v reflect.Value) {
 	d.w.Write(closeParenBytes)

 	// Display pointer information.
-	if !d.cs.DisablePointerAddresses && len(pointerChain) > 0 {
+	if len(pointerChain) > 0 {
 		d.w.Write(openParenBytes)
 		for i, addr := range pointerChain {
 			if i > 0 {
@@ -282,13 +282,13 @@ func (d *dumpState) dump(v reflect.Value) {
 	case reflect.Map, reflect.String:
 		valueLen = v.Len()
 	}
-	if valueLen != 0 || !d.cs.DisableCapacities && valueCap != 0 {
+	if valueLen != 0 || valueCap != 0 {
 		d.w.Write(openParenBytes)
 		if valueLen != 0 {
 			d.w.Write(lenEqualsBytes)
 			printInt(d.w, int64(valueLen), 10)
 		}
-		if !d.cs.DisableCapacities && valueCap != 0 {
+		if valueCap != 0 {
 			if valueLen != 0 {
 				d.w.Write(spaceBytes)
 			}
--- a/Godeps/_workspace/src/github.com/davecgh/go-spew/spew/format.go
+++ b/Godeps/_workspace/src/github.com/davecgh/go-spew/spew/format.go
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2013-2016 Dave Collins <dave@davec.name>
+ * Copyright (c) 2013 Dave Collins <dave@davec.name>
 *
 * Permission to use, copy, modify, and distribute this software for any
 * purpose with or without fee is hereby granted, provided that the above
--- a/Godeps/_workspace/src/github.com/davecgh/go-spew/spew/spew.go
+++ b/Godeps/_workspace/src/github.com/davecgh/go-spew/spew/spew.go
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2013-2016 Dave Collins <dave@davec.name>
+ * Copyright (c) 2013 Dave Collins <dave@davec.name>
 *
 * Permission to use, copy, modify, and distribute this software for any
 * purpose with or without fee is hereby granted, provided that the above
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/.gitignore
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/.gitignore
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/.travis.yml
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/.travis.yml
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/CMakeLists.txt
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/CMakeLists.txt
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/MANIFEST.in
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/MANIFEST.in
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/Makefile
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/Makefile
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/README.md
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/README.md
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/Vagrantfile
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/Vagrantfile
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/appveyor.yml
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/appveyor.yml
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/ethash.go
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/ethash.go
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/ethash_opencl.go
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/ethash_opencl.go
@@ -0,0 +1,628 @@
+// Copyright 2014 The go-ethereum Authors
+// This file is part of the go-ethereum library.
+//
+// The go-ethereum library is free software: you can redistribute it and/or modify
+// it under the terms of the GNU Lesser General Public License as published by
+// the Free Software Foundation, either version 3 of the License, or
+// (at your option) any later version.
+//
+// The go-ethereum library is distributed in the hope that it will be useful,
+// but WITHOUT ANY WARRANTY; without even the implied warranty of
+// MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+// GNU Lesser General Public License for more details.
+//
+// You should have received a copy of the GNU Lesser General Public License
+// along with the go-ethereum library. If not, see <http://www.gnu.org/licenses/>.
+
+// +build opencl
+
+package ethash
+
+//#cgo LDFLAGS: -w
+//#include <stdint.h>
+//#include <string.h>
+//#include "src/libethash/internal.h"
+import "C"
+
+import (
+	crand "crypto/rand"
+	"encoding/binary"
+	"fmt"
+	"math"
+	"math/big"
+	mrand "math/rand"
+	"strconv"
+	"strings"
+	"sync"
+	"sync/atomic"
+	"time"
+	"unsafe"
+
+	"github.com/Gustav-Simonsson/go-opencl/cl"
+	"github.com/ethereum/go-ethereum/common"
+	"github.com/ethereum/go-ethereum/pow"
+)
+
+/*
+
+  This code have two main entry points:
+
+  1. The initCL(...)  function configures one or more OpenCL device
+     (for now only GPU) and loads the Ethash DAG onto device memory
+
+  2. The Search(...) function loads a Ethash nonce into device(s) memory and
+     executes the Ethash OpenCL kernel.
+
+  Throughout the code, we refer to "host memory" and "device memory".
+  For most systems (e.g. regular PC GPU miner) the host memory is RAM and
+  device memory is the GPU global memory (e.g. GDDR5).
+
+  References mentioned in code comments:
+
+  1. https://github.com/ethereum/wiki/wiki/Ethash
+  2. https://github.com/ethereum/cpp-ethereum/blob/develop/libethash-cl/ethash_cl_miner.cpp
+  3. https://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/
+  4. http://amd-dev.wpengine.netdna-cdn.com/wordpress/media/2013/12/AMD_OpenCL_Programming_User_Guide.pdf
+
+*/
+
+type OpenCLDevice struct {
+	deviceId int
+	device   *cl.Device
+	openCL11 bool // OpenCL version 1.1 and 1.2 are handled a bit different
+	openCL12 bool
+
+	dagBuf        *cl.MemObject // Ethash full DAG in device mem
+	headerBuf     *cl.MemObject // Hash of block-to-mine in device mem
+	searchBuffers []*cl.MemObject
+
+	searchKernel *cl.Kernel
+	hashKernel   *cl.Kernel
+
+	queue         *cl.CommandQueue
+	ctx           *cl.Context
+	workGroupSize int
+
+	nonceRand *mrand.Rand // seeded by crypto/rand, see comments where it's initialised
+	result    common.Hash
+}
+
+type OpenCLMiner struct {
+	mu sync.Mutex
+
+	ethash *Ethash // Ethash full DAG & cache in host mem
+
+	deviceIds []int
+	devices   []*OpenCLDevice
+
+	dagSize uint64
+
+	hashRate int32 // Go atomics & uint64 have some issues; int32 is supported on all platforms
+}
+
+type pendingSearch struct {
+	bufIndex   uint32
+	startNonce uint64
+}
+
+const (
+	SIZEOF_UINT32 = 4
+
+	// See [1]
+	ethashMixBytesLen = 128
+	ethashAccesses    = 64
+
+	// See [4]
+	workGroupSize    = 32 // must be multiple of 8
+	maxSearchResults = 63
+	searchBufSize    = 2
+	globalWorkSize   = 1024 * 256
+)
+
+func NewCL(deviceIds []int) *OpenCLMiner {
+	ids := make([]int, len(deviceIds))
+	copy(ids, deviceIds)
+	return &OpenCLMiner{
+		ethash:    New(),
+		dagSize:   0, // to see if we need to update DAG.
+		deviceIds: ids,
+	}
+}
+
+func PrintDevices() {
+	fmt.Println("=============================================")
+	fmt.Println("============ OpenCL Device Info =============")
+	fmt.Println("=============================================")
+
+	var found []*cl.Device
+
+	platforms, err := cl.GetPlatforms()
+	if err != nil {
+		fmt.Println("Plaform error (check your OpenCL installation):", err)
+		return
+	}
+
+	for i, p := range platforms {
+		fmt.Println("Platform id             ", i)
+		fmt.Println("Platform Name           ", p.Name())
+		fmt.Println("Platform Vendor         ", p.Vendor())
+		fmt.Println("Platform Version        ", p.Version())
+		fmt.Println("Platform Extensions     ", p.Extensions())
+		fmt.Println("Platform Profile        ", p.Profile())
+		fmt.Println("")
+
+		devices, err := cl.GetDevices(p, cl.DeviceTypeGPU)
+		if err != nil {
+			fmt.Println("Device error (check your GPU drivers) :", err)
+			return
+		}
+
+		for _, d := range devices {
+			fmt.Println("Device OpenCL id        ", i)
+			fmt.Println("Device id for mining    ", len(found))
+			fmt.Println("Device Name             ", d.Name())
+			fmt.Println("Vendor                  ", d.Vendor())
+			fmt.Println("Version                 ", d.Version())
+			fmt.Println("Driver version          ", d.DriverVersion())
+			fmt.Println("Address bits            ", d.AddressBits())
+			fmt.Println("Max clock freq          ", d.MaxClockFrequency())
+			fmt.Println("Global mem size         ", d.GlobalMemSize())
+			fmt.Println("Max constant buffer size", d.MaxConstantBufferSize())
+			fmt.Println("Max mem alloc size      ", d.MaxMemAllocSize())
+			fmt.Println("Max compute units       ", d.MaxComputeUnits())
+			fmt.Println("Max work group size     ", d.MaxWorkGroupSize())
+			fmt.Println("Max work item sizes     ", d.MaxWorkItemSizes())
+			fmt.Println("=============================================")
+
+			found = append(found, d)
+		}
+	}
+	if len(found) == 0 {
+		fmt.Println("Found no GPU(s). Check that your OS can see the GPU(s)")
+	} else {
+		var idsFormat string
+		for i := 0; i < len(found); i++ {
+			idsFormat += strconv.Itoa(i)
+			if i != len(found)-1 {
+				idsFormat += ","
+			}
+		}
+		fmt.Printf("Found %v devices. Benchmark first GPU:       geth gpubench 0\n", len(found))
+		fmt.Printf("Mine using all GPUs:                        geth --minegpu %v\n", idsFormat)
+	}
+}
+
+// See [2]. We basically do the same here, but the Go OpenCL bindings
+// are at a slightly higher abtraction level.
+func InitCL(blockNum uint64, c *OpenCLMiner) error {
+	platforms, err := cl.GetPlatforms()
+	if err != nil {
+		return fmt.Errorf("Plaform error: %v\nCheck your OpenCL installation and then run geth gpuinfo", err)
+	}
+
+	var devices []*cl.Device
+	for _, p := range platforms {
+		ds, err := cl.GetDevices(p, cl.DeviceTypeGPU)
+		if err != nil {
+			return fmt.Errorf("Devices error: %v\nCheck your GPU drivers and then run geth gpuinfo", err)
+		}
+		for _, d := range ds {
+			devices = append(devices, d)
+		}
+	}
+
+	pow := New()
+	_ = pow.getDAG(blockNum)     // generates DAG if we don't have it
+	pow.Light.getCache(blockNum) // and cache
+
+	c.ethash = pow
+	dagSize := uint64(C.ethash_get_datasize(C.uint64_t(blockNum)))
+	c.dagSize = dagSize
+
+	for _, id := range c.deviceIds {
+		if id > len(devices)-1 {
+			return fmt.Errorf("Device id not found. See available device ids with: geth gpuinfo")
+		} else {
+			err := initCLDevice(id, devices[id], c)
+			if err != nil {
+				return err
+			}
+		}
+	}
+	if len(c.devices) == 0 {
+		return fmt.Errorf("No GPU devices found")
+	}
+	return nil
+}
+
+func initCLDevice(deviceId int, device *cl.Device, c *OpenCLMiner) error {
+	devMaxAlloc := uint64(device.MaxMemAllocSize())
+	devGlobalMem := uint64(device.GlobalMemSize())
+
+	// TODO: more fine grained version logic
+	if device.Version() == "OpenCL 1.0" {
+		fmt.Println("Device OpenCL version not supported: ", device.Version())
+		return fmt.Errorf("opencl version not supported")
+	}
+
+	var cl11, cl12 bool
+	if device.Version() == "OpenCL 1.1" {
+		cl11 = true
+	}
+	if device.Version() == "OpenCL 1.2" {
+		cl12 = true
+	}
+
+	// log warnings but carry on; some device drivers report inaccurate values
+	if c.dagSize > devGlobalMem {
+		fmt.Printf("WARNING: device memory may be insufficient: %v. DAG size: %v.\n", devGlobalMem, c.dagSize)
+	}
+
+	if c.dagSize > devMaxAlloc {
+		fmt.Printf("WARNING: DAG size (%v) larger than device max memory allocation size (%v).\n", c.dagSize, devMaxAlloc)
+		fmt.Printf("You probably have to export GPU_MAX_ALLOC_PERCENT=95\n")
+	}
+
+	fmt.Printf("Initialising device %v: %v\n", deviceId, device.Name())
+
+	context, err := cl.CreateContext([]*cl.Device{device})
+	if err != nil {
+		return fmt.Errorf("failed creating context: %v", err)
+	}
+
+	// TODO: test running with CL_QUEUE_PROFILING_ENABLE for profiling?
+	queue, err := context.CreateCommandQueue(device, 0)
+	if err != nil {
+		return fmt.Errorf("command queue err: %v", err)
+	}
+
+	// See [4] section 3.2 and [3] "clBuildProgram".
+	// The OpenCL kernel code is compiled at run-time.
+	kvs := make(map[string]string, 4)
+	kvs["GROUP_SIZE"] = strconv.FormatUint(workGroupSize, 10)
+	kvs["DAG_SIZE"] = strconv.FormatUint(c.dagSize/ethashMixBytesLen, 10)
+	kvs["ACCESSES"] = strconv.FormatUint(ethashAccesses, 10)
+	kvs["MAX_OUTPUTS"] = strconv.FormatUint(maxSearchResults, 10)
+	kernelCode := replaceWords(kernel, kvs)
+
+	program, err := context.CreateProgramWithSource([]string{kernelCode})
+	if err != nil {
+		return fmt.Errorf("program err: %v", err)
+	}
+
+	/* if using AMD OpenCL impl, you can set this to debug on x86 CPU device.
+	   see AMD OpenCL programming guide section 4.2
+
+	   export in shell before running:
+	   export AMD_OCL_BUILD_OPTIONS_APPEND="-g -O0"
+	   export CPU_MAX_COMPUTE_UNITS=1
+
+	buildOpts := "-g -cl-opt-disable"
+
+	*/
+	buildOpts := ""
+	err = program.BuildProgram([]*cl.Device{device}, buildOpts)
+	if err != nil {
+		return fmt.Errorf("program build err: %v", err)
+	}
+
+	var searchKernelName, hashKernelName string
+	searchKernelName = "ethash_search"
+	hashKernelName = "ethash_hash"
+
+	searchKernel, err := program.CreateKernel(searchKernelName)
+	hashKernel, err := program.CreateKernel(hashKernelName)
+	if err != nil {
+		return fmt.Errorf("kernel err: %v", err)
+	}
+
+	// TODO: when this DAG size appears, patch the Go bindings
+	// (context.go) to work with uint64 as size_t
+	if c.dagSize > math.MaxInt32 {
+		fmt.Println("DAG too large for allocation.")
+		return fmt.Errorf("DAG too large for alloc")
+	}
+
+	// TODO: patch up Go bindings to work with size_t, will overflow if > maxint32
+	// TODO: fuck. shit's gonna overflow around 2017-06-09 12:17:02
+	dagBuf := *(new(*cl.MemObject))
+	dagBuf, err = context.CreateEmptyBuffer(cl.MemReadOnly, int(c.dagSize))
+	if err != nil {
+		return fmt.Errorf("allocating dag buf failed: %v", err)
+	}
+
+	// write DAG to device mem
+	dagPtr := unsafe.Pointer(c.ethash.Full.current.ptr.data)
+	_, err = queue.EnqueueWriteBuffer(dagBuf, true, 0, int(c.dagSize), dagPtr, nil)
+	if err != nil {
+		return fmt.Errorf("writing to dag buf failed: %v", err)
+	}
+
+	searchBuffers := make([]*cl.MemObject, searchBufSize)
+	for i := 0; i < searchBufSize; i++ {
+		searchBuff, err := context.CreateEmptyBuffer(cl.MemWriteOnly, (1+maxSearchResults)*SIZEOF_UINT32)
+		if err != nil {
+			return fmt.Errorf("search buffer err: %v", err)
+		}
+		searchBuffers[i] = searchBuff
+	}
+
+	headerBuf, err := context.CreateEmptyBuffer(cl.MemReadOnly, 32)
+	if err != nil {
+		return fmt.Errorf("header buffer err: %v", err)
+	}
+
+	// Unique, random nonces are crucial for mining efficieny.
+	// While we do not need cryptographically secure PRNG for nonces,
+	// we want to have uniform distribution and minimal repetition of nonces.
+	// We could guarantee strict uniqueness of nonces by generating unique ranges,
+	// but a int64 seed from crypto/rand should be good enough.
+	// we then use math/rand for speed and to avoid draining OS entropy pool
+	seed, err := crand.Int(crand.Reader, big.NewInt(math.MaxInt64))
+	if err != nil {
+		return err
+	}
+	nonceRand := mrand.New(mrand.NewSource(seed.Int64()))
+
+	deviceStruct := &OpenCLDevice{
+		deviceId: deviceId,
+		device:   device,
+		openCL11: cl11,
+		openCL12: cl12,
+
+		dagBuf:        dagBuf,
+		headerBuf:     headerBuf,
+		searchBuffers: searchBuffers,
+
+		searchKernel: searchKernel,
+		hashKernel:   hashKernel,
+
+		queue: queue,
+		ctx:   context,
+
+		workGroupSize: workGroupSize,
+
+		nonceRand: nonceRand,
+	}
+	c.devices = append(c.devices, deviceStruct)
+
+	return nil
+}
+
+func (c *OpenCLMiner) Search(block pow.Block, stop <-chan struct{}, index int) (uint64, []byte) {
+	c.mu.Lock()
+	newDagSize := uint64(C.ethash_get_datasize(C.uint64_t(block.NumberU64())))
+	if newDagSize > c.dagSize {
+		// TODO: clean up buffers from previous DAG?
+		err := InitCL(block.NumberU64(), c)
+		if err != nil {
+			fmt.Println("OpenCL init error: ", err)
+			return 0, []byte{0}
+		}
+	}
+	defer c.mu.Unlock()
+
+	// Avoid unneeded OpenCL initialisation if we received stop while running InitCL
+	select {
+	case <-stop:
+		return 0, []byte{0}
+	default:
+	}
+
+	headerHash := block.HashNoNonce()
+	diff := block.Difficulty()
+	target256 := new(big.Int).Div(maxUint256, diff)
+	target64 := new(big.Int).Rsh(target256, 192).Uint64()
+	var zero uint32 = 0
+
+	d := c.devices[index]
+
+	_, err := d.queue.EnqueueWriteBuffer(d.headerBuf, false, 0, 32, unsafe.Pointer(&headerHash[0]), nil)
+	if err != nil {
+		fmt.Println("Error in Search clEnqueueWriterBuffer : ", err)
+		return 0, []byte{0}
+	}
+
+	for i := 0; i < searchBufSize; i++ {
+		_, err := d.queue.EnqueueWriteBuffer(d.searchBuffers[i], false, 0, 4, unsafe.Pointer(&zero), nil)
+		if err != nil {
+			fmt.Println("Error in Search clEnqueueWriterBuffer : ", err)
+			return 0, []byte{0}
+		}
+	}
+
+	// wait for all search buffers to complete
+	err = d.queue.Finish()
+	if err != nil {
+		fmt.Println("Error in Search clFinish : ", err)
+		return 0, []byte{0}
+	}
+
+	err = d.searchKernel.SetArg(1, d.headerBuf)
+	if err != nil {
+		fmt.Println("Error in Search clSetKernelArg : ", err)
+		return 0, []byte{0}
+	}
+
+	err = d.searchKernel.SetArg(2, d.dagBuf)
+	if err != nil {
+		fmt.Println("Error in Search clSetKernelArg : ", err)
+		return 0, []byte{0}
+	}
+
+	err = d.searchKernel.SetArg(4, target64)
+	if err != nil {
+		fmt.Println("Error in Search clSetKernelArg : ", err)
+		return 0, []byte{0}
+	}
+	err = d.searchKernel.SetArg(5, uint32(math.MaxUint32))
+	if err != nil {
+		fmt.Println("Error in Search clSetKernelArg : ", err)
+		return 0, []byte{0}
+	}
+
+	// wait on this before returning
+	var preReturnEvent *cl.Event
+	if d.openCL12 {
+		preReturnEvent, err = d.ctx.CreateUserEvent()
+		if err != nil {
+			fmt.Println("Error in Search create CL user event : ", err)
+			return 0, []byte{0}
+		}
+	}
+
+	pending := make([]pendingSearch, 0, searchBufSize)
+	var p *pendingSearch
+	searchBufIndex := uint32(0)
+	var checkNonce uint64
+	loops := int64(0)
+	prevHashRate := int32(0)
+	start := time.Now().UnixNano()
+	// we grab a single random nonce and sets this as argument to the kernel search function
+	// the device will then add each local threads gid to the nonce, creating a unique nonce
+	// for each device computing unit executing in parallel
+	initNonce := uint64(d.nonceRand.Int63())
+	for nonce := initNonce; ; nonce += uint64(globalWorkSize) {
+		select {
+		case <-stop:
+
+			/*
+				if d.openCL12 {
+					err = cl.WaitForEvents([]*cl.Event{preReturnEvent})
+					if err != nil {
+						fmt.Println("Error in Search WaitForEvents: ", err)
+					}
+				}
+			*/
+
+			atomic.AddInt32(&c.hashRate, -prevHashRate)
+			return 0, []byte{0}
+		default:
+		}
+
+		if (loops % (1 << 7)) == 0 {
+			elapsed := time.Now().UnixNano() - start
+			// TODO: verify if this is correct hash rate calculation
+			hashes := (float64(1e9) / float64(elapsed)) * float64(loops*1024*256)
+			hashrateDiff := int32(hashes) - prevHashRate
+			prevHashRate = int32(hashes)
+			atomic.AddInt32(&c.hashRate, hashrateDiff)
+		}
+		loops++
+
+		err = d.searchKernel.SetArg(0, d.searchBuffers[searchBufIndex])
+		if err != nil {
+			fmt.Println("Error in Search clSetKernelArg : ", err)
+			return 0, []byte{0}
+		}
+		err = d.searchKernel.SetArg(3, nonce)
+		if err != nil {
+			fmt.Println("Error in Search clSetKernelArg : ", err)
+			return 0, []byte{0}
+		}
+
+		// execute kernel
+		_, err := d.queue.EnqueueNDRangeKernel(
+			d.searchKernel,
+			[]int{0},
+			[]int{globalWorkSize},
+			[]int{d.workGroupSize},
+			nil)
+		if err != nil {
+			fmt.Println("Error in Search clEnqueueNDRangeKernel : ", err)
+			return 0, []byte{0}
+		}
+
+		pending = append(pending, pendingSearch{bufIndex: searchBufIndex, startNonce: nonce})
+		searchBufIndex = (searchBufIndex + 1) % searchBufSize
+
+		if len(pending) == searchBufSize {
+			p = &(pending[searchBufIndex])
+			cres, _, err := d.queue.EnqueueMapBuffer(d.searchBuffers[p.bufIndex], true,
+				cl.MapFlagRead, 0, (1+maxSearchResults)*SIZEOF_UINT32,
+				nil)
+			if err != nil {
+				fmt.Println("Error in Search clEnqueueMapBuffer: ", err)
+				return 0, []byte{0}
+			}
+
+			results := cres.ByteSlice()
+			nfound := binary.LittleEndian.Uint32(results)
+			nfound = uint32(math.Min(float64(nfound), float64(maxSearchResults)))
+			// OpenCL returns the offsets from the start nonce
+			for i := uint32(0); i < nfound; i++ {
+				lo := (i + 1) * SIZEOF_UINT32
+				hi := (i + 2) * SIZEOF_UINT32
+				upperNonce := uint64(binary.LittleEndian.Uint32(results[lo:hi]))
+				checkNonce = p.startNonce + upperNonce
+				if checkNonce != 0 {
+					// We verify that the nonce is indeed a solution by
+					// executing the Ethash verification function (on the CPU).
+					cache := c.ethash.Light.getCache(block.NumberU64())
+					ok, mixDigest, result := cache.compute(c.dagSize, headerHash, checkNonce)
+
+					// TODO: return result first
+					if ok && result.Big().Cmp(target256) <= 0 {
+						_, err = d.queue.EnqueueUnmapMemObject(d.searchBuffers[p.bufIndex], cres, nil)
+						if err != nil {
+							fmt.Println("Error in Search clEnqueueUnmapMemObject: ", err)
+						}
+						if d.openCL12 {
+							err = cl.WaitForEvents([]*cl.Event{preReturnEvent})
+							if err != nil {
+								fmt.Println("Error in Search WaitForEvents: ", err)
+							}
+						}
+						return checkNonce, mixDigest.Bytes()
+					}
+					_, err := d.queue.EnqueueWriteBuffer(d.searchBuffers[p.bufIndex], false, 0, 4, unsafe.Pointer(&zero), nil)
+					if err != nil {
+						fmt.Println("Error in Search cl: EnqueueWriteBuffer", err)
+						return 0, []byte{0}
+					}
+				}
+			}
+			_, err = d.queue.EnqueueUnmapMemObject(d.searchBuffers[p.bufIndex], cres, nil)
+			if err != nil {
+				fmt.Println("Error in Search clEnqueueUnMapMemObject: ", err)
+				return 0, []byte{0}
+			}
+			pending = append(pending[:searchBufIndex], pending[searchBufIndex+1:]...)
+		}
+	}
+	if d.openCL12 {
+		err := cl.WaitForEvents([]*cl.Event{preReturnEvent})
+		if err != nil {
+			fmt.Println("Error in Search clWaitForEvents: ", err)
+			return 0, []byte{0}
+		}
+	}
+	return 0, []byte{0}
+}
+
+func (c *OpenCLMiner) Verify(block pow.Block) bool {
+	return c.ethash.Light.Verify(block)
+}
+func (c *OpenCLMiner) GetHashrate() int64 {
+	return int64(atomic.LoadInt32(&c.hashRate))
+}
+func (c *OpenCLMiner) Turbo(on bool) {
+	// This is GPU mining. Always be turbo.
+}
+
+func replaceWords(text string, kvs map[string]string) string {
+	for k, v := range kvs {
+		text = strings.Replace(text, k, v, -1)
+	}
+	return text
+}
+
+func logErr(err error) {
+	if err != nil {
+		fmt.Println("Error in OpenCL call:", err)
+	}
+}
+
+func argErr(err error) error {
+	return fmt.Errorf("arg err: %v", err)
+}
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/ethash_opencl_kernel_go_str.go
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/ethash_opencl_kernel_go_str.go
@@ -0,0 +1,600 @@
+package ethash
+
+/*  DO NOT EDIT!!!
+
+    This code is version controlled at
+    https://github.com/ethereum/cpp-ethereum/blob/develop/libethash-cl/ethash_cl_miner_kernel.cl
+
+    If needed change it there first, then copy over here.
+*/
+
+const kernel = `
+// author Tim Hughes <tim@twistedfury.com>
+// Tested on Radeon HD 7850
+// Hashrate: 15940347 hashes/s
+// Bandwidth: 124533 MB/s
+// search kernel should fit in <= 84 VGPRS (3 wavefronts)
+
+#define THREADS_PER_HASH (128 / 16)
+#define HASHES_PER_LOOP (GROUP_SIZE / THREADS_PER_HASH)
+
+#define FNV_PRIME	0x01000193
+
+__constant uint2 const Keccak_f1600_RC[24] = {
+	(uint2)(0x00000001, 0x00000000),
+	(uint2)(0x00008082, 0x00000000),
+	(uint2)(0x0000808a, 0x80000000),
+	(uint2)(0x80008000, 0x80000000),
+	(uint2)(0x0000808b, 0x00000000),
+	(uint2)(0x80000001, 0x00000000),
+	(uint2)(0x80008081, 0x80000000),
+	(uint2)(0x00008009, 0x80000000),
+	(uint2)(0x0000008a, 0x00000000),
+	(uint2)(0x00000088, 0x00000000),
+	(uint2)(0x80008009, 0x00000000),
+	(uint2)(0x8000000a, 0x00000000),
+	(uint2)(0x8000808b, 0x00000000),
+	(uint2)(0x0000008b, 0x80000000),
+	(uint2)(0x00008089, 0x80000000),
+	(uint2)(0x00008003, 0x80000000),
+	(uint2)(0x00008002, 0x80000000),
+	(uint2)(0x00000080, 0x80000000),
+	(uint2)(0x0000800a, 0x00000000),
+	(uint2)(0x8000000a, 0x80000000),
+	(uint2)(0x80008081, 0x80000000),
+	(uint2)(0x00008080, 0x80000000),
+	(uint2)(0x80000001, 0x00000000),
+	(uint2)(0x80008008, 0x80000000),
+};
+
+void keccak_f1600_round(uint2* a, uint r, uint out_size)
+{
+   #if !__ENDIAN_LITTLE__
+	for (uint i = 0; i != 25; ++i)
+		a[i] = a[i].yx;
+   #endif
+
+	uint2 b[25];
+	uint2 t;
+
+	// Theta
+	b[0] = a[0] ^ a[5] ^ a[10] ^ a[15] ^ a[20];
+	b[1] = a[1] ^ a[6] ^ a[11] ^ a[16] ^ a[21];
+	b[2] = a[2] ^ a[7] ^ a[12] ^ a[17] ^ a[22];
+	b[3] = a[3] ^ a[8] ^ a[13] ^ a[18] ^ a[23];
+	b[4] = a[4] ^ a[9] ^ a[14] ^ a[19] ^ a[24];
+	t = b[4] ^ (uint2)(b[1].x << 1 | b[1].y >> 31, b[1].y << 1 | b[1].x >> 31);
+	a[0] ^= t;
+	a[5] ^= t;
+	a[10] ^= t;
+	a[15] ^= t;
+	a[20] ^= t;
+	t = b[0] ^ (uint2)(b[2].x << 1 | b[2].y >> 31, b[2].y << 1 | b[2].x >> 31);
+	a[1] ^= t;
+	a[6] ^= t;
+	a[11] ^= t;
+	a[16] ^= t;
+	a[21] ^= t;
+	t = b[1] ^ (uint2)(b[3].x << 1 | b[3].y >> 31, b[3].y << 1 | b[3].x >> 31);
+	a[2] ^= t;
+	a[7] ^= t;
+	a[12] ^= t;
+	a[17] ^= t;
+	a[22] ^= t;
+	t = b[2] ^ (uint2)(b[4].x << 1 | b[4].y >> 31, b[4].y << 1 | b[4].x >> 31);
+	a[3] ^= t;
+	a[8] ^= t;
+	a[13] ^= t;
+	a[18] ^= t;
+	a[23] ^= t;
+	t = b[3] ^ (uint2)(b[0].x << 1 | b[0].y >> 31, b[0].y << 1 | b[0].x >> 31);
+	a[4] ^= t;
+	a[9] ^= t;
+	a[14] ^= t;
+	a[19] ^= t;
+	a[24] ^= t;
+
+	// Rho Pi
+	b[0] = a[0];
+	b[10] = (uint2)(a[1].x << 1 | a[1].y >> 31, a[1].y << 1 | a[1].x >> 31);
+	b[7] = (uint2)(a[10].x << 3 | a[10].y >> 29, a[10].y << 3 | a[10].x >> 29);
+	b[11] = (uint2)(a[7].x << 6 | a[7].y >> 26, a[7].y << 6 | a[7].x >> 26);
+	b[17] = (uint2)(a[11].x << 10 | a[11].y >> 22, a[11].y << 10 | a[11].x >> 22);
+	b[18] = (uint2)(a[17].x << 15 | a[17].y >> 17, a[17].y << 15 | a[17].x >> 17);
+	b[3] = (uint2)(a[18].x << 21 | a[18].y >> 11, a[18].y << 21 | a[18].x >> 11);
+	b[5] = (uint2)(a[3].x << 28 | a[3].y >> 4, a[3].y << 28 | a[3].x >> 4);
+	b[16] = (uint2)(a[5].y << 4 | a[5].x >> 28, a[5].x << 4 | a[5].y >> 28);
+	b[8] = (uint2)(a[16].y << 13 | a[16].x >> 19, a[16].x << 13 | a[16].y >> 19);
+	b[21] = (uint2)(a[8].y << 23 | a[8].x >> 9, a[8].x << 23 | a[8].y >> 9);
+	b[24] = (uint2)(a[21].x << 2 | a[21].y >> 30, a[21].y << 2 | a[21].x >> 30);
+	b[4] = (uint2)(a[24].x << 14 | a[24].y >> 18, a[24].y << 14 | a[24].x >> 18);
+	b[15] = (uint2)(a[4].x << 27 | a[4].y >> 5, a[4].y << 27 | a[4].x >> 5);
+	b[23] = (uint2)(a[15].y << 9 | a[15].x >> 23, a[15].x << 9 | a[15].y >> 23);
+	b[19] = (uint2)(a[23].y << 24 | a[23].x >> 8, a[23].x << 24 | a[23].y >> 8);
+	b[13] = (uint2)(a[19].x << 8 | a[19].y >> 24, a[19].y << 8 | a[19].x >> 24);
+	b[12] = (uint2)(a[13].x << 25 | a[13].y >> 7, a[13].y << 25 | a[13].x >> 7);
+	b[2] = (uint2)(a[12].y << 11 | a[12].x >> 21, a[12].x << 11 | a[12].y >> 21);
+	b[20] = (uint2)(a[2].y << 30 | a[2].x >> 2, a[2].x << 30 | a[2].y >> 2);
+	b[14] = (uint2)(a[20].x << 18 | a[20].y >> 14, a[20].y << 18 | a[20].x >> 14);
+	b[22] = (uint2)(a[14].y << 7 | a[14].x >> 25, a[14].x << 7 | a[14].y >> 25);
+	b[9] = (uint2)(a[22].y << 29 | a[22].x >> 3, a[22].x << 29 | a[22].y >> 3);
+	b[6] = (uint2)(a[9].x << 20 | a[9].y >> 12, a[9].y << 20 | a[9].x >> 12);
+	b[1] = (uint2)(a[6].y << 12 | a[6].x >> 20, a[6].x << 12 | a[6].y >> 20);
+
+	// Chi
+	a[0] = bitselect(b[0] ^ b[2], b[0], b[1]);
+	a[1] = bitselect(b[1] ^ b[3], b[1], b[2]);
+	a[2] = bitselect(b[2] ^ b[4], b[2], b[3]);
+	a[3] = bitselect(b[3] ^ b[0], b[3], b[4]);
+	if (out_size >= 4)
+	{
+		a[4] = bitselect(b[4] ^ b[1], b[4], b[0]);
+		a[5] = bitselect(b[5] ^ b[7], b[5], b[6]);
+		a[6] = bitselect(b[6] ^ b[8], b[6], b[7]);
+		a[7] = bitselect(b[7] ^ b[9], b[7], b[8]);
+		a[8] = bitselect(b[8] ^ b[5], b[8], b[9]);
+		if (out_size >= 8)
+		{
+			a[9] = bitselect(b[9] ^ b[6], b[9], b[5]);
+			a[10] = bitselect(b[10] ^ b[12], b[10], b[11]);
+			a[11] = bitselect(b[11] ^ b[13], b[11], b[12]);
+			a[12] = bitselect(b[12] ^ b[14], b[12], b[13]);
+			a[13] = bitselect(b[13] ^ b[10], b[13], b[14]);
+			a[14] = bitselect(b[14] ^ b[11], b[14], b[10]);
+			a[15] = bitselect(b[15] ^ b[17], b[15], b[16]);
+			a[16] = bitselect(b[16] ^ b[18], b[16], b[17]);
+			a[17] = bitselect(b[17] ^ b[19], b[17], b[18]);
+			a[18] = bitselect(b[18] ^ b[15], b[18], b[19]);
+			a[19] = bitselect(b[19] ^ b[16], b[19], b[15]);
+			a[20] = bitselect(b[20] ^ b[22], b[20], b[21]);
+			a[21] = bitselect(b[21] ^ b[23], b[21], b[22]);
+			a[22] = bitselect(b[22] ^ b[24], b[22], b[23]);
+			a[23] = bitselect(b[23] ^ b[20], b[23], b[24]);
+			a[24] = bitselect(b[24] ^ b[21], b[24], b[20]);
+		}
+	}
+
+	// Iota
+	a[0] ^= Keccak_f1600_RC[r];
+
+   #if !__ENDIAN_LITTLE__
+	for (uint i = 0; i != 25; ++i)
+		a[i] = a[i].yx;
+   #endif
+}
+
+void keccak_f1600_no_absorb(ulong* a, uint in_size, uint out_size, uint isolate)
+{
+	for (uint i = in_size; i != 25; ++i)
+	{
+		a[i] = 0;
+	}
+#if __ENDIAN_LITTLE__
+	a[in_size] ^= 0x0000000000000001;
+	a[24-out_size*2] ^= 0x8000000000000000;
+#else
+	a[in_size] ^= 0x0100000000000000;
+	a[24-out_size*2] ^= 0x0000000000000080;
+#endif
+
+	// Originally I unrolled the first and last rounds to interface
+	// better with surrounding code, however I haven't done this
+	// without causing the AMD compiler to blow up the VGPR usage.
+	uint r = 0;
+	do
+	{
+		// This dynamic branch stops the AMD compiler unrolling the loop
+		// and additionally saves about 33% of the VGPRs, enough to gain another
+		// wavefront. Ideally we'd get 4 in flight, but 3 is the best I can
+		// massage out of the compiler. It doesn't really seem to matter how
+		// much we try and help the compiler save VGPRs because it seems to throw
+		// that information away, hence the implementation of keccak here
+		// doesn't bother.
+		if (isolate)
+		{
+			keccak_f1600_round((uint2*)a, r++, 25);
+		}
+	}
+	while (r < 23);
+
+	// final round optimised for digest size
+	keccak_f1600_round((uint2*)a, r++, out_size);
+}
+
+#define copy(dst, src, count) for (uint i = 0; i != count; ++i) { (dst)[i] = (src)[i]; }
+
+#define countof(x) (sizeof(x) / sizeof(x[0]))
+
+uint fnv(uint x, uint y)
+{
+	return x * FNV_PRIME ^ y;
+}
+
+uint4 fnv4(uint4 x, uint4 y)
+{
+	return x * FNV_PRIME ^ y;
+}
+
+uint fnv_reduce(uint4 v)
+{
+	return fnv(fnv(fnv(v.x, v.y), v.z), v.w);
+}
+
+typedef union
+{
+	ulong ulongs[32 / sizeof(ulong)];
+	uint uints[32 / sizeof(uint)];
+} hash32_t;
+
+typedef union
+{
+	ulong ulongs[64 / sizeof(ulong)];
+	uint4 uint4s[64 / sizeof(uint4)];
+} hash64_t;
+
+typedef union
+{
+	uint uints[128 / sizeof(uint)];
+	uint4 uint4s[128 / sizeof(uint4)];
+} hash128_t;
+
+hash64_t init_hash(__constant hash32_t const* header, ulong nonce, uint isolate)
+{
+	hash64_t init;
+	uint const init_size = countof(init.ulongs);
+	uint const hash_size = countof(header->ulongs);
+
+	// sha3_512(header .. nonce)
+	ulong state[25];
+	copy(state, header->ulongs, hash_size);
+	state[hash_size] = nonce;
+	keccak_f1600_no_absorb(state, hash_size + 1, init_size, isolate);
+
+	copy(init.ulongs, state, init_size);
+	return init;
+}
+
+uint inner_loop_chunks(uint4 init, uint thread_id, __local uint* share, __global hash128_t const* g_dag, __global hash128_t const* g_dag1, __global hash128_t const* g_dag2, __global hash128_t const* g_dag3, uint isolate)
+{
+	uint4 mix = init;
+
+	// share init0
+	if (thread_id == 0)
+		*share = mix.x;
+	barrier(CLK_LOCAL_MEM_FENCE);
+	uint init0 = *share;
+
+	uint a = 0;
+	do
+	{
+		bool update_share = thread_id == (a/4) % THREADS_PER_HASH;
+
+		#pragma unroll
+		for (uint i = 0; i != 4; ++i)
+		{
+			if (update_share)
+			{
+				uint m[4] = { mix.x, mix.y, mix.z, mix.w };
+				*share = fnv(init0 ^ (a+i), m[i]) % DAG_SIZE;
+			}
+			barrier(CLK_LOCAL_MEM_FENCE);
+
+			mix = fnv4(mix, *share>=3 * DAG_SIZE / 4 ? g_dag3[*share - 3 * DAG_SIZE / 4].uint4s[thread_id] : *share>=DAG_SIZE / 2 ? g_dag2[*share - DAG_SIZE / 2].uint4s[thread_id] : *share>=DAG_SIZE / 4 ? g_dag1[*share - DAG_SIZE / 4].uint4s[thread_id]:g_dag[*share].uint4s[thread_id]);
+		}
+	} while ((a += 4) != (ACCESSES & isolate));
+
+	return fnv_reduce(mix);
+}
+
+
+
+uint inner_loop(uint4 init, uint thread_id, __local uint* share, __global hash128_t const* g_dag, uint isolate)
+{
+	uint4 mix = init;
+
+	// share init0
+	if (thread_id == 0)
+		*share = mix.x;
+	barrier(CLK_LOCAL_MEM_FENCE);
+	uint init0 = *share;
+
+	uint a = 0;
+	do
+	{
+		bool update_share = thread_id == (a/4) % THREADS_PER_HASH;
+
+		#pragma unroll
+		for (uint i = 0; i != 4; ++i)
+		{
+			if (update_share)
+			{
+				uint m[4] = { mix.x, mix.y, mix.z, mix.w };
+				*share = fnv(init0 ^ (a+i), m[i]) % DAG_SIZE;
+			}
+			barrier(CLK_LOCAL_MEM_FENCE);
+
+			mix = fnv4(mix, g_dag[*share].uint4s[thread_id]);
+		}
+	}
+	while ((a += 4) != (ACCESSES & isolate));
+
+	return fnv_reduce(mix);
+}
+
+
+hash32_t final_hash(hash64_t const* init, hash32_t const* mix, uint isolate)
+{
+	ulong state[25];
+
+	hash32_t hash;
+	uint const hash_size = countof(hash.ulongs);
+	uint const init_size = countof(init->ulongs);
+	uint const mix_size = countof(mix->ulongs);
+
+	// keccak_256(keccak_512(header..nonce) .. mix);
+	copy(state, init->ulongs, init_size);
+	copy(state + init_size, mix->ulongs, mix_size);
+	keccak_f1600_no_absorb(state, init_size+mix_size, hash_size, isolate);
+
+	// copy out
+	copy(hash.ulongs, state, hash_size);
+	return hash;
+}
+
+hash32_t compute_hash_simple(
+	__constant hash32_t const* g_header,
+	__global hash128_t const* g_dag,
+	ulong nonce,
+	uint isolate
+	)
+{
+	hash64_t init = init_hash(g_header, nonce, isolate);
+
+	hash128_t mix;
+	for (uint i = 0; i != countof(mix.uint4s); ++i)
+	{
+		mix.uint4s[i] = init.uint4s[i % countof(init.uint4s)];
+	}
+
+	uint mix_val = mix.uints[0];
+	uint init0 = mix.uints[0];
+	uint a = 0;
+	do
+	{
+		uint pi = fnv(init0 ^ a, mix_val) % DAG_SIZE;
+		uint n = (a+1) % countof(mix.uints);
+
+		#pragma unroll
+		for (uint i = 0; i != countof(mix.uints); ++i)
+		{
+			mix.uints[i] = fnv(mix.uints[i], g_dag[pi].uints[i]);
+			mix_val = i == n ? mix.uints[i] : mix_val;
+		}
+	}
+	while (++a != (ACCESSES & isolate));
+
+	// reduce to output
+	hash32_t fnv_mix;
+	for (uint i = 0; i != countof(fnv_mix.uints); ++i)
+	{
+		fnv_mix.uints[i] = fnv_reduce(mix.uint4s[i]);
+	}
+
+	return final_hash(&init, &fnv_mix, isolate);
+}
+
+typedef union
+{
+	struct
+	{
+		hash64_t init;
+		uint pad; // avoid lds bank conflicts
+	};
+	hash32_t mix;
+} compute_hash_share;
+
+
+hash32_t compute_hash(
+	__local compute_hash_share* share,
+	__constant hash32_t const* g_header,
+	__global hash128_t const* g_dag,
+	ulong nonce,
+	uint isolate
+	)
+{
+	uint const gid = get_global_id(0);
+
+	// Compute one init hash per work item.
+	hash64_t init = init_hash(g_header, nonce, isolate);
+
+	// Threads work together in this phase in groups of 8.
+	uint const thread_id = gid % THREADS_PER_HASH;
+	uint const hash_id = (gid % GROUP_SIZE) / THREADS_PER_HASH;
+
+	hash32_t mix;
+	uint i = 0;
+	do
+	{
+		// share init with other threads
+		if (i == thread_id)
+			share[hash_id].init = init;
+		barrier(CLK_LOCAL_MEM_FENCE);
+
+		uint4 thread_init = share[hash_id].init.uint4s[thread_id % (64 / sizeof(uint4))];
+		barrier(CLK_LOCAL_MEM_FENCE);
+
+		uint thread_mix = inner_loop(thread_init, thread_id, share[hash_id].mix.uints, g_dag, isolate);
+
+		share[hash_id].mix.uints[thread_id] = thread_mix;
+		barrier(CLK_LOCAL_MEM_FENCE);
+
+		if (i == thread_id)
+			mix = share[hash_id].mix;
+		barrier(CLK_LOCAL_MEM_FENCE);
+	}
+	while (++i != (THREADS_PER_HASH & isolate));
+
+	return final_hash(&init, &mix, isolate);
+}
+
+
+hash32_t compute_hash_chunks(
+	__local compute_hash_share* share,
+	__constant hash32_t const* g_header,
+	__global hash128_t const* g_dag,
+	__global hash128_t const* g_dag1,
+	__global hash128_t const* g_dag2,
+	__global hash128_t const* g_dag3,
+	ulong nonce,
+	uint isolate
+	)
+{
+	uint const gid = get_global_id(0);
+
+	// Compute one init hash per work item.
+	hash64_t init = init_hash(g_header, nonce, isolate);
+
+	// Threads work together in this phase in groups of 8.
+	uint const thread_id = gid % THREADS_PER_HASH;
+	uint const hash_id = (gid % GROUP_SIZE) / THREADS_PER_HASH;
+
+	hash32_t mix;
+	uint i = 0;
+	do
+	{
+		// share init with other threads
+		if (i == thread_id)
+			share[hash_id].init = init;
+		barrier(CLK_LOCAL_MEM_FENCE);
+
+		uint4 thread_init = share[hash_id].init.uint4s[thread_id % (64 / sizeof(uint4))];
+		barrier(CLK_LOCAL_MEM_FENCE);
+
+		uint thread_mix = inner_loop_chunks(thread_init, thread_id, share[hash_id].mix.uints, g_dag, g_dag1, g_dag2, g_dag3, isolate);
+
+		share[hash_id].mix.uints[thread_id] = thread_mix;
+		barrier(CLK_LOCAL_MEM_FENCE);
+
+		if (i == thread_id)
+			mix = share[hash_id].mix;
+		barrier(CLK_LOCAL_MEM_FENCE);
+	}
+	while (++i != (THREADS_PER_HASH & isolate));
+
+	return final_hash(&init, &mix, isolate);
+}
+
+__attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1)))
+__kernel void ethash_hash_simple(
+	__global hash32_t* g_hashes,
+	__constant hash32_t const* g_header,
+	__global hash128_t const* g_dag,
+	ulong start_nonce,
+	uint isolate
+	)
+{
+	uint const gid = get_global_id(0);
+	g_hashes[gid] = compute_hash_simple(g_header, g_dag, start_nonce + gid, isolate);
+}
+
+__attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1)))
+__kernel void ethash_search_simple(
+	__global volatile uint* restrict g_output,
+	__constant hash32_t const* g_header,
+	__global hash128_t const* g_dag,
+	ulong start_nonce,
+	ulong target,
+	uint isolate
+	)
+{
+	uint const gid = get_global_id(0);
+	hash32_t hash = compute_hash_simple(g_header, g_dag, start_nonce + gid, isolate);
+
+	if (hash.ulongs[countof(hash.ulongs)-1] < target)
+	{
+		uint slot = min(convert_uint(MAX_OUTPUTS), convert_uint(atomic_inc(&g_output[0]) + 1));
+		g_output[slot] = gid;
+	}
+}
+
+
+__attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1)))
+__kernel void ethash_hash(
+	__global hash32_t* g_hashes,
+	__constant hash32_t const* g_header,
+	__global hash128_t const* g_dag,
+	ulong start_nonce,
+	uint isolate
+	)
+{
+	__local compute_hash_share share[HASHES_PER_LOOP];
+
+	uint const gid = get_global_id(0);
+	g_hashes[gid] = compute_hash(share, g_header, g_dag, start_nonce + gid, isolate);
+}
+
+__attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1)))
+__kernel void ethash_search(
+	__global volatile uint* restrict g_output,
+	__constant hash32_t const* g_header,
+	__global hash128_t const* g_dag,
+	ulong start_nonce,
+	ulong target,
+	uint isolate
+	)
+{
+	__local compute_hash_share share[HASHES_PER_LOOP];
+
+	uint const gid = get_global_id(0);
+	hash32_t hash = compute_hash(share, g_header, g_dag, start_nonce + gid, isolate);
+
+	if (as_ulong(as_uchar8(hash.ulongs[0]).s76543210) < target)
+	{
+		uint slot = min((uint)MAX_OUTPUTS, atomic_inc(&g_output[0]) + 1);
+		g_output[slot] = gid;
+	}
+}
+
+__attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1)))
+__kernel void ethash_hash_chunks(
+	__global hash32_t* g_hashes,
+	__constant hash32_t const* g_header,
+	__global hash128_t const* g_dag,
+	__global hash128_t const* g_dag1,
+	__global hash128_t const* g_dag2,
+	__global hash128_t const* g_dag3,
+	ulong start_nonce,
+	uint isolate
+	)
+{
+	__local compute_hash_share share[HASHES_PER_LOOP];
+
+	uint const gid = get_global_id(0);
+	g_hashes[gid] = compute_hash_chunks(share, g_header, g_dag, g_dag1, g_dag2, g_dag3,start_nonce + gid, isolate);
+}
+
+__attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1)))
+__kernel void ethash_search_chunks(
+	__global volatile uint* restrict g_output,
+	__constant hash32_t const* g_header,
+	__global hash128_t const* g_dag,
+	__global hash128_t const* g_dag1,
+	__global hash128_t const* g_dag2,
+	__global hash128_t const* g_dag3,
+	ulong start_nonce,
+	ulong target,
+	uint isolate
+	)
+{
+	__local compute_hash_share share[HASHES_PER_LOOP];
+
+	uint const gid = get_global_id(0);
+	hash32_t hash = compute_hash_chunks(share, g_header, g_dag, g_dag1, g_dag2, g_dag3, start_nonce + gid, isolate);
+
+	if (as_ulong(as_uchar8(hash.ulongs[0]).s76543210) < target)
+	{
+		uint slot = min(convert_uint(MAX_OUTPUTS), convert_uint(atomic_inc(&g_output[0]) + 1));
+		g_output[slot] = gid;
+	}
+}
+`
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/ethashc.go
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/ethashc.go
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/setup.py
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/setup.py
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/CMakeLists.txt
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/CMakeLists.txt
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/compiler.h
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/compiler.h
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/data_sizes.h
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/data_sizes.h
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/endian.h
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/endian.h
@@ -19,7 +19,7 @@
  # define BYTE_ORDER    LITTLE_ENDIAN
 #elif defined( __QNXNTO__ ) && defined( __BIGENDIAN__ )
  # define BIG_ENDIAN 1234
-  # define BYTE_ORDER BIG_ENDIAN
+  # define BYTE_ORDER    BIG_ENDIAN
 #else
 # include <endian.h>
 #endif
@@ -59,20 +59,21 @@

 #define fix_endian32(dst_, src_) dst_ = ethash_swap_u32(src_)
 #define fix_endian32_same(val_) val_ = ethash_swap_u32(val_)
-#define fix_endian64(dst_, src_) dst_ = ethash_swap_u64(src_)
+#define fix_endian64(dst_, src_) dst_ = ethash_swap_u64(src_
 #define fix_endian64_same(val_) val_ = ethash_swap_u64(val_)
-#define fix_endian_arr32(arr_, size_) \
-  do { \
-    for (unsigned i_ = 0; i_ < (size_); ++i_) { \
-      arr_[i_] = ethash_swap_u32(arr_[i_]); \
-    } \
-  } while (0)
-#define fix_endian_arr64(arr_, size_) \
-  do { \
-    for (unsigned i_ = 0; i_ < (size_); ++i_) { \
-      arr_[i_] = ethash_swap_u64(arr_[i_]); \
-    } \
-  } while (0)
+#define fix_endian_arr32(arr_, size_)			\
+	do {										\
+	for (unsigned i_ = 0; i_ < (size_), ++i_) { \
+		arr_[i_] = ethash_swap_u32(arr_[i_]);	\
+	}											\
+	while (0)
+#define fix_endian_arr64(arr_, size_)			\
+	do {										\
+	for (unsigned i_ = 0; i_ < (size_), ++i_) { \
+		arr_[i_] = ethash_swap_u64(arr_[i_]);	\
+	}											\
+	while (0)									\
+
 #else
 # error "endian not supported"
 #endif // BYTE_ORDER
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/ethash.h
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/ethash.h
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/fnv.h
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/fnv.h
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/internal.c
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/internal.c
@@ -257,7 +257,7 @@ static bool ethash_hash(
 void ethash_quick_hash(
 	ethash_h256_t* return_hash,
 	ethash_h256_t const* header_hash,
-	uint64_t nonce,
+	uint64_t const nonce,
 	ethash_h256_t const* mix_hash
 )
 {
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/internal.h
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/internal.h
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/io.c
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/io.c
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/io.h
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/io.h
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/io_posix.c
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/io_posix.c
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/io_win32.c
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/io_win32.c
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/mmap.h
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/mmap.h
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/mmap_win32.c
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/mmap_win32.c
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/sha3.c
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/sha3.c
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/sha3.h
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/sha3.h
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/sha3_cryptopp.cpp
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/sha3_cryptopp.cpp
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/sha3_cryptopp.h
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/sha3_cryptopp.h
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/util.h
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/util.h
--- a/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/util_win32.c
+++ b/Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/util_win32.c
--- a/Godeps/_workspace/src/github.com/fatih/color/.travis.yml
+++ b/Godeps/_workspace/src/github.com/fatih/color/.travis.yml
@@ -0,0 +1,3 @@
+language: go
+go: 1.3
+
--- a/Godeps/_workspace/src/github.com/fatih/color/LICENSE.md
+++ b/Godeps/_workspace/src/github.com/fatih/color/LICENSE.md
--- a/Godeps/_workspace/src/github.com/fatih/color/README.md
+++ b/Godeps/_workspace/src/github.com/fatih/color/README.md
@@ -2,10 +2,7 @@



-Color lets you use colorized outputs in terms of [ANSI Escape
-Codes](http://en.wikipedia.org/wiki/ANSI_escape_code#Colors) in Go (Golang). It
-has support for Windows too! The API can be used in several ways, pick one that
-suits you.
+Color lets you use colorized outputs in terms of [ANSI Escape Codes](http://en.wikipedia.org/wiki/ANSI_escape_code#Colors) in Go (Golang). It has support for Windows too! The API can be used in several ways, pick one that suits you.



@@ -81,8 +78,8 @@ info := color.New(color.FgWhite, color.BgGreen).SprintFunc()
 fmt.Printf("This %s rocks!\n", info("package"))

 // Use helper functions
-fmt.Println("This", color.RedString("warning"), "should be not neglected.")
-fmt.Printf("%v %v\n", color.GreenString("Info:"), "an important message.")
+fmt.Printf("This", color.RedString("warning"), "should be not neglected.")
+fmt.Printf(color.GreenString("Info:"), "an important message." )

 // Windows supported too! Just don't forget to change the output to color.Output
 fmt.Fprintf(color.Output, "Windows support: %s", color.GreenString("PASS"))
@@ -146,7 +143,7 @@ c.Println("This prints again cyan...")
 ## Credits

 * [Fatih Arslan](https://github.com/fatih)
- * Windows support via @mattn: [colorable](https://github.com/mattn/go-colorable)
+ * Windows support via @shiena: [ansicolor](https://github.com/shiena/ansicolor)

 ## License

--- a/Godeps/_workspace/src/github.com/fatih/color/color.go
+++ b/Godeps/_workspace/src/github.com/fatih/color/color.go
@@ -5,7 +5,6 @@ import (
 	"os"
 	"strconv"
 	"strings"
-	"sync"

 	"github.com/mattn/go-colorable"
 	"github.com/mattn/go-isatty"
@@ -313,111 +312,91 @@ func boolPtr(v bool) *bool {
 	return &v
 }

-// colorsCache is used to reduce the count of created Color objects and
-// allows to reuse already created objects with required Attribute.
-var colorsCache = make(map[Attribute]*Color)
+// Black is an convenient helper function to print with black foreground. A
+// newline is appended to format by default.
+func Black(format string, a ...interface{}) { printColor(format, FgBlack, a...) }

-var colorsCacheMu = new(sync.Mutex) // protects colorsCache
+// Red is an convenient helper function to print with red foreground. A
+// newline is appended to format by default.
+func Red(format string, a ...interface{}) { printColor(format, FgRed, a...) }

-func getCachedColor(p Attribute) *Color {
-	colorsCacheMu.Lock()
-	defer colorsCacheMu.Unlock()
+// Green is an convenient helper function to print with green foreground. A
+// newline is appended to format by default.
+func Green(format string, a ...interface{}) { printColor(format, FgGreen, a...) }

-	c, ok := colorsCache[p]
-	if !ok {
-		c = New(p)
-		colorsCache[p] = c
-	}
+// Yellow is an convenient helper function to print with yellow foreground.
+// A newline is appended to format by default.
+func Yellow(format string, a ...interface{}) { printColor(format, FgYellow, a...) }

-	return c
-}
+// Blue is an convenient helper function to print with blue foreground. A
+// newline is appended to format by default.
+func Blue(format string, a ...interface{}) { printColor(format, FgBlue, a...) }

-func colorPrint(format string, p Attribute, a ...interface{}) {
-	c := getCachedColor(p)
+// Magenta is an convenient helper function to print with magenta foreground.
+// A newline is appended to format by default.
+func Magenta(format string, a ...interface{}) { printColor(format, FgMagenta, a...) }

+// Cyan is an convenient helper function to print with cyan foreground. A
+// newline is appended to format by default.
+func Cyan(format string, a ...interface{}) { printColor(format, FgCyan, a...) }
+
+// White is an convenient helper function to print with white foreground. A
+// newline is appended to format by default.
+func White(format string, a ...interface{}) { printColor(format, FgWhite, a...) }
+
+func printColor(format string, p Attribute, a ...interface{}) {
 	if !strings.HasSuffix(format, "\n") {
 		format += "\n"
 	}

-	if len(a) == 0 {
-		c.Print(format)
-	} else {
-		c.Printf(format, a...)
-	}
+	c := &Color{params: []Attribute{p}}
+	c.Printf(format, a...)
 }

-func colorString(format string, p Attribute, a ...interface{}) string {
-	c := getCachedColor(p)
-
-	if len(a) == 0 {
-		return c.SprintFunc()(format)
-	}
-
-	return c.SprintfFunc()(format, a...)
-}
-
-// Black is an convenient helper function to print with black foreground. A
-// newline is appended to format by default.
-func Black(format string, a ...interface{}) { colorPrint(format, FgBlack, a...) }
-
-// Red is an convenient helper function to print with red foreground. A
-// newline is appended to format by default.
-func Red(format string, a ...interface{}) { colorPrint(format, FgRed, a...) }
-
-// Green is an convenient helper function to print with green foreground. A
-// newline is appended to format by default.
-func Green(format string, a ...interface{}) { colorPrint(format, FgGreen, a...) }
-
-// Yellow is an convenient helper function to print with yellow foreground.
-// A newline is appended to format by default.
-func Yellow(format string, a ...interface{}) { colorPrint(format, FgYellow, a...) }
-
-// Blue is an convenient helper function to print with blue foreground. A
-// newline is appended to format by default.
-func Blue(format string, a ...interface{}) { colorPrint(format, FgBlue, a...) }
-
-// Magenta is an convenient helper function to print with magenta foreground.
-// A newline is appended to format by default.
-func Magenta(format string, a ...interface{}) { colorPrint(format, FgMagenta, a...) }
-
-// Cyan is an convenient helper function to print with cyan foreground. A
-// newline is appended to format by default.
-func Cyan(format string, a ...interface{}) { colorPrint(format, FgCyan, a...) }
-
-// White is an convenient helper function to print with white foreground. A
-// newline is appended to format by default.
-func White(format string, a ...interface{}) { colorPrint(format, FgWhite, a...) }
-
 // BlackString is an convenient helper function to return a string with black
 // foreground.
-func BlackString(format string, a ...interface{}) string { return colorString(format, FgBlack, a...) }
+func BlackString(format string, a ...interface{}) string {
+	return New(FgBlack).SprintfFunc()(format, a...)
+}

 // RedString is an convenient helper function to return a string with red
 // foreground.
-func RedString(format string, a ...interface{}) string { return colorString(format, FgRed, a...) }
+func RedString(format string, a ...interface{}) string {
+	return New(FgRed).SprintfFunc()(format, a...)
+}

 // GreenString is an convenient helper function to return a string with green
 // foreground.
-func GreenString(format string, a ...interface{}) string { return colorString(format, FgGreen, a...) }
+func GreenString(format string, a ...interface{}) string {
+	return New(FgGreen).SprintfFunc()(format, a...)
+}

 // YellowString is an convenient helper function to return a string with yellow
 // foreground.
-func YellowString(format string, a ...interface{}) string { return colorString(format, FgYellow, a...) }
+func YellowString(format string, a ...interface{}) string {
+	return New(FgYellow).SprintfFunc()(format, a...)
+}

 // BlueString is an convenient helper function to return a string with blue
 // foreground.
-func BlueString(format string, a ...interface{}) string { return colorString(format, FgBlue, a...) }
+func BlueString(format string, a ...interface{}) string {
+	return New(FgBlue).SprintfFunc()(format, a...)
+}

 // MagentaString is an convenient helper function to return a string with magenta
 // foreground.
 func MagentaString(format string, a ...interface{}) string {
-	return colorString(format, FgMagenta, a...)
+	return New(FgMagenta).SprintfFunc()(format, a...)
 }

 // CyanString is an convenient helper function to return a string with cyan
 // foreground.
-func CyanString(format string, a ...interface{}) string { return colorString(format, FgCyan, a...) }
+func CyanString(format string, a ...interface{}) string {
+	return New(FgCyan).SprintfFunc()(format, a...)
+}

 // WhiteString is an convenient helper function to return a string with white
 // foreground.
-func WhiteString(format string, a ...interface{}) string { return colorString(format, FgWhite, a...) }
+func WhiteString(format string, a ...interface{}) string {
+	return New(FgWhite).SprintfFunc()(format, a...)
+}
--- a/Godeps/_workspace/src/github.com/fatih/color/doc.go
+++ b/Godeps/_workspace/src/github.com/fatih/color/doc.go
--- a/Godeps/_workspace/src/github.com/gizak/termui/.gitignore
+++ b/Godeps/_workspace/src/github.com/gizak/termui/.gitignore
@@ -23,4 +23,3 @@ _testmain.go
 *.test
 *.prof
 .DS_Store
-/vendor
--- a/Godeps/_workspace/src/github.com/gizak/termui/.travis.yml
+++ b/Godeps/_workspace/src/github.com/gizak/termui/.travis.yml
--- a/Godeps/_workspace/src/github.com/gizak/termui/LICENSE
+++ b/Godeps/_workspace/src/github.com/gizak/termui/LICENSE
--- a/Godeps/_workspace/src/github.com/gizak/termui/README.md
+++ b/Godeps/_workspace/src/github.com/gizak/termui/README.md
@@ -12,8 +12,6 @@ Now version v2 has arrived! It brings new event system, new theme system, new `B

 	go get -u github.com/gizak/termui

-It is recommanded to use locked deps by using [glide](https://glide.sh): move to `termui` src directory then run `glide up`.
-
 For the compatible reason, you can choose to install the legacy version of `termui`:

 	go get gopkg.in/gizak/termui.v1
@@ -26,7 +24,7 @@ To use `termui`, the very first thing you may want to know is how to manage layo

 __Absolute layout__

-Each widget has an underlying block structure which basically is a box model. It has border, label and padding properties. A border of a widget can be chosen to hide or display (with its border label), you can pick a different front/back colour for the border as well. To display such a widget at a specific location in terminal window, you need to assign `.X`, `.Y`, `.Height`, `.Width` values for each widget before sending it to `.Render`. Let's demonstrate these by a code snippet:
+Each widget has an underlying block structure which basically is a box model. It has border, label and padding properties. A border of a widget can be chosen to hide or display (with its border label), you can pick a different front/back colour for the border as well. To display such a widget at a specific location in terminal window, you need to assign `.X`, `.Y`, `.Height`, `.Width` values for each widget before send it to `.Render`. Let's demonstrate these by a code snippet:

 `````go
 	import ui "github.com/gizak/termui" // <- ui shortcut, optional
@@ -67,7 +65,7 @@ __Grid layout:__

 <img src="./_example/grid.gif" alt="grid" width="60%">

-Grid layout uses [12 columns grid system](http://www.w3schools.com/bootstrap/bootstrap_grid_system.asp) with expressive syntax. To use `Grid`, all we need to do is build a widget tree consisting of `Row`s and `Col`s (Actually a `Col` is also a `Row` but with a widget endpoint attached).
+Grid layout uses [12 columns grid system](http://www.w3schools.com/bootstrap/bootstrap_grid_system.asp) with expressive syntax. To use `Grid`, all we need to do is build a widget tree consisting of `Row`s and Cols (Actually a Col is also a `Row` but with a widget endpoint attached).

 ```go
 	import ui "github.com/gizak/termui"
--- a/Godeps/_workspace/src/github.com/gizak/termui/barchart.go
+++ b/Godeps/_workspace/src/github.com/gizak/termui/barchart.go
@@ -11,7 +11,7 @@ import "fmt"
   bc := termui.NewBarChart()
   data := []int{3, 2, 5, 3, 9, 5}
   bclabels := []string{"S0", "S1", "S2", "S3", "S4", "S5"}
-   bc.BorderLabel = "Bar Chart"
+   bc.Border.Label = "Bar Chart"
   bc.Data = data
   bc.Width = 26
   bc.Height = 10
@@ -29,7 +29,6 @@ type BarChart struct {
 	DataLabels []string
 	BarWidth   int
 	BarGap     int
-	CellChar   rune
 	labels     [][]rune
 	dataNum    [][]rune
 	numBar     int
@@ -45,7 +44,6 @@ func NewBarChart() *BarChart {
 	bc.TextColor = ThemeAttr("barchart.text.fg")
 	bc.BarGap = 1
 	bc.BarWidth = 3
-	bc.CellChar = ' '
 	return bc
 }

@@ -89,27 +87,16 @@ func (bc *BarChart) Buffer() Buffer {
 	for i := 0; i < bc.numBar && i < len(bc.Data) && i < len(bc.DataLabels); i++ {
 		h := int(float64(bc.Data[i]) / bc.scale)
 		oftX := i * (bc.BarWidth + bc.BarGap)
-
-		barBg := bc.Bg
-		barFg := bc.BarColor
-
-		if bc.CellChar == ' ' {
-			barBg = bc.BarColor
-			barFg = ColorDefault
-			if bc.BarColor == ColorDefault { // the same as above
-				barBg |= AttrReverse
-			}
-		}
-
 		// plot bar
 		for j := 0; j < bc.BarWidth; j++ {
 			for k := 0; k < h; k++ {
 				c := Cell{
-					Ch: bc.CellChar,
-					Bg: barBg,
-					Fg: barFg,
+					Ch: ' ',
+					Bg: bc.BarColor,
+				}
+				if bc.BarColor == ColorDefault { // when color is default, space char treated as transparent!
+					c.Bg |= AttrReverse
 				}
-
 				x := bc.innerArea.Min.X + i*(bc.BarWidth+bc.BarGap) + j
 				y := bc.innerArea.Min.Y + bc.innerArea.Dy() - 2 - k
 				buf.Set(x, y, c)
@@ -133,9 +120,11 @@ func (bc *BarChart) Buffer() Buffer {
 			c := Cell{
 				Ch: bc.dataNum[i][j],
 				Fg: bc.NumColor,
-				Bg: barBg,
+				Bg: bc.BarColor,
+			}
+			if bc.BarColor == ColorDefault { // the same as above
+				c.Bg |= AttrReverse
 			}
-
 			if h == 0 {
 				c.Bg = bc.Bg
 			}
--- a/Godeps/_workspace/src/github.com/gizak/termui/block.go
+++ b/Godeps/_workspace/src/github.com/gizak/termui/block.go
--- a/Godeps/_workspace/src/github.com/gizak/termui/block_common.go
+++ b/Godeps/_workspace/src/github.com/gizak/termui/block_common.go
--- a/Godeps/_workspace/src/github.com/gizak/termui/block_windows.go
+++ b/Godeps/_workspace/src/github.com/gizak/termui/block_windows.go
--- a/Show More
+++ b/Show More
Author	SHA1	Message	Date
Jeffrey Wilcke	f213a9d8e8	VERSION, cmd/geth: bumped version 1.4.3	2016-05-10 13:46:03 +02:00
Jeffrey Wilcke	6826b2040f	[release/1.4.3] miner: fixed pending state by not shutting down update loop (cherry picked from commit `a824c3f02f`)	2016-05-10 13:46:03 +02:00
Jeffrey Wilcke	2a3657d8d9	VERSION, cmd/geth: bumped version 1.4.2	2016-05-10 13:46:02 +02:00
Jeffrey Wilcke	290e851f57	VERSION, cmd/geth: bumped version 1.4.2	2016-05-09 22:06:27 +02:00
Jeffrey Wilcke	8f96d66241	Merge branch 'develop' into release/1.4	2016-05-09 22:04:40 +02:00
Jeffrey Wilcke	4b9de75623	Merge branch 'develop' into release/1.4	2016-05-03 14:05:30 +02:00
Jeffrey Wilcke	d52a693f80	Merge branch 'develop' into release/1.4 Conflicts: cmd/geth/main.go	2016-05-03 13:52:58 +02:00
Jeffrey Wilcke	8241fa5227	VERSION, cmd/geth: bumped version	2016-04-19 18:15:04 +02:00