Cuda_Image_Filter.ipynb

{
  "cells": [
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "PqvyY4x0V333"
      },
      "source": [
        "# Download Images and Libraries"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "yCOxT2eEyQU9",
        "outputId": "7319d8d8-a505-4d74-cced-a76ef2a7cf78"
      },
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "--2023-01-08 12:46:47--  https://raw.githubusercontent.com/nothings/stb/master/stb_image.h\n",
            "Resolving raw.githubusercontent.com (raw.githubusercontent.com)... 185.199.108.133, 185.199.109.133, 185.199.110.133, ...\n",
            "Connecting to raw.githubusercontent.com (raw.githubusercontent.com)|185.199.108.133|:443... connected.\n",
            "HTTP request sent, awaiting response... 200 OK\n",
            "Length: 279339 (273K) [text/plain]\n",
            "Saving to: ‘stb_image.h’\n",
            "\n",
            "stb_image.h         100%[===================>] 272.79K  --.-KB/s    in 0.004s  \n",
            "\n",
            "2023-01-08 12:46:47 (73.0 MB/s) - ‘stb_image.h’ saved [279339/279339]\n",
            "\n",
            "--2023-01-08 12:46:47--  https://raw.githubusercontent.com/nothings/stb/master/stb_image_write.h\n",
            "Resolving raw.githubusercontent.com (raw.githubusercontent.com)... 185.199.108.133, 185.199.109.133, 185.199.110.133, ...\n",
            "Connecting to raw.githubusercontent.com (raw.githubusercontent.com)|185.199.108.133|:443... connected.\n",
            "HTTP request sent, awaiting response... 200 OK\n",
            "Length: 71221 (70K) [text/plain]\n",
            "Saving to: ‘stb_image_write.h’\n",
            "\n",
            "stb_image_write.h   100%[===================>]  69.55K  --.-KB/s    in 0.001s  \n",
            "\n",
            "2023-01-08 12:46:47 (83.2 MB/s) - ‘stb_image_write.h’ saved [71221/71221]\n",
            "\n",
            "--2023-01-08 12:46:47--  https://raw.githubusercontent.com/frischzenger/yolo-windows/master/3rdparty/include/pthread.h\n",
            "Resolving raw.githubusercontent.com (raw.githubusercontent.com)... 185.199.108.133, 185.199.109.133, 185.199.110.133, ...\n",
            "Connecting to raw.githubusercontent.com (raw.githubusercontent.com)|185.199.108.133|:443... connected.\n",
            "HTTP request sent, awaiting response... 200 OK\n",
            "Length: 42499 (42K) [text/plain]\n",
            "Saving to: ‘pthread.h’\n",
            "\n",
            "pthread.h           100%[===================>]  41.50K  --.-KB/s    in 0s      \n",
            "\n",
            "2023-01-08 12:46:48 (82.2 MB/s) - ‘pthread.h’ saved [42499/42499]\n",
            "\n",
            "--2023-01-08 12:46:48--  https://raw.githubusercontent.com/NathanBruggeman/Images_Cuda/9cb838f170745b866d13da5476affb7f51425e33/img0.png\n",
            "Resolving raw.githubusercontent.com (raw.githubusercontent.com)... 185.199.108.133, 185.199.109.133, 185.199.110.133, ...\n",
            "Connecting to raw.githubusercontent.com (raw.githubusercontent.com)|185.199.108.133|:443... connected.\n",
            "HTTP request sent, awaiting response... 200 OK\n",
            "Length: 4354 (4.3K) [image/png]\n",
            "Saving to: ‘/content/Images/img0.png’\n",
            "\n",
            "img0.png            100%[===================>]   4.25K  --.-KB/s    in 0s      \n",
            "\n",
            "2023-01-08 12:46:48 (50.5 MB/s) - ‘/content/Images/img0.png’ saved [4354/4354]\n",
            "\n",
            "--2023-01-08 12:46:48--  https://raw.githubusercontent.com/NathanBruggeman/Images_Cuda/9cb838f170745b866d13da5476affb7f51425e33/img1.png\n",
            "Resolving raw.githubusercontent.com (raw.githubusercontent.com)... 185.199.110.133, 185.199.108.133, 185.199.109.133, ...\n",
            "Connecting to raw.githubusercontent.com (raw.githubusercontent.com)|185.199.110.133|:443... connected.\n",
            "HTTP request sent, awaiting response... 200 OK\n",
            "Length: 1947165 (1.9M) [image/png]\n",
            "Saving to: ‘/content/Images/img1.png’\n",
            "\n",
            "img1.png            100%[===================>]   1.86M  --.-KB/s    in 0.008s  \n",
            "\n",
            "2023-01-08 12:46:49 (221 MB/s) - ‘/content/Images/img1.png’ saved [1947165/1947165]\n",
            "\n",
            "--2023-01-08 12:46:49--  https://raw.githubusercontent.com/NathanBruggeman/Images_Cuda/main/img2.png\n",
            "Resolving raw.githubusercontent.com (raw.githubusercontent.com)... 185.199.108.133, 185.199.109.133, 185.199.110.133, ...\n",
            "Connecting to raw.githubusercontent.com (raw.githubusercontent.com)|185.199.108.133|:443... connected.\n",
            "HTTP request sent, awaiting response... 200 OK\n",
            "Length: 1622809 (1.5M) [image/png]\n",
            "Saving to: ‘/content/Images/img2.png’\n",
            "\n",
            "img2.png            100%[===================>]   1.55M  --.-KB/s    in 0.007s  \n",
            "\n",
            "2023-01-08 12:46:49 (213 MB/s) - ‘/content/Images/img2.png’ saved [1622809/1622809]\n",
            "\n",
            "--2023-01-08 12:46:49--  https://raw.githubusercontent.com/NathanBruggeman/Images_Cuda/main/img3.png\n",
            "Resolving raw.githubusercontent.com (raw.githubusercontent.com)... 185.199.110.133, 185.199.109.133, 185.199.111.133, ...\n",
            "Connecting to raw.githubusercontent.com (raw.githubusercontent.com)|185.199.110.133|:443... connected.\n",
            "HTTP request sent, awaiting response... 200 OK\n",
            "Length: 2465326 (2.4M) [image/png]\n",
            "Saving to: ‘/content/Images/img3.png’\n",
            "\n",
            "img3.png            100%[===================>]   2.35M  --.-KB/s    in 0.01s   \n",
            "\n",
            "2023-01-08 12:46:50 (239 MB/s) - ‘/content/Images/img3.png’ saved [2465326/2465326]\n",
            "\n",
            "--2023-01-08 12:46:50--  https://raw.githubusercontent.com/NathanBruggeman/Images_Cuda/main/img4.png\n",
            "Resolving raw.githubusercontent.com (raw.githubusercontent.com)... 185.199.110.133, 185.199.108.133, 185.199.109.133, ...\n",
            "Connecting to raw.githubusercontent.com (raw.githubusercontent.com)|185.199.110.133|:443... connected.\n",
            "HTTP request sent, awaiting response... 200 OK\n",
            "Length: 1817087 (1.7M) [image/png]\n",
            "Saving to: ‘/content/Images/img4.png’\n",
            "\n",
            "img4.png            100%[===================>]   1.73M  --.-KB/s    in 0.009s  \n",
            "\n",
            "2023-01-08 12:46:50 (199 MB/s) - ‘/content/Images/img4.png’ saved [1817087/1817087]\n",
            "\n",
            "--2023-01-08 12:46:50--  https://raw.githubusercontent.com/NathanBruggeman/Images_Cuda/main/img5.png\n",
            "Resolving raw.githubusercontent.com (raw.githubusercontent.com)... 185.199.110.133, 185.199.108.133, 185.199.109.133, ...\n",
            "Connecting to raw.githubusercontent.com (raw.githubusercontent.com)|185.199.110.133|:443... connected.\n",
            "HTTP request sent, awaiting response... 200 OK\n",
            "Length: 897772 (877K) [image/png]\n",
            "Saving to: ‘/content/Images/img5.png’\n",
            "\n",
            "img5.png            100%[===================>] 876.73K  --.-KB/s    in 0.005s  \n",
            "\n",
            "2023-01-08 12:46:50 (171 MB/s) - ‘/content/Images/img5.png’ saved [897772/897772]\n",
            "\n",
            "--2023-01-08 12:46:50--  https://raw.githubusercontent.com/NathanBruggeman/Images_Cuda/main/img6.png\n",
            "Resolving raw.githubusercontent.com (raw.githubusercontent.com)... 185.199.108.133, 185.199.109.133, 185.199.110.133, ...\n",
            "Connecting to raw.githubusercontent.com (raw.githubusercontent.com)|185.199.108.133|:443... connected.\n",
            "HTTP request sent, awaiting response... 200 OK\n",
            "Length: 6649783 (6.3M) [application/octet-stream]\n",
            "Saving to: ‘/content/Images/img6.png’\n",
            "\n",
            "img6.png            100%[===================>]   6.34M  --.-KB/s    in 0.05s   \n",
            "\n",
            "2023-01-08 12:46:51 (134 MB/s) - ‘/content/Images/img6.png’ saved [6649783/6649783]\n",
            "\n",
            "--2023-01-08 12:46:51--  https://raw.githubusercontent.com/NathanBruggeman/Images_Cuda/main/img7.png\n",
            "Resolving raw.githubusercontent.com (raw.githubusercontent.com)... 185.199.108.133, 185.199.109.133, 185.199.110.133, ...\n",
            "Connecting to raw.githubusercontent.com (raw.githubusercontent.com)|185.199.108.133|:443... connected.\n",
            "HTTP request sent, awaiting response... 200 OK\n",
            "Length: 70000223 (67M) [application/octet-stream]\n",
            "Saving to: ‘/content/Images/img7.png’\n",
            "\n",
            "img7.png            100%[===================>]  66.76M   295MB/s    in 0.2s    \n",
            "\n",
            "2023-01-08 12:46:56 (295 MB/s) - ‘/content/Images/img7.png’ saved [70000223/70000223]\n",
            "\n",
            "--2023-01-08 12:46:56--  https://raw.githubusercontent.com/NathanBruggeman/Images_Cuda/main/img8.png\n",
            "Resolving raw.githubusercontent.com (raw.githubusercontent.com)... 185.199.108.133, 185.199.109.133, 185.199.110.133, ...\n",
            "Connecting to raw.githubusercontent.com (raw.githubusercontent.com)|185.199.108.133|:443... connected.\n",
            "HTTP request sent, awaiting response... 200 OK\n",
            "Length: 79905229 (76M) [application/octet-stream]\n",
            "Saving to: ‘/content/Images/img8.png’\n",
            "\n",
            "img8.png            100%[===================>]  76.20M   326MB/s    in 0.2s    \n",
            "\n",
            "2023-01-08 12:47:01 (326 MB/s) - ‘/content/Images/img8.png’ saved [79905229/79905229]\n",
            "\n",
            "--2023-01-08 12:47:01--  https://raw.githubusercontent.com/NathanBruggeman/Images_Cuda/main/img9.png\n",
            "Resolving raw.githubusercontent.com (raw.githubusercontent.com)... 185.199.108.133, 185.199.109.133, 185.199.110.133, ...\n",
            "Connecting to raw.githubusercontent.com (raw.githubusercontent.com)|185.199.108.133|:443... connected.\n",
            "HTTP request sent, awaiting response... 200 OK\n",
            "Length: 44156792 (42M) [application/octet-stream]\n",
            "Saving to: ‘/content/Images/img9.png’\n",
            "\n",
            "img9.png            100%[===================>]  42.11M  --.-KB/s    in 0.1s    \n",
            "\n",
            "2023-01-08 12:47:04 (432 MB/s) - ‘/content/Images/img9.png’ saved [44156792/44156792]\n",
            "\n"
          ]
        }
      ],
      "source": [
        "!mkdir Images\n",
        "!mkdir Output_Images\n",
        "!mkdir Output_Images/Convolution\n",
        "!mkdir Output_Images/Pooling\n",
        "\n",
        "!wget https://raw.githubusercontent.com/nothings/stb/master/stb_image.h\n",
        "!wget https://raw.githubusercontent.com/nothings/stb/master/stb_image_write.h\n",
        "!wget https://raw.githubusercontent.com/frischzenger/yolo-windows/master/3rdparty/include/pthread.h\n",
        "\n",
        "!wget -P /content/Images https://raw.githubusercontent.com/NathanBruggeman/Images_Cuda/9cb838f170745b866d13da5476affb7f51425e33/img0.png\n",
        "!wget -P /content/Images https://raw.githubusercontent.com/NathanBruggeman/Images_Cuda/9cb838f170745b866d13da5476affb7f51425e33/img1.png\n",
        "!wget -P /content/Images https://raw.githubusercontent.com/NathanBruggeman/Images_Cuda/main/img2.png\n",
        "!wget -P /content/Images https://raw.githubusercontent.com/NathanBruggeman/Images_Cuda/main/img3.png\n",
        "!wget -P /content/Images https://raw.githubusercontent.com/NathanBruggeman/Images_Cuda/main/img4.png\n",
        "!wget -P /content/Images https://raw.githubusercontent.com/NathanBruggeman/Images_Cuda/main/img5.png\n",
        "!wget -P /content/Images https://raw.githubusercontent.com/NathanBruggeman/Images_Cuda/main/img6.png\n",
        "!wget -P /content/Images https://raw.githubusercontent.com/NathanBruggeman/Images_Cuda/main/img7.png\n",
        "!wget -P /content/Images https://raw.githubusercontent.com/NathanBruggeman/Images_Cuda/main/img8.png\n",
        "!wget -P /content/Images https://raw.githubusercontent.com/NathanBruggeman/Images_Cuda/main/img9.png"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "otqZ4wNpWALa"
      },
      "source": [
        "# Write and compile the ImageFilter_Sequential.c\n",
        "## This program filters 10 images sequential."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "-GGK5tdGvvUn",
        "outputId": "4063b1b6-a0ba-4eea-af71-b26d2e3fed94"
      },
      "outputs": [
        {
          "name": "stdout",
          "output_type": "stream",
          "text": [
            "Overwriting ImageFilter_Sequential.c\n"
          ]
        }
      ],
      "source": [
        "%%writefile ImageFilter_Sequential.c\n",
        "#include <stdio.h>\n",
        "#include <stdlib.h>\n",
        "#define STB_IMAGE_IMPLEMENTATION\n",
        "#define STB_IMAGE_WRITE_IMPLEMENTATION\n",
        "#include \"stb_image.h\"\n",
        "#include \"stb_image_write.h\"\n",
        "#include <string.h>\n",
        "#include <time.h>\n",
        "\n",
        "typedef struct Pixel\n",
        "{\n",
        "    unsigned char r, g, b, a;\n",
        "} Pixel;\n",
        "\n",
        "void ConvertImageToGrayCpu(unsigned char *originalImage, unsigned char *imageDataGrayscale, int width, int height);\n",
        "void convolveImage(unsigned char *imageDataGrayscale, unsigned char *imageDataConvolution, int width, int height);\n",
        "void minPooling(unsigned char *originalImage, unsigned char *minPoolingImage, int width, int height);\n",
        "void maxPooling(unsigned char *originalImage, unsigned char *maxPoolingImage, int width, int height);\n",
        "\n",
        "int main()\n",
        "{\n",
        "    clock_t timer_start, timer_end;\n",
        "    timer_start = clock(); // Start the timer\n",
        "\n",
        "    for (int i = 0; i < 10; i++)\n",
        "    {\n",
        "        char inputImage[64] = \"Images/img\";\n",
        "        char outputImageConvolution[64] = \"Output_Images/Convolution/img\";\n",
        "        char outputImagePoolingMax[64] = \"Output_Images/Pooling/imgMax\";\n",
        "        char outputImagePoolingMin[64] = \"Output_Images/Pooling/imgMin\";\n",
        "        char imageNumber[3];\n",
        "\n",
        "        // Make a complete path for the input (inputImage + imageNumber + extention)\n",
        "        sprintf(imageNumber, \"%d\", i); // Get the number of the current iteration and convert it to char imageNumber\n",
        "        strcat(inputImage, imageNumber);\n",
        "        strcat(inputImage, \".png\");\n",
        "\n",
        "        ////////////////\n",
        "        // Open Image //\n",
        "        ////////////////\n",
        "        // Open the image and allocate memory for each process\n",
        "        int width, height, componentCount;\n",
        "        printf(\"Loading %s.\\r\\n\", inputImage);\n",
        "        unsigned char *originalImage = stbi_load(inputImage, &width, &height, &componentCount, 4); // Saves original image\n",
        "        if (!originalImage)\n",
        "        {\n",
        "            printf(\"Failed to open image!\\r\\n\");\n",
        "            stbi_image_free(originalImage);\n",
        "            return -1;\n",
        "        }\n",
        "\n",
        "        if (width % 32 || height % 32)\n",
        "        {\n",
        "            // NOTE: Leaked memory of \"imageData\"\n",
        "            printf(\"Width and/or Height is not dividable by 32!\\r\\n\");\n",
        "            stbi_image_free(originalImage);\n",
        "            return -1;\n",
        "        }\n",
        "\n",
        "        unsigned char *imageData = (unsigned char *)malloc(width * height * 4);           // Saves grayscale image\n",
        "        unsigned char *imageDataTest = (unsigned char *)malloc(width * height * 4);       // Saves output image\n",
        "        unsigned char *imageDataMinPooling = (unsigned char *)malloc(width * height * 4); // Saves Min pooling image\n",
        "        unsigned char *imageDataMaxPooling = (unsigned char *)malloc(width * height * 4); // Saves Max pooling image\n",
        "\n",
        "        ///////////////////////\n",
        "        // Convert Grayscale //\n",
        "        ///////////////////////\n",
        "        // Convert image to grayscale on CPU\n",
        "        printf(\"Processing image grayscale.\\r\\n\");\n",
        "        // Process image on cpu\n",
        "        ConvertImageToGrayCpu(originalImage, imageData, width, height);\n",
        "        printf(\"Done\\r\\n\");\n",
        "\n",
        "        ////////////////////\n",
        "        // Convolve image //\n",
        "        ////////////////////\n",
        "        // Make a complete path for the convolved image (outputImage convolution + imageNumber + extention)\n",
        "        strcat(outputImageConvolution, imageNumber);\n",
        "        strcat(outputImageConvolution, \".png\");\n",
        "\n",
        "        // Convolve the image on CPU\n",
        "        printf(\"Convolving image.\\r\\n\");\n",
        "        convolveImage(imageData, imageDataTest, width, height);\n",
        "        printf(\"Done\\r\\n\");\n",
        "\n",
        "        // Write convolved image to disk\n",
        "        printf(\"Writing image to disk\\r\\n\");\n",
        "        stbi_write_png(outputImageConvolution, width, height, 4, imageDataTest, 4 * width);\n",
        "        printf(\"Done\\r\\n\");\n",
        "\n",
        "        /////////////////\n",
        "        // Min Pooling //\n",
        "        /////////////////\n",
        "        // Make a complete path for max pooling\n",
        "        sprintf(imageNumber, \"%d\", i); // Get the number of the current iteration and convert it to char imageNumber\n",
        "        strcat(outputImagePoolingMin, imageNumber);\n",
        "        strcat(outputImagePoolingMin, \".png\");\n",
        "        minPooling(originalImage, imageDataMinPooling, width, height);\n",
        "        // Write image back to disk\n",
        "        stbi_write_png(outputImagePoolingMin, width / 2, height / 2, 4, imageDataMinPooling, 4 * (width / 2));\n",
        "\n",
        "        /////////////////\n",
        "        // Max Pooling //\n",
        "        /////////////////\n",
        "        sprintf(imageNumber, \"%d\", i); // Get the number of the current iteration and convert it to char imageNumber\n",
        "        strcat(outputImagePoolingMax, imageNumber);\n",
        "        strcat(outputImagePoolingMax, \".png\");\n",
        "        maxPooling(originalImage, imageDataMaxPooling, width, height);\n",
        "\n",
        "        // Write image back to disk\n",
        "        stbi_write_png(outputImagePoolingMax, width / 2, height / 2, 4, imageDataMaxPooling, 4 * (width / 2));\n",
        "\n",
        "        /////////////////\n",
        "        // Free memory //\n",
        "        /////////////////\n",
        "        stbi_image_free(originalImage);\n",
        "        free(imageData);\n",
        "        free(imageDataTest);\n",
        "        free(imageDataMinPooling);\n",
        "        free(imageDataMaxPooling);\n",
        "\n",
        "        printf(\"\\r\\nNext Image\\r\\n\");\n",
        "    }\n",
        "\n",
        "    ///////////////\n",
        "    // End timer //\n",
        "    ///////////////\n",
        "    timer_end = clock(); // end the timer\n",
        "    double time_spent = (double)(timer_end - timer_start) / CLOCKS_PER_SEC;\n",
        "    printf(\"\\nProgram time: %.3fs\\n\", time_spent);\n",
        "\n",
        "    return 0;\n",
        "}\n",
        "\n",
        "void ConvertImageToGrayCpu(unsigned char *originalImage, unsigned char *imageDataGrayscale, int width, int height)\n",
        "{\n",
        "    for (int y = 0; y < height; y++)\n",
        "    {\n",
        "        for (int x = 0; x < width; x++)\n",
        "        {\n",
        "            Pixel *ptrPixel = (Pixel *)&imageDataGrayscale[y * width * 4 + 4 * x];\n",
        "            Pixel *ptrPixelOriginal = (Pixel *)&originalImage[y * width * 4 + 4 * x];\n",
        "            unsigned char pixelValue = (unsigned char)(ptrPixelOriginal->r * 0.2126f + ptrPixelOriginal->g * 0.7152f + ptrPixelOriginal->b * 0.0722f);\n",
        "            ptrPixel->r = pixelValue;\n",
        "            ptrPixel->g = pixelValue;\n",
        "            ptrPixel->b = pixelValue;\n",
        "            ptrPixel->a = 255;\n",
        "        }\n",
        "    }\n",
        "}\n",
        "\n",
        "void convolveImage(unsigned char *imageDataGrayscale, unsigned char *imageDataConvolution, int width, int height)\n",
        "{\n",
        "    int kernel[3][3] =\n",
        "        {\n",
        "            {1, 0, -1},\n",
        "            {1, 0, -1},\n",
        "            {1, 0, -1}};\n",
        "\n",
        "    int pixels[3][3] = {0};\n",
        "    int finalPixel = 0;\n",
        "\n",
        "    for (int y = 0; y < height - 2; y++)\n",
        "    {\n",
        "        for (int x = 0; x < width - 2; x++)\n",
        "        {\n",
        "            for (int i = 0; i <= 2; i++)\n",
        "            {\n",
        "                Pixel *ptrPixel = (Pixel *)&imageDataGrayscale[(y * width * 4 + 4 * x) + i * 4];\n",
        "\n",
        "                pixels[0][i] = ptrPixel->r * kernel[0][i];\n",
        "            }\n",
        "\n",
        "            for (int i = 0; i <= 2; i++)\n",
        "            {\n",
        "                Pixel *ptrPixel = (Pixel *)&imageDataGrayscale[(y * width * 4 + 4 * x) + width * 4 + i * 4];\n",
        "\n",
        "                pixels[1][i] = ptrPixel->r * kernel[1][i];\n",
        "            }\n",
        "\n",
        "            for (int i = 0; i <= 2; i++)\n",
        "            {\n",
        "                Pixel *ptrPixel = (Pixel *)&imageDataGrayscale[(y * width * 4 + 4 * x) + (2 * width * 4) + i * 4];\n",
        "\n",
        "                pixels[2][i] = ptrPixel->r * kernel[2][i];\n",
        "            }\n",
        "\n",
        "            finalPixel = (pixels[0][0] + pixels[0][1] + pixels[0][2] + pixels[1][0] + pixels[1][1] + pixels[1][2] + pixels[2][0] + pixels[2][1] + pixels[2][2]) / 9;\n",
        "\n",
        "            Pixel *ptrPixel = (Pixel *)&imageDataConvolution[(y * width * 4 + 4 * x)];\n",
        "            ptrPixel->r = finalPixel;\n",
        "            ptrPixel->g = finalPixel;\n",
        "            ptrPixel->b = finalPixel;\n",
        "            ptrPixel->a = 255;\n",
        "        }\n",
        "    }\n",
        "}\n",
        "\n",
        "void minPooling(unsigned char *originalImage, unsigned char *minPoolingImage, int width, int height)\n",
        "{\n",
        "    int counter = 0;\n",
        "\n",
        "    // Iterate over the image in 2x2 blocks\n",
        "    for (int y = 0; y < height; y += 2)\n",
        "    {\n",
        "        for (int x = 0; x < width; x += 2)\n",
        "        {\n",
        "            // For each channel, find the maximum value in the 2x2 block\n",
        "            for (int c = 0; c < 4; c++)\n",
        "            {\n",
        "                Pixel *ptrPixelMinPooling = (Pixel *)&minPoolingImage[counter];\n",
        "                unsigned char min = 255;\n",
        "                for (int dy = 0; dy < 2; dy++)\n",
        "                {\n",
        "                    for (int dx = 0; dx < 2; dx++)\n",
        "                    {\n",
        "                        // Calculate the index of the current pixel in the 1D array\n",
        "                        int index = (y + dy) * width * 4 + (x + dx) * 4 + c;\n",
        "                        unsigned char value = originalImage[index];\n",
        "                        min = (value < min) ? value : min;\n",
        "                    }\n",
        "                }\n",
        "                // Store the minimum value in the result array\n",
        "                ptrPixelMinPooling->r = min;\n",
        "                ptrPixelMinPooling->g = min;\n",
        "                ptrPixelMinPooling->b = min;\n",
        "                ptrPixelMinPooling->a = min;\n",
        "                counter++;\n",
        "            }\n",
        "        }\n",
        "    }\n",
        "}\n",
        "\n",
        "void maxPooling(unsigned char *originalImage, unsigned char *maxPoolingImage, int width, int height)\n",
        "{\n",
        "    int counter = 0;\n",
        "\n",
        "    // Iterate over the image in 2x2 blocks\n",
        "    for (int y = 0; y < height; y += 2)\n",
        "    {\n",
        "        for (int x = 0; x < width; x += 2)\n",
        "        {\n",
        "            // For each channel, find the maximum value in the 2x2 block\n",
        "            for (int c = 0; c < 4; c++)\n",
        "            {\n",
        "                Pixel *ptrPixelMaxPooling = (Pixel *)&maxPoolingImage[counter];\n",
        "                unsigned char max = 0;\n",
        "                for (int dy = 0; dy < 2; dy++)\n",
        "                {\n",
        "                    for (int dx = 0; dx < 2; dx++)\n",
        "                    {\n",
        "                        // Calculate the index of the current pixel in the 1D array\n",
        "                        int index = (y + dy) * width * 4 + (x + dx) * 4 + c;\n",
        "                        unsigned char value = originalImage[index];\n",
        "                        max = (value > max) ? value : max;\n",
        "                    }\n",
        "                }\n",
        "                // Store the maximum value in the result array\n",
        "                ptrPixelMaxPooling->r = max;\n",
        "                ptrPixelMaxPooling->g = max;\n",
        "                ptrPixelMaxPooling->b = max;\n",
        "                ptrPixelMaxPooling->a = max;\n",
        "                counter++;\n",
        "            }\n",
        "        }\n",
        "    }\n",
        "}\n",
        "\n",
        "\n"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "6NcXx8mVySrg"
      },
      "outputs": [],
      "source": [
        "!g++ ImageFilter_Sequential.c -o ImageFilter_Sequential"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "Z-t_1Bu7yYdu",
        "outputId": "8de0d0c2-a2ec-4518-9507-9973b31b4267"
      },
      "outputs": [
        {
          "name": "stdout",
          "output_type": "stream",
          "text": [
            "Loading Images/img0.png.\r\n",
            "Processing image grayscale.\r\n",
            "Done\r\n",
            "Convolving image.\r\n",
            "Done\r\n",
            "Writing image to disk\r\n",
            "Done\r\n",
            "\r\n",
            "Next Image\r\n",
            "Loading Images/img1.png.\r\n",
            "Processing image grayscale.\n",
            "Done\n",
            "Convolving image.\n",
            "Done\n",
            "Writing image to disk\n",
            "Done\n",
            "\n",
            "Next Image\n",
            "Loading Images/img2.png.\n",
            "Processing image grayscale.\n",
            "Done\n",
            "Convolving image.\n",
            "Done\n",
            "Writing image to disk\n",
            "Done\n",
            "\n",
            "Next Image\n",
            "Loading Images/img3.png.\n",
            "Processing image grayscale.\n",
            "Done\n",
            "Convolving image.\n",
            "Done\n",
            "Writing image to disk\n",
            "Done\n",
            "\n",
            "Next Image\n",
            "Loading Images/img4.png.\n",
            "Processing image grayscale.\n",
            "Done\n",
            "Convolving image.\n",
            "Done\n",
            "Writing image to disk\n",
            "Done\n",
            "\n",
            "Next Image\n",
            "Loading Images/img5.png.\n",
            "Processing image grayscale.\n",
            "Done\n",
            "Convolving image.\n",
            "Done\n",
            "Writing image to disk\n",
            "Done\n",
            "\n",
            "Next Image\n",
            "Loading Images/img6.png.\n",
            "Processing image grayscale.\n",
            "Done\n",
            "Convolving image.\n",
            "Done\n",
            "Writing image to disk\n",
            "Done\n",
            "\n",
            "Next Image\n",
            "Loading Images/img7.png.\n",
            "Processing image grayscale.\n",
            "Done\n",
            "Convolving image.\n",
            "Done\n",
            "Writing image to disk\n",
            "Done\n",
            "\n",
            "Next Image\n",
            "Loading Images/img8.png.\n",
            "Processing image grayscale.\n",
            "Done\n",
            "Convolving image.\n",
            "Done\n",
            "Writing image to disk\n",
            "Done\n",
            "\n",
            "Next Image\n",
            "Loading Images/img9.png.\n",
            "Processing image grayscale.\n",
            "Done\n",
            "Convolving image.\n",
            "Done\n",
            "Writing image to disk\n",
            "Done\n",
            "\n",
            "Next Image\n",
            "\n",
            "Program time: 223.028s\n"
          ]
        }
      ],
      "source": [
        "!./ImageFilter_Sequential"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "2XH8W8_vWUXG"
      },
      "source": [
        "# Write and compile the ImageFilter_Threads.c\n",
        "## This program filters 10 images using threads."
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "Fx4nDOHyWg1A",
        "outputId": "2b421d9a-e5ee-48a9-dc4e-ff4c15ebe57d"
      },
      "outputs": [
        {
          "name": "stdout",
          "output_type": "stream",
          "text": [
            "Overwriting ImageFilter_Threads.c\n"
          ]
        }
      ],
      "source": [
        "%%writefile ImageFilter_Threads.c\n",
        "#include <stdio.h>\n",
        "#include <stdlib.h>\n",
        "#include <pthread.h>\n",
        "#include <time.h>\n",
        "#define STB_IMAGE_IMPLEMENTATION\n",
        "#define STB_IMAGE_WRITE_IMPLEMENTATION\n",
        "#include \"stb_image.h\"\n",
        "#include \"stb_image_write.h\"\n",
        "\n",
        "// The amount of pictures you have.\n",
        "// Relies on the fact you name your images correct (ie img11.png).\n",
        "#define NUM_THREADS 9\n",
        "\n",
        "// Used to store which thread we are in.\n",
        "int threadNumber = 0;\n",
        "\n",
        "typedef struct Pixel\n",
        "{\n",
        "    unsigned char r, g, b, a;\n",
        "} Pixel;\n",
        "\n",
        "// Function to convert the image from colour to grayscale.\n",
        "void ConvertImageToGrayCpu(unsigned char *originalImage, unsigned char *imageRGBA, int width, int height);\n",
        "void convolveImage(unsigned char *imageRGBA, unsigned char *imageTest, int width, int height);\n",
        "void minPooling(unsigned char *originalImage, unsigned char *minPoolingImage, int width, int height);\n",
        "void maxPooling(unsigned char *originalImage, unsigned char *maxPoolingImage, int width, int height);\n",
        "\n",
        "// This function runs all the threads\n",
        "void *runThreads(void *vargp)\n",
        "{\n",
        "    clock_t timer_start, timer_end;\n",
        "    timer_start = clock(); // Start the timer\n",
        "    // Get which thread we are in, and do ++ for the next thread.\n",
        "    int threadId = threadNumber++;\n",
        "\n",
        "    char inputImage[64] = \"Images/img\";\n",
        "    char outputImageConvolution[64] = \"Output_Images/Convolution/img\";\n",
        "    char outputImagePoolingMax[64] = \"Output_Images/Pooling/imgMax\";\n",
        "    char outputImagePoolingMin[64] = \"Output_Images/Pooling/imgMin\";\n",
        "    char imageNumber[3];\n",
        "\n",
        "    // Make a complete path for the input (inputImage + imageNumber + extention)\n",
        "    sprintf(imageNumber, \"%d\", threadId); // Get the number of the current iteration and convert it to char imageNumber\n",
        "    strcat(inputImage, imageNumber);\n",
        "    strcat(inputImage, \".png\");\n",
        "\n",
        "    // Open image\n",
        "    int width, height, componentCount;\n",
        "    unsigned char *originalImage = stbi_load(inputImage, &width, &height, &componentCount, 4); // Saves original image\n",
        "    if (!originalImage)\n",
        "    {\n",
        "        printf(\"Failed to open Image\\r\\n\");\n",
        "        exit(-1);\n",
        "    }\n",
        "\n",
        "    // Validate image sizes\n",
        "    if (width % 32 || height % 32)\n",
        "    {\n",
        "        // NOTE: Leaked memory of \"imageData\"\n",
        "        printf(\"Width and/or Height is not dividable by 32! ( %s )\\r\\n\", inputImage);\n",
        "        exit(-1);\n",
        "    }\n",
        "\n",
        "    unsigned char *imageData = (unsigned char *)malloc(width * height * 4);           // Saves grayscale image\n",
        "    unsigned char *imageDataTest = (unsigned char *)malloc(width * height * 4);       // Saves output image\n",
        "    unsigned char *imageDataMinPooling = (unsigned char *)malloc(width * height * 4); // Saves Min pooling image\n",
        "    unsigned char *imageDataMaxPooling = (unsigned char *)malloc(width * height * 4); // Saves Max pooling image\n",
        "\n",
        "    // Process image on cpu\n",
        "    ConvertImageToGrayCpu(originalImage, imageData, width, height);\n",
        "\n",
        "    ////////////////////\n",
        "    // Convolve image //\n",
        "    ////////////////////\n",
        "    // Process image on cpu\n",
        "    convolveImage(imageData, imageDataTest, width, height);\n",
        "\n",
        "    // Make a complete path for the convolved image (outputImage convolution + imageNumber + extention)\n",
        "    strcat(outputImageConvolution, imageNumber);\n",
        "    strcat(outputImageConvolution, \".png\");\n",
        "\n",
        "    // Write image back to disk\n",
        "    stbi_write_png(outputImageConvolution, width, height, 4, imageDataTest, 4 * width);\n",
        "\n",
        "    /////////////////\n",
        "    // Min Pooling //\n",
        "    /////////////////\n",
        "    // Make a complete path for max pooling\n",
        "    sprintf(imageNumber, \"%d\", threadId); // Get the number of the current iteration and convert it to char imageNumber\n",
        "    strcat(outputImagePoolingMin, imageNumber);\n",
        "    strcat(outputImagePoolingMin, \".png\");\n",
        "\n",
        "    minPooling(originalImage, imageDataMinPooling, width, height);\n",
        "\n",
        "    // Write image back to disk\n",
        "    stbi_write_png(outputImagePoolingMin, width / 2, height / 2, 4, imageDataMinPooling, 4 * (width / 2));\n",
        "\n",
        "    /////////////////\n",
        "    // Max Pooling //\n",
        "    /////////////////\n",
        "    sprintf(imageNumber, \"%d\", threadId); // Get the number of the current iteration and convert it to char imageNumber\n",
        "    strcat(outputImagePoolingMax, imageNumber);\n",
        "    strcat(outputImagePoolingMax, \".png\");\n",
        "\n",
        "    maxPooling(originalImage, imageDataMaxPooling, width, height);\n",
        "\n",
        "    // Write image back to disk\n",
        "    stbi_write_png(outputImagePoolingMax, width / 2, height / 2, 4, imageDataMaxPooling, 4 * (width / 2));\n",
        "\n",
        "    stbi_image_free(originalImage);\n",
        "    free(imageData);\n",
        "    free(imageDataTest);\n",
        "    free(imageDataMinPooling);\n",
        "    free(imageDataMaxPooling);\n",
        "\n",
        "    timer_end = clock(); // end the timer\n",
        "    double time_spent = (double)(timer_end - timer_start) / CLOCKS_PER_SEC;\n",
        "    printf(\"\\n----------\\n%s took %.3fs\\n----------\\n\",inputImage, time_spent);\n",
        "    return NULL;\n",
        "}\n",
        "\n",
        "int main()\n",
        "{\n",
        "    pthread_t tid;\n",
        "\n",
        "    // Make NUM_THREADS amount of threads.\n",
        "    for (int i = 0; i <= NUM_THREADS; i++)\n",
        "    {\n",
        "        pthread_create(&tid, NULL, runThreads, NULL);\n",
        "    }\n",
        "\n",
        "    pthread_exit(NULL);\n",
        "    return 0;\n",
        "}\n",
        "\n",
        "void ConvertImageToGrayCpu(unsigned char *originalImage, unsigned char *imageRGBA, int width, int height)\n",
        "{\n",
        "    for (int y = 0; y < height; y++)\n",
        "    {\n",
        "        for (int x = 0; x < width; x++)\n",
        "        {\n",
        "            Pixel *ptrPixel = (Pixel *)&imageRGBA[y * width * 4 + 4 * x];\n",
        "            Pixel *ptrPixelOriginal = (Pixel *)&originalImage[y * width * 4 + 4 * x];\n",
        "            unsigned char pixelValue = (unsigned char)(ptrPixelOriginal->r * 0.2126f + ptrPixelOriginal->g * 0.7152f + ptrPixelOriginal->b * 0.0722f);\n",
        "            ptrPixel->r = pixelValue;\n",
        "            ptrPixel->g = pixelValue;\n",
        "            ptrPixel->b = pixelValue;\n",
        "            ptrPixel->a = 255;\n",
        "        }\n",
        "    }\n",
        "}\n",
        "\n",
        "void convolveImage(unsigned char *imageRGBA, unsigned char *imageTest, int width, int height)\n",
        "{\n",
        "    int kernel[3][3] =\n",
        "        {\n",
        "            {1, 0, -1},\n",
        "            {1, 0, -1},\n",
        "            {1, 0, -1}};\n",
        "\n",
        "    int pixels[3][3] = {0};\n",
        "    int finalPixel = 0;\n",
        "\n",
        "    for (int y = 0; y < height - 2; y++)\n",
        "    {\n",
        "        for (int x = 0; x < width - 2; x++)\n",
        "        {\n",
        "            for (int i = 0; i <= 2; i++)\n",
        "            {\n",
        "                Pixel *ptrPixel = (Pixel *)&imageRGBA[(y * width * 4 + 4 * x) + i * 4];\n",
        "\n",
        "                pixels[0][i] = ptrPixel->r * kernel[0][i];\n",
        "            }\n",
        "\n",
        "            for (int i = 0; i <= 2; i++)\n",
        "            {\n",
        "                Pixel *ptrPixel = (Pixel *)&imageRGBA[(y * width * 4 + 4 * x) + width * 4 + i * 4];\n",
        "\n",
        "                pixels[1][i] = ptrPixel->r * kernel[1][i];\n",
        "            }\n",
        "\n",
        "            for (int i = 0; i <= 2; i++)\n",
        "            {\n",
        "                Pixel *ptrPixel = (Pixel *)&imageRGBA[(y * width * 4 + 4 * x) + (2 * width * 4) + i * 4];\n",
        "\n",
        "                pixels[2][i] = ptrPixel->r * kernel[2][i];\n",
        "            }\n",
        "\n",
        "            finalPixel = (pixels[0][0] + pixels[0][1] + pixels[0][2] + pixels[1][0] + pixels[1][1] + pixels[1][2] + pixels[2][0] + pixels[2][1] + pixels[2][2]) / 9;\n",
        "\n",
        "            Pixel *ptrPixel = (Pixel *)&imageTest[(y * width * 4 + 4 * x)];\n",
        "            ptrPixel->r = finalPixel;\n",
        "            ptrPixel->g = finalPixel;\n",
        "            ptrPixel->b = finalPixel;\n",
        "            ptrPixel->a = 255;\n",
        "        }\n",
        "    }\n",
        "}\n",
        "\n",
        "void minPooling(unsigned char *originalImage, unsigned char *minPoolingImage, int width, int height)\n",
        "{\n",
        "    int counter = 0;\n",
        "\n",
        "    // Iterate over the image in 2x2 blocks\n",
        "    for (int y = 0; y < height; y += 2)\n",
        "    {\n",
        "        for (int x = 0; x < width; x += 2)\n",
        "        {\n",
        "            // For each channel, find the maximum value in the 2x2 block\n",
        "            for (int c = 0; c < 4; c++)\n",
        "            {\n",
        "                Pixel *ptrPixelMinPooling = (Pixel *)&minPoolingImage[counter];\n",
        "                unsigned char min = 255;\n",
        "                for (int dy = 0; dy < 2; dy++)\n",
        "                {\n",
        "                    for (int dx = 0; dx < 2; dx++)\n",
        "                    {\n",
        "                        // Calculate the index of the current pixel in the 1D array\n",
        "                        int index = (y + dy) * width * 4 + (x + dx) * 4 + c;\n",
        "                        unsigned char value = originalImage[index];\n",
        "                        min = (value < min) ? value : min;\n",
        "                    }\n",
        "                }\n",
        "                // Store the minimum value in the result array\n",
        "                ptrPixelMinPooling->r = min;\n",
        "                ptrPixelMinPooling->g = min;\n",
        "                ptrPixelMinPooling->b = min;\n",
        "                ptrPixelMinPooling->a = min;\n",
        "                counter++;\n",
        "            }\n",
        "        }\n",
        "    }\n",
        "}\n",
        "\n",
        "void maxPooling(unsigned char *originalImage, unsigned char *maxPoolingImage, int width, int height)\n",
        "{\n",
        "    int counter = 0;\n",
        "\n",
        "    // Iterate over the image in 2x2 blocks\n",
        "    for (int y = 0; y < height; y += 2)\n",
        "    {\n",
        "        for (int x = 0; x < width; x += 2)\n",
        "        {\n",
        "            // For each channel, find the maximum value in the 2x2 block\n",
        "            for (int c = 0; c < 4; c++)\n",
        "            {\n",
        "                Pixel *ptrPixelMaxPooling = (Pixel *)&maxPoolingImage[counter];\n",
        "                unsigned char max = 0;\n",
        "                for (int dy = 0; dy < 2; dy++)\n",
        "                {\n",
        "                    for (int dx = 0; dx < 2; dx++)\n",
        "                    {\n",
        "                        // Calculate the index of the current pixel in the 1D array\n",
        "                        int index = (y + dy) * width * 4 + (x + dx) * 4 + c;\n",
        "                        unsigned char value = originalImage[index];\n",
        "                        max = (value > max) ? value : max;\n",
        "                    }\n",
        "                }\n",
        "                // Store the maximum value in the result array\n",
        "                ptrPixelMaxPooling->r = max;\n",
        "                ptrPixelMaxPooling->g = max;\n",
        "                ptrPixelMaxPooling->b = max;\n",
        "                ptrPixelMaxPooling->a = max;\n",
        "                counter++;\n",
        "            }\n",
        "        }\n",
        "    }\n",
        "}"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "vxXnrVlwWyoc"
      },
      "outputs": [],
      "source": [
        "!g++ -pthread ImageFilter_Threads.c -o ImageFilter_Threads"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "id": "ypuVzGLhXdGZ"
      },
      "outputs": [],
      "source": [
        "!./ImageFilter_Threads"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "LfqZ8jPjgL5J"
      },
      "source": [
        "# Write and compile the ImageFilter.cu\n",
        "## This program filters 10 accelerated with CUDA."
      ]
    },
    {
      "cell_type": "markdown",
      "source": [
        "Het programma loopt door de grote afbeeldingen tegen de gebruikers limiten van colab aan bij het filteren van 10 afbeeldingen. Om het programma toch te laten werken kan de define worden aangepast naar een maximum van 4 images."
      ],
      "metadata": {
        "id": "SyHCrhgCyFDT"
      }
    },
    {
      "cell_type": "code",
      "execution_count": null,
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "ILKOX9YOqX-S",
        "outputId": "1fd2f541-8f1b-4c23-de52-637b2a46d535"
      },
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Writing ImageFilter_CUDA.cu\n"
          ]
        }
      ],
      "source": [
        "%%writefile ImageFilter_CUDA.cu\n",
        "#include <stdio.h>\n",
        "#include <stdlib.h>\n",
        "#include <time.h>\n",
        "#include \"cuda.h\"\n",
        "#include \"cuda_runtime.h\"\n",
        "#define STB_IMAGE_IMPLEMENTATION\n",
        "#define STB_IMAGE_WRITE_IMPLEMENTATION\n",
        "#include \"stb_image.h\"\n",
        "#include \"stb_image_write.h\"\n",
        "\n",
        "#define NUMBER_OF_IMAGES 4 // (Standaard 10 maar door limiten van colab werkt het maar tot 4)\n",
        "\n",
        "typedef struct Pixel\n",
        "{\n",
        "    unsigned char r, g, b, a;\n",
        "} Pixel;\n",
        "\n",
        "__global__ void ConvertImageToGrayCpu(unsigned char *originalImage, unsigned char *imageDataGrayscale, int width, int height);\n",
        "__global__ void convolveImage(unsigned char *imageDataGrayscale, unsigned char *imageDataConvolution, int width, int height);\n",
        "__global__ void minPooling(unsigned char *originalImage, unsigned char *minPoolingImage, int width, int height);\n",
        "__global__ void maxPooling(unsigned char *originalImage, unsigned char *maxPoolingImage, int width, int height);\n",
        "\n",
        "int main(int argc, char **argv)\n",
        "{\n",
        "    clock_t timer_start, timer_end, timer_start_process, timer_end_process;\n",
        "    timer_start = clock();\n",
        "\n",
        "    size_t threadsPerBlock = 128;\n",
        "    size_t numberOfBlocks = 32;\n",
        "\n",
        "    cudaStream_t stream[10];\n",
        "\n",
        "    for (int i = 0; i < 10; i++)\n",
        "    {\n",
        "        cudaStreamCreate(&stream[i]);\n",
        "    }\n",
        "\n",
        "    printf(\"Building filepaths\\r\\n\");\n",
        "\n",
        "    const char *inputFileName[10] =\n",
        "        {\n",
        "            \"Images/img0.png\",\n",
        "            \"Images/img1.png\",\n",
        "            \"Images/img2.png\",\n",
        "            \"Images/img3.png\",\n",
        "            \"Images/img4.png\",\n",
        "            \"Images/img5.png\",\n",
        "            \"Images/img6.png\",\n",
        "            \"Images/img7.png\",\n",
        "            \"Images/img8.png\",\n",
        "            \"Images/img9.png\",\n",
        "        };\n",
        "\n",
        "    // Build output filename\n",
        "    const char *fileNameOutConvolution[10] =\n",
        "        {\n",
        "            \"Output_Images/Convolution/OutputConvolution0.png\",\n",
        "            \"Output_Images/Convolution/OutputConvolution1.png\",\n",
        "            \"Output_Images/Convolution/OutputConvolution2.png\",\n",
        "            \"Output_Images/Convolution/OutputConvolution3.png\",\n",
        "            \"Output_Images/Convolution/OutputConvolution4.png\",\n",
        "            \"Output_Images/Convolution/OutputConvolution5.png\",\n",
        "            \"Output_Images/Convolution/OutputConvolution6.png\",\n",
        "            \"Output_Images/Convolution/OutputConvolution7.png\",\n",
        "            \"Output_Images/Convolution/OutputConvolution8.png\",\n",
        "            \"Output_Images/Convolution/OutputConvolution9.png\",\n",
        "        };\n",
        "\n",
        "    const char *fileNameOutMinPooling[10] =\n",
        "        {\n",
        "            \"Output_Images/Pooling/OutputMinPooling0.png\",\n",
        "            \"Output_Images/Pooling/OutputMinPooling1.png\",\n",
        "            \"Output_Images/Pooling/OutputMinPooling2.png\",\n",
        "            \"Output_Images/Pooling/OutputMinPooling3.png\",\n",
        "            \"Output_Images/Pooling/OutputMinPooling4.png\",\n",
        "            \"Output_Images/Pooling/OutputMinPooling5.png\",\n",
        "            \"Output_Images/Pooling/OutputMinPooling6.png\",\n",
        "            \"Output_Images/Pooling/OutputMinPooling7.png\",\n",
        "            \"Output_Images/Pooling/OutputMinPooling8.png\",\n",
        "            \"Output_Images/Pooling/OutputMinPooling9.png\",\n",
        "        };\n",
        "\n",
        "    const char *fileNameOutMaxPooling[10] =\n",
        "        {\n",
        "            \"Output_Images/Pooling/OutputMaxPooling0.png\",\n",
        "            \"Output_Images/Pooling/OutputMaxPooling1.png\",\n",
        "            \"Output_Images/Pooling/OutputMaxPooling2.png\",\n",
        "            \"Output_Images/Pooling/OutputMaxPooling3.png\",\n",
        "            \"Output_Images/Pooling/OutputMaxPooling4.png\",\n",
        "            \"Output_Images/Pooling/OutputMaxPooling5.png\",\n",
        "            \"Output_Images/Pooling/OutputMaxPooling6.png\",\n",
        "            \"Output_Images/Pooling/OutputMaxPooling7.png\",\n",
        "            \"Output_Images/Pooling/OutputMaxPooling8.png\",\n",
        "            \"Output_Images/Pooling/OutputMaxPooling9.png\",\n",
        "        };\n",
        "\n",
        "    int width[10], height[10], componentCount[10], size[10];\n",
        "\n",
        "    unsigned char *originalImageHost[10];\n",
        "    unsigned char *imageDataConvolutionHost[10]; // Saves output image\n",
        "    unsigned char *imageDataMinPoolingHost[10];  // Saves Min pooling image\n",
        "    unsigned char *imageDataMaxPoolingHost[10];  // Saves Max pooling image\n",
        "    unsigned char *originalImage[10];            // Saves the original image on host\n",
        "    unsigned char *imageDataGrayscale[10];       // Saves the grayscale image on device\n",
        "    unsigned char *imageDataConvolution[10];     // Saves the convolved image\n",
        "    unsigned char *imageDataMinPooling[10];      // Saves the min pooled image\n",
        "    unsigned char *imageDataMaxPooling[10];      // Saves the max pooled image\n",
        "\n",
        "    printf(\"Done\\r\\n\");\n",
        "\n",
        "    printf(\"Loading png files\\r\\n\");\n",
        "\n",
        "    for (int i = 0; i < NUMBER_OF_IMAGES; i++)\n",
        "    {\n",
        "        originalImageHost[i] = stbi_load(inputFileName[i], &width[i], &height[i], &componentCount[i], 4);\n",
        "\n",
        "        size[i] = height[i] * width[i] * 4;\n",
        "\n",
        "        // Saves output image\n",
        "        imageDataConvolutionHost[i] = (unsigned char *)malloc(size[i]);\n",
        "\n",
        "        // Saves Min pooling image\n",
        "        imageDataMinPoolingHost[i] = (unsigned char *)malloc(size[i]);\n",
        "\n",
        "        // Saves Max pooling image\n",
        "        imageDataMaxPoolingHost[i] = (unsigned char *)malloc(size[i]);\n",
        "\n",
        "        cudaMalloc(&originalImage[i], size[i]);\n",
        "        cudaMalloc(&imageDataGrayscale[i], size[i]);\n",
        "        cudaMalloc(&imageDataConvolution[i], size[i]);\n",
        "        cudaMalloc(&imageDataMinPooling[i], size[i]);\n",
        "        cudaMalloc(&imageDataMaxPooling[i], size[i]);\n",
        "\n",
        "        cudaMemcpy(originalImage[i], originalImageHost[i], size[i], cudaMemcpyHostToDevice);\n",
        "    }\n",
        "\n",
        "    printf(\"Done\\r\\n\");\n",
        "\n",
        "    timer_start_process = clock();\n",
        "\n",
        "    // Process grayscale\n",
        "    printf(\"Processing images grayscale\\r\\n\");\n",
        "    for (int i = 0; i < NUMBER_OF_IMAGES; i++)\n",
        "    {\n",
        "        ConvertImageToGrayCpu<<<numberOfBlocks, threadsPerBlock, i, stream[i]>>>(originalImage[i], imageDataGrayscale[i], width[i], height[i]);\n",
        "    }\n",
        "    cudaDeviceSynchronize();\n",
        "    printf(\"Done\\r\\n\");\n",
        "\n",
        "    // Process convolution\n",
        "    printf(\"Processing image convolution\\r\\n\");\n",
        "    for (int i = 0; i < NUMBER_OF_IMAGES; i++)\n",
        "    {\n",
        "        convolveImage<<<numberOfBlocks, threadsPerBlock, i, stream[i]>>>(imageDataGrayscale[i], imageDataConvolution[i], width[i], height[i]);\n",
        "    }\n",
        "    cudaDeviceSynchronize();\n",
        "    printf(\"Done\\r\\n\");\n",
        "\n",
        "    // Process min pooling\n",
        "    printf(\"Processing images minimum pooling\\r\\n\");\n",
        "    for (int i = 0; i < NUMBER_OF_IMAGES; i++)\n",
        "    {\n",
        "        dim3 gridDim(width[i] / 2, height[i] / 2);\n",
        "        dim3 blockDim(2, 2);\n",
        "        // minPooling<<<numberOfBlocks, threadsPerBlock, i, stream[i]>>>(originalImage[i], imageDataMinPooling[i], width[i], height[i]);\n",
        "        minPooling<<<gridDim, blockDim, i, stream[i]>>>(originalImage[i], imageDataMinPooling[i], width[i], height[i]);\n",
        "    }\n",
        "    cudaDeviceSynchronize();\n",
        "    printf(\"Done\\r\\n\");\n",
        "\n",
        "    // Process max pooling\n",
        "    printf(\"Processing image maximum pooling\\r\\n\");\n",
        "    for (int i = 0; i < NUMBER_OF_IMAGES; i++)\n",
        "    {\n",
        "        dim3 gridDim(width[i] / 2, height[i] / 2);\n",
        "        dim3 blockDim(2, 2);\n",
        "        // maxPooling<<<numberOfBlocks, threadsPerBlock, i, stream[i]>>>(originalImage[i], imageDataMaxPooling[i], width[i], height[i]);\n",
        "        maxPooling<<<gridDim, blockDim, i, stream[i]>>>(originalImage[i], imageDataMaxPooling[i], width[i], height[i]);\n",
        "    }\n",
        "    cudaDeviceSynchronize();\n",
        "    printf(\"Done\\r\\n\");\n",
        "\n",
        "    timer_end_process = clock();\n",
        "\n",
        "    // Writing Convolved images\n",
        "    printf(\"Writing convolved png to disk\\r\\n\");\n",
        "    for (int i = 0; i < NUMBER_OF_IMAGES; i++)\n",
        "    {\n",
        "        printf(\"Image %d of %d\\r\\n\", i + 1, 10);\n",
        "        cudaMemcpy(imageDataConvolutionHost[i], imageDataConvolution[i], size[i], cudaMemcpyDeviceToHost);\n",
        "        stbi_write_png(fileNameOutConvolution[i], width[i] - 2, height[i] - 2, 4, imageDataConvolutionHost[i], 4 * width[i]);\n",
        "    }\n",
        "    printf(\"Done\\r\\n\");\n",
        "\n",
        "    // Writing min pooled images\n",
        "    printf(\"Writing min pooling png to disk\\r\\n\");\n",
        "    for (int i = 0; i < NUMBER_OF_IMAGES; i++)\n",
        "    {\n",
        "        printf(\"Image %d of %d\\r\\n\", i + 1, 10);\n",
        "        cudaMemcpy(imageDataMinPoolingHost[i], imageDataMinPooling[i], size[i], cudaMemcpyDeviceToHost);\n",
        "        stbi_write_png(fileNameOutMinPooling[i], width[i] / 2, height[i] / 2, 4, imageDataMinPoolingHost[i], 4 * (width[i] / 2));\n",
        "    }\n",
        "    printf(\"Done\\r\\n\");\n",
        "\n",
        "    // Writing max pooled images\n",
        "    printf(\"Writing max pooling png to disk\\r\\n\");\n",
        "    for (int i = 0; i < NUMBER_OF_IMAGES; i++)\n",
        "    {\n",
        "        printf(\"Image %d of %d\\r\\n\", i + 1, 10);\n",
        "        cudaMemcpy(imageDataMaxPoolingHost[i], imageDataMaxPooling[i], size[i], cudaMemcpyDeviceToHost);\n",
        "        stbi_write_png(fileNameOutMaxPooling[i], width[i] / 2, height[i] / 2, 4, imageDataMaxPoolingHost[i], 4 * (width[i] / 2));\n",
        "    }\n",
        "    printf(\"Done\\r\\n\");\n",
        "\n",
        "    // Free memory and destroy streams\n",
        "    for (int i = 0; i < NUMBER_OF_IMAGES; i++)\n",
        "    {\n",
        "        stbi_image_free(originalImageHost[i]);\n",
        "\n",
        "        free(imageDataConvolutionHost[i]);\n",
        "        free(imageDataMinPoolingHost[i]);\n",
        "        free(imageDataMaxPoolingHost[i]);\n",
        "\n",
        "        cudaFree(originalImage[i]);\n",
        "        cudaFree(imageDataConvolution[i]);\n",
        "        cudaFree(imageDataMinPooling[i]);\n",
        "        cudaFree(imageDataMaxPooling[i]);\n",
        "\n",
        "        cudaStreamDestroy(stream[i]);\n",
        "    }\n",
        "\n",
        "    timer_end = clock(); // end the timer\n",
        "    double time_spent = (double)(timer_end - timer_start) / CLOCKS_PER_SEC;\n",
        "    printf(\"\\nTotal program time: %.3fs\\n\", time_spent);\n",
        "\n",
        "    double time_spent_process = (double)(timer_end_process - timer_start_process) / CLOCKS_PER_SEC;\n",
        "    printf(\"\\nProcessing time (CUDA Kernels): %.3fs\\n\", time_spent_process);\n",
        "\n",
        "    return 0;\n",
        "}\n",
        "\n",
        "__global__ void ConvertImageToGrayCpu(unsigned char *originalImage, unsigned char *imageDataGrayscale, int width, int height)\n",
        "{\n",
        "    int idx = (threadIdx.x + blockIdx.x * blockDim.x) * 4;\n",
        "    int gridStride = blockDim.x * gridDim.x;\n",
        "    int totalPixels = width * height * 4;\n",
        "\n",
        "    for (int x = idx; x < totalPixels; x += gridStride)\n",
        "    {\n",
        "        Pixel *ptrPixel = (Pixel *)&imageDataGrayscale[x];\n",
        "        Pixel *ptrPixelOriginal = (Pixel *)&originalImage[x];\n",
        "        unsigned char pixelValue = (unsigned char)(ptrPixelOriginal->r * 0.2126f + ptrPixelOriginal->g * 0.7152f + ptrPixelOriginal->b * 0.0722f);\n",
        "        ptrPixel->r = pixelValue;\n",
        "        ptrPixel->g = pixelValue;\n",
        "        ptrPixel->b = pixelValue;\n",
        "        ptrPixel->a = 255;\n",
        "    }\n",
        "}\n",
        "\n",
        "__global__ void convolveImage(unsigned char *imageDataGrayscale, unsigned char *imageDataConvolution, int width, int height)\n",
        "{\n",
        "    int idx = (threadIdx.x + blockIdx.x * blockDim.x) * 4;\n",
        "    int gridStride = blockDim.x * gridDim.x;\n",
        "    int totalPixels = width * height * 4;\n",
        "\n",
        "    int kernel[3][3] =\n",
        "        {\n",
        "            {1, 0, -1},\n",
        "            {1, 0, -1},\n",
        "            {1, 0, -1}};\n",
        "\n",
        "    int pixels[3][3] = {0};\n",
        "    int finalPixel = 0;\n",
        "\n",
        "    for (int x = idx; x < totalPixels - 2; x += gridStride)\n",
        "    {\n",
        "        for (int i = 0; i <= 2; i++)\n",
        "        {\n",
        "            Pixel *ptrPixel = (Pixel *)&imageDataGrayscale[x + i * 4];\n",
        "\n",
        "            pixels[0][i] = ptrPixel->r * kernel[0][i];\n",
        "        }\n",
        "\n",
        "        for (int i = 0; i <= 2; i++)\n",
        "        {\n",
        "            Pixel *ptrPixel = (Pixel *)&imageDataGrayscale[x + width * 4 + i * 4];\n",
        "\n",
        "            pixels[1][i] = ptrPixel->r * kernel[1][i];\n",
        "        }\n",
        "\n",
        "        for (int i = 0; i <= 2; i++)\n",
        "        {\n",
        "            Pixel *ptrPixel = (Pixel *)&imageDataGrayscale[x + (2 * width * 4) + i * 4];\n",
        "\n",
        "            pixels[2][i] = ptrPixel->r * kernel[2][i];\n",
        "        }\n",
        "\n",
        "        finalPixel = (pixels[0][0] + pixels[0][1] + pixels[0][2] + pixels[1][0] + pixels[1][1] + pixels[1][2] + pixels[2][0] + pixels[2][1] + pixels[2][2]) / 9;\n",
        "\n",
        "        Pixel *ptrPixel = (Pixel *)&imageDataConvolution[x];\n",
        "        ptrPixel->r = finalPixel;\n",
        "        ptrPixel->g = finalPixel;\n",
        "        ptrPixel->b = finalPixel;\n",
        "        ptrPixel->a = 255;\n",
        "    }\n",
        "}\n",
        "\n",
        "__global__ void minPooling(unsigned char *originalImage, unsigned char *minPoolingImage, int width, int height)\n",
        "{\n",
        "    // Calculate the 2D block index\n",
        "    int blockY = blockIdx.y;\n",
        "    int blockX = blockIdx.x;\n",
        "\n",
        "    // Calculate the starting position of the 2x2 block\n",
        "    int y = blockY * 2;\n",
        "    int x = blockX * 2;\n",
        "\n",
        "    // Calculate the index of the current pixel in the 1D arrays\n",
        "    int indexOut = blockY * width / 2 * 4 + blockX * 4;\n",
        "\n",
        "    // For each channel, find the minimum value in the 2x2 block\n",
        "    for (int c = 0; c < 4; c++)\n",
        "    {\n",
        "        unsigned char min = 255;\n",
        "        for (int dy = 0; dy < 2; dy++)\n",
        "        {\n",
        "            for (int dx = 0; dx < 2; dx++)\n",
        "            {\n",
        "                // Calculate the index of the current pixel in the 1D array\n",
        "                int index = (y + dy) * width * 4 + (x + dx) * 4 + c;\n",
        "                unsigned char value = originalImage[index];\n",
        "                min = (value < min) ? value : min;\n",
        "            }\n",
        "        }\n",
        "        // Store the minimum value in the result array\n",
        "        minPoolingImage[indexOut + c] = min;\n",
        "    }\n",
        "}\n",
        "\n",
        "__global__ void maxPooling(unsigned char *originalImage, unsigned char *minPoolingImage, int width, int height)\n",
        "{\n",
        "    // Calculate the 2D block index\n",
        "    int blockY = blockIdx.y;\n",
        "    int blockX = blockIdx.x;\n",
        "\n",
        "    // Calculate the starting position of the 2x2 block\n",
        "    int y = blockY * 2;\n",
        "    int x = blockX * 2;\n",
        "\n",
        "    // Calculate the index of the current pixel in the 1D arrays\n",
        "    int indexOut = blockY * width / 2 * 4 + blockX * 4;\n",
        "\n",
        "    // For each channel, find the minimum value in the 2x2 block\n",
        "    for (int c = 0; c < 4; c++)\n",
        "    {\n",
        "        unsigned char min = 0;\n",
        "        for (int dy = 0; dy < 2; dy++)\n",
        "        {\n",
        "            for (int dx = 0; dx < 2; dx++)\n",
        "            {\n",
        "                // Calculate the index of the current pixel in the 1D array\n",
        "                int index = (y + dy) * width * 4 + (x + dx) * 4 + c;\n",
        "                unsigned char value = originalImage[index];\n",
        "                min = (value > min) ? value : min;\n",
        "            }\n",
        "        }\n",
        "        // Store the minimum value in the result array\n",
        "        minPoolingImage[indexOut + c] = min;\n",
        "    }\n",
        "}\n",
        "\n"
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "!nvcc ImageFilter_CUDA.cu -o ImageFilter_CUDA --generate-line-info"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "DSuaXagjy51R",
        "outputId": "e98ec0e9-cc87-4de3-addf-808bc4d09435"
      },
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "stb_image.h(4207): warning: variable \"old_limit\" was set but never used\n",
            "\n",
            "stb_image.h(5121): warning: variable \"idata_limit_old\" was set but never used\n",
            "\n",
            "stb_image.h(6894): warning: variable \"out_size\" was set but never used\n",
            "\n",
            "stb_image.h(6895): warning: variable \"delays_size\" was set but never used\n",
            "\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "!nvprof --print-gpu-trace ./ImageFilter_CUDA"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "LizHoTeCydYv",
        "outputId": "ccbb17a6-ad0b-48fa-d8fd-f6ade4e4af38"
      },
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "==1219== NVPROF is profiling process 1219, command: ./ImageFilter_CUDA\n",
            "Building filepaths\n",
            "Done\n",
            "Loading png files\n",
            "Done\n",
            "Processing images grayscale\n",
            "Done\n",
            "Processing image convolution\n",
            "Done\n",
            "Processing images minimum pooling\n",
            "Done\n",
            "Processing image maximum pooling\n",
            "Done\n",
            "Writing convolved png to disk\n",
            "Image 1 of 10\n",
            "Image 2 of 10\n",
            "Image 3 of 10\n",
            "Image 4 of 10\n",
            "Done\n",
            "Writing min pooling png to disk\n",
            "Image 1 of 10\n",
            "Image 2 of 10\n",
            "Image 3 of 10\n",
            "Image 4 of 10\n",
            "Done\n",
            "Writing max pooling png to disk\n",
            "Image 1 of 10\n",
            "Image 2 of 10\n",
            "Image 3 of 10\n",
            "Image 4 of 10\n",
            "Done\n",
            "\n",
            "Total program time: 8.696s\n",
            "\n",
            "Processing time (CUDA Kernels): 0.034s\n",
            "==1219== Profiling application: ./ImageFilter_CUDA\n",
            "==1219== Profiling result:\n",
            "   Start  Duration            Grid Size      Block Size     Regs*    SSMem*    DSMem*      Size  Throughput  SrcMemType  DstMemType           Device   Context    Stream  Name\n",
            "525.97ms  3.4880us                    -               -         -         -         -  16.000KB  4.3747GB/s    Pageable      Device     Tesla T4 (0)         1         7  [CUDA memcpy HtoD]\n",
            "785.25ms  8.6600ms                    -               -         -         -         -  39.063MB  4.4049GB/s    Pageable      Device     Tesla T4 (0)         1         7  [CUDA memcpy HtoD]\n",
            "855.50ms  517.50us                    -               -         -         -         -  3.0000MB  5.6613GB/s    Pageable      Device     Tesla T4 (0)         1         7  [CUDA memcpy HtoD]\n",
            "983.05ms  1.4897ms                    -               -         -         -         -  7.5703MB  4.9626GB/s    Pageable      Device     Tesla T4 (0)         1         7  [CUDA memcpy HtoD]\n",
            "984.58ms  7.8080us             (32 1 1)       (128 1 1)        17        0B        0B         -           -           -           -     Tesla T4 (0)         1        13  ConvertImageToGrayCpu(unsigned char*, unsigned char*, int, int) [145]\n",
            "984.59ms  6.6385ms             (32 1 1)       (128 1 1)        17        0B        1B         -           -           -           -     Tesla T4 (0)         1        14  ConvertImageToGrayCpu(unsigned char*, unsigned char*, int, int) [146]\n",
            "984.60ms  542.01us             (32 1 1)       (128 1 1)        17        0B        2B         -           -           -           -     Tesla T4 (0)         1        15  ConvertImageToGrayCpu(unsigned char*, unsigned char*, int, int) [147]\n",
            "984.61ms  1.3444ms             (32 1 1)       (128 1 1)        17        0B        3B         -           -           -           -     Tesla T4 (0)         1        16  ConvertImageToGrayCpu(unsigned char*, unsigned char*, int, int) [148]\n",
            "991.27ms  7.0070us             (32 1 1)       (128 1 1)        18        0B        0B         -           -           -           -     Tesla T4 (0)         1        13  convolveImage(unsigned char*, unsigned char*, int, int) [150]\n",
            "991.28ms  6.5780ms             (32 1 1)       (128 1 1)        18        0B        1B         -           -           -           -     Tesla T4 (0)         1        14  convolveImage(unsigned char*, unsigned char*, int, int) [151]\n",
            "991.29ms  557.43us             (32 1 1)       (128 1 1)        18        0B        2B         -           -           -           -     Tesla T4 (0)         1        15  convolveImage(unsigned char*, unsigned char*, int, int) [152]\n",
            "991.29ms  1.3665ms             (32 1 1)       (128 1 1)        18        0B        3B         -           -           -           -     Tesla T4 (0)         1        16  convolveImage(unsigned char*, unsigned char*, int, int) [153]\n",
            "997.88ms  9.6960us            (32 32 1)         (2 2 1)        16        0B        0B         -           -           -           -     Tesla T4 (0)         1        13  minPooling(unsigned char*, unsigned char*, int, int) [155]\n",
            "997.89ms  8.2163ms        (1600 1600 1)         (2 2 1)        16        0B        1B         -           -           -           -     Tesla T4 (0)         1        14  minPooling(unsigned char*, unsigned char*, int, int) [156]\n",
            "1.00611s  641.88us          (384 512 1)         (2 2 1)        16        0B        2B         -           -           -           -     Tesla T4 (0)         1        15  minPooling(unsigned char*, unsigned char*, int, int) [157]\n",
            "1.00674s  1.5946ms          (608 816 1)         (2 2 1)        16        0B        3B         -           -           -           -     Tesla T4 (0)         1        16  minPooling(unsigned char*, unsigned char*, int, int) [158]\n",
            "1.00836s  8.9280us            (32 32 1)         (2 2 1)        16        0B        0B         -           -           -           -     Tesla T4 (0)         1        13  maxPooling(unsigned char*, unsigned char*, int, int) [160]\n",
            "1.00837s  8.2415ms        (1600 1600 1)         (2 2 1)        16        0B        1B         -           -           -           -     Tesla T4 (0)         1        14  maxPooling(unsigned char*, unsigned char*, int, int) [161]\n",
            "1.01661s  644.53us          (384 512 1)         (2 2 1)        16        0B        2B         -           -           -           -     Tesla T4 (0)         1        15  maxPooling(unsigned char*, unsigned char*, int, int) [162]\n",
            "1.01725s  1.5973ms          (608 816 1)         (2 2 1)        16        0B        3B         -           -           -           -     Tesla T4 (0)         1        16  maxPooling(unsigned char*, unsigned char*, int, int) [163]\n",
            "1.01892s  3.3280us                    -               -         -         -         -  16.000KB  4.5850GB/s      Device    Pageable     Tesla T4 (0)         1         7  [CUDA memcpy DtoH]\n",
            "1.02101s  26.071ms                    -               -         -         -         -  39.063MB  1.4632GB/s      Device    Pageable     Tesla T4 (0)         1         7  [CUDA memcpy DtoH]\n",
            "4.72120s  463.64us                    -               -         -         -         -  3.0000MB  6.3189GB/s      Device    Pageable     Tesla T4 (0)         1         7  [CUDA memcpy DtoH]\n",
            "5.07280s  2.1179ms                    -               -         -         -         -  7.5703MB  3.4907GB/s      Device    Pageable     Tesla T4 (0)         1         7  [CUDA memcpy DtoH]\n",
            "5.69251s  3.2320us                    -               -         -         -         -  16.000KB  4.7212GB/s      Device    Pageable     Tesla T4 (0)         1         7  [CUDA memcpy DtoH]\n",
            "5.69326s  25.721ms                    -               -         -         -         -  39.063MB  1.4831GB/s      Device    Pageable     Tesla T4 (0)         1         7  [CUDA memcpy DtoH]\n",
            "6.69726s  470.20us                    -               -         -         -         -  3.0000MB  6.2307GB/s      Device    Pageable     Tesla T4 (0)         1         7  [CUDA memcpy DtoH]\n",
            "6.84467s  1.9548ms                    -               -         -         -         -  7.5703MB  3.7819GB/s      Device    Pageable     Tesla T4 (0)         1         7  [CUDA memcpy DtoH]\n",
            "7.26798s  3.2320us                    -               -         -         -         -  16.000KB  4.7212GB/s      Device    Pageable     Tesla T4 (0)         1         7  [CUDA memcpy DtoH]\n",
            "7.26875s  25.988ms                    -               -         -         -         -  39.063MB  1.4679GB/s      Device    Pageable     Tesla T4 (0)         1         7  [CUDA memcpy DtoH]\n",
            "8.30683s  488.86us                    -               -         -         -         -  3.0000MB  5.9930GB/s      Device    Pageable     Tesla T4 (0)         1         7  [CUDA memcpy DtoH]\n",
            "8.45294s  4.4649ms                    -               -         -         -         -  7.5703MB  1.6558GB/s      Device    Pageable     Tesla T4 (0)         1         7  [CUDA memcpy DtoH]\n",
            "\n",
            "Regs: Number of registers used per CUDA thread. This number includes registers used internally by the CUDA driver and/or tools and can be more than what the compiler shows.\n",
            "SSMem: Static shared memory allocated per CUDA block.\n",
            "DSMem: Dynamic shared memory allocated per CUDA block.\n",
            "SrcMemType: The type of source memory accessed by memory operation/copy\n",
            "DstMemType: The type of destination memory accessed by memory operation/copy\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "!nsys profile ImageFilter_CUDA"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "-nmf-YIzzXAK",
        "outputId": "6fa2803b-d27f-4ddc-886a-fe4c935912d2"
      },
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Warning: LBR backtrace method is not supported on this platform. DWARF backtrace method will be used.\n",
            "Collecting data...\n",
            "Building filepaths\n",
            "Done\n",
            "Loading png files\n",
            "Done\n",
            "Processing images grayscale\n",
            "Done\n",
            "Processing image convolution\n",
            "Done\n",
            "Processing images minimum pooling\n",
            "Done\n",
            "Processing image maximum pooling\n",
            "Done\n",
            "Writing convolved png to disk\n",
            "Image 1 of 10\n",
            "Image 2 of 10\n",
            "Image 3 of 10\n",
            "Image 4 of 10\n",
            "Done\n",
            "Writing min pooling png to disk\n",
            "Image 1 of 10\n",
            "Image 2 of 10\n",
            "Image 3 of 10\n",
            "Image 4 of 10\n",
            "Done\n",
            "Writing max pooling png to disk\n",
            "Image 1 of 10\n",
            "Image 2 of 10\n",
            "Image 3 of 10\n",
            "Image 4 of 10\n",
            "Done\n",
            "\n",
            "Total program time: 9.839s\n",
            "\n",
            "Processing time (CUDA Kernels): 0.034s\n",
            "Processing events...\n",
            "Capturing symbol files...\n",
            "Saving temporary \"/tmp/nsys-report-faeb-951e-168c-7876.qdstrm\" file to disk...\n",
            "Creating final output files...\n",
            "\n",
            "Processing [==============================================================100%]\n",
            "Saved report file to \"/tmp/nsys-report-faeb-951e-168c-7876.qdrep\"\n",
            "Report file moved to \"/content/report1.qdrep\"\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "!ncu --set full -o ImageFilter_CUDA ./ImageFilter_CUDA"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "9EJil7-KzkAr",
        "outputId": "f4d961bf-9314-4013-dfe3-e2d61cf1bdab"
      },
      "execution_count": null,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "==PROF== Connected to process 1431 (/content/ImageFilter_CUDA)\n",
            "Building filepaths\n",
            "Done\n",
            "Loading png files\n",
            "Done\n",
            "Processing images grayscale\n",
            "==PROF== Profiling \"ConvertImageToGrayCpu\" - 1: 0%....50%....100% - 32 passes\n",
            "==PROF== Profiling \"ConvertImageToGrayCpu\" - 2: 0%....50%....100% - 32 passes\n",
            "==PROF== Profiling \"ConvertImageToGrayCpu\" - 3: 0%....50%....100% - 32 passes\n",
            "==PROF== Profiling \"ConvertImageToGrayCpu\" - 4: 0%....50%....100% - 32 passes\n",
            "Done\n",
            "Processing image convolution\n",
            "==PROF== Profiling \"convolveImage\" - 5: 0%....50%....100% - 32 passes\n",
            "==PROF== Profiling \"convolveImage\" - 6: 0%....50%....100% - 32 passes\n",
            "==PROF== Profiling \"convolveImage\" - 7: 0%....50%....100% - 32 passes\n",
            "==PROF== Profiling \"convolveImage\" - 8: 0%....50%....100% - 32 passes\n",
            "Done\n",
            "Processing images minimum pooling\n",
            "==PROF== Profiling \"minPooling\" - 9: 0%....50%....100% - 32 passes\n",
            "==PROF== Profiling \"minPooling\" - 10: 0%....50%....100% - 32 passes\n",
            "==PROF== Profiling \"minPooling\" - 11: 0%....50%....100% - 32 passes\n",
            "==PROF== Profiling \"minPooling\" - 12: 0%....50%....100% - 32 passes\n",
            "Done\n",
            "Processing image maximum pooling\n",
            "==PROF== Profiling \"maxPooling\" - 13: 0%....50%....100% - 32 passes\n",
            "==PROF== Profiling \"maxPooling\" - 14: 0%....50%....100% - 32 passes\n",
            "==PROF== Profiling \"maxPooling\" - 15: 0%....50%....100% - 32 passes\n",
            "==PROF== Profiling \"maxPooling\" - 16: 0%....50%....100% - 32 passes\n",
            "Done\n",
            "Writing convolved png to disk\n",
            "Image 1 of 10\n",
            "Image 2 of 10\n",
            "Image 3 of 10\n",
            "Image 4 of 10\n",
            "Done\n",
            "Writing min pooling png to disk\n",
            "Image 1 of 10\n",
            "Image 2 of 10\n",
            "Image 3 of 10\n",
            "Image 4 of 10\n",
            "Done\n",
            "Writing max pooling png to disk\n",
            "Image 1 of 10\n",
            "Image 2 of 10\n",
            "Image 3 of 10\n",
            "Image 4 of 10\n",
            "Done\n",
            "\n",
            "Total program time: 29.047s\n",
            "\n",
            "Processing time (CUDA Kernels): 20.003s\n",
            "==PROF== Disconnected from process 1431\n",
            "==PROF== Report: /content/ImageFilter_CUDA.ncu-rep\n"
          ]
        }
      ]
    }
  ],
  "metadata": {
    "accelerator": "GPU",
    "colab": {
      "provenance": []
    },
    "gpuClass": "standard",
    "kernelspec": {
      "display_name": "Python 3",
      "name": "python3"
    },
    "language_info": {
      "name": "python"
    }
  },
  "nbformat": 4,
  "nbformat_minor": 0
}