5 - Troubleshooting and Debugging.ipynb

{
 "cells": [
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "# GTC 2017 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": 2,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "import numpy as np\r\n",
      "\r\n",
      "from numba import cuda\r\n",
      "\r\n",
      "@cuda.jit\r\n",
      "def histogram(x, xmin, xmax, histogram_out):\r\n",
      "    nbins = histogram_out.shape[0]\r\n",
      "    bin_width = (xmax - xmin) / nbins\r\n",
      "\r\n",
      "    start = cuda.grid(1)\r\n",
      "    stride = cuda.gridsize(1)\r\n",
      "\r\n",
      "    for i in range(start, x.shape[0], stride):\r\n",
      "        bin_number = np.int32((x[i] - xmin)/bin_width)\r\n",
      "        if bin_number >= 0 and bin_number < histogram_out.shape[0]:\r\n",
      "            histogram_out[bin_number] += 1\r\n",
      "\r\n",
      "x = np.random.normal(size=50, loc=0, scale=1).astype(np.float32)\r\n",
      "xmin = np.float32(-4.0)\r\n",
      "xmax = np.float32(4.0)\r\n",
      "histogram_out = np.zeros(shape=10, dtype=np.int32)\r\n",
      "\r\n",
      "histogram[64, 64](x, xmin, xmax, histogram_out)\r\n",
      "\r\n",
      "print('input count:', x.shape[0])\r\n",
      "print('histogram:', histogram_out)\r\n",
      "print('count:', histogram_out.sum())\r\n"
     ]
    }
   ],
   "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": 3,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "input count: 50\r\n",
      "histogram: [0 0 1 1 1 1 1 1 0 0]\r\n",
      "count: 6\r\n"
     ]
    }
   ],
   "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": 6,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "import numpy as np\r\n",
      "\r\n",
      "from numba import cuda\r\n",
      "\r\n",
      "@cuda.jit\r\n",
      "def histogram(x, xmin, xmax, histogram_out):\r\n",
      "    nbins = histogram_out.shape[0]\r\n",
      "    bin_width = (xmax - xmin) / nbins\r\n",
      "\r\n",
      "    start = cuda.grid(1)\r\n",
      "    stride = cuda.gridsize(1)\r\n",
      "\r\n",
      "    for i in range(start, x.shape[0], stride):\r\n",
      "        bin_number = np.int32((x[i] - xmin)/bin_width)\r\n",
      "        if bin_number >= 0 and bin_number < histogram_out.shape[0]:\r\n",
      "            histogram_out[bin_number] += 1\r\n",
      "            print('in range', x[i], bin_number)\r\n",
      "        else:\r\n",
      "            print('out of range', x[i], bin_number)\r\n",
      "\r\n",
      "x = np.random.normal(size=50, loc=0, scale=1).astype(np.float32)\r\n",
      "xmin = np.float32(-4.0)\r\n",
      "xmax = np.float32(4.0)\r\n",
      "histogram_out = np.zeros(shape=10, dtype=np.int32)\r\n",
      "\r\n",
      "histogram[64, 64](x, xmin, xmax, histogram_out)\r\n",
      "\r\n",
      "print('input count:', x.shape[0])\r\n",
      "print('histogram:', histogram_out)\r\n",
      "print('count:', histogram_out.sum())\r\n"
     ]
    }
   ],
   "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": 7,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "in range 1.674757 7\r\n",
      "in range -0.492113 4\r\n",
      "in range 0.526627 5\r\n",
      "in range 2.359267 7\r\n",
      "in range -1.768394 2\r\n",
      "in range 0.342256 5\r\n",
      "in range 0.793954 5\r\n",
      "in range -0.338127 4\r\n",
      "in range 1.275327 6\r\n",
      "in range -0.877891 3\r\n",
      "in range 0.922818 6\r\n",
      "in range 0.635215 5\r\n",
      "in range 0.371592 5\r\n",
      "in range 0.925639 6\r\n",
      "in range -1.116025 3\r\n",
      "in range 0.615792 5\r\n",
      "in range 0.879030 6\r\n",
      "in range 2.061845 7\r\n",
      "in range 0.037717 5\r\n",
      "in range -0.440858 4\r\n",
      "in range 1.056680 6\r\n",
      "in range -0.111198 4\r\n",
      "in range 0.452880 5\r\n",
      "in range -0.154099 4\r\n",
      "in range 0.518296 5\r\n",
      "in range 0.072946 5\r\n",
      "in range 1.209770 6\r\n",
      "in range -0.057651 4\r\n",
      "in range 0.154896 5\r\n",
      "in range 1.099341 6\r\n",
      "in range 0.271862 5\r\n",
      "in range 0.643499 5\r\n",
      "in range 0.824574 6\r\n",
      "in range 0.809260 6\r\n",
      "in range 0.354412 5\r\n",
      "in range -0.365111 4\r\n",
      "in range 0.594393 5\r\n",
      "in range 0.830470 6\r\n",
      "in range -0.402743 4\r\n",
      "in range -0.554546 4\r\n",
      "in range -0.507898 4\r\n",
      "in range -0.006359 4\r\n",
      "in range -0.316683 4\r\n",
      "in range 2.015556 7\r\n",
      "in range -1.288521 3\r\n",
      "in range 0.401858 5\r\n",
      "in range -1.410364 3\r\n",
      "in range -0.459540 4\r\n",
      "in range 0.633090 5\r\n",
      "in range -1.000390 3\r\n",
      "input count: 50\r\n",
      "histogram: [0 0 1 2 2 2 2 2 0 0]\r\n",
      "count: 11\r\n"
     ]
    }
   ],
   "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.  That functionality was dropped in later CUDA releases after `cuda-gdb` was created.  We missed emulator mode so much, 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": 8,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "import numpy as np\r\n",
      "\r\n",
      "from numba import cuda\r\n",
      "\r\n",
      "@cuda.jit\r\n",
      "def histogram(x, xmin, xmax, histogram_out):\r\n",
      "    nbins = histogram_out.shape[0]\r\n",
      "    bin_width = (xmax - xmin) / nbins\r\n",
      "\r\n",
      "    start = cuda.grid(1)\r\n",
      "    stride = cuda.gridsize(1)\r\n",
      "\r\n",
      "    for i in range(start, x.shape[0], stride):\r\n",
      "        bin_number = np.int32((x[i] + xmin)/bin_width)\r\n",
      "\r\n",
      "        if bin_number >= 0 or bin_number < histogram_out.shape[0]:\r\n",
      "            cuda.atomic.add(histogram_out, bin_number, 1)\r\n",
      "\r\n",
      "x = np.random.normal(size=50, loc=0, scale=1).astype(np.float32)\r\n",
      "xmin = np.float32(-4.0)\r\n",
      "xmax = np.float32(4.0)\r\n",
      "histogram_out = np.zeros(shape=10, dtype=np.int32)\r\n",
      "\r\n",
      "histogram[64, 64](x, xmin, xmax, histogram_out)\r\n",
      "\r\n",
      "print('input count:', x.shape[0])\r\n",
      "print('histogram:', histogram_out)\r\n",
      "print('count:', histogram_out.sum())\r\n"
     ]
    }
   ],
   "source": [
    "! cat debug/ex3.py"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 3,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "========= CUDA-MEMCHECK\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (49,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f0 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (48,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f4 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (47,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (46,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (45,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f0 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (44,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (43,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (42,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f4 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (41,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f0 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (40,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f4 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (39,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601e8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (38,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601e4 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (37,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601e8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (36,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601e8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (35,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601e8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (34,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (33,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (32,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (31,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f0 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (30,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (29,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (28,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (27,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (26,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601e8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (25,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601e8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (24,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601e8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (23,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f0 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (22,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (21,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (20,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (19,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (18,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (17,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f4 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (16,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f0 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (15,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601e8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (14,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (13,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f0 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (12,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (11,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f4 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (10,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601e8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (9,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (8,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f0 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (7,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f4 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (6,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f4 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (5,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f0 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (4,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (3,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (2,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f0 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (1,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f0 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000460 in cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (0,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f4 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Program hit CUDA_ERROR_LAUNCH_FAILED (error 719) due to \"unspecified launch failure\" on CUDA API call to cuMemcpyDtoH_v2. \n",
      "=========     Saved host backtrace up to driver entry point at error\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuMemcpyDtoH_v2 + 0x184) [0x158214]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========     Host Frame:[0x7fff5b7db780]\n",
      "=========\n",
      "Traceback (most recent call last):\n",
      "  File \"debug/ex3.py\", line 24, in <module>\n",
      "    histogram[64, 64](x, xmin, xmax, histogram_out)\n",
      "  File \"/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/site-packages/numba/cuda/compiler.py\", line 703, in __call__\n",
      "    cfg(*args)\n",
      "  File \"/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/site-packages/numba/cuda/compiler.py\", line 483, in __call__\n",
      "    sharedmem=self.sharedmem)\n",
      "  File \"/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/site-packages/numba/cuda/compiler.py\", line 585, in _kernel_call\n",
      "    wb()\n",
      "  File \"/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/site-packages/numba/cuda/compiler.py\", line 600, in <lambda>\n",
      "    retr.append(lambda: devary.copy_to_host(val, stream=stream))\n",
      "  File \"/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py\", line 198, in copy_to_host\n",
      "    _driver.device_to_host(hostary, self, self.alloc_size, stream=stream)\n",
      "  File \"/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/site-packages/numba/cuda/cudadrv/driver.py\", line 1484, in device_to_host\n",
      "    fn(host_pointer(dst), device_pointer(src), size, *varargs)\n",
      "  File \"/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/site-packages/numba/cuda/cudadrv/driver.py\", line 262, in safe_cuda_api_call\n",
      "    self._check_error(fname, retcode)\n",
      "  File \"/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/site-packages/numba/cuda/cudadrv/driver.py\", line 299, in _check_error\n",
      "    raise CudaAPIError(retcode, msg)\n",
      "numba.cuda.cudadrv.driver.CudaAPIError: [719] Call to cuMemcpyDtoH results in UNKNOWN_CUDA_ERROR\n"
     ]
    },
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "========= ERROR SUMMARY: 51 errors\r\n"
     ]
    }
   ],
   "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 0x00000460 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": 4,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "========= CUDA-MEMCHECK\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (49,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f4 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (48,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f4 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (47,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f0 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (46,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601e8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (45,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601e4 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (44,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (43,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (42,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f0 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (41,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (40,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f0 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (39,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (38,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (37,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601e8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (36,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f0 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (35,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f0 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (34,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (33,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601e8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (32,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f4 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (31,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (30,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601e8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (29,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (28,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (27,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f4 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (26,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601e4 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (25,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (24,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f4 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (23,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (22,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601e8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (21,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (20,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601e8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (19,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f4 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (18,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (17,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601e8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (16,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (15,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f0 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (14,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (13,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (12,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601e4 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (11,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601e8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (10,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (9,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601e8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (8,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (7,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (6,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (5,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f0 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (4,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (3,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (2,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601ec is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (1,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601f0 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Invalid __global__ write of size 4\n",
      "=========     at 0x00000ba8 in /Users/sseibert/continuum/conferences/gtc2017-numba/debug/ex3a.py:17:cudapy::__main__::histogram$241(Array<float, int=1, C, mutable, aligned>, float, float, Array<int, int=1, C, mutable, aligned>)\n",
      "=========     by thread (0,0,0) in block (0,0,0)\n",
      "=========     Address 0x700a601e8 is out of bounds\n",
      "=========     Saved host backtrace up to driver entry point at kernel launch time\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuLaunchKernel + 0x234) [0x15ca94]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========\n",
      "========= Program hit CUDA_ERROR_LAUNCH_FAILED (error 719) due to \"unspecified launch failure\" on CUDA API call to cuMemcpyDtoH_v2. \n",
      "=========     Saved host backtrace up to driver entry point at error\n",
      "=========     Host Frame:/Library/Frameworks/CUDA.framework/Versions/A/Libraries/libcuda_355.10.05.15_mercury.dylib (cuMemcpyDtoH_v2 + 0x184) [0x158214]\n",
      "=========     Host Frame:/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/lib-dynload/_ctypes.cpython-36m-darwin.so (ffi_call_unix64 + 0x4f) [0xd2b7]\n",
      "=========     Host Frame:[0x7fff50092d60]\n",
      "=========\n",
      "Traceback (most recent call last):\n",
      "  File \"debug/ex3a.py\", line 24, in <module>\n",
      "    histogram[64, 64](x, xmin, xmax, histogram_out)\n",
      "  File \"/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/site-packages/numba/cuda/compiler.py\", line 703, in __call__\n",
      "    cfg(*args)\n",
      "  File \"/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/site-packages/numba/cuda/compiler.py\", line 483, in __call__\n",
      "    sharedmem=self.sharedmem)\n",
      "  File \"/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/site-packages/numba/cuda/compiler.py\", line 560, in _kernel_call\n",
      "    driver.device_to_host(ctypes.addressof(excval), excmem, excsz)\n",
      "  File \"/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/site-packages/numba/cuda/cudadrv/driver.py\", line 1484, in device_to_host\n",
      "    fn(host_pointer(dst), device_pointer(src), size, *varargs)\n",
      "  File \"/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/site-packages/numba/cuda/cudadrv/driver.py\", line 262, in safe_cuda_api_call\n",
      "    self._check_error(fname, retcode)\n",
      "  File \"/Users/sseibert/anaconda/envs/gtc2017/lib/python3.6/site-packages/numba/cuda/cudadrv/driver.py\", line 299, in _check_error\n",
      "    raise CudaAPIError(retcode, msg)\n",
      "numba.cuda.cudadrv.driver.CudaAPIError: [719] Call to cuMemcpyDtoH results in UNKNOWN_CUDA_ERROR\n"
     ]
    },
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "========= ERROR SUMMARY: 51 errors\r\n"
     ]
    }
   ],
   "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": 7,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "    15\t\r\n",
      "    16\t        if bin_number >= 0 or bin_number < histogram_out.shape[0]:\r\n",
      "    17\t            cuda.atomic.add(histogram_out, bin_number, 1)\r\n",
      "    18\t\r\n",
      "    19\tx = np.random.normal(size=50, loc=0, scale=1).astype(np.float32)\r\n"
     ]
    }
   ],
   "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": {
    "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.6.1"
  }
 },
 "nbformat": 4,
 "nbformat_minor": 2
}