05_DirectivesOpenACC/05p2_openACC_miniWeather_Tutorial.ipynb

{
 "cells": [
  {
   "cell_type": "markdown",
   "id": "36cc985c-ebf2-4a54-9611-90234bb70d39",
   "metadata": {
    "slideshow": {
     "slide_type": "slide"
    },
    "tags": []
   },
   "source": [
    "![NCAR UCAR Logo](../NCAR_CISL_NSF_banner.jpeg)\n",
    "# Directive Based Programming with OpenACC and MiniWeather, Part 2\n",
    "\n",
    "By: Daniel Howard [dhoward@ucar.edu](mailto:dhoward@ucar.edu), Consulting Services Group, CISL & NCAR \n",
    "\n",
    "Date: April 14th 2022"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "dee8029a-c48d-482c-8907-a2e4cb1e6588",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "In this notebook we explore the mini-app [MiniWeather](https://github.com/mrnorman/miniWeather) to present techniques and code examples using MiniWeather for using OpenACC to develop for GPUs. Extending from [Part 1](05_openACC_miniWeather_Tutorial.ipynb), we will cover:\n",
    "1. Detailing of OpenACC API Directives\n",
    "    * Data Constructs - `!$acc data` & `!$acc end data` plus `!$acc enter/exit data`\n",
    "    * Routine Directives and Other Clauses - `!$acc routine gang/worker/vector/seq`\n",
    "    * Async and Wait Directives - `!$acc async()` & `!$acc wait()`\n",
    "    * OpenACC API Runtime Library Routines\n",
    "2. Interfacing OpenACC with CUDA\n",
    "3. Advanced OpenACC Optimization Techniques"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "dfe31bf2-bfbd-41f4-b8f7-27e000315671",
   "metadata": {
    "slideshow": {
     "slide_type": "slide"
    },
    "tags": []
   },
   "source": [
    " Head to the [NCAR JupyterHub portal](https://jupyterhub.hpc.ucar.edu/stable) and __start a JupyterHub session on Casper login__ (or batch nodes using 1 CPU, no GPUs) and open the notebook in `05_DirectivesOpenACC/05p2_openACC_miniWeather_Tutorial.ipynb`. Be sure to clone (if needed) and update/pull the NCAR GPU_workshop directory.\n",
    "\n",
    "```shell\n",
    "# Use the JupyterHub GitHub GUI on the left panel or the below shell commands\n",
    "git clone git@github.com:NCAR/GPU_workshop.git\n",
    "git pull\n",
    "```\n",
    "\n",
    "# Workshop Etiquette\n",
    "* Please mute yourself and turn off video during the session.\n",
    "* Questions may be submitted in the chat and will be answered when appropriate. You may also raise your hand, unmute, and ask questions during Q&A at the end of the presentation.\n",
    "* By participating, you are agreeing to [UCAR’s Code of Conduct](https://www.ucar.edu/who-we-are/ethics-integrity/codes-conduct/participants)\n",
    "* Recordings & other material will be archived & shared publicly.\n",
    "* Feel free to follow up with the GPU workshop team via Slack or submit support requests to [support.ucar.edu](https://support.ucar.edu)\n",
    "    * Office Hours: Asynchronous support via [Slack](https://ncargpuusers.slack.com) or schedule a time with an organizer"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "ec2680d6-17f5-4647-b4c1-671c73d6c7a8",
   "metadata": {
    "slideshow": {
     "slide_type": "slide"
    },
    "tags": []
   },
   "source": [
    "## Notebook Setup\n",
    "Set the `PROJECT` code to a currently active project, ie `UCIS0004` for the GPU workshop, and `QUEUE` to the appropriate routing queue depending on if during a live workshop session (`gpuworkshop`), during weekday 8am to 5:30pm MT (`gpudev`), or all other times (`casper`). Due to limited shared GPU resources, please use `GPU_TYPE=gp100` during the workshop. Otherwise, set `GPU_TYPE=v100` (required for `gpudev`) for independent work. See [Casper queue documentation](https://arc.ucar.edu/knowledge_base/72581396#StartingCasperjobswithPBS-Concurrentresourcelimits) for more info.  "
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "46f5252f-e638-4667-b922-93b171cdafd4",
   "metadata": {
    "slideshow": {
     "slide_type": "fragment"
    },
    "tags": []
   },
   "outputs": [],
   "source": [
    "export PROJECT=UCIS0004\n",
    "export QUEUE=gpudev\n",
    "export GPU_TYPE=v100"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "66468d40-070a-4277-b725-7d87f61576a6",
   "metadata": {
    "slideshow": {
     "slide_type": "fragment"
    },
    "tags": []
   },
   "source": [
    "Test that MiniWeather builds correctly below. This build uses the already refactored and complete [miniWeather_mpi_openacc.F90](fortran/miniWeather_mpi_openacc.F90) source file as well as the CPU only [miniWeather_mpi.F90](fortran/miniWeather_mpi.F90) source file which serves as a basis for later exercises."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "cf0f5191-b837-4826-b7aa-5fb385ffe0b3",
   "metadata": {
    "slideshow": {
     "slide_type": "fragment"
    },
    "tags": []
   },
   "outputs": [],
   "source": [
    "cd fortran/build\n",
    "source cmake_casper_nvhpc.sh\n",
    "cd ../..\n",
    "# After running this, there will be the executables `mpi`, `mpi_test`, `openacc` and `openacc_test` in \"fortran/build\""
   ]
  },
  {
   "cell_type": "markdown",
   "id": "516ad65a-6111-4758-b44b-5da406586d55",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "Lastly, `make` and run the new [miniWeather_mpi_exercise2.F90](fortran/miniWeather_mpi_exercise2.F90) program to establish a performance baseline."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "14bd959f-9f42-46e5-94d2-dad722fa46a9",
   "metadata": {
    "slideshow": {
     "slide_type": "fragment"
    },
    "tags": []
   },
   "outputs": [],
   "source": [
    "make -C fortran/build openacc_test_ex2"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "d1a93c22-5d38-4cf2-b21f-93f754d02b45",
   "metadata": {},
   "outputs": [],
   "source": [
    "cd fortran/build\n",
    "qcmd -A $PROJECT -q $QUEUE -l select=1:ncpus=1:ngpus=1 -l gpu_type=$GPU_TYPE -l walltime=60 -v NVCOMPILER_ACC_TIME=1 -- \\\n",
    "$PWD/check_output.sh $PWD/openacc_test_ex2 1e-13 4.5e-5\n",
    "cd ../.."
   ]
  },
  {
   "cell_type": "markdown",
   "id": "0ac238d3-ee59-48a7-9755-06bd74c79a8c",
   "metadata": {
    "slideshow": {
     "slide_type": "slide"
    },
    "tags": []
   },
   "source": [
    "## Data Locaility and Data Movement Bandwidth\n",
    "Previously, our work with MiniWeather utilized the `-gpu=managed` flag which setup a unified memory environment allowing us to not have to worry about data movement. However, unless the code has been designed to keep data resident on the GPU, __managed memory typically will not allow for optimal performance__.\n",
    "\n",
    "Recalling the design of heterogeneous GPU accelerated systems, the following diagram shows limits of data movement by the size of the arrows between each memory space."
   ]
  },
  {
   "cell_type": "markdown",
   "id": "1baf2b2a-0531-4f82-86ba-9e90fec14642",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "![GPU Memory Movement](img/GPUmemoryMove.png)\n",
    "\n",
    "* CPU to GPU I/O PCIe Bus: __15.75 GB/s__ for PCIe Gen3 V100s (31.5 GB/s for PCIe Gen4 capable A100s)\n",
    "* CPU to High Capacity DDR4-2666 Memory (Casper): __127.8 GB/s__ per CPU socket across 6 memory channels, 21.3 GB/s each\n",
    "* GPU to GPU via NVLink: __300 GB/s__ for V100s (600 GB/s for A100s)\n",
    "* GPU HBM2 (High Bandwidth Memory 2):  __900 GB/s__ for V100s (1,555 GB/s for the 40GB A100s)"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "3ec9935d-90d5-4a62-bfa2-5ca2461e6c34",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "## CPU to GPU Data Movement is Costly!\n",
    "* The bandwidth transfer rate for CPU to GPU data movement is the slowest in present day heterogeneous systems\n",
    "* Data transfers between the CPU and GPU should be minimized and ideally avoided\n",
    "\n",
    "Using OpenACC's directives `!$acc data`, `!$acc enter/exit data`, and `!$acc update ...` allow you to manage directly residency of data across the distinct CPU and GPU memory spaces."
   ]
  },
  {
   "cell_type": "markdown",
   "id": "563c0a2d-84af-4738-9464-8e5eb9b87ff4",
   "metadata": {
    "slideshow": {
     "slide_type": "slide"
    },
    "tags": []
   },
   "source": [
    "## Data Directives\n",
    "Reviewing `-Minfo=accel` output from last session, many data clauses were implicitly specified by the compiler despite using `-gpu=managed`. Data clauses are below:\n",
    "\n",
    "|OpenACC|OpenMP|Description|\n",
    "|---|---|---|\n",
    "|__Data Clauses__|   | Specifies data movement between CPU & GPU in parallel/data regions |\n",
    "| `create([zero]:vars)` | `alloc(vars)` | Allocates memory on target device for data object, optionally initializes to __zero__ values |\n",
    "| `copy(vars)` | `map(tofrom:vars)` | Allocates memory if needed and copies data at region entry/exit |\n",
    "| `copyin(vars)` | `map(to:vars)` | Allocates memory if needed and copies data at region entry |\n",
    "| `copyout([zero]:vars)` | `map(from:vars)` | Allocates memory if needed and copies data object at region exit, optionally initializes to __zero__ values |\n",
    "| `present(vars)` | `assert(omp_target_is_present(vars))` | Indicates that a data object is already present on GPU. Previous clauses include an implicit `present()` such that if true, that clause's action will not be performed |\n",
    "| `delete(vars)` | `dealloc(vars)` | Deallocates memory on target device for data object |\n",
    "| `deviceptr(vars)` | `is_device_ptr(vars)` | Declares device pointers such that data does not need to be moved/allocated |\n",
    "| `attach(vars)` | N/A | Increments the attachment counter for a pointer |\n",
    "| `detach(vars)` | N/A | Decrements the attachment counter for a pointer |\n",
    "| `finalize` | N/A | Sets structured or dynamic reference counter to 0 and forces action of `copyout()`, `detach()`, or `delete()` |"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "17dfebfc-bdfb-420e-a9ff-e7bd87969418",
   "metadata": {
    "slideshow": {
     "slide_type": "fragment"
    },
    "tags": []
   },
   "source": [
    "Each clause can reference a list of variables, ie `vars`=`var1,var2,...` or subsets of variable arrays. For example, `copy(a,b(10:20))` copies from CPU to GPU at region entry and from GPU to CPU at region exit both the variable `a` (full array by default) and only the slice of `b` indexed from 10 through 20."
   ]
  },
  {
   "cell_type": "markdown",
   "id": "2f6d2b02-8858-472e-8e05-36953ddad9de",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "## Scope of Data Directives and Data Regions\n",
    "\n",
    "All data clauses can only be used in the scope of a ...\n",
    "* __Structured Data Region__ or localized regions of code\n",
    "    * __Compute constructs regions__ like `!$acc kernels/parallel/serial ...` & `!$acc end kernels/parallel/serial`\n",
    "    * __Data constructs regions__ like `!$acc data ...` & `!$acc end data`\n",
    "    * __Implicit data regions__ of a function, subroutine, or program dependent on where `!$acc declare ...` directive is used\n",
    "    * Uses structured reference counters\n",
    "* __Unstructured Data Region__ or global program execution\n",
    "    * In a __data directive__ like `!$acc enter data ...` or `!$acc exit data ...`\n",
    "    * Uses dynamic reference counters"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "5326d5d2-5123-4898-a065-6b644779c4ea",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "__Unstructured Data Regions__ are preferred particularly when dealing with object oriented codes where the programmer may need to manage data locality across multiple modules, subroutines/functions, and source files.\n",
    "\n",
    "__Specific clauses can only be used in certain contexts__. For example:\n",
    "* `present()` cannot be used with `!$acc enter/exit data ...`\n",
    "* `delete()` can only be used with `!$acc exit data ...`\n",
    "* See specifications for each region construct in [OpenACC 2.7 Quick Reference Guide](https://www.openacc.org/sites/default/files/inline-files/API%20Guide%202.7.pdf) for more details\n",
    "\n",
    "Note: We will not cover `deviceptr()`, `attach()`, and `dettach()` in this notebook. These clauses are typically useful for CUDA libraries interoperaility and CUDA aware MPI with OpenACC."
   ]
  },
  {
   "cell_type": "markdown",
   "id": "f5f538b6-db2f-484c-8422-669292f38df5",
   "metadata": {
    "slideshow": {
     "slide_type": "slide"
    },
    "tags": []
   },
   "source": [
    "## Visualizing Data Directives with Data and Compute Regions\n",
    "\n",
    "This code and following visualizations highlight the data locality of variables across __CPU memory space, the black line,__ and the __GPU memory space, the green line__. To note, there is __always an implied data region with any compute construct__.\n",
    "\n",
    "```fortran\n",
    "!$acc parallel \n",
    "{!!! data region !!!\n",
    "    {!** parallel region **!\n",
    "    !$acc loop\n",
    "    do i = 1,1000\n",
    "        flux(i) = ...\n",
    "    end do\n",
    "    }!** end parallel region **!\n",
    "}!!! end data region !!!\n",
    "!$acc end parallel \n",
    "```"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "67bff34a-fa1e-4ddf-b5a5-4d6c25c8102c",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "## Implicit Data Region with a Parallel Region\n",
    "\n",
    "<img src=\"img/DataParLegend.png\" alt=\"Data Management Legend\" style=\"width:300px;\"/>\n",
    "\n",
    "Thanks go to Pierre-François Lavallée and Thibaut Véry from IDRIS, France for this and following images inspiraton. See their PRACE course [Introduction to OpenACC and OpenMP GPU](http://www.idris.fr/media/formations/openacc/gpu_directives.pdf) presented July 2019."
   ]
  },
  {
   "cell_type": "markdown",
   "id": "e500a90a-1675-4106-86c6-e55876084b7f",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "## OpenACC Program without Data Managment\n",
    "<img src=\"img/NoDataConstructs_repeats.png\" alt=\"Naive Data Movement\" style=\"width:600px;\"/>\n",
    "\n",
    "Without specifiying data movement, data transfers/allocations are performed by default at each compute region as determined by provided clauses or compiler (see `-Minfo=accel` output). For a program with many time steps, this could be 1,000s of unecessary transfers!"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "2e3bbbab-00eb-4565-a863-6f1f6d5f4820",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "## OpenACC Program with Structured Data Region\n",
    "<img src=\"img/OptimizedDataConstructs.png\" alt=\"Optimized Data Management\" style=\"width:600px;\"/>\n",
    "\n",
    "By specifying data movement over an encapsulating data region, data actions are done only at region entry and/or exit."
   ]
  },
  {
   "cell_type": "markdown",
   "id": "fca27369-5238-4f43-ad04-bac7399f8758",
   "metadata": {
    "slideshow": {
     "slide_type": "slide"
    },
    "tags": []
   },
   "source": [
    "## EXERCISE: Specify Variables in Data Regions around Time Step Loop\n",
    "Let's return to MiniWeather. We will now be using a new exercise source file [miniWeather_mpi_exercise2.F90](fortran/miniWeather_mpi_exercise2.F90#L128). This file has the `!$acc parallel ...` directives completed from the previous session. Now, we want to add a data region around the main time step loop section of the code near __Line 128__ (see `TODO: MANAGE GPU DATA`). This will look like:\n",
    "\n",
    "```fortran\n",
    "!$acc data ... ! Structured data region\n",
    "    !$acc kernels\n",
    "        ...\n",
    "    !$acc end kernels\n",
    "    call f_has_kernels()\n",
    "!$acc end data\n",
    "\n",
    "!!!!! OR !!!!!\n",
    "\n",
    "!$acc enter data ... ! Unstructured data directive\n",
    "    !$acc kernels\n",
    "        ...\n",
    "    !$acc end kernels\n",
    "    call f_has_kernels()\n",
    "!$acc exit data ... ! Unstructured data directive\n",
    "```\n",
    "\n",
    "Only consider usage of `copy()`, `copyin()`, `copyout()`, and `create()`. Hints are provided (`CTRL`/`CMD` + `F` `GPU_data`) to help define which clauses to use and variables you need to consider. Note: Variables listed in `GPU_data` sections are used across the `reductions()`, `output()`, and `perform_timestep()` subroutines called within the main time step loop. "
   ]
  },
  {
   "cell_type": "markdown",
   "id": "e7212e7b-c590-4367-a3b3-6b2b46491585",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "1. __Which clause do you use to ensure that data is availabile on the GPU and updates the same data on CPU after GPU work is done?__\n",
    "2. __Which clauses are optimal towards minimizing data movement in MiniWeather?__\n",
    "3. __If you used an unstructured data region, did this prevent data movement for the structured compute kernels? Why or why not?__\n",
    "\n",
    "Once you are satisfied with your changes, `make` and run the new executable and see how that has impacted performance."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "fae4a416-da63-4788-be5a-7e59b6eb28da",
   "metadata": {
    "slideshow": {
     "slide_type": "fragment"
    },
    "tags": []
   },
   "outputs": [],
   "source": [
    "make -C fortran/build openacc_test_ex2"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "94408b77-8c83-4d5e-ada1-54324d358ef1",
   "metadata": {
    "slideshow": {
     "slide_type": "fragment"
    },
    "tags": []
   },
   "outputs": [],
   "source": [
    "cd fortran/build\n",
    "qcmd -A $PROJECT -q $QUEUE -l select=1:ncpus=1:ngpus=1 -l gpu_type=$GPU_TYPE -l walltime=60 -v NVCOMPILER_ACC_TIME=1 -- \\\n",
    "$PWD/check_output.sh $PWD/openacc_test_ex2 1e-13 4.5e-5\n",
    "cd ../.."
   ]
  },
  {
   "cell_type": "markdown",
   "id": "f337de16-137b-4de6-a6a8-52c36d8da082",
   "metadata": {
    "slideshow": {
     "slide_type": "slide"
    },
    "tags": []
   },
   "source": [
    "## Data Directives - `!$acc update ...` and the `present()` Clause\n",
    "Unfortunately, the previous exercise introduced incorrect results. In some cases, you need to guarantee the movement of data between the host and device in order to ensure correct data is being processed. This uses the `!$acc update ...` directive:\n",
    "\n",
    "* `!$acc update host(...)` or `!$acc update self(...)` - Copies listed data variables from the GPU device to the host CPU.\n",
    "* `!$acc update device(...)` - Copies listed data variables from the host CPU to the device GPU.\n",
    "\n",
    "The `update` directive is useful if you are processing data on the host side after the data has been modified on the device side or vice versa. Note that `!$acc update ...` cannot be added within a compute construct.\n",
    "\n",
    "The `present()` clause will prevent an implicit data directive from being used on a parallel region. It makes clear to the compiler the availability of data on the device and only checks that data is present on the device. It also can make the code more readable to the programmer. Whether or not the `present()` data is in sync with the data on the CPU depends on recent computations or usage of `!$acc update ...`."
   ]
  },
  {
   "cell_type": "markdown",
   "id": "b80a796a-b8a9-4919-b249-e55395f7ff35",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "## Example Diagram for Data Movement Clauses\n",
    "![Example Data Movement](img/OpenACC_DataExample.png)"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "7891ab3e-8631-4c7a-8a13-53928b3ade24",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "## Reference Counters and Attachment Counters\n",
    "\n",
    "All data clauses also utilize a __reference counter__ for device variables or __attachment counter__ for device pointers. There are separate counters for __structured__ and __unstructured__ data regions. Only when a variable's counter increments from 0 or is reduced to 0 does a data action get performed. Example:\n",
    "\n",
    "1. Enter first data region with `copy(a,b)` clause. `a` & `b` copied from CPU to GPU, __reference counters__ increment 0 -> 1\n",
    "2. Kernel inside data region with `copy(a,b)` clause. No copy action but __reference counters__ increment 1 -> 2\n",
    "3. Kernel exits and no copy is performed. __Reference counters__ decrement from 2->1\n",
    "4. First data region exits, `a` & `b` copied out and deallocated on GPU, __reference counters__ decrement 1 -> 0.\n",
    "\n",
    "Essentially, counters track whether data is already _present_ or _not present_ on the GPU, ie _present_ = _true_ if _counter_ > 0 and informs data movement actions. This can avoid unnecessarily moving data at each kernel. To ensure a data action is performed, the use of `!$acc update ...` is recommended.\n",
    "\n",
    "Another option is adding `finalize` for unstructured data regions only, ie `!$acc exit data ...` with `delete()`, `detach()`, or `copyout()`. This forces the __reference counter__ to 0."
   ]
  },
  {
   "cell_type": "markdown",
   "id": "8b1fa4d5-d69c-4da2-aa75-38993dceb9cf",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "## OpenACC and Reference Counters\n",
    "<img src=\"img/OptimizedDataConstructs_references.png\" alt=\"Optimized Data Management with References\" style=\"width:600px;\"/>\n",
    "\n",
    "Even when compute regions specify data movement clauses,  data actions only happen when the counter increments from 0 -> 1 or decrements 1 -> 0. In this example, the variable `a` only has a data action performed on it at at data region entry and exit."
   ]
  },
  {
   "cell_type": "markdown",
   "id": "07202de7-95f5-436e-8da3-4401d5ad2740",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "## OpenACC and Update Directives\n",
    "<img src=\"img/OptimizedDataConstructs_update.png\" alt=\"Optimized Data Management with Updates\" style=\"width:600px;\"/>\n",
    "\n",
    "To ensure a data operation takes place, use `!$acc update host()` or `!$acc update device()` outside a parallel compute region. Another option: add `finalize` to clauses `copyout()`, `delete()`, or `detach()` in any structured or unstructured data region. `finalize` will set the reference counter to 0."
   ]
  },
  {
   "cell_type": "markdown",
   "id": "e8edaf4f-3007-44b0-9e70-4f003990888d",
   "metadata": {
    "slideshow": {
     "slide_type": "slide"
    },
    "tags": []
   },
   "source": [
    "## EXERCISE: Specify Update Directives for MPI Processes\n",
    "Again using the exercise source file [miniWeather_mpi_exercise2.F90](fortran/miniWeather_mpi_exercise2.F90#L440), we want to add a update directives at the following locations to manage the send and receive buffers as well as manage I/O output appropriately. Search for `TODO: UPDATE DATA` or go to:\n",
    "\n",
    "* MPI Buffers\n",
    "    - [] Line 448\n",
    "    - [] Line 460\n",
    "* Output I/O\n",
    "    - [] Line 828 - hints provided\n",
    "\n",
    "1. Look at the buffer variables that need to be passed between MPI tasks. __When do the data objects need to be updated on the host and device when doing an MPI send and MPI receive?__"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "ddfd86fd-953f-4798-8881-60d392d02a31",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "Once you are satisfied with your changes, `make` and run the new executable and review how that has impacted performance."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "9ea87343-c20e-4af2-aa45-5b1061c7dfd4",
   "metadata": {
    "slideshow": {
     "slide_type": "fragment"
    },
    "tags": []
   },
   "outputs": [],
   "source": [
    "make -C fortran/build openacc_test_ex2"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "026171bb-7a88-49d7-8ddd-fd03278b840d",
   "metadata": {
    "slideshow": {
     "slide_type": "fragment"
    },
    "tags": []
   },
   "outputs": [],
   "source": [
    "cd fortran/build\n",
    "qcmd -A $PROJECT -q $QUEUE -l select=1:ncpus=1:ngpus=1 -l gpu_type=$GPU_TYPE -l walltime=60 -v NVCOMPILER_ACC_TIME=1 -- \\\n",
    "$PWD/check_output.sh $PWD/openacc_test_ex2_update 1e-13 4.5e-5\n",
    "cd ../.."
   ]
  },
  {
   "cell_type": "markdown",
   "id": "fe31dbea-b3af-43a1-9fb7-1b6efc596e0a",
   "metadata": {
    "slideshow": {
     "slide_type": "slide"
    },
    "tags": []
   },
   "source": [
    "## Asynchronous Execution with `async()` and `wait()` Clauses\n",
    "One of the more important advanced OpenACC features is __asynchronous execution__. This allows you to interface with __CUDA Streams__ and manage directly the scheduling queues the GPU uses to run kernels on the device in a defined sequence.\n",
    "\n",
    "* __Synchronous Execution__ - Enabled by default for all directives\n",
    "    * Host must wait for parallel compute regions and data regions/updates to complete\n",
    "    * All operations and parallel task units must run in a serialized sequence\n",
    "* __Asynchronous Execution__ - Must be specified for every compatible directive\n",
    "    * Host can perform work while the device GPU is also performing work\n",
    "    * Data transfers can be scheduled and start processing before the data is needed while other work is performed\n",
    "    * Multiple parallel compute units can be scheduled on device to run in an uninterrupted pipeline. Can avoid kernel startup/spindown scheduling costs before being forced to wait/synchronize"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "d608663b-9ee5-474a-bb44-0cf00609b7e4",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "So far, the CPU and GPU have performed __synchronous execution__. The CPU must wait to synchronize with the completion of compute or data movement work on the GPU. This often creates time gaps between compute kernels on the GPU (say during data movement or kernel initialization) where the hardware is not being utilized.\n",
    "\n",
    "<img src=\"img/async_Pipelining.png\" alt=\"Asynchronous Pipelining\" style=\"width:600px;\"/>\n",
    "\n",
    "If multiple asynchronous queues are used and there are available compute resources, it's also possible to overlap compute kernels with each other in addition to data movement kernels. __Serialized execution of parallel kernels is not required.__"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "699909e5-7d70-4506-ae2d-43e5b063b908",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "## Usage of `async()` and `wait()` Clauses\n",
    "The `async` clause may appear on a `parallel`, `serial`, `kernels`, or `data` construct, or an `enter data`, `exit data`, `update`, or `wait` directive. The `wait` clause can appear on the same (except itself).\n",
    "\n",
    "* `async(n)` - Launches work asynchronously in queue `n`. If `n` is omitted, the default queue is selected.\n",
    "* `wait(n,m,...)` - Blocks host (or blocks queue `m` if paired with `async(m)`) until all prior operations in queues `n,m,...` have completed. If an argument is omitted, all queues must complete.\n",
    "\n",
    "The optional argument `n` must be an integer value. An essentially limitless number of queues/streams are able to be created. However, only 48 CUDA stream contexts (16 pre-Volta) will be allowed to run concurrently.\n",
    "\n",
    "Complex compute kernel task dependencies can be designed using clever queue scheduling to fully saturate the GPU. However, limited time does not permit us to review this here. Instead, please refer to section 7.1 of the [OpenACC Best Practices Programming Guide, May 2021](../reference/OpenACC-Best-Practices-Programming-Guide_May2021.pdf) or more briefly in Steve Abbott's [Advanced OpenACC Lecture Slides](http://icl.cs.utk.edu/classes/cosc462/2017/pdf/OpenACC_3.pdf) given at University of Tennessee's Innovation Computing Laboratory."
   ]
  },
  {
   "cell_type": "markdown",
   "id": "8754d997-1989-44bf-aa0f-2558058356b9",
   "metadata": {
    "slideshow": {
     "slide_type": "slide"
    },
    "tags": []
   },
   "source": [
    "## EXERCISE: Specify `async` Compute Kernels and Place `wait` Barriers in MiniWeather\n",
    "Again using the exercise source file [miniWeather_mpi_exercise2.F90](fortran/miniWeather_mpi_exercise2.F90#L440), we want to add `async` and `wait` at appropriate places in the code. Search for `TODO: ASYNC ME` or go to:\n",
    "\n",
    "- [] Line 251\n",
    "- [] Line 304\n",
    "- [] Line 337\n",
    "- [] Line 366\n",
    "- [] Line 404\n",
    "- [] Line 435\n",
    "- [] Line 449\n",
    "- [] Line 461\n",
    "- [] Line 481\n",
    "- [] Line 503"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "a3c7f3ae-0c65-4ed6-b9b7-92dce8c9cb9b",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "1. __Is it possible to set any of MiniWeather's compute kernels in thier own queues or do dependencies between kernels prevent this?__\n",
    "2. How did __asynchronous execution__ impact performance? __Review previous incremental changes to MiniWeather and record below the change in CPU time.__ (Note: Select cell and press `M` or double click this cell to edit the table)\n",
    "3. Enabling asynchronous execution only improved performance slightly. __What additional factors are likely contributing to this modest improvement?__ This is best answered using a profiler via `nsys profile -t openacc ./openacc_test_ex2` but feel free to speculate.\n",
    "\n",
    "| MiniWeather Edits | CPU Time (s) |\n",
    "|---|---|\n",
    "| BaseLine (on V100) |               26.1405 |\n",
    "| Data regions (d_mass wrong) |      XX      |\n",
    "| Update directives (d_mass fixed) | XX      |\n",
    "| Async and wait |                   XX      |"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "285851df-4ad9-499b-9c1f-a68511750c1c",
   "metadata": {
    "slideshow": {
     "slide_type": "fragment"
    },
    "tags": []
   },
   "outputs": [],
   "source": [
    "make -C fortran/build openacc_test_ex2"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "42c5f762-159e-4d00-9bfd-95310516827e",
   "metadata": {
    "slideshow": {
     "slide_type": "fragment"
    },
    "tags": []
   },
   "outputs": [],
   "source": [
    "cd fortran/build\n",
    "qcmd -A $PROJECT -q $QUEUE -l select=1:ncpus=1:ngpus=1 -l gpu_type=$GPU_TYPE -l walltime=60 -v NVCOMPILER_ACC_TIME=1 -- \\\n",
    "$PWD/check_output.sh $PWD/openacc_test_ex2 1e-13 4.5e-5\n",
    "cd ../.."
   ]
  },
  {
   "cell_type": "markdown",
   "id": "64b44cdf-ed4e-4fce-8585-d69942e75ddc",
   "metadata": {
    "slideshow": {
     "slide_type": "slide"
    },
    "tags": []
   },
   "source": [
    "#### __Remaining Concepts Will Be Discussed with Time Remaining or Provided for Independent Review__\n",
    "Some great discussions on these advanced concepts are covered in John Urbanic's June 2021 GPU Programming Lectures [Advanced OpenACC](https://www.psc.edu/wp-content/uploads/2021/06/Advanced_OpenACC.pdf) and [Using OpenACC With CUDA Libraries](https://www.psc.edu/wp-content/uploads/2021/06/OpenACC_Using_OpenACC_with_CUDA-Libraries.pdf). Nonetheless, most optimizations using OpenACC at this level should include profiling as part of any __development cycle__. We will cover usage of the Nsight Systems and Nsight Compute profilers at later sessions.\n",
    "\n",
    "<img src=\"img/development-cycle.png\" alt=\"OpenACC Development Cycle\" style=\"width:500px;\"/>"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "020300ce-176c-4d57-a101-516147f77a25",
   "metadata": {
    "slideshow": {
     "slide_type": "slide"
    },
    "tags": []
   },
   "source": [
    "## Routines and Functions Called from a GPU Kernel - `!$acc routine ...` Directive\n",
    "Compute regions are able to call external functions from within the kernel. However, the external function needs to be declared to the compiler that it is a GPU kernel.\n",
    "\n",
    "This is done using the `!$acc routine ...` directive. This directive should be placed in the specification part of the routine it should apply to or you may place it in the calling routine's/module's specification section as `!$acc routine(name)` where `name` is the routine which it applies. Required with any `!$acc routine ...` directive is one of four clauses:\n",
    "\n",
    "* `seq` - Specifies the routine should run in `seq` mode and is not parallel, should be run sequentially\n",
    "* `vector` - Specifies the routine should run in `vector` mode and contains a `!$acc loop vector` parallel construct or calls another vector routine\n",
    "* `worker` - Specifies the routine should run in `worker` mode and contains a `!$acc loop worker` parallel construct or calls another worker routine\n",
    "* `gang` - Specifies the routine should run in `gang` mode and contains a `!$acc loop gang` parallel construct or calls another gang routine\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "96907e6a-5033-48c8-8205-8ef10f099406",
   "metadata": {
    "slideshow": {
     "slide_type": "slide"
    },
    "tags": []
   },
   "source": [
    "## EXERCISE: Convert Inlined `sample_ellipse_cosine()` to Function Call\n",
    "\n",
    "In [miniWeather_mpi_exercise2.F90](fortran/miniWeather_mpi_exercise2.F90#L266), you can try specifying a GPU routine by searching for `TODO: TRY GPU ROUTINE OPTION` or going to __Lines 264 and 769__. As a comment, compilers often have non-optimal and sometimes even broken support for `!$acc routine ...`. Thus, perhaps consider instead inlining required routines into each relevant parallel region. "
   ]
  },
  {
   "cell_type": "markdown",
   "id": "c9754e1c-db15-4624-af7d-e7b0fe931f07",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "`make` and run the new executable and evaluate if this changed performance in any way.\n",
    "\n",
    "1. __Which parallel mode (`seq`, `vector`, `worker`, or `gang`) is required to be set for the function `sample_ellipse_cosine`?__\n",
    "2. __Do you see any other algorithms/groups of operations, perhaps at a `vector` level or higher, you might write into its own subroutine/function? Try it out in MiniWeather!__"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "0d7bf017-aa51-4a12-b91e-1a73ccccfd9b",
   "metadata": {
    "slideshow": {
     "slide_type": "fragment"
    },
    "tags": []
   },
   "outputs": [],
   "source": [
    "make -C fortran/build openacc_test_ex2"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "7232cccb-cd86-4f73-a2b7-5d91a1711f0d",
   "metadata": {
    "slideshow": {
     "slide_type": "fragment"
    },
    "tags": []
   },
   "outputs": [],
   "source": [
    "cd fortran/build\n",
    "qcmd -A $PROJECT -q $QUEUE -l select=1:ncpus=1:ngpus=1 -l gpu_type=$GPU_TYPE -l walltime=60 -v NVCOMPILER_ACC_TIME=1 -- \\\n",
    "$PWD/check_output.sh $PWD/openacc_test_ex2 1e-13 4.5e-5\n",
    "cd ../.."
   ]
  },
  {
   "cell_type": "markdown",
   "id": "1d95284e-c732-459a-adf5-a431e1d92baf",
   "metadata": {
    "slideshow": {
     "slide_type": "slide"
    },
    "tags": []
   },
   "source": [
    "## OpenACC API Routines\n",
    "OpenACC provides access to runtime library routines via `use openacc` module in FORTRAN (or `#include <openacc.h>` header in C/C++). This allows executing many features of OpenACC without directives, as well as a few important features not available through directives. Two important routines are below. These are most relevant for __Multi-GPU Programming__ using __CUDA aware MPI__ and will be discussed in future sessions.\n",
    "\n",
    "1. `acc_get_num_devices(acc_device_t dev_type)` - Returns the integer value of the number of devices connected to the host. \n",
    "    * Use `dev_type`=`acc_device_nvidia` to enumerate NVIDIA devices.\n",
    "2. `acc_set_device_num(int idev, acc_device_t dev_type)` - Sets for the runtime which device to use, where 0 <= `idev` < `dev_num`.\n",
    "\n",
    "See page 15 of the [OpenACC API Guide, v2.7](../reference/OpenACC-API-Guide_v2.7.pdf) or page 89 of [OpenACC Full Specification, v3.2](../reference/OpenACC-Full-Specification_v3.2.pdf) for a complete description of all possible API routines."
   ]
  },
  {
   "cell_type": "markdown",
   "id": "6e0a8f3c-f0ce-4dcf-8120-87f60d3259c2",
   "metadata": {
    "slideshow": {
     "slide_type": "slide"
    },
    "tags": []
   },
   "source": [
    "## OpenACC with CUDA and GPU Accelerated Libraries\n",
    "There is functionality to use manually developed CUDA kernels or CUDA libraries alongside OpenACC. This is particularly useful if you have decided to develop optimal CUDA code for a particularly expensive kernel or need to interface with a highly performant CUDA Accelerated Library such as cuBLAS or cuFFT. Note however this will likely impact the portability of your code with non-NVIDIA devices.\n",
    "\n",
    "This use case is effectively explained in John Urbanic's [Using OpenACC With CUDA Libraries](https://www.psc.edu/wp-content/uploads/2021/06/OpenACC_Using_OpenACC_with_CUDA-Libraries.pdf) presentation. Nonetheless, below are the most important concepts to remember:\n",
    "\n",
    "* `deviceptr(vars)` - Declares in a data or compute construct that `vars` already refer to device pointers and need not be allocated or moved in device memory.\n",
    "    * Allows using `vars` within a compute construct that typically originated from a `cudaMalloc()` routine and were used in a CUDA kernel or CUDA library.\n",
    "* `!$acc host_data use_device(vars)` - Specifies within a host compute region that `vars` already exist on the device and replaces references to `vars` with their respective device pointer. \n",
    "    * Allows calling a CUDA kernel or library from the host outside a parallel construct."
   ]
  },
  {
   "cell_type": "markdown",
   "id": "97d54cd9-5ec2-47e3-9b90-cf3a368cfae1",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "## OpenACC cuTENSOR Example\n",
    "\n",
    "Here is an example use case for the cuTENSOR library using OpenACC. The complete FORTRAN can be found in NVIDIA's HPC SDK Documentation [Using cuTENSOR from OpenACC Host Code](https://docs.nvidia.com/hpc-sdk/compilers/fortran-cuda-interfaces/#cflib-tensor-oacc-host). Notably, this code is portable with and without OpenACC using `!@acc`.\n",
    "```fortran\n",
    "!@acc use openacc\n",
    "!@acc use cutensorex\n",
    "integer, parameter :: ni=1280, nj=1024, nk=960, ntimes=1\n",
    "real(8) :: a(ni,nk), b(nk,nj), c(ni,nj), d(ni,nj)\n",
    "\n",
    "call random_number(a)\n",
    "call random_number(b)\n",
    "a = dble(int(4.0d0*a - 2.0d0))\n",
    "b = dble(int(8.0d0*b - 4.0d0))\n",
    "c = 2.0; d = 0.0\n",
    "\n",
    "!$acc enter data copyin(a,b,c) create(d)\n",
    "!@acc istat = cutensorExSetStream(acc_get_cuda_stream(acc_async_sync))\n",
    "!$acc host_data use_device(a,b,c,d)\n",
    "do nt = 1, ntimes\n",
    "  d = c + matmul(a,b)\n",
    "end do\n",
    "!$acc end host_data\n",
    "!$acc update host(d)\n",
    "```"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "6f6ee01d-0133-449a-b827-ad2efddbae4a",
   "metadata": {
    "slideshow": {
     "slide_type": "slide"
    },
    "tags": []
   },
   "source": [
    "## Advanced Optimizations with `cache`, `tile`, and `gang/worker/vector` Clauses\n",
    "The compiler typically does a good job choosing close to optimal arrangements of `gang`, `worker`, and `vector` parallel execution. You can review compiler choices in `-Minfo=accel` output, basic profiling via `NVCOMPILER_ACC_TIME=1`, or NSight Systems/Compute profiling. However, further optimization can be explored by manualy specifying optimization parameters. __This process requires significant experimentation with some intuition__. Clauses not previously discussed are detailed below:\n",
    "\n",
    "* `cache()` - When included inside and at the top of a loop, specifies array elements or subarrays that should be fetched into the highest level of cache, ie the shared memory of the SM assigned to each gang.\n",
    "* `tile(n,m,...)` - Splits loops into two loops, outer _tile_ loops and inner _element_ loops. Listed integer arguments specify the size of each tile and correspond to each tightly nested loop. The first `n` corresponds to the innermost loop. If `n`=`*`, the compiler has freedom to choose tile size.\n",
    "    * A `vector` clause is applied to _element_ loops \n",
    "    * A `gang` clause is applied to _tile_ loops\n",
    "    * A `worker` clause is applied to _element_ loops or to _tile_ loops if the `vector` clause is also present"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "826c6583-950d-4c9c-9381-6d941547a6b6",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "### Usage of `cache()`\n",
    "If you have an array or subarray of data to be accessed frequently within a compute kernel, the `cache()` directive can direct the compiler to keep that portion of memory nearby in a fast L1 shared memory cache for each gang. The size of this cache is limited and dependent on the hardware (V100 - 96kB, A100 - 160kB) but can increase performance if configured correctly. Example below:\n",
    "```fortran\n",
    "!$acc parallel loop gang vector\n",
    "do i = 2, n-1\n",
    "  !$acc cache(a(i−1:i+1))\n",
    "  lower = i−1\n",
    "  upper = i+1\n",
    "  sum = 0;\n",
    "  !$acc loop seq\n",
    "  do j = lower, upper\n",
    "    sum = sum + a(i);\n",
    "  end do\n",
    "end do\n",
    "!$acc end parallel\n",
    "```\n",
    "Feel free to review Lashgar's and Baniasadi's [Efficient Implementation of OpenACC cache Directive on NVIDIA GPUs](http://www.ahmado.com/profile/lashgar/files/17ijhpcn.pdf) in 2017 IJHPCN for a more in depth discussion from a compiler perspective."
   ]
  },
  {
   "cell_type": "markdown",
   "id": "563304a7-81ba-4d8e-889f-f2ce0e10bbf6",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "### Target Specific Devices when Specifying Parallelism\n",
    "When tuning performance using the following clauses, OpenACC should target specific hardware or hardware types. To promote portability, consider pairing these clauses with `device_type(d_type)` where `d_type`=`acc_device_nvidia`,`acc_device_radeon`,`acc_device_host`, etc depending on what OpenACC specification is implemented and what the compiler recognizes. \n",
    "\n",
    "For example, `!$acc loop device_type(acc_device_nvidia) vector_length(1024) device_type(acc_device_radeon) vector_length(128)` will apply distinct clauses for NVIDIA and AMD devices respectively.\n",
    "\n",
    "If neither device type is attached, then the compiler chooses its own arrangement of levels of parallelism."
   ]
  },
  {
   "cell_type": "markdown",
   "id": "1eeb89c3-674b-46a8-8897-694d5a94c3c6",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "### Usage of `tile()`\n",
    "For multidimensional loops, `tile()` splits loops into blocks to distribute work across a device. This is essentially opposite to `collapse()` which exposes parallelism and unrolls loops into the largest possible units. \n",
    "\n",
    "Where `collapse()` can often negatively impact data locality, `tile()` should be used to increase data locality or encourage __memory coalescing__, improving performance in some cases. Lastly, tiles can be executed simultaneously across distinct gangs."
   ]
  },
  {
   "cell_type": "markdown",
   "id": "5648d26d-8520-4e9b-a795-964ab8f731df",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "<img src=\"img/tile.png\" alt=\"OpenACC Tile\" style=\"width:400px;\"/>\n",
    "\n",
    "```fortran\n",
    "!$acc kernels loop tile(2,2)\n",
    "do x = 1, 4\n",
    "  do y = 1, 4\n",
    "    a(x,y) = a(x,y) + 1\n",
    "  end do\n",
    "end do\n",
    "```"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "9229aad0-518c-479b-a73a-b2653f7973b4",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "### Usage of `gang/worker/vector`\n",
    "Modifying levels of parallelism is more an art than a science, requiring intuition and experimentation to find the optimal arrangement for a given kernel, which also depends on kernel problem size. __Sometimes, refactoring the kernel algorithm is recommended instead__. Here are some main points:\n",
    "\n",
    "* `gang` parallism is often best for outer loops and `vector` for inner loops.\n",
    "* `vector` should be of lengths 32 or multiples of 32 (NVIDIA) since each SIMT warp is length 32.\n",
    "* The number/length of parallelism level should correlate with expected size of a loop and should not be greater than the number of steps in that loop.\n",
    "* `num_workers` * `vector_length` is the number of threads in each __CUDA block__, max 1,024 (NVIDIA).\n",
    "* More often than not, it is best to let the compiler decide the number of gangs.\n",
    "* Effective tuning can often be done by only adjusting the `vector_length`.\n",
    "* Usage of a `worker` loop is very rare and often reserved for when `vector_length` is small, ie less than 128.\n",
    "* __Coallescing memory access__ in `vector` loops is likely most important consideration, ie threads access contiguous memory elements in matrix collumns in FORTRAN (rows in C/C++)."
   ]
  },
  {
   "cell_type": "markdown",
   "id": "d433d1e4-4274-48df-b3ff-c2220063e3c7",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "### Coallesced Memory Example in FORTRAN\n",
    "```fortran\n",
    "!$acc parallel         ! j-i ordering allows coallesced memory access in FORTRAN\n",
    "!$acc loop gang\n",
    "do j = 1, ny\n",
    "    !$acc loop vector\n",
    "    do i = 1, nx\n",
    "        A(i,j) = B(i,j) + C(i,j)\n",
    "    end do\n",
    "end do\n",
    "!$acc end parallel\n",
    "```\n",
    "The below example is uncoalesced after switching the `i/j` in the loops (this is opposite in C/C++):\n",
    "```fortran\n",
    "!$acc parallel         ! i-j ordering DOES NOT allow coallesced memory access in FORTRAN\n",
    "!$acc loop gang\n",
    "do i = 1, nx\n",
    "    !$acc loop vector\n",
    "    do j = 1, ny\n",
    "        A(i,j) = B(i,j) + C(i,j)\n",
    "    end do\n",
    "end do\n",
    "!$acc end parallel\n",
    "```"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "9c22c3c0-47d7-4968-8fa7-9a89fffbd4bf",
   "metadata": {
    "slideshow": {
     "slide_type": "slide"
    },
    "tags": []
   },
   "source": [
    "## EXERCISE - Manually Tune Performance of Select MiniWeather Kernels\n",
    "Try out some of the previous performance tuning directives in [miniWeather_mpi_exercise2.F90](fortran/miniWeather_mpi_exercise2.F90). \n",
    "\n",
    "It is recommended to modify user parameters such that `_NX=1024`, `_NZ=512`, and `_SIM_TIME=10` either: \n",
    "1. Near the head of the [CMakeLists.txt](fortran/CMakeLists.txt). Compilations of test executables WILL NOT BE affected.\n",
    "2. At __[Line 57](fortran/miniWeather_mpi_exercise2.F90#L57)__ of the exercise 2 source file. Compilations of test executables WILL BE affected.\n",
    "\n",
    "This sets up a sufficiently large problem to fill GPU SMs and the short simulation time promotes rapid profile driven development. You might also try even larger domain sizes but runtime will increase."
   ]
  },
  {
   "cell_type": "markdown",
   "id": "20df9157-4974-4398-9cc4-b6f81cb5e031",
   "metadata": {
    "slideshow": {
     "slide_type": "subslide"
    },
    "tags": []
   },
   "source": [
    "__Focus on just one kernel at a time__ (choose a random kernel or select the most costly kernel given timing output via `NVCOMPILER_ACC_TIME=1` in last few exercises) and __note below the line number of the kernel__. `make` and run the new executable and record how performance changes with each modification.\n",
    "\n",
    "1. __Can you beat the performance of the `collapse()` clause and the automatic choices made by the compiler?__\n",
    "2. __Share your results in Slack (specify kernel you chose) and see what others were able to achieve.__\n",
    "\n",
    "| MiniWeather Kernel L### | CPU Time (s) |\n",
    "|---|---|\n",
    "| BaseLine (on V100) - collapse() | XX      |\n",
    "| clause - gang/vector            | XX      |\n",
    "| clause - tile(#,#,#)            | XX      |\n",
    "| clause - tile(\\*,\\*,\\*)         | XX      |\n",
    "| clause - vector_length(XX)      | XX      |\n",
    "| ...                             | XX      |"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "39f67ef9-cdf2-4e6f-a99e-8ac34af97420",
   "metadata": {
    "slideshow": {
     "slide_type": "fragment"
    },
    "tags": []
   },
   "outputs": [],
   "source": [
    "make -C fortran/build openacc_ex2"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "eb6005f9-0201-4fa6-9eef-86e85b8ceb80",
   "metadata": {
    "slideshow": {
     "slide_type": "fragment"
    },
    "tags": []
   },
   "outputs": [],
   "source": [
    "cd fortran/build\n",
    "qcmd -A $PROJECT -q $QUEUE -l select=1:ncpus=1:ngpus=1 -l gpu_type=$GPU_TYPE -l walltime=60 -v NVCOMPILER_ACC_TIME=1 -- \\\n",
    "$PWD/check_output.sh $PWD/openacc_ex2 1e-13 4.5e-5\n",
    "cd ../.."
   ]
  },
  {
   "cell_type": "markdown",
   "id": "b2ed5230-9e44-49a4-9ec2-c30662474704",
   "metadata": {
    "slideshow": {
     "slide_type": "slide"
    },
    "tags": []
   },
   "source": [
    "## Verify Correctness and Display MiniWeather Output\n",
    "Results of the program may be viewed by initiating a ssh X session with Casper on a terminal `ssh -Y [username]@casper.ucar.edu` then running `module load ncview` and `ncview output.nc` on the output file that should now be in your `$HOME` directory or the folder where you ran the MiniWeather executable from.\n",
    "\n",
    "Alternatively, use the Python script [ncplot_MiniWeather.ipynb](ncplot_MiniWeather.ipynb) to display and save image/GIF results within JupyterHub.\n",
    "\n",
    "Visualization can be an useful tool to verify correctness of recent changes to the code as well. However, if you want a quantitative method to verify correctness, we will discuss the use of [PCAST](https://developer.nvidia.com/blog/detecting-divergence-using-pcast-to-compare-gpu-to-cpu-results/), Parallel Compiler Assisted Software Testing, in a later session. This compares results between GPU and CPU versions of the code running simultaneously or against a \"golden\" baseline file."
   ]
  },
  {
   "cell_type": "markdown",
   "id": "86d33a14-9f6a-40e9-a8af-fec71068d73f",
   "metadata": {
    "slideshow": {
     "slide_type": "slide"
    },
    "tags": []
   },
   "source": [
    "## Suggested Resources\n",
    "* Matt Norman's [A Practical Introduction to GPU Refactoring in FORTRAN with Directives for Climate](https://github.com/mrnorman/miniWeather/wiki/A-Practical-Introduction-to-GPU-Refactoring-in-Fortran-with-Directives-for-Climate)\n",
    "* May 2021, [OpenACC Programming and Best Practices Guide](../reference/OpenACC-Best-Practices-Programming-Guide_May2021.pdf) and [Github](https://github.com/OpenACC/openacc-best-practices-guide)\n",
    "* [OpenACC 2.7 Quick Reference Guide](../reference/OpenACC-API-Guide_v2.7.pdf)\n",
    "* Official [OpenACC 3.2 Full Standard Specification](../reference/OpenACC-Full-Specification_v3.2.pdf) - Not all updated features are implemented yet by compatible compilers\n",
    "* If you want to dive deep into lower level control and optimization of GPU performance, check out Oak Ridge National Lab's [CUDA Training Series](https://olcf.ornl.gov/cuda-training-series/https://olcf.ornl.gov/cuda-training-series/)."
   ]
  }
 ],
 "metadata": {
  "kernelspec": {
   "display_name": "Bash",
   "language": "bash",
   "name": "bash"
  },
  "language_info": {
   "codemirror_mode": "shell",
   "file_extension": ".sh",
   "mimetype": "text/x-sh",
   "name": "bash"
  }
 },
 "nbformat": 4,
 "nbformat_minor": 5
}