CudaColab.ipynb

{
  "cells": [
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "Ktm3JkI7vrnR"
      },
      "source": [
        "# CUDA on Colab\n",
        "\n",
        "This notebook, based on an example from Nvidia, shows how to check the GPU status of your Colab notebook, check out a github repository containing your c++ code, and compile it using either g++ for CPU or nvcc for GPU. and run it.\n",
        "\n",
        "Not yet covered, profiling.\n",
        "\n",
        "Author: Evelyn Mitchell\n",
        "Source Repository: https://github.com/evelynmitchell/cuda-on-colab\n",
        "Date: 2023-12-04"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "This will time the execution of each cell, and print the time at the end of the cell."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {},
      "outputs": [],
      "source": [
        "!pip install ipython-autotime\n",
        "%load-ext autotime"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "6Fsj6QjYwFtx"
      },
      "source": [
        "The nvidia-smi cli tells you about your GPU. The sample outputs for different types of GPUs or TPUs follow."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "MJ-Pv2U5IKTJ"
      },
      "outputs": [],
      "source": [
        "!nvidia-smi"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "00ulMfjWJMUH"
      },
      "source": [
        "A100 GPU\n",
        "```\n",
        "\n",
        "```"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "IpxfIJiUI_c9"
      },
      "source": [
        "V100 GPU\n",
        "```\n",
        "Mon Dec  4 18:42:36 2023       \n",
        "+-----------------------------------------------------------------------------+\n",
        "| NVIDIA-SMI 525.105.17   Driver Version: 525.105.17   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 V100-SXM2...  Off  | 00000000:00:04.0 Off |                    0 |\n",
        "| N/A   31C    P0    23W / 300W |      0MiB / 16384MiB |      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",
        "|  No running processes found                                                 |\n",
        "+-----------------------------------------------------------------------------+\n",
        "```"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "2tTpRQF-IikI"
      },
      "source": [
        "T4 TPU\n",
        "```\n",
        "Mon Dec  4 18:40:38 2023       \n",
        "+-----------------------------------------------------------------------------+\n",
        "| NVIDIA-SMI 525.105.17   Driver Version: 525.105.17   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   48C    P8    11W /  70W |      0MiB / 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",
        "|  No running processes found                                                 |\n",
        "+-----------------------------------------------------------------------------+\n",
        "```"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "DlfDZjClfzwb"
      },
      "source": [
        "# libcuda Driver\n",
        "\n",
        "If you install some libraries from source, such as Triton, you may lose the libcuda driver which is already installed in colab, when you uninstall triton, to install from source. Following the diagnosis and fix in [6] we will find out if the library is installed, then make sure it is in our execution path.\n",
        "\n",
        "This will show up as:\n",
        "```\n",
        "libcuda.so cannot found\n",
        "```\n",
        "\n",
        "To check if the cuda library is availanble run:\n",
        "```\n",
        "!ldconfig -p |grep libcuda\n",
        "```\n",
        "Which should show a result like\n",
        "```\n",
        "libcudart.so.11.0 (libc6,x86-64) => /usr/local/cuda/targets/x86_64-linux/lib/libcudart.so.11.0\n",
        "libcudart.so (libc6,x86-64) => /usr/local/cuda/targets/x86_64-linux/lib/libcudart.so\n",
        "```\n",
        "Note that ```libcuda.so``` is not listed.\n",
        "\n",
        "[6] (https://github.com/pytorch/pytorch/issues/107960#issuecomment-1709589190)"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "koCseLsGfxr7"
      },
      "outputs": [],
      "source": [
        "!ldconfig -p | grep libcuda"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "9Xo8tBZ9hf01"
      },
      "source": [
        "To find the path to ```libcuda.so``` run\n",
        "```\n",
        "find /usr -name 'libcuda.so'\n",
        "```\n",
        "Which should output something similar to:\n",
        "```\n",
        "/usr/local/cuda-11.8/compat/libcuda.so\n",
        "/usr/local/cuda-11.8/targets/x86_64-linux/lib/stubs/libcuda.so\n",
        "/usr/lib64-nvidia/libcuda.so\n",
        "```\n",
        "The version numbers may be different."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "AHTH84sVhKhL"
      },
      "outputs": [],
      "source": [
        "!find /usr -name 'libcuda.so'"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "5XggaJnTi2B2"
      },
      "source": [
        "We have the same issue as in [6], in that the ```stubs``` path is incorrect, so we will apply the fix, which is to add ```/usr/lib64-nvidia/libcuda.so``` to our shared libraries with:\n",
        "```\n",
        "ldconfig /usr/lib64-nvidia\n",
        "```"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "IrQFIwZfjShm"
      },
      "outputs": [],
      "source": [
        "!ldconfig /usr/lib64-nvidia"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "DxWt0YnijVYT"
      },
      "source": [
        "And  then verify"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "RedH0NsWja0o"
      },
      "outputs": [],
      "source": [
        "!ldconfig -p | grep libcuda"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "-Y4n7t8Fjiwl"
      },
      "source": [
        "```\n",
        "\tlibcudart.so.11.0 (libc6,x86-64) => /usr/local/cuda/targets/x86_64-linux/lib/libcudart.so.11.0\n",
        "\tlibcudart.so (libc6,x86-64) => /usr/local/cuda/targets/x86_64-linux/lib/libcudart.so\n",
        "\tlibcudadebugger.so.1 (libc6,x86-64) => /usr/lib64-nvidia/libcudadebugger.so.1\n",
        "\tlibcuda.so.1 (libc6,x86-64) => /usr/lib64-nvidia/libcuda.so.1\n",
        "\tlibcuda.so (libc6,x86-64) => /usr/lib64-nvidia/libcuda.so\n",
        "  ```\n",
        "  that libcuda.so shows up in the list of shared libraries."
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "XN1nWSaHKUVf"
      },
      "source": [
        "# C++ for CUDA\n",
        "Install the c++ build chain, which should be already available on colab."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "DKBGPAMsJ4Ji"
      },
      "outputs": [],
      "source": [
        "!apt install build-essential"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "3-VVuTSdwVRJ"
      },
      "source": [
        "The GPU compiler for c++ from Nvidia is called nvcc, and is already installed on Colab, as is build-essential, which provides g++ as well."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "pZQLuKHZJ0hW"
      },
      "outputs": [],
      "source": [
        "!nvcc --version"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "gjUMM5ygKngM"
      },
      "source": [
        "## Get the code\n",
        "This notebook will show the files inline, and you can also checkout the repository containing the c++ files to compile.\n",
        "```\n",
        "!git clone https://github.com/evelynmitchell/cuda-on-colab\n",
        "```"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "The simple c++ example of adding the elements of two arrays, without gpu."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {},
      "outputs": [],
      "source": [
        "%%file /tmp/simple.cpp\n",
        "#include <iostream>\n",
        "#include <math.h>\n",
        "\n",
        "// function to add the elements of two arrays\n",
        "void add(int n, float *x, float *y)\n",
        "{\n",
        "  for (int i = 0; i < n; i++)\n",
        "      y[i] = x[i] + y[i];\n",
        "}\n",
        "\n",
        "int main(void)\n",
        "{\n",
        "  int N = 1<<20; // 1M elements\n",
        "\n",
        "  float *x = new float[N];\n",
        "  float *y = new float[N];\n",
        "\n",
        "  // initialize x and y arrays on the host\n",
        "  for (int i = 0; i < N; i++) {\n",
        "    x[i] = 1.0f;\n",
        "    y[i] = 2.0f;\n",
        "  }\n",
        "\n",
        "  // Run kernel on 1M elements on the CPU\n",
        "  add(N, x, y);\n",
        "\n",
        "  // Check for errors (all values should be 3.0f)\n",
        "  float maxError = 0.0f;\n",
        "  for (int i = 0; i < N; i++)\n",
        "    maxError = fmax(maxError, fabs(y[i]-3.0f));\n",
        "  std::cout << \"Max error: \" << maxError << std::endl;\n",
        "\n",
        "  // Free memory\n",
        "  delete [] x;\n",
        "  delete [] y;\n",
        "\n",
        "  return 0;\n",
        "}"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "z4LGcpR3Nm91"
      },
      "source": [
        "## Build the code for CPU"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "QHSDLMN4Nmjg"
      },
      "outputs": [],
      "source": [
        "# compile the code checked out from the repoository to\n",
        "# create the binary /content/cuda-on-colab/src/simple\n",
        "# add the executable bit, and then run it.\n",
        "#!g++ /content/cuda-on-colab/src/simple.cpp -o simple\n",
        "# !chmod +x ./simple\n",
        "# !./simple\n",
        "\n",
        "# compile the code in the cell to create the binary /tmp/simple, \n",
        "#  and then run it.\n",
        "! g++ /tmp/simple.cpp -o /tmp/simple && /tmp/simple\n",
        "\n"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "3myEKLJzShQn"
      },
      "source": [
        "## Compile to a CUDA kernel\n",
        "\n",
        "Adding the  ```__global__``` specifier to a function indicates it will be compiled to a CUDA kernel and run on a GPU processor.\n",
        "\n",
        "This code fails when it's compiled due to an error in how it is called. The error and fix follow this section."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {},
      "outputs": [],
      "source": [
        "%%file /tmp/simple_cuda.cu\n",
        "#include <iostream>\n",
        "#include <math.h>\n",
        "\n",
        "// CUDA Kernel function to add the elements of two arrays on the GPU\n",
        "__global__\n",
        "void add(int n, float *x, float *y)\n",
        "{\n",
        "  for (int i = 0; i < n; i++)\n",
        "      y[i] = x[i] + y[i];\n",
        "}\n",
        "\n",
        "int main(void)\n",
        "{\n",
        "  int N = 1<<20; // 1M elements\n",
        "\n",
        "  float *x = new float[N];\n",
        "  float *y = new float[N];\n",
        "\n",
        "  // initialize x and y arrays on the host\n",
        "  for (int i = 0; i < N; i++) {\n",
        "    x[i] = 1.0f;\n",
        "    y[i] = 2.0f;\n",
        "  }\n",
        "\n",
        "  // Run kernel on 1M elements on the CPU\n",
        "  add(N, x, y);\n",
        "\n",
        "  // Check for errors (all values should be 3.0f)\n",
        "  float maxError = 0.0f;\n",
        "  for (int i = 0; i < N; i++)\n",
        "    maxError = fmax(maxError, fabs(y[i]-3.0f));\n",
        "  std::cout << \"Max error: \" << maxError << std::endl;\n",
        "\n",
        "  // Free memory\n",
        "  delete [] x;\n",
        "  delete [] y;\n",
        "\n",
        "  return 0;\n",
        "}"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "rYgRPRX7Sbuf"
      },
      "outputs": [],
      "source": [
        "# from the repository\n",
        "# !nvcc /content/cuda-on-colab/src/simple_cuda.cu -o simple_cuda && /tmp/simple_cuda\n",
        "\n",
        "# from the cell\n",
        "! nvcc /tmp/simple_cuda.cu -o /tmp/simple_cuda && /tmp/simple_cuda"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "K0QUba4rrIgO"
      },
      "source": [
        "## Configure kernel launch\n",
        "\n",
        "The error from the prior version of the compilation \"__global__ function call must be configured\" is corrected by adding kernel launch parameters <<<gridsize,blocksize>>> to the function."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {},
      "outputs": [],
      "source": [
        "%%file /tmp/simple_cuda_kernel_launch.cu\n",
        "#include <iostream>\n",
        "#include <math.h>\n",
        "\n",
        "// CUDA Kernel function to add the elements of two arrays on the GPU\n",
        "\n",
        "__global__\n",
        "void add(int n, float *x, float *y)\n",
        "{\n",
        "  for (int i = 0; i < n; i++)\n",
        "      y[i] = x[i] + y[i];\n",
        "}\n",
        "\n",
        "int main(void)\n",
        "{\n",
        "  int N = 1<<20; // 1M elements\n",
        "\n",
        "  float *x = new float[N];\n",
        "  float *y = new float[N];\n",
        "\n",
        "  // initialize x and y arrays on the host\n",
        "  for (int i = 0; i < N; i++) {\n",
        "    x[i] = 1.0f;\n",
        "    y[i] = 2.0f;\n",
        "  }\n",
        "\n",
        "  // Run kernel on 1M elements on the GPU\n",
        "  // <<< (gridsize), (blocksize) >>>\n",
        "  // <<<1,1>>> means 1 block with 1 thread\n",
        "  add<<<1,1>>>(N, x, y);\n",
        "\n",
        "  // Check for errors (all values should be 3.0f)\n",
        "  float maxError = 0.0f;\n",
        "  for (int i = 0; i < N; i++)\n",
        "    maxError = fmax(maxError, fabs(y[i]-3.0f));\n",
        "  std::cout << \"Max error: \" << maxError << std::endl;\n",
        "\n",
        "  // Free memory\n",
        "  delete [] x;\n",
        "  delete [] y;\n",
        "\n",
        "  return 0;\n",
        "}"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "pbGjUBb6rGdn"
      },
      "outputs": [],
      "source": [
        "# from the repository\n",
        "# !nvcc /content/cuda-on-colab/src/simple_cuda_kernel_launch.cu -o simple_cuda_kernal_launch\n",
        "\n",
        "# from the cell\n",
        "!nvcc /tmp/simple_cuda_kernel_launch.cu -o /tmp/simple_cuda_kernal_launch && /tmp/simple_cuda_kernal_launch\n"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "## Configure kernel threads\n"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {},
      "outputs": [],
      "source": [
        "%%file /tmp/simple_cuda_kernel_threads.cu\n",
        "#include <iostream>\n",
        "#include <math.h>\n",
        "\n",
        "// CUDA Kernel function to add the elements of two arrays on the GPU\n",
        "\n",
        "__global__\n",
        "void add(int n, float *x, float *y)\n",
        "{\n",
        "  for (int i = 0; i < n; i++)\n",
        "      y[i] = x[i] + y[i];\n",
        "}\n",
        "\n",
        "int main(void)\n",
        "{\n",
        "  int N = 1<<20; // 1M elements\n",
        "\n",
        "  float *x = new float[N];\n",
        "  float *y = new float[N];\n",
        "\n",
        "  // initialize x and y arrays on the host\n",
        "  for (int i = 0; i < N; i++) {\n",
        "    x[i] = 1.0f;\n",
        "    y[i] = 2.0f;\n",
        "  }\n",
        "\n",
        "  // Run kernel on 1M elements on the GPU\n",
        "  // <<< (gridsize), (blocksize) >>>\n",
        "  // <<<1,1>>> means 1 block with 1 thread\n",
        "  // \"CUDA GPUs run kernels using blocks of threads that are a multiple of \n",
        "  // 32 in size, so 256 threads is a reasonable size to choose.\"\"\n",
        "  add<<<1,256>>>(N, x, y);\n",
        "\n",
        "  // Check for errors (all values should be 3.0f)\n",
        "  float maxError = 0.0f;\n",
        "  for (int i = 0; i < N; i++)\n",
        "    maxError = fmax(maxError, fabs(y[i]-3.0f));\n",
        "  std::cout << \"Max error: \" << maxError << std::endl;\n",
        "\n",
        "  // Free memory\n",
        "  delete [] x;\n",
        "  delete [] y;\n",
        "\n",
        "  return 0;\n",
        "}"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {},
      "outputs": [],
      "source": [
        "# from the repository\n",
        "# !nvcc /content/cuda-on-colab/src/simple_cuda_kernel_threads.cu -o simple_cuda_kernal_threads\n",
        "\n",
        "# from the cell\n",
        "!nvcc /tmp/simple_cuda_kernel_threads.cu -o /tmp/simple_cuda_kernal_threads && /tmp/simple_cuda_kernal_threads\n"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "## Profile the CUDA code\n",
        "\n",
        "nvprof is the nvidia profiler for CUDA code. "
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "OeOa2oOjZYhm"
      },
      "outputs": [],
      "source": [
        "# %cd /content/cuda-on-colab to run from the repository\n",
        "# %cd /tmp to run from the cell\n",
        "!nvprof /tmp/simple_cuda_kernel_launch"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "## Memory profiling\n",
        "\n",
        "nvprof can also be used to profile memory usage. First we compile an\n",
        "example that uses a lot of memory, then we profile it."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {},
      "outputs": [],
      "source": [
        "%%file /tmp/simple_cuda_memory_alloc.cu\n",
        "#include <iostream>\n",
        "#include <math.h>\n",
        "// Kernel function to add the elements of two arrays\n",
        "__global__\n",
        "void add(int n, float *x, float *y)\n",
        "{\n",
        "  for (int i = 0; i < n; i++)\n",
        "    y[i] = x[i] + y[i];\n",
        "}\n",
        "\n",
        "int main(void)\n",
        "{\n",
        "  int N = 1<<20;\n",
        "  float *x, *y;\n",
        "\n",
        "  // Allocate Unified Memory – accessible from CPU or GPU\n",
        "  cudaMallocManaged(&x, N*sizeof(float));\n",
        "  cudaMallocManaged(&y, N*sizeof(float));\n",
        "\n",
        "  // initialize x and y arrays on the host\n",
        "  for (int i = 0; i < N; i++) {\n",
        "    x[i] = 1.0f;\n",
        "    y[i] = 2.0f;\n",
        "  }\n",
        "\n",
        "  // Run kernel on 1M elements on the GPU\n",
        "  add<<<1, 1>>>(N, x, y);\n",
        "\n",
        "  // Wait for GPU to finish before accessing on host\n",
        "  cudaDeviceSynchronize();\n",
        "\n",
        "  // Check for errors (all values should be 3.0f)\n",
        "  float maxError = 0.0f;\n",
        "  for (int i = 0; i < N; i++)\n",
        "    maxError = fmax(maxError, fabs(y[i]-3.0f));\n",
        "  std::cout << \"Max error: \" << maxError << std::endl;\n",
        "\n",
        "  // Free memory\n",
        "  cudaFree(x);\n",
        "  cudaFree(y);\n",
        "  \n",
        "  return 0;\n",
        "}"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "0_fXXIOJiye_"
      },
      "outputs": [],
      "source": [
        "# from the repository\n",
        "# !nvcc /content/cuda-on-colab/src/simple_cuda_memory_alloc.cu -o simple_cuda_memory_alloc\n",
        "\n",
        "# from the cell\n",
        "!nvcc /tmp/simple_cuda_memory_alloc.cu -o /tmp/simple_cuda_memory_alloc && /tmp/simple_cuda_memory_alloc \n"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "5ozN41O-jCAi"
      },
      "outputs": [],
      "source": [
        "# run the executable with nvprof\n",
        "!nvprof /tmp/simple_cuda_memory_alloc"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": []
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "The profiling output will look like this:\n",
        "\n",
        "```\n",
        "==968== NVPROF is profiling process 968, command: /tmp/simple_cuda_memory_alloc\n",
        "Max error: 0\n",
        "==968== Profiling application: /tmp/simple_cuda_memory_alloc\n",
        "==968== Profiling result:\n",
        "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
        " GPU activities:  100.00%  100.03ms         1  100.03ms  100.03ms  100.03ms  add(int, float*, float*)\n",
        "      API calls:   61.56%  161.55ms         2  80.777ms  56.205us  161.50ms  cudaMallocManaged\n",
        "                   38.12%  100.04ms         1  100.04ms  100.04ms  100.04ms  cudaDeviceSynchronize\n",
        "                    0.18%  462.05us         2  231.02us  228.77us  233.27us  cudaFree\n",
        "                    0.08%  197.51us         1  197.51us  197.51us  197.51us  cudaLaunchKernel\n",
        "                    0.05%  133.89us       114  1.1740us     134ns  51.337us  cuDeviceGetAttribute\n",
        "                    0.01%  14.762us         1  14.762us  14.762us  14.762us  cuDeviceGetName\n",
        "                    0.00%  5.3400us         1  5.3400us  5.3400us  5.3400us  cuDeviceTotalMem\n",
        "                    0.00%  5.1380us         1  5.1380us  5.1380us  5.1380us  cuDeviceGetPCIBusId\n",
        "                    0.00%  1.6230us         3     541ns     160ns  1.1080us  cuDeviceGetCount\n",
        "                    0.00%  1.2110us         2     605ns     223ns     988ns  cuDeviceGet\n",
        "                    0.00%     355ns         1     355ns     355ns     355ns  cuModuleGetLoadingMode\n",
        "                    0.00%     226ns         1     226ns     226ns     226ns  cuDeviceGetUuid\n",
        "\n",
        "==968== Unified Memory profiling result:\n",
        "Device \"Tesla V100-SXM2-16GB (0)\"\n",
        "   Count  Avg Size  Min Size  Max Size  Total Size  Total Time  Name\n",
        "      48  170.67KB  4.0000KB  0.9961MB  8.000000MB  837.4650us  Host To Device\n",
        "      24  170.67KB  4.0000KB  0.9961MB  4.000000MB  356.8290us  Device To Host\n",
        "      12         -         -         -           -  3.111407ms  Gpu page fault groups\n",
        "Total CPU Page faults: 36\n",
        "```"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "Implement an RGB to grayscale conversion kernel that matches the reference implementation.\n",
        "The kernel should convert square RGB images with even sizes to grayscale using the standard coefficients:\n",
        "Y = 0.2989 R + 0.5870 G + 0.1140 B\n",
        "\n",
        "Input: RGB tensor of shape (H, W, 3) with values in [0, 1]\n",
        "Output: Grayscale tensor of shape (H, W) with values in [0, 1]\n"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {},
      "outputs": [],
      "source": [
        "import triton\n",
        "import torch\n",
        "\n",
        "import triton.language as tl\n",
        "\n",
        "@triton.jit\n",
        "def rgb_to_grayscale_kernel(input_ptr, output_ptr, H, W):\n",
        "    pid = tl.program_id(0)\n",
        "    row = pid // W\n",
        "    col = pid % W\n",
        "\n",
        "    if row < H and col < W:\n",
        "        r = tl.load(input_ptr + (row * W + col) * 3 + 0)\n",
        "        g = tl.load(input_ptr + (row * W + col) * 3 + 1)\n",
        "        b = tl.load(input_ptr + (row * W + col) * 3 + 2)\n",
        "        gray = 0.2989 * r + 0.5870 * g + 0.1140 * b\n",
        "        tl.store(output_ptr + row * W + col, gray)\n",
        "\n",
        "def rgb_to_grayscale(input_tensor):\n",
        "    H, W, _ = input_tensor.shape\n",
        "    output_tensor = torch.empty((H, W), dtype=input_tensor.dtype, device=input_tensor.device)\n",
        "    grid = (H * W,)\n",
        "    rgb_to_grayscale_kernel[grid](input_tensor, output_tensor, H, W)\n",
        "    return output_tensor\n",
        "\n",
        "# Example usage\n",
        "input_tensor = torch.rand((4, 4, 3), dtype=torch.float32, device='cuda')\n",
        "output_tensor = rgb_to_grayscale(input_tensor)\n",
        "print(output_tensor)"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {},
      "outputs": [],
      "source": []
    }
  ],
  "metadata": {
    "accelerator": "GPU",
    "colab": {
      "gpuType": "V100",
      "machine_shape": "hm",
      "private_outputs": true,
      "provenance": []
    },
    "kernelspec": {
      "display_name": "Python 3",
      "name": "python3"
    },
    "language_info": {
      "name": "python"
    }
  },
  "nbformat": 4,
  "nbformat_minor": 0
}