5 - Troubleshooting and Debugging.ipynb

{
 "cells": [
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "# GTC 2019 Numba Tutorial Notebook 5: Troubleshooting and Debugging\n",
    "\n",
    "## Note about the Terminal\n",
    "\n",
    "Debugging is an important part of programming.  Unfortuntely, it is pretty difficult to debug CUDA kernels directly in the Jupyter notebook for a variety of reasons, so this notebook will show terminal commands by executing Jupyter notebook cells using the shell.  These shell commands will appear in notebook cells with the command line prefixed by `!`. When applying the debug methods described in this notebook, you will likely run the commands in the terminal directly.\n",
    "\n",
    "## Printing\n",
    "\n",
    "A common debugging strategy is printing to the console.  Numba supports printing from CUDA kernels, with some restrictions.  Note that output printed from a CUDA kernel will not be captured by Jupyter, so you will need to debug with a script you can run from the terminal.\n",
    "\n",
    "Let's look at a CUDA kernel with a bug:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "! cat debug/ex1.py"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "collapsed": true
   },
   "source": [
    "When we run this code to histogram 50 values, we see the histogram is not getting 50 entries: "
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "! python debug/ex1.py"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "*(You might have already spotted the mistake, but let's pretend we don't know the answer.)*\n",
    "\n",
    "We hypothesize that maybe a bin calculation error is causing many of the histogram entries to appear out of range.  Let's add some printing around the `if` statement to show us what is going on:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "! cat debug/ex1a.py"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "This kernel will print every value and bin number it calculates.  Looking at one of the print statements, we see that `print` supports constant strings, and scalar values:\n",
    "\n",
    "``` python\n",
    "print('in range', x[i], bin_number)\n",
    "```\n",
    "\n",
    "String substitution (using C printf syntax or the newer `format()` syntax) is not supported.  If we run this script we see:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "! python debug/ex1a.py"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Scanning down that output, we see that all 50 values should be in range.  Clearly we have some kind of race condition updating the histogram.  In fact, the culprit line is:\n",
    "\n",
    "``` python\n",
    "histogram_out[bin_number] += 1\n",
    "```\n",
    "\n",
    "which should be (as you may have seen in a previous exercise)\n",
    "\n",
    "``` python\n",
    "cuda.atomic.add(histogram_out, bin_number, 1)\n",
    "```"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "collapsed": true
   },
   "source": [
    "## CUDA Simulator\n",
    "\n",
    "Back in the early days of CUDA, `nvcc` had an \"emulator\" mode that would execute CUDA code on the CPU for debugging.  That functionality was dropped in later CUDA releases after `cuda-gdb` was created.  There isn't a debugger for CUDA+Python, so Numba includes a \"CUDA simulator\" in Numba that runs your CUDA code with the Python interpreter on the host CPU.  This allows you to debug the logic of your code using Python modules and functions that would otherwise be not allowed by the compile.\n",
    "\n",
    "A very common use case is to start the Python debugger inside one thread of a CUDA kernel:\n",
    "``` python\n",
    "import numpy as np\n",
    "\n",
    "from numba import cuda\n",
    "\n",
    "@cuda.jit\n",
    "def histogram(x, xmin, xmax, histogram_out):\n",
    "    nbins = histogram_out.shape[0]\n",
    "    bin_width = (xmax - xmin) / nbins\n",
    "\n",
    "    start = cuda.grid(1)\n",
    "    stride = cuda.gridsize(1)\n",
    "\n",
    "    ### DEBUG FIRST THREAD\n",
    "    if start == 0:\n",
    "        from pdb import set_trace; set_trace()\n",
    "    ###\n",
    "\n",
    "    for i in range(start, x.shape[0], stride):\n",
    "        bin_number = np.int32((x[i] + xmin)/bin_width)\n",
    "\n",
    "        if bin_number >= 0 and bin_number < histogram_out.shape[0]:\n",
    "            cuda.atomic.add(histogram_out, bin_number, 1)\n",
    "\n",
    "x = np.random.normal(size=50, loc=0, scale=1).astype(np.float32)\n",
    "xmin = np.float32(-4.0)\n",
    "xmax = np.float32(4.0)\n",
    "histogram_out = np.zeros(shape=10, dtype=np.int32)\n",
    "\n",
    "histogram[64, 64](x, xmin, xmax, histogram_out)\n",
    "\n",
    "print('input count:', x.shape[0])\n",
    "print('histogram:', histogram_out)\n",
    "print('count:', histogram_out.sum())\n",
    "```\n",
    "\n",
    "This code allows a debug session like the following to take place:\n",
    "```\n",
    "(gtc2017) 0179-sseibert:gtc2017-numba sseibert$ NUMBA_ENABLE_CUDASIM=1 python debug/ex2.py\n",
    "> /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex2.py(18)histogram()\n",
    "-> for i in range(start, x.shape[0], stride):\n",
    "(Pdb) n\n",
    "> /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex2.py(19)histogram()\n",
    "-> bin_number = np.int32((x[i] + xmin)/bin_width)\n",
    "(Pdb) n\n",
    "> /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex2.py(21)histogram()\n",
    "-> if bin_number >= 0 and bin_number < histogram_out.shape[0]:\n",
    "(Pdb) p bin_number, x[i]\n",
    "(-6, -1.4435024)\n",
    "(Pdb) p x[i], xmin, bin_width\n",
    "(-1.4435024, -4.0, 0.80000000000000004)\n",
    "(Pdb) p (x[i] - xmin) / bin_width\n",
    "3.1956219673156738\n",
    "(Pdb) q\n",
    "```"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "collapsed": true
   },
   "source": [
    "## CUDA Memcheck\n",
    "\n",
    "Another common error occurs when a CUDA kernel has an invalid memory access, typically caused by running off the end of an array.  The full CUDA toolkit from NVIDIA (not the `cudatoolkit` conda package) contain a utility called `cuda-memcheck` that can check for a wide range of memory access mistakes in CUDA code.\n",
    "\n",
    "Let's debug the following code:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "! cat debug/ex3.py"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "! cuda-memcheck python debug/ex3.py"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "The output of `cuda-memcheck` is clearly showing a problem with our histogram function:\n",
    "```\n",
    "========= Invalid __global__ write of size 4\n",
    "=========     at 0x00000548 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
    "```\n",
    "But we don't know which line it is.  To get better error information, we can turn \"debug\" mode on when compiling the kernel, by changing the kernel to look like this:\n",
    "``` python\n",
    "@cuda.jit(debug=True)\n",
    "def histogram(x, xmin, xmax, histogram_out):\n",
    "    nbins = histogram_out.shape[0]\n",
    "```"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "! cuda-memcheck python debug/ex3a.py"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Now we get an error message that includes a source file and line number: `ex3a.py:17`."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "! cat -n debug/ex3a.py | grep -C 2 \"17\""
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "At this point, we might realize that our if statement incorrect has an `or` instead of an `and`.\n",
    "\n",
    "`cuda-memcheck` has different modes for detecting different kinds of problems (similar to `valgrind` for debugging CPU memory access errors).  Take a look at the documentation for more information: http://docs.nvidia.com/cuda/cuda-memcheck/"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "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.7.1"
  }
 },
 "nbformat": 4,
 "nbformat_minor": 2
}