session-2/exercises/Typing Exercises with output.ipynb

{
 "cells": [
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "# Numba Typing Exercises\n",
    "\n",
    "This notebook provides some exposition of Numba's typing mechanisms, and how to deal with some of the issues you may encounter with typing. It covers:\n",
    "\n",
    "* How to display the typing of functions and understand the output,\n",
    "* Examination of different typings of the same function,\n",
    "* How to understand and fix typing errors,\n",
    "* Some CUDA-specific issues related to performance and occupancy.\n",
    "\n",
    "This notebook as published in the Git repository includes all the output from a previous run of the notebook - this is because some of the output (e.g. temporary variable names, register counts, etc.) may vary slightly with different versions of Numba, or different CUDA toolkits. It is suggested to clear all the output and work through the notebook, and keep the version with output provided for a reference in case there appears to be a discrepancy between the output of Numba and the description given in the text.\n",
    "\n",
    "We'll begin with importing some required packages. We use the `@njit` decorator for CPU-targeted examples, which is shorthand for `@jit(nopython=True)` - the nopython mode typing has more requirements than the object mode typing, which leads to better performance, so it is a better choice for learning about typing. The `@cuda.jit` decorator is used for the CUDA-targeted examples."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 1,
   "metadata": {},
   "outputs": [],
   "source": [
    "from numba import njit, cuda\n",
    "import numpy as np"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## Inspecting the typing\n",
    "\n",
    "Throughout this notebook we will use `inspect_types()` extensively to inspect the results of the typing algorithm. We'll start with a very simple example:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 2,
   "metadata": {},
   "outputs": [],
   "source": [
    "@njit\n",
    "def f(a, b):\n",
    "    return a + b"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Let's call this function with a pair of `float32`s to force a typing:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 3,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "3.0"
      ]
     },
     "execution_count": 3,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "f(np.float32(1), np.float32(2))"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Now we'll inspect the typing for this call:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 4,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "f (float32, float32)\n",
      "--------------------------------------------------------------------------------\n",
      "# File: <ipython-input-2-7136c8bb3c8c>\n",
      "# --- LINE 1 --- \n",
      "\n",
      "@njit\n",
      "\n",
      "# --- LINE 2 --- \n",
      "\n",
      "def f(a, b):\n",
      "\n",
      "    # --- LINE 3 --- \n",
      "    # label 0\n",
      "    #   a = arg(0, name=a)  :: float32\n",
      "    #   b = arg(1, name=b)  :: float32\n",
      "    #   $6binary_add.2 = a + b  :: float32\n",
      "    #   del b\n",
      "    #   del a\n",
      "    #   $8return_value.3 = cast(value=$6binary_add.2)  :: float32\n",
      "    #   del $6binary_add.2\n",
      "    #   return $8return_value.3\n",
      "\n",
      "    return a + b\n",
      "\n",
      "\n",
      "================================================================================\n"
     ]
    }
   ],
   "source": [
    "f.inspect_types()"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "The output of `inspect_types()` is a printout of the function's source code annotated with the Numba IR for each line, and the type of each IR node. Note that `del` nodes have no type, as they simply delete an existing variable.\n",
    "\n",
    "Types are separated from IR with a double colon. In one example from above:\n",
    "\n",
    "```\n",
    "$8return_value.3 = cast(value=$6binary_add.2)  :: float32\n",
    "```\n",
    "\n",
    "the type of `$return_value.3` (which is also the type returned by `cast(value=$6binary_add.2)` is `float32`.\n",
    "\n",
    "## An example with branching\n",
    "\n",
    "When a variable takes a value from multiple different control flow paths (i.e. *branches*), a unification is needed to determine a type that is suitable for representing the types across all the different control flow path. We can explore unification using a simple function with a branch in it:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 5,
   "metadata": {},
   "outputs": [],
   "source": [
    "@njit\n",
    "def select(a, b, c):\n",
    "    if c:\n",
    "        ret = a\n",
    "    else:\n",
    "        ret = b\n",
    "    return ret"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "We'll start by calling the function with `a` and `b` both as `float32` for a first example:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 6,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "1.0"
      ]
     },
     "execution_count": 6,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "select(np.float32(1), np.float32(2), True)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "If we inspect the typing we get:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 7,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "select (float32, float32, bool)\n",
      "--------------------------------------------------------------------------------\n",
      "# File: <ipython-input-5-edbb7033ccdc>\n",
      "# --- LINE 1 --- \n",
      "\n",
      "@njit\n",
      "\n",
      "# --- LINE 2 --- \n",
      "\n",
      "def select(a, b, c):\n",
      "\n",
      "    # --- LINE 3 --- \n",
      "    # label 0\n",
      "    #   a = arg(0, name=a)  :: float32\n",
      "    #   b = arg(1, name=b)  :: float32\n",
      "    #   c = arg(2, name=c)  :: bool\n",
      "    #   branch c, 6, 12\n",
      "\n",
      "    if c:\n",
      "\n",
      "        # --- LINE 4 --- \n",
      "        # label 6\n",
      "        #   del c\n",
      "        #   del b\n",
      "        #   ret = a  :: float32\n",
      "        #   del a\n",
      "        #   jump 16\n",
      "\n",
      "        ret = a\n",
      "\n",
      "    # --- LINE 5 --- \n",
      "\n",
      "    else:\n",
      "\n",
      "        # --- LINE 6 --- \n",
      "        # label 12\n",
      "        #   del c\n",
      "        #   del a\n",
      "        #   ret.1 = b  :: float32\n",
      "        #   del b\n",
      "\n",
      "        ret = b\n",
      "\n",
      "    # --- LINE 7 --- \n",
      "    #   jump 16\n",
      "    # label 16\n",
      "    #   ret.2 = phi(incoming_values=[Var(ret.1, <ipython-input-5-edbb7033ccdc>:6), Var(ret, <ipython-input-5-edbb7033ccdc>:4)], incoming_blocks=[12, 6])  :: float32\n",
      "    #   del ret.1\n",
      "    #   del ret\n",
      "    #   $18return_value.1 = cast(value=ret.2)  :: float32\n",
      "    #   del ret.2\n",
      "    #   return $18return_value.1\n",
      "\n",
      "    return ret\n",
      "\n",
      "\n",
      "================================================================================\n"
     ]
    }
   ],
   "source": [
    "select.inspect_types()"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "We see that where the value of a variable can come from two separate branches, there is a *phi node*: `ret.2 = phi(incoming_values=...)`. The `incoming_values` track the different sources of this variable - in this example, `ret` from the `if` side of the branch, and `ret.1` from the `else` side of the branch.\n",
    "\n",
    "The type of the phi node (`float32` in this case) is the type resulting from unification of the types of all the incoming values."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Another typing of the branching function\n",
    "\n",
    "If we call the function with a `float32` and a `float64`, we get another typing:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 8,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "1.0"
      ]
     },
     "execution_count": 8,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "select(np.float32(1), np.float64(2), True)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Now let's inspect types again. This time, there will be two sets of typings - one for the `(float32, float32, boolean)` call earlier, and another for the `(float32, float64, boolean)` call in the previous cell. If we call `inspect_types()` with no arguments, it will print out the typings for all sets of argument types that have been seen so far. In order to focus on just the case we are interested in, we can pass the `signature` keyword argument with a tuple of Numba types to get the typing for a specific set of argument types. Numba types are imported from `numba` - for a comprehensive list of them, see [Types and signatures](http://numba.pydata.org/numba-doc/latest/reference/types.html#types-and-signatures) in the Numba documentation."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 9,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "select (float32, float64, bool)\n",
      "--------------------------------------------------------------------------------\n",
      "# File: <ipython-input-5-edbb7033ccdc>\n",
      "# --- LINE 1 --- \n",
      "\n",
      "@njit\n",
      "\n",
      "# --- LINE 2 --- \n",
      "\n",
      "def select(a, b, c):\n",
      "\n",
      "    # --- LINE 3 --- \n",
      "    # label 0\n",
      "    #   a = arg(0, name=a)  :: float32\n",
      "    #   b = arg(1, name=b)  :: float64\n",
      "    #   c = arg(2, name=c)  :: bool\n",
      "    #   branch c, 6, 12\n",
      "\n",
      "    if c:\n",
      "\n",
      "        # --- LINE 4 --- \n",
      "        # label 6\n",
      "        #   del c\n",
      "        #   del b\n",
      "        #   ret = a  :: float32\n",
      "        #   del a\n",
      "        #   jump 16\n",
      "\n",
      "        ret = a\n",
      "\n",
      "    # --- LINE 5 --- \n",
      "\n",
      "    else:\n",
      "\n",
      "        # --- LINE 6 --- \n",
      "        # label 12\n",
      "        #   del c\n",
      "        #   del a\n",
      "        #   ret.1 = b  :: float64\n",
      "        #   del b\n",
      "\n",
      "        ret = b\n",
      "\n",
      "    # --- LINE 7 --- \n",
      "    #   jump 16\n",
      "    # label 16\n",
      "    #   ret.2 = phi(incoming_values=[Var(ret.1, <ipython-input-5-edbb7033ccdc>:6), Var(ret, <ipython-input-5-edbb7033ccdc>:4)], incoming_blocks=[12, 6])  :: float64\n",
      "    #   del ret.1\n",
      "    #   del ret\n",
      "    #   $18return_value.1 = cast(value=ret.2)  :: float64\n",
      "    #   del ret.2\n",
      "    #   return $18return_value.1\n",
      "\n",
      "    return ret\n",
      "\n",
      "\n",
      "================================================================================\n"
     ]
    }
   ],
   "source": [
    "from numba import float32, float64, boolean\n",
    "select.inspect_types(signature=(float32, float64, boolean))"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Here we see the types from each branch (`ret = a  :: float32` and `ret.1 = b  :: float64`) have unified to `float64` at the phi node.\n",
    "\n",
    "### Failing unification\n",
    "\n",
    "Sometimes unification can fail. If we try to choose between a tuple and a scalar:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 10,
   "metadata": {},
   "outputs": [
    {
     "ename": "TypingError",
     "evalue": "Failed in nopython mode pipeline (step: nopython frontend)\nCannot unify float64 and UniTuple(int64 x 2) for 'ret.2', defined at <ipython-input-5-edbb7033ccdc> (7)\n\nFile \"<ipython-input-5-edbb7033ccdc>\", line 7:\ndef select(a, b, c):\n    <source elided>\n        ret = b\n    return ret\n    ^\n\n[1] During: typing of assignment at <ipython-input-5-edbb7033ccdc> (7)\n\nFile \"<ipython-input-5-edbb7033ccdc>\", line 7:\ndef select(a, b, c):\n    <source elided>\n        ret = b\n    return ret\n    ^\n",
     "output_type": "error",
     "traceback": [
      "\u001b[0;31m---------------------------------------------------------------------------\u001b[0m",
      "\u001b[0;31mTypingError\u001b[0m                               Traceback (most recent call last)",
      "\u001b[0;32m<ipython-input-10-7125d7a99dcc>\u001b[0m in \u001b[0;36m<module>\u001b[0;34m\u001b[0m\n\u001b[0;32m----> 1\u001b[0;31m \u001b[0mselect\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;36m1\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;36m2\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;36m3.0\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;32mFalse\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m",
      "\u001b[0;32m~/numbadev/numba/numba/core/dispatcher.py\u001b[0m in \u001b[0;36m_compile_for_args\u001b[0;34m(self, *args, **kws)\u001b[0m\n\u001b[1;32m    399\u001b[0m                 \u001b[0me\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mpatch_message\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mmsg\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    400\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m--> 401\u001b[0;31m             \u001b[0merror_rewrite\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0me\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m'typing'\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m    402\u001b[0m         \u001b[0;32mexcept\u001b[0m \u001b[0merrors\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mUnsupportedError\u001b[0m \u001b[0;32mas\u001b[0m \u001b[0me\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    403\u001b[0m             \u001b[0;31m# Something unsupported is present in the user code, add help info\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;32m~/numbadev/numba/numba/core/dispatcher.py\u001b[0m in \u001b[0;36merror_rewrite\u001b[0;34m(e, issue_type)\u001b[0m\n\u001b[1;32m    342\u001b[0m                 \u001b[0;32mraise\u001b[0m \u001b[0me\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    343\u001b[0m             \u001b[0;32melse\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m--> 344\u001b[0;31m                 \u001b[0mreraise\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mtype\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0me\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0me\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;32mNone\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m    345\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    346\u001b[0m         \u001b[0margtypes\u001b[0m \u001b[0;34m=\u001b[0m \u001b[0;34m[\u001b[0m\u001b[0;34m]\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;32m~/numbadev/numba/numba/core/utils.py\u001b[0m in \u001b[0;36mreraise\u001b[0;34m(tp, value, tb)\u001b[0m\n\u001b[1;32m     78\u001b[0m         \u001b[0mvalue\u001b[0m \u001b[0;34m=\u001b[0m \u001b[0mtp\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m     79\u001b[0m     \u001b[0;32mif\u001b[0m \u001b[0mvalue\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0m__traceback__\u001b[0m \u001b[0;32mis\u001b[0m \u001b[0;32mnot\u001b[0m \u001b[0mtb\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m---> 80\u001b[0;31m         \u001b[0;32mraise\u001b[0m \u001b[0mvalue\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mwith_traceback\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mtb\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m     81\u001b[0m     \u001b[0;32mraise\u001b[0m \u001b[0mvalue\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m     82\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;31mTypingError\u001b[0m: Failed in nopython mode pipeline (step: nopython frontend)\nCannot unify float64 and UniTuple(int64 x 2) for 'ret.2', defined at <ipython-input-5-edbb7033ccdc> (7)\n\nFile \"<ipython-input-5-edbb7033ccdc>\", line 7:\ndef select(a, b, c):\n    <source elided>\n        ret = b\n    return ret\n    ^\n\n[1] During: typing of assignment at <ipython-input-5-edbb7033ccdc> (7)\n\nFile \"<ipython-input-5-edbb7033ccdc>\", line 7:\ndef select(a, b, c):\n    <source elided>\n        ret = b\n    return ret\n    ^\n"
     ]
    }
   ],
   "source": [
    "select((1, 2), 3.0, False)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "The typing fails at unification: `Cannot unify int64 and UniTuple(int64 x 2) for 'ret.2'`.\n",
    "\n",
    "When a typing error occurs, we can debug the propagation of type information by setting the environment variable `NUMBA_DEBUG_TYPEINFER` to `1`, or setting `numba.config.DEBUG_TYPEINFER` to `True`. It helps to also dump the Numba IR to understand the results of propagation better, so we should also set `numba.config.DUMP_IR` to `True` (or use the corresponding environment variable `NUMBA_DUMP_IR`). The debug output won't appear in the Jupyter notebook, but we can get the output by re-running this example as an external script:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 11,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "--------------------------------IR DUMP: select---------------------------------\n",
      "label 0:\n",
      "    a = arg(0, name=a)                       ['a']\n",
      "    b = arg(1, name=b)                       ['b']\n",
      "    c = arg(2, name=c)                       ['c']\n",
      "    branch c, 6, 12                          ['c']\n",
      "label 6:\n",
      "    ret = a                                  ['a', 'ret']\n",
      "    jump 16                                  []\n",
      "label 12:\n",
      "    ret = b                                  ['b', 'ret']\n",
      "    jump 16                                  []\n",
      "label 16:\n",
      "    $18return_value.1 = cast(value=ret)      ['$18return_value.1', 'ret']\n",
      "    return $18return_value.1                 ['$18return_value.1']\n",
      "\n",
      "--------------------------------IR DUMP: select---------------------------------\n",
      "label 0:\n",
      "    a = arg(0, name=a)                       ['a']\n",
      "    b = arg(1, name=b)                       ['b']\n",
      "    c = arg(2, name=c)                       ['c']\n",
      "    branch c, 6, 12                          ['c']\n",
      "label 6:\n",
      "    ret = a                                  ['a', 'ret']\n",
      "    jump 16                                  []\n",
      "label 12:\n",
      "    ret = b                                  ['b', 'ret']\n",
      "    jump 16                                  []\n",
      "label 16:\n",
      "    $18return_value.1 = cast(value=ret)      ['$18return_value.1', 'ret']\n",
      "    return $18return_value.1                 ['$18return_value.1']\n",
      "\n",
      "-----------------------------------propagate------------------------------------\n",
      "---- type variables ----\n",
      "[$18return_value.1 := float64,\n",
      " a := UniTuple(int64 x 2),\n",
      " arg.a := UniTuple(int64 x 2),\n",
      " arg.b := float64,\n",
      " arg.c := bool,\n",
      " b := float64,\n",
      " c := bool,\n",
      " ret := UniTuple(int64 x 2),\n",
      " ret.1 := float64,\n",
      " ret.2 := float64]\n",
      "-----------------------------------propagate------------------------------------\n",
      "---- type variables ----\n",
      "[$18return_value.1 := float64,\n",
      " a := UniTuple(int64 x 2),\n",
      " arg.a := UniTuple(int64 x 2),\n",
      " arg.b := float64,\n",
      " arg.c := bool,\n",
      " b := float64,\n",
      " c := bool,\n",
      " ret := UniTuple(int64 x 2),\n",
      " ret.1 := float64,\n",
      " ret.2 := float64]\n"
     ]
    },
    {
     "name": "stderr",
     "output_type": "stream",
     "text": [
      "Traceback (most recent call last):\n",
      "  File \"<stdin>\", line 15, in <module>\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/dispatcher.py\", line 401, in _compile_for_args\n",
      "    error_rewrite(e, 'typing')\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/dispatcher.py\", line 344, in error_rewrite\n",
      "    reraise(type(e), e, None)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/utils.py\", line 80, in reraise\n",
      "    raise value.with_traceback(tb)\n",
      "numba.core.errors.TypingError: Failed in nopython mode pipeline (step: nopython frontend)\n",
      "Cannot unify float64 and UniTuple(int64 x 2) for 'ret.2', defined at <stdin> (11)\n",
      "\n",
      "File \"<stdin>\", line 11:\n",
      "<source missing, REPL/exec in use?>\n",
      "\n",
      "[1] During: typing of assignment at <stdin> (11)\n",
      "\n",
      "File \"<stdin>\", line 11:\n",
      "<source missing, REPL/exec in use?>\n",
      "\n"
     ]
    },
    {
     "ename": "CalledProcessError",
     "evalue": "Command 'b'from numba import njit\\nfrom numba import config\\n\\n\\n@njit\\ndef select(a, b, c):\\n    if c:\\n        ret = a\\n    else:\\n        ret = b\\n    return ret\\n\\nconfig.DEBUG_TYPEINFER = True\\nconfig.DUMP_IR=True\\nselect((1, 2), 3.0, False)\\n'' returned non-zero exit status 1.",
     "output_type": "error",
     "traceback": [
      "\u001b[0;31m---------------------------------------------------------------------------\u001b[0m",
      "\u001b[0;31mCalledProcessError\u001b[0m                        Traceback (most recent call last)",
      "\u001b[0;32m<ipython-input-11-c57ef22c677f>\u001b[0m in \u001b[0;36m<module>\u001b[0;34m\u001b[0m\n\u001b[0;32m----> 1\u001b[0;31m \u001b[0mget_ipython\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mrun_cell_magic\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m'script'\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m'python'\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m'from numba import njit\\nfrom numba import config\\n\\n\\n@njit\\ndef select(a, b, c):\\n    if c:\\n        ret = a\\n    else:\\n        ret = b\\n    return ret\\n\\nconfig.DEBUG_TYPEINFER = True\\nconfig.DUMP_IR=True\\nselect((1, 2), 3.0, False)\\n'\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m",
      "\u001b[0;32m~/miniconda3/envs/numba/lib/python3.8/site-packages/IPython/core/interactiveshell.py\u001b[0m in \u001b[0;36mrun_cell_magic\u001b[0;34m(self, magic_name, line, cell)\u001b[0m\n\u001b[1;32m   2360\u001b[0m             \u001b[0;32mwith\u001b[0m \u001b[0mself\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mbuiltin_trap\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m   2361\u001b[0m                 \u001b[0margs\u001b[0m \u001b[0;34m=\u001b[0m \u001b[0;34m(\u001b[0m\u001b[0mmagic_arg_s\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mcell\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m-> 2362\u001b[0;31m                 \u001b[0mresult\u001b[0m \u001b[0;34m=\u001b[0m \u001b[0mfn\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m*\u001b[0m\u001b[0margs\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m**\u001b[0m\u001b[0mkwargs\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m   2363\u001b[0m             \u001b[0;32mreturn\u001b[0m \u001b[0mresult\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m   2364\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;32m<decorator-gen-110>\u001b[0m in \u001b[0;36mshebang\u001b[0;34m(self, line, cell)\u001b[0m\n",
      "\u001b[0;32m~/miniconda3/envs/numba/lib/python3.8/site-packages/IPython/core/magic.py\u001b[0m in \u001b[0;36m<lambda>\u001b[0;34m(f, *a, **k)\u001b[0m\n\u001b[1;32m    185\u001b[0m     \u001b[0;31m# but it's overkill for just that one bit of state.\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    186\u001b[0m     \u001b[0;32mdef\u001b[0m \u001b[0mmagic_deco\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0marg\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m--> 187\u001b[0;31m         \u001b[0mcall\u001b[0m \u001b[0;34m=\u001b[0m \u001b[0;32mlambda\u001b[0m \u001b[0mf\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m*\u001b[0m\u001b[0ma\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m**\u001b[0m\u001b[0mk\u001b[0m\u001b[0;34m:\u001b[0m \u001b[0mf\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m*\u001b[0m\u001b[0ma\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m**\u001b[0m\u001b[0mk\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m    188\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    189\u001b[0m         \u001b[0;32mif\u001b[0m \u001b[0mcallable\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0marg\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;32m~/miniconda3/envs/numba/lib/python3.8/site-packages/IPython/core/magics/script.py\u001b[0m in \u001b[0;36mshebang\u001b[0;34m(self, line, cell)\u001b[0m\n\u001b[1;32m    243\u001b[0m             \u001b[0msys\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mstderr\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mflush\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    244\u001b[0m         \u001b[0;32mif\u001b[0m \u001b[0margs\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mraise_error\u001b[0m \u001b[0;32mand\u001b[0m \u001b[0mp\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mreturncode\u001b[0m\u001b[0;34m!=\u001b[0m\u001b[0;36m0\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m--> 245\u001b[0;31m             \u001b[0;32mraise\u001b[0m \u001b[0mCalledProcessError\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mp\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mreturncode\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mcell\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0moutput\u001b[0m\u001b[0;34m=\u001b[0m\u001b[0mout\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mstderr\u001b[0m\u001b[0;34m=\u001b[0m\u001b[0merr\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m    246\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    247\u001b[0m     \u001b[0;32mdef\u001b[0m \u001b[0m_run_script\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mself\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mp\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mcell\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mto_close\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;31mCalledProcessError\u001b[0m: Command 'b'from numba import njit\\nfrom numba import config\\n\\n\\n@njit\\ndef select(a, b, c):\\n    if c:\\n        ret = a\\n    else:\\n        ret = b\\n    return ret\\n\\nconfig.DEBUG_TYPEINFER = True\\nconfig.DUMP_IR=True\\nselect((1, 2), 3.0, False)\\n'' returned non-zero exit status 1."
     ]
    }
   ],
   "source": [
    "%%script python\n",
    "from numba import njit\n",
    "from numba import config\n",
    "\n",
    "\n",
    "@njit\n",
    "def select(a, b, c):\n",
    "    if c:\n",
    "        ret = a\n",
    "    else:\n",
    "        ret = b\n",
    "    return ret\n",
    "\n",
    "config.DEBUG_TYPEINFER = True\n",
    "config.DUMP_IR=True\n",
    "select((1, 2), 3.0, False)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Numba has printed a dump of the type of all variables after each propagation step. The type inference happens on the IR in [Static Single Assignment (SSA) form](https://en.wikipedia.org/wiki/Static_single_assignment_form) so the names of variables after propagation carry a \"version number\" - e.g. `ret`, `ret.1`, `ret.2`, etc - however, the IR dump does not presently print the version numbers of each variable, so it can be a little tricky to work out which variable each versioned variable refers to.\n",
    "\n",
    "The different versions of the variable make up the set that is being unified, so we can see that the variable `ret` has a set of `{UniTuple(int64 x 2), float64, float64}` from its versions `ret`, `ret.1` and `ret.2`.\n",
    "\n",
    "A general strategy for debugging typing issues is to examine the changes in the types of variables at each propagate step, to determine how a typing error is occurring.\n",
    "\n",
    "### Exercises\n",
    "\n",
    "Execute the code in the following cell, and try to locate the typing of `x` in the output. Try to understand the message accompanying the `TypingError` (which begins with `Invalid use of Function(...`). You may find it easier to run this example on the terminal to avoid a lot of scrolling through a frame in the IPython notebook."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 12,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "----------------------------IR DUMP: array_vs_scalar----------------------------\n",
      "label 0:\n",
      "    $2load_global.0 = global(np: <module 'numpy' from '/home/gmarkall/miniconda3/envs/numba/lib/python3.8/site-packages/numpy/__init__.py'>) ['$2load_global.0']\n",
      "    $4load_method.1 = getattr(value=$2load_global.0, attr=zeros) ['$2load_global.0', '$4load_method.1']\n",
      "    $const6.2 = const(int, 20)               ['$const6.2']\n",
      "    $8call_method.3 = call $4load_method.1($const6.2, func=$4load_method.1, args=[Var($const6.2, <stdin>:9)], kws=(), vararg=None) ['$4load_method.1', '$8call_method.3', '$const6.2']\n",
      "    x = $8call_method.3                      ['$8call_method.3', 'x']\n",
      "    $const12.4 = const(int, 10)              ['$const12.4']\n",
      "    $const16.6 = const(int, 0)               ['$const16.6']\n",
      "    x[$const16.6] = $const12.4               ['$const12.4', '$const16.6', 'x']\n",
      "    $const20.7 = const(int, 20)              ['$const20.7']\n",
      "    $const_0 = const(int, 0)                 ['$const_0']\n",
      "    $const_1 = const(int, 1)                 ['$const_1']\n",
      "    $const24.9 = build_tuple(items=[Var($const_0, <stdin>:11), Var($const_1, <stdin>:11)]) ['$const24.9', '$const_0', '$const_1']\n",
      "    x[$const24.9] = $const20.7               ['$const20.7', '$const24.9', 'x']\n",
      "    $const28.10 = const(NoneType, None)      ['$const28.10']\n",
      "    $30return_value.11 = cast(value=$const28.10) ['$30return_value.11', '$const28.10']\n",
      "    return $30return_value.11                ['$30return_value.11']\n",
      "\n",
      "______________________________________________________________________\n",
      "REWRITING (RewriteConstSetitems):\n",
      "    $2load_global.0 = global(np: <module 'numpy' from '/home/gmarkall/miniconda3/envs/numba/lib/python3.8/site-packages/numpy/__init__.py'>) ['$2load_global.0']\n",
      "    $4load_method.1 = getattr(value=$2load_global.0, attr=zeros) ['$2load_global.0', '$4load_method.1']\n",
      "    $const6.2 = const(int, 20)               ['$const6.2']\n",
      "    $8call_method.3 = call $4load_method.1($const6.2, func=$4load_method.1, args=[Var($const6.2, <stdin>:9)], kws=(), vararg=None) ['$4load_method.1', '$8call_method.3', '$const6.2']\n",
      "    x = $8call_method.3                      ['$8call_method.3', 'x']\n",
      "    $const12.4 = const(int, 10)              ['$const12.4']\n",
      "    $const16.6 = const(int, 0)               ['$const16.6']\n",
      "    x[$const16.6] = $const12.4               ['$const12.4', '$const16.6', 'x']\n",
      "    $const20.7 = const(int, 20)              ['$const20.7']\n",
      "    $const_0 = const(int, 0)                 ['$const_0']\n",
      "    $const_1 = const(int, 1)                 ['$const_1']\n",
      "    $const24.9 = build_tuple(items=[Var($const_0, <stdin>:11), Var($const_1, <stdin>:11)]) ['$const24.9', '$const_0', '$const_1']\n",
      "    x[$const24.9] = $const20.7               ['$const20.7', '$const24.9', 'x']\n",
      "    $const28.10 = const(NoneType, None)      ['$const28.10']\n",
      "    $30return_value.11 = cast(value=$const28.10) ['$30return_value.11', '$const28.10']\n",
      "    return $30return_value.11                ['$30return_value.11']\n",
      "____________________________________________________________\n",
      "    $2load_global.0 = global(np: <module 'numpy' from '/home/gmarkall/miniconda3/envs/numba/lib/python3.8/site-packages/numpy/__init__.py'>) ['$2load_global.0']\n",
      "    $4load_method.1 = getattr(value=$2load_global.0, attr=zeros) ['$2load_global.0', '$4load_method.1']\n",
      "    $const6.2 = const(int, 20)               ['$const6.2']\n",
      "    $8call_method.3 = call $4load_method.1($const6.2, func=$4load_method.1, args=[Var($const6.2, <stdin>:9)], kws=(), vararg=None) ['$4load_method.1', '$8call_method.3', '$const6.2']\n",
      "    x = $8call_method.3                      ['$8call_method.3', 'x']\n",
      "    $const12.4 = const(int, 10)              ['$const12.4']\n",
      "    $const16.6 = const(int, 0)               ['$const16.6']\n",
      "    x[0] = $const12.4                        ['$const12.4', '$const16.6', 'x']\n",
      "    $const20.7 = const(int, 20)              ['$const20.7']\n",
      "    $const_0 = const(int, 0)                 ['$const_0']\n",
      "    $const_1 = const(int, 1)                 ['$const_1']\n",
      "    $const24.9 = build_tuple(items=[Var($const_0, <stdin>:11), Var($const_1, <stdin>:11)]) ['$const24.9', '$const_0', '$const_1']\n",
      "    x[(0, 1)] = $const20.7                   ['$const20.7', '$const24.9', 'x']\n",
      "    $const28.10 = const(NoneType, None)      ['$const28.10']\n",
      "    $30return_value.11 = cast(value=$const28.10) ['$30return_value.11', '$const28.10']\n",
      "    return $30return_value.11                ['$30return_value.11']\n",
      "______________________________________________________________________\n",
      "----------------------------IR DUMP: array_vs_scalar----------------------------\n",
      "label 0:\n",
      "    $2load_global.0 = global(np: <module 'numpy' from '/home/gmarkall/miniconda3/envs/numba/lib/python3.8/site-packages/numpy/__init__.py'>) ['$2load_global.0']\n",
      "    $4load_method.1 = getattr(value=$2load_global.0, attr=zeros) ['$2load_global.0', '$4load_method.1']\n",
      "    $const6.2 = const(int, 20)               ['$const6.2']\n",
      "    $8call_method.3 = call $4load_method.1($const6.2, func=$4load_method.1, args=[Var($const6.2, <stdin>:9)], kws=(), vararg=None) ['$4load_method.1', '$8call_method.3', '$const6.2']\n",
      "    x = $8call_method.3                      ['$8call_method.3', 'x']\n",
      "    $const12.4 = const(int, 10)              ['$const12.4']\n",
      "    $const16.6 = const(int, 0)               ['$const16.6']\n",
      "    x[0] = $const12.4                        ['$const12.4', '$const16.6', 'x']\n",
      "    $const20.7 = const(int, 20)              ['$const20.7']\n",
      "    $const_0 = const(int, 0)                 ['$const_0']\n",
      "    $const_1 = const(int, 1)                 ['$const_1']\n",
      "    $const24.9 = build_tuple(items=[Var($const_0, <stdin>:11), Var($const_1, <stdin>:11)]) ['$const24.9', '$const_0', '$const_1']\n",
      "    x[(0, 1)] = $const20.7                   ['$const20.7', '$const24.9', 'x']\n",
      "    $const28.10 = const(NoneType, None)      ['$const28.10']\n",
      "    $30return_value.11 = cast(value=$const28.10) ['$30return_value.11', '$const28.10']\n",
      "    return $30return_value.11                ['$30return_value.11']\n",
      "\n",
      "-----------------------------------propagate------------------------------------\n",
      "---- type variables ----\n",
      "[$2load_global.0 := Module(<module 'numpy' from '/home/gmarkall/miniconda3/envs/numba/lib/python3.8/site-packages/numpy/__init__.py'>),\n",
      " $30return_value.11 := none,\n",
      " $4load_method.1 := Function(<built-in function zeros>),\n",
      " $8call_method.3 := array(float64, 1d, C),\n",
      " $const12.4 := Literal[int](10),\n",
      " $const16.6 := Literal[int](0),\n",
      " $const20.7 := Literal[int](20),\n",
      " $const24.9 := Tuple(Literal[int](0), Literal[int](1)),\n",
      " $const28.10 := none,\n",
      " $const6.2 := Literal[int](20),\n",
      " $const_0 := Literal[int](0),\n",
      " $const_1 := Literal[int](1),\n",
      " x := array(float64, 1d, C)]\n",
      "-----------------------------------propagate------------------------------------\n",
      "---- type variables ----\n",
      "[$2load_global.0 := Module(<module 'numpy' from '/home/gmarkall/miniconda3/envs/numba/lib/python3.8/site-packages/numpy/__init__.py'>),\n",
      " $30return_value.11 := none,\n",
      " $4load_method.1 := Function(<built-in function zeros>),\n",
      " $8call_method.3 := array(float64, 1d, C),\n",
      " $const12.4 := Literal[int](10),\n",
      " $const16.6 := Literal[int](0),\n",
      " $const20.7 := Literal[int](20),\n",
      " $const24.9 := Tuple(Literal[int](0), Literal[int](1)),\n",
      " $const28.10 := none,\n",
      " $const6.2 := Literal[int](20),\n",
      " $const_0 := Literal[int](0),\n",
      " $const_1 := Literal[int](1),\n",
      " x := array(float64, 1d, C)]\n"
     ]
    },
    {
     "name": "stderr",
     "output_type": "stream",
     "text": [
      "Traceback (most recent call last):\n",
      "  File \"<stdin>\", line 13, in <module>\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/dispatcher.py\", line 401, in _compile_for_args\n",
      "    error_rewrite(e, 'typing')\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/dispatcher.py\", line 344, in error_rewrite\n",
      "    reraise(type(e), e, None)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/utils.py\", line 80, in reraise\n",
      "    raise value.with_traceback(tb)\n",
      "numba.core.errors.TypingError: Failed in nopython mode pipeline (step: nopython frontend)\n",
      "Invalid use of Function(<built-in function setitem>) with argument(s) of type(s): (array(float64, 1d, C), Tuple(Literal[int](0), Literal[int](1)), Literal[int](20))\n",
      " * parameterized\n",
      "In definition 0:\n",
      "    All templates rejected with literals.\n",
      "In definition 1:\n",
      "    All templates rejected without literals.\n",
      "In definition 2:\n",
      "    All templates rejected with literals.\n",
      "In definition 3:\n",
      "    All templates rejected without literals.\n",
      "In definition 4:\n",
      "    All templates rejected with literals.\n",
      "In definition 5:\n",
      "    All templates rejected without literals.\n",
      "In definition 6:\n",
      "    All templates rejected with literals.\n",
      "In definition 7:\n",
      "    All templates rejected without literals.\n",
      "In definition 8:\n",
      "    TypeError: cannot index array(float64, 1d, C) with 2 indices: Tuple(Literal[int](0), Literal[int](1))\n",
      "    raised from /home/gmarkall/numbadev/numba/numba/core/typing/arraydecl.py:84\n",
      "In definition 9:\n",
      "    TypeError: cannot index array(float64, 1d, C) with 2 indices: UniTuple(int64 x 2)\n",
      "    raised from /home/gmarkall/numbadev/numba/numba/core/typing/arraydecl.py:84\n",
      "This error is usually caused by passing an argument of a type that is unsupported by the named function.\n",
      "[1] During: typing of staticsetitem at <stdin> (11)\n",
      "\n",
      "File \"<stdin>\", line 11:\n",
      "<source missing, REPL/exec in use?>\n",
      "\n"
     ]
    },
    {
     "ename": "CalledProcessError",
     "evalue": "Command 'b'from numba import njit\\nfrom numba import config\\nimport numpy as np\\nconfig.DEBUG_TYPEINFER = True\\nconfig.DUMP_IR = True\\n\\n@njit\\ndef array_vs_scalar():\\n    x = np.zeros(20)\\n    x[0] = 10\\n    x[0, 1] = 20\\n\\narray_vs_scalar()\\n'' returned non-zero exit status 1.",
     "output_type": "error",
     "traceback": [
      "\u001b[0;31m---------------------------------------------------------------------------\u001b[0m",
      "\u001b[0;31mCalledProcessError\u001b[0m                        Traceback (most recent call last)",
      "\u001b[0;32m<ipython-input-12-8d9539e19d8e>\u001b[0m in \u001b[0;36m<module>\u001b[0;34m\u001b[0m\n\u001b[0;32m----> 1\u001b[0;31m \u001b[0mget_ipython\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mrun_cell_magic\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m'script'\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m'python'\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m'from numba import njit\\nfrom numba import config\\nimport numpy as np\\nconfig.DEBUG_TYPEINFER = True\\nconfig.DUMP_IR = True\\n\\n@njit\\ndef array_vs_scalar():\\n    x = np.zeros(20)\\n    x[0] = 10\\n    x[0, 1] = 20\\n\\narray_vs_scalar()\\n'\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m",
      "\u001b[0;32m~/miniconda3/envs/numba/lib/python3.8/site-packages/IPython/core/interactiveshell.py\u001b[0m in \u001b[0;36mrun_cell_magic\u001b[0;34m(self, magic_name, line, cell)\u001b[0m\n\u001b[1;32m   2360\u001b[0m             \u001b[0;32mwith\u001b[0m \u001b[0mself\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mbuiltin_trap\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m   2361\u001b[0m                 \u001b[0margs\u001b[0m \u001b[0;34m=\u001b[0m \u001b[0;34m(\u001b[0m\u001b[0mmagic_arg_s\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mcell\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m-> 2362\u001b[0;31m                 \u001b[0mresult\u001b[0m \u001b[0;34m=\u001b[0m \u001b[0mfn\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m*\u001b[0m\u001b[0margs\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m**\u001b[0m\u001b[0mkwargs\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m   2363\u001b[0m             \u001b[0;32mreturn\u001b[0m \u001b[0mresult\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m   2364\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;32m<decorator-gen-110>\u001b[0m in \u001b[0;36mshebang\u001b[0;34m(self, line, cell)\u001b[0m\n",
      "\u001b[0;32m~/miniconda3/envs/numba/lib/python3.8/site-packages/IPython/core/magic.py\u001b[0m in \u001b[0;36m<lambda>\u001b[0;34m(f, *a, **k)\u001b[0m\n\u001b[1;32m    185\u001b[0m     \u001b[0;31m# but it's overkill for just that one bit of state.\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    186\u001b[0m     \u001b[0;32mdef\u001b[0m \u001b[0mmagic_deco\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0marg\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m--> 187\u001b[0;31m         \u001b[0mcall\u001b[0m \u001b[0;34m=\u001b[0m \u001b[0;32mlambda\u001b[0m \u001b[0mf\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m*\u001b[0m\u001b[0ma\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m**\u001b[0m\u001b[0mk\u001b[0m\u001b[0;34m:\u001b[0m \u001b[0mf\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m*\u001b[0m\u001b[0ma\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m**\u001b[0m\u001b[0mk\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m    188\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    189\u001b[0m         \u001b[0;32mif\u001b[0m \u001b[0mcallable\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0marg\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;32m~/miniconda3/envs/numba/lib/python3.8/site-packages/IPython/core/magics/script.py\u001b[0m in \u001b[0;36mshebang\u001b[0;34m(self, line, cell)\u001b[0m\n\u001b[1;32m    243\u001b[0m             \u001b[0msys\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mstderr\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mflush\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    244\u001b[0m         \u001b[0;32mif\u001b[0m \u001b[0margs\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mraise_error\u001b[0m \u001b[0;32mand\u001b[0m \u001b[0mp\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mreturncode\u001b[0m\u001b[0;34m!=\u001b[0m\u001b[0;36m0\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m--> 245\u001b[0;31m             \u001b[0;32mraise\u001b[0m \u001b[0mCalledProcessError\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mp\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mreturncode\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mcell\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0moutput\u001b[0m\u001b[0;34m=\u001b[0m\u001b[0mout\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mstderr\u001b[0m\u001b[0;34m=\u001b[0m\u001b[0merr\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m    246\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    247\u001b[0m     \u001b[0;32mdef\u001b[0m \u001b[0m_run_script\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mself\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mp\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mcell\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mto_close\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;31mCalledProcessError\u001b[0m: Command 'b'from numba import njit\\nfrom numba import config\\nimport numpy as np\\nconfig.DEBUG_TYPEINFER = True\\nconfig.DUMP_IR = True\\n\\n@njit\\ndef array_vs_scalar():\\n    x = np.zeros(20)\\n    x[0] = 10\\n    x[0, 1] = 20\\n\\narray_vs_scalar()\\n'' returned non-zero exit status 1."
     ]
    }
   ],
   "source": [
    "%%script python\n",
    "from numba import njit\n",
    "from numba import config\n",
    "import numpy as np\n",
    "config.DEBUG_TYPEINFER = True\n",
    "config.DUMP_IR = True\n",
    "\n",
    "@njit\n",
    "def array_vs_scalar():\n",
    "    x = np.zeros(20)\n",
    "    x[0] = 10\n",
    "    x[0, 1] = 20\n",
    "\n",
    "array_vs_scalar()"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "This is an example of a function unsupported on the CUDA target. Numba tries to implement this function using the `array_sum_impl` internal function, which you will see in the output. Try to determine which function is unsupported (in the message beginning with `Use of unsupported NumPy function...`) and locate the call to it in the IR for `array_sum_impl`."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 13,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "------------------------------IR DUMP: sum_reduce-------------------------------\n",
      "label 0:\n",
      "    x = arg(0, name=x)                       ['x']\n",
      "    $4load_method.1 = getattr(value=x, attr=sum) ['$4load_method.1', 'x']\n",
      "    $6call_method.2 = call $4load_method.1(func=$4load_method.1, args=[], kws=(), vararg=None) ['$4load_method.1', '$6call_method.2']\n",
      "    $const10.4 = const(int, 0)               ['$const10.4']\n",
      "    x[$const10.4] = $6call_method.2          ['$6call_method.2', '$const10.4', 'x']\n",
      "    $const14.5 = const(NoneType, None)       ['$const14.5']\n",
      "    $16return_value.6 = cast(value=$const14.5) ['$16return_value.6', '$const14.5']\n",
      "    return $16return_value.6                 ['$16return_value.6']\n",
      "\n",
      "______________________________________________________________________\n",
      "REWRITING (RewriteConstSetitems):\n",
      "    x = arg(0, name=x)                       ['x']\n",
      "    $4load_method.1 = getattr(value=x, attr=sum) ['$4load_method.1', 'x']\n",
      "    $6call_method.2 = call $4load_method.1(func=$4load_method.1, args=[], kws=(), vararg=None) ['$4load_method.1', '$6call_method.2']\n",
      "    $const10.4 = const(int, 0)               ['$const10.4']\n",
      "    x[$const10.4] = $6call_method.2          ['$6call_method.2', '$const10.4', 'x']\n",
      "    $const14.5 = const(NoneType, None)       ['$const14.5']\n",
      "    $16return_value.6 = cast(value=$const14.5) ['$16return_value.6', '$const14.5']\n",
      "    return $16return_value.6                 ['$16return_value.6']\n",
      "____________________________________________________________\n",
      "    x = arg(0, name=x)                       ['x']\n",
      "    $4load_method.1 = getattr(value=x, attr=sum) ['$4load_method.1', 'x']\n",
      "    $6call_method.2 = call $4load_method.1(func=$4load_method.1, args=[], kws=(), vararg=None) ['$4load_method.1', '$6call_method.2']\n",
      "    $const10.4 = const(int, 0)               ['$const10.4']\n",
      "    x[0] = $6call_method.2                   ['$6call_method.2', '$const10.4', 'x']\n",
      "    $const14.5 = const(NoneType, None)       ['$const14.5']\n",
      "    $16return_value.6 = cast(value=$const14.5) ['$16return_value.6', '$const14.5']\n",
      "    return $16return_value.6                 ['$16return_value.6']\n",
      "______________________________________________________________________\n",
      "------------------------------IR DUMP: sum_reduce-------------------------------\n",
      "label 0:\n",
      "    x = arg(0, name=x)                       ['x']\n",
      "    $4load_method.1 = getattr(value=x, attr=sum) ['$4load_method.1', 'x']\n",
      "    $6call_method.2 = call $4load_method.1(func=$4load_method.1, args=[], kws=(), vararg=None) ['$4load_method.1', '$6call_method.2']\n",
      "    $const10.4 = const(int, 0)               ['$const10.4']\n",
      "    x[0] = $6call_method.2                   ['$6call_method.2', '$const10.4', 'x']\n",
      "    $const14.5 = const(NoneType, None)       ['$const14.5']\n",
      "    $16return_value.6 = cast(value=$const14.5) ['$16return_value.6', '$const14.5']\n",
      "    return $16return_value.6                 ['$16return_value.6']\n",
      "\n",
      "-----------------------------------propagate------------------------------------\n",
      "---- type variables ----\n",
      "[$16return_value.6 := none,\n",
      " $4load_method.1 := BoundFunction(array.sum for array(float64, 1d, C)),\n",
      " $6call_method.2 := float64,\n",
      " $const10.4 := Literal[int](0),\n",
      " $const14.5 := none,\n",
      " arg.x := array(float64, 1d, C),\n",
      " x := array(float64, 1d, C)]\n",
      "-----------------------------------propagate------------------------------------\n",
      "---- type variables ----\n",
      "[$16return_value.6 := none,\n",
      " $4load_method.1 := BoundFunction(array.sum for array(float64, 1d, C)),\n",
      " $6call_method.2 := float64,\n",
      " $const10.4 := Literal[int](0),\n",
      " $const14.5 := none,\n",
      " arg.x := array(float64, 1d, C),\n",
      " x := array(float64, 1d, C)]\n",
      "---------------------------------Variable types---------------------------------\n",
      "{'$16return_value.6': none,\n",
      " '$4load_method.1': BoundFunction(array.sum for array(float64, 1d, C)),\n",
      " '$6call_method.2': float64,\n",
      " '$const10.4': Literal[int](0),\n",
      " '$const14.5': none,\n",
      " 'arg.x': array(float64, 1d, C),\n",
      " 'x': array(float64, 1d, C)}\n",
      "----------------------------------Return type-----------------------------------\n",
      "none\n",
      "-----------------------------------Call types-----------------------------------\n",
      "{call $4load_method.1(func=$4load_method.1, args=[], kws=(), vararg=None): () -> float64,\n",
      " x[0] = $6call_method.2: (array(float64, 1d, C), int64, float64) -> none}\n",
      "-------------------IR DUMP: array_sum.<locals>.array_sum_impl-------------------\n",
      "label 0:\n",
      "    arr = arg(0, name=arr)                   ['arr']\n",
      "    $2load_deref.0 = freevar(zero: 0.0)      ['$2load_deref.0']\n",
      "    c = $2load_deref.0                       ['$2load_deref.0', 'c']\n",
      "    $6load_global.1 = global(np: <module 'numpy' from '/home/gmarkall/miniconda3/envs/numba/lib/python3.8/site-packages/numpy/__init__.py'>) ['$6load_global.1']\n",
      "    $8load_method.2 = getattr(value=$6load_global.1, attr=nditer) ['$6load_global.1', '$8load_method.2']\n",
      "    $12call_method.4 = call $8load_method.2(arr, func=$8load_method.2, args=[Var(arr, arraymath.py:166)], kws=(), vararg=None) ['$12call_method.4', '$8load_method.2', 'arr']\n",
      "    $14get_iter.5 = getiter(value=$12call_method.4) ['$12call_method.4', '$14get_iter.5']\n",
      "    $phi16.0 = $14get_iter.5                 ['$14get_iter.5', '$phi16.0']\n",
      "    jump 16                                  []\n",
      "label 16:\n",
      "    $16for_iter.1 = iternext(value=$phi16.0) ['$16for_iter.1', '$phi16.0']\n",
      "    $16for_iter.2 = pair_first(value=$16for_iter.1) ['$16for_iter.1', '$16for_iter.2']\n",
      "    $16for_iter.3 = pair_second(value=$16for_iter.1) ['$16for_iter.1', '$16for_iter.3']\n",
      "    $phi18.1 = $16for_iter.2                 ['$16for_iter.2', '$phi18.1']\n",
      "    branch $16for_iter.3, 18, 34             ['$16for_iter.3']\n",
      "label 18:\n",
      "    v = $phi18.1                             ['$phi18.1', 'v']\n",
      "    $24load_method.4 = getattr(value=v, attr=item) ['$24load_method.4', 'v']\n",
      "    $26call_method.5 = call $24load_method.4(func=$24load_method.4, args=[], kws=(), vararg=None) ['$24load_method.4', '$26call_method.5']\n",
      "    $28inplace_add.6 = inplace_binop(fn=<built-in function iadd>, immutable_fn=<built-in function add>, lhs=c, rhs=$26call_method.5, static_lhs=Undefined, static_rhs=Undefined) ['$26call_method.5', '$28inplace_add.6', 'c']\n",
      "    c = $28inplace_add.6                     ['$28inplace_add.6', 'c']\n",
      "    jump 16                                  []\n",
      "label 34:\n",
      "    $36return_value.1 = cast(value=c)        ['$36return_value.1', 'c']\n",
      "    return $36return_value.1                 ['$36return_value.1']\n",
      "\n",
      "-------------------IR DUMP: array_sum.<locals>.array_sum_impl-------------------\n",
      "label 0:\n",
      "    arr = arg(0, name=arr)                   ['arr']\n",
      "    $2load_deref.0 = freevar(zero: 0.0)      ['$2load_deref.0']\n",
      "    c = $2load_deref.0                       ['$2load_deref.0', 'c']\n",
      "    $6load_global.1 = global(np: <module 'numpy' from '/home/gmarkall/miniconda3/envs/numba/lib/python3.8/site-packages/numpy/__init__.py'>) ['$6load_global.1']\n",
      "    $8load_method.2 = getattr(value=$6load_global.1, attr=nditer) ['$6load_global.1', '$8load_method.2']\n",
      "    $12call_method.4 = call $8load_method.2(arr, func=$8load_method.2, args=[Var(arr, arraymath.py:166)], kws=(), vararg=None) ['$12call_method.4', '$8load_method.2', 'arr']\n",
      "    $14get_iter.5 = getiter(value=$12call_method.4) ['$12call_method.4', '$14get_iter.5']\n",
      "    $phi16.0 = $14get_iter.5                 ['$14get_iter.5', '$phi16.0']\n",
      "    jump 16                                  []\n",
      "label 16:\n",
      "    $16for_iter.1 = iternext(value=$phi16.0) ['$16for_iter.1', '$phi16.0']\n",
      "    $16for_iter.2 = pair_first(value=$16for_iter.1) ['$16for_iter.1', '$16for_iter.2']\n",
      "    $16for_iter.3 = pair_second(value=$16for_iter.1) ['$16for_iter.1', '$16for_iter.3']\n",
      "    $phi18.1 = $16for_iter.2                 ['$16for_iter.2', '$phi18.1']\n",
      "    branch $16for_iter.3, 18, 34             ['$16for_iter.3']\n",
      "label 18:\n",
      "    v = $phi18.1                             ['$phi18.1', 'v']\n",
      "    $24load_method.4 = getattr(value=v, attr=item) ['$24load_method.4', 'v']\n",
      "    $26call_method.5 = call $24load_method.4(func=$24load_method.4, args=[], kws=(), vararg=None) ['$24load_method.4', '$26call_method.5']\n",
      "    $28inplace_add.6 = inplace_binop(fn=<built-in function iadd>, immutable_fn=<built-in function add>, lhs=c, rhs=$26call_method.5, static_lhs=Undefined, static_rhs=Undefined) ['$26call_method.5', '$28inplace_add.6', 'c']\n",
      "    c = $28inplace_add.6                     ['$28inplace_add.6', 'c']\n",
      "    jump 16                                  []\n",
      "label 34:\n",
      "    $36return_value.1 = cast(value=c)        ['$36return_value.1', 'c']\n",
      "    return $36return_value.1                 ['$36return_value.1']\n",
      "\n",
      "-----------------------------------propagate------------------------------------\n",
      "---- type variables ----\n",
      "[$12call_method.4 := <undecided>,\n",
      " $14get_iter.5 := <undecided>,\n",
      " $16for_iter.1 := <undecided>,\n",
      " $16for_iter.2 := <undecided>,\n",
      " $24load_method.4 := <undecided>,\n",
      " $26call_method.5 := <undecided>,\n",
      " $28inplace_add.6 := <undecided>,\n",
      " $2load_deref.0 := float64,\n",
      " $36return_value.1 := float64,\n",
      " $6load_global.1 := Module(<module 'numpy' from '/home/gmarkall/miniconda3/envs/numba/lib/python3.8/site-packages/numpy/__init__.py'>),\n",
      " $8load_method.2 := <undecided>,\n",
      " $phi16.0 := <undecided>,\n",
      " $phi18.1 := <undecided>,\n",
      " arg.arr := array(float64, 1d, C),\n",
      " arr := array(float64, 1d, C),\n",
      " c := float64,\n",
      " c.1 := float64,\n",
      " c.2 := float64,\n",
      " v := <undecided>]\n",
      "-----------------------------------propagate------------------------------------\n",
      "---- type variables ----\n",
      "[$12call_method.4 := <undecided>,\n",
      " $14get_iter.5 := <undecided>,\n",
      " $16for_iter.1 := <undecided>,\n",
      " $16for_iter.2 := <undecided>,\n",
      " $24load_method.4 := <undecided>,\n",
      " $26call_method.5 := <undecided>,\n",
      " $28inplace_add.6 := <undecided>,\n",
      " $2load_deref.0 := float64,\n",
      " $36return_value.1 := float64,\n",
      " $6load_global.1 := Module(<module 'numpy' from '/home/gmarkall/miniconda3/envs/numba/lib/python3.8/site-packages/numpy/__init__.py'>),\n",
      " $8load_method.2 := <undecided>,\n",
      " $phi16.0 := <undecided>,\n",
      " $phi18.1 := <undecided>,\n",
      " arg.arr := array(float64, 1d, C),\n",
      " arr := array(float64, 1d, C),\n",
      " c := float64,\n",
      " c.1 := float64,\n",
      " c.2 := float64,\n",
      " v := <undecided>]\n"
     ]
    },
    {
     "name": "stderr",
     "output_type": "stream",
     "text": [
      "Traceback (most recent call last):\n",
      "  File \"<stdin>\", line 12, in <module>\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/cuda/compiler.py\", line 758, in __call__\n",
      "    kernel = self.specialize(*args)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/cuda/compiler.py\", line 769, in specialize\n",
      "    kernel = self.compile(argtypes)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/cuda/compiler.py\", line 784, in compile\n",
      "    kernel = compile_kernel(self.py_func, argtypes,\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/compiler_lock.py\", line 32, in _acquire_compile_lock\n",
      "    return func(*args, **kwargs)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/cuda/compiler.py\", line 57, in compile_kernel\n",
      "    cres = compile_cuda(pyfunc, types.void, args, debug=debug, inline=inline)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/compiler_lock.py\", line 32, in _acquire_compile_lock\n",
      "    return func(*args, **kwargs)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/cuda/compiler.py\", line 40, in compile_cuda\n",
      "    cres = compiler.compile_extra(typingctx=typingctx,\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/compiler.py\", line 568, in compile_extra\n",
      "    return pipeline.compile_extra(func)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/compiler.py\", line 339, in compile_extra\n",
      "    return self._compile_bytecode()\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/compiler.py\", line 401, in _compile_bytecode\n",
      "    return self._compile_core()\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/compiler.py\", line 381, in _compile_core\n",
      "    raise e\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/compiler.py\", line 372, in _compile_core\n",
      "    pm.run(self.state)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/compiler_machinery.py\", line 341, in run\n",
      "    raise patched_exception\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/compiler_machinery.py\", line 332, in run\n",
      "    self._runPass(idx, pass_inst, state)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/compiler_lock.py\", line 32, in _acquire_compile_lock\n",
      "    return func(*args, **kwargs)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/compiler_machinery.py\", line 291, in _runPass\n",
      "    mutated |= check(pss.run_pass, internal_state)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/compiler_machinery.py\", line 264, in check\n",
      "    mangled = func(compiler_state)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/typed_passes.py\", line 442, in run_pass\n",
      "    NativeLowering().run_pass(state)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/typed_passes.py\", line 370, in run_pass\n",
      "    lower.lower()\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/lowering.py\", line 179, in lower\n",
      "    self.lower_normal_function(self.fndesc)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/lowering.py\", line 227, in lower_normal_function\n",
      "    entry_block_tail = self.lower_function_body()\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/lowering.py\", line 253, in lower_function_body\n",
      "    self.lower_block(block)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/lowering.py\", line 267, in lower_block\n",
      "    self.lower_inst(inst)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/lowering.py\", line 364, in lower_inst\n",
      "    val = self.lower_assign(ty, inst)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/lowering.py\", line 538, in lower_assign\n",
      "    return self.lower_expr(ty, value)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/lowering.py\", line 1062, in lower_expr\n",
      "    res = self.lower_call(resty, expr)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/lowering.py\", line 798, in lower_call\n",
      "    res = self._lower_call_normal(fnty, expr, signature)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/lowering.py\", line 1033, in _lower_call_normal\n",
      "    res = impl(self.builder, argvals, self.loc)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/base.py\", line 1155, in __call__\n",
      "    res = self._imp(self._context, builder, self._sig, args, loc=loc)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/base.py\", line 1185, in wrapper\n",
      "    return fn(*args, **kwargs)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/np/arraymath.py\", line 171, in array_sum\n",
      "    res = context.compile_internal(builder, array_sum_impl, sig, args,\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/base.py\", line 875, in compile_internal\n",
      "    cres = self.compile_subroutine(builder, impl, sig, locals)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/base.py\", line 860, in compile_subroutine\n",
      "    cres = self._compile_subroutine_no_cache(builder, impl, sig,\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/base.py\", line 831, in _compile_subroutine_no_cache\n",
      "    cres = compiler.compile_internal(self.typing_context, self,\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/compiler.py\", line 642, in compile_internal\n",
      "    return pipeline.compile_extra(func)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/compiler.py\", line 339, in compile_extra\n",
      "    return self._compile_bytecode()\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/compiler.py\", line 401, in _compile_bytecode\n",
      "    return self._compile_core()\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/compiler.py\", line 381, in _compile_core\n",
      "    raise e\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/compiler.py\", line 372, in _compile_core\n",
      "    pm.run(self.state)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/compiler_machinery.py\", line 341, in run\n",
      "    raise patched_exception\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/compiler_machinery.py\", line 332, in run\n",
      "    self._runPass(idx, pass_inst, state)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/compiler_lock.py\", line 32, in _acquire_compile_lock\n",
      "    return func(*args, **kwargs)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/compiler_machinery.py\", line 291, in _runPass\n",
      "    mutated |= check(pss.run_pass, internal_state)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/compiler_machinery.py\", line 264, in check\n",
      "    mangled = func(compiler_state)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/typed_passes.py\", line 92, in run_pass\n",
      "    typemap, return_type, calltypes = type_inference_stage(\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/typed_passes.py\", line 70, in type_inference_stage\n",
      "    infer.propagate(raise_errors=raise_errors)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/typeinfer.py\", line 986, in propagate\n",
      "    raise errors[0]\n",
      "numba.core.errors.TypingError: Failed in nopython mode pipeline (step: nopython mode backend)\n",
      "Failed in nopython mode pipeline (step: nopython frontend)\n",
      "Use of unsupported NumPy function 'numpy.nditer' or unsupported use of the function.\n",
      "\n",
      "File \"../../../../numba/numba/np/arraymath.py\", line 167:\n",
      "    def array_sum_impl(arr):\n",
      "        <source elided>\n",
      "        c = zero\n",
      "        for v in np.nditer(arr):\n",
      "        ^\n",
      "\n",
      "[1] During: typing of get attribute at /home/gmarkall/numbadev/numba/numba/np/arraymath.py (167)\n",
      "\n",
      "File \"../../../../numba/numba/np/arraymath.py\", line 167:\n",
      "    def array_sum_impl(arr):\n",
      "        <source elided>\n",
      "        c = zero\n",
      "        for v in np.nditer(arr):\n",
      "        ^\n",
      "\n",
      "[1] During: lowering \"$6call_method.2 = call $4load_method.1(func=$4load_method.1, args=[], kws=(), vararg=None)\" at <stdin> (9)\n"
     ]
    },
    {
     "ename": "CalledProcessError",
     "evalue": "Command 'b'from numba import cuda\\nfrom numba import config\\nimport numpy as np\\nconfig.DEBUG_TYPEINFER = True\\nconfig.DUMP_IR = True\\n\\n@cuda.jit\\ndef sum_reduce(x):\\n    x[0] = x.sum()\\n\\nx = np.ones(10)\\nsum_reduce(x)\\n'' returned non-zero exit status 1.",
     "output_type": "error",
     "traceback": [
      "\u001b[0;31m---------------------------------------------------------------------------\u001b[0m",
      "\u001b[0;31mCalledProcessError\u001b[0m                        Traceback (most recent call last)",
      "\u001b[0;32m<ipython-input-13-ff3a8e67780e>\u001b[0m in \u001b[0;36m<module>\u001b[0;34m\u001b[0m\n\u001b[0;32m----> 1\u001b[0;31m \u001b[0mget_ipython\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mrun_cell_magic\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m'script'\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m'python'\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m'from numba import cuda\\nfrom numba import config\\nimport numpy as np\\nconfig.DEBUG_TYPEINFER = True\\nconfig.DUMP_IR = True\\n\\n@cuda.jit\\ndef sum_reduce(x):\\n    x[0] = x.sum()\\n\\nx = np.ones(10)\\nsum_reduce(x)\\n'\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m",
      "\u001b[0;32m~/miniconda3/envs/numba/lib/python3.8/site-packages/IPython/core/interactiveshell.py\u001b[0m in \u001b[0;36mrun_cell_magic\u001b[0;34m(self, magic_name, line, cell)\u001b[0m\n\u001b[1;32m   2360\u001b[0m             \u001b[0;32mwith\u001b[0m \u001b[0mself\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mbuiltin_trap\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m   2361\u001b[0m                 \u001b[0margs\u001b[0m \u001b[0;34m=\u001b[0m \u001b[0;34m(\u001b[0m\u001b[0mmagic_arg_s\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mcell\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m-> 2362\u001b[0;31m                 \u001b[0mresult\u001b[0m \u001b[0;34m=\u001b[0m \u001b[0mfn\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m*\u001b[0m\u001b[0margs\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m**\u001b[0m\u001b[0mkwargs\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m   2363\u001b[0m             \u001b[0;32mreturn\u001b[0m \u001b[0mresult\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m   2364\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;32m<decorator-gen-110>\u001b[0m in \u001b[0;36mshebang\u001b[0;34m(self, line, cell)\u001b[0m\n",
      "\u001b[0;32m~/miniconda3/envs/numba/lib/python3.8/site-packages/IPython/core/magic.py\u001b[0m in \u001b[0;36m<lambda>\u001b[0;34m(f, *a, **k)\u001b[0m\n\u001b[1;32m    185\u001b[0m     \u001b[0;31m# but it's overkill for just that one bit of state.\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    186\u001b[0m     \u001b[0;32mdef\u001b[0m \u001b[0mmagic_deco\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0marg\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m--> 187\u001b[0;31m         \u001b[0mcall\u001b[0m \u001b[0;34m=\u001b[0m \u001b[0;32mlambda\u001b[0m \u001b[0mf\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m*\u001b[0m\u001b[0ma\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m**\u001b[0m\u001b[0mk\u001b[0m\u001b[0;34m:\u001b[0m \u001b[0mf\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m*\u001b[0m\u001b[0ma\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m**\u001b[0m\u001b[0mk\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m    188\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    189\u001b[0m         \u001b[0;32mif\u001b[0m \u001b[0mcallable\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0marg\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;32m~/miniconda3/envs/numba/lib/python3.8/site-packages/IPython/core/magics/script.py\u001b[0m in \u001b[0;36mshebang\u001b[0;34m(self, line, cell)\u001b[0m\n\u001b[1;32m    243\u001b[0m             \u001b[0msys\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mstderr\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mflush\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    244\u001b[0m         \u001b[0;32mif\u001b[0m \u001b[0margs\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mraise_error\u001b[0m \u001b[0;32mand\u001b[0m \u001b[0mp\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mreturncode\u001b[0m\u001b[0;34m!=\u001b[0m\u001b[0;36m0\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m--> 245\u001b[0;31m             \u001b[0;32mraise\u001b[0m \u001b[0mCalledProcessError\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mp\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mreturncode\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mcell\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0moutput\u001b[0m\u001b[0;34m=\u001b[0m\u001b[0mout\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mstderr\u001b[0m\u001b[0;34m=\u001b[0m\u001b[0merr\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m    246\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    247\u001b[0m     \u001b[0;32mdef\u001b[0m \u001b[0m_run_script\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mself\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mp\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mcell\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mto_close\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;31mCalledProcessError\u001b[0m: Command 'b'from numba import cuda\\nfrom numba import config\\nimport numpy as np\\nconfig.DEBUG_TYPEINFER = True\\nconfig.DUMP_IR = True\\n\\n@cuda.jit\\ndef sum_reduce(x):\\n    x[0] = x.sum()\\n\\nx = np.ones(10)\\nsum_reduce(x)\\n'' returned non-zero exit status 1."
     ]
    }
   ],
   "source": [
    "%%script python\n",
    "from numba import cuda\n",
    "from numba import config\n",
    "import numpy as np\n",
    "config.DEBUG_TYPEINFER = True\n",
    "config.DUMP_IR = True\n",
    "\n",
    "@cuda.jit\n",
    "def sum_reduce(x):\n",
    "    x[0] = x.sum()\n",
    "\n",
    "x = np.ones(10)\n",
    "sum_reduce(x)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## Branch Elimination\n",
    "\n",
    "Sometimes Numba can eliminate code from dead branches, if it can determine that the branch will never run for a given set of argument types - this can avoid a unification error that would otherwise have occurred if Numba could not eliminate these dead branches. The next example demonstrates this capability when it does work, and also when it doesn't."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 14,
   "metadata": {},
   "outputs": [],
   "source": [
    "@njit\n",
    "def branch_elim_example(a, b, cond):\n",
    "    if cond is None:\n",
    "        return a\n",
    "    else:\n",
    "        return b"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "This call, where `cond` is `None`, succeeds due to the elision of the `else` branch:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 15,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "1"
      ]
     },
     "execution_count": 15,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "branch_elim_example(1, (1, 2), None)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "In the following call branch elimination fails, forcing an attempt to unify two things that cannot be unified:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 16,
   "metadata": {},
   "outputs": [
    {
     "ename": "TypingError",
     "evalue": "Failed in nopython mode pipeline (step: nopython frontend)\nCan't unify return type from the following types: UniTuple(int64 x 2), int64\nReturn of: IR name '$16return_value.1', type 'UniTuple(int64 x 2)', location: \nFile \"<ipython-input-14-483ae52ee5e6>\", line 6:\ndef branch_elim_example(a, b, cond):\n    <source elided>\n    else:\n        return b\n        ^\nReturn of: IR name '$12return_value.1', type 'int64', location: \nFile \"<ipython-input-14-483ae52ee5e6>\", line 4:\ndef branch_elim_example(a, b, cond):\n    <source elided>\n    if cond is None:\n        return a\n        ^",
     "output_type": "error",
     "traceback": [
      "\u001b[0;31m---------------------------------------------------------------------------\u001b[0m",
      "\u001b[0;31mTypingError\u001b[0m                               Traceback (most recent call last)",
      "\u001b[0;32m<ipython-input-16-d72526aa5562>\u001b[0m in \u001b[0;36m<module>\u001b[0;34m\u001b[0m\n\u001b[0;32m----> 1\u001b[0;31m \u001b[0mbranch_elim_example\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;36m1\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m(\u001b[0m\u001b[0;36m1\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;36m2\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;32mTrue\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m",
      "\u001b[0;32m~/numbadev/numba/numba/core/dispatcher.py\u001b[0m in \u001b[0;36m_compile_for_args\u001b[0;34m(self, *args, **kws)\u001b[0m\n\u001b[1;32m    399\u001b[0m                 \u001b[0me\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mpatch_message\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mmsg\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    400\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m--> 401\u001b[0;31m             \u001b[0merror_rewrite\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0me\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m'typing'\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m    402\u001b[0m         \u001b[0;32mexcept\u001b[0m \u001b[0merrors\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mUnsupportedError\u001b[0m \u001b[0;32mas\u001b[0m \u001b[0me\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    403\u001b[0m             \u001b[0;31m# Something unsupported is present in the user code, add help info\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;32m~/numbadev/numba/numba/core/dispatcher.py\u001b[0m in \u001b[0;36merror_rewrite\u001b[0;34m(e, issue_type)\u001b[0m\n\u001b[1;32m    342\u001b[0m                 \u001b[0;32mraise\u001b[0m \u001b[0me\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    343\u001b[0m             \u001b[0;32melse\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m--> 344\u001b[0;31m                 \u001b[0mreraise\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mtype\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0me\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0me\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;32mNone\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m    345\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    346\u001b[0m         \u001b[0margtypes\u001b[0m \u001b[0;34m=\u001b[0m \u001b[0;34m[\u001b[0m\u001b[0;34m]\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;32m~/numbadev/numba/numba/core/utils.py\u001b[0m in \u001b[0;36mreraise\u001b[0;34m(tp, value, tb)\u001b[0m\n\u001b[1;32m     78\u001b[0m         \u001b[0mvalue\u001b[0m \u001b[0;34m=\u001b[0m \u001b[0mtp\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m     79\u001b[0m     \u001b[0;32mif\u001b[0m \u001b[0mvalue\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0m__traceback__\u001b[0m \u001b[0;32mis\u001b[0m \u001b[0;32mnot\u001b[0m \u001b[0mtb\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m---> 80\u001b[0;31m         \u001b[0;32mraise\u001b[0m \u001b[0mvalue\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mwith_traceback\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mtb\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m     81\u001b[0m     \u001b[0;32mraise\u001b[0m \u001b[0mvalue\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m     82\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;31mTypingError\u001b[0m: Failed in nopython mode pipeline (step: nopython frontend)\nCan't unify return type from the following types: UniTuple(int64 x 2), int64\nReturn of: IR name '$16return_value.1', type 'UniTuple(int64 x 2)', location: \nFile \"<ipython-input-14-483ae52ee5e6>\", line 6:\ndef branch_elim_example(a, b, cond):\n    <source elided>\n    else:\n        return b\n        ^\nReturn of: IR name '$12return_value.1', type 'int64', location: \nFile \"<ipython-input-14-483ae52ee5e6>\", line 4:\ndef branch_elim_example(a, b, cond):\n    <source elided>\n    if cond is None:\n        return a\n        ^"
     ]
    }
   ],
   "source": [
    "branch_elim_example(1, (1, 2), True)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "The following cell contains the same function and call, but run with `%%script` so you can inspect the IR and typing if you wish."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 17,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "--------------------------IR DUMP: branch_elim_example--------------------------\n",
      "label 0:\n",
      "    a = arg(0, name=a)                       ['a']\n",
      "    b = arg(1, name=b)                       ['b']\n",
      "    cond = arg(2, name=cond)                 ['cond']\n",
      "    $const4.1 = const(NoneType, None)        ['$const4.1']\n",
      "    $6compare_op.2 = cond is $const4.1       ['$6compare_op.2', '$const4.1', 'cond']\n",
      "    branch $6compare_op.2, 10, 14            ['$6compare_op.2']\n",
      "label 10:\n",
      "    $12return_value.1 = cast(value=a)        ['$12return_value.1', 'a']\n",
      "    return $12return_value.1                 ['$12return_value.1']\n",
      "label 14:\n",
      "    $16return_value.1 = cast(value=b)        ['$16return_value.1', 'b']\n",
      "    return $16return_value.1                 ['$16return_value.1']\n",
      "\n",
      "--------------------------IR DUMP: branch_elim_example--------------------------\n",
      "label 0:\n",
      "    a = arg(0, name=a)                       ['a']\n",
      "    b = arg(1, name=b)                       ['b']\n",
      "    cond = arg(2, name=cond)                 ['cond']\n",
      "    $const4.1 = const(NoneType, None)        ['$const4.1']\n",
      "    $6compare_op.2 = cond is $const4.1       ['$6compare_op.2', '$const4.1', 'cond']\n",
      "    branch $6compare_op.2, 10, 14            ['$6compare_op.2']\n",
      "label 10:\n",
      "    $12return_value.1 = cast(value=a)        ['$12return_value.1', 'a']\n",
      "    return $12return_value.1                 ['$12return_value.1']\n",
      "label 14:\n",
      "    $16return_value.1 = cast(value=b)        ['$16return_value.1', 'b']\n",
      "    return $16return_value.1                 ['$16return_value.1']\n",
      "\n",
      "-----------------------------------propagate------------------------------------\n",
      "---- type variables ----\n",
      "[$12return_value.1 := int64,\n",
      " $16return_value.1 := UniTuple(int64 x 2),\n",
      " $6compare_op.2 := bool,\n",
      " $const4.1 := none,\n",
      " a := int64,\n",
      " arg.a := int64,\n",
      " arg.b := UniTuple(int64 x 2),\n",
      " arg.cond := bool,\n",
      " b := UniTuple(int64 x 2),\n",
      " cond := bool]\n",
      "-----------------------------------propagate------------------------------------\n",
      "---- type variables ----\n",
      "[$12return_value.1 := int64,\n",
      " $16return_value.1 := UniTuple(int64 x 2),\n",
      " $6compare_op.2 := bool,\n",
      " $const4.1 := none,\n",
      " a := int64,\n",
      " arg.a := int64,\n",
      " arg.b := UniTuple(int64 x 2),\n",
      " arg.cond := bool,\n",
      " b := UniTuple(int64 x 2),\n",
      " cond := bool]\n"
     ]
    },
    {
     "name": "stderr",
     "output_type": "stream",
     "text": [
      "Traceback (most recent call last):\n",
      "  File \"<stdin>\", line 14, in <module>\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/dispatcher.py\", line 401, in _compile_for_args\n",
      "    error_rewrite(e, 'typing')\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/dispatcher.py\", line 344, in error_rewrite\n",
      "    reraise(type(e), e, None)\n",
      "  File \"/home/gmarkall/numbadev/numba/numba/core/utils.py\", line 80, in reraise\n",
      "    raise value.with_traceback(tb)\n",
      "numba.core.errors.TypingError: Failed in nopython mode pipeline (step: nopython frontend)\n",
      "Can't unify return type from the following types: UniTuple(int64 x 2), int64\n",
      "Return of: IR name '$12return_value.1', type 'int64', location: \n",
      "File \"<stdin>\", line 10:\n",
      "<source missing, REPL/exec in use?>\n",
      "Return of: IR name '$16return_value.1', type 'UniTuple(int64 x 2)', location: \n",
      "File \"<stdin>\", line 12:\n",
      "<source missing, REPL/exec in use?>\n"
     ]
    },
    {
     "ename": "CalledProcessError",
     "evalue": "Command 'b'from numba import njit\\nfrom numba import config\\nimport numpy as np\\nconfig.DEBUG_TYPEINFER = True\\nconfig.DUMP_IR = True\\n\\n@njit\\ndef branch_elim_example(a, b, cond):\\n    if cond is None:\\n        return a\\n    else:\\n        return b\\n    \\nbranch_elim_example(1, (1, 2), True)\\n'' returned non-zero exit status 1.",
     "output_type": "error",
     "traceback": [
      "\u001b[0;31m---------------------------------------------------------------------------\u001b[0m",
      "\u001b[0;31mCalledProcessError\u001b[0m                        Traceback (most recent call last)",
      "\u001b[0;32m<ipython-input-17-1e2a3a37bb5e>\u001b[0m in \u001b[0;36m<module>\u001b[0;34m\u001b[0m\n\u001b[0;32m----> 1\u001b[0;31m \u001b[0mget_ipython\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mrun_cell_magic\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m'script'\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m'python'\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m'from numba import njit\\nfrom numba import config\\nimport numpy as np\\nconfig.DEBUG_TYPEINFER = True\\nconfig.DUMP_IR = True\\n\\n@njit\\ndef branch_elim_example(a, b, cond):\\n    if cond is None:\\n        return a\\n    else:\\n        return b\\n    \\nbranch_elim_example(1, (1, 2), True)\\n'\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m",
      "\u001b[0;32m~/miniconda3/envs/numba/lib/python3.8/site-packages/IPython/core/interactiveshell.py\u001b[0m in \u001b[0;36mrun_cell_magic\u001b[0;34m(self, magic_name, line, cell)\u001b[0m\n\u001b[1;32m   2360\u001b[0m             \u001b[0;32mwith\u001b[0m \u001b[0mself\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mbuiltin_trap\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m   2361\u001b[0m                 \u001b[0margs\u001b[0m \u001b[0;34m=\u001b[0m \u001b[0;34m(\u001b[0m\u001b[0mmagic_arg_s\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mcell\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m-> 2362\u001b[0;31m                 \u001b[0mresult\u001b[0m \u001b[0;34m=\u001b[0m \u001b[0mfn\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m*\u001b[0m\u001b[0margs\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m**\u001b[0m\u001b[0mkwargs\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m   2363\u001b[0m             \u001b[0;32mreturn\u001b[0m \u001b[0mresult\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m   2364\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;32m<decorator-gen-110>\u001b[0m in \u001b[0;36mshebang\u001b[0;34m(self, line, cell)\u001b[0m\n",
      "\u001b[0;32m~/miniconda3/envs/numba/lib/python3.8/site-packages/IPython/core/magic.py\u001b[0m in \u001b[0;36m<lambda>\u001b[0;34m(f, *a, **k)\u001b[0m\n\u001b[1;32m    185\u001b[0m     \u001b[0;31m# but it's overkill for just that one bit of state.\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    186\u001b[0m     \u001b[0;32mdef\u001b[0m \u001b[0mmagic_deco\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0marg\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m--> 187\u001b[0;31m         \u001b[0mcall\u001b[0m \u001b[0;34m=\u001b[0m \u001b[0;32mlambda\u001b[0m \u001b[0mf\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m*\u001b[0m\u001b[0ma\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m**\u001b[0m\u001b[0mk\u001b[0m\u001b[0;34m:\u001b[0m \u001b[0mf\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m*\u001b[0m\u001b[0ma\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m**\u001b[0m\u001b[0mk\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m    188\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    189\u001b[0m         \u001b[0;32mif\u001b[0m \u001b[0mcallable\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0marg\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;32m~/miniconda3/envs/numba/lib/python3.8/site-packages/IPython/core/magics/script.py\u001b[0m in \u001b[0;36mshebang\u001b[0;34m(self, line, cell)\u001b[0m\n\u001b[1;32m    243\u001b[0m             \u001b[0msys\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mstderr\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mflush\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    244\u001b[0m         \u001b[0;32mif\u001b[0m \u001b[0margs\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mraise_error\u001b[0m \u001b[0;32mand\u001b[0m \u001b[0mp\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mreturncode\u001b[0m\u001b[0;34m!=\u001b[0m\u001b[0;36m0\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m--> 245\u001b[0;31m             \u001b[0;32mraise\u001b[0m \u001b[0mCalledProcessError\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mp\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mreturncode\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mcell\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0moutput\u001b[0m\u001b[0;34m=\u001b[0m\u001b[0mout\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mstderr\u001b[0m\u001b[0;34m=\u001b[0m\u001b[0merr\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m    246\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    247\u001b[0m     \u001b[0;32mdef\u001b[0m \u001b[0m_run_script\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mself\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mp\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mcell\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mto_close\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;31mCalledProcessError\u001b[0m: Command 'b'from numba import njit\\nfrom numba import config\\nimport numpy as np\\nconfig.DEBUG_TYPEINFER = True\\nconfig.DUMP_IR = True\\n\\n@njit\\ndef branch_elim_example(a, b, cond):\\n    if cond is None:\\n        return a\\n    else:\\n        return b\\n    \\nbranch_elim_example(1, (1, 2), True)\\n'' returned non-zero exit status 1."
     ]
    }
   ],
   "source": [
    "%%script python\n",
    "from numba import njit\n",
    "from numba import config\n",
    "import numpy as np\n",
    "config.DEBUG_TYPEINFER = True\n",
    "config.DUMP_IR = True\n",
    "\n",
    "@njit\n",
    "def branch_elim_example(a, b, cond):\n",
    "    if cond is None:\n",
    "        return a\n",
    "    else:\n",
    "        return b\n",
    "    \n",
    "branch_elim_example(1, (1, 2), True)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### General summary of Branch Elimination\n",
    "\n",
    "* Branch elimination can sometimes remove dead code and prevent unification errors.\n",
    "* In practive if you find that some calls fail to unify, then branch elimination may be involved.\n",
    "\n",
    "# CUDA-specific issues\n",
    "\n",
    "This section looks at a few issues where performance on CUDA can be impacted due to the typing. These are:\n",
    "\n",
    "* Widening unification\n",
    "* Widening arithmetic, and its propagation\n",
    "* The typing of integer arithmetic\n",
    "* Register usage control"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## Widening unification\n",
    "\n",
    "Unification of types can result in a type that is larger than any of the types from the set that was unified. This first example uses the CPU target because it makes for a simpler example, but the general idea of widening unification applies to the CUDA target as well."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 18,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "1.0"
      ]
     },
     "execution_count": 18,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "@njit\n",
    "def select(a, b, threshold, value):\n",
    "    if threshold < value:\n",
    "        r = a\n",
    "    else:\n",
    "        r = b\n",
    "    return r\n",
    "\n",
    "a = np.float32(1)\n",
    "b = np.int32(2)\n",
    "select(a, b, 10, 11)  # Call with (float32, int32, int64, int64)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "After the call, we can inspect the typing:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 19,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "select (float32, int32, int64, int64)\n",
      "--------------------------------------------------------------------------------\n",
      "# File: <ipython-input-18-0f89bde90efd>\n",
      "# --- LINE 1 --- \n",
      "\n",
      "@njit\n",
      "\n",
      "# --- LINE 2 --- \n",
      "\n",
      "def select(a, b, threshold, value):\n",
      "\n",
      "    # --- LINE 3 --- \n",
      "    # label 0\n",
      "    #   a = arg(0, name=a)  :: float32\n",
      "    #   b = arg(1, name=b)  :: int32\n",
      "    #   threshold = arg(2, name=threshold)  :: int64\n",
      "    #   value = arg(3, name=value)  :: int64\n",
      "    #   $6compare_op.2 = threshold < value  :: bool\n",
      "    #   del value\n",
      "    #   del threshold\n",
      "    #   branch $6compare_op.2, 10, 16\n",
      "\n",
      "    if threshold < value:\n",
      "\n",
      "        # --- LINE 4 --- \n",
      "        # label 10\n",
      "        #   del b\n",
      "        #   del $6compare_op.2\n",
      "        #   r = a  :: float32\n",
      "        #   del a\n",
      "        #   jump 20\n",
      "\n",
      "        r = a\n",
      "\n",
      "    # --- LINE 5 --- \n",
      "\n",
      "    else:\n",
      "\n",
      "        # --- LINE 6 --- \n",
      "        # label 16\n",
      "        #   del a\n",
      "        #   del $6compare_op.2\n",
      "        #   r.1 = b  :: int32\n",
      "        #   del b\n",
      "\n",
      "        r = b\n",
      "\n",
      "    # --- LINE 7 --- \n",
      "    #   jump 20\n",
      "    # label 20\n",
      "    #   r.2 = phi(incoming_values=[Var(r.1, <ipython-input-18-0f89bde90efd>:6), Var(r, <ipython-input-18-0f89bde90efd>:4)], incoming_blocks=[16, 10])  :: float64\n",
      "    #   del r.1\n",
      "    #   del r\n",
      "    #   $22return_value.1 = cast(value=r.2)  :: float64\n",
      "    #   del r.2\n",
      "    #   return $22return_value.1\n",
      "\n",
      "    return r\n",
      "\n",
      "\n",
      "================================================================================\n"
     ]
    }
   ],
   "source": [
    "select.inspect_types()"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Exercises\n",
    "\n",
    "Try to determine from the typing:\n",
    "\n",
    "* Try to determine the return type from the typing output. What was it? \n",
    "* Why was this type chosen instead of one of the types in the set?\n",
    "* Fix the above code so that the return type is no wider than any of the input types.\n",
    "\n",
    "## Width of constants\n",
    "\n",
    "The default width of constants and the propagation of their width can have an effect on the typing that results in slower code due to the use of double precision units, and increased register usage. We will build up an example step-by-step to see the impact on the propagated types and the knock-on effects on the LLVM IR and PTX code. \n",
    "\n",
    "We begin with a very simple example, where we assign a constant to an array element:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 20,
   "metadata": {},
   "outputs": [],
   "source": [
    "from numba import void\n",
    "\n",
    "@cuda.jit(void(float32[:]))\n",
    "def assign_constant(x):\n",
    "    x[0] = 2.0"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Now let's see the typing:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 21,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "_ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE (array(float32, 1d, A),)\n",
      "--------------------------------------------------------------------------------\n",
      "# File: <ipython-input-20-f028b2d99a7d>\n",
      "# --- LINE 3 --- \n",
      "\n",
      "@cuda.jit(void(float32[:]))\n",
      "\n",
      "# --- LINE 4 --- \n",
      "\n",
      "def assign_constant(x):\n",
      "\n",
      "    # --- LINE 5 --- \n",
      "    # label 0\n",
      "    #   x = arg(0, name=x)  :: array(float32, 1d, A)\n",
      "    #   $const2.0 = const(float, 2.0)  :: float64\n",
      "    #   $const6.2 = const(int, 0)  :: Literal[int](0)\n",
      "    #   x[0] = $const2.0\n",
      "    #   del x\n",
      "    #   del $const6.2\n",
      "    #   del $const2.0\n",
      "    #   $const10.3 = const(NoneType, None)  :: none\n",
      "    #   $12return_value.4 = cast(value=$const10.3)  :: none\n",
      "    #   del $const10.3\n",
      "    #   return $12return_value.4\n",
      "\n",
      "    x[0] = 2.0\n",
      "\n",
      "\n",
      "================================================================================\n"
     ]
    }
   ],
   "source": [
    "assign_constant.inspect_types()"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "The constant has a type of `float64`. Now let's look at what LLVM does with that, by viewing the LLVM IR after LLVM optimizations:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 22,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "source_filename = \"<string>\"\n",
      "target datalayout = \"e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64\"\n",
      "target triple = \"nvptx64-nvidia-cuda\"\n",
      "\n",
      "@\"_ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE__errcode__\" = local_unnamed_addr global i32 0\n",
      "@\"_ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE__tidx__\" = local_unnamed_addr global i32 0\n",
      "@\"_ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE__ctaidx__\" = local_unnamed_addr global i32 0\n",
      "@\"_ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE__tidy__\" = local_unnamed_addr global i32 0\n",
      "@\"_ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE__ctaidy__\" = local_unnamed_addr global i32 0\n",
      "@\"_ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE__tidz__\" = local_unnamed_addr global i32 0\n",
      "@\"_ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE__ctaidz__\" = local_unnamed_addr global i32 0\n",
      "@\"_ZN08NumbaEnv8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE\" = common local_unnamed_addr global i8* null\n",
      "\n",
      "; Function Attrs: norecurse nounwind writeonly\n",
      "define void @\"_ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE\"(i8* nocapture readnone %.1, i8* nocapture readnone %.2, i64 %.3, i64 %.4, float* nocapture %.5, i64 %.6, i64 %.7) local_unnamed_addr #0 {\n",
      ".9:\n",
      "  %.12 = alloca i8*, align 8\n",
      "  %.17 = call i32 @\"_ZN8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE\"(i8** nonnull %.12, i8* %.1, i8* %.2, i64 %.3, i64 %.4, float* %.5, i64 %.6, i64 %.7)\n",
      "  ret void\n",
      "}\n",
      "\n",
      "; Function Attrs: norecurse nounwind writeonly\n",
      "define linkonce_odr i32 @\"_ZN8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE\"(i8** nocapture %.ret, i8* nocapture readnone %arg.x.0, i8* nocapture readnone %arg.x.1, i64 %arg.x.2, i64 %arg.x.3, float* nocapture %arg.x.4, i64 %arg.x.5.0, i64 %arg.x.6.0) local_unnamed_addr #0 {\n",
      "entry:\n",
      "  store float 2.000000e+00, float* %arg.x.4, align 4\n",
      "  store i8* null, i8** %.ret, align 8\n",
      "  ret i32 0\n",
      "}\n",
      "\n",
      "attributes #0 = { norecurse nounwind writeonly }\n",
      "\n",
      "!nvvm.annotations = !{!0}\n",
      "!nvvmir.version = !{!1}\n",
      "\n",
      "!0 = !{void (i8*, i8*, i64, i64, float*, i64, i64)* @\"_ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE\", !\"kernel\", i32 1}\n",
      "!1 = !{i32 1, i32 2, i32 2, i32 0}\n",
      "\n"
     ]
    }
   ],
   "source": [
    "print(assign_constant.inspect_llvm())"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "It turns out that the LLVM optimizer was able to convert this back to a 32-bit constant: `store float 2.000000e+00, float* %arg.x.4, align 4`.\n",
    "\n",
    "We see a similar width in the PTX:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 23,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "//\n",
      "// Generated by NVIDIA NVVM Compiler\n",
      "//\n",
      "// Compiler Build ID: CL-27506705\n",
      "// Cuda compilation tools, release 10.2, V10.2.89\n",
      "// Based on LLVM 3.4svn\n",
      "//\n",
      "\n",
      ".version 6.5\n",
      ".target sm_70\n",
      ".address_size 64\n",
      "\n",
      "\t// .globl\t_ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE\n",
      ".visible .global .align 4 .u32 _ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE__errcode__;\n",
      ".visible .global .align 4 .u32 _ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE__tidx__;\n",
      ".visible .global .align 4 .u32 _ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE__ctaidx__;\n",
      ".visible .global .align 4 .u32 _ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE__tidy__;\n",
      ".visible .global .align 4 .u32 _ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE__ctaidy__;\n",
      ".visible .global .align 4 .u32 _ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE__tidz__;\n",
      ".visible .global .align 4 .u32 _ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE__ctaidz__;\n",
      ".common .global .align 8 .u64 _ZN08NumbaEnv8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE;\n",
      "\n",
      ".visible .entry _ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE(\n",
      "\t.param .u64 _ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE_param_0,\n",
      "\t.param .u64 _ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE_param_1,\n",
      "\t.param .u64 _ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE_param_2,\n",
      "\t.param .u64 _ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE_param_3,\n",
      "\t.param .u64 _ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE_param_4,\n",
      "\t.param .u64 _ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE_param_5,\n",
      "\t.param .u64 _ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE_param_6\n",
      ")\n",
      "{\n",
      "\t.reg .b32 \t%r<2>;\n",
      "\t.reg .b64 \t%rd<3>;\n",
      "\n",
      "\n",
      "\tld.param.u64 \t%rd1, [_ZN6cudapy8__main__19assign_constant$248E5ArrayIfLi1E1A7mutable7alignedE_param_4];\n",
      "\tcvta.to.global.u64 \t%rd2, %rd1;\n",
      "\tmov.u32 \t%r1, 1073741824;\n",
      "\tst.global.u32 \t[%rd2], %r1;\n",
      "\tret;\n",
      "}\n",
      "\n",
      "\n",
      "\u0000\n"
     ]
    }
   ],
   "source": [
    "print(assign_constant.inspect_asm())"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Correspondingly, we have `mov.u32 \t%r1, 1073741824;`. So far, so good.\n",
    "\n",
    "### Increasing complexity slightly - in-place addition\n",
    "\n",
    "Now let's build up the example a little - instead of assigning a constant, we add a constant to the array element:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 24,
   "metadata": {},
   "outputs": [],
   "source": [
    "@cuda.jit(void(float32[:]))\n",
    "def add_constant(x):\n",
    "    x[0] += 2.0"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "If we inspect the types, we see:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 25,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "_ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE (array(float32, 1d, A),)\n",
      "--------------------------------------------------------------------------------\n",
      "# File: <ipython-input-24-25f8fdc40e3e>\n",
      "# --- LINE 1 --- \n",
      "\n",
      "@cuda.jit(void(float32[:]))\n",
      "\n",
      "# --- LINE 2 --- \n",
      "\n",
      "def add_constant(x):\n",
      "\n",
      "    # --- LINE 3 --- \n",
      "    # label 0\n",
      "    #   x = arg(0, name=x)  :: array(float32, 1d, A)\n",
      "    #   $const4.1 = const(int, 0)  :: Literal[int](0)\n",
      "    #   $8binary_subscr.4 = static_getitem(value=x, index=0, index_var=$const4.1)  :: float32\n",
      "    #   $const10.5 = const(float, 2.0)  :: float64\n",
      "    #   $12inplace_add.6 = inplace_binop(fn=<built-in function iadd>, immutable_fn=<built-in function add>, lhs=$8binary_subscr.4, rhs=$const10.5, static_lhs=Undefined, static_rhs=Undefined)  :: float64\n",
      "    #   del $const10.5\n",
      "    #   del $8binary_subscr.4\n",
      "    #   x[0] = $12inplace_add.6\n",
      "    #   del x\n",
      "    #   del $const4.1\n",
      "    #   del $12inplace_add.6\n",
      "    #   $const18.7 = const(NoneType, None)  :: none\n",
      "    #   $20return_value.8 = cast(value=$const18.7)  :: none\n",
      "    #   del $const18.7\n",
      "    #   return $20return_value.8\n",
      "\n",
      "    x[0] += 2.0\n",
      "\n",
      "\n",
      "================================================================================\n"
     ]
    }
   ],
   "source": [
    "add_constant.inspect_types()"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Again the typing of the constant is `float64`, and also the addition of the `float32` and `float64` (`$8binary_subscr.4` plus `$const10.5` stored in `$12inplace_add.6`) results in a `float64`.\n",
    "\n",
    "But does the addition result in a 64-bit operation in the LLVM IR?"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 26,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "source_filename = \"<string>\"\n",
      "target datalayout = \"e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64\"\n",
      "target triple = \"nvptx64-nvidia-cuda\"\n",
      "\n",
      "@\"_ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE__errcode__\" = local_unnamed_addr global i32 0\n",
      "@\"_ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE__tidx__\" = local_unnamed_addr global i32 0\n",
      "@\"_ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE__ctaidx__\" = local_unnamed_addr global i32 0\n",
      "@\"_ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE__tidy__\" = local_unnamed_addr global i32 0\n",
      "@\"_ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE__ctaidy__\" = local_unnamed_addr global i32 0\n",
      "@\"_ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE__tidz__\" = local_unnamed_addr global i32 0\n",
      "@\"_ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE__ctaidz__\" = local_unnamed_addr global i32 0\n",
      "@\"_ZN08NumbaEnv8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE\" = common local_unnamed_addr global i8* null\n",
      "\n",
      "; Function Attrs: norecurse nounwind\n",
      "define void @\"_ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE\"(i8* nocapture readnone %.1, i8* nocapture readnone %.2, i64 %.3, i64 %.4, float* nocapture %.5, i64 %.6, i64 %.7) local_unnamed_addr #0 {\n",
      ".9:\n",
      "  %.12 = alloca i8*, align 8\n",
      "  store i8* null, i8** %.12, align 8\n",
      "  %.17 = call i32 @\"_ZN8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE\"(i8** nonnull %.12, i8* %.1, i8* %.2, i64 %.3, i64 %.4, float* %.5, i64 %.6, i64 %.7)\n",
      "  ret void\n",
      "}\n",
      "\n",
      "; Function Attrs: norecurse nounwind\n",
      "define linkonce_odr i32 @\"_ZN8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE\"(i8** nocapture %.ret, i8* nocapture readnone %arg.x.0, i8* nocapture readnone %arg.x.1, i64 %arg.x.2, i64 %arg.x.3, float* nocapture %arg.x.4, i64 %arg.x.5.0, i64 %arg.x.6.0) local_unnamed_addr #0 {\n",
      "entry:\n",
      "  %.4957 = load float, float* %arg.x.4, align 4\n",
      "  %.99 = fadd float %.4957, 2.000000e+00\n",
      "  store float %.99, float* %arg.x.4, align 4\n",
      "  store i8* null, i8** %.ret, align 8\n",
      "  ret i32 0\n",
      "}\n",
      "\n",
      "attributes #0 = { norecurse nounwind }\n",
      "\n",
      "!nvvm.annotations = !{!0}\n",
      "!nvvmir.version = !{!1}\n",
      "\n",
      "!0 = !{void (i8*, i8*, i64, i64, float*, i64, i64)* @\"_ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE\", !\"kernel\", i32 1}\n",
      "!1 = !{i32 1, i32 2, i32 2, i32 0}\n",
      "\n"
     ]
    }
   ],
   "source": [
    "print(add_constant.inspect_llvm())"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "No! Again the LLVM optimizer has managed to reduce this to a 32-bit operation: `fadd float %.4957, 2.000000e+00`.\n",
    "\n",
    "The PTX corresponds:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 27,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "//\n",
      "// Generated by NVIDIA NVVM Compiler\n",
      "//\n",
      "// Compiler Build ID: CL-27506705\n",
      "// Cuda compilation tools, release 10.2, V10.2.89\n",
      "// Based on LLVM 3.4svn\n",
      "//\n",
      "\n",
      ".version 6.5\n",
      ".target sm_70\n",
      ".address_size 64\n",
      "\n",
      "\t// .globl\t_ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE\n",
      ".visible .global .align 4 .u32 _ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE__errcode__;\n",
      ".visible .global .align 4 .u32 _ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE__tidx__;\n",
      ".visible .global .align 4 .u32 _ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE__ctaidx__;\n",
      ".visible .global .align 4 .u32 _ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE__tidy__;\n",
      ".visible .global .align 4 .u32 _ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE__ctaidy__;\n",
      ".visible .global .align 4 .u32 _ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE__tidz__;\n",
      ".visible .global .align 4 .u32 _ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE__ctaidz__;\n",
      ".common .global .align 8 .u64 _ZN08NumbaEnv8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE;\n",
      "\n",
      ".visible .entry _ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE(\n",
      "\t.param .u64 _ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE_param_0,\n",
      "\t.param .u64 _ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE_param_1,\n",
      "\t.param .u64 _ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE_param_2,\n",
      "\t.param .u64 _ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE_param_3,\n",
      "\t.param .u64 _ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE_param_4,\n",
      "\t.param .u64 _ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE_param_5,\n",
      "\t.param .u64 _ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE_param_6\n",
      ")\n",
      "{\n",
      "\t.reg .f32 \t%f<3>;\n",
      "\t.reg .b64 \t%rd<3>;\n",
      "\n",
      "\n",
      "\tld.param.u64 \t%rd1, [_ZN6cudapy8__main__16add_constant$249E5ArrayIfLi1E1A7mutable7alignedE_param_4];\n",
      "\tcvta.to.global.u64 \t%rd2, %rd1;\n",
      "\tld.global.f32 \t%f1, [%rd2];\n",
      "\tadd.f32 \t%f2, %f1, 0f40000000;\n",
      "\tst.global.f32 \t[%rd2], %f2;\n",
      "\tret;\n",
      "}\n",
      "\n",
      "\n",
      "\u0000\n"
     ]
    }
   ],
   "source": [
    "print(add_constant.inspect_asm())"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "As expected, we see `add.f32 \t%f2, %f1, 0f40000000;`.\n",
    "\n",
    "### Bringing in another addition\n",
    "\n",
    "As well as adding a constant, we'll now add another array element:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 28,
   "metadata": {},
   "outputs": [],
   "source": [
    "@cuda.jit(void(float32[:], float32[:]))\n",
    "def add_constant_2(x, y):\n",
    "    x[0] += y[0] + 2.0"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "We would expect the IR to contain more `float64` operations:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 29,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "_ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE (array(float32, 1d, A), array(float32, 1d, A))\n",
      "--------------------------------------------------------------------------------\n",
      "# File: <ipython-input-28-aa462724f951>\n",
      "# --- LINE 1 --- \n",
      "\n",
      "@cuda.jit(void(float32[:], float32[:]))\n",
      "\n",
      "# --- LINE 2 --- \n",
      "\n",
      "def add_constant_2(x, y):\n",
      "\n",
      "    # --- LINE 3 --- \n",
      "    # label 0\n",
      "    #   x = arg(0, name=x)  :: array(float32, 1d, A)\n",
      "    #   y = arg(1, name=y)  :: array(float32, 1d, A)\n",
      "    #   $const4.1 = const(int, 0)  :: Literal[int](0)\n",
      "    #   $8binary_subscr.4 = static_getitem(value=x, index=0, index_var=$const4.1)  :: float32\n",
      "    #   $const12.6 = const(int, 0)  :: Literal[int](0)\n",
      "    #   $14binary_subscr.7 = static_getitem(value=y, index=0, index_var=$const12.6)  :: float32\n",
      "    #   del y\n",
      "    #   del $const12.6\n",
      "    #   $const16.8 = const(float, 2.0)  :: float64\n",
      "    #   $18binary_add.9 = $14binary_subscr.7 + $const16.8  :: float64\n",
      "    #   del $const16.8\n",
      "    #   del $14binary_subscr.7\n",
      "    #   $20inplace_add.10 = inplace_binop(fn=<built-in function iadd>, immutable_fn=<built-in function add>, lhs=$8binary_subscr.4, rhs=$18binary_add.9, static_lhs=Undefined, static_rhs=Undefined)  :: float64\n",
      "    #   del $8binary_subscr.4\n",
      "    #   del $18binary_add.9\n",
      "    #   x[0] = $20inplace_add.10\n",
      "    #   del x\n",
      "    #   del $const4.1\n",
      "    #   del $20inplace_add.10\n",
      "    #   $const26.11 = const(NoneType, None)  :: none\n",
      "    #   $28return_value.12 = cast(value=$const26.11)  :: none\n",
      "    #   del $const26.11\n",
      "    #   return $28return_value.12\n",
      "\n",
      "    x[0] += y[0] + 2.0\n",
      "\n",
      "\n",
      "================================================================================\n"
     ]
    }
   ],
   "source": [
    "add_constant_2.inspect_types()"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "What happens this time in the LLVM IR? Let's see:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 30,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "source_filename = \"<string>\"\n",
      "target datalayout = \"e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64\"\n",
      "target triple = \"nvptx64-nvidia-cuda\"\n",
      "\n",
      "@\"_ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE__errcode__\" = local_unnamed_addr global i32 0\n",
      "@\"_ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE__tidx__\" = local_unnamed_addr global i32 0\n",
      "@\"_ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE__ctaidx__\" = local_unnamed_addr global i32 0\n",
      "@\"_ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE__tidy__\" = local_unnamed_addr global i32 0\n",
      "@\"_ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE__ctaidy__\" = local_unnamed_addr global i32 0\n",
      "@\"_ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE__tidz__\" = local_unnamed_addr global i32 0\n",
      "@\"_ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE__ctaidz__\" = local_unnamed_addr global i32 0\n",
      "@\"_ZN08NumbaEnv8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE\" = common local_unnamed_addr global i8* null\n",
      "\n",
      "; Function Attrs: norecurse nounwind\n",
      "define void @\"_ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE\"(i8* nocapture readnone %.1, i8* nocapture readnone %.2, i64 %.3, i64 %.4, float* nocapture %.5, i64 %.6, i64 %.7, i8* nocapture readnone %.8, i8* nocapture readnone %.9, i64 %.10, i64 %.11, float* nocapture readonly %.12, i64 %.13, i64 %.14) local_unnamed_addr #0 {\n",
      ".16:\n",
      "  %.21 = alloca i8*, align 8\n",
      "  store i8* null, i8** %.21, align 8\n",
      "  %.28 = call i32 @\"_ZN8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE\"(i8** nonnull %.21, i8* %.1, i8* %.2, i64 %.3, i64 %.4, float* %.5, i64 %.6, i64 %.7, i8* %.8, i8* %.9, i64 %.10, i64 %.11, float* %.12, i64 %.13, i64 %.14)\n",
      "  ret void\n",
      "}\n",
      "\n",
      "; Function Attrs: norecurse nounwind\n",
      "define linkonce_odr i32 @\"_ZN8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE\"(i8** nocapture %.ret, i8* nocapture readnone %arg.x.0, i8* nocapture readnone %arg.x.1, i64 %arg.x.2, i64 %arg.x.3, float* nocapture %arg.x.4, i64 %arg.x.5.0, i64 %arg.x.6.0, i8* nocapture readnone %arg.y.0, i8* nocapture readnone %arg.y.1, i64 %arg.y.2, i64 %arg.y.3, float* nocapture readonly %arg.y.4, i64 %arg.y.5.0, i64 %arg.y.6.0) local_unnamed_addr #0 {\n",
      "entry:\n",
      "  %.61 = load float, float* %arg.x.4, align 4\n",
      "  %.99 = load float, float* %arg.y.4, align 4\n",
      "  %.112 = fpext float %.99 to double\n",
      "  %.113 = fadd double %.112, 2.000000e+00\n",
      "  %.123 = fpext float %.61 to double\n",
      "  %.124 = fadd double %.113, %.123\n",
      "  %.164 = fptrunc double %.124 to float\n",
      "  store float %.164, float* %arg.x.4, align 4\n",
      "  store i8* null, i8** %.ret, align 8\n",
      "  ret i32 0\n",
      "}\n",
      "\n",
      "attributes #0 = { norecurse nounwind }\n",
      "\n",
      "!nvvm.annotations = !{!0}\n",
      "!nvvmir.version = !{!1}\n",
      "\n",
      "!0 = !{void (i8*, i8*, i64, i64, float*, i64, i64, i8*, i8*, i64, i64, float*, i64, i64)* @\"_ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE\", !\"kernel\", i32 1}\n",
      "!1 = !{i32 1, i32 2, i32 2, i32 0}\n",
      "\n"
     ]
    }
   ],
   "source": [
    "print(add_constant_2.inspect_llvm())"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Instead of operations on 32-bit floats, we now see casts (`fpext` / `fptrunc`) between 32- and 64-bit values, and operations on 64-bit values (`fadd double`). This time, the optimizer couldn't save us!\n",
    "\n",
    "NVVM doesn't help us in this case either:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 31,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "//\n",
      "// Generated by NVIDIA NVVM Compiler\n",
      "//\n",
      "// Compiler Build ID: CL-27506705\n",
      "// Cuda compilation tools, release 10.2, V10.2.89\n",
      "// Based on LLVM 3.4svn\n",
      "//\n",
      "\n",
      ".version 6.5\n",
      ".target sm_70\n",
      ".address_size 64\n",
      "\n",
      "\t// .globl\t_ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE\n",
      ".visible .global .align 4 .u32 _ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE__errcode__;\n",
      ".visible .global .align 4 .u32 _ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE__tidx__;\n",
      ".visible .global .align 4 .u32 _ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE__ctaidx__;\n",
      ".visible .global .align 4 .u32 _ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE__tidy__;\n",
      ".visible .global .align 4 .u32 _ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE__ctaidy__;\n",
      ".visible .global .align 4 .u32 _ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE__tidz__;\n",
      ".visible .global .align 4 .u32 _ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE__ctaidz__;\n",
      ".common .global .align 8 .u64 _ZN08NumbaEnv8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE;\n",
      "\n",
      ".visible .entry _ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE(\n",
      "\t.param .u64 _ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE_param_0,\n",
      "\t.param .u64 _ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE_param_1,\n",
      "\t.param .u64 _ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE_param_2,\n",
      "\t.param .u64 _ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE_param_3,\n",
      "\t.param .u64 _ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE_param_4,\n",
      "\t.param .u64 _ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE_param_5,\n",
      "\t.param .u64 _ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE_param_6,\n",
      "\t.param .u64 _ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE_param_7,\n",
      "\t.param .u64 _ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE_param_8,\n",
      "\t.param .u64 _ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE_param_9,\n",
      "\t.param .u64 _ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE_param_10,\n",
      "\t.param .u64 _ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE_param_11,\n",
      "\t.param .u64 _ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE_param_12,\n",
      "\t.param .u64 _ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE_param_13\n",
      ")\n",
      "{\n",
      "\t.reg .f32 \t%f<4>;\n",
      "\t.reg .f64 \t%fd<5>;\n",
      "\t.reg .b64 \t%rd<5>;\n",
      "\n",
      "\n",
      "\tld.param.u64 \t%rd1, [_ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE_param_4];\n",
      "\tld.param.u64 \t%rd2, [_ZN6cudapy8__main__19add_constant_2$2410E5ArrayIfLi1E1A7mutable7alignedE5ArrayIfLi1E1A7mutable7alignedE_param_11];\n",
      "\tcvta.to.global.u64 \t%rd3, %rd2;\n",
      "\tcvta.to.global.u64 \t%rd4, %rd1;\n",
      "\tld.global.f32 \t%f1, [%rd4];\n",
      "\tld.global.f32 \t%f2, [%rd3];\n",
      "\tcvt.f64.f32\t%fd1, %f2;\n",
      "\tadd.f64 \t%fd2, %fd1, 0d4000000000000000;\n",
      "\tcvt.f64.f32\t%fd3, %f1;\n",
      "\tadd.f64 \t%fd4, %fd3, %fd2;\n",
      "\tcvt.rn.f32.f64\t%f3, %fd4;\n",
      "\tst.global.f32 \t[%rd4], %f3;\n",
      "\tret;\n",
      "}\n",
      "\n",
      "\n",
      "\u0000\n"
     ]
    }
   ],
   "source": [
    "print(add_constant_2.inspect_asm())"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Similarly we see casts (e.g. `cvt.f64.f32`) and operations on 64-bit values (e.g. `add.f64`).\n",
    "\n",
    "### Exercise:\n",
    "\n",
    "* Fix the typing of the `add_constant_2` function with an appropriate cast.\n",
    "* Re-run the inspection of the typing, LLVM, and PTX to verify that the width of operations is reduced.\n",
    "\n",
    "## Register usage\n",
    "\n",
    "We can find out the register usage of the kernel from its `regs` attribute:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 32,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "8"
      ]
     },
     "execution_count": 32,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "add_constant_2._func.get().attrs.regs"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "With the original typing, this gives 8 registers on my setup. With the \"corrected\" typing, fewer registers are needed - 6 in my case. In general, reducing the width of operations reduces register usage and can increase occupancy.\n",
    "\n",
    "## Controlling register usage by parameter\n",
    "\n",
    "The `max_registers` keyword argument of the `@cuda.jit` decorator can also be used to limit register usage, which can be helpful if the limit reducing register usage via code changes has been hit.\n",
    "\n",
    "This only has an effect for kernels of a minimum level of complexity - the following is about the size of the simplest example for which it can be seen to take effect:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 33,
   "metadata": {},
   "outputs": [],
   "source": [
    "@cuda.jit\n",
    "def busy_arithmetic(x, y, a):\n",
    "    a = y[0]\n",
    "    b = 2.0\n",
    "    c = y[1] / 6\n",
    "    d = y[2] % 8\n",
    "    e = y[3] * y[4]\n",
    "    for i in range(a):\n",
    "        a += 2\n",
    "        b -= c\n",
    "        e *= d\n",
    "        x[0] += a * b + c * d - e\n",
    "\n",
    "x = np.empty(32, dtype=np.float32)\n",
    "y = np.empty(32, dtype=np.float32)\n",
    "kernel = busy_arithmetic.specialize(x, y, 5)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Note here we used the `specialize()` function of the CUDA-jitted kernel - this can be used to give us a compiled kernel with a typing for a particular set of arguments without launching a kernel. This is convenient when we only want to experiment with a particular typing of a function.\n",
    "\n",
    "Let's examine the register usage of the kernel:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 34,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "36"
      ]
     },
     "execution_count": 34,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "kernel._func.get().attrs.regs"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Now if we redefine the kernel with the `max_registers` keyword argument and inspect the register usage:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 35,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "24"
      ]
     },
     "execution_count": 35,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "@cuda.jit(max_registers=24)\n",
    "def busy_arithmetic_maxreg_24(x, y, a):\n",
    "    a = y[0]\n",
    "    b = 2.0\n",
    "    c = y[1] / 6\n",
    "    d = y[2] % 8\n",
    "    e = y[3] * y[4]\n",
    "    for i in range(a):\n",
    "        a += 2\n",
    "        b -= c\n",
    "        e *= d\n",
    "        x[0] += a * b + c * d - e\n",
    "        \n",
    "kernel_maxreg_24 = busy_arithmetic_maxreg_24.specialize(x, y, 5)\n",
    "kernel_maxreg_24._func.get().attrs.regs"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "We see that the register usage is reduced to the level we requested. However, the `max_registers` kwarg places no commitment on the optimizer, so it may not be honored. For example:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 36,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "24"
      ]
     },
     "execution_count": 36,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "@cuda.jit(max_registers=20)\n",
    "def busy_arithmetic_maxreg_20(x, y, a):\n",
    "    a = y[0]\n",
    "    b = 2.0\n",
    "    c = y[1] / 6\n",
    "    d = y[2] % 8\n",
    "    e = y[3] * y[4]\n",
    "    for i in range(a):\n",
    "        a += 2\n",
    "        b -= c\n",
    "        e *= d\n",
    "        x[0] += a * b + c * d - e\n",
    "        \n",
    "kernel_maxreg_20 = busy_arithmetic_maxreg_20.specialize(x, y, 5)\n",
    "kernel_maxreg_20._func.get().attrs.regs"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "The register usage was reduced, but only to 24, which was the minimum achievable.\n",
    "\n",
    "## Integer arithmetic width\n",
    "\n",
    "Numba strongly prefers using `int64` values for all integer arithmetic. Let's consider an example:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 37,
   "metadata": {},
   "outputs": [],
   "source": [
    "@cuda.jit\n",
    "def index_computation(x):\n",
    "    i = cuda.grid(1)                     # int32\n",
    "\n",
    "    if i < x.shape[0]:                   # x.shape[0] will be int64\n",
    "        for j in range(3):               # range_iter_int64\n",
    "            x[i, j] = (i * 2) + (j * 3)  # int64 computations\n",
    "\n",
    "x = np.zeros((1024, 3), dtype=np.int32)\n",
    "kernel = index_computation.specialize(x)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Now if we inspect the typing:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 38,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "_ZN6cudapy8__main__22index_computation$2414E5ArrayIiLi2E1C7mutable7alignedE (array(int32, 2d, C),)\n",
      "--------------------------------------------------------------------------------\n",
      "# File: <ipython-input-37-7e0a8164e397>\n",
      "# --- LINE 1 --- \n",
      "\n",
      "@cuda.jit\n",
      "\n",
      "# --- LINE 2 --- \n",
      "\n",
      "def index_computation(x):\n",
      "\n",
      "    # --- LINE 3 --- \n",
      "    # label 0\n",
      "    #   x = arg(0, name=x)  :: array(int32, 2d, C)\n",
      "    #   $2load_global.0 = global(cuda: <module 'numba.cuda' from '/home/gmarkall/numbadev/numba/numba/cuda/__init__.py'>)  :: Module(<module 'numba.cuda' from '/home/gmarkall/numbadev/numba/numba/cuda/__init__.py'>)\n",
      "    #   $4load_method.1 = getattr(value=$2load_global.0, attr=grid)  :: Macro(<class 'numba.cuda.cudadecl.Cuda_grid'>)\n",
      "    #   del $4load_method.1\n",
      "    #   del $2load_global.0\n",
      "    #   $const6.2 = const(int, 1)  :: Literal[int](1)\n",
      "    #   $8call_method.3 = call ptx.grid.1d($const6.2, func=ptx.grid.1d, args=[Var($const6.2, <ipython-input-37-7e0a8164e397>:3)], kws=(), vararg=None)  :: (int64,) -> int32\n",
      "    #   del $const6.2\n",
      "    #   i = $8call_method.3  :: int32\n",
      "    #   del $8call_method.3\n",
      "\n",
      "    i = cuda.grid(1)                     # int32\n",
      "\n",
      "# --- LINE 4 --- \n",
      "\n",
      "\n",
      "\n",
      "    # --- LINE 5 --- \n",
      "    #   $16load_attr.6 = getattr(value=x, attr=shape)  :: UniTuple(int64 x 2)\n",
      "    #   $const18.7 = const(int, 0)  :: Literal[int](0)\n",
      "    #   $20binary_subscr.8 = static_getitem(value=$16load_attr.6, index=0, index_var=$const18.7)  :: int64\n",
      "    #   del $const18.7\n",
      "    #   del $16load_attr.6\n",
      "    #   $22compare_op.9 = i < $20binary_subscr.8  :: bool\n",
      "    #   del $20binary_subscr.8\n",
      "    #   branch $22compare_op.9, 26, 64\n",
      "\n",
      "    if i < x.shape[0]:                   # x.shape[0] will be int64\n",
      "\n",
      "        # --- LINE 6 --- \n",
      "        # label 26\n",
      "        #   del $22compare_op.9\n",
      "        #   $26load_global.0 = global(range: <class 'range'>)  :: Function(<class 'range'>)\n",
      "        #   $const28.1 = const(int, 3)  :: Literal[int](3)\n",
      "        #   $30call_function.2 = call $26load_global.0($const28.1, func=$26load_global.0, args=[Var($const28.1, <ipython-input-37-7e0a8164e397>:6)], kws=(), vararg=None)  :: (int64,) -> range_state_int64\n",
      "        #   del $const28.1\n",
      "        #   del $26load_global.0\n",
      "        #   $32get_iter.3 = getiter(value=$30call_function.2)  :: range_iter_int64\n",
      "        #   del $30call_function.2\n",
      "        #   $phi34.0 = $32get_iter.3  :: range_iter_int64\n",
      "        #   del $32get_iter.3\n",
      "        #   jump 34\n",
      "        # label 34\n",
      "        #   $34for_iter.1 = iternext(value=$phi34.0)  :: pair<int64, bool>\n",
      "        #   $34for_iter.2 = pair_first(value=$34for_iter.1)  :: int64\n",
      "        #   $34for_iter.3 = pair_second(value=$34for_iter.1)  :: bool\n",
      "        #   del $34for_iter.1\n",
      "        #   $phi36.1 = $34for_iter.2  :: int64\n",
      "        #   del $34for_iter.2\n",
      "        #   branch $34for_iter.3, 36, 64\n",
      "        # label 36\n",
      "        #   del $34for_iter.3\n",
      "        #   j = $phi36.1  :: int64\n",
      "        #   del $phi36.1\n",
      "\n",
      "        for j in range(3):               # range_iter_int64\n",
      "\n",
      "            # --- LINE 7 --- \n",
      "            #   $const40.3 = const(int, 2)  :: Literal[int](2)\n",
      "            #   $42binary_multiply.4 = i * $const40.3  :: int64\n",
      "            #   del $const40.3\n",
      "            #   $const46.6 = const(int, 3)  :: Literal[int](3)\n",
      "            #   $48binary_multiply.7 = j * $const46.6  :: int64\n",
      "            #   del $const46.6\n",
      "            #   $50binary_add.8 = $42binary_multiply.4 + $48binary_multiply.7  :: int64\n",
      "            #   del $48binary_multiply.7\n",
      "            #   del $42binary_multiply.4\n",
      "            #   $58build_tuple.12 = build_tuple(items=[Var(i, <ipython-input-37-7e0a8164e397>:3), Var(j, <ipython-input-37-7e0a8164e397>:6)])  :: Tuple(int32, int64)\n",
      "            #   del j\n",
      "            #   x[$58build_tuple.12] = $50binary_add.8  :: (array(int32, 2d, C), UniTuple(int64 x 2), int64) -> none\n",
      "            #   del $58build_tuple.12\n",
      "            #   del $50binary_add.8\n",
      "            #   jump 34\n",
      "            # label 64\n",
      "            #   del x\n",
      "            #   del i\n",
      "            #   del $phi36.1\n",
      "            #   del $phi34.0\n",
      "            #   del $34for_iter.3\n",
      "            #   del $22compare_op.9\n",
      "            #   $const64.0 = const(NoneType, None)  :: none\n",
      "            #   $66return_value.1 = cast(value=$const64.0)  :: none\n",
      "            #   del $const64.0\n",
      "            #   return $66return_value.1\n",
      "\n",
      "            x[i, j] = (i * 2) + (j * 3)  # int64 computations\n",
      "\n",
      "\n",
      "================================================================================\n"
     ]
    }
   ],
   "source": [
    "kernel.inspect_types()"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "We see that most of the arithmetic happens using `int64` values, and the range iterates over `int64` (the `range_iter_int64` type).\n",
    "\n",
    "We can attempt to reduce the width of arithmetic operations using casts, but it requires a lot of casts:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 39,
   "metadata": {},
   "outputs": [],
   "source": [
    "from numba import int32\n",
    "\n",
    "@cuda.jit\n",
    "def index_computation_int32(x):\n",
    "    i = cuda.grid(1)                     # int32\n",
    "\n",
    "    if i < int32(x.shape[0]):            # Attempt to compare using int32 arithmetic\n",
    "        for j in range(int32(3)):        # Force iteration over int32 - a range_iter_int32\n",
    "            x[i, j] = int32(int32(int32(i) * int32(2))\n",
    "                            + int32(int32(j) * int32(3)))\n",
    "                                         # Attempt to make all constants and operations int32\n",
    "\n",
    "kernel_int32 = index_computation_int32.specialize(x)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "If we have been successful, we should see a reduced register usage for the `index_computation_int32` kernel:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 40,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "12"
      ]
     },
     "execution_count": 40,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "kernel._func.get().attrs.regs"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 41,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "14"
      ]
     },
     "execution_count": 41,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "kernel_int32._func.get().attrs.regs"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "We have actually made things worse! Often it is better to try not to reduce the width of `int64` operations, because it results in a mix of `int32` and `int64` values, which ends up requiring more registers.\n",
    "\n",
    "Exercises:\n",
    "\n",
    "* Inspect the IR, LLVM, and PTX to see where `int64` computations remain in `index_computation_int32`.\n",
    "\n",
    "# Summary\n",
    "\n",
    "Throughout the course of this notebook, we have:\n",
    "\n",
    "* Seen how to use `inspect_types()` to view the typing of jitted functions\n",
    "* Examined *phi nodes* and looked at the unification of types at phi nodes.\n",
    "* Seen how calls with different argument types result in different specialisations of a function, that have different typings.\n",
    "* Examined typing errors:\n",
    "  * Unification failures, and how to determine what failed to unify\n",
    "  * Use of a variable with inconsistent typing throughout the function (e.g. 1D array vs. 2D array)\n",
    "  * Use of unsupported functions, or functions implemented using unsupported functions in the CUDA target.\n",
    "* Seen an example of branch elimination, and how it sometimes succeeds in allowing typings with arguments that could otherwise have resulted in unification errors.\n",
    "* Looked at CUDA-specific issues, mainly related to register usage:\n",
    "  * When widening unification occurs, and how to prevent it.\n",
    "  * When widening arithmetic occurs, and how to avoid it for floating point types.\n",
    "  * How integer arithmetic strongly prefers `int64`, and how it can be counterproductive to try to reduce it to `int32` and narrower types.\n",
    "* Seen how to control register usage using the `max_registers` keyword argument."
   ]
  }
 ],
 "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.8.2"
  }
 },
 "nbformat": 4,
 "nbformat_minor": 4
}