Newer
Older
Andreas Klöckner
committed
"signature": "sha256:81f3deed7cdc26b0fc756b3ee1eb6e8f9b1be96304ddfc6ff484d223c2b8a942"
},
"nbformat": 3,
"nbformat_minor": 0,
"worksheets": [
{
"cells": [
{
"cell_type": "code",
"collapsed": false,
"input": [
"from __future__ import division\n",
"import pyopencl as cl\n",
"import pyopencl.array"
],
"language": "python",
"metadata": {},
Andreas Klöckner
committed
"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": [
Andreas Klöckner
committed
"%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",
Andreas Klöckner
committed
"[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"
]
}
],
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"-----\n",
"\n",
"Define an OpenCL kernel using the `%%cl_kernel` magic:"
]
},
{
"cell_type": "code",
"collapsed": false,
"input": [
Alex Rothberg
committed
"%%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": [],
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
},
{
"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": [],
},
{
"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",
Andreas Klöckner
committed
"<pyopencl._cl.Event at 0x7fc14f3fdf30>"
},
{
"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": [],
}
],
"metadata": {}
}
]
}