{
 "metadata": {
  "name": "",
  "signature": "sha256:81f3deed7cdc26b0fc756b3ee1eb6e8f9b1be96304ddfc6ff484d223c2b8a942"
 },
 "nbformat": 3,
 "nbformat_minor": 0,
 "worksheets": [
  {
   "cells": [
    {
     "cell_type": "code",
     "collapsed": false,
     "input": [
      "from __future__ import division\n",
      "import numpy as np\n",
      "import pyopencl as cl\n",
      "import pyopencl.array"
     ],
     "language": "python",
     "metadata": {},
     "outputs": [
      {
       "output_type": "stream",
       "stream": "stderr",
       "text": [
        "/usr/lib/python2.7/pkgutil.py:186: ImportWarning: Not importing directory '/usr/lib/python2.7/dist-packages/enthought': missing __init__.py\n",
        "  file, filename, etc = imp.find_module(subname, path)\n"
       ]
      }
     ],
     "prompt_number": 1
    },
    {
     "cell_type": "markdown",
     "metadata": {},
     "source": [
      "Load the PyOpenCL IPython extension:"
     ]
    },
    {
     "cell_type": "code",
     "collapsed": false,
     "input": [
      "%load_ext pyopencl.ipython_ext"
     ],
     "language": "python",
     "metadata": {},
     "outputs": [],
     "prompt_number": 3
    },
    {
     "cell_type": "markdown",
     "metadata": {},
     "source": [
      "Create an OpenCL context and a command queue:"
     ]
    },
    {
     "cell_type": "code",
     "collapsed": false,
     "input": [
      "ctx = cl.create_some_context(interactive=True)\n",
      "queue = cl.CommandQueue(ctx)"
     ],
     "language": "python",
     "metadata": {},
     "outputs": [
      {
       "output_type": "stream",
       "stream": "stdout",
       "text": [
        "Choose platform:\n",
        "[0] <pyopencl.Platform 'AMD Accelerated Parallel Processing' at 0x7fc14f1b0080>\n",
        "[1] <pyopencl.Platform 'Intel(R) OpenCL' at 0x32aed00>\n"
       ]
      },
      {
       "name": "stdout",
       "output_type": "stream",
       "stream": "stdout",
       "text": [
        "Choice [0]:0\n"
       ]
      },
      {
       "output_type": "stream",
       "stream": "stdout",
       "text": [
        "Set the environment variable PYOPENCL_CTX='0' to avoid being asked again.\n"
       ]
      }
     ],
     "prompt_number": 5
    },
    {
     "cell_type": "markdown",
     "metadata": {},
     "source": [
      "-----\n",
      "\n",
      "Define an OpenCL kernel using the `%%cl_kernel` magic:"
     ]
    },
    {
     "cell_type": "code",
     "collapsed": false,
     "input": [
      "%%cl_kernel -o \"-cl-fast-relaxed-math\"\n",
      "\n",
      "__kernel void sum_vector(__global const float *a,\n",
      "__global const float *b, __global float *c)\n",
      "{\n",
      "  int gid = get_global_id(0);\n",
      "  c[gid] = a[gid] + b[gid];\n",
      "}"
     ],
     "language": "python",
     "metadata": {},
     "outputs": [],
     "prompt_number": 6
    },
    {
     "cell_type": "markdown",
     "metadata": {},
     "source": [
      "This looks for `cl_ctx` or `ctx` in the user namespace to find a PyOpenCL context.\n",
      "\n",
      "Kernel names are automatically injected into the user namespace, so we can just use `sum_vector` from Python below.\n",
      "\n",
      "Now create some data to work on:"
     ]
    },
    {
     "cell_type": "code",
     "collapsed": false,
     "input": [
      "n = 10000\n",
      "\n",
      "a = cl.array.empty(queue, n, dtype=np.float32)\n",
      "a.fill(15)\n",
      "\n",
      "b_host = np.random.randn(n).astype(np.float32)\n",
      "b = cl.array.to_device(queue, b_host)\n",
      "\n",
      "c = cl.array.empty_like(a)"
     ],
     "language": "python",
     "metadata": {},
     "outputs": [],
     "prompt_number": 7
    },
    {
     "cell_type": "markdown",
     "metadata": {},
     "source": [
      "Run the kernel:"
     ]
    },
    {
     "cell_type": "code",
     "collapsed": false,
     "input": [
      "sum_vector(queue, (n,), None, a.data, b.data, c.data)"
     ],
     "language": "python",
     "metadata": {},
     "outputs": [
      {
       "metadata": {},
       "output_type": "pyout",
       "prompt_number": 8,
       "text": [
        "<pyopencl._cl.Event at 0x7fc14f3fdf30>"
       ]
      }
     ],
     "prompt_number": 8
    },
    {
     "cell_type": "markdown",
     "metadata": {},
     "source": [
      "Check the result using `numpy` operations:"
     ]
    },
    {
     "cell_type": "code",
     "collapsed": false,
     "input": [
      "assert (c.get() == b_host + 15).all()"
     ],
     "language": "python",
     "metadata": {},
     "outputs": [],
     "prompt_number": 9
    }
   ],
   "metadata": {}
  }
 ]
}