Compare commits

..

1 Commits

Author SHA1 Message Date
4d8dd20b06 check if threadIdx.x > 0 && threadIdx.y > 0 2023-06-27 20:38:59 +02:00
3 changed files with 104 additions and 221 deletions

View File

@@ -1,40 +1,30 @@
cmake_minimum_required(VERSION 3.18) cmake_minimum_required(VERSION 3.18)
project(TV_Denoising_CUDA) project(TV_Denoising_CUDA)
# Set tool chain # Find CUDA
set(VCPKG_DIR "E:/programming/vcpkg") enable_language(CUDA)
set(VCPKG_INSTALLED_DIR "${VCPKG_DIR}/installed")
set(CMAKE_TOOLCHAIN_FILE "${VCPKG_DIR}/scripts/buildsystems/vcpkg.cmake" CACHE STRING "Vcpkg toolchain file") # Find OpenCV
set(VCPKG_INSTALLED_DIR "E:/programming/vcpkg/installed")
# Find GTK3 set(OpenCV_DIR "${VCPKG_INSTALLED_DIR}/x64-windows/share/opencv2")
find_package(PkgConfig REQUIRED) find_package(OpenCV REQUIRED)
pkg_check_modules(GTK3 REQUIRED gtk+-3.0)
# Set CUDA flags and properties
# Find OpenCV set(CUDA_SEPARABLE_COMPILATION ON)
set(OpenCV_DIR "${VCPKG_INSTALLED_DIR}/x64-windows/share/opencv2") set(CUDA_PROPAGATE_HOST_FLAGS OFF)
find_package(OpenCV REQUIRED)
# Add the CUDA source files
# Find CUDA file(GLOB CUDA_SOURCE_FILES "src/*.cu")
enable_language(CUDA) set_source_files_properties(${CUDA_SOURCE_FILES} PROPERTIES CUDA_SOURCE_PROPERTY_FORMAT OBJ)
# Set CUDA flags and properties # Add the C++ source files
set(CUDA_SEPARABLE_COMPILATION ON) file(GLOB CPP_SOURCE_FILES "src/*.cpp")
set(CUDA_PROPAGATE_HOST_FLAGS OFF)
# Set the include directories
# Add the CUDA source files include_directories(${OpenCV_DIR} ${CUDA_INCLUDE_DIRS} "include")
file(GLOB CUDA_SOURCE_FILES "src/*.cu")
set_source_files_properties(${CUDA_SOURCE_FILES} PROPERTIES CUDA_SOURCE_PROPERTY_FORMAT OBJ) # Create the executable
add_executable(TV_Denoising_CUDA ${CPP_SOURCE_FILES} ${CUDA_SOURCE_FILES})
# Add the C++ source files
file(GLOB CPP_SOURCE_FILES "src/*.cpp") # Link CUDA libraries
target_link_libraries(TV_Denoising_CUDA ${CUDA_LIBRARIES})
# Set the include directories
include_directories(${OpenCV_DIR} ${CUDA_INCLUDE_DIRS} ${GTK3_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} ${GTK3_LIBRARIES})
target_compile_options(TV_Denoising_CUDA PUBLIC ${GTK3_CFLAGS_OTHER})

View File

