Week8/Conv1D_Tiled.ipynb

{
  "nbformat": 4,
  "nbformat_minor": 0,
  "metadata": {
    "colab": {
      "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/Week8/Conv1D_Tiled.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": "b296ae82-f840-4232-d156-f1d4a2f1f631"
      },
      "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-qlw1cedb\n",
            "  Running command git clone --filter=blob:none --quiet https://github.com/andreinechaev/nvcc4jupyter.git /tmp/pip-req-build-qlw1cedb\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=4304 sha256=7599776fbb7d0c2e49b139df96ae3b4523eb280f3202cb2916501f636fcefcb5\n",
            "  Stored in directory: /tmp/pip-ephem-wheel-cache-uthtw17u/wheels/f3/08/cc/e2b5b0e1c92df07dbb50a6f024a68ce090f5e7b2316b41756d\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",
        "\n",
        "inline cudaError_t checkCudaErr(cudaError_t err, const char* msg) {\n",
        "  if (err != cudaSuccess) {\n",
        "    fprintf(stderr, \"CUDA Runtime error at %s: %s\\n\", msg, cudaGetErrorString(err));\n",
        "  }\n",
        "  return err;\n",
        "}\n",
        "\n",
        "#define MASK_WIDTH 5\n",
        "#define TILE_SIZE MASK_WIDTH/2\n",
        "\n",
        "__constant__ float M[MASK_WIDTH]; \n",
        "\n",
        "__global__ void convolution_1D_basic_kernel(float* N, float* P,  int Width)\n",
        "{\n",
        "  int i = blockIdx.x*blockDim.x+threadIdx.x;  \n",
        "  float Pvalue=0; \n",
        "  int N_start_point = i-(MASK_WIDTH/2); \n",
        " \n",
        "  for (int j=0; j<MASK_WIDTH; j++) \n",
        "  {\n",
        "      if(N_start_point+j>=0 && N_start_point+j< Width) \n",
        "      { \n",
        "        Pvalue += N[N_start_point+j]*M[j]; \n",
        "      } \n",
        "  }\n",
        "  P[i]=Pvalue; \n",
        "} \n",
        "\n",
        "__global__ void convolution_1D_basic_kernel_tiled(float* N, float* P, int Width)\n",
        "{\n",
        "  int i = blockIdx.x*blockDim.x+threadIdx.x;  \n",
        " __shared__ float N_ds[TILE_SIZE + MASK_WIDTH - 1];\n",
        "\n",
        "  int n = TILE_SIZE;\n",
        " \n",
        "  int halo_index_left = (blockIdx.x - 1) * blockDim.x + threadIdx.x;\n",
        "  if(threadIdx.x >= (blockDim.x - n))\n",
        "  {\n",
        "      N_ds[threadIdx.x - (blockDim.x - n)] = (halo_index_left < 0)? 0: N[halo_index_left];\n",
        "  }\n",
        "  \n",
        "  N_ds[n + threadIdx.x] = N[i];\n",
        " \n",
        "  int halo_index_right = (blockIdx.x + 1) * blockDim.x + threadIdx.x;\n",
        " \n",
        "  if(threadIdx.x < n)\n",
        "  {\n",
        "      N_ds[n+ blockDim.x + threadIdx.x] = (halo_index_right >= Width)? 0: N[halo_index_right];\n",
        "  }\n",
        " \n",
        "  __syncthreads();\n",
        " \n",
        "  float Pvalue = 0;\n",
        "  for (int j=0; j<MASK_WIDTH; j++) \n",
        "  {\n",
        "     Pvalue += N_ds[threadIdx.x + j]*M[j];       \n",
        "  }\n",
        "  P[i]=Pvalue; \n",
        "}\n",
        " \n",
        "int main() { \n",
        "\t\tfloat   *a,  *c=0, *tiled_c=0; \n",
        "    const int N = 16; \n",
        "    const int threadsPerBlock = 4;\n",
        "\n",
        "    float h_M[MASK_WIDTH]={3,4,5,4,3};\n",
        "\n",
        "    // Allocate Unified Memory -- accessible from CPU or GPU\n",
        "    checkCudaErr(cudaMallocManaged(&a, N*sizeof(float)), \"cudaMallocManaged a\");\n",
        "    checkCudaErr(cudaMallocManaged(&c, N*sizeof(float)), \"cudaMallocManaged c\");\n",
        "    checkCudaErr(cudaMallocManaged(&tiled_c, N*sizeof(float)), \"cudaMallocManaged tiled_c\");\n",
        "     \n",
        "    // fill in the memory with data\n",
        "    for (int i=0; i<N; i++) {\n",
        "        a[i] = i+1;\n",
        "        c[i] = 0; \n",
        "        tiled_c[i] = 0;\n",
        "    }  \n",
        " \n",
        "    cudaMemcpyToSymbol(M,h_M,MASK_WIDTH*sizeof(float));\n",
        "    const int blocksPerGrid =  (N / threadsPerBlock);\n",
        " \n",
        "    // Prefetch the data to the GPU\n",
        "    int device = -1;\n",
        "    cudaGetDevice(&device);\n",
        "    cudaMemPrefetchAsync(a, N*sizeof(float), device, NULL);\n",
        "    cudaMemPrefetchAsync(M, MASK_WIDTH*sizeof(float), device, NULL);\n",
        "\n",
        "    //lets time the conv1D kernel\n",
        "    cudaEvent_t start, stop;\n",
        "    float gpu_elapsed_time_ms=0;\n",
        "    float gpu_elapsed_time_tiled_ms=0;\n",
        "\n",
        "    cudaEventCreate(&start);\n",
        "    cudaEventCreate(&stop);\n",
        "\n",
        "    // start to count execution time \n",
        "    cudaEventRecord(start, 0);\n",
        "    convolution_1D_basic_kernel<<<blocksPerGrid,threadsPerBlock>>>(a, c, N);\n",
        "    \n",
        "    cudaDeviceSynchronize();\n",
        " \n",
        "    // time counting terminate\n",
        "    cudaEventRecord(stop, 0);\n",
        "    cudaEventSynchronize(stop);\n",
        "\n",
        "    // compute time elapse on GPU computing\n",
        "    cudaEventElapsedTime(&gpu_elapsed_time_ms, start, stop);\n",
        "\n",
        "    //call the tiled Conv1D function\n",
        "    // start to count execution time \n",
        "    cudaEventRecord(start, 0);\n",
        "    convolution_1D_basic_kernel_tiled<<<blocksPerGrid,threadsPerBlock>>>(a, tiled_c, N);\n",
        "    \n",
        "    cudaDeviceSynchronize();\n",
        " \n",
        "    // time counting terminate\n",
        "    cudaEventRecord(stop, 0);\n",
        "    cudaEventSynchronize(stop);\n",
        "\n",
        "    // compute time elapse on GPU computing\n",
        "    cudaEventElapsedTime(&gpu_elapsed_time_tiled_ms, start, stop);\n",
        "\n",
        "    //output the result\n",
        "    printf(\"Results:\\nConv1D: %3.3f msecs: \\nConv1D Tiled: %3.3f msecs\\n\", gpu_elapsed_time_ms, gpu_elapsed_time_tiled_ms);\n",
        "    printf(\" p[i]  | conv[i]| conv_tiled[i]\\n\");\n",
        "    printf(\"-------+--------+--------------\\n\");\n",
        "    for(int i=0; i<N; ++i)\n",
        "      printf(\" p[%2d] | %6.2f | %6.2f\\n\", i, c[i], tiled_c[i]);\n",
        "    \n",
        "    printf(\"-------+--------+--------------\\n\");\n",
        "\n",
        "    // free memory on the gpu side\n",
        "    checkCudaErr( cudaFree( a ) , \"cudaFree a\"); \n",
        "    checkCudaErr( cudaFree( c ) , \"cudaFree c\");\n",
        "    checkCudaErr( cudaFree( tiled_c ) , \"cudaFree tiled_c\");\n",
        "\t\tcheckCudaErr( cudaDeviceReset(), \"cudaDeviceReset\");\n",
        "\n",
        "\t\treturn 0;\n",
        "}"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "GYZ627HO-XKd",
        "outputId": "97cce0a4-9a90-4c0a-e5a4-e54dad8f3a0d"
      },
      "execution_count": 33,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Results:\n",
            "Conv1D: 0.030 msecs: \n",
            "Conv1D Tiled: 0.012 msecs\n",
            " p[i]  | conv[i]| conv_tiled[i]\n",
            "-------+--------+--------------\n",
            " p[ 0] |  22.00 |  22.00\n",
            " p[ 1] |  38.00 |  38.00\n",
            " p[ 2] |  57.00 |  57.00\n",
            " p[ 3] |  76.00 |  76.00\n",
            " p[ 4] |  95.00 |  95.00\n",
            " p[ 5] | 114.00 | 114.00\n",
            " p[ 6] | 133.00 | 133.00\n",
            " p[ 7] | 152.00 | 152.00\n",
            " p[ 8] | 171.00 | 171.00\n",
            " p[ 9] | 190.00 | 190.00\n",
            " p[10] | 209.00 | 209.00\n",
            " p[11] | 228.00 | 228.00\n",
            " p[12] | 247.00 | 247.00\n",
            " p[13] | 266.00 | 266.00\n",
            " p[14] | 234.00 | 234.00\n",
            " p[15] | 182.00 | 182.00\n",
            "-------+--------+--------------\n",
            "\n"
          ]
        }
      ]
    }
  ]
}