colab-notebook.ipynb

{
  "cells": [
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "LcZ68UjXQ6om"
      },
      "source": [
        "## Exploring CUDA in Colab\n",
        "\n",
        "In this Colab notebook, we will explore CUDA examples.\n",
        "\n",
        "Google provides free access to an Nvidia Tesla T4 GPU on Colab that can be chosen by navigating to Runtime->Change runtime type."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 1,
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "h0IwwmjUqLa5",
        "outputId": "0d15380d-b747-4ce7-fdc4-44cd3c88c356"
      },
      "outputs": [
        {
          "name": "stdout",
          "output_type": "stream",
          "text": [
            "Mounted at /content/drive\n"
          ]
        }
      ],
      "source": [
        "from google.colab import drive\n",
        "drive.mount('/content/drive')\n",
        "# from google.colab import drive\n",
        "# drive.flush_and_unmount()"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "h-t1EiYWwEDM"
      },
      "source": [
        "Nvidia's cuda compiler driver `nvcc`  comes pre-installed in Colab's file system, as can be checked by typing `!nvcc --version`"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 2,
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "Nd3OPf-JTfP2",
        "outputId": "7d633c98-62c2-480e-9c7b-822418c1e4eb"
      },
      "outputs": [
        {
          "name": "stdout",
          "output_type": "stream",
          "text": [
            "nvcc: NVIDIA (R) Cuda compiler driver\n",
            "Copyright (c) 2005-2023 NVIDIA Corporation\n",
            "Built on Tue_Aug_15_22:02:13_PDT_2023\n",
            "Cuda compilation tools, release 12.2, V12.2.140\n",
            "Build cuda_12.2.r12.2/compiler.33191640_0\n"
          ]
        }
      ],
      "source": [
        "!nvcc --version"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "1cTEZjH1wJR0"
      },
      "source": [
        "Subsequently, we install and load the `nvcc4jupyter` extension as follows:"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 5,
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "4ibX5wv1TfMP",
        "outputId": "fba7dbf4-6fc0-4977-b3b7-b40480349e36"
      },
      "outputs": [
        {
          "name": "stdout",
          "output_type": "stream",
          "text": [
            "Collecting git+https://github.com/andreinechaev/nvcc4jupyter.git\n",
            "  Cloning https://github.com/andreinechaev/nvcc4jupyter.git to /tmp/pip-req-build-689dld_q\n",
            "  Running command git clone --filter=blob:none --quiet https://github.com/andreinechaev/nvcc4jupyter.git /tmp/pip-req-build-689dld_q\n",
            "  Resolved https://github.com/andreinechaev/nvcc4jupyter.git to commit 801584cceb559adc54e828ebe9b385c5f53fe70f\n",
            "  Installing build dependencies ... \u001b[?25l\u001b[?25hdone\n",
            "  Getting requirements to build wheel ... \u001b[?25l\u001b[?25hdone\n",
            "  Preparing metadata (pyproject.toml) ... \u001b[?25l\u001b[?25hdone\n",
            "Building wheels for collected packages: nvcc4jupyter\n",
            "  Building wheel for nvcc4jupyter (pyproject.toml) ... \u001b[?25l\u001b[?25hdone\n",
            "  Created wheel for nvcc4jupyter: filename=nvcc4jupyter-1.2.1-py3-none-any.whl size=10743 sha256=3a570d3ab74e9051a615aa78d54722c78a69a0c83a1d407fccde14c37a30e96c\n",
            "  Stored in directory: /tmp/pip-ephem-wheel-cache-fm3_pvlk/wheels/a8/b9/18/23f8ef71ceb0f63297dd1903aedd067e6243a68ea756d6feea\n",
            "Successfully built nvcc4jupyter\n",
            "Installing collected packages: nvcc4jupyter\n",
            "Successfully installed nvcc4jupyter-1.2.1\n",
            "Detected platform \"Colab\". Running its setup...\n",
            "Source files will be saved in \"/tmp/tmpzsp8jg5n\".\n"
          ]
        }
      ],
      "source": [
        "!pip install git+https://github.com/andreinechaev/nvcc4jupyter.git\n",
        "%load_ext nvcc4jupyter"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "-eWmlLmywwBx"
      },
      "source": [
        "We are now ready to test CUDA using the `%%cuda` cell magic"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 6,
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "KinHeGyQTe-O",
        "outputId": "9417308d-e336-406d-d525-40023c353176"
      },
      "outputs": [
        {
          "name": "stdout",
          "output_type": "stream",
          "text": [
            "Congratulations, CUDA is up and running!\n",
            "\n"
          ]
        }
      ],
      "source": [
        "%%cuda\n",
        "#include <iostream>\n",
        "int main()\n",
        "{\n",
        "\tstd::cout << \"Congratulations, CUDA is up and running!\\n\";\n",
        "\treturn 0;\n",
        "}\n"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "D6BPQQRtzJQl"
      },
      "source": [
        "For context, a single Tesla T4 GPU has 40 SMs with 64 cuda cores each (hence 2560 cuda cores in total). With warp-scheduling and latency hiding, a single SM can support a maximum of 1024 concurrent threads (32 warps each with 32 threads). Hopefully we can test the full power of T4 by the end of this notebook!"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "rDQW_LzL3Vt2"
      },
      "source": [
        "Let us test a CUDA program that squares every number in the array [0,1, ... 64]"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 7,
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "collapsed": true,
        "id": "xJKUZLC7Te70",
        "outputId": "6d86e531-59af-495f-86f4-fa19fb5bed5d"
      },
      "outputs": [
        {
          "name": "stdout",
          "output_type": "stream",
          "text": [
            "0.000000\t1.000000\t4.000000\t9.000000\n",
            "16.000000\t25.000000\t36.000000\t49.000000\n",
            "64.000000\t81.000000\t100.000000\t121.000000\n",
            "144.000000\t169.000000\t196.000000\t225.000000\n",
            "256.000000\t289.000000\t324.000000\t361.000000\n",
            "400.000000\t441.000000\t484.000000\t529.000000\n",
            "576.000000\t625.000000\t676.000000\t729.000000\n",
            "784.000000\t841.000000\t900.000000\t961.000000\n",
            "1024.000000\t1089.000000\t1156.000000\t1225.000000\n",
            "1296.000000\t1369.000000\t1444.000000\t1521.000000\n",
            "1600.000000\t1681.000000\t1764.000000\t1849.000000\n",
            "1936.000000\t2025.000000\t2116.000000\t2209.000000\n",
            "2304.000000\t2401.000000\t2500.000000\t2601.000000\n",
            "2704.000000\t2809.000000\t2916.000000\t3025.000000\n",
            "3136.000000\t3249.000000\t3364.000000\t3481.000000\n",
            "3600.000000\t3721.000000\t3844.000000\t3969.000000\n",
            "\n"
          ]
        }
      ],
      "source": [
        "%%cuda\n",
        "#include <stdio.h>\n",
        "\n",
        "//a simple kernel that squares every element in an array\n",
        "__global__ void square(float * d_in, float * d_out) {\n",
        "    int idx = threadIdx.x;\n",
        "    float f = d_in[idx];\n",
        "    d_out[idx] = f * f;\n",
        "}\n",
        "\n",
        "//cpu program flow\n",
        "int main(int argc, char ** argv) {\n",
        "    const int ARRAY_SIZE = 64;\n",
        "    const int ARRAY_BYTES = ARRAY_SIZE * sizeof(float);\n",
        "\n",
        "    //generate input array, declare output array on cpu (host)\n",
        "    float h_in[ARRAY_SIZE];\n",
        "    for (int i = 0; i < ARRAY_SIZE; i++) {\n",
        "        h_in[i] = float(i);\n",
        "    }\n",
        "    float h_out[ARRAY_SIZE];\n",
        "\n",
        "    //declare gpu memory pointers, allocate gpu memory (device)\n",
        "    float * d_in;\n",
        "    float * d_out;\n",
        "    cudaMalloc((void **) &d_in, ARRAY_BYTES);\n",
        "    cudaMalloc((void **) &d_out, ARRAY_BYTES);\n",
        "\n",
        "    //transfer input from host to device\n",
        "    cudaMemcpy(d_in, h_in, ARRAY_BYTES, cudaMemcpyHostToDevice);\n",
        "    //launch kernel on device\n",
        "    square<<<1, ARRAY_SIZE>>>(d_in, d_out);\n",
        "    //transfer output from device to host\n",
        "    cudaMemcpy(h_out, d_out, ARRAY_BYTES, cudaMemcpyDeviceToHost);\n",
        "\n",
        "    //print output\n",
        "    for (int i = 0; i < ARRAY_SIZE; i++) {\n",
        "        printf(\"%f\", h_out[i]);\n",
        "        printf(((i%4) != 3) ? \"\\t\" : \"\\n\");\n",
        "    }\n",
        "\n",
        "    //free gpu memory\n",
        "    cudaFree(d_in);\n",
        "    cudaFree(d_out);\n",
        "\n",
        "    return 0;\n",
        "}"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "jVUWGoAU3juO"
      },
      "source": [
        "Okay, so CUDA indeed works as we expect it to with respect to the standard program flow! (in sequence: declare CPU i/o memory, initialize inputs in CPU, allocate GPU i/o memory, copy inputs to GPU, compute in GPU, copy outputs back to CPU). We are now prepared to take on more complex tasks!"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 3,
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "6fYMG4EiTe5t",
        "outputId": "aa7d4836-6845-4637-9093-0f33ddd81b5c"
      },
      "outputs": [
        {
          "name": "stdout",
          "output_type": "stream",
          "text": [
            "Collecting colab-xterm\n",
            "  Downloading colab_xterm-0.2.0-py3-none-any.whl (115 kB)\n",
            "\u001b[?25l     \u001b[90m━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━\u001b[0m \u001b[32m0.0/115.6 kB\u001b[0m \u001b[31m?\u001b[0m eta \u001b[36m-:--:--\u001b[0m\r\u001b[2K     \u001b[91m━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━\u001b[0m\u001b[91m╸\u001b[0m\u001b[90m━\u001b[0m \u001b[32m112.6/115.6 kB\u001b[0m \u001b[31m3.6 MB/s\u001b[0m eta \u001b[36m0:00:01\u001b[0m\r\u001b[2K     \u001b[90m━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━\u001b[0m \u001b[32m115.6/115.6 kB\u001b[0m \u001b[31m1.9 MB/s\u001b[0m eta \u001b[36m0:00:00\u001b[0m\n",
            "\u001b[?25hRequirement already satisfied: ptyprocess~=0.7.0 in /usr/local/lib/python3.10/dist-packages (from colab-xterm) (0.7.0)\n",
            "Requirement already satisfied: tornado>5.1 in /usr/local/lib/python3.10/dist-packages (from colab-xterm) (6.3.3)\n",
            "Installing collected packages: colab-xterm\n",
            "Successfully installed colab-xterm-0.2.0\n",
            "env: TERM=xterm\n"
          ]
        }
      ],
      "source": [
        "!pip install colab-xterm\n",
        "%load_ext colabxterm\n",
        "%env TERM=xterm"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "collapsed": true,
        "id": "os8xU7p9XK1I"
      },
      "outputs": [],
      "source": [
        "%xterm"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "xjmtukvWQvlh"
      },
      "source": [
        " ## CS344 Intro to Parallel Programming (Problem Sets)\n",
        "\n",
        "The problem sets of this course are image processing assignments using CUDA and OpenCV. Running these assignments in colab needs using `colab-xterm` whose installation is described in the previous snippets. Once installed, we navigate into a copy of [this fork of CS344](https://github.com/Adeemj/cs344) and follow instructions in the README to install and build OpenCV in the repo folder.\n",
        "\n",
        "After the above is done, for each problem we solve (by writing a kernel in `student_func.cu`), we run `make` and `./HWk input.jpg` where \"k\" indexes Problem Set k whose folder we are in). Here we present a summary of our results.\n",
        "\n",
        "We successfully solved and ran Problem 1. Problem 1 asks us to write a simple kernel that converts color images to greyscale images by weighted averaging the RGB channels in each pixel as output = 0.299f * R + 0.587f * G + 0.114f * B. We launch the kernel below as a parallel computation of `numRows` blocks each with `numCols` threads in one dimension.\n",
        "\n",
        "```\n",
        "__global__\n",
        "void rgba_to_greyscale(const uchar4* const rgbaImage, unsigned char* const greyImage, int numRows, int numCols)\n",
        "{\n",
        "  int idx = threadIdx.x + (blockIdx.x * numCols);\n",
        "  float val = (.299f * rgbaImage[idx].x) + (.587f * rgbaImage[idx].y) + (.114f * rgbaImage[idx].z);\n",
        "  greyImage[idx] = val;\n",
        "}\n",
        "\n",
        "void your_rgba_to_greyscale(const uchar4 * const h_rgbaImage, uchar4 * const d_rgbaImage, unsigned char* const d_greyImage, size_t numRows, size_t numCols)\n",
        "{\n",
        "  const dim3 blockSize(numCols, 1, 1);  \n",
        "  const dim3 gridSize(numRows, 1, 1);  \n",
        "  rgba_to_greyscale<<<gridSize, blockSize>>>(d_rgbaImage, d_greyImage, numRows, numCols);\n",
        "  cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError());\n",
        "}\n",
        "```"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "Bp8wilBZY6ss"
      },
      "source": [
        "Input:\n",
        "\n",
        "<img src=\"https://drive.google.com/uc?export=view&id=1lLxx9JxLuMZqrKufokpC17HTWRhJFH5I\" width=\"500px\">\n",
        "\n",
        "Output:\n",
        "\n",
        "<img src=\"https://drive.google.com/uc?export=view&id=1X_VABMZxBYLyrOmhPGOIuRVFMOZ9xifx\" width=\"500px\">\n",
        "\n",
        "Reference:\n",
        "\n",
        "<img src=\"https://drive.google.com/uc?export=view&id=1fiFTfDfZwF11ctiQdEtSWYHWfZ9BUXPy\" width=\"500px\">\n",
        "\n",
        "Difference (between reference and output):\n",
        "\n",
        "<img src=\"https://drive.google.com/uc?export=view&id=1ZdOVvEiBo79Yh9KTtnxGCnnGp_CxybhZ\" width=\"500px\">"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "fA4CU1jNhP4r"
      },
      "source": [
        "\n",
        "We notice a small number of isolated white pixels in the difference (like stars in the night sky!). It is unclear what the basis of this miniscule error is. Perhaps they were introduced intentionally, but not sure!\n"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "6rJwVfEetzAM"
      },
      "source": [
        "In Problem 2, we are asked to implement [Gaussian Blurring](https://en.wikipedia.org/wiki/Gaussian_blur) of the same image using a 9x9 filter. This is a slightly more involved task. We omit the full code here for brevity, and focus on the results."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "9474NTOPvk0M"
      },
      "source": [
        "xterm window:\n",
        "\n",
        "<img src=\"https://drive.google.com/uc?export=view&id=1NVQErePZuTnsZtBvU_GIucgIxt7b4IY0\" width=\"500px\">\n",
        "\n",
        "Input:\n",
        "\n",
        "<img src=\"https://drive.google.com/uc?export=view&id=1lLxx9JxLuMZqrKufokpC17HTWRhJFH5I\" width=\"500px\">\n",
        "\n",
        "Output:\n",
        "\n",
        "<img src=\"https://drive.google.com/uc?export=view&id=13Vy9seUNQU5Gzv1YDWIYm_s1a6_1bBop\" width=\"500px\">\n",
        "\n",
        "Reference:\n",
        "\n",
        "<img src=\"https://drive.google.com/uc?export=view&id=13SSI5x_-CnaVXyF988lskSjDRhgkEKIM\" width=\"500px\">\n",
        "\n",
        "Difference (between reference and output):\n",
        "\n",
        "<img src=\"https://drive.google.com/uc?export=view&id=13OlDGjoC3ECy9T61kmcN_-XULU9A5if1\" width=\"500px\">\n"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "mXAOaAlUx8OC"
      },
      "source": [
        "We got it perfectly correct this time. Basically a 9x9 (9 is the default filter width they choose) matrix of weights is centered at each pixel and averaged over its neighborhood. The main `gaussian_blur` kernel was written as follows:\n",
        "```\n",
        "__global__\n",
        "void gaussian_blur(const unsigned char* const inputChannel,\n",
        "                   unsigned char* const outputChannel,\n",
        "                   int numRows, int numCols,\n",
        "                   const float* const filter, const int filterWidth)\n",
        "{\n",
        "  // Calculate thread position\n",
        "  int col = blockIdx.x * blockDim.x + threadIdx.x;\n",
        "  int row = blockIdx.y * blockDim.y + threadIdx.y;\n",
        "  \n",
        "  if (col < numCols && row < numRows)\n",
        "  {\n",
        "    float result = 0.0f;\n",
        "    // For each value in the filter\n",
        "    for (int i = 0; i < filterWidth; i++)\n",
        "    {\n",
        "      for (int j = 0; j < filterWidth; j++)\n",
        "      {\n",
        "        // Find the global image position for this filter position\n",
        "        // Clamp to boundary of the image\n",
        "        int iRow = min(max(row + i - filterWidth/2, 0), numRows-1);\n",
        "        int iCol = min(max(col + j - filterWidth/2, 0), numCols-1);\n",
        "        \n",
        "        result += inputChannel[iRow*numCols + iCol] * filter[i*filterWidth + j];\n",
        "      }\n",
        "    }\n",
        "    // Write the new value to the output\n",
        "    outputChannel[row*numCols + col] = result;\n",
        "  }\n",
        "  ```"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "TULMR03mRFk4"
      },
      "source": [
        "---\n",
        "### References  \n",
        "[1] <a id=\"ref-1\"></a> [github.com/Adeemj/cs344](https://github.com/Adeemj/cs344)  \n"
      ]
    }
  ],
  "metadata": {
    "accelerator": "GPU",
    "colab": {
      "gpuType": "T4",
      "provenance": []
    },
    "kernelspec": {
      "display_name": "Python 3",
      "name": "python3"
    },
    "language_info": {
      "name": "python"
    }
  },
  "nbformat": 4,
  "nbformat_minor": 0
}