energy/01_Code_Optimizations_for_Energy.ipynb

{
  "cells": [
    {
      "cell_type": "markdown",
      "id": "7d24968e",
      "metadata": {
        "id": "7d24968e"
      },
      "source": [
        "# Energy Efficient GPU Computing - Tutorial\n",
        "\n",
        "## Hands-on: Code Optimizations\n",
        "\n",
        "Welcome to the second hands-on of the Energy Efficient GPU Computing tutorial. In this hands-on exercise, we will learn about code optimizations and look at their impact on energy efficiency.\n",
        "\n",
        "We again use Kernel Tuner to easily benchmark different implementations of our GPU kernels. To make it available in this notebook, run the following cell:"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "id": "a8bfafdb",
      "metadata": {
        "id": "a8bfafdb"
      },
      "outputs": [],
      "source": [
        "%%capture\n",
        "%pip install kernel-tuner[tutorial]==1.0.0b5"
      ]
    },
    {
      "cell_type": "markdown",
      "id": "cef70e1d",
      "metadata": {
        "id": "cef70e1d"
      },
      "source": [
        "And then import the modules we need:"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "id": "fdfe0872",
      "metadata": {
        "id": "fdfe0872"
      },
      "outputs": [],
      "source": [
        "import numpy as np\n",
        "import kernel_tuner as kt\n",
        "from kernel_tuner.observers.nvml import NVMLObserver\n",
        "from kernel_tuner.util import get_best_config"
      ]
    },
    {
      "cell_type": "markdown",
      "id": "18628d44",
      "metadata": {
        "id": "18628d44"
      },
      "source": [
        "The main example of a code optimization that we'll look at in this hands-on exercise is a simple case where we can apply *Kernel Fusion*. Kernel Fusion is the process of merging or 'fusing' multiple kernels into a single kernel.\n",
        "\n",
        "The operation that we would like to perform is to compute a dot product of two vectors a and b of equal length. The computation consists of two steps. First, all elements in a and b are point-wise multiplied, then these values are summed together to a single value. These two computations can be implemented by different kernels, or as we will see, fused into one.\n",
        "\n",
        "For our first CUDA kernel, let's take a look at a kernel that applies a point-wise multiplication of two vectors:"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "id": "f1c332c5",
      "metadata": {
        "id": "f1c332c5"
      },
      "outputs": [],
      "source": [
        "%%writefile vector_mul_kernel.cu\n",
        "\n",
        "__global__ void vector_mul(float * c, float * a, float * b, int n) {\n",
        "    int i = (blockIdx.x * blockDim.x) + threadIdx.x;\n",
        "    if (i < n) {\n",
        "        c[i] = a[i] * b[i];\n",
        "    }\n",
        "}"
      ]
    },
    {
      "cell_type": "markdown",
      "id": "32234e49",
      "metadata": {
        "id": "32234e49"
      },
      "source": [
        "To compute a full dot product of two vectors, we also have to sum the result. So let's take a simple CUDA kernel to compute the sum of a large vector."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "id": "a8a1309d",
      "metadata": {
        "id": "a8a1309d"
      },
      "outputs": [],
      "source": [
        "%%writefile sum_kernel.cu\n",
        "\n",
        "__global__ void sum(float *result, float *X, int n) {\n",
        "    __shared__ float cache[block_size_x];\n",
        "    int i = threadIdx.x + blockIdx.x * blockDim.x;\n",
        "    float temp = 0.0f;\n",
        "    for (; i<n; i+= blockDim.x * gridDim.x) {\n",
        "        temp += X[i];\n",
        "    }\n",
        "    cache[threadIdx.x] = temp;\n",
        "    __syncthreads();\n",
        "    for (int s=block_size_x/2; s>0; s/=2) {\n",
        "        if (threadIdx.x < s) {\n",
        "            cache[threadIdx.x] += cache[threadIdx.x + s];\n",
        "        }\n",
        "        __syncthreads();\n",
        "    }\n",
        "    if (threadIdx.x == 0) {\n",
        "        atomicAdd(result, cache[0]);\n",
        "    }\n",
        "}"
      ]
    },
    {
      "cell_type": "markdown",
      "id": "f91c1693",
      "metadata": {
        "id": "f91c1693"
      },
      "source": [
        "Now, we use Kernel Tuner to measure the execution time and energy consumption of both kernels."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "id": "639fae3e",
      "metadata": {
        "id": "639fae3e"
      },
      "outputs": [],
      "source": [
        "# the size of the vectors\n",
        "size = 100000000\n",
        "\n",
        "# all the kernel input and output data need to use numpy data types,\n",
        "# note that we explicitly state that these arrays should consist of\n",
        "# 32 bit floating-point values, to match our kernel source code\n",
        "a = np.random.randn(size).astype(np.float32)\n",
        "b = np.random.randn(size).astype(np.float32)\n",
        "c = np.zeros_like(b)\n",
        "r = np.zeros(1, dtype=np.float32)\n",
        "n = np.int32(size)\n",
        "\n",
        "# Setup the tunable parameters\n",
        "tune_params = dict(block_size_x=[16, 32, 64, 128, 256, 512, 1024])\n",
        "\n",
        "# Setup the NVMLObserver to measure energy\n",
        "nvmlobserver = NVMLObserver([\"nvml_energy\"])\n",
        "metrics = dict(energy=lambda p:p[\"nvml_energy\"])\n",
        "\n",
        "# Let's call tune_kernel to start the tuning process.\n",
        "results_mul, env = kt.tune_kernel(\"vector_mul\", \"vector_mul_kernel.cu\", size, [c, a, b, n], tune_params, lang=\"cupy\",\n",
        "                                  observers=[nvmlobserver], metrics=metrics)\n",
        "\n",
        "# Let's call tune_kernel to start the tuning process.\n",
        "results_sum, env = kt.tune_kernel(\"sum\", \"sum_kernel.cu\", size, [r, c, n], tune_params, lang=\"cupy\",\n",
        "                                 observers=[nvmlobserver], metrics=metrics)"
      ]
    },
    {
      "cell_type": "markdown",
      "id": "7a8717ed",
      "metadata": {
        "id": "7a8717ed"
      },
      "source": [
        "Now if we create a single kernel that in one go applies a point-wise multiplication followed by a sum reduction, we effectively implement a dot product kernel. This kernel can be seen as the fusion of the previous two kernels.\n",
        "\n",
        "**TODO**: Complete the missing line of code in the kernel below! Then, run the cell to store the code to a file."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "id": "95d4c3aa",
      "metadata": {
        "id": "95d4c3aa"
      },
      "outputs": [],
      "source": [
        "|%%writefile dot_kernel.cu\n",
        "\n",
        "__global__ void dot(float *result, float *a, float *b, int n) {\n",
        "    __shared__ float cache[block_size_x];\n",
        "    int i = threadIdx.x + blockIdx.x * blockDim.x;\n",
        "    float temp = 0.0f;\n",
        "    for (; i<n; i+= blockDim.x * gridDim.x) {\n",
        "        temp += ...; //TODO: write code here\n",
        "    }\n",
        "    cache[threadIdx.x] = temp;\n",
        "    __syncthreads();\n",
        "    for (int s=block_size_x/2; s>0; s/=2) {\n",
        "        if (threadIdx.x < s) {\n",
        "            cache[threadIdx.x] += cache[threadIdx.x + s];\n",
        "        }\n",
        "        __syncthreads();\n",
        "    }\n",
        "    if (threadIdx.x == 0) {\n",
        "        atomicAdd(result, cache[0]);\n",
        "    }\n",
        "}"
      ]
    },
    {
      "cell_type": "markdown",
      "id": "ed6bd28d",
      "metadata": {
        "id": "ed6bd28d"
      },
      "source": [
        "Now that we've fused the mul and sum kernels into a dot kernel, let's have a look at the impact on the time and energy of our computations:"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "id": "8526fdbf",
      "metadata": {
        "id": "8526fdbf"
      },
      "outputs": [],
      "source": [
        "# TODO: create the argument list\n",
        "# hint: we want to compute the dot product of a and b and store the result in r\n",
        "# hint2: look at the CUDA kernel code above to see the kernel arguments\n",
        "dot_args = [ ]\n",
        "\n",
        "if not dot_args:\n",
        "    print(\"Error: First write the argument list dot_args!\")\n",
        "else:\n",
        "    # Let's call tune_kernel to start the tuning process.\n",
        "    results_dot, env = kt.tune_kernel(\"dot\", \"dot_kernel.cu\", size, dot_args, tune_params, lang=\"cupy\",\n",
        "                                      observers=[nvmlobserver], metrics=metrics)"
      ]
    },
    {
      "cell_type": "markdown",
      "id": "a6baa3ef",
      "metadata": {
        "id": "a6baa3ef"
      },
      "source": [
        "OK, now that we've collected our measurements. Let's take a look at the energy needed to complete our dot product computation, using two separate kernels, versus a single-kernel approach:"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "id": "9d616380",
      "metadata": {
        "id": "9d616380"
      },
      "outputs": [],
      "source": [
        "# The best kernel configurations in terms of energy for mul and sum were:\n",
        "mul_best = get_best_config(results_mul, \"energy\")\n",
        "print('Energy of the best mul kernel:', mul_best[\"energy\"], \"Joule\")\n",
        "sum_best = get_best_config(results_sum, \"energy\")\n",
        "print('Energy of the best sum kernel:', sum_best[\"energy\"], \"Joule\")\n",
        "\n",
        "# The total amount of energy for the two separate kernels:\n",
        "total = mul_best[\"energy\"] + sum_best[\"energy\"]\n",
        "print('Energy used in total by using separate kernels:', total, \"Joule\")\n",
        "\n",
        "# Energy used by the fused kernel:\n",
        "dot_best = get_best_config(results_dot, \"energy\")\n",
        "print('Energy of the best dot kernel:', dot_best[\"energy\"], \"Joule\")"
      ]
    },
    {
      "cell_type": "markdown",
      "id": "d8d296bf",
      "metadata": {
        "id": "d8d296bf"
      },
      "source": [
        "Can you think of why the approach with two separate kernels uses so much more energy than a single kernel?\n",
        "\n",
        "Hint: Think about the total amount of data that is loaded from and stored to global memory by the two separate kernels, and the same for the fused kernel.\n",
        "\n",
        "**That's it! You've successfully completed the second hands-on!**"
      ]
    }
  ],
  "metadata": {
    "kernelspec": {
      "display_name": "Python 3",
      "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.12"
    },
    "colab": {
      "provenance": [],
      "gpuType": "T4"
    },
    "accelerator": "GPU"
  },
  "nbformat": 4,
  "nbformat_minor": 5
}