60_python_cuda_intro.ipynb

{
  "cells": [
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "view-in-github",
        "colab_type": "text"
      },
      "source": [
        "<a href=\"https://colab.research.google.com/github/ggruszczynski/gpu_colab/blob/main/60_python_cuda_intro.ipynb\" target=\"_parent\"><img src=\"https://colab.research.google.com/assets/colab-badge.svg\" alt=\"Open In Colab\"/></a>"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "9qF1uQPZBNA3"
      },
      "source": [
        "# Python + cuda\n",
        "\n",
        "Let us repeat the previous exercises in python."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 5,
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "vt3YAkBnBW-f",
        "outputId": "aaea4ba4-0e40-4a12-8407-97e5af629a9c"
      },
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Tue Apr 18 16:17:38 2023       \n",
            "+-----------------------------------------------------------------------------+\n",
            "| NVIDIA-SMI 525.85.12    Driver Version: 525.85.12    CUDA Version: 12.0     |\n",
            "|-------------------------------+----------------------+----------------------+\n",
            "| GPU  Name        Persistence-M| Bus-Id        Disp.A | Volatile Uncorr. ECC |\n",
            "| Fan  Temp  Perf  Pwr:Usage/Cap|         Memory-Usage | GPU-Util  Compute M. |\n",
            "|                               |                      |               MIG M. |\n",
            "|===============================+======================+======================|\n",
            "|   0  Tesla T4            Off  | 00000000:00:04.0 Off |                    0 |\n",
            "| N/A   67C    P0    29W /  70W |    619MiB / 15360MiB |      0%      Default |\n",
            "|                               |                      |                  N/A |\n",
            "+-------------------------------+----------------------+----------------------+\n",
            "                                                                               \n",
            "+-----------------------------------------------------------------------------+\n",
            "| Processes:                                                                  |\n",
            "|  GPU   GI   CI        PID   Type   Process name                  GPU Memory |\n",
            "|        ID   ID                                                   Usage      |\n",
            "|=============================================================================|\n",
            "|    0   N/A  N/A      1893      C   nvidia-smi                        616MiB |\n",
            "+-----------------------------------------------------------------------------+\n"
          ]
        }
      ],
      "source": [
        "!nvidia-smi"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 6,
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "nHQPUQyGfXlv",
        "outputId": "ae5aa145-543d-4160-86be-4443efbd6462"
      },
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Number of elements: 67108864 \n",
            "Memory size of array element in [MB]: 536.870912\n"
          ]
        }
      ],
      "source": [
        "from numba import cuda\n",
        "from numba import jit\n",
        "import numpy as np\n",
        "from numba import vectorize, int32, int64, float32, float64\n",
        "import matplotlib.pyplot as plt\n",
        "\n",
        "%matplotlib inline\n",
        "\n",
        "N = 2**26\n",
        "x = np.arange(N, dtype=np.float64) # [0...N] on the host\n",
        "\n",
        "print(f\"Number of elements: {N} \\nMemory size of array element in [MB]: {x.nbytes/1E6}\")\n"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "j55C2R4pkD5M"
      },
      "source": [
        "## Reduction"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 7,
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "H-inIbgFgpPj",
        "outputId": "dbc1d960-d181-4677-d6e3-c93c54332b10"
      },
      "outputs": [
        {
          "output_type": "stream",
          "name": "stderr",
          "text": [
            "/usr/local/lib/python3.9/dist-packages/numba/cuda/dispatcher.py:488: NumbaPerformanceWarning: Grid size 64 will likely result in GPU under-utilization due to low occupancy.\n",
            "  warn(NumbaPerformanceWarning(msg))\n",
            "/usr/local/lib/python3.9/dist-packages/numba/cuda/cudadrv/devicearray.py:885: NumbaPerformanceWarning: Host array used in CUDA kernel will incur copy overhead to/from device.\n",
            "  warn(NumbaPerformanceWarning(msg))\n",
            "/usr/local/lib/python3.9/dist-packages/numba/cuda/dispatcher.py:488: NumbaPerformanceWarning: Grid size 1 will likely result in GPU under-utilization due to low occupancy.\n",
            "  warn(NumbaPerformanceWarning(msg))\n"
          ]
        }
      ],
      "source": [
        "# reference: https://numba.pydata.org/numba-doc/dev/cuda/reduction.html\n",
        "\n",
        "@cuda.reduce\n",
        "def sum_reduce(a, b):\n",
        "    return a + b\n",
        "\n",
        "expect = x.sum()      # numpy sum reduction\n",
        "got = sum_reduce(x)   # cuda sum reduction\n",
        "assert expect == got\n"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 8,
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "aTVuw5m9gxfi",
        "outputId": "c438c258-deca-4d25-b12d-48dd69b4cd9a"
      },
      "outputs": [
        {
          "output_type": "stream",
          "name": "stderr",
          "text": [
            "/usr/local/lib/python3.9/dist-packages/numba/cuda/dispatcher.py:488: NumbaPerformanceWarning: Grid size 64 will likely result in GPU under-utilization due to low occupancy.\n",
            "  warn(NumbaPerformanceWarning(msg))\n",
            "/usr/local/lib/python3.9/dist-packages/numba/cuda/cudadrv/devicearray.py:885: NumbaPerformanceWarning: Host array used in CUDA kernel will incur copy overhead to/from device.\n",
            "  warn(NumbaPerformanceWarning(msg))\n",
            "/usr/local/lib/python3.9/dist-packages/numba/cuda/dispatcher.py:488: NumbaPerformanceWarning: Grid size 1 will likely result in GPU under-utilization due to low occupancy.\n",
            "  warn(NumbaPerformanceWarning(msg))\n"
          ]
        }
      ],
      "source": [
        "#Lambda functions can also be used here:\n",
        "sum_reduce_lam = cuda.reduce(lambda a, b: a + b)\n",
        "\n",
        "expect = x.sum()      # numpy sum reduction\n",
        "got = sum_reduce_lam(x)   # cuda sum reduction\n",
        "assert expect == got"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 9,
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "sRekYJDwg0_l",
        "outputId": "efdda7d0-6ad7-4e64-bd0c-d303f007f0b8"
      },
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "44.2 ms ± 246 µs per loop (mean ± std. dev. of 7 runs, 10 loops each)\n"
          ]
        }
      ],
      "source": [
        "%timeit x.sum()    # NumPy on CPU"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 10,
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "xM-z1feZg4g1",
        "outputId": "623fc5fa-9434-47a2-f6ed-2c0beb080ece"
      },
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "258 ms ± 30.6 ms per loop (mean ± std. dev. of 7 runs, 1 loop each)\n"
          ]
        }
      ],
      "source": [
        "%timeit sum_reduce_lam(x) # Numba on GPU - data from host"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 11,
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "ZkL2PuC7kYFp",
        "outputId": "4b3f12df-909f-45cd-85bd-07de25f94bf6"
      },
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "2.49 ms ± 17.5 µs per loop (mean ± std. dev. of 7 runs, 100 loops each)\n"
          ]
        }
      ],
      "source": [
        "d_x = cuda.to_device(x)\n",
        "%timeit sum_reduce_lam(d_x) # Numba on GPU - prefetched data"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "BXs5A0DQESu-"
      },
      "source": []
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "u1LqTI_TEawo"
      },
      "source": [
        "## SAXPY\n",
        "\n",
        "**SAXPY** stands for “Single-Precision A·X Plus Y”. It is a function in the standard Basic Linear Algebra Subroutines (BLAS) library."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 12,
      "metadata": {
        "id": "AASMcWZbsQMv"
      },
      "outputs": [],
      "source": [
        "a = 10\n",
        "y = np.copy(x)\n",
        "\n",
        "d_a = cuda.to_device(a) # Copy of a on the device\n",
        "d_x = cuda.to_device(x) # Copy of x on the device\n",
        "d_y = cuda.to_device(y) # Copy of y on the device\n",
        "d_out = cuda.device_array_like(d_x) # Like np.array_like, but for device arrays\n"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 13,
      "metadata": {
        "id": "DR4IAy8Df_Sq"
      },
      "outputs": [],
      "source": [
        "\n",
        "@vectorize(['float64(int64, float64, float64)'], target='cuda') # Type signature and target are required for the GPU\n",
        "def add_ufunc(a, x, y):\n",
        "    return a*x + y"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 14,
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "HiYG5WgHnMdU",
        "outputId": "5b3d7b7a-0760-4825-e50a-4a2a1155523c"
      },
      "outputs": [
        {
          "output_type": "execute_result",
          "data": {
            "text/plain": [
              "True"
            ]
          },
          "metadata": {},
          "execution_count": 14
        }
      ],
      "source": [
        "expect = a*x + y      # numpy sum reduction\n",
        "got = add_ufunc(d_a, d_x, d_y)   # cuda sum reduction\n",
        "# assert expect == got\n",
        "\n",
        "np.allclose(expect, got)"
      ]
    }
  ],
  "metadata": {
    "accelerator": "GPU",
    "colab": {
      "name": "example50_thrust.ipynb",
      "provenance": [],
      "include_colab_link": true
    },
    "gpuClass": "standard",
    "kernelspec": {
      "display_name": "Python 3 (ipykernel)",
      "language": "python",
      "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.8.10"
    }
  },
  "nbformat": 4,
  "nbformat_minor": 0
}