Add Cuda kernel, to perform denoising

.gitignore

@@ -27,8 +27,27 @@
 *.a
 *.lib
 
+# Build results
+[Dd]ebug/
+[Dd]ebugPublic/
+[Rr]elease/
+[Rr]eleases/
+x64/
+x86/
+[Ww][Ii][Nn]32/
+[Aa][Rr][Mm]/
+[Aa][Rr][Mm]64/
+bld/
+[Bb]in/
+[Oo]bj/
+[Ll]og/
+[Ll]ogs/
+out/
+build/
+
 # Executables
 *.exe
 *.out
 *.app
 
+.vscode/*

CMakeLists.txt

@@ -0,0 +1,30 @@
+cmake_minimum_required(VERSION 3.18)
+project(TV_Denoising_CUDA)
+
+# Find CUDA
+enable_language(CUDA)
+
+# Find OpenCV
+set(VCPKG_INSTALLED_DIR "E:/programming/vcpkg/installed")
+set(OpenCV_DIR "${VCPKG_INSTALLED_DIR}/x64-windows/share/opencv2")
+find_package(OpenCV REQUIRED)
+
+# Set CUDA flags and properties
+set(CUDA_SEPARABLE_COMPILATION ON)
+set(CUDA_PROPAGATE_HOST_FLAGS OFF)
+
+# Add the CUDA source files
+file(GLOB CUDA_SOURCE_FILES "src/*.cu")
+set_source_files_properties(${CUDA_SOURCE_FILES} PROPERTIES CUDA_SOURCE_PROPERTY_FORMAT OBJ)
+
+# Add the C++ source files
+file(GLOB CPP_SOURCE_FILES "src/*.cpp")
+
+# Set the include directories
+include_directories(${OpenCV_DIR} ${CUDA_INCLUDE_DIRS} "include")
+
+# Create the executable
+add_executable(TV_Denoising_CUDA ${CPP_SOURCE_FILES} ${CUDA_SOURCE_FILES})
+
+# Link CUDA libraries
+target_link_libraries(TV_Denoising_CUDA ${CUDA_LIBRARIES})

README.md

@@ -1,3 +1,30 @@
 # tv-image-denoising
 
-Using the TV algorithm to denoise images with c++ and cuda.
+Using the TV algorithm to denoise images with c++ and cuda.
+
+## Requirements
+    g++
+    vcpkg
+    cmake
+    opencv2
+    cuda
+## Build
+```bash
+# Generate build files
+cmake -B build .
+# Build
+cmake --build build
+```
+## Run
+```bash
+./TV_Denoising_CUDA <input_image> <output_image> <lambda> <iterations>
+```
+
+## Original Exercise:
+7. Image denoising (up to 50+10+20=80%+)
+- [x] Use and implement a total variation image denoising method for the GPU. The input should be a noisy image, processed on the GPU by solving the energy minimization problem, then the output image should be displayed or saved to the disk. (50%) 
+
+- [ ] You can gain extra points for providing a simple GUI that allows the user to tune the parameters of the denoising algorithm, and immediately display the results! (+10%)
+
+- [ ] Additional extra points can be achieved by also submitting a comparative evaluation of yours and other denoising techniques, applied to various inputs. I.e., use metrics such as SSIM, NCC, MAD, etc. to compare results to ground truth images. (+20%)
+(Plus points: The more techniques you compare or the better quality that your figures,charts & diagrams are, you may be rewarded by further extra points.)

include/tv_denoising.hpp

@@ -0,0 +1,12 @@
+/*
+Author: Vargha Csongor Csaba
+Created: 2023-06-25 10:23:33
+*/
+#ifndef TV_DENOISING_H
+#define TV_DENOISING_H
+
+#include <opencv2/opencv.hpp>
+
+extern "C" void TVDenoising(cv::Mat& image, float lambda, int maxIterations);
+
+#endif // TV_DENOISING_H

src/main.cpp

@@ -0,0 +1,61 @@
+/*
+Author: Vargha Csongor Csaba
+Created: 2023-06-25 10:10:13
+
+Description:
+    This file contains the main function for the TV image denoising cli tool.
+    It reads an image file, denoises it using the TV denoising algorithm,
+    and saves the denoised image to a file.
+    You can run it with the following command:
+        ./TV_Denoising_CUDA <input_image> <output_image> <lambda> <iterations>
+    where:
+        - <input_image> is the path to the input image file you want to denoise.
+        - <output_image> is the path to the output denoised image file.
+        - <lambda> is the regularization parameter for TV denoising (optional, default: 0.02).
+        - <iterations> is the number of iterations for TV denoising (optional, default: 10).
+*/
+#include <iostream>
+#include <opencv2/opencv.hpp>
+#include "tv_denoising.hpp"
+
+int main(int argc, char** argv)
+{
+    // Check if the required arguments are provided
+    if (argc < 3)
+    {
+        std::cerr << "Usage: ./TV_Denoising_CUDA <input_image> <output_image> [<lambda>] [<iterations>]" << std::endl;
+        return 1;
+    }
+
+    // Read the input arguments
+    std::string inputImagePath = argv[1];
+    std::string outputImagePath = argv[2];
+    float lambda = 0.02;
+    int iterations = 10;
+
+    // Check if optional arguments are provided and update the corresponding variables
+    if (argc >= 4)
+        lambda = std::stof(argv[3]);
+    if (argc >= 5)
+        iterations = std::stoi(argv[4]);
+
+    // Read the input image
+    cv::Mat image = cv::imread(inputImagePath, cv::IMREAD_GRAYSCALE);
+
+    // Check if the image was successfully loaded
+    if (image.empty())
+    {
+        std::cerr << "Failed to read the input image." << std::endl;
+        return 1;
+    }
+
+    // Perform TV denoising
+    TVDenoising(image, lambda, iterations);
+
+    // Display and save the denoised image
+    cv::imshow("Denoised Image", image);
+    cv::waitKey(0);
+    cv::imwrite(outputImagePath, image);
+
+    return 0;
+}

src/tv_denoising.cu

@@ -0,0 +1,77 @@
+#include <cuda.h>
+#include <cuda_runtime.h>
+#include <device_launch_parameters.h>
+#include <opencv2/opencv.hpp>
+
+#define BLOCK_SIZE_X 16
+#define BLOCK_SIZE_Y 16
+
+__global__ void tvDenoisingKernel(float* image, int width, int height, float lambda, int maxIterations)
+{
+    // Calculate the global thread index
+    int col = blockIdx.x * blockDim.x + threadIdx.x;
+    int row = blockIdx.y * blockDim.y + threadIdx.y;
+    int index = row * width + col;
+
+     // Declare shared memory arrays
+     __shared__ float gradientX[BLOCK_SIZE_X][BLOCK_SIZE_Y];
+     __shared__ float gradientY[BLOCK_SIZE_X][BLOCK_SIZE_Y];
+     __shared__ float updatedImage[BLOCK_SIZE_X][BLOCK_SIZE_Y];
+ 
+     // Perform TV denoising iteratively
+     for (int iteration = 0; iteration < maxIterations; ++iteration)
+     {
+         // Calculate the gradients using central differences
+        gradientX[threadIdx.x][threadIdx.y] = image[index + 1] - image[index - 1];
+        gradientY[threadIdx.x][threadIdx.y] = image[index + width] - image[index - width];
+
+        // Synchronize threads to ensure all gradient calculations are complete
+        __syncthreads();
+
+        // Apply TV denoising update rule
+        updatedImage[threadIdx.x][threadIdx.y] = image[index] + lambda * (
+            gradientX[threadIdx.x][threadIdx.y] - gradientX[threadIdx.x - 1][threadIdx.y] +
+            gradientY[threadIdx.x][threadIdx.y] - gradientY[threadIdx.x][threadIdx.y - 1]
+        );
+
+        // Update the global image array with the updated pixel values
+        image[index] = updatedImage[threadIdx.x][threadIdx.y];
+
+        // Synchronize threads to ensure all image updates are complete
+        __syncthreads();
+     }
+}
+
+extern "C" void TVDenoising(cv::Mat& image, float lambda, int maxIterations)
+{
+    // Convert the image to float precision
+    cv::Mat floatImage;
+    image.convertTo(floatImage, CV_32F);
+
+    // Get image dimensions
+    int width = image.cols;
+    int height = image.rows;
+
+    // Calculate the number of blocks and threads per block
+    dim3 blockSize(16, 16);
+    dim3 gridSize((width + blockSize.x - 1) / blockSize.x, (height + blockSize.y - 1) / blockSize.y);
+
+    // Allocate GPU memory for the image
+    float* d_image;
+    cudaMalloc(&d_image, width * height * sizeof(float));
+
+    // Copy the image data from host to device
+    cudaMemcpy(d_image, floatImage.ptr<float>(0), width * height * sizeof(float), cudaMemcpyHostToDevice);
+
+    // Invoke the TV denoising kernel
+    tvDenoisingKernel<<<gridSize, blockSize>>>(d_image, width, height, lambda, maxIterations);
+
+    // Copy the denoised image data back from device to host
+    cudaMemcpy(floatImage.ptr<float>(0), d_image, width * height * sizeof(float), cudaMemcpyDeviceToHost);
+
+    // Convert the denoised image back to the original data type
+    floatImage.convertTo(image, image.type());
+
+    // Free the GPU memory
+    cudaFree(d_image);
+}