{
 "cells": [
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "cc7d0709",
   "metadata": {
    "collapsed": false,
    "jupyter": {
     "outputs_hidden": false
    }
   },
   "outputs": [],
   "source": [
    "from __future__ import division\n",
    "import numpy as np\n",
    "import pyopencl as cl\n",
    "import pyopencl.array"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "8ac8d7bb",
   "metadata": {},
   "source": [
    "Load the PyOpenCL IPython extension:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "7023ca2f",
   "metadata": {
    "collapsed": false,
    "jupyter": {
     "outputs_hidden": false
    }
   },
   "outputs": [],
   "source": [
    "%load_ext pyopencl.ipython_ext"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "9544b53c",
   "metadata": {},
   "source": [
    "Create an OpenCL context and a command queue:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "fac17999",
   "metadata": {
    "collapsed": false,
    "jupyter": {
     "outputs_hidden": false
    }
   },
   "outputs": [],
   "source": [
    "ctx = cl.create_some_context(interactive=True)\n",
    "queue = cl.CommandQueue(ctx)"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "a29daf04",
   "metadata": {},
   "source": [
    "-----\n",
    "\n",
    "Define an OpenCL kernel using the `%%cl_kernel` magic:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "65c7e6c9",
   "metadata": {
    "collapsed": false,
    "jupyter": {
     "outputs_hidden": false
    }
   },
   "outputs": [],
   "source": [
    "%%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",
    "}"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "cfb57357",
   "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",
   "execution_count": null,
   "id": "1d80ff38",
   "metadata": {
    "collapsed": false,
    "jupyter": {
     "outputs_hidden": false
    }
   },
   "outputs": [],
   "source": [
    "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)"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "61fccb61",
   "metadata": {},
   "source": [
    "Run the kernel:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "2ba991b3",
   "metadata": {
    "collapsed": false,
    "jupyter": {
     "outputs_hidden": false
    }
   },
   "outputs": [],
   "source": [
    "sum_vector(queue, (n,), None, a.data, b.data, c.data)"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "11a55b38",
   "metadata": {},
   "source": [
    "Check the result using `numpy` operations:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "ee3560c1",
   "metadata": {
    "collapsed": false,
    "jupyter": {
     "outputs_hidden": false
    }
   },
   "outputs": [],
   "source": [
    "assert (c.get() == b_host + 15).all()"
   ]
  }
 ],
 "metadata": {
  "kernelspec": {
   "display_name": "Python 3 (ipykernel)",
   "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.10.5"
  }
 },
 "nbformat": 4,
 "nbformat_minor": 5
}