hands-on/cuda/00_Kernel_Tuner_Introduction.ipynb

{
  "cells": [
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "AmH1n4971xm0"
      },
      "source": [
        "# Kernel Tuner Tutorial\n",
        "\n",
        "## Introduction Hands-on\n",
        "\n",
        "Welcome to the first hands-on of the Kernel Tuner tutorial. In this hands-on exercise we will learn how to **install** and access Kernel Tuner from Python, and **tune** our first CUDA kernel.\n",
        "\n",
        "To install the latest version of `kernel_tuner` we use `pip`, but before installing it we need to have a working CUDA installation available, and be sure to have both `numpy`  and `cupy` installed."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "7LBBP5n8EV-H"
      },
      "outputs": [],
      "source": [
        "%pip install kernel_tuner"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "cR2rHfIa4e7f"
      },
      "source": [
        "After installing all necessary packages, we can import `numpy` and `kernel_tuner`."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "As8FyYq3EtoK"
      },
      "outputs": [],
      "source": [
        "import numpy as np\n",
        "import kernel_tuner as kt"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "MLe_BXzR5PTk"
      },
      "source": [
        "Before using Kernel Tuner, we will create a text file containing the code of the CUDA kernel that we are going to use in this hands-on.\n",
        "\n",
        "This simple kernel is called `vector_add` and computes the elementwise sum of two vectors of size `n`."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "xgnQk_bmEx_F"
      },
      "outputs": [],
      "source": [
        "%%writefile vector_add_kernel.cu\n",
        "\n",
        "__global__ void vector_add(float * c, float * a, float * b, int n) {\n",
        "    int i = (blockIdx.x * blockDim.x) + threadIdx.x;\n",
        "    \n",
        "    if ( i < n ) {\n",
        "        c[i] = a[i] + b[i];\n",
        "    }\n",
        "}"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "cbkMDIgO7V9h"
      },
      "source": [
        "The execution of the cell above created the file `vector_add_kernel.cu` containing the source code of our kernel.\n",
        "\n",
        "We can now use Kernel Tuner to execute the code on the GPU; please read carefully both code and comments to become familiar with how Kernel Tuner works.\n",
        "\n",
        "For more details refer to the [API](https://KernelTuner.github.io/kernel_tuner/stable/user-api.html)."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "cZPTL1kjE-OS"
      },
      "outputs": [],
      "source": [
        "# the size of the vectors\n",
        "size = 1000000\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",
        "n = np.int32(size)\n",
        "\n",
        "# now we combine these variables in an argument list, which matches\n",
        "# the order and types of the function arguments of our CUDA kernel\n",
        "args = [c, a, b, n]\n",
        "\n",
        "# the next step is to create a dictionary to tell Kernel Tuner about\n",
        "# the tunable parameters in our code and what values these may take\n",
        "tune_params = dict()\n",
        "\n",
        "# finally, we call tune_kernel to start the tuning process. To do so,\n",
        "# we pass \n",
        "#    the name of the kernel we'd like to tune, in our case: \"vector_add\", \n",
        "#    the name of the file containing our source code,\n",
        "#    the problem_size that our kernel operates on\n",
        "#    the argument list to call our kernel function\n",
        "#    the dictionary with tunable parameters\n",
        "results, env = kt.tune_kernel(\"vector_add\", \"vector_add_kernel.cu\", size, args, tune_params, lang=\"cupy\")"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "IF1HCWx4Dzp8"
      },
      "source": [
        "\n",
        "\n",
        "The `tune_kernel` function returns two outputs that we saved as `results` and `env`:\n",
        "\n",
        "* `results` is a list of dictionaries, each containing detailed information about the configurations that have been benchmarked;\n",
        "* `env` is a dictionary that stores information about the hardware and software environment in which this experiment took place; it is recommended to store this information along with the benchmark results.\n",
        "\n",
        "We can also print the content of `results` to have a look at the output."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "LHlbScKUGir3"
      },
      "outputs": [],
      "source": [
        "print(f\"Number of configurations: {len(results)}\")\n",
        "print(results)"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "NbvRO_gzFOcU"
      },
      "source": [
        "If you paid attention to the output of Kernel Tuner, you may have noticed the following message:\n",
        "\n",
        "`UserWarning: None of the tunable parameters specify thread block dimensions!`\n",
        "\n",
        "The reason for this message is that we have not provided any tuning parameter, we have only executed the code.\n",
        "\n",
        "To tune `vector_add` we need to specify the values of a tunable parameter in `tune_params`."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "nqpzVDFfGt-9"
      },
      "outputs": [],
      "source": [
        "# to benchmark different configurations for the number of threads in a\n",
        "# thread block, we can use the special tuning parameter \"block_size_x\" and\n",
        "# add the values to test to it\n",
        "tune_params[\"block_size_x\"] = []\n",
        "\n",
        "# finally, we call tune_kernel to start the tuning process\n",
        "results, env = kt.tune_kernel(\"vector_add\", \"vector_add_kernel.cu\", size, args, tune_params, lang=\"cupy\")"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "ImKnDOweHCtL"
      },
      "source": [
        "Printing the content of `results` again we should be able to see the results of all benchmarked configurations."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "gYhIjq_pHMJN"
      },
      "outputs": [],
      "source": [
        "print(f\"Number of configurations: {len(results)}\")\n",
        "print(results)"
      ]
    },
    {
      "cell_type": "code",
      "source": [
        ""
      ],
      "metadata": {
        "id": "tz4WDvMBumPE"
      },
      "execution_count": null,
      "outputs": []
    }
  ],
  "metadata": {
    "accelerator": "GPU",
    "colab": {
      "collapsed_sections": [],
      "name": "00-Kernel_Tuner-Introduction.ipynb",
      "provenance": []
    },
    "kernelspec": {
      "display_name": "Python 3",
      "name": "python3"
    },
    "language_info": {
      "name": "python"
    }
  },
  "nbformat": 4,
  "nbformat_minor": 0
}