From 732c99403a11398ad4ee0f0700dd649f81e4b5a8 Mon Sep 17 00:00:00 2001 From: Stephen Seo Date: Fri, 26 Nov 2021 23:29:48 +0900 Subject: [PATCH] Impl OpenCL Dithering with Blue-Noise May need some cleanup (like adding comments). --- src/image.cc | 231 +++++++++++++++++++++++++++++++++++++++++++ src/image.h | 32 ++++++ src/main.cc | 21 ++++ src/opencl_handle.cc | 77 +++++++++++++++ src/opencl_handle.h | 14 ++- 5 files changed, 374 insertions(+), 1 deletion(-) diff --git a/src/image.cc b/src/image.cc index 6ee3abf..b1e0256 100644 --- a/src/image.cc +++ b/src/image.cc @@ -3,11 +3,14 @@ #include #include #include +#include #include #include #include +const char *Image::opencl_kernel_ = nullptr; + Image::Image() : data_(), width_(0), height_(0), is_grayscale_(true) {} Image::Image(const char *filename) : Image(std::string(filename)) {} @@ -657,3 +660,231 @@ void Image::DecodePPM(const std::string &filename) { << filename << '"' << std::endl; } } + +uint8_t Image::ColorToGray(uint8_t red, uint8_t green, uint8_t blue) { + // values taken from Wikipedia article about conversion of color to grayscale + double y_linear = 0.2126 * (red / 255.0) + 0.7152 * (green / 255.0) + + 0.0722 * (blue / 255.0); + + if (y_linear <= 0.0031308) { + return std::round((12.92 * y_linear) * 255.0); + } else { + return std::round((1.055 * std::pow(y_linear, 1 / 2.4) - 0.055) * 255.0); + } +} + +std::unique_ptr Image::ToGrayscale() const { + if (IsGrayscale()) { + return std::unique_ptr(new Image(*this)); + } + + std::unique_ptr grayscale_image = std::unique_ptr(new Image{}); + grayscale_image->width_ = this->width_; + grayscale_image->height_ = this->height_; + grayscale_image->data_.resize(width_ * height_); + + for (unsigned int i = 0; i < width_ * height_; ++i) { + grayscale_image->data_.at(i) = + ColorToGray(this->data_.at(i * 4), this->data_.at(i * 4 + 1), + this->data_.at(i * 4 + 2)); + } + + return grayscale_image; +} + +std::unique_ptr Image::ToDitheredWithBlueNoise(Image *blue_noise) { + if (!blue_noise->IsGrayscale()) { + std::cout << "ERROR ToDitheredWithBlueNoise: blue_noise is not grayscale" + << std::endl; + return {}; + } + + auto grayscale_image = ToGrayscale(); + if (!grayscale_image) { + std::cout << "ERROR ToDitheredWithBlueNoise: Failed to get grayscale Image" + << std::endl; + return {}; + } + auto opencl_handle = GetOpenCLHandle(); + if (!opencl_handle) { + std::cout << "ERROR ToDitheredWithBlueNoise: Failed to get OpenCLHandle" + << std::endl; + return {}; + } + + // set up kernel and buffers + auto kid = + opencl_handle->CreateKernelFromSource(GetDitheringKernel(), "Dither"); + if (kid == 0) { + std::cout << "ERROR ToDitheredWithBlueNoise: Failed to create OpenCL Kernel" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + + auto input_buffer_id = opencl_handle->CreateKernelBuffer( + kid, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + grayscale_image->data_.size(), grayscale_image->data_.data()); + if (input_buffer_id == 0) { + std::cout << "ERROR ToDitheredWithBlueNoise: Failed to set input buffer" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + + auto output_buffer_id = opencl_handle->CreateKernelBuffer( + kid, CL_MEM_WRITE_ONLY, grayscale_image->data_.size(), nullptr); + if (output_buffer_id == 0) { + std::cout << "ERROR ToDitheredWithBlueNoise: Failed to set output buffer" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + + auto blue_noise_buffer_id = opencl_handle->CreateKernelBuffer( + kid, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, blue_noise->data_.size(), + (void *)blue_noise->data_.data()); + if (blue_noise_buffer_id == 0) { + std::cout + << "ERROR ToDitheredWithBlueNoise: Failed to set blue-noise buffer" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + + // assign buffers/data to kernel parameters + if (!opencl_handle->AssignKernelBuffer(kid, 0, input_buffer_id)) { + std::cout << "ERROR ToDitheredWithBlueNoise: Failed to set parameter 0" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + if (!opencl_handle->AssignKernelBuffer(kid, 1, blue_noise_buffer_id)) { + std::cout << "ERROR ToDitheredWithBlueNoise: Failed to set parameter 1" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + if (!opencl_handle->AssignKernelBuffer(kid, 2, output_buffer_id)) { + std::cout << "ERROR ToDitheredWithBlueNoise: Failed to set parameter 2" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + unsigned int width = grayscale_image->GetWidth(); + if (!opencl_handle->AssignKernelArgument(kid, 3, sizeof(unsigned int), + &width)) { + std::cout << "ERROR ToDitheredWithBlueNoise: Failed to set parameter 3" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + unsigned int height = grayscale_image->GetHeight(); + if (!opencl_handle->AssignKernelArgument(kid, 4, sizeof(unsigned int), + &height)) { + std::cout << "ERROR ToDitheredWithBlueNoise: Failed to set parameter 4" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + unsigned int blue_noise_size = blue_noise->GetWidth(); + if (!opencl_handle->AssignKernelArgument(kid, 5, sizeof(unsigned int), + &blue_noise_size)) { + std::cout << "ERROR ToDitheredWithBlueNoise: Failed to set parameter 5" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + std::srand(std::time(nullptr)); + unsigned int blue_noise_offset = std::rand() % blue_noise_size; + if (!opencl_handle->AssignKernelArgument(kid, 6, sizeof(unsigned int), + &blue_noise_offset)) { + std::cout << "ERROR ToDitheredWithBlueNoise: Failed to set parameter 6" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + + // auto global_work_sizes = opencl_handle->GetGlobalWorkSize(kid); + auto work_group_size = opencl_handle->GetWorkGroupSize(kid); + std::cout << "Got work_group_size == " << work_group_size << std::endl; + + // auto max_work_group_size = opencl_handle->GetDeviceMaxWorkGroupSize(); + // std::cout << "Got max_work_group_size == " << max_work_group_size + // << std::endl; + + std::size_t work_group_size_0 = std::sqrt(work_group_size); + std::size_t work_group_size_1 = work_group_size_0; + + while (work_group_size_0 > 1 && width % work_group_size_0 != 0) { + --work_group_size_0; + } + while (work_group_size_1 > 1 && height % work_group_size_1 != 0) { + --work_group_size_1; + } + + std::cout << "Using WIDTHxHEIGHT: " << width << "x" << height + << " with work_group_sizes: " << work_group_size_0 << "x" + << work_group_size_1 << std::endl; + + if (!opencl_handle->ExecuteKernel2D(kid, width, height, work_group_size_0, + work_group_size_1, true)) { + std::cout << "ERROR ToDitheredWithBlueNoise: Failed to execute Kernel" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + + if (!opencl_handle->GetBufferData(kid, output_buffer_id, + grayscale_image->GetSize(), + grayscale_image->data_.data())) { + std::cout + << "ERROR ToDitheredWithBlueNoise: Failed to get output buffer data" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + + opencl_handle->CleanupAllKernels(); + return grayscale_image; +} + +const char *Image::GetDitheringKernel() { + if (opencl_kernel_ == nullptr) { + opencl_kernel_ = + "unsigned int BN_INDEX(\n" + "unsigned int x,\n" + "unsigned int y,\n" + "unsigned int o,\n" + "unsigned int bn_size) {\n" + "return (o + x + y * bn_size) % (bn_size * bn_size);\n" + "}\n" + "\n" + "__kernel void Dither(\n" + "__global const unsigned char *input,\n" + "__global const unsigned char *blue_noise,\n" + "__global unsigned char *output,\n" + "const unsigned int input_width,\n" + "const unsigned int input_height,\n" + "const unsigned int blue_noise_size,\n" + "const unsigned int blue_noise_offset) {\n" + "unsigned int idx = get_global_id(0);\n" + "unsigned int idy = get_global_id(1);\n" + "unsigned int b_i = BN_INDEX(idx, idy, blue_noise_offset, " + "blue_noise_size);\n" + "unsigned int input_index = idx + idy * input_width;\n" + "output[input_index] = input[input_index] > blue_noise[b_i] ? 255 : " + "0;\n" + "}\n"; + } + + return opencl_kernel_; +} + +OpenCLHandle::Ptr Image::GetOpenCLHandle() { + if (!opencl_handle_) { + opencl_handle_ = OpenCLContext::GetHandle(); + } + + return opencl_handle_; +} diff --git a/src/image.h b/src/image.h index ed4701e..ac4fd87 100644 --- a/src/image.h +++ b/src/image.h @@ -2,9 +2,12 @@ #define IGPUP_DITHERING_PROJECT_IMAGE_H_ #include +#include #include #include +#include "opencl_handle.h" + class Image { public: Image(); @@ -82,7 +85,36 @@ class Image { /// Same as SaveAsPPM() bool SaveAsPPM(const char *filename, bool overwrite, bool packed = true); + /// Converts rgb to gray with luminance-preserving algorithm + static uint8_t ColorToGray(uint8_t red, uint8_t green, uint8_t blue); + + /*! + * \brief Returns a grayscale version of the Image. + * + * Using std::optional would be ideal here, but this program is aiming for + * compatibility up to C++11, and std::optional was made available in C++17. + * + * \return A std::unique_ptr holding an Image on success, empty otherwise. + */ + std::unique_ptr ToGrayscale() const; + + /*! + * \brief Returns a grayscaled and dithered version of the current Image. + * + * \return A std::unique_ptr holding an Image on success, empty otherwise. + */ + std::unique_ptr ToDitheredWithBlueNoise(Image *blue_noise); + + /// Returns the Dithering Kernel function as a C string + static const char *GetDitheringKernel(); + + /// Returns the OpenCLHandle::Ptr instance + OpenCLHandle::Ptr GetOpenCLHandle(); + private: + static const char *opencl_kernel_; + OpenCLHandle::Ptr opencl_handle_; + /// Internally holds rgba std::vector data_; unsigned int width_; unsigned int height_; diff --git a/src/main.cc b/src/main.cc index c424458..b3c25cc 100644 --- a/src/main.cc +++ b/src/main.cc @@ -1,8 +1,29 @@ +#include + #include "image.h" int main(int argc, char **argv) { // Image image("testin.ppm"); // image.SaveAsPNG("testout.png", true); + Image input("input.png"); + if (!input.IsValid()) { + std::cout << "ERROR: input.png is invalid" << std::endl; + return 1; + } + + Image bluenoise("bluenoise.png"); + if (!bluenoise.IsValid()) { + std::cout << "ERROR: bluenoise.png is invalid" << std::endl; + return 1; + } + + auto output = input.ToDitheredWithBlueNoise(&bluenoise); + if (!output || !output->IsValid()) { + std::cout << "ERROR: output Image is invalid" << std::endl; + return 1; + } + output->SaveAsPNG("output.png", true); + return 0; } diff --git a/src/opencl_handle.cc b/src/opencl_handle.cc index c2a4443..786e64b 100644 --- a/src/opencl_handle.cc +++ b/src/opencl_handle.cc @@ -379,6 +379,31 @@ std::size_t OpenCLContext::OpenCLHandle::GetWorkGroupSize(KernelID kernel_id) { return size; } +std::size_t OpenCLContext::OpenCLHandle::GetDeviceMaxWorkGroupSize() { + if (!IsValid()) { + std::cout << "ERROR: OpenCLContext is not initialized" << std::endl; + return 0; + } + auto context_ptr = opencl_ptr_.lock(); + if (!context_ptr) { + std::cout << "ERROR: OpenCLHandle::GetDeviceMaxWorkGroupSize: " + "OpenCLContext is not initialized" + << std::endl; + return 0; + } + std::size_t value; + cl_int err_num = + clGetDeviceInfo(context_ptr->device_id_, CL_DEVICE_MAX_WORK_GROUP_SIZE, + sizeof(std::size_t), &value, nullptr); + if (err_num != CL_SUCCESS) { + std::cout << "ERROR: OpenCLHandle::GetDeviceMaxWorkGroupSize: " + "Failed to get max work group size" + << std::endl; + } + + return value; +} + bool OpenCLContext::OpenCLHandle::ExecuteKernel(KernelID kernel_id, std::size_t global_work_size, std::size_t local_work_size, @@ -426,6 +451,58 @@ bool OpenCLContext::OpenCLHandle::ExecuteKernel(KernelID kernel_id, return true; } +bool OpenCLContext::OpenCLHandle::ExecuteKernel2D( + KernelID kernel_id, std::size_t global_work_size_0, + std::size_t global_work_size_1, std::size_t local_work_size_0, + std::size_t local_work_size_1, bool is_blocking) { + if (!IsValid()) { + std::cout << "ERROR: OpenCLContext is not initialized" << std::endl; + return false; + } + auto context_ptr = opencl_ptr_.lock(); + if (!context_ptr) { + std::cout << "ERROR: OpenCLHandle::ExecuteKernel2D: OpenCLContext is not " + "initialized" + << std::endl; + return false; + } + + auto kernel_iter = kernels_.find(kernel_id); + if (kernel_iter == kernels_.end()) { + std::cout << "ERROR: OpenCLHandle::ExecuteKernel2D: Invalid kernel_id" + << std::endl; + return false; + } + + std::size_t global_work_size[2] = {global_work_size_0, global_work_size_1}; + std::size_t local_work_size[2] = {local_work_size_0, local_work_size_1}; + cl_event event; + cl_int err_num = clEnqueueNDRangeKernel( + context_ptr->queue_, kernel_iter->second.kernel_, 2, nullptr, + global_work_size, local_work_size, 0, nullptr, &event); + if (err_num != CL_SUCCESS) { + std::cout + << "ERROR: OpenCLHandle::ExecuteKernel2D: Failed to execute kernel" + << " (" << err_num << ")" << std::endl; + return false; + } + + if (is_blocking) { + err_num = clWaitForEvents(1, &event); + if (err_num != CL_SUCCESS) { + std::cout << "WARNING: OpenCLHandle::ExecuteKernel2D: Explicit wait on " + "kernel failed" + << " (" << err_num << ")" << std::endl; + clReleaseEvent(event); + return false; + } + } + + clReleaseEvent(event); + + return true; +} + bool OpenCLContext::OpenCLHandle::GetBufferData(KernelID kernel_id, BufferID buffer_id, std::size_t out_size, diff --git a/src/opencl_handle.h b/src/opencl_handle.h index 6a22c1e..b93d914 100644 --- a/src/opencl_handle.h +++ b/src/opencl_handle.h @@ -141,6 +141,8 @@ class OpenCLContext { */ std::size_t GetWorkGroupSize(KernelID kernel_id); + std::size_t GetDeviceMaxWorkGroupSize(); + /*! * \brief Executes the kernel with the given kernel_id * @@ -149,6 +151,16 @@ class OpenCLContext { bool ExecuteKernel(KernelID kernel_id, std::size_t global_work_size, std::size_t local_work_size, bool is_blocking); + /*! + * \brief Executes the kernel with the given kernel_id + * + * \return true on success + */ + bool ExecuteKernel2D(KernelID kernel_id, std::size_t global_work_size_0, + std::size_t global_work_size_1, + std::size_t local_work_size_0, + std::size_t local_work_size_1, bool is_blocking); + /*! * \brief Copies device memory to data_out * @@ -215,7 +227,7 @@ class OpenCLContext { OpenCLContext &operator=(OpenCLContext &&other) = delete; /// Returns a OpenCLHandle wrapped in a std::shared_ptr - OpenCLHandle::Ptr GetHandle(); + static OpenCLHandle::Ptr GetHandle(); private: OpenCLContext();