hands-on/cuda/01_Kernel_Tuner_Getting_Started.ipynb

{
  "cells": [
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "O0oxTslc7KWc"
      },
      "source": [
        "# Kernel Tuner Tutorial\n",
        "\n",
        "## Getting Started Hands-on\n",
        "\n",
        "In this hands-on we will look at two features of Kernel Tuner that have been recently introduced to you: **tunable grid dimensions** and **user defined metrics**.\n",
        "\n",
        "But first, if you have not done it already, it is time to install and import kernel_tuner and its dependencies."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "Ieqcoy4_3BDX"
      },
      "outputs": [],
      "source": [
        "%pip install kernel_tuner\n",
        "\n",
        "import numpy as np\n",
        "import kernel_tuner as kt\n",
        "import collections"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "kd9z8p1VM4yG"
      },
      "source": [
        "To introduce these concepts we will use a modified vector add kernel.\n",
        "\n",
        "This kernel computes the same result as the kernel in the previous hands-on, i.e. the elementwise sum of two vectors of size `n`, but each thread can compute more than one element.\n",
        "\n",
        "The content of the cell is written to the `vector_add_tiled.cu` file, and you only need to execute this cell once as this hands-on does not require to change the implementation of the kernel."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "SJa6qNHOTnZK"
      },
      "outputs": [],
      "source": [
        "%%writefile vector_add_tiled.cu\n",
        "\n",
        "__global__ void vector_add(float * c, float * a, float * b, int n) {\n",
        "    int i = (blockIdx.x * blockDim.x * tiling_factor) + (threadIdx.x * tiling_factor);\n",
        "\n",
        "    if ( (i + tiling_factor) <= n ) {\n",
        "        #pragma unroll\n",
        "        for ( int item = 0; item < tiling_factor; item++ ) {\n",
        "            c[i + item] = a[i + item] + b[i + item];\n",
        "        }\n",
        "    }\n",
        "}"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "ANUVMtorYIad"
      },
      "source": [
        "Before running the code we need to allocate memory and add some tuning parameters."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "NFB85iP0aUBx"
      },
      "outputs": [],
      "source": [
        "size = 1_000_000\n",
        "\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",
        "n = np.int32(size)\n",
        "\n",
        "args = [c, a, b, n]\n",
        "\n",
        "tune_params = collections.OrderedDict()\n",
        "tune_params[\"block_size_x\"] = [2**i for i in range(0, 11)]\n",
        "tune_params[\"tiling_factor\"] = [i for i in range(1, 11)]"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "X9kh32ixcssa"
      },
      "source": [
        "Normally, Kernel Tuner computes the grid size of our CUDA kernel automatically, based on problem size and number of threads per block (``block_size_x``). However, this is not possible for cases where other tunable parameters (i.e. `tiling_factor`) also affect the grid size.\n",
        "\n",
        "It is your responsibility to tell Kernel Tuner to work with **tunable grid dimensions**. To do so, you can define a Python list containing the names of the tunable parameters that should be used to compute the grid dimensions from the problem size."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "1aQ0AqT7eEqI"
      },
      "outputs": [],
      "source": [
        "# EXERCISE 1: Provide a list of tunable parameter names that divide the grid dimenions\n",
        "grid_div_x = []"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "SziiWji4fJ5o"
      },
      "source": [
        "Execution time is important, but not always the most relevant metric for many users. Because of this, Kernel Tuner allows to work with **user defined metrics** that are computed within and then returned by `tune_kernel`.\n",
        "\n",
        "Metrics are passed to Kernel Tuner as `lambda` functions contained in an ordered dictionary, with the key of the entry being the name of the metric itself. The order is important because it is allowed for metrics to build on earlier defined metrics.\n",
        "\n",
        "It is your responsibility to define one or more metrics and then tune the provided kernel. Possible user defined metrics in this case are the number of operations per second, or memory bandwidth.\n"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "NbJwYvWsffiI"
      },
      "outputs": [],
      "source": [
        "# First we create an OrderedDict. Actually, in newer Python versions all dictionaries are ordered.\n",
        "metrics = collections.OrderedDict()\n",
        "\n",
        "# Now we define our first metric. In this case, we want the performance of our kernel to\n",
        "# be computed in billions of floating-point operations per second.\n",
        "metrics[\"Performance (GFLOP/s)\"] = lambda p: (n / 1e9) / (p[\"time\"] / 1e3)\n",
        "# Let's unpack what the above line means:. We've created a lambda function that\n",
        "# takes an argument 'p' that contains the results collected by Kernel Tuner.\n",
        "# Our function should return the performance in GFLOP/s of this specific code\n",
        "# variant of our kernel.\n",
        "# Because 'n' is the size of our array, and equal to the number of floating-point additions\n",
        "# our kernel performs, we start with dividing n by one billion (1e9).\n",
        "# Kernel Tuner measures the execution time of our kernel in miliseconds. So, to arrive\n",
        "# at the execution time in seconds, we divide the execution time by a thousand.\n",
        "\n",
        "# EXERCISE 2: Define a user-defined metric for the achieved memory bandwith (throughput)\n",
        "# of our vector_add kernel, use \"Throughput (GB/s)\" as the key.\n",
        "# Because the vector_add kernel reads twice as much data as it writes, it is OK\n",
        "# to only consider the bandwidth required for the input data.\n",
        "# Think of how to express the througput in gigabytes per second of our kernel."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "Q2JQBidxfftX"
      },
      "source": [
        "Now we are ready to pass these additional arguments to the `tune_kernel` function as documented in Kernel Tuner's [API](https://KernelTuner.github.io/kernel_tuner/stable/user-api.html)."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "I-bDAqozfscV"
      },
      "outputs": [],
      "source": [
        "if not grid_div_x:\n",
        "    print(\"Error: first setup grid_div_x (Exercise 1)\")\n",
        "elif \"Throughput (GB/s)\" not in metrics:\n",
        "    print(\"Error: first define a metric for the throughput (Exercise 2)\")\n",
        "\n",
        "# Call the tuner\n",
        "# Mostly the same as before, but now we also pass:\n",
        "#    grid_div_x, to tell Kernel Tuner how to compute the grid dimensions\n",
        "#    metrics, a dictionary with user-defined metrics\n",
        "else:\n",
        "    results, env = kt.tune_kernel(\"vector_add\", \"vector_add_tiled.cu\", size, args, tune_params,\n",
        "                                  grid_div_x=grid_div_x,\n",
        "                                  metrics=metrics,\n",
        "                                  lang=\"cupy\")\n",
        "if results:\n",
        "    print(f\"Number of configurations: {len(results)}\")"
      ]
    },
    {
      "cell_type": "code",
      "source": [],
      "metadata": {
        "id": "icW_d44lweGO"
      },
      "execution_count": null,
      "outputs": []
    }
  ],
  "metadata": {
    "colab": {
      "name": "01-Kernel_Tuner-Getting_Started.ipynb",
      "provenance": []
    },
    "kernelspec": {
      "display_name": "Python 3",
      "name": "python3"
    },
    "language_info": {
      "name": "python"
    },
    "accelerator": "GPU"
  },
  "nbformat": 4,
  "nbformat_minor": 0
}