Week8/PrefixSum_WorkEfficient.ipynb

{
  "nbformat": 4,
  "nbformat_minor": 0,
  "metadata": {
    "colab": {
      "name": "PrefixSum_WorkEfficient.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/Week8/PrefixSum_WorkEfficient.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": "a3f4f4bd-8178-4b00-eecd-3ee991af318d"
      },
      "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-7ox39abj\n",
            "  Running command git clone --filter=blob:none --quiet https://github.com/andreinechaev/nvcc4jupyter.git /tmp/pip-req-build-7ox39abj\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=94d952ccd00763de9ee9e62f29e2935209e0dc3127e66c7c2a46d2a6a7634d2b\n",
            "  Stored in directory: /tmp/pip-ephem-wheel-cache-hbamko7e/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",
        "#define SECTION_SIZE 4 \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",
        "//CPU version\n",
        "void sequential_scan(int* x, int* y, int N) \n",
        "{\n",
        "  y[0]=x[0]; \n",
        "  for (int i=1; i < N; i++)\n",
        "  {\n",
        "    y[i]= y [i-1] + x[i]; \n",
        "  }\n",
        "}\n",
        "\n",
        "//Inefficient exclusive scan kernel\n",
        "__global__ void work_inefficient_exc_scan_kernel(int *X, int *Y, int N) {\n",
        "  __shared__ int XY[SECTION_SIZE]; \n",
        "  int i = blockIdx.x*blockDim.x + threadIdx.x; \n",
        "  if (i < N && threadIdx.x != 0) \n",
        "  {\n",
        "    XY[ threadIdx.x ] = X[i-1];\n",
        "  } \n",
        "  else \n",
        "  {\n",
        "    XY[ threadIdx.x ] = 0;\n",
        "  }\n",
        "\n",
        "  // the code below performs iterative scan on XY \n",
        "  for (unsigned int stride=1; stride<=threadIdx.x; stride*= 2) \n",
        "  { \n",
        "    __syncthreads(); \n",
        "    XY[threadIdx.x] += XY[threadIdx.x-stride]; \n",
        "  }\n",
        "  \n",
        "  Y[i] = XY[threadIdx.x];\n",
        "}\n",
        "\n",
        "//Efficient exclusive scan kernel\n",
        "__global__ void work_efficient_exc_scan_kernel(int *X, int *Y, int N) {\n",
        "  __shared__ int XY[SECTION_SIZE]; \n",
        "  int i = blockIdx.x*blockDim.x + threadIdx.x; \n",
        "  if (i < N && threadIdx.x != 0) \n",
        "  {\n",
        "    XY[ threadIdx.x ] = X[i-1];\n",
        "  } \n",
        "  else \n",
        "  {\n",
        "    XY[ threadIdx.x ] = 0;\n",
        "  }\n",
        "\n",
        "  // the code below performs iterative scan on XY \n",
        "  for (unsigned int stride=1; stride< blockDim.x; stride*= 2) \n",
        "  { \n",
        "    __syncthreads(); \n",
        "    int index = (threadIdx.x + 1)*2*stride - 1;\n",
        "    if(index < blockDim.x )\n",
        "      XY[index] += XY[index-stride]; \n",
        "  }\n",
        "\n",
        "  for (int stride = SECTION_SIZE/2; stride>0; stride/=2)\n",
        "  {\n",
        "    __syncthreads();\n",
        "    int index = (threadIdx.x + 1)*2*stride - 1;\n",
        "    if(index + stride < 2*blockDim.x)\n",
        "      XY[index + stride] += XY[index]; \n",
        "  }\n",
        "   \n",
        "  __syncthreads();\n",
        "  Y[i] = XY[threadIdx.x];\n",
        "}\n",
        "\n",
        "//Inefficient inclusive scan kernel\n",
        "__global__ void work_inefficient_inc_scan_kernel(int *X, int *Y, int N) {\n",
        "  __shared__ int XY[SECTION_SIZE]; \n",
        "  int i = blockIdx.x*blockDim.x + threadIdx.x; \n",
        "  if (i < N) { \n",
        "    XY[threadIdx.x] = X[i]; \n",
        "  }\n",
        "\n",
        "  // the code below performs iterative scan on XY \n",
        "  for (unsigned int stride=1; stride<=threadIdx.x; stride*= 2) \n",
        "  { \n",
        "    __syncthreads(); \n",
        "    XY[threadIdx.x] += XY[threadIdx.x-stride]; \n",
        "  }\n",
        "  \n",
        "  Y[i] = XY[threadIdx.x];\n",
        "}\n",
        "   \n",
        "//Efficient inclusive scan kernel\n",
        "__global__ void work_efficient_inc_scan_kernel(int *X, int *Y, int N) {\n",
        "  __shared__ int XY[SECTION_SIZE]; \n",
        "  int i = blockIdx.x*blockDim.x + threadIdx.x; \n",
        "  if (i < N) { \n",
        "    XY[threadIdx.x] = X[i]; \n",
        "  }\n",
        "\n",
        "  // the code below performs iterative scan on XY \n",
        "  for (unsigned int stride=1; stride< blockDim.x; stride*= 2) \n",
        "  { \n",
        "    __syncthreads(); \n",
        "    int index = (threadIdx.x + 1)*2*stride - 1;\n",
        "    if(index < blockDim.x)\n",
        "      XY[index] += XY[index-stride]; \n",
        "  }\n",
        " \n",
        "  /*for (int stride = SECTION_SIZE/4; stride>0; stride/=2)\n",
        "  {\n",
        "    __syncthreads();\n",
        "    int index = (threadIdx.x + 1) *2*stride - 1;\n",
        "    if(index + stride < blockDim.x)\n",
        "      XY[index + stride] += XY[index]; \n",
        "  }*/\n",
        "  for (int stride = SECTION_SIZE/2; stride>0; stride/=2)\n",
        "  {\n",
        "    __syncthreads();\n",
        "    int index = (threadIdx.x + 1)*2*stride - 1;\n",
        "    if(index + stride < 2*blockDim.x)\n",
        "      XY[index + stride] += XY[index]; \n",
        "  }\n",
        "  \n",
        "  __syncthreads();\n",
        "  Y[i] = XY[threadIdx.x];\n",
        "}\n",
        "\n",
        "int main() \n",
        "{ \n",
        "\t\tint   *X, *Y; \n",
        "    const int N = 32; \n",
        "    const int threadsPerBlock = SECTION_SIZE;\n",
        "    const int blocksPerGrid =  (N / threadsPerBlock);\n",
        "\n",
        "    // Allocate Unified Memory -- accessible from CPU or GPU\n",
        "    checkCudaErr(cudaMallocManaged(&X, N*sizeof(int)), \"cudaMallocManaged1\");\n",
        "    checkCudaErr(cudaMallocManaged(&Y, N*sizeof(int)), \"cudaMallocManaged2\"); \n",
        "     \n",
        "    // fill in the memory with data\n",
        "    for (int i=0; i<N; i++) \n",
        "    {\n",
        "        X[i] = i+1;\n",
        "        Y[i] = 0; \n",
        "    }  \n",
        "  \n",
        "    // Prefetch the data to the GPU\n",
        "    int device = -1;\n",
        "    cudaGetDevice(&device);\n",
        "    cudaMemPrefetchAsync(X, N*sizeof(int), device, NULL);\n",
        "    cudaMemPrefetchAsync(Y, N*sizeof(int), device, NULL);\n",
        "    \n",
        "    cudaEvent_t start, stop;\n",
        "    float cpu_elapsed_time_ms=0,  \n",
        "          gpu1_elapsed_time_ms=0, \n",
        "          gpu2_elapsed_time_ms=0,\n",
        "          gpu3_elapsed_time_ms=0,\n",
        "          gpu4_elapsed_time_ms=0;\n",
        " \n",
        "    cudaEventCreate(&start);\n",
        "    cudaEventCreate(&stop);\n",
        " \n",
        "    ////lets time the CPU code\n",
        "    cudaEventRecord(start, 0);\n",
        "      sequential_scan(X, Y, N);\n",
        "    cudaEventRecord(stop, 0);\n",
        "    cudaEventSynchronize(stop);\n",
        "    cudaEventElapsedTime(&cpu_elapsed_time_ms, start, stop);\n",
        "\n",
        "    if(N<10) {\n",
        "      puts(\"Results on CPU: \");\n",
        "      puts(\"X[i]\\t| Y[i]\");\n",
        "      puts(\"--------+-------\");\n",
        "\n",
        "      for(int i=0; i<N; ++i)\n",
        "        printf(\"%3d\\t| %3d\\n\", X[i], Y[i]);\n",
        "    }\n",
        " \n",
        "    //reset Y for GPU\n",
        "    for (int i=0; i<N; i++) \n",
        "    {         \n",
        "        Y[i] = 0; \n",
        "    }  \n",
        "\n",
        "    //lets time the inefficient inclusive scan GPU code  \n",
        "    cudaEventRecord(start, 0); \n",
        "      work_inefficient_inc_scan_kernel<<<1,N>>>(X, Y, N);    \n",
        "    cudaEventRecord(stop, 0);\n",
        "    cudaEventSynchronize(stop); \n",
        "    cudaEventElapsedTime(&gpu1_elapsed_time_ms, start, stop);\n",
        "\n",
        "    if(N<10) {\n",
        "      puts(\"Results on GPU (inclusive inefficient): \");\n",
        "      puts(\"X[i]\\t| Y[i]\");\n",
        "      puts(\"--------+-------\");\n",
        "\n",
        "      for(int i=0; i<N; ++i)\n",
        "        printf(\"%3d\\t| %3d\\n\", X[i], Y[i]);\n",
        "    }\n",
        " \n",
        "    //reset Y for GPU\n",
        "    for (int i=0; i<N; i++) \n",
        "    {         \n",
        "        Y[i] = 0; \n",
        "    }  \n",
        " \n",
        "    //lets time the inefficient exclusive scan GPU code  \n",
        "    cudaEventRecord(start, 0);\n",
        "      work_inefficient_exc_scan_kernel<<<1,N>>>(X, Y, N);    \n",
        "    cudaEventRecord(stop, 0);\n",
        "    cudaEventSynchronize(stop); \n",
        "    cudaEventElapsedTime(&gpu2_elapsed_time_ms, start, stop);\n",
        " \n",
        "    if(N<10) {\n",
        "      puts(\"----------------\");\n",
        "      puts(\"Results on GPU (exclusive inefficient): \");\n",
        "      puts(\"X[i]\\t| Y[i]\");\n",
        "      puts(\"--------+-------\");\n",
        "\n",
        "      for(int i=0; i<N; ++i)\n",
        "        printf(\"%3d\\t| %3d\\n\", X[i], Y[i]);\n",
        "    }\n",
        "\n",
        "    //reset Y for GPU\n",
        "    for (int i=0; i<N; i++) \n",
        "    {         \n",
        "        Y[i] = 0; \n",
        "    }  \n",
        " \n",
        "    //lets time the efficient inclusive scan GPU code  \n",
        "    cudaEventRecord(start, 0);\n",
        "      work_inefficient_inc_scan_kernel<<<1,N>>>(X, Y, N);    \n",
        "    cudaEventRecord(stop, 0);\n",
        "    cudaEventSynchronize(stop); \n",
        "    cudaEventElapsedTime(&gpu3_elapsed_time_ms, start, stop);\n",
        "\n",
        "    if(N<10) {\n",
        "      puts(\"----------------\");\n",
        "      puts(\"Results on GPU (inclusive efficient): \");\n",
        "      puts(\"X[i]\\t| Y[i]\");\n",
        "      puts(\"--------+-------\");\n",
        "\n",
        "      for(int i=0; i<N; ++i)\n",
        "        printf(\"%3d\\t| %3d\\n\", X[i], Y[i]);\n",
        "    }\n",
        "\n",
        "    //reset Y for GPU\n",
        "    for (int i=0; i<N; i++) \n",
        "    {         \n",
        "        Y[i] = 0; \n",
        "    }  \n",
        "\n",
        "    //lets time the efficient exclusive scan GPU code  \n",
        "    cudaEventRecord(start, 0); \n",
        "      work_efficient_exc_scan_kernel<<<1,N>>>(X, Y, N);    \n",
        "    cudaEventRecord(stop, 0);\n",
        "    cudaEventSynchronize(stop); \n",
        "    cudaEventElapsedTime(&gpu4_elapsed_time_ms, start, stop);\n",
        "\n",
        "    if(N<10) {\n",
        "      puts(\"----------------\");\n",
        "      puts(\"Results on GPU (exclusive efficient): \");\n",
        "      puts(\"X[i]\\t| Y[i]\");\n",
        "      puts(\"--------+-------\");\n",
        "\n",
        "      for(int i=0; i<N; ++i)\n",
        "        printf(\"%3d\\t| %3d\\n\", X[i], Y[i]);\n",
        "    } \n",
        "\n",
        "    //output the result\n",
        "    puts(\"Performance Results:\");\n",
        "    printf(\"CPU Time: %3.3f msecs\\n\", cpu_elapsed_time_ms);\n",
        "    printf(\"Inefficient Scans: GPU Time (inclusive scan): %3.3f, GPU (exclusive scan): %3.3f\\n\",gpu1_elapsed_time_ms, gpu2_elapsed_time_ms);\n",
        "    printf(\"Efficient Scans: GPU Time (inclusive scan): %3.3f, GPU (exclusive scan): %3.3f\\n\",gpu3_elapsed_time_ms, gpu4_elapsed_time_ms);\n",
        "   \n",
        "\n",
        "    // free memory on the gpu side\n",
        "    checkCudaErr( cudaFree( X ) , \"cudaFree1\");\n",
        "    checkCudaErr( cudaFree( Y ) , \"cudaFree2\"); \n",
        "\t\tcheckCudaErr( cudaDeviceReset(), \"cudaDeviceReset\");\n",
        "\n",
        "\t\treturn 0;\n",
        "}"
      ],
      "metadata": {
        "id": "GYZ627HO-XKd",
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "outputId": "95c947f3-7eef-489f-f063-1683a038b7de"
      },
      "execution_count": 2,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Performance Results:\n",
            "CPU Time: 0.003 msecs\n",
            "Inefficient Scans: GPU Time (inclusive scan): 0.168, GPU (exclusive scan): 0.092\n",
            "Efficient Scans: GPU Time (inclusive scan): 0.071, GPU (exclusive scan): 0.080\n",
            "\n"
          ]
        }
      ]
    }
  ]
}