diff --git a/cuda-blur/Makefile b/cuda-blur/Makefile
new file mode 100644
index 0000000..3f811e3
--- /dev/null
+++ b/cuda-blur/Makefile
@@ -0,0 +1,61 @@
+# Define build directories
+BUILD_DIR := build
+OBJ_DIR := $(BUILD_DIR)/obj
+BIN_DIR := $(BUILD_DIR)/bin
+LIB_DIR := $(BUILD_DIR)/lib
+
+# Include and library paths
+INCLUDES := -Iexternal/lib/opencv/include/opencv4 \
+            -Iinclude \
+            -I/usr/local/cuda/include
+
+LIBS := -Lexternal/lib/opencv/lib \
+        -L/usr/local/cuda/lib64 \
+        -lopencv_core \
+        -lopencv_highgui \
+        -lopencv_imgproc \
+        -lopencv_imgcodecs \
+        -lcudart \
+        -Wl,-rpath,external/lib/opencv/lib
+
+NVCC_FLAGS := -std=c++17 -Xcudafe --diag_suppress=611 $(INCLUDES)
+CXX := g++
+NVCC := nvcc
+
+# Source files
+SRCS := src/main.cu \
+        src/utils.cpp \
+        src/rgb_to_grayscale.cu \
+        src/blur_image.cu
+
+# Object files
+OBJS := $(patsubst src/%.cu,$(OBJ_DIR)/%.o,$(filter %.cu,$(SRCS))) \
+        $(patsubst src/%.cpp,$(OBJ_DIR)/%.o,$(filter %.cpp,$(SRCS)))
+
+# Output binary
+TARGET := $(BIN_DIR)/main
+
+# Default target
+all: $(TARGET)
+
+# Create build directories
+$(OBJ_DIR) $(BIN_DIR) $(LIB_DIR):
+	@mkdir -p $@
+
+# Compile object files
+$(OBJ_DIR)/%.o: src/%.cu | $(OBJ_DIR)
+	$(NVCC) $(NVCC_FLAGS) -c $< -o $@
+
+$(OBJ_DIR)/%.o: src/%.cpp | $(OBJ_DIR)
+	$(NVCC) $(NVCC_FLAGS) -c $< -o $@
+
+# Link the final executable
+$(TARGET): $(OBJS) | $(BIN_DIR)
+	$(CXX) $^ $(INCLUDES) $(LIBS) -o $@
+
+# Cleanup
+clean:
+	rm -rf $(BUILD_DIR)
+
+.PHONY: all clean
+
diff --git a/cuda-blur/README.md b/cuda-blur/README.md
new file mode 100644
index 0000000..29f0639
--- /dev/null
+++ b/cuda-blur/README.md
@@ -0,0 +1,22 @@
+# Instructions
+
+Install opencv
+
+```bash
+./external/install-opencv.sh
+```
+
+Download image
+
+```bash
+./scripts/download-img.sh
+```
+
+Finally compile and run
+
+
+```bash
+TASK=compile_exec ./scripts/compile-run.sh
+
+./build/bin/main img/raw_img.jpeg
+```
\ No newline at end of file
diff --git a/cuda-blur/external/install-opencv.sh b/cuda-blur/external/install-opencv.sh
new file mode 100755
index 0000000..355b31a
--- /dev/null
+++ b/cuda-blur/external/install-opencv.sh
@@ -0,0 +1,28 @@
+#!/bin/bash
+
+THIS_DIR=$(dirname "$(realpath "$0")")
+ROOT_DIR=$(dirname ${THIS_DIR})
+OPENCV_VERSION=4.10.0
+
+
+# library installed in this directory/lib
+LIBDIR=${THIS_DIR}/lib
+
+# download and untar
+wget https://github.com/opencv/opencv/archive/refs/tags/${OPENCV_VERSION}.tar.gz -O ${THIS_DIR}/opencv.tar.gz
+cd ${THIS_DIR} && tar -xzf ${THIS_DIR}/opencv.tar.gz
+
+# build the library
+cd ${THIS_DIR}/opencv-${OPENCV_VERSION}
+mkdir -p build && cd build
+
+cmake -D CMAKE_BUILD_TYPE=Release \
+	-D CMAKE_INSTALL_PREFIX=$LIBDIR/opencv \
+	-D BUILD_EXAMPLES=ON ..
+
+make -j$(nproc)
+make install
+
+# remove temporary files
+cd ${THIS_DIR} && rm -rf opencv-${OPENCV_VERSION}
+cd ${THIS_DIR} && rm -rf opencv.tar.gz
\ No newline at end of file
diff --git a/cuda-blur/include/blur_image.cuh b/cuda-blur/include/blur_image.cuh
new file mode 100644
index 0000000..7afd603
--- /dev/null
+++ b/cuda-blur/include/blur_image.cuh
@@ -0,0 +1,7 @@
+#ifndef BLUR_IMAGE_CUH
+#define BLUR_IMAGE_CUH
+
+__global__ void blur_kernel(unsigned char* input, unsigned char *output, int width, int height);
+void img_to_blur(unsigned char* input, int width, int height);
+
+#endif
diff --git a/cuda-blur/include/rgb_to_grayscale.cuh b/cuda-blur/include/rgb_to_grayscale.cuh
new file mode 100644
index 0000000..60e8ef7
--- /dev/null
+++ b/cuda-blur/include/rgb_to_grayscale.cuh
@@ -0,0 +1,8 @@
+#ifndef RGB_TO_GRAYSCALE_CUH
+#define RGB_TO_GRAYSCALE_CUH
+
+__global__ void rgb_to_grayscale_kernel(const unsigned char* rgb, unsigned char* gray, int width, int height);
+void img_to_gray(unsigned char* input, int width, int height);
+void cpu_grayscale(unsigned char* input, int width, int height);
+
+#endif
\ No newline at end of file
diff --git a/cuda-blur/include/utils.h b/cuda-blur/include/utils.h
new file mode 100644
index 0000000..21e8990
--- /dev/null
+++ b/cuda-blur/include/utils.h
@@ -0,0 +1,6 @@
+#ifndef UTILS_H
+#define UTILS_H
+
+int getBitsPerChannel(int depth);
+
+#endif
diff --git a/cuda-blur/scripts/compile-run.sh b/cuda-blur/scripts/compile-run.sh
new file mode 100755
index 0000000..0f920d0
--- /dev/null
+++ b/cuda-blur/scripts/compile-run.sh
@@ -0,0 +1,92 @@
+#!/bin/bash
+
+THIS_DIR=$(dirname "$(realpath "$0")")
+ROOT_DIR=$(dirname ${THIS_DIR})
+
+recreate_dirs(){
+    # removing build directory
+    echo "Removing ${ROOT_DIR}/build and recreating..."
+    rm -rf ${ROOT_DIR}/build
+    mkdir ${ROOT_DIR}/build
+
+    # creating directories for the build
+    mkdir ${ROOT_DIR}/build/obj
+    mkdir ${ROOT_DIR}/build/bin
+    mkdir ${ROOT_DIR}/build/lib
+}
+
+compile_exec(){
+    recreate_dirs
+
+    # compile to objects
+    echo "Compiling objects for executable..."
+
+    nvcc -std=c++17 \
+        -Xcudafe --diag_suppress=611 \
+        -I${ROOT_DIR}/external/lib/opencv/include/opencv4 \
+        -I${ROOT_DIR}/include \
+        -c ${ROOT_DIR}/src/main.cu \
+        -o ${ROOT_DIR}/build/obj/main.o
+
+    nvcc -std=c++17 \
+        -Xcudafe --diag_suppress=611 \
+        -I${ROOT_DIR}/external/lib/opencv/include/opencv4 \
+        -I${ROOT_DIR}/include \
+        -I/usr/local/cuda/include \
+        -L/usr/local/cuda/lib64 \
+        -c ${ROOT_DIR}/src/utils.cpp \
+        -o ${ROOT_DIR}/build/obj/utils.o
+
+    nvcc -std=c++17 \
+        -Xcudafe --diag_suppress=611 \
+        -I${ROOT_DIR}/external/lib/opencv/include/opencv4 \
+        -I${ROOT_DIR}/include \
+        -I/usr/local/cuda/include \
+        -L/usr/local/cuda/lib64 \
+        -c ${ROOT_DIR}/src/rgb_to_grayscale.cu \
+        -o ${ROOT_DIR}/build/obj/rgb_to_grayscale.o
+
+    nvcc -std=c++17 \
+        -Xcudafe --diag_suppress=611 \
+        -I${ROOT_DIR}/external/lib/opencv/include/opencv4 \
+        -I${ROOT_DIR}/include \
+        -I/usr/local/cuda/include \
+        -L/usr/local/cuda/lib64 \
+        -c ${ROOT_DIR}/src/blur_image.cu \
+        -o ${ROOT_DIR}/build/obj/blur_image.o
+
+
+    # link all the objects
+    g++ ${ROOT_DIR}/build/obj/main.o \
+        ${ROOT_DIR}/build/obj/rgb_to_grayscale.o \
+        ${ROOT_DIR}/build/obj/blur_image.o \
+        ${ROOT_DIR}/build/obj/utils.o \
+        -I${ROOT_DIR}/external/lib/opencv/include/opencv4 \
+        -I/usr/local/cuda/include \
+        -L${ROOT_DIR}/external/lib/opencv/lib \
+        -L/usr/local/cuda/lib64 \
+        -lopencv_core \
+        -lopencv_highgui \
+        -lopencv_imgproc \
+        -lopencv_imgcodecs \
+        -lcudart \
+        -Wl,-rpath,${ROOT_DIR}/external/lib/opencv/lib \
+        -o ${ROOT_DIR}/build/bin/main
+}
+
+
+croak(){
+    echo "[ERROR] $*" > /dev/stderr
+    exit 1
+}
+
+main(){
+    if [[ -z "$TASK" ]]; then
+        croak "No TASK specified."
+    fi
+    echo "[INFO] running $TASK $*"
+    $TASK "$@"
+}
+
+main "$@"
+
diff --git a/cuda-blur/scripts/download-img.sh b/cuda-blur/scripts/download-img.sh
new file mode 100755
index 0000000..8eea05a
--- /dev/null
+++ b/cuda-blur/scripts/download-img.sh
@@ -0,0 +1,7 @@
+#!/bin/bash
+
+THIS_DIR=$(dirname "$(realpath "$0")")
+ROOT_DIR=$(dirname ${THIS_DIR})
+
+mkdir ${ROOT_DIR}/img
+wget "https://upload.wikimedia.org/wikipedia/commons/2/28/20100723_Miyajima_4904.jpg" -O ${ROOT_DIR}/img/raw_img.jpeg
\ No newline at end of file
diff --git a/cuda-blur/src/blur_image.cu b/cuda-blur/src/blur_image.cu
new file mode 100644
index 0000000..af5edf9
--- /dev/null
+++ b/cuda-blur/src/blur_image.cu
@@ -0,0 +1,61 @@
+#include <opencv2/opencv.hpp>
+#include <iostream>
+#include <cuda_runtime.h>
+#include "blur_image.cuh"
+
+#define BLUR_SIZE 10
+
+__global__ void blur_kernel(unsigned char* input, unsigned char *output, int width, int height) {
+    int col = blockIdx.x * blockDim.x + threadIdx.x;
+    int row = blockIdx.y * blockDim.y + threadIdx.y;
+
+    if (col < width && row < height) {
+        int pixVal[3] = {0, 0, 0};
+        int pixels = 0;
+
+        for(int blurRow=-BLUR_SIZE; blurRow<BLUR_SIZE+1; blurRow++){
+            for(int blurCol=-BLUR_SIZE; blurCol<BLUR_SIZE+1; blurCol++){
+                int curRow = row + blurRow;
+                int curCol = col + blurCol;
+
+                if (curRow >= 0 && curRow < height && curCol >= 0 && curCol < width) {
+                    int offset = (curRow * width + curCol) * 3;
+                    pixVal[0] += input[offset];     // R
+                    pixVal[1] += input[offset + 1]; // G
+                    pixVal[2] += input[offset + 2]; // B
+                    pixels++;
+                }
+            }
+        }
+        int out_offset = (row * width + col) * 3;
+        output[out_offset]     = static_cast<unsigned char>(pixVal[0] / pixels);
+        output[out_offset + 1] = static_cast<unsigned char>(pixVal[1] / pixels);
+        output[out_offset + 2] = static_cast<unsigned char>(pixVal[2] / pixels);
+    }
+}
+
+void img_to_blur(unsigned char* input, int width, int height) {
+    unsigned char* d_rgb;
+    unsigned char* d_blur;
+
+    int n_pixels = width * height * 3;
+    unsigned char* h_blur = new unsigned char[n_pixels];
+
+    cudaMalloc(&d_rgb, n_pixels * sizeof(unsigned char));
+    cudaMalloc(&d_blur, n_pixels * sizeof(unsigned char));
+    cudaMemcpy(d_rgb, input, n_pixels * sizeof(unsigned char), cudaMemcpyHostToDevice);
+
+    dim3 threadsPerBlock(16, 16);
+    dim3 blocksPerGrid((height + 15) / 16, (width + 15) / 16);
+    blur_kernel<<<blocksPerGrid, threadsPerBlock>>>(d_rgb, d_blur, height, width);
+    cudaDeviceSynchronize();
+
+    cudaMemcpy(h_blur, d_blur, n_pixels * sizeof(unsigned char), cudaMemcpyDeviceToHost);
+    cudaFree(d_rgb);
+    cudaFree(d_blur);
+
+    cv::Mat blur_image(width, height, CV_8UC3, h_blur);
+    cv::imwrite("blur_output_cuda.jpg", blur_image);
+
+    delete[] h_blur;
+}
\ No newline at end of file
diff --git a/cuda-blur/src/main.cu b/cuda-blur/src/main.cu
new file mode 100644
index 0000000..940df03
--- /dev/null
+++ b/cuda-blur/src/main.cu
@@ -0,0 +1,53 @@
+#include <opencv2/opencv.hpp>
+#include <iostream>
+
+#include "blur_image.cuh"
+#include "rgb_to_grayscale.cuh"
+#include "utils.h"
+
+
+int main(int argc, char* argv[]) {
+    // Check if the user provided an argument
+    if (argc != 2) {
+        std::cerr << "Usage: " << argv[0] << " <image_path>" << std::endl;
+        return -1;
+    }
+
+    // Get the image path from the command-line argument
+    std::string imagePath = argv[1];
+
+    // Read the image
+    cv::Mat image = cv::imread(imagePath);
+
+    // Check if the image was successfully loaded
+    if (image.empty()) {
+        std::cerr << "Error: Unable to load image at " << imagePath << std::endl;
+        return -1;
+    }
+
+    // cast image to 1d array
+    int width = image.size().width;
+    int height = image.size().height;
+    int channels = image.channels();
+    int bitsPerChannel = getBitsPerChannel(image.depth());
+    int num_pixels = width * height;
+
+    // print size of image in bytes
+    std::cout << "Image size: " << image.total() << " bytes" << std::endl;
+    std::cout << "Image size: " << image.total() / 1024 << " KB" << std::endl;
+    std::cout << "Image size: " << image.total() / 1024 / 1024 << " MB" << std::endl;
+
+    // print image shape
+    std::cout << "Image shape: " << width << "x" << height << "x" << channels << std::endl;
+    std::cout << "Number of pixels: " << num_pixels << std::endl;
+    std::cout << "Bits per channel per pixel: " << bitsPerChannel << std::endl;
+    std::cout << "Bytes per channel per pixel: " << bitsPerChannel / 8 << std::endl;
+
+    unsigned char* image_data = image.data;
+
+    cpu_grayscale(image_data, height, width);
+    img_to_gray(image_data, height, width);
+    img_to_blur(image_data, height, width);
+
+    return 0;
+}
\ No newline at end of file
diff --git a/cuda-blur/src/rgb_to_grayscale.cu b/cuda-blur/src/rgb_to_grayscale.cu
new file mode 100644
index 0000000..3373ee1
--- /dev/null
+++ b/cuda-blur/src/rgb_to_grayscale.cu
@@ -0,0 +1,58 @@
+#include <opencv2/opencv.hpp>
+#include <iostream>
+#include <cuda_runtime.h>
+#include "rgb_to_grayscale.cuh"
+
+
+__global__ void rgb_to_grayscale_kernel(const unsigned char* rgb, unsigned char* gray, int width, int height) {
+    int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    int total_pixels = width * height;
+
+    if (idx < total_pixels) {
+        int rgb_idx = idx * 3;
+        unsigned char r = rgb[rgb_idx];
+        unsigned char g = rgb[rgb_idx + 1];
+        unsigned char b = rgb[rgb_idx + 2];
+
+        gray[idx] = static_cast<unsigned char>(0.299f * r + 0.587f * g + 0.114f * b);
+    }
+}
+
+
+void img_to_gray(unsigned char* input, int width, int height) {
+    unsigned char* image_data_gray = new unsigned char[width * height];
+
+    unsigned char* d_rgb;
+    unsigned char* d_gray;
+    unsigned char* h_gray = new unsigned char[width * height];
+
+    cudaMalloc(&d_rgb, width * height * 3);
+    cudaMalloc(&d_gray, width * height);
+
+    cudaMemcpy(d_rgb, input, width * height * 3, cudaMemcpyHostToDevice);
+
+    int threads = 256;
+    int blocks = (width * height + threads - 1) / threads;
+    rgb_to_grayscale_kernel<<<blocks, threads>>>(d_rgb, d_gray, width, height);
+    cudaDeviceSynchronize();
+
+    cudaMemcpy(h_gray, d_gray, width * height, cudaMemcpyDeviceToHost);
+    cudaFree(d_rgb);
+    cudaFree(d_gray);
+
+    cv::Mat gray_image(width, height, CV_8UC1, h_gray);
+    cv::imwrite("gray_output_cuda.jpg", gray_image);
+    delete[] h_gray;    
+}
+
+void cpu_grayscale(unsigned char* input, int width, int height) {
+    unsigned char* image_data_gray = new unsigned char[width * height];
+
+    for (int i = 0; i < width * height; i++) {
+        image_data_gray[i] = 0.299 * input[i * 3] + 0.587 * input[i * 3 + 1] + 0.114 * input[i * 3 + 2];
+    }
+
+    cv::Mat gray_image(width, height, CV_8UC1, image_data_gray);
+    cv::imwrite("gray_output_cpu.jpg", gray_image);
+    delete[] image_data_gray;
+}
\ No newline at end of file
diff --git a/cuda-blur/src/utils.cpp b/cuda-blur/src/utils.cpp
new file mode 100644
index 0000000..bbb35d0
--- /dev/null
+++ b/cuda-blur/src/utils.cpp
@@ -0,0 +1,15 @@
+#include <opencv2/opencv.hpp>
+#include <iostream>
+
+int getBitsPerChannel(int depth) {
+    switch (depth) {
+        case CV_8U:
+        case CV_8S:  return 8;
+        case CV_16U:
+        case CV_16S: return 16;
+        case CV_32S:
+        case CV_32F: return 32;
+        case CV_64F: return 64;
+        default:     return -1; // unknown type
+    }
+}
\ No newline at end of file