Week10/SpMV_Formats.ipynb

{
  "nbformat": 4,
  "nbformat_minor": 0,
  "metadata": {
    "colab": {
      "name": "SpMV_Formats.ipynb",
      "provenance": [],
      "include_colab_link": true
    },
    "kernelspec": {
      "name": "python3",
      "display_name": "Python 3"
    },
    "language_info": {
      "name": "python"
    },
    "accelerator": "GPU"
  },
  "cells": [
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "view-in-github",
        "colab_type": "text"
      },
      "source": [
        "<a href=\"https://colab.research.google.com/github/mmmovania/CUDA_Spring2023/blob/main/Week10/SpMV_Formats.ipynb\" target=\"_parent\"><img src=\"https://colab.research.google.com/assets/colab-badge.svg\" alt=\"Open In Colab\"/></a>"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 1,
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "OB2MY3st-VUm",
        "outputId": "b7b870c4-ab8f-4f3e-f5e3-b8ab691860f4"
      },
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Looking in indexes: https://pypi.org/simple, https://us-python.pkg.dev/colab-wheels/public/simple/\n",
            "Collecting git+https://github.com/andreinechaev/nvcc4jupyter.git\n",
            "  Cloning https://github.com/andreinechaev/nvcc4jupyter.git to /tmp/pip-req-build-yd7dgvab\n",
            "  Running command git clone --filter=blob:none --quiet https://github.com/andreinechaev/nvcc4jupyter.git /tmp/pip-req-build-yd7dgvab\n",
            "  Resolved https://github.com/andreinechaev/nvcc4jupyter.git to commit aac710a35f52bb78ab34d2e52517237941399eff\n",
            "  Preparing metadata (setup.py) ... \u001b[?25l\u001b[?25hdone\n",
            "Building wheels for collected packages: NVCCPlugin\n",
            "  Building wheel for NVCCPlugin (setup.py) ... \u001b[?25l\u001b[?25hdone\n",
            "  Created wheel for NVCCPlugin: filename=NVCCPlugin-0.0.2-py3-none-any.whl size=4305 sha256=942dbf3699e25ba9717b5dd71d6d5c605ad70b12727e62ffc46e22bf53182f7b\n",
            "  Stored in directory: /tmp/pip-ephem-wheel-cache-s_ispawi/wheels/db/c1/1f/a2bb07bbb4a1ce3c43921252aeafaa6205f08637e292496f04\n",
            "Successfully built NVCCPlugin\n",
            "Installing collected packages: NVCCPlugin\n",
            "Successfully installed NVCCPlugin-0.0.2\n",
            "created output directory at /content/src\n",
            "Out bin /content/result.out\n"
          ]
        }
      ],
      "source": [
        "!pip install git+https://github.com/andreinechaev/nvcc4jupyter.git\n",
        "%load_ext nvcc_plugin"
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "%%cu\n",
        "#include <stdio.h>\n",
        " \n",
        "inline cudaError_t checkCudaErr(cudaError_t err, const char* msg) {\n",
        "\tif (err != cudaSuccess) {\n",
        "\t\tfprintf(stderr, \"CUDA Runtime error at %s: %s\\n\", msg, cudaGetErrorString(err));\n",
        "\t}\n",
        "\treturn err;\n",
        "}\n",
        "\n",
        "// Parallel SpMV using CSR format\n",
        "__global__ void spmv_csr(int num_rows, int *data, int *col_index, int* row_ptr, int *x, int *y)\n",
        " {\n",
        "    int row = threadIdx.x + blockIdx.x * blockDim.x;\n",
        "\t\t\n",
        "\t\tif( row < num_rows)\n",
        "\t\t{\n",
        "\t\t\t\tint dotProduct = 0;\n",
        "        const int row_start = row_ptr[row];\n",
        "        const int row_end = row_ptr[row + 1];\n",
        "        \n",
        "        for (int j = row_start; j < row_end; j++) \n",
        "\t\t\t\t{\n",
        "            dotProduct += data[j] * x[col_index[j]];\n",
        "        }\n",
        "        \n",
        "        y[row] = dotProduct;\t\n",
        "    }  \n",
        "} \n",
        "\n",
        "// Parallel SpMV using ELL format\n",
        "__global__ void spmv_ell(int num_rows, int *data, int *col_index, int num_elements, int *x, int *y)\n",
        " { \n",
        "\t\tint row = threadIdx.x + blockIdx.x * blockDim.x;\n",
        "\t\tif( row < num_rows)\n",
        "\t\t{\n",
        "\t\t\t\tint dotProduct = 0;\n",
        "         \n",
        "        for (int j = 0; j < num_elements ; j++) \n",
        "\t\t\t\t{\n",
        "            dotProduct += data[row+ j*num_rows] * x[col_index[row + j*num_rows]];\n",
        "        }\n",
        "        \n",
        "        y[row] = dotProduct;\t \n",
        "    }  \n",
        "}\n",
        "\n",
        "// Parallel SpMV using COO format\n",
        "__global__ void spmv_coo(int num_rows, int *data, int *col_index, int* row_index,int num_elements, int *x, int *y)\n",
        " {\n",
        "    int row = threadIdx.x + blockIdx.x * blockDim.x;\n",
        "\t\t\n",
        "\t\tif( row < num_rows)\n",
        "\t\t{\n",
        "\t\t\t\tint dotProduct = 0;\n",
        "    \n",
        "        for (int j = 0; j < num_elements; j++)  \n",
        "\t\t\t\t{\n",
        "            dotProduct += data[row_index[j]] * x[col_index[j]];\n",
        "        }\n",
        "        \n",
        "        y[row] = dotProduct;\t\n",
        "    }  \n",
        "} \n",
        "\n",
        "int main() \n",
        "{ \n",
        "\t//data for CSR and COO formats\n",
        "\tint data[7] = {3, 1, 2, 4, 1, 1, 1}; \n",
        "\tint row_ptr[5] = {0, 2, 2, 5, 7};\n",
        "\tint col_index[7] = {0, 2, 1, 2, 3, 0, 3};\n",
        "\tint row_index[7] = {0, 0, 2, 2, 2, 3, 3};\n",
        "\t\n",
        "\t//data for ELL format\n",
        "\tint dataPaddedTransposed[12]={3, 0, 2, 1, 1, 0, 4, 1, 0, 0, 1, 0 };\n",
        "\tint col_indexPaddedTransposed[12] = { 0, 0, 1, 0, 2, 0, 2, 3, 0, 0, 3, 0};\n",
        "\t \n",
        "\tconst int N = 4;\n",
        "\tint* dataIn;\n",
        "\tint* dataInELL;\n",
        "\tint* colIndex;\n",
        "\tint* rowIndex;\n",
        "\tint* colIndexELL;\n",
        "\tint* rowPtr;\n",
        "\tint* x;\n",
        "\tint* y;\n",
        "\n",
        "\t// Allocate Unified Memory -- accessible from CPU or GPU\n",
        "\tcheckCudaErr(cudaMallocManaged(&x, N*sizeof(int)), \"cudaMallocManaged1\");\n",
        "\tcheckCudaErr(cudaMallocManaged(&y, N*sizeof(int)), \"cudaMallocManaged2\");\n",
        "\tcheckCudaErr(cudaMallocManaged(&dataIn, 7*sizeof(int)), \"cudaMallocManaged3\");\n",
        "\tcheckCudaErr(cudaMallocManaged(&colIndex, 7*sizeof(int)), \"cudaMallocManaged4\");\n",
        "\tcheckCudaErr(cudaMallocManaged(&rowPtr, 5*sizeof(int)), \"cudaMallocManaged5\");\n",
        "\tcheckCudaErr(cudaMallocManaged(&dataInELL, 12*sizeof(int)), \"cudaMallocManaged6\");\n",
        "\tcheckCudaErr(cudaMallocManaged(&colIndexELL, 12*sizeof(int)), \"cudaMallocManaged7\");\n",
        "\tcheckCudaErr(cudaMallocManaged(&rowIndex, 7*sizeof(int)), \"cudaMallocManaged8\");\n",
        "\n",
        "\t// fill in the memory with data\n",
        "\tfor (int i=0; i<12; i++) {\n",
        "\t\tif(i<N)\n",
        "\t\t{\n",
        "\t\t\tx[i] = i+1;\n",
        "\t\t\ty[i] = 0;\n",
        "\t\t} \t\n",
        "\t\tif(i<5)\n",
        "\t\t\trowPtr[i] = row_ptr[i];\n",
        "\t\t\n",
        "\t\tif(i<7)\n",
        "\t\t{ \n",
        "\t\t\tdataIn[i] = data[i];\n",
        "\t\t\tcolIndex[i] = col_index[i];\n",
        "\t\t\trowIndex[i] = row_index[i];\n",
        "\t\t}\n",
        "\n",
        "\t\tdataInELL[i] = dataPaddedTransposed[i];\t \n",
        "\t\tcolIndexELL[i] = col_indexPaddedTransposed[i];\n",
        "\t} \n",
        "\n",
        "  int blocksPerGrid = 1;\n",
        "\tint threadsPerBlock = 32;\n",
        " \n",
        "\t// call the CSR kernel\n",
        "  spmv_csr<<<blocksPerGrid, threadsPerBlock>>>(N, dataIn, colIndex, rowPtr, x, y);     \n",
        "\tcudaDeviceSynchronize();\n",
        "\tcudaError_t err = cudaGetLastError();\n",
        "\tif(err != 0)\n",
        "\t\t printf(\"CUDA Error: %s\", cudaGetErrorString(err));\n",
        "  \n",
        "\tputs(\"Output using CSR Format: \");\n",
        "\tfor(int i=0; i<N; ++i)\n",
        " \t{ \n",
        "     printf(\"x[i]: %d, y[i]: %d\\n\", x[i], y[i]);\n",
        "\t\t y[i] = 0; //preparing for next kernel \n",
        "\t}\n",
        "\t\n",
        "\t//now lets try the ELL format\n",
        "\tspmv_ell<<<blocksPerGrid, threadsPerBlock>>>(N, dataInELL, colIndexELL, 12, x, y);       \n",
        "\tcudaDeviceSynchronize();\n",
        "\terr = cudaGetLastError();\n",
        "\tif(err != 0)\n",
        "\t\t printf(\"CUDA Error: %s\", cudaGetErrorString(err));\n",
        "  \n",
        "\tputs(\"\\nOutput using ELL Format: \");\n",
        "\tfor(int i=0; i<N; ++i)\n",
        " \t{ \n",
        "     printf(\"x[i]: %d, y[i]: %d\\n\", x[i], y[i]);\n",
        "\t\t y[i] = 0; //preparing for next kernel\n",
        "\t}\n",
        "\t\n",
        "\t//now lets try the COO format\n",
        "\tspmv_coo<<<blocksPerGrid, threadsPerBlock>>>(N, dataIn, colIndex, rowIndex, 7, x, y);       \n",
        "\tcudaDeviceSynchronize();\n",
        "\terr = cudaGetLastError();\n",
        "\tif(err != 0)\n",
        "\t\t printf(\"CUDA Error: %s\", cudaGetErrorString(err));\n",
        "  \n",
        "\tputs(\"\\nOutput using COO Format: \");\n",
        "\tfor(int i=0; i<N; ++i)\n",
        " \t{ \n",
        "     printf(\"x[i]: %d, y[i]: %d\\n\", x[i], y[i]); \n",
        "\t}\n",
        "\n",
        "\t// free memory on the gpu side\n",
        "\tcheckCudaErr( cudaFree( x ), \"cudaFree1\");\n",
        "\tcheckCudaErr( cudaFree( y ), \"cudaFree2\");\n",
        "\tcheckCudaErr( cudaFree( dataIn ), \"cudaFree3\"); \n",
        "\tcheckCudaErr( cudaFree( colIndex ), \"cudaFree4\");\n",
        "\tcheckCudaErr( cudaFree( rowPtr ), \"cudaFree5\");\n",
        "\tcheckCudaErr( cudaFree( dataInELL ), \"cudaFree6\");\n",
        "\tcheckCudaErr( cudaFree( colIndexELL ), \"cudaFree7\");\n",
        "\tcheckCudaErr( cudaDeviceReset(), \"cudaDeviceReset\");\n",
        "\n",
        "\treturn 0;\n",
        "}"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "GYZ627HO-XKd",
        "outputId": "9c959445-68b5-4ac3-a81f-d8a39e1b3788"
      },
      "execution_count": 2,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Output using CSR Format: \n",
            "x[i]: 1, y[i]: 6\n",
            "x[i]: 2, y[i]: 0\n",
            "x[i]: 3, y[i]: 20\n",
            "x[i]: 4, y[i]: 5\n",
            "\n",
            "Output using ELL Format: \n",
            "x[i]: 1, y[i]: 6\n",
            "x[i]: 2, y[i]: 0\n",
            "x[i]: 3, y[i]: 20\n",
            "x[i]: 4, y[i]: 5\n",
            "\n",
            "Output using COO Format: \n",
            "x[i]: 1, y[i]: 50\n",
            "x[i]: 2, y[i]: 50\n",
            "x[i]: 3, y[i]: 50\n",
            "x[i]: 4, y[i]: 50\n",
            "\n"
          ]
        }
      ]
    }
  ]
}