{
  "cells": [
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "# PyOpenCL: An exercise"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 1,
      "metadata": {
        "collapsed": false
      },
      "outputs": [],
      "source": [
        "import pyopencl as cl\n",
        "import numpy as np\n",
        "import numpy.linalg as la\n",
        "import pyopencl.array\n",
        "import pyopencl.clrandom"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "Change the code below to:\n",
        "    \n",
        "* Compute $c_i = a_ib_i$\n",
        "* Use work groups of $16\\times 16$ items\n",
        "* Benchmark $1\\times 1$ workgroups against $16\\times 16$ workgroups\n",
        "\n",
        "  * Use `time()` from the `time` module. (i.e. `import time`)\n",
        "  * Use `queue.finish()`."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 2,
      "metadata": {
        "collapsed": false
      },
      "outputs": [],
      "source": [
        "ctx = cl.create_some_context()\n",
        "queue = cl.CommandQueue(ctx)"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 3,
      "metadata": {
        "collapsed": false
      },
      "outputs": [],
      "source": [
        "a = np.random.rand(1024, 1024).astype(np.float32)"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 4,
      "metadata": {
        "collapsed": false
      },
      "outputs": [],
      "source": [
        "prg = cl.Program(ctx, \"\"\"\n",
        "    __kernel void twice(__global float *a)\n",
        "    {\n",
        "      int gid0 = get_global_id(0);\n",
        "      int gid1 = get_global_id(1);\n",
        "      int i = gid1 * 1024 + gid0;\n",
        "      a[i] = 2*a[i];\n",
        "    }\n",
        "    \"\"\").build()\n",
        "twice = prg.twice"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 5,
      "metadata": {
        "collapsed": false
      },
      "outputs": [
        {
          "data": {
            "text/plain": [
              "<pyopencl.cffi_cl.Event at 0x7f679a3145f8>"
            ]
          },
          "execution_count": 5,
          "metadata": {},
          "output_type": "execute_result"
        }
      ],
      "source": [
        "a_dev = cl.array.to_device(queue, a)\n",
        "twice(queue, a_dev.shape, None, a_dev.data)"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 6,
      "metadata": {
        "collapsed": false
      },
      "outputs": [
        {
          "name": "stdout",
          "output_type": "stream",
          "text": [
            "0.0 591.347\n"
          ]
        }
      ],
      "source": [
        "print(la.norm(a_dev.get() - 2*a), la.norm(a))"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "collapsed": false
      },
      "outputs": [],
      "source": []
    }
  ],
  "metadata": {
    "kernelspec": {
      "display_name": "Python 3",
      "language": "python",
      "name": "python3"
    },
    "language_info": {
      "codemirror_mode": {
        "name": "ipython",
        "version": 3
      },
      "file_extension": ".py",
      "mimetype": "text/x-python",
      "name": "python",
      "nbconvert_exporter": "python",
      "pygments_lexer": "ipython3",
      "version": "3.5.0+"
    }
  },
  "nbformat": 4,
  "nbformat_minor": 0
}