@@ -1,176 +1,61 @@
/* /*
Author: Vargha Csongor Csaba Author: Vargha Csongor Csaba
Created: 2023-06-25 10:10:13 Created: 2023-06-25 10:10:13
Description: Description:
This file contains the main function for the TV image denoising cli tool. 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, It reads an image file, denoises it using the TV denoising algorithm,
and saves the denoised image to a file. and saves the denoised image to a file.
You can run it with the following command: You can run it with the following command:
./TV_Denoising_CUDA <input_image> <output_image> <lambda> <iterations> ./TV_Denoising_CUDA <input_image> <output_image> <lambda> <iterations>
where: where:
- <input_image> is the path to the input image file you want to denoise. - <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. - <output_image> is the path to the output denoised image file.
- <lambda> is the regularization parameter for TV denoising (optional, default: 0.02). - <lambda> is the regularization parameter for TV denoising (optional, default: 0.02).
- <iterations> is the number of iterations for TV denoising (optional, default: 10). - <iterations> is the number of iterations for TV denoising (optional, default: 10).
*/ */
#include <iostream> #include <iostream>
#include <gtk/gtk.h> #include <opencv2/opencv.hpp>
#include <opencv2/opencv.hpp> #include "tv_denoising.hpp"
#include "tv_denoising.hpp"
int main(int argc, char** argv)
// Global variables for GUI elements {
GtkLabel* originalImageLabel; // Check if the required arguments are provided
GtkLabel* resultImageLabel; if (argc < 3)
GtkEntry* inputImagePathEntry; {
GtkEntry* outputImagePathEntry; std::cerr << "Usage: ./TV_Denoising_CUDA <input_image> <output_image> [<lambda>] [<iterations>]" << std::endl;
GtkScale* lambdaScale; return 1;
GtkScale* iterationsScale; }
static void updateDenoisedImage(GtkWidget* widget, gpointer data) // Read the input arguments
{ std::string inputImagePath = argv[1];
std::string inputImagePath = gtk_entry_get_text(inputImagePathEntry); std::string outputImagePath = argv[2];
std::string outputImagePath = gtk_entry_get_text(outputImagePathEntry); float lambda = 0.02;
float lambda = gtk_range_get_value(GTK_RANGE(lambdaScale)) / 100.0; int iterations = 10;
int iterations = gtk_range_get_value(GTK_RANGE(iterationsScale));
// Check if optional arguments are provided and update the corresponding variables
// Read the input image if (argc >= 4)
cv::Mat image = cv::imread(inputImagePath, cv::IMREAD_GRAYSCALE); lambda = std::stof(argv[3]);
if (argc >= 5)
// Check if the image was successfully loaded iterations = std::stoi(argv[4]);
if (image.empty())
{ // Read the input image
GtkWidget* dialog = gtk_message_dialog_new(NULL, GTK_DIALOG_MODAL, GTK_MESSAGE_ERROR, GTK_BUTTONS_OK, cv::Mat image = cv::imread(inputImagePath, cv::IMREAD_GRAYSCALE);
"Failed to read the input image.");
gtk_dialog_run(GTK_DIALOG(dialog)); // Check if the image was successfully loaded
gtk_widget_destroy(dialog); if (image.empty())
return; {
} std::cerr << "Failed to read the input image." << std::endl;
return 1;
// Perform TV denoising }
TVDenoising(image, lambda, iterations);
// Perform TV denoising
// Display the denoised image TVDenoising(image, lambda, iterations);
GdkPixbuf* pixbuf = gdk_pixbuf_new_from_data(image.data, GDK_COLORSPACE_RGB, FALSE, 8, image.cols, image.rows,
image.step, NULL, NULL); // Display and save the denoised image
gtk_image_set_from_pixbuf(GTK_IMAGE(resultImageLabel), pixbuf); cv::imshow("Denoised Image", image);
g_object_unref(pixbuf); cv::waitKey(0);
cv::imwrite(outputImagePath, image);
// Save the denoised image
cv::imwrite(outputImagePath, image); return 0;
} }
static void openInputImage(GtkWidget* widget, gpointer data)
{
GtkWidget* dialog = gtk_file_chooser_dialog_new("Select Input Image", GTK_WINDOW(data), GTK_FILE_CHOOSER_ACTION_OPEN,
"Cancel", GTK_RESPONSE_CANCEL, "Open", GTK_RESPONSE_ACCEPT, NULL);
if (gtk_dialog_run(GTK_DIALOG(dialog)) == GTK_RESPONSE_ACCEPT)
{
char* filename = gtk_file_chooser_get_filename(GTK_FILE_CHOOSER(dialog));
gtk_entry_set_text(inputImagePathEntry, filename);
g_free(filename);
}
gtk_widget_destroy(dialog);
}
static void openOutputImage(GtkWidget* widget, gpointer data)
{
GtkWidget* dialog = gtk_file_chooser_dialog_new("Select Output Image", GTK_WINDOW(data), GTK_FILE_CHOOSER_ACTION_SAVE,
"Cancel", GTK_RESPONSE_CANCEL, "Save", GTK_RESPONSE_ACCEPT, NULL);
if (gtk_dialog_run(GTK_DIALOG(dialog)) == GTK_RESPONSE_ACCEPT)
{
char* filename = gtk_file_chooser_get_filename(GTK_FILE_CHOOSER(dialog));
gtk_entry_set_text(outputImagePathEntry, filename);
g_free(filename);
}
gtk_widget_destroy(dialog);
}
int main(int argc, char** argv)
{
// Initialize GTK
gtk_init(&argc, &argv);
// Create the main window
GtkWidget* window = gtk_window_new(GTK_WINDOW_TOPLEVEL);
gtk_window_set_title(GTK_WINDOW(window), "TV Image Denoising");
gtk_window_set_default_size(GTK_WINDOW(window), 800, 600);
g_signal_connect(window, "destroy", G_CALLBACK(gtk_main_quit), NULL);
// Create the main vertical box layout
GtkWidget* vbox = gtk_box_new(GTK_ORIENTATION_VERTICAL, 5);
gtk_container_add(GTK_CONTAINER(window), vbox);
// Create the original image label
originalImageLabel = GTK_LABEL(gtk_label_new("Original Image"));
gtk_box_pack_start(GTK_BOX(vbox), GTK_WIDGET(originalImageLabel), FALSE, FALSE, 0);
// Create the result image label
resultImageLabel = GTK_LABEL(gtk_label_new("Result Image"));
gtk_box_pack_start(GTK_BOX(vbox), GTK_WIDGET(resultImageLabel), FALSE, FALSE, 0);
// Create the input image path input field
GtkWidget* inputImageHBox = gtk_box_new(GTK_ORIENTATION_HORIZONTAL, 5);
gtk_box_pack_start(GTK_BOX(vbox), inputImageHBox, FALSE, FALSE, 0);
GtkWidget* inputImageLabel = gtk_label_new("Input Image:");
gtk_box_pack_start(GTK_BOX(inputImageHBox), inputImageLabel, FALSE, FALSE, 0);
inputImagePathEntry = GTK_ENTRY(gtk_entry_new());
gtk_box_pack_start(GTK_BOX(inputImageHBox), GTK_WIDGET(inputImagePathEntry), TRUE, TRUE, 0);
GtkWidget* inputBrowseButton = gtk_button_new_with_label("Browse");
g_signal_connect(inputBrowseButton, "clicked", G_CALLBACK(openInputImage), window);
gtk_box_pack_start(GTK_BOX(inputImageHBox), inputBrowseButton, FALSE, FALSE, 0);
// Create the output image path input field
GtkWidget* outputImageHBox = gtk_box_new(GTK_ORIENTATION_HORIZONTAL, 5);
gtk_box_pack_start(GTK_BOX(vbox), outputImageHBox, FALSE, FALSE, 0);
GtkWidget* outputImageLabel = gtk_label_new("Output Image:");
gtk_box_pack_start(GTK_BOX(outputImageHBox), outputImageLabel, FALSE, FALSE, 0);
outputImagePathEntry = GTK_ENTRY(gtk_entry_new());
gtk_box_pack_start(GTK_BOX(outputImageHBox), GTK_WIDGET(outputImagePathEntry), TRUE, TRUE, 0);
GtkWidget* outputBrowseButton = gtk_button_new_with_label("Browse");
g_signal_connect(outputBrowseButton, "clicked", G_CALLBACK(openOutputImage), window);
gtk_box_pack_start(GTK_BOX(outputImageHBox), outputBrowseButton, FALSE, FALSE, 0);
// Create the lambda scale
GtkWidget* lambdaHBox = gtk_box_new(GTK_ORIENTATION_HORIZONTAL, 5);
gtk_box_pack_start(GTK_BOX(vbox), lambdaHBox, FALSE, FALSE, 0);
GtkWidget* lambdaLabel = gtk_label_new("Lambda:");
gtk_box_pack_start(GTK_BOX(lambdaHBox), lambdaLabel, FALSE, FALSE, 0);
lambdaScale = GTK_SCALE(gtk_scale_new_with_range(GTK_ORIENTATION_HORIZONTAL, 0, 100, 1));
gtk_scale_set_value_pos(lambdaScale, GTK_POS_RIGHT);
gtk_scale_set_digits(lambdaScale, 2);
gtk_box_pack_start(GTK_BOX(lambdaHBox), GTK_WIDGET(lambdaScale), TRUE, TRUE, 0);
// Create the iterations scale
GtkWidget* iterationsHBox = gtk_box_new(GTK_ORIENTATION_HORIZONTAL, 5);
gtk_box_pack_start(GTK_BOX(vbox), iterationsHBox, FALSE, FALSE, 0);
GtkWidget* iterationsLabel = gtk_label_new("Iterations:");
gtk_box_pack_start(GTK_BOX(iterationsHBox), iterationsLabel, FALSE, FALSE, 0);
iterationsScale = GTK_SCALE(gtk_scale_new_with_range(GTK_ORIENTATION_HORIZONTAL, 1, 100, 1));
gtk_scale_set_value_pos(iterationsScale, GTK_POS_RIGHT);
gtk_scale_set_digits(iterationsScale, 0);
gtk_box_pack_start(GTK_BOX(iterationsHBox), GTK_WIDGET(iterationsScale), TRUE, TRUE, 0);
// Create the denoise button
GtkWidget* denoiseButton = gtk_button_new_with_label("Denoise");
g_signal_connect(denoiseButton, "clicked", G_CALLBACK(updateDenoisedImage), NULL);
gtk_box_pack_start(GTK_BOX(vbox), denoiseButton, FALSE, FALSE, 0);
// Show all widgets
gtk_widget_show_all(window);
// Run the GTK main loop
gtk_main();
return 0;
}

View File

@@ -21,7 +21,7 @@ __global__ void tvDenoisingKernel(float* image, int width, int height, float lam
// Perform TV denoising iteratively // Perform TV denoising iteratively
for (int iteration = 0; iteration < maxIterations; ++iteration) for (int iteration = 0; iteration < maxIterations; ++iteration)
{ {
// Calculate the gradients using central differences // Calculate the gradients using central differences
gradientX[threadIdx.x][threadIdx.y] = image[index + 1] - image[index - 1]; gradientX[threadIdx.x][threadIdx.y] = image[index + 1] - image[index - 1];
gradientY[threadIdx.x][threadIdx.y] = image[index + width] - image[index - width]; gradientY[threadIdx.x][threadIdx.y] = image[index + width] - image[index - width];
@@ -29,10 +29,12 @@ __global__ void tvDenoisingKernel(float* image, int width, int height, float lam
__syncthreads(); __syncthreads();
// Apply TV denoising update rule // Apply TV denoising update rule
updatedImage[threadIdx.x][threadIdx.y] = image[index] + lambda * ( if (threadIdx.x > 0 && threadIdx.y > 0) {
gradientX[threadIdx.x][threadIdx.y] - gradientX[threadIdx.x - 1][threadIdx.y] + updatedImage[threadIdx.x][threadIdx.y] = image[index] + lambda * (
gradientY[threadIdx.x][threadIdx.y] - gradientY[threadIdx.x][threadIdx.y - 1] 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 // Update the global image array with the updated pixel values
image[index] = updatedImage[threadIdx.x][threadIdx.y]; image[index] = updatedImage[threadIdx.x][threadIdx.y];
@@ -65,6 +67,12 @@ extern "C" void TVDenoising(cv::Mat& image, float lambda, int maxIterations)
// Invoke the TV denoising kernel // Invoke the TV denoising kernel
tvDenoisingKernel<<<gridSize, blockSize>>>(d_image, width, height, lambda, maxIterations); tvDenoisingKernel<<<gridSize, blockSize>>>(d_image, width, height, lambda, maxIterations);
// Check for errors during kernel launch
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess)
{
printf("Error: %s\n", cudaGetErrorString(err));
}
// Copy the denoised image data back from device to host // Copy the denoised image data back from device to host
cudaMemcpy(floatImage.ptr<float>(0), d_image, width * height * sizeof(float), cudaMemcpyDeviceToHost); cudaMemcpy(floatImage.ptr<float>(0), d_image, width * height * sizeof(float), cudaMemcpyDeviceToHost);