Streaming_and_Visual_Profiling.ipynb

{
 "cells": [
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "<h1><div align=\"center\">Asynchronous Streaming, and Visual Profiling with CUDA C/C++</div></h1>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "![CUDA](./images/CUDA_Logo.jpg)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "The CUDA toolkit ships with the **Nsight Systems**, a powerful GUI application to support the development of accelerated CUDA applications. Nsight Systems generates a graphical timeline of an accelerated application, with detailed information about CUDA API calls, kernel execution, memory activity, and the use of **CUDA streams**.\n",
    "\n",
    "In this lab, you will be using the Nsight Systems timeline to guide you in optimizing accelerated applications. Additionally, you will learn some intermediate CUDA programming techniques to support your work: **unmanaged memory allocation and migration**; **pinning**, or **page-locking** host memory; and **non-default concurrent CUDA streams**.\n",
    "\n",
    "At the end of this lab, you will be presented with an assessment, to accelerate and optimize a simple n-body particle simulator, which will allow you to demonstrate the skills you have developed during this course. Those of you who are able to accelerate the simulator while maintaining its correctness, will be granted a certification as proof of your competency."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## Prerequisites\n",
    "\n",
    "To get the most out of this lab you should already be able to:\n",
    "\n",
    "- Write, compile, and run C/C++ programs that both call CPU functions and launch GPU kernels.\n",
    "- Control parallel thread hierarchy using execution configuration.\n",
    "- Refactor serial loops to execute their iterations in parallel on a GPU.\n",
    "- Allocate and free CUDA Unified Memory.\n",
    "- Understand the behavior of Unified Memory with regard to page faulting and data migrations.\n",
    "- Use asynchronous memory prefetching to reduce page faults and data migrations."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## Objectives\n",
    "\n",
    "By the time you complete this lab you will be able to:\n",
    "\n",
    "- Use **Nsight Systems** to visually profile the timeline of GPU-accelerated CUDA applications.\n",
    "- Use Nsight Systems to identify, and exploit, optimization opportunities in GPU-accelerated CUDA applications.\n",
    "- Utilize CUDA streams for concurrent kernel execution in accelerated applications.\n",
    "- (**Optional Advanced Content**) Use manual device memory allocation, including allocating pinned memory, in order to asynchronously transfer data in concurrent CUDA streams."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## Running Nsight Systems"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "For this interactive lab environment, we have set up a remote desktop you can access from your browser, where you will be able to launch and use Nsight Systems.\n",
    "\n",
    "You will begin by creating a report file for an already-existing vector addition program, after which you will be walked through a series of steps to open this report file in Nsight Systems, and to make the visual experience nice."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Generate Report File\n",
    "\n",
    "[`01-vector-add.cu`](../edit/01-vector-add/01-vector-add.cu) (<-------- click on these links to source files to edit them in the browser) contains a working, accelerated, vector addition application. Use the code execution cell directly below (you can execute it, and any of the code execution cells in this lab by `CTRL` + clicking it) to compile and run it. You should see a message printed that indicates it was successful."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 1,
   "metadata": {
    "scrolled": true
   },
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "Success! All values calculated correctly.\r\n"
     ]
    }
   ],
   "source": [
    "!nvcc -o vector-add-no-prefetch 01-vector-add/01-vector-add.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Next, use `nsys profile --stats=true` to create a report file that you will be able to open in the Nsight Systems visual profiler. Here we use the `-o` flag to give the report file a memorable name:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 2,
   "metadata": {
    "scrolled": true
   },
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "Warning: LBR backtrace method is not supported on this platform. DWARF backtrace method will be used.\n",
      "WARNING: The command line includes a target application therefore the CPU context-switch scope has been set to process-tree.\n",
      "Collecting data...\n",
      "Success! All values calculated correctly.\n",
      "Processing events...\n",
      "Saving temporary \"/tmp/nsys-report-6a5a-6eb4-e9e5-8890.qdstrm\" file to disk...\n",
      "\n",
      "Creating final output files...\n",
      "Processing [==============================================================100%]\n",
      "Saved report file to \"/tmp/nsys-report-6a5a-6eb4-e9e5-8890.qdrep\"\n",
      "Exporting 10235 events: [=================================================100%]\n",
      "\n",
      "Exported successfully to\n",
      "/tmp/nsys-report-6a5a-6eb4-e9e5-8890.sqlite\n",
      "\n",
      "\n",
      "CUDA API Statistics:\n",
      "\n",
      " Time(%)  Total Time (ns)  Num Calls    Average     Minimum    Maximum           Name         \n",
      " -------  ---------------  ---------  -----------  ---------  ---------  ---------------------\n",
      "    61.9        246083765          3   82027921.7      17255  246026906  cudaMallocManaged    \n",
      "    32.5        129021842          1  129021842.0  129021842  129021842  cudaDeviceSynchronize\n",
      "     5.6         22384952          3    7461650.7    6701513    8922765  cudaFree             \n",
      "     0.0            45634          1      45634.0      45634      45634  cudaLaunchKernel     \n",
      "\n",
      "\n",
      "\n",
      "CUDA Kernel Statistics:\n",
      "\n",
      " Time(%)  Total Time (ns)  Instances    Average     Minimum    Maximum                      Name                    \n",
      " -------  ---------------  ---------  -----------  ---------  ---------  -------------------------------------------\n",
      "   100.0        129012278          1  129012278.0  129012278  129012278  addVectorsInto(float*, float*, float*, int)\n",
      "\n",
      "\n",
      "\n",
      "CUDA Memory Operation Statistics (by time):\n",
      "\n",
      " Time(%)  Total Time (ns)  Operations  Average  Minimum  Maximum              Operation            \n",
      " -------  ---------------  ----------  -------  -------  -------  ---------------------------------\n",
      "    78.8         78911017        8314   9491.3     2143   128189  [CUDA Unified Memory memcpy HtoD]\n",
      "    21.2         21192428         768  27594.3     1567   159742  [CUDA Unified Memory memcpy DtoH]\n",
      "\n",
      "\n",
      "\n",
      "CUDA Memory Operation Statistics (by size in KiB):\n",
      "\n",
      "   Total     Operations  Average  Minimum  Maximum               Operation            \n",
      " ----------  ----------  -------  -------  --------  ---------------------------------\n",
      " 393216.000        8314   47.296    4.000   764.000  [CUDA Unified Memory memcpy HtoD]\n",
      " 131072.000         768  170.667    4.000  1020.000  [CUDA Unified Memory memcpy DtoH]\n",
      "\n",
      "\n",
      "\n",
      "Operating System Runtime API Statistics:\n",
      "\n",
      " Time(%)  Total Time (ns)  Num Calls   Average    Minimum   Maximum        Name     \n",
      " -------  ---------------  ---------  ----------  -------  ---------  --------------\n",
      "    83.3       1330490805         72  18479039.0    32582  100127978  poll          \n",
      "     8.8        140558156         63   2231081.8    15224   20538800  sem_timedwait \n",
      "     6.3         99891506        676    147768.5     1070   18788038  ioctl         \n",
      "     1.6         24916487         94    265069.0     1407    8868714  mmap          \n",
      "     0.1          1469485         82     17920.5     4713      28255  open64        \n",
      "     0.0           214257          3     71419.0    69421      74578  fgets         \n",
      "     0.0           154636          4     38659.0    35417      46043  pthread_create\n",
      "     0.0           117941         25      4717.6     1563      25048  fopen         \n",
      "     0.0            77800         11      7072.7     4800      11352  write         \n",
      "     0.0            41449         11      3768.1     1443       5929  munmap        \n",
      "     0.0            28638          5      5727.6     3301       8493  open          \n",
      "     0.0            26997         18      1499.8     1113       3778  fclose        \n",
      "     0.0            22073          6      3678.8     1091       9657  fgetc         \n",
      "     0.0            21042         13      1618.6     1047       2401  read          \n",
      "     0.0            14895          3      4965.0     1663       7476  fread         \n",
      "     0.0            10453          7      1493.3     1002       3690  fcntl         \n",
      "     0.0             9649          2      4824.5     4477       5172  socket        \n",
      "     0.0             7203          1      7203.0     7203       7203  pipe2         \n",
      "     0.0             6243          1      6243.0     6243       6243  connect       \n",
      "     0.0             2535          1      2535.0     2535       2535  bind          \n",
      "     0.0             1461          1      1461.0     1461       1461  listen        \n",
      "\n",
      "Report file moved to \"/dli/task/vector-add-no-prefetch-report.qdrep\"\n",
      "Report file moved to \"/dli/task/vector-add-no-prefetch-report.sqlite\"\n",
      "\n"
     ]
    }
   ],
   "source": [
    "!nsys profile --stats=true -o vector-add-no-prefetch-report ./vector-add-no-prefetch"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Open the Remote Desktop"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Run the next cell to generate a link to the remote desktop. Then, read the instructions that follow in the notebook."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 3,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "application/javascript": [
       "var port = ((window.location.port == 80) ? \"\" : (\":\"+window.location.port));\n",
       "var url = 'http://' + window.location.hostname + port + '/nsight/vnc.html?resize=scale';\n",
       "let a = document.createElement('a');\n",
       "a.setAttribute('href', url)\n",
       "a.setAttribute('target', '_blank')\n",
       "a.innerText = 'Click to open remote desktop'\n",
       "element.append(a);\n"
      ],
      "text/plain": [
       "<IPython.core.display.Javascript object>"
      ]
     },
     "metadata": {},
     "output_type": "display_data"
    }
   ],
   "source": [
    "%%js\n",
    "var port = ((window.location.port == 80) ? \"\" : (\":\"+window.location.port));\n",
    "var url = 'http://' + window.location.hostname + port + '/nsight/vnc.html?resize=scale';\n",
    "let a = document.createElement('a');\n",
    "a.setAttribute('href', url)\n",
    "a.setAttribute('target', '_blank')\n",
    "a.innerText = 'Click to open remote desktop'\n",
    "element.append(a);"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "After clicking the _Connect_ button you will be asked for a password, which is `nvidia`."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Open Nsight Systems"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "To open Nsight Systems, double-click the \"NVIDIA Nsight Systems\" icon on the remote desktop."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "![open nsight](images/open-nsight-sys.png)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Enable Usage Reporting"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "When prompted, click \"Yes\" to enable usage reporting:"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "![enable usage](images/enable_usage.png)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Select GPU Rows on Top"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "When prompted, select _GPU Rows on Top_ and then click _Okay_."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "![gpu)_rows_on_top](images/gpu_on_top.png)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Open the Report File"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Open this report file by visiting _File_ -> _Open_ from the Nsight Systems menu and select `vector-add-no-prefetch-report.qdrep`:"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "![open-report](images/open-report.png)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Ignore Warnings/Errors"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "You can close and ignore any warnings or errors you see, which are just a result of our particular remote desktop environment:"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "![ignore errors](images/ignore-error.png)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Make More Room for the Timelines"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "To make your experience nicer, full-screen the profiler, close the _Project Explorer_ and hide the *Events View*:"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "![make nice](images/make-nice.png)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Your screen should now look like this:"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "![now nice](images/now-nice.png)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Expand the CUDA Unified Memory Timelines"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Next, expand the _CUDA_ -> _Unified memory_ and _Context_ timelines, and close the _Threads_ timelines:"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "![open memory](images/open-memory.png)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Observe Many Memory Transfers"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "From a glance you can see that your application is taking about 1 second to run, and that also, during the time when the `addVectorsInto` kernel is running, that there is a lot of UM memory activity:"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "![memory and kernel](images/memory-and-kernel.png)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Zoom into the memory timelines to see more clearly all the small memory transfers being caused by the on-demand memory page faults. A couple tips:\n",
    "\n",
    "1. You can zoom in and out at any point of the timeline by holding `CTRL` while scrolling your mouse/trackpad\n",
    "2. You can zoom into any section by click + dragging a rectangle around it, and then selecting _Zoom in_"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Here's an example of zooming in to see the many small memory transfers:"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "![many transfers](images/many-transfers.png)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## Comparing Code Refactors Iteratively with Nsight Systems"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Now that you have Nsight Systems up and running and are comfortable moving around the timelines, you will be profiling a series of programs that were iteratively improved using techniques already familiar to you. Each time you profile, information in the timeline will give information supporting how you should next modify your code. Doing this will further increase your understanding of how various CUDA programming techniques affect application performance."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Exercise: Compare the Timelines of Prefetching vs. Non-Prefetching\n",
    "\n",
    "[`01-vector-add-prefetch-solution.cu`](../edit/01-vector-add/solutions/01-vector-add-prefetch-solution.cu) refactors the vector addition application from above so that the 3 vectors needed by its `addVectorsInto` kernel are asynchronously prefetched to the active GPU device prior to launching the kernel (using [`cudaMemPrefetchAsync`](http://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html#group__CUDART__MEMORY_1ge8dc9199943d421bc8bc7f473df12e42)). Open the source code and identify where in the application these changes were made.\n",
    "\n",
    "After reviewing the changes, compile and run the refactored application using the code execution cell directly below. You should see its success message printed."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 4,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "Success! All values calculated correctly.\r\n"
     ]
    }
   ],
   "source": [
    "!nvcc -o vector-add-prefetch 01-vector-add/solutions/01-vector-add-prefetch-solution.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Now create a report file for this version of the application:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 5,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "Warning: LBR backtrace method is not supported on this platform. DWARF backtrace method will be used.\n",
      "WARNING: The command line includes a target application therefore the CPU context-switch scope has been set to process-tree.\n",
      "Collecting data...\n",
      "Success! All values calculated correctly.\n",
      "Processing events...\n",
      "Saving temporary \"/tmp/nsys-report-02b7-473d-4ff0-86b9.qdstrm\" file to disk...\n",
      "\n",
      "Creating final output files...\n",
      "Processing [==============================================================100%]\n",
      "Saved report file to \"/tmp/nsys-report-02b7-473d-4ff0-86b9.qdrep\"\n",
      "Exporting 2126 events: [==================================================100%]\n",
      "\n",
      "Exported successfully to\n",
      "/tmp/nsys-report-02b7-473d-4ff0-86b9.sqlite\n",
      "\n",
      "\n",
      "CUDA API Statistics:\n",
      "\n",
      " Time(%)  Total Time (ns)  Num Calls   Average    Minimum    Maximum           Name         \n",
      " -------  ---------------  ---------  ----------  --------  ---------  ---------------------\n",
      "    73.9        260466014          3  86822004.7     18188  260391147  cudaMallocManaged    \n",
      "    16.7         58973167          1  58973167.0  58973167   58973167  cudaDeviceSynchronize\n",
      "     6.4         22708163          3   7569387.7   6837370    8952196  cudaFree             \n",
      "     2.9         10368727          3   3456242.3      6106   10236373  cudaMemPrefetchAsync \n",
      "     0.0            33007          1     33007.0     33007      33007  cudaLaunchKernel     \n",
      "\n",
      "\n",
      "\n",
      "CUDA Kernel Statistics:\n",
      "\n",
      " Time(%)  Total Time (ns)  Instances   Average   Minimum  Maximum                     Name                    \n",
      " -------  ---------------  ---------  ---------  -------  -------  -------------------------------------------\n",
      "   100.0          1697611          1  1697611.0  1697611  1697611  addVectorsInto(float*, float*, float*, int)\n",
      "\n",
      "\n",
      "\n",
      "CUDA Memory Operation Statistics (by time):\n",
      "\n",
      " Time(%)  Total Time (ns)  Operations  Average   Minimum  Maximum              Operation            \n",
      " -------  ---------------  ----------  --------  -------  -------  ---------------------------------\n",
      "    75.7         65658107         192  341969.3   339868   344220  [CUDA Unified Memory memcpy HtoD]\n",
      "    24.3         21105829         768   27481.5     1631   160062  [CUDA Unified Memory memcpy DtoH]\n",
      "\n",
      "\n",
      "\n",
      "CUDA Memory Operation Statistics (by size in KiB):\n",
      "\n",
      "   Total     Operations  Average   Minimum   Maximum               Operation            \n",
      " ----------  ----------  --------  --------  --------  ---------------------------------\n",
      " 393216.000         192  2048.000  2048.000  2048.000  [CUDA Unified Memory memcpy HtoD]\n",
      " 131072.000         768   170.667     4.000  1020.000  [CUDA Unified Memory memcpy DtoH]\n",
      "\n",
      "\n",
      "\n",
      "Operating System Runtime API Statistics:\n",
      "\n",
      " Time(%)  Total Time (ns)  Num Calls   Average    Minimum   Maximum        Name     \n",
      " -------  ---------------  ---------  ----------  -------  ---------  --------------\n",
      "    79.8       1274388745         69  18469402.1     2354  100130974  poll          \n",
      "     9.9        158140955        681    232218.7     1039   22553208  ioctl         \n",
      "     7.8        125032904         57   2193559.7    13781   20810303  sem_timedwait \n",
      "     1.6         25390711         94    270113.9     1488    8895883  mmap          \n",
      "     0.8         12179618          2   6089809.0    41011   12138607  sem_wait      \n",
      "     0.1          1668057         82     20342.2     4819      40370  open64        \n",
      "     0.0           216883          3     72294.3    70550      74923  fgets         \n",
      "     0.0           199047          5     39809.4    32239      53739  pthread_create\n",
      "     0.0           124619         25      4984.8     1740      19616  fopen         \n",
      "     0.0           112647         12      9387.3     4251      13515  write         \n",
      "     0.0            50935         11      4630.5     1828      13674  munmap        \n",
      "     0.0            28969          5      5793.8     3550       9457  open          \n",
      "     0.0            28270         18      1570.6     1104       4745  fclose        \n",
      "     0.0            23903         17      1406.1     1117       3800  fcntl         \n",
      "     0.0            21582         13      1660.2     1072       2530  read          \n",
      "     0.0            21571          6      3595.2     1104       9659  fgetc         \n",
      "     0.0            13691          1     13691.0    13691      13691  pipe2         \n",
      "     0.0            10171          2      5085.5     4670       5501  socket        \n",
      "     0.0             9623          3      3207.7     2083       4027  fread         \n",
      "     0.0             5704          1      5704.0     5704       5704  connect       \n",
      "     0.0             2150          1      2150.0     2150       2150  bind          \n",
      "     0.0             1550          1      1550.0     1550       1550  listen        \n",
      "\n",
      "Report file moved to \"/dli/task/vector-add-prefetch-report.qdrep\"\n",
      "Report file moved to \"/dli/task/vector-add-prefetch-report.sqlite\"\n",
      "\n"
     ]
    }
   ],
   "source": [
    "!nsys profile --stats=true -o vector-add-prefetch-report ./vector-add-prefetch"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Open the report in Nsight Systems, leaving the previous report open for comparison.\n",
    "\n",
    "- How does the execution time compare to that of the `addVectorsInto` kernel prior to adding asynchronous prefetching?\n",
    "- Locate `cudaMemPrefetchAsync` in the *CUDA API* section of the timeline.\n",
    "- How have the memory transfers changed?"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "\n",
    "### Exercise: Profile Refactor with Launch Init in Kernel\n",
    "\n",
    "In the previous iteration of the vector addition application, the vector data is being initialized on the CPU, and therefore needs to be migrated to the GPU before the `addVectorsInto` kernel can operate on it.\n",
    "\n",
    "The next iteration of the application, [01-init-kernel-solution.cu](../edit/02-init-kernel/solutions/01-init-kernel-solution.cu), the application has been refactored to initialize the data in parallel on the GPU.\n",
    "\n",
    "Since the initialization now takes place on the GPU, prefetching has been done prior to initialization, rather than prior to the vector addition work. Review the source code to identify where these changes have been made.\n",
    "\n",
    "After reviewing the changes, compile and run the refactored application using the code execution cell directly below. You should see its success message printed."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 6,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "Success! All values calculated correctly.\r\n"
     ]
    }
   ],
   "source": [
    "!nvcc -o init-kernel 02-init-kernel/solutions/01-init-kernel-solution.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Now create a report file for this version of the application:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 7,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "Warning: LBR backtrace method is not supported on this platform. DWARF backtrace method will be used.\n",
      "WARNING: The command line includes a target application therefore the CPU context-switch scope has been set to process-tree.\n",
      "Collecting data...\n",
      "Success! All values calculated correctly.\n",
      "Processing events...\n",
      "Saving temporary \"/tmp/nsys-report-5f32-ace4-b6e6-efee.qdstrm\" file to disk...\n",
      "\n",
      "Creating final output files...\n",
      "Processing [==============================================================100%]\n",
      "Saved report file to \"/tmp/nsys-report-5f32-ace4-b6e6-efee.qdrep\"\n",
      "Exporting 1859 events: [==================================================100%]\n",
      "\n",
      "Exported successfully to\n",
      "/tmp/nsys-report-5f32-ace4-b6e6-efee.sqlite\n",
      "\n",
      "\n",
      "CUDA API Statistics:\n",
      "\n",
      " Time(%)  Total Time (ns)  Num Calls   Average    Minimum   Maximum           Name         \n",
      " -------  ---------------  ---------  ----------  -------  ---------  ---------------------\n",
      "    91.5        252124091          3  84041363.7    22287  252048304  cudaMallocManaged    \n",
      "     6.4         17516020          3   5838673.3   829781   15599715  cudaFree             \n",
      "     1.3          3558920          1   3558920.0  3558920    3558920  cudaDeviceSynchronize\n",
      "     0.8          2160427          3    720142.3   689193     759977  cudaMemPrefetchAsync \n",
      "     0.0            52789          4     13197.3     5045      35027  cudaLaunchKernel     \n",
      "\n",
      "\n",
      "\n",
      "CUDA Kernel Statistics:\n",
      "\n",
      " Time(%)  Total Time (ns)  Instances   Average   Minimum  Maximum                     Name                    \n",
      " -------  ---------------  ---------  ---------  -------  -------  -------------------------------------------\n",
      "    52.3          1867305          3   622435.0   617145   627960  initWith(float, float*, int)               \n",
      "    47.7          1704011          1  1704011.0  1704011  1704011  addVectorsInto(float*, float*, float*, int)\n",
      "\n",
      "\n",
      "\n",
      "CUDA Memory Operation Statistics (by time):\n",
      "\n",
      " Time(%)  Total Time (ns)  Operations  Average  Minimum  Maximum              Operation            \n",
      " -------  ---------------  ----------  -------  -------  -------  ---------------------------------\n",
      "   100.0         21213610         768  27621.9     1630   160062  [CUDA Unified Memory memcpy DtoH]\n",
      "\n",
      "\n",
      "\n",
      "CUDA Memory Operation Statistics (by size in KiB):\n",
      "\n",
      "   Total     Operations  Average  Minimum  Maximum               Operation            \n",
      " ----------  ----------  -------  -------  --------  ---------------------------------\n",
      " 131072.000         768  170.667    4.000  1020.000  [CUDA Unified Memory memcpy DtoH]\n",
      "\n",
      "\n",
      "\n",
      "Operating System Runtime API Statistics:\n",
      "\n",
      " Time(%)  Total Time (ns)  Num Calls   Average    Minimum   Maximum        Name     \n",
      " -------  ---------------  ---------  ----------  -------  ---------  --------------\n",
      "    74.2        572547650         32  17892114.1    24139  100131583  poll          \n",
      "    14.1        108530709        679    159839.0     1024   18504371  ioctl         \n",
      "     8.8         68114580         27   2522762.2    21853   20563888  sem_timedwait \n",
      "     2.6         20412497         94    217154.2     1420   15532158  mmap          \n",
      "     0.2          1543347         82     18821.3     5164      34251  open64        \n",
      "     0.0           215139          3     71713.0    69654      75122  fgets         \n",
      "     0.0           159600          4     39900.0    34073      45112  pthread_create\n",
      "     0.0           132350         25      5294.0     1558      25807  fopen         \n",
      "     0.0            82525         11      7502.3     4169      13477  write         \n",
      "     0.0            45493         11      4135.7     1508       8834  munmap        \n",
      "     0.0            33023          5      6604.6     4063       9975  open          \n",
      "     0.0            28753         18      1597.4     1081       5567  fclose        \n",
      "     0.0            28142         13      2164.8     1317       4375  read          \n",
      "     0.0            23888          6      3981.3     1045      10151  fgetc         \n",
      "     0.0            21889          4      5472.3     2005      12049  fread         \n",
      "     0.0            14733          8      1841.6     1063       6533  fcntl         \n",
      "     0.0            13630          2      6815.0     6462       7168  socket        \n",
      "     0.0             8312          1      8312.0     8312       8312  pipe2         \n",
      "     0.0             7776          1      7776.0     7776       7776  connect       \n",
      "     0.0             3028          1      3028.0     3028       3028  bind          \n",
      "     0.0             1572          1      1572.0     1572       1572  listen        \n",
      "\n",
      "Report file moved to \"/dli/task/init-kernel-report.qdrep\"\n",
      "Report file moved to \"/dli/task/init-kernel-report.sqlite\"\n",
      "\n"
     ]
    }
   ],
   "source": [
    "!nsys profile --stats=true -o init-kernel-report ./init-kernel"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Open the new report file in Nsight Systems and do the following:\n",
    "\n",
    "- Compare the application and `addVectorsInto` run times to the previous version of the application, how did they change?\n",
    "- Look at the *Kernels* section of the timeline. Which of the two kernels (`addVectorsInto` and the initialization kernel) is taking up the majority of the time on the GPU?\n",
    "- Which of the following does your application contain?\n",
    "  - Data Migration (HtoD)\n",
    "  - Data Migration (DtoH)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Exercise: Profile Refactor with Asynchronous Prefetch Back to the Host\n",
    "\n",
    "Currently, the vector addition application verifies the work of the vector addition kernel on the host. The next refactor of the application, [01-prefetch-check-solution.cu](../edit/04-prefetch-check/solutions/01-prefetch-check-solution.cu), asynchronously prefetches the data back to the host for verification.\n",
    "\n",
    "After reviewing the changes, compile and run the refactored application using the code execution cell directly below. You should see its success message printed."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 8,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "Success! All values calculated correctly.\r\n"
     ]
    }
   ],
   "source": [
    "!nvcc -o prefetch-to-host 04-prefetch-check/solutions/01-prefetch-check-solution.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Now create a report file for this version of the application:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 9,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "Warning: LBR backtrace method is not supported on this platform. DWARF backtrace method will be used.\n",
      "WARNING: The command line includes a target application therefore the CPU context-switch scope has been set to process-tree.\n",
      "Collecting data...\n",
      "Success! All values calculated correctly.\n",
      "Processing events...\n",
      "Saving temporary \"/tmp/nsys-report-ccc2-31d8-5adf-38a3.qdstrm\" file to disk...\n",
      "\n",
      "Creating final output files...\n",
      "Processing [==============================================================100%]\n",
      "Saved report file to \"/tmp/nsys-report-ccc2-31d8-5adf-38a3.qdrep\"\n",
      "Exporting 1158 events: [==================================================100%]\n",
      "\n",
      "Exported successfully to\n",
      "/tmp/nsys-report-ccc2-31d8-5adf-38a3.sqlite\n",
      "\n",
      "\n",
      "CUDA API Statistics:\n",
      "\n",
      " Time(%)  Total Time (ns)  Num Calls   Average    Minimum   Maximum           Name         \n",
      " -------  ---------------  ---------  ----------  -------  ---------  ---------------------\n",
      "    81.2        252404214          3  84134738.0    18510  252336626  cudaMallocManaged    \n",
      "    14.1         43878374          4  10969593.5   663019   41769636  cudaMemPrefetchAsync \n",
      "     3.4         10660288          3   3553429.3   825844    8735064  cudaFree             \n",
      "     1.2          3720615          1   3720615.0  3720615    3720615  cudaDeviceSynchronize\n",
      "     0.0            47051          4     11762.8     4583      30360  cudaLaunchKernel     \n",
      "\n",
      "\n",
      "\n",
      "CUDA Kernel Statistics:\n",
      "\n",
      " Time(%)  Total Time (ns)  Instances   Average   Minimum  Maximum                     Name                    \n",
      " -------  ---------------  ---------  ---------  -------  -------  -------------------------------------------\n",
      "    52.3          1868937          3   622979.0   619928   624664  initWith(float, float*, int)               \n",
      "    47.7          1705355          1  1705355.0  1705355  1705355  addVectorsInto(float*, float*, float*, int)\n",
      "\n",
      "\n",
      "\n",
      "CUDA Memory Operation Statistics (by time):\n",
      "\n",
      " Time(%)  Total Time (ns)  Operations  Average   Minimum  Maximum              Operation            \n",
      " -------  ---------------  ----------  --------  -------  -------  ---------------------------------\n",
      "   100.0         20436902          64  319326.6   319068   320444  [CUDA Unified Memory memcpy DtoH]\n",
      "\n",
      "\n",
      "\n",
      "CUDA Memory Operation Statistics (by size in KiB):\n",
      "\n",
      "   Total     Operations  Average   Minimum   Maximum               Operation            \n",
      " ----------  ----------  --------  --------  --------  ---------------------------------\n",
      " 131072.000          64  2048.000  2048.000  2048.000  [CUDA Unified Memory memcpy DtoH]\n",
      "\n",
      "\n",
      "\n",
      "Operating System Runtime API Statistics:\n",
      "\n",
      " Time(%)  Total Time (ns)  Num Calls   Average    Minimum   Maximum        Name     \n",
      " -------  ---------------  ---------  ----------  -------  ---------  --------------\n",
      "    66.1        431214396         27  15970903.6     4223  100128871  poll          \n",
      "    23.0        149919620        684    219180.7     1008   41692932  ioctl         \n",
      "     8.5         55186369         24   2299432.0    21688   20879828  sem_timedwait \n",
      "     2.1         13671847         94    145445.2     1398    8672313  mmap          \n",
      "     0.2          1574155         82     19197.0     6394      35079  open64        \n",
      "     0.0           218311          3     72770.3    69399      76987  fgets         \n",
      "     0.0           184190          4     46047.5    36837      54460  pthread_create\n",
      "     0.0           152442         25      6097.7     1623      25245  fopen         \n",
      "     0.0            96930         11      8811.8     4477      15038  write         \n",
      "     0.0            50963         12      4246.9     1589       7939  munmap        \n",
      "     0.0            46298          5      9259.6     4515      16069  open          \n",
      "     0.0            31548         18      1752.7     1149       5936  fclose        \n",
      "     0.0            30225         13      2325.0     1444       3626  read          \n",
      "     0.0            25703          6      4283.8     1085      12284  fgetc         \n",
      "     0.0            19723          2      9861.5     8124      11599  socket        \n",
      "     0.0            18406         12      1533.8     1022       5967  fcntl         \n",
      "     0.0            14393          4      3598.3     2264       5370  fread         \n",
      "     0.0            10377          1     10377.0    10377      10377  pipe2         \n",
      "     0.0             9986          1      9986.0     9986       9986  connect       \n",
      "     0.0             3235          1      3235.0     3235       3235  bind          \n",
      "     0.0             1879          1      1879.0     1879       1879  listen        \n",
      "\n",
      "Report file moved to \"/dli/task/prefetch-to-host-report.qdrep\"\n",
      "Report file moved to \"/dli/task/prefetch-to-host-report.sqlite\"\n",
      "\n"
     ]
    }
   ],
   "source": [
    "!nsys profile --stats=true -o prefetch-to-host-report ./prefetch-to-host"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Open this report file in Nsight Systems, and do the following:\n",
    "\n",
    "- Use the *Unified Memory* section of the timeline to compare and contrast the *Data Migration (DtoH)* events before and after adding prefetching back to the CPU."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## Concurrent CUDA Streams"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "You are now going to learn about a new concept, **CUDA Streams**. After an introduction to them, you will return to using Nsight Systems to better evaluate their impact on your application's performance."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "The following slides present upcoming material visually, at a high level. Click through the slides before moving on to more detailed coverage of their topics in following sections."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 10,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/html": [
       "\n",
       "<div align=\"center\"><iframe src=\"https://view.officeapps.live.com/op/view.aspx?src=https://developer.download.nvidia.com/training/courses/C-AC-01-V1/embedded/task3/NVVP-Streams-1.pptx\" width=\"800px\" height=\"500px\" frameborder=\"0\"></iframe></div>\n"
      ],
      "text/plain": [
       "<IPython.core.display.HTML object>"
      ]
     },
     "metadata": {},
     "output_type": "display_data"
    }
   ],
   "source": [
    "%%HTML\n",
    "\n",
    "<div align=\"center\"><iframe src=\"https://view.officeapps.live.com/op/view.aspx?src=https://developer.download.nvidia.com/training/courses/C-AC-01-V1/embedded/task3/NVVP-Streams-1.pptx\" width=\"800px\" height=\"500px\" frameborder=\"0\"></iframe></div>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "In CUDA programming, a **stream** is a series of commands that execute in order. In CUDA applications, kernel execution, as well as some memory transfers, occur within CUDA streams. Up until this point in time, you have not been interacting explicitly with CUDA streams, but in fact, your CUDA code has been executing its kernels inside of a stream called *the default stream*.\n",
    "\n",
    "CUDA programmers can create and utilize non-default CUDA streams in addition to the default stream, and in doing so, perform multiple operations, such as executing multiple kernels, concurrently, in different streams. Using multiple streams can add an additional layer of parallelization to your accelerated applications, and offers many more opportunities for application optimization."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Rules Governing the Behavior of CUDA Streams\n",
    "\n",
    "There are a few rules, concerning the behavior of CUDA streams, that should be learned in order to utilize them effectively:\n",
    "\n",
    "- Operations within a given stream occur in order.\n",
    "- Operations in different non-default streams are not guaranteed to operate in any specific order relative to each other.\n",
    "- The default stream is blocking and will both wait for all other streams to complete before running, and, will block other streams from running until it completes."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Creating, Utilizing, and Destroying Non-Default CUDA Streams\n",
    "\n",
    "The following code snippet demonstrates how to create, utilize, and destroy a non-default CUDA stream. You will note, that to launch a CUDA kernel in a non-default CUDA stream, the stream must be passed as the optional 4th argument of the execution configuration. Up until now you have only utilized the first 2 arguments of the execution configuration:\n",
    "\n",
    "```cpp\n",
    "cudaStream_t stream;       // CUDA streams are of type `cudaStream_t`.\n",
    "cudaStreamCreate(&stream); // Note that a pointer must be passed to `cudaCreateStream`.\n",
    "\n",
    "someKernel<<<number_of_blocks, threads_per_block, 0, stream>>>(); // `stream` is passed as 4th EC argument.\n",
    "\n",
    "cudaStreamDestroy(stream); // Note that a value, not a pointer, is passed to `cudaDestroyStream`.\n",
    "```\n",
    "\n",
    "Outside the scope of this lab, but worth mentioning, is the optional 3rd argument of the execution configuration. This argument allows programmers to supply the number of bytes in **shared memory** (an advanced topic that will not be covered presently) to be dynamically allocated per block for this kernel launch. The default number of bytes allocated to shared memory per block is `0`, and for the remainder of the lab, you will be passing `0` as this value, in order to expose the 4th argument, which is of immediate interest:"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Exercise: Predict Default Stream Behavior\n",
    "\n",
    "The [01-print-numbers](../edit/05-stream-intro/01-print-numbers.cu) application has a very simple `printNumber` kernel which accepts an integer and prints it. The kernel is only being executed with a single thread inside a single block. However, it is being executed 5 times, using a for-loop, and passing each launch the number of the for-loop's iteration.\n",
    "\n",
    "Compile and run [01-print-numbers](../edit/05-stream-intro/01-print-numbers.cu) using the code execution block below. You should see the numbers `0` through `4` printed."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 11,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "0\r\n",
      "1\r\n",
      "2\r\n",
      "3\r\n",
      "4\r\n"
     ]
    }
   ],
   "source": [
    "!nvcc -o print-numbers 05-stream-intro/01-print-numbers.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Knowing that by default kernels are executed in the default stream, would you expect that the 5 launches of the `print-numbers` program executed serially, or in parallel? You should be able to mention two features of the default stream to support your answer. Create a report file in the cell below and open it in Nsight Systems to confirm your answer."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 12,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "Warning: LBR backtrace method is not supported on this platform. DWARF backtrace method will be used.\n",
      "WARNING: The command line includes a target application therefore the CPU context-switch scope has been set to process-tree.\n",
      "Collecting data...\n",
      "0\n",
      "1\n",
      "2\n",
      "3\n",
      "4\n",
      "Processing events...\n",
      "Saving temporary \"/tmp/nsys-report-e76f-0d9b-67ab-9e93.qdstrm\" file to disk...\n",
      "\n",
      "Creating final output files...\n",
      "Processing [==============================================================100%]\n",
      "Saved report file to \"/tmp/nsys-report-e76f-0d9b-67ab-9e93.qdrep\"\n",
      "Exporting 1028 events: [==================================================100%]\n",
      "\n",
      "Exported successfully to\n",
      "/tmp/nsys-report-e76f-0d9b-67ab-9e93.sqlite\n",
      "\n",
      "\n",
      "CUDA API Statistics:\n",
      "\n",
      " Time(%)  Total Time (ns)  Num Calls   Average    Minimum   Maximum           Name         \n",
      " -------  ---------------  ---------  ----------  -------  ---------  ---------------------\n",
      "    99.9        231590227          5  46318045.4     3972  231570322  cudaLaunchKernel     \n",
      "     0.1           274340          1    274340.0   274340     274340  cudaDeviceSynchronize\n",
      "\n",
      "\n",
      "\n",
      "CUDA Kernel Statistics:\n",
      "\n",
      " Time(%)  Total Time (ns)  Instances  Average  Minimum  Maximum        Name      \n",
      " -------  ---------------  ---------  -------  -------  -------  ----------------\n",
      "   100.0           274364          5  54872.8    53087    61695  printNumber(int)\n",
      "\n",
      "\n",
      "\n",
      "Operating System Runtime API Statistics:\n",
      "\n",
      " Time(%)  Total Time (ns)  Num Calls   Average    Minimum   Maximum        Name     \n",
      " -------  ---------------  ---------  ----------  -------  ---------  --------------\n",
      "    67.2        230846288         14  16489020.6    24046  100129940  poll          \n",
      "    30.9        106382468        668    159255.2     1101   18435339  ioctl         \n",
      "     0.9          3145982         87     36160.7     1479     984888  mmap          \n",
      "     0.6          1893699         82     23093.9     7394      43232  open64        \n",
      "     0.2           631128         11     57375.3    14440     384134  sem_timedwait \n",
      "     0.1           217219          3     72406.3    69462      75986  fgets         \n",
      "     0.0           170027          4     42506.8    34711      50220  pthread_create\n",
      "     0.0           165936         25      6637.4     1585      26231  fopen         \n",
      "     0.0           103666         12      8638.8     4567      15562  write         \n",
      "     0.0            36806          5      7361.2     4820      10487  open          \n",
      "     0.0            31476          7      4496.6     1399       9667  munmap        \n",
      "     0.0            30858         18      1714.3     1153       5721  fclose        \n",
      "     0.0            24766          6      4127.7     1070      11018  fgetc         \n",
      "     0.0            24370         13      1874.6     1013       2926  read          \n",
      "     0.0            19169          2      9584.5     7529      11640  socket        \n",
      "     0.0            17964          2      8982.0     5776      12188  fread         \n",
      "     0.0            13661          9      1517.9     1003       4603  fcntl         \n",
      "     0.0             8429          1      8429.0     8429       8429  pipe2         \n",
      "     0.0             7592          1      7592.0     7592       7592  connect       \n",
      "     0.0             2728          1      2728.0     2728       2728  bind          \n",
      "     0.0             1577          1      1577.0     1577       1577  listen        \n",
      "\n",
      "Report file moved to \"/dli/task/print-numbers-report.qdrep\"\n",
      "Report file moved to \"/dli/task/print-numbers-report.sqlite\"\n",
      "\n"
     ]
    }
   ],
   "source": [
    "!nsys profile --stats=true -o print-numbers-report ./print-numbers"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Exercise: Implement Concurrent CUDA Streams\n",
    "\n",
    "Both because all 5 kernel launches occurred in the same stream, you should not be surprised to have seen that the 5 kernels executed serially. Additionally you could make the case that because the default stream is blocking, each launch of the kernel would wait to complete before the next launch, and this is also true.\n",
    "\n",
    "Refactor [01-print-numbers](../edit/05-stream-intro/01-print-numbers.cu) so that each kernel launch occurs in its own non-default stream. Be sure to destroy the streams you create after they are no longer needed. Compile and run the refactored code with the code execution cell directly below. You should still see the numbers `0` through `4` printed, though not necessarily in ascending order. Refer to [the solution](../edit/05-stream-intro/solutions/01-print-numbers-solution.cu) if you get stuck."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 13,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "0\r\n",
      "1\r\n",
      "2\r\n",
      "3\r\n",
      "4\r\n"
     ]
    }
   ],
   "source": [
    "!nvcc -o print-numbers-in-streams 05-stream-intro/01-print-numbers.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Now that you are using 5 different non-default streams for each of the 5 kernel launches, do you expect that they will run serially or in parallel? In addition to what you now know about streams, take into account how trivial the `printNumber` kernel is, meaning, even if you predict parallel runs, will the speed at which one kernel will complete allow for complete overlap?\n",
    "\n",
    "After hypothesizing, open a new report file in Nsight Systems to view its actual behavior. You should notice that now, there are additional rows in the _CUDA_ section for each of the non-default streams you created:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 14,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "Warning: LBR backtrace method is not supported on this platform. DWARF backtrace method will be used.\n",
      "WARNING: The command line includes a target application therefore the CPU context-switch scope has been set to process-tree.\n",
      "Collecting data...\n",
      "0\n",
      "1\n",
      "2\n",
      "3\n",
      "4\n",
      "Processing events...\n",
      "Saving temporary \"/tmp/nsys-report-80b0-dda7-f5cd-d58e.qdstrm\" file to disk...\n",
      "\n",
      "Creating final output files...\n",
      "Processing [==============================================================100%]\n",
      "Saved report file to \"/tmp/nsys-report-80b0-dda7-f5cd-d58e.qdrep\"\n",
      "Exporting 1031 events: [==================================================100%]\n",
      "\n",
      "Exported successfully to\n",
      "/tmp/nsys-report-80b0-dda7-f5cd-d58e.sqlite\n",
      "\n",
      "\n",
      "CUDA API Statistics:\n",
      "\n",
      " Time(%)  Total Time (ns)  Num Calls   Average    Minimum   Maximum           Name         \n",
      " -------  ---------------  ---------  ----------  -------  ---------  ---------------------\n",
      "    99.9        232188000          5  46437600.0     3983  232166656  cudaLaunchKernel     \n",
      "     0.1           274624          1    274624.0   274624     274624  cudaDeviceSynchronize\n",
      "\n",
      "\n",
      "\n",
      "CUDA Kernel Statistics:\n",
      "\n",
      " Time(%)  Total Time (ns)  Instances  Average  Minimum  Maximum        Name      \n",
      " -------  ---------------  ---------  -------  -------  -------  ----------------\n",
      "   100.0           275516          5  55103.2    53183    62335  printNumber(int)\n",
      "\n",
      "\n",
      "\n",
      "Operating System Runtime API Statistics:\n",
      "\n",
      " Time(%)  Total Time (ns)  Num Calls   Average    Minimum   Maximum        Name     \n",
      " -------  ---------------  ---------  ----------  -------  ---------  --------------\n",
      "    67.1        231071908         14  16505136.3    24669  100131150  poll          \n",
      "    31.0        106625128        669    159379.9     1088   18340350  ioctl         \n",
      "     0.9          3074756         87     35342.0     1324     942128  mmap          \n",
      "     0.4          1499785         82     18290.1     5654      30185  open64        \n",
      "     0.3           982950         11     89359.1    22281     659979  sem_timedwait \n",
      "     0.1           216922          3     72307.3    69507      75987  fgets         \n",
      "     0.1           177391          4     44347.8    35740      52742  pthread_create\n",
      "     0.0           138453         25      5538.1     1542      25362  fopen         \n",
      "     0.0            93945         12      7828.8     4387      12563  write         \n",
      "     0.0            65779          9      7308.8     1866      33186  munmap        \n",
      "     0.0            43884          7      6269.1     1070      16570  fgetc         \n",
      "     0.0            34449          5      6889.8     4479       9616  open          \n",
      "     0.0            31744         14      2267.4     1342       3974  read          \n",
      "     0.0            28994         18      1610.8     1144       5509  fclose        \n",
      "     0.0            15567          2      7783.5     6599       8968  socket        \n",
      "     0.0            13499          7      1928.4     1027       6238  fcntl         \n",
      "     0.0             9418          2      4709.0     3816       5602  fread         \n",
      "     0.0             8740          1      8740.0     8740       8740  connect       \n",
      "     0.0             7586          1      7586.0     7586       7586  pipe2         \n",
      "     0.0             2759          1      2759.0     2759       2759  bind          \n",
      "     0.0             1730          1      1730.0     1730       1730  listen        \n",
      "\n",
      "Report file moved to \"/dli/task/print-numbers-in-streams-report.qdrep\"\n",
      "Report file moved to \"/dli/task/print-numbers-in-streams-report.sqlite\"\n",
      "\n"
     ]
    }
   ],
   "source": [
    "!nsys profile --stats=true -o print-numbers-in-streams-report print-numbers-in-streams"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "![streams print](images/streams-print.png)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Exercise: Use Streams for Concurrent Data Initialization Kernels\n",
    "\n",
    "The vector addition application you have been working with, [01-prefetch-check-solution.cu](../edit/04-prefetch-check/solutions/01-prefetch-check-solution.cu), currently launches an initialization kernel 3 times - once each for each of the 3 vectors needing initialization for the `vectorAdd` kernel. Refactor it to launch each of the 3 initialization kernel launches in their own non-default stream. You should still see the success message print when compiling and running with the code execution cell below. Refer to [the solution](../edit/06-stream-init/solutions/01-stream-init-solution.cu) if you get stuck."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 15,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "Success! All values calculated correctly.\r\n"
     ]
    }
   ],
   "source": [
    "!nvcc -o init-in-streams 04-prefetch-check/solutions/01-prefetch-check-solution.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Open a report in Nsight Systems to confirm that your 3 initialization kernel launches are running in their own non-default streams, with some degree of concurrent overlap."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 16,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "Warning: LBR backtrace method is not supported on this platform. DWARF backtrace method will be used.\n",
      "WARNING: The command line includes a target application therefore the CPU context-switch scope has been set to process-tree.\n",
      "Collecting data...\n",
      "Success! All values calculated correctly.\n",
      "Processing events...\n",
      "Saving temporary \"/tmp/nsys-report-8dc5-ecf0-b876-18b7.qdstrm\" file to disk...\n",
      "\n",
      "Creating final output files...\n",
      "Processing [==============================================================100%]\n",
      "Saved report file to \"/tmp/nsys-report-8dc5-ecf0-b876-18b7.qdrep\"\n",
      "Exporting 1159 events: [==================================================100%]\n",
      "\n",
      "Exported successfully to\n",
      "/tmp/nsys-report-8dc5-ecf0-b876-18b7.sqlite\n",
      "\n",
      "\n",
      "CUDA API Statistics:\n",
      "\n",
      " Time(%)  Total Time (ns)  Num Calls   Average    Minimum   Maximum           Name         \n",
      " -------  ---------------  ---------  ----------  -------  ---------  ---------------------\n",
      "    80.9        244573306          3  81524435.3    19891  244517730  cudaMallocManaged    \n",
      "    14.4         43671906          4  10917976.5   668161   41576927  cudaMemPrefetchAsync \n",
      "     3.5         10527865          3   3509288.3   842234    8773636  cudaFree             \n",
      "     1.2          3561958          1   3561958.0  3561958    3561958  cudaDeviceSynchronize\n",
      "     0.0            41950          4     10487.5     4543      26473  cudaLaunchKernel     \n",
      "\n",
      "\n",
      "\n",
      "CUDA Kernel Statistics:\n",
      "\n",
      " Time(%)  Total Time (ns)  Instances   Average   Minimum  Maximum                     Name                    \n",
      " -------  ---------------  ---------  ---------  -------  -------  -------------------------------------------\n",
      "    52.3          1869641          3   623213.7   615769   628600  initWith(float, float*, int)               \n",
      "    47.7          1702219          1  1702219.0  1702219  1702219  addVectorsInto(float*, float*, float*, int)\n",
      "\n",
      "\n",
      "\n",
      "CUDA Memory Operation Statistics (by time):\n",
      "\n",
      " Time(%)  Total Time (ns)  Operations  Average   Minimum  Maximum              Operation            \n",
      " -------  ---------------  ----------  --------  -------  -------  ---------------------------------\n",
      "   100.0         20436450          64  319319.5   319036   320412  [CUDA Unified Memory memcpy DtoH]\n",
      "\n",
      "\n",
      "\n",
      "CUDA Memory Operation Statistics (by size in KiB):\n",
      "\n",
      "   Total     Operations  Average   Minimum   Maximum               Operation            \n",
      " ----------  ----------  --------  --------  --------  ---------------------------------\n",
      " 131072.000          64  2048.000  2048.000  2048.000  [CUDA Unified Memory memcpy DtoH]\n",
      "\n",
      "\n",
      "\n",
      "Operating System Runtime API Statistics:\n",
      "\n",
      " Time(%)  Total Time (ns)  Num Calls   Average    Minimum   Maximum        Name     \n",
      " -------  ---------------  ---------  ----------  -------  ---------  --------------\n",
      "    67.3        440777481         28  15742052.9    23517  100130959  poll          \n",
      "    21.9        143514488        682    210431.8     1015   41519041  ioctl         \n",
      "     8.4         54711008         23   2378739.5    20803   20541787  sem_timedwait \n",
      "     2.0         13355760         94    142082.6     1431    8719833  mmap          \n",
      "     0.3          1674280         82     20418.0     4831      37871  open64        \n",
      "     0.0           213829          3     71276.3    69211      74445  fgets         \n",
      "     0.0           148427          4     37106.8    32971      41242  pthread_create\n",
      "     0.0           116171         25      4646.8     1548      20856  fopen         \n",
      "     0.0            90438         11      8221.6     4346      13884  write         \n",
      "     0.0            40755         11      3705.0     1866       5626  munmap        \n",
      "     0.0            27540         18      1530.0     1109       4290  fclose        \n",
      "     0.0            27466          5      5493.2     3306       8088  open          \n",
      "     0.0            23602         16      1475.1     1024       4901  fcntl         \n",
      "     0.0            22748         13      1749.8     1069       2634  read          \n",
      "     0.0            21690          6      3615.0     1122       9576  fgetc         \n",
      "     0.0            12134          4      3033.5     1844       4366  fread         \n",
      "     0.0            10615          2      5307.5     4884       5731  socket        \n",
      "     0.0             7620          1      7620.0     7620       7620  pipe2         \n",
      "     0.0             6242          1      6242.0     6242       6242  connect       \n",
      "     0.0             2160          1      2160.0     2160       2160  bind          \n",
      "     0.0             1494          1      1494.0     1494       1494  listen        \n",
      "\n",
      "Report file moved to \"/dli/task/init-in-streams-report.qdrep\"\n",
      "Report file moved to \"/dli/task/init-in-streams-report.sqlite\"\n",
      "\n"
     ]
    }
   ],
   "source": [
    "!nsys profile --stats=true -o init-in-streams-report ./init-in-streams"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## Summary\n",
    "\n",
    "At this point in the lab you are able to:\n",
    "\n",
    "- Use the **Nsight Systems** to visually profile the timeline of GPU-accelerated CUDA applications.\n",
    "- Use Nsight Systems to identify, and exploit, optimization opportunities in GPU-accelerated CUDA applications.\n",
    "- Utilize CUDA streams for concurrent kernel execution in accelerated applications.\n",
    "\n",
    "At this point in time you have a wealth of fundamental tools and techniques for accelerating CPU-only applications, and for then optimizing those accelerated applications. In the final exercise, you will have a chance to apply everything that you've learned to accelerate an [n-body](https://en.wikipedia.org/wiki/N-body_problem) simulator, which predicts the individual motions of a group of objects interacting with each other gravitationally."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## Final Exercise: Accelerate and Optimize an N-Body Simulator\n",
    "\n",
    "An [n-body](https://en.wikipedia.org/wiki/N-body_problem) simulator predicts the individual motions of a group of objects interacting with each other gravitationally. [01-nbody.cu](../edit/09-nbody/01-nbody.cu) contains a simple, though working, n-body simulator for bodies moving through 3 dimensional space.\n",
    "\n",
    "In its current CPU-only form, this application takes about 5 seconds to run on 4096 particles, and **20 minutes** to run on 65536 particles. Your task is to GPU accelerate the program, retaining the correctness of the simulation."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Considerations to Guide Your Work\n",
    "\n",
    "Here are some things to consider before beginning your work:\n",
    "\n",
    "- Especially for your first refactors, the logic of the application, the `bodyForce` function in particular, can and should remain largely unchanged: focus on accelerating it as easily as possible.\n",
    "- The code base contains a for-loop inside `main` for integrating the interbody forces calculated by `bodyForce` into the positions of the bodies in the system. This integration both needs to occur after `bodyForce` runs, and, needs to complete before the next call to `bodyForce`. Keep this in mind when choosing how and where to parallelize.\n",
    "- Use a **profile driven** and iterative approach.\n",
    "- You are not required to add error handling to your code, but you might find it helpful, as you are responsible for your code working correctly.\n",
    "\n",
    "**Have Fun!**"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Use this cell to compile the nbody simulator. Although it is initially a CPU-only application, is does accurately simulate the positions of the particles."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 17,
   "metadata": {
    "scrolled": true
   },
   "outputs": [],
   "source": [
    "!nvcc -std=c++11 -o nbody 09-nbody/01-nbody.cu"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "It is highly recommended you use the profiler to assist your work. Execute the following cell to generate a report file:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 18,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "Warning: LBR backtrace method is not supported on this platform. DWARF backtrace method will be used.\n",
      "WARNING: The command line includes a target application therefore the CPU context-switch scope has been set to process-tree.\n",
      "Collecting data...\n",
      "0.041 Billion Interactions / second\n",
      "Processing events...\n",
      "Saving temporary \"/tmp/nsys-report-ec9f-8f8a-c6f2-19fc.qdstrm\" file to disk...\n",
      "\n",
      "Creating final output files...\n",
      "Processing [==============================================================100%]\n",
      "Saved report file to \"/tmp/nsys-report-ec9f-8f8a-c6f2-19fc.qdrep\"\n",
      "Exporting 39 events: [====================================================100%]\n",
      "\n",
      "Exported successfully to\n",
      "/tmp/nsys-report-ec9f-8f8a-c6f2-19fc.sqlite\n",
      "\n",
      "\n",
      "Operating System Runtime API Statistics:\n",
      "\n",
      " Time(%)  Total Time (ns)  Num Calls  Average  Minimum  Maximum   Name  \n",
      " -------  ---------------  ---------  -------  -------  -------  -------\n",
      "    32.5            80266          1  80266.0    80266    80266  writev \n",
      "    30.3            74685          2  37342.5     7774    66911  fopen64\n",
      "    21.7            53575          1  53575.0    53575    53575  read   \n",
      "    15.5            38284          2  19142.0     2357    35927  fclose \n",
      "\n",
      "Report file moved to \"/dli/task/nbody-report.qdrep\"\n",
      "Report file moved to \"/dli/task/nbody-report.sqlite\"\n",
      "\n"
     ]
    }
   ],
   "source": [
    "!nsys profile --stats=true --force-overwrite=true -o nbody-report ./nbody"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Here we import a function that will run your `nbody` simulator against a various number of particles, checking for performance and accuracy."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 19,
   "metadata": {},
   "outputs": [],
   "source": [
    "from assessment import run_assessment"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Execute the following cell to run and assess `nbody`:"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 20,
   "metadata": {
    "scrolled": true
   },
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "Running nbody simulator with 4096 bodies\n",
      "----------------------------------------\n",
      "\n",
      "Application should run faster than 0.9s\n",
      "Your application ran in: 4.1226s\n",
      "Your application is not yet fast enough\n"
     ]
    }
   ],
   "source": [
    "run_assessment()"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## Generate a Certificate"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "If you passed the assessment, please return to the course page (shown below) and click the \"ASSESS TASK\" button, which will generate your certificate for the course."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "![run_assessment](./images/run_assessment.png)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## Advanced Content\n",
    "\n",
    "The following sections, for those of you with time and interest, introduce more intermediate techniques involving some manual device memory management, and using non-default streams to overlap kernel execution and memory copies.\n",
    "\n",
    "After learning about each of the techniques below, try to further optimize your nbody simulation using these techniques."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## Manual Device Memory Allocation and Copying\n",
    "\n",
    "While `cudaMallocManaged` and `cudaMemPrefetchAsync` are performant, and greatly simplify memory migration, sometimes it can be worth it to use more manual methods for memory allocation. This is particularly true when it is known that data will only be accessed on the device or host, and the cost of migrating data can be reclaimed in exchange for the fact that no automatic on-demand migration is needed.\n",
    "\n",
    "Additionally, using manual device memory management can allow for the use of non-default streams for overlapping data transfers with computational work. In this section you will learn some basic manual device memory allocation and copy techniques, before extending these techniques to overlap data copies with computational work. \n",
    "\n",
    "Here are some CUDA commands for manual device memory management:\n",
    "\n",
    "- `cudaMalloc` will allocate memory directly to the active GPU. This prevents all GPU page faults. In exchange, the pointer it returns is not available for access by host code.\n",
    "- `cudaMallocHost` will allocate memory directly to the CPU. It also \"pins\" the memory, or page locks it, which will allow for asynchronous copying of the memory to and from a GPU. Too much pinned memory can interfere with CPU performance, so use it only with intention. Pinned memory should be freed with `cudaFreeHost`.\n",
    "- `cudaMemcpy` can copy (not transfer) memory, either from host to device or from device to host."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Manual Device Memory Management Example\n",
    "\n",
    "Here is a snippet of code that demonstrates the use of the above CUDA API calls.\n",
    "\n",
    "```cpp\n",
    "int *host_a, *device_a;        // Define host-specific and device-specific arrays.\n",
    "cudaMalloc(&device_a, size);   // `device_a` is immediately available on the GPU.\n",
    "cudaMallocHost(&host_a, size); // `host_a` is immediately available on CPU, and is page-locked, or pinned.\n",
    "\n",
    "initializeOnHost(host_a, N);   // No CPU page faulting since memory is already allocated on the host.\n",
    "\n",
    "// `cudaMemcpy` takes the destination, source, size, and a CUDA-provided variable for the direction of the copy.\n",
    "cudaMemcpy(device_a, host_a, size, cudaMemcpyHostToDevice);\n",
    "\n",
    "kernel<<<blocks, threads, 0, someStream>>>(device_a, N);\n",
    "\n",
    "// `cudaMemcpy` can also copy data from device to host.\n",
    "cudaMemcpy(host_a, device_a, size, cudaMemcpyDeviceToHost);\n",
    "\n",
    "verifyOnHost(host_a, N);\n",
    "\n",
    "cudaFree(device_a);\n",
    "cudaFreeHost(host_a);          // Free pinned memory like this.\n",
    "```"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Exercise: Manually Allocate Host and Device Memory\n",
    "\n",
    "The most recent iteration of the vector addition application, [01-stream-init-solution](../edit/06-stream-init/solutions/01-stream-init-solution.cu), is using `cudaMallocManaged` to allocate managed memory first used on the device by the initialization kernels, then on the device by the vector add kernel, and then by the host, where the memory is automatically transferred, for verification. This is a sensible approach, but it is worth experimenting with some manual device memory allocation and copying to observe its impact on the application's performance.\n",
    "\n",
    "Refactor the [01-stream-init-solution](../edit/06-stream-init/solutions/01-stream-init-solution.cu) application to **not** use `cudaMallocManaged`. In order to do this you will need to do the following:\n",
    "\n",
    "- Replace calls to `cudaMallocManaged` with `cudaMalloc`.\n",
    "- Create an additional vector that will be used for verification on the host. This is required since the memory allocated with `cudaMalloc` is not available to the host. Allocate this host vector with `cudaMallocHost`.\n",
    "- After the `addVectorsInto` kernel completes, use `cudaMemcpy` to copy the vector with the addition results, into the host vector you created with `cudaMallocHost`.\n",
    "- Use `cudaFreeHost` to free the memory allocated with `cudaMallocHost`.\n",
    "\n",
    "Refer to [the solution](../edit/07-manual-malloc/solutions/01-manual-malloc-solution.cu) if you get stuck."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "!nvcc -o vector-add-manual-alloc 06-stream-init/solutions/01-stream-init-solution.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "After completing the refactor, open a report in Nsight Systems, and use the timeline to do the following:\n",
    "\n",
    "- Notice that there is no longer a *Unified Memory* section of the timeline.\n",
    "- Comparing this timeline to that of the previous refactor, compare the run times of `cudaMalloc` in the current application vs. `cudaMallocManaged` in the previous.\n",
    "- Notice how in the current application, work on the initialization kernels does not start until a later time than it did in the previous iteration. Examination of the timeline will show the difference is the time taken by `cudaMallocHost`. This clearly points out the difference between memory transfers, and memory copies. When copying memory, as you are doing presently, the data will exist in 2 different places in the system. In the current case, the allocation of the 4th host-only vector incurs a small cost in performance, compared to only allocating 3 vectors in the previous iteration."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## Using Streams to Overlap Data Transfers and Code Execution\n",
    "\n",
    "The following slides present upcoming material visually, at a high level. Click through the slides before moving on to more detailed coverage of their topics in following sections."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "%%HTML\n",
    "\n",
    "<div align=\"center\"><iframe src=\"https://view.officeapps.live.com/op/view.aspx?src=https://developer.download.nvidia.com/training/courses/C-AC-01-V1/embedded/task3/NVVP-Streams-3.pptx\" width=\"800px\" height=\"500px\" frameborder=\"0\"></iframe></div>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "In addition to `cudaMemcpy` is `cudaMemcpyAsync` which can asynchronously copy memory either from host to device or from device to host as long as the host memory is pinned, which can be done by allocating it with `cudaMallocHost`.\n",
    "\n",
    "Similar to kernel execution, `cudaMemcpyAsync` is only asynchronous by default with respect to the host. It executes, by default, in the default stream and therefore is a blocking operation with regard to other CUDA operations occurring on the GPU. The `cudaMemcpyAsync` function, however, takes as an optional 5th argument, a non-default stream. By passing it a non-default stream, the memory transfer can be concurrent to other CUDA operations occurring in other non-default streams.\n",
    "\n",
    "A common and useful pattern is to use a combination of pinned host memory, asynchronous memory copies in non-default streams, and kernel executions in non-default streams, to overlap memory transfers with kernel execution.\n",
    "\n",
    "In the following example, rather than wait for the entire memory copy to complete before beginning work on the kernel, segments of the required data are copied and worked on, with each copy/work segment running in its own non-default stream. Using this technique, work on parts of the data can begin while memory transfers for later segments occur concurrently. Extra care must be taken when using this technique to calculate segment-specific values for the number of operations, and the offset location inside arrays, as shown here:\n",
    "\n",
    "```cpp\n",
    "int N = 2<<24;\n",
    "int size = N * sizeof(int);\n",
    "\n",
    "int *host_array;\n",
    "int *device_array;\n",
    "\n",
    "cudaMallocHost(&host_array, size);               // Pinned host memory allocation.\n",
    "cudaMalloc(&device_array, size);                 // Allocation directly on the active GPU device.\n",
    "\n",
    "initializeData(host_array, N);                   // Assume this application needs to initialize on the host.\n",
    "\n",
    "const int numberOfSegments = 4;                  // This example demonstrates slicing the work into 4 segments.\n",
    "int segmentN = N / numberOfSegments;             // A value for a segment's worth of `N` is needed.\n",
    "size_t segmentSize = size / numberOfSegments;    // A value for a segment's worth of `size` is needed.\n",
    "\n",
    "// For each of the 4 segments...\n",
    "for (int i = 0; i < numberOfSegments; ++i)\n",
    "{\n",
    "  // Calculate the index where this particular segment should operate within the larger arrays.\n",
    "  segmentOffset = i * segmentN;\n",
    "\n",
    "  // Create a stream for this segment's worth of copy and work.\n",
    "  cudaStream_t stream;\n",
    "  cudaStreamCreate(&stream);\n",
    "  \n",
    "  // Asynchronously copy segment's worth of pinned host memory to device over non-default stream.\n",
    "  cudaMemcpyAsync(&device_array[segmentOffset],  // Take care to access correct location in array.\n",
    "                  &host_array[segmentOffset],    // Take care to access correct location in array.\n",
    "                  segmentSize,                   // Only copy a segment's worth of memory.\n",
    "                  cudaMemcpyHostToDevice,\n",
    "                  stream);                       // Provide optional argument for non-default stream.\n",
    "                  \n",
    "  // Execute segment's worth of work over same non-default stream as memory copy.\n",
    "  kernel<<<number_of_blocks, threads_per_block, 0, stream>>>(&device_array[segmentOffset], segmentN);\n",
    "  \n",
    "  // `cudaStreamDestroy` will return immediately (is non-blocking), but will not actually destroy stream until\n",
    "  // all stream operations are complete.\n",
    "  cudaStreamDestroy(stream);\n",
    "}\n",
    "```"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Exercise: Overlap Kernel Execution and Memory Copy Back to Host\n",
    "\n",
    "The most recent iteration of the vector addition application, [01-manual-malloc-solution.cu](../edit/07-manual-malloc/solutions/01-manual-malloc-solution.cu), is currently performing all of its vector addition work on the GPU before copying the memory back to the host for verification.\n",
    "\n",
    "Refactor [01-manual-malloc-solution.cu](../edit/07-manual-malloc/solutions/01-manual-malloc-solution.cu) to perform the vector addition in 4 segments, in non-default streams, so that asynchronous memory copies can begin before waiting for all vector addition work to complete. Refer to [the solution](../edit/08-overlap-xfer/solutions/01-overlap-xfer-solution.cu) if you get stuck."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "!nvcc -o vector-add-manual-alloc 07-manual-malloc/solutions/01-manual-malloc-solution.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "After completing the refactor, open a report in Nsight Systems, and use the timeline to do the following:\n",
    "\n",
    "- Note when the device to host memory transfers begin, is it before or after all kernel work has completed?\n",
    "- Notice that the 4 memory copy segments themselves do not overlap. Even in separate non-default streams, only one memory transfer in a given direction (DtoH here) at a time can occur simultaneously. The performance gains here are in the ability to start the transfers earlier than otherwise, and it is not hard to imagine in an application where a less trivial amount of work was being done compared to a simple addition operation, that the memory copies would not only start earlier, but also overlap with kernel execution."
   ]
  }
 ],
 "metadata": {
  "kernelspec": {
   "display_name": "Python 3 (ipykernel)",
   "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.9.7"
  }
 },
 "nbformat": 4,
 "nbformat_minor": 1
}