{
  "cells": [
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "# Practice: Generating a Simple Kernel\n",
        "\n",
        "The purpose of this practice problem is to generate a simple kernel that applies a user-supplied expression to every entry of an array. Implement a class `ExpressionKernel` that can be used as shown in the test at the end of this notebook."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 2,
      "metadata": {
        "collapsed": false
      },
      "outputs": [],
      "source": [
        "import numpy as np\n",
        "import numpy.linalg as la\n",
        "\n",
        "import pyopencl as cl\n",
        "import pyopencl.array\n",
        "import pyopencl.clmath\n",
        "import pyopencl.clrandom"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 3,
      "metadata": {
        "collapsed": false
      },
      "outputs": [],
      "source": [
        "\n",
        "class ExpressionKernel:\n",
        "    def __init__(self, cl_context, expression):\n",
        "        # ...\n",
        "        pass\n",
        "    \n",
        "    def __call__(self, queue, ary):\n",
        "        # ...\n",
        "        pass"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 4,
      "metadata": {
        "collapsed": false
      },
      "outputs": [],
      "source": [
        "# Solution\n",
        "\n",
        "class ExpressionKernel:\n",
        "    def __init__(self, cl_context, expression):\n",
        "        src = \"\"\"\n",
        "            kernel void apply(__global double *out, global double *in)\n",
        "            {\n",
        "                int i = get_global_id(0);\n",
        "                double x = in[i];\n",
        "                out[i] = RESULT;\n",
        "            }\n",
        "            \"\"\"\n",
        "\n",
        "        from pymbolic.mapper.c_code import CCodeMapper\n",
        "        ccm = CCodeMapper()\n",
        "        src = src.replace(\"RESULT\", ccm(expression))\n",
        "        self.prg = cl.Program(cl_context, src).build()\n",
        "        self.knl = self.prg.apply\n",
        "\n",
        "    def __call__(self, queue, ary):\n",
        "        result = cl.array.empty_like(ary)\n",
        "        self.knl(queue, ary.shape, None, result.data, ary.data)\n",
        "        return result"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "To test our implementation, we create a context and an array full of random numbers:"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 5,
      "metadata": {
        "collapsed": false
      },
      "outputs": [],
      "source": [
        "cl_context = cl.create_some_context()\n",
        "queue = cl.CommandQueue(cl_context)\n",
        "\n",
        "ary = cl.clrandom.rand(queue, 500, dtype=np.float64)"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 6,
      "metadata": {
        "collapsed": false
      },
      "outputs": [
        {
          "name": "stdout",
          "output_type": "stream",
          "text": [
            "9.76586150045e-16\n"
          ]
        }
      ],
      "source": [
        "\n",
        "from pymbolic import var\n",
        "\n",
        "x = var(\"x\")\n",
        "eknl = ExpressionKernel(cl_context, var(\"sqrt\")(1-x**2))\n",
        "\n",
        "result = eknl(queue, ary)\n",
        "\n",
        "diff = result - cl.clmath.sqrt(1-ary**2)\n",
        "print(la.norm(diff.get()))"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "collapsed": true
      },
      "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.1+"
    }
  },
  "nbformat": 4,
  "nbformat_minor": 0
}