hands-on/cuda/02_Kernel_Tuner_Intermediate.ipynb

{
  "cells": [
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "-y3BG43Gkw61"
      },
      "source": [
        "# Kernel Tuner Tutorial\n",
        "\n",
        "## Intermediate Hands-on\n",
        "\n",
        "In this hands-on we will look at three features of Kernel Tuner that have been recently introduced to you: **search space restrictions**, **caching**, and **output verification**.\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": "TboeDlEAgqp0"
      },
      "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": "wdSglXdizKW5"
      },
      "source": [
        "To work with these features we will use a matrix multiplication kernel.\n",
        "\n",
        "Matrix multiplication is one of the most well-known and widely-used linear algebra operations, and is frequently used to demonstrate the high-performance computing capabilities of GPUs. As such, matrix multiplication presents a familiar starting point for many GPU programmers. More information about matrix multiplication can be found on [Wikipedia](https://en.wikipedia.org/wiki/Matrix_multiplication).\n",
        "\n",
        "The following cell contains the code of a matrix multiply kernel using shared memory. The content of the cell is written to the `matmul_shared.cu` file, and you only need to execute the cell once as this hands-on does not require to change the implementation of the kernel.\n",
        "\n",
        "This kernel assumes that the width and height of the matrices `A`, `B`, and `C` is equal to `WIDTH`, which is known at compile time. Of course, you'll want a more flexible solution in reality, but this is just an example kernel to demonstrate how to use Kernel Tuner."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "JJRLQecvoc1_"
      },
      "outputs": [],
      "source": [
        "%%writefile matmul_shared.cu\n",
        "\n",
        "#define WIDTH 512\n",
        "\n",
        "__global__ void matmul_kernel(float *C, float *A, float *B) {\n",
        "\n",
        "    __shared__ float sA[block_size_y][block_size_x];\n",
        "    __shared__ float sB[block_size_y][block_size_x];\n",
        "\n",
        "    int tx = threadIdx.x;\n",
        "    int ty = threadIdx.y;\n",
        "    int x = blockIdx.x * block_size_x + tx;\n",
        "    int y = blockIdx.y * block_size_y + ty;\n",
        "\n",
        "    float sum = 0.0;\n",
        "    int k,kb;\n",
        "\n",
        "    for (k=0; k<WIDTH; k+=block_size_x) {\n",
        "        __syncthreads();\n",
        "        sA[ty][tx] = A[y*WIDTH+k+tx];\n",
        "        sB[ty][tx] = B[(k+ty)*WIDTH+x];\n",
        "        __syncthreads();\n",
        "\n",
        "        for (kb=0; kb<block_size_x; kb++) {\n",
        "            sum += sA[ty][kb] * sB[kb][tx];\n",
        "        }\n",
        "\n",
        "    }\n",
        "\n",
        "    C[y*WIDTH+x] = sum;\n",
        "}"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "ftZM8ogLzt71"
      },
      "source": [
        "Before running the code we need to allocate input and output matrices, and add some tuning parameters."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "5K1BmJhdzuO2"
      },
      "outputs": [],
      "source": [
        "# matrix width needs to match the value in the kernel source\n",
        "problem_size = (512, 512)\n",
        "\n",
        "A = np.random.randn(*problem_size).astype(np.float32)\n",
        "B = np.random.randn(*problem_size).astype(np.float32)\n",
        "C = np.zeros_like(A)\n",
        "\n",
        "args = [C, A, B]\n",
        "\n",
        "tune_params = collections.OrderedDict()\n",
        "tune_params[\"block_size_x\"] = [2**i for i in range(0, 11)]\n",
        "tune_params[\"block_size_y\"] = [2**i for i in range(0, 11)]"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "5WuLDI7BlQ2c"
      },
      "source": [
        "It is now your turn to add some **search space restrictions**. You are free to add all the restrictions you want, but there is one in particular that is required for the kernel to produce correct results: the shape of the thread block needs to be **exactly** a square.\n",
        "\n",
        "Remember that restrictions are specified as either a Python list containing strings, each string being one restriction, or as a callable object that returns `True` if the configuration is valid and `False` otherwise."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "hTI0HEDkmHas"
      },
      "outputs": [],
      "source": [
        "# EXERCISE 1: Define the required search space restriction for the matrix multiplication kernel\n",
        "restrict = []"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "v9qQzbuBlWcf"
      },
      "source": [
        "To enable the **caching** of intermediate results during tuning, Kernel Tuner needs to know the name of the cache file. The name can be specified as a string, to which Kernel Tuner automatically adds the `.json` extension if not specified."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "_R9F5guoaMAd"
      },
      "outputs": [],
      "source": [
        "# define a string containing the cache file name\n",
        "cache_name = \"my_cache_file.json\""
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "hT9NLyIrZiWl"
      },
      "source": [
        "Do not forget to pass the restrictions to the `tune_kernel` function and enable caching 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": "LgrPl5Vc0cCe"
      },
      "outputs": [],
      "source": [
        "if not restrict:\n",
        "    print(\"Error: you must first define a search space restriction! (Exercise 1)\")\n",
        "\n",
        "# Call the tuner with the restricted search space\n",
        "else:\n",
        "    results, env = kt.tune_kernel(\"matmul_kernel\", \"matmul_shared.cu\",\n",
        "                                  problem_size, args, tune_params, restrictions=restrict,\n",
        "                                  cache=cache_name, verbose=True, lang=\"cupy\")\n",
        "\n",
        "    print(f\"Number of configurations: {len(results)}\")\n"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "7WKhURhGmOnS"
      },
      "source": [
        "### Output verification\n",
        "\n",
        "There are times, like with this matrix multiplication kernel, when some tuning configurations may produce wrong results.\n",
        "\n",
        "It is important to catch this as soon as possible, and Kernel Tuner allows to pass to the `tune_kernel` function a reference answer to which the results produced by all configuration are compared against.\n",
        "\n",
        "The reference answer is a Python list that matches in size and order the argument list provided to the kernel (`args` in our case), with `None` for all elements for which a comparison is not needed. In case of working with floating point values, Kernel Tuner allows also to specify a tolerance value.\n",
        "\n",
        "Again refer to the [API](https://KernelTuner.github.io/kernel_tuner/stable/user-api.html) for more information."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "Bqdn4OYdmRI4"
      },
      "outputs": [],
      "source": [
        "# compute the reference result, e.g. by using NumPy\n",
        "reference = A.dot(B)\n",
        "\n",
        "# EXERCISE 2: Correctly construct the answer list required by Kernel Tuner\n",
        "answer = []\n"
      ]
    },
    {
      "cell_type": "markdown",
      "source": [
        "Now, we are ready to call the tuner again with output verification enabled."
      ],
      "metadata": {
        "id": "qi-b2b2rGjvp"
      }
    },
    {
      "cell_type": "code",
      "source": [
        "if not answer:\n",
        "    print(\"Error: you must first setup the answer list correctly! (Exercise 2)\")\n",
        "\n",
        "# Call the tuner with output verification enabled\n",
        "else:\n",
        "    results, env = kt.tune_kernel(\"matmul_kernel\", \"matmul_shared.cu\",\n",
        "                             problem_size, args, tune_params, restrictions=restrict,\n",
        "                             answer=answer, lang=\"cupy\", atol=1e-4)\n",
        "\n",
        "    print(f\"Number of configurations: {len(results)}\")"
      ],
      "metadata": {
        "id": "8zAZb6ThGicZ"
      },
      "execution_count": null,
      "outputs": []
    }
  ],
  "metadata": {
    "accelerator": "GPU",
    "colab": {
      "name": "02-Kernel_Tuner-Intermediate.ipynb",
      "provenance": []
    },
    "kernelspec": {
      "display_name": "Python 3",
      "name": "python3"
    },
    "language_info": {
      "name": "python"
    }
  },
  "nbformat": 4,
  "nbformat_minor": 0
}