fix typos and formatting

Author: benvanwerkhoven

tutorial/grid3d.ipynb

@@ -6,7 +6,7 @@
    "source": [
     "# 3D Grid on GPU with Kernel Tuner\n",
     "\n",
-    "In this tutrial we are going to see how to map a series of Gaussian functions, each located at a different point on a 3D a grid. We are going to optimize the kernel of the GPU code and compare its performance with the CPU implementation. \n",
+    "In this tutorial we are going to see how to map a series of Gaussian functions, each located at a different point on a 3D a grid. We are going to optimize the GPU code and compare its performance with the CPU implementation. \n",
     "\n",
     "<div class=\"alert alert-info\">\n",
     "\n",
@@ -21,13 +21,13 @@
    "source": [
     "## Let's start on the CPU\n",
     "\n",
-    "Before dwelving on the GPU implementation, let's start with a simple CPU implementation of the problem. The problem at hand is to compute the values of the following function \n",
+    "Before delving into the GPU implementation, let's start with a simple CPU implementation of the problem. The problem at hand is to compute the values of the following function \n",
     "\n",
     "\\begin{equation} \\nonumber\n",
     "f = \\sum_{i=1}^{N}\\exp\\left(-\\beta \\sqrt{(x-x_i)^2+(y-y_i)^2+(z-z_i)^2}\\right)\n",
     "\\end{equation}\n",
     "\n",
-    "on a 3d grid. The $x$, $y$ and $z$ vectors contain here the coordinate of the points in the cartesian space. We can define a simple python function that computes the value of the function $f$ for one given Gaussian.  "
+    "on a 3d grid. The $x$, $y$ and $z$ vectors contain the coordinate of the points in the Cartesian space. We can define a simple Python function that computes the value of the function $f$ for one given Gaussian. Don't forget to execute all the code cells, like the one below, as you read through this notebook by selecting the cell and pressing *shift+enter*."
    ]
   },
   {
@@ -54,15 +54,17 @@
    "cell_type": "markdown",
    "metadata": {},
    "source": [
-    "For a given center, this function returns the values of the corresponding Gaussian function mapped on the 3d Grid. The grid points are here defined by the variables `xgrid`, `ygrid` and `zgrid`. These variables are themselves 3D grid obtained as we wil see in an instant with the `numpy.meshgrid` function. \n",
+    "For a given center, this function returns the values of the corresponding Gaussian function mapped on the 3D grid. The grid points are here defined by the variables `xgrid`, `ygrid` and `zgrid`. These variables are themselves 3D grids obtained, as we will see in an instant, using the `numpy.meshgrid` function. \n",
     "\n",
-    "To use this function we simply have to create the x,y and z grids. Since we want to later on send these vectors on the GPU we define them as float. For simplicity we here select the interval $[-1:1]$ to define the grid. We use $n=256$ grid points in order to have a sufficiently large probelm without requiring too long calculations. We then create meshgrids to be passed to the function above. We define here 100 gaussian centers that are randomly distributed within the grid space."
+    "To use this function we simply have to create the grid, defined by the vectors x, y, and z. Since we want to later on send these vectors to the GPU we define them as 32-bit floats. For simplicity, we here select the interval $[-1:1]$ to define our grid. We use $n=256$ grid points in order to have a sufficiently large problem without requiring too long calculations. We then create meshgrids to be passed to the function above. We define here 100 gaussian centers that are randomly distributed within the 3D space."
    ]
   },
   {
    "cell_type": "code",
    "execution_count": 2,
-   "metadata": {},
+   "metadata": {
+    "collapsed": false
+   },
    "outputs": [
     {
      "name": "stdout",
@@ -88,7 +90,6 @@
     "# centers\n",
     "npts = 100\n",
     "center = (-1 + 2*np.random.rand(npts,3)).astype(np.float32)\n",
-    "center = center.astype(np.float32)\n",
     "\n",
     "# compute the grid and time the operation\n",
     "t0 = time()\n",
@@ -101,18 +102,18 @@
    "cell_type": "markdown",
    "metadata": {},
    "source": [
-    "Depending on your hardware it might take a few seconds for the calculations of the function to finish."
+    "Depending on your hardware it might take a few seconds for the calculations above to finish."
    ]
   },
   {
    "cell_type": "markdown",
    "metadata": {},
    "source": [
-    "## Let's go on GPU\n",
+    "## Let's move to the GPU\n",
     "\n",
-    "Let's see now how that will look like on the GPU. We first write a kernel template that does the same calculation as the function above. As yp can see see below, the variables *block_size_x, block_size_y* and *block_size_z* are npt explicitly defined here. These variables define how many thread will run simultaneously on the GPU and are the main parameters that the kernel tuner will optimize.Therefore during the tuning phase, the kernel tuner will automatically insert **#define** statements for these parameters at the top of the kernel code. So for now we don't have to sepcify their values. \n",
+    "Let's see now how that will look like on the GPU. We first write a kernel that does the same calculation as the above function. As you can see see below, the variables `block_size_x`, `block_size_y` and `block_size_z` are not yet defined here. These variables are used to set the number of threads per thread block on the GPU and are the main parameters that we will optimize in this tutorial. During tuning, Kernel Tuner will automatically insert `#define` statements for these parameters at the top of the kernel code. So for now we don't have to specify their values. \n",
     "\n",
-    "The dimensions of the problem, called here  *nx, ny, nz*, are parameters of the template. We are going to see in the following how to use this template to generate a kernel code with specific dimensions."
+    "The dimensions of the problem `nx`, `ny`, and `nz`, are the number of grid points in the x, y, and z dimensions. We can again use Kernel Tuner to insert these parameters into the code."
    ]
   },
   {
@@ -127,7 +128,7 @@
     "# several parameters are available\n",
     "# block sizes : bx, by, bz \n",
     "# dimensions  : nx, ny, nz\n",
-    "kernel_code_template = \"\"\"\n",
+    "kernel_code = \"\"\"\n",
     "#include <math.h>\n",
     "\n",
     "// a simple gaussian function\n",
@@ -146,14 +147,14 @@
     "    int y = threadIdx.y + block_size_y * blockIdx.y;\n",
     "    int z = threadIdx.z + block_size_z * blockIdx.z;\n",
     "\n",
-    "    if ( ( x < %(nx)s ) && (y < %(ny)s) && (z < %(nz)s) )\n",
+    "    if ( ( x < nx ) && (y < ny) && (z < nz) )\n",
     "    {\n",
     "\n",
     "        float dx = xvect[x]-x0;\n",
     "        float dy = yvect[y]-y0;\n",
     "        float dz = zvect[z]-z0;\n",
     "        float d = sqrt(dx*dx + dy*dy + dz*dz);\n",
-    "        out[y * %(nx)s * %(nz)s + x * %(nz)s + z] = f(d);\n",
+    "        out[y * nx * nz + x * nz + z] = f(d);\n",
     "    }\n",
     "}\n",
     "\"\"\""
@@ -165,12 +166,14 @@
    "source": [
     "### Tune the kernel\n",
     "\n",
-    "We can now use the kernel tuner to optimize the block sizes on our GPU. To do so we define the tune_params dictionary that assigns to each block size the values we want the kernel tuner to explore. We also define a list containing the arguments of the CUDA function (AddGrid) above. Since we only want to optimize the performance of the kernel we only consider here one center in the middle of the grid. Note that the kernel tuner needs either numpy.ndarray or numpy.scalar as argumenst of the kernel. Hence we need to be specific on the types of the gaussians positions. "
+    "We can now use the tuner to optimize the thread block dimensions on our GPU. To do so we define the tunable parameters of our kernel using the `tune_params` dictionary, which assigns to each block size the values we want the tuner to explore. We also use the tunable parameters to insert the domain dimensions `nx`, `ny`, and `nz`.\n",
+    "\n",
+    "We also define a list containing the arguments of the CUDA function (AddGrid) above. Since we only want to optimize the performance of the kernel we only consider here one center in the middle of the grid. Note that Kernel Tuner needs either `numpy.ndarray` or `numpy.scalar` as arguments of the kernel. Hence we need to be specific on the types of the Gaussians positions. "
    ]
   },
   {
    "cell_type": "code",
-   "execution_count": 4,
+   "execution_count": null,
    "metadata": {
     "collapsed": true,
     "scrolled": true
@@ -185,6 +188,9 @@
     "tune_params['block_size_x'] = [2,4,8,16,32]\n",
     "tune_params['block_size_y'] = [2,4,8,16,32]\n",
     "tune_params['block_size_z'] = [2,4,8,16,32]\n",
+    "tune_params['nx'] = [n]\n",
+    "tune_params['ny'] = [n]\n",
+    "tune_params['nz'] = [n]\n",
     "\n",
     "# define the final grid\n",
     "grid = np.zeros_like(xgrid)\n",
@@ -201,13 +207,15 @@
    "cell_type": "markdown",
    "metadata": {},
    "source": [
-    "We then generate the kernel code by replacing the dimensions of the problem by their values. This is  simply done by using the kernel template to replace the variables *nx,ny,nz*, that are written *%(nx)s, %(ny)s, %(nz)s* in the template by their values. As mentionned earlier, the kernel tuner will automatically insert #define statements at the top of the kernel to define the block sizes so we don't need to specify them here. Then we simply call the tune_kernel function. "
+    "As mentioned earlier, the tuner will automatically insert `#define` statements at the top of the kernel to define the block sizes and domain dimensions, so we don't need to specify them here. Then, we simply call the `tune_kernel` function. "
    ]
   },
   {
    "cell_type": "code",
    "execution_count": 5,
-   "metadata": {},
+   "metadata": {
+    "collapsed": false
+   },
    "outputs": [
     {
      "name": "stdout",
@@ -309,18 +317,17 @@
     }
    ],
    "source": [
-    "# generate the kernel code from the template\n",
-    "kernel_code = kernel_code_template % {'nx' : n, 'ny': n, 'nz' : n}\n",
-    "\n",
     "# call the kernel tuner\n",
-    "result = tune_kernel('AddGrid', kernel_code,problem_size,args,tune_params)"
+    "result = tune_kernel('AddGrid', kernel_code, problem_size, args, tune_params)"
    ]
   },
   {
    "cell_type": "markdown",
    "metadata": {},
    "source": [
-    "The kernel explore all the possible combinations of tunable parameters (here only the block size). For each of them the kernel automatically introduces **#define** statements that specify the block sizes, compile the code and time the execution. At the end of the the run the tuner output the optimal conbination of the tunable parameters. As you can see the range of performances is quite large. With our GPU (GeForce GTX 1080 Ti) we obtained a maximum time of 5.30 ms and minimum one of 0.84 ms. Optimising the kernel allows dividing the execution time by a factor 6 !"
+    "The `tune_kernel` function explores all the possible combinations of tunable parameters (here only the block size). For each possible kernel configuration, the tuner compiles the code and its measures execution time (by default using 7 iterations). At the end of the the run, the `tune_kernel` outputs the optimal combination of the tunable parameters. But the measured execution time of all benchmarked kernels is also returned by `tune_kernel` for programmatic access to the data.\n",
+    "\n",
+    "As you can see the range of performances is quite large. With our GPU (GeForce GTX 1080 Ti) we obtained a maximum time of 5.30 ms and minimum one of 0.84 ms. The performance of the kernel varies by a factor 6 depending on the thread block size!"
    ]
   },
   {
@@ -344,17 +351,17 @@
     "import pycuda.autoinit\n",
     "\n",
     "# optimal values of the block size\n",
-    "block = [4,2,16]\n",
+    "block = [4, 2, 16]\n",
     "\n",
     "# corresponding grid size\n",
-    "grid_dim = [ int(np.ceil(n/b)) for b,n in zip(block,problem_size)]"
+    "grid_dim = [int(np.ceil(n/b)) for b, n in zip(block, problem_size)]"
    ]
   },
   {
    "cell_type": "markdown",
    "metadata": {},
    "source": [
-    "Before using the kernel we need to specify the block size in its definition. There are different ways of doing this, we here simply replace the *block_size_x,block_size_y* and *block_size_z* by their values determined by the tuner.Inorder to do that we create a dictionary that associates the name of the block size and their values and simply make the substitution. Once the block size are specified, we can compile the kernel ourselves and get the function."
+    "Before using the kernel we need to specify the block size in its definition. There are different ways of doing this, we here simply replace the `block_size_x`, `block_size_y` and `block_size_z` by their values determined by the tuner. In order to do that we create a dictionary that associates the name of the block size and their values and simply make the substitution. Once the block size are specified, we can compile the kernel ourselves and get the function."
    ]
   },
   {
@@ -370,6 +377,9 @@
     "fixed_params['block_size_x'] = block[0]\n",
     "fixed_params['block_size_y'] = block[1]\n",
     "fixed_params['block_size_z'] = block[2]\n",
+    "fixed_params['nx'] = n\n",
+    "fixed_params['ny'] = n\n",
+    "fixed_params['nz'] = n\n",
     "\n",
     "for k,v in fixed_params.items():\n",
     "    kernel_code = kernel_code.replace(k,str(v))\n",
@@ -389,7 +399,9 @@
   {
    "cell_type": "code",
    "execution_count": 8,
-   "metadata": {},
+   "metadata": {
+    "collapsed": false
+   },
    "outputs": [
     {
      "name": "stdout",
@@ -441,7 +453,7 @@
    "name": "python",
    "nbconvert_exporter": "python",
    "pygments_lexer": "ipython3",
-   "version": "3.6.1"
+   "version": "3.5.1"
   }
  },
  "nbformat": 4,