Week9/Histogram_Strategy_4.ipynb

{
  "nbformat": 4,
  "nbformat_minor": 0,
  "metadata": {
    "colab": {
      "name": "Histogram_Strategy_4.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/Week9/Histogram_Strategy_4.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": "82770334-a443-46e1-eb02-c2bf5bcef2ef"
      },
      "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-ytzmxlk1\n",
            "  Running command git clone --filter=blob:none --quiet https://github.com/andreinechaev/nvcc4jupyter.git /tmp/pip-req-build-ytzmxlk1\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=a4064de64ef3535123ad34d6d587f5e2e9cfe3ac98000b89ff5778b00a2514e9\n",
            "  Stored in directory: /tmp/pip-ephem-wheel-cache-n6e3_w1h/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",
        "#include <cstdlib> //rand() function\n",
        "\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",
        "void histogram_CPU(int *data, const int N, int* histogram, const int Nbins)\n",
        "{\n",
        "\tfor(int i=0; i<N; ++i)\n",
        "\t{\n",
        "\t\thistogram[data[i]]++;\t\t\n",
        "\t}\n",
        "}\n",
        "\n",
        "__global__ void histogram_GPU_1( int *data, const int N, int *histo) \n",
        "{\n",
        "\tint i = threadIdx.x + blockIdx.x * blockDim.x;    \n",
        "\tint section_size = (N-1)/(blockDim.x *gridDim.x)+1; \n",
        "\tint start = i*section_size; //All threads handle blockDim.x * gridDim.x consecutive elements \n",
        "\tfor (int k = 0; k < section_size; k++) \n",
        "\t{\n",
        "\t\tif (start+k < N) \n",
        "\t\t{ \n",
        "\t\t\tatomicAdd(&(histo[data[start+k]]), 1); \n",
        "\t\t} \n",
        "\t}  \n",
        "}\n",
        "\n",
        "__global__ void histogram_GPU_2( int *data, const int N, int *histo) \n",
        "{\n",
        "\tint tid = threadIdx.x + blockIdx.x * blockDim.x;\n",
        "  \n",
        "\t//All threads handle blockDim.x * gridDim.x consecutive elements in each \n",
        "\t//iteration\n",
        "\n",
        "\tfor(int i=tid ; i < N ; i += blockDim.x*gridDim.x )\n",
        "  {      \n",
        "\t\t atomicAdd(&(histo[data[i]]), 1); \t\t \n",
        "\t}  \n",
        "}\n",
        "\n",
        "__global__ void histogram_GPU_3( int *data, const int N, int *histo, const int num_bins) \n",
        "{\n",
        "\tint tid = threadIdx.x + blockIdx.x * blockDim.x;\n",
        "  \n",
        "\t//Privatized bins\n",
        "\textern __shared__ int histo_s[];\n",
        "\t\n",
        "\tfor(int binIdx = threadIdx.x ; binIdx < num_bins ; binIdx += blockDim.x )\n",
        "     histo_s[ binIdx ] = 0u;\n",
        "\t\n",
        "\t__syncthreads();\n",
        "\n",
        "\t//Histogram\n",
        "\tfor(int i = tid ; i < N ; i += blockDim.x*gridDim.x )\n",
        "\t\tatomicAdd(&(histo_s[data[i]]), 1);\n",
        "\n",
        "  __syncthreads();\n",
        "\n",
        "\t//Commit to global memory (merge step)\n",
        "  for(int binIdx = threadIdx.x ; binIdx < num_bins ; binIdx += blockDim.x )\n",
        "\t\tatomicAdd(&(histo[binIdx]), histo_s[binIdx]);\t \n",
        "}\n",
        "\n",
        "__global__ void histogram_GPU_4( int *data, const int N, int *histo, const int num_bins) \n",
        "{\n",
        "\tint tid = threadIdx.x + blockIdx.x * blockDim.x;\n",
        "  \n",
        "\t//Privatized bins\n",
        "\textern __shared__ int histo_s[];\n",
        "\t\n",
        "\tfor(int binIdx = threadIdx.x ; binIdx < num_bins ; binIdx += blockDim.x )\n",
        "     histo_s[ binIdx ] = 0;\n",
        "\t\n",
        "\t__syncthreads();\n",
        "  \n",
        "  int prev_index = -1;\n",
        "\tint accumulator = 0;\n",
        "  \t\n",
        "\t//Histogram\n",
        "\tfor(int i = tid ; i < N ; i += blockDim.x*gridDim.x )\n",
        "  {\n",
        "\t\tint curr_index = data[i];\n",
        "\t\tif ( curr_index != prev_index) \n",
        "\t  {\n",
        "\t\t\t if(accumulator >0)\n",
        "\t\t\t \t  atomicAdd(&(histo_s[curr_index]), accumulator);\t\t\t \n",
        "\t\t\t accumulator = 1;\n",
        "\t\t\t prev_index = curr_index;\t\t\t\t\t\t\n",
        "\t\t} \n",
        "\t\telse \n",
        "\t\t{\n",
        "\t\t\taccumulator++;\n",
        "\t\t}\n",
        "\t}\t \n",
        "\t\n",
        "  __syncthreads();\n",
        "\n",
        "\t//Commit to global memory (merge step)\n",
        "  for(int binIdx = threadIdx.x ; binIdx < num_bins ; binIdx += blockDim.x  )\n",
        "\t\tatomicAdd(&(histo[binIdx]), histo_s[binIdx]);\t \n",
        "}\n",
        "\n",
        "int main() \n",
        "{ \n",
        "\tint *data;\n",
        "\tint *histogram;\n",
        "\tint *histogramGPU; \n",
        "\tconst int Nbins = 256;\n",
        "\tconst int N = 1024 * 1024;\n",
        "\tconst int threadsPerBlock = 256;\n",
        "\n",
        "\t#define imin(a,b) (a<b?a:b)\n",
        "\n",
        "\tconst int blocksPerGrid =  imin( 32, (N+threadsPerBlock-1) / threadsPerBlock );\n",
        " \n",
        "\t// Allocate Unified Memory -- accessible from CPU or GPU\n",
        "\tcheckCudaErr(cudaMallocManaged(&data, N*sizeof(int)), \"cudaMallocManaged1 data\");\n",
        "\tcheckCudaErr(cudaMallocManaged(&histogram, Nbins*sizeof(int)), \"cudaMallocManaged2 histogram\"); \n",
        "\tcheckCudaErr(cudaMallocManaged(&histogramGPU, Nbins*sizeof(int)), \"cudaMallocManaged2 histogramGPU\"); \n",
        "\n",
        "\t// fill in the data with random values between 0-255\n",
        "\tfor (int i=0; i<N; i++) {\n",
        "\t\tdata[i] = (rand() % Nbins); \n",
        "\t} \n",
        "\n",
        "\t// some events to count the execution time\n",
        "  cudaEvent_t start, stop;\n",
        "\tfloat cpu_elapsed_time_ms, gpu_1_elapsed_time_ms, gpu_2_elapsed_time_ms,\n",
        "\t\t\t\tgpu_3_elapsed_time_ms, gpu_4_elapsed_time_ms; \n",
        "\n",
        "  cudaEventCreate(&start);\n",
        "  cudaEventCreate(&stop);\n",
        "\t\n",
        "\t//start to count execution time of CPU version\n",
        "\tcudaEventRecord(start, 0);\n",
        " \n",
        "  //calculate histogram on the CPU\n",
        "\thistogram_CPU(data, N, histogram, Nbins);\n",
        "\n",
        "  // time counting terminate\n",
        "  cudaEventRecord(stop, 0);\n",
        "\tcudaEventSynchronize(stop);\n",
        "\n",
        "  //compute time elapsed on CPU \n",
        "  cudaEventElapsedTime(&cpu_elapsed_time_ms, start, stop);\n",
        "  printf(\"Histogram(CPU) - Time: %f ms.\\n\", cpu_elapsed_time_ms);\n",
        "\n",
        " \t//start to count execution time of GPU version\n",
        "\tcudaEventRecord(start, 0);\n",
        "\n",
        "\t//launch kernel\n",
        "\thistogram_GPU_1<<<blocksPerGrid,threadsPerBlock>>>( data, N, histogramGPU );\n",
        "\n",
        "\t// time counting terminate\n",
        "  cudaEventRecord(stop, 0);\n",
        "\tcudaEventSynchronize(stop);\n",
        "\n",
        "  // compute time elapse on GPU computing\n",
        "  cudaEventElapsedTime(&gpu_1_elapsed_time_ms, start, stop);\n",
        "  printf(\"Histogram_1(GPU) - Time: %f ms.\\n\", gpu_1_elapsed_time_ms);\n",
        "\t\n",
        "\t//clear old gpu histogram for histogram calc uisng strategy 2\n",
        "\tfor (int i=0; i<Nbins; i++) {\n",
        "\t\thistogramGPU[i] =0; \n",
        "\t} \n",
        "\n",
        "\t//start to count execution time of GPU version\n",
        "\tcudaEventRecord(start, 0);\n",
        "\n",
        "\t//launch kernel\n",
        "\thistogram_GPU_2<<<blocksPerGrid,threadsPerBlock>>>( data, N, histogramGPU );\n",
        "\n",
        "\t// time counting terminate\n",
        "  cudaEventRecord(stop, 0);\n",
        "\tcudaEventSynchronize(stop);\n",
        "\n",
        "  // compute time elapse on GPU computing\n",
        "  cudaEventElapsedTime(&gpu_2_elapsed_time_ms, start, stop);\n",
        "  printf(\"Histogram_2(GPU) - Time: %f ms.\\n\", gpu_2_elapsed_time_ms);\n",
        "\n",
        "  //clear old gpu histogram for histogram calc uisng strategy 3\n",
        "\tfor (int i=0; i<Nbins; i++) {\n",
        "\t\thistogramGPU[i] =0; \n",
        "\t} \n",
        "\n",
        "\t//start to count execution time of GPU version\n",
        "\tcudaEventRecord(start, 0);\n",
        "\n",
        "\t//launch kernel\n",
        "\thistogram_GPU_3<<<blocksPerGrid,threadsPerBlock, Nbins*sizeof(int)>>>( data, N, histogramGPU, Nbins );\n",
        "\n",
        "\t// time counting terminate\n",
        "  cudaEventRecord(stop, 0);\n",
        "\tcudaEventSynchronize(stop);\n",
        "\n",
        "  // compute time elapse on GPU computing\n",
        "  cudaEventElapsedTime(&gpu_3_elapsed_time_ms, start, stop);\n",
        "  printf(\"Histogram_3(GPU) - Time: %f ms.\\n\", gpu_3_elapsed_time_ms);\n",
        "\n",
        "  //clear old gpu histogram for histogram calc uisng strategy 4\n",
        "\tfor (int i=0; i<Nbins; i++) {\n",
        "\t\thistogramGPU[i] =0; \n",
        "\t} \n",
        "\n",
        "\t//start to count execution time of GPU version\n",
        "\tcudaEventRecord(start, 0);\n",
        "\n",
        "\t//launch kernel\n",
        "\thistogram_GPU_4<<<blocksPerGrid,threadsPerBlock, Nbins*sizeof(int)>>>( data, N, histogramGPU, Nbins );\n",
        "  \t   \n",
        "\t// time counting terminate\n",
        "  cudaEventRecord(stop, 0);\n",
        "\tcudaEventSynchronize(stop);\n",
        "\n",
        "  // compute time elapse on GPU computing\n",
        "  cudaEventElapsedTime(&gpu_4_elapsed_time_ms, start, stop);\n",
        "  printf(\"Histogram_4(GPU) - Time: %f ms.\\n\", gpu_4_elapsed_time_ms);\n",
        "\n",
        "\tint allOK = 1;\n",
        " \n",
        "  printf(\"+------------------+-------------------+\\n\");\n",
        "  printf(\"| CPU Histogram[i] |  GPU Histogram[i] |\\n\");\n",
        "\tprintf(\"+------------------+-------------------+\\n\");\n",
        "\t\n",
        "  //compare results on CPU and GPU\n",
        "\tfor(int i=0 ; i<Nbins; ++i)\n",
        " \t{\n",
        "\t\t printf(\"|      %5d       |      %5d        |\\n\", histogram[i], histogramGPU[i]);\n",
        "      \n",
        "\t\t if(histogram[i]!=histogramGPU[i])\n",
        "\t\t {\n",
        "\t\t    printf(\"|\\t Histogram mismatch!!!         |\\n\");\n",
        "\t\t\t\tallOK=0;\n",
        "\t\t\t\tbreak;\n",
        "\t\t }\n",
        "\t}\n",
        "  printf(\"+------------------+-------------------+\\n\");\n",
        "\t\n",
        "\tif(allOK==1)\n",
        "  \tprintf(\"Results correct on both CPU and GPU\");\n",
        "\n",
        "\t// free memory on the gpu side\n",
        "\tcheckCudaErr( cudaFree( data ) , \"cudaFree1\");\n",
        "\tcheckCudaErr( cudaFree( histogram ) , \"cudaFree2\"); \n",
        "\tcheckCudaErr( cudaFree( histogramGPU ) , \"cudaFree3\"); \n",
        "\tcheckCudaErr( cudaDeviceReset(), \"cudaDeviceReset\");\n",
        "\n",
        "\treturn 0;\n",
        "}"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "GYZ627HO-XKd",
        "outputId": "8d6163c0-a7ba-47d9-d412-3630ec0d2aef"
      },
      "execution_count": 2,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Histogram(CPU) - Time: 2.768736 ms.\n",
            "Histogram_1(GPU) - Time: 2.186752 ms.\n",
            "Histogram_2(GPU) - Time: 0.790688 ms.\n",
            "Histogram_3(GPU) - Time: 0.149952 ms.\n",
            "Histogram_4(GPU) - Time: 0.172224 ms.\n",
            "+------------------+-------------------+\n",
            "| CPU Histogram[i] |  GPU Histogram[i] |\n",
            "+------------------+-------------------+\n",
            "|       4083       |       4049        |\n",
            "|\t Histogram mismatch!!!         |\n",
            "+------------------+-------------------+\n",
            "\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [],
      "metadata": {
        "id": "x7sxTPKWeSpW"
      },
      "execution_count": 2,
      "outputs": []
    }
  ]
}