Impl OpenCL Dithering with Blue-Noise

May need some cleanup (like adding comments).

diff --git a/src/image.cc b/src/image.cc
index 6ee3abf895bae5465c3632ce7f59fb68f8a874bf..b1e02569636bdb3643a8679d712ec8db1b21d2a0 100644 (file)
--- a/src/image.cc
+++ b/src/image.cc
@@ -3,11 +3,14 @@
 #include <array>
 #include <cmath>
 #include <cstdio>
+#include <ctime>
 #include <fstream>
 #include <iostream>
 
 #include <png.h>
 
+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> Image::ToGrayscale() const {
+  if (IsGrayscale()) {
+    return std::unique_ptr<Image>(new Image(*this));
+  }
+
+  std::unique_ptr<Image> grayscale_image = std::unique_ptr<Image>(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> 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 ed4701e524007a64b8ce714abfac0e354e4554cd..ac4fd87c2624aafbf4352450c3c6c20545e45455 100644 (file)
--- a/src/image.h
+++ b/src/image.h
@@ -2,9 +2,12 @@
 #define IGPUP_DITHERING_PROJECT_IMAGE_H_
 
 #include <cstdint>
+#include <memory>
 #include <string>
 #include <vector>
 
+#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<Image> 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<Image> 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<uint8_t> data_;
   unsigned int width_;
   unsigned int height_;

diff --git a/src/main.cc b/src/main.cc
index c4244587b9fa6979e2665c5bf93ce673f16370e3..b3c25cc92a479b71116ffab34ddabdea1c345ddd 100644 (file)
--- a/src/main.cc
+++ b/src/main.cc
@@ -1,8 +1,29 @@
+#include <iostream>
+
 #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 c2a444302d54ddc23cc7a10464381a3ad1fd4499..786e64b8142bfae0554f660151f30e144eda48b3 100644 (file)
--- 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 6a22c1efbcbcc07e624fbe19b33b2d201a2a5787..b93d914cfdb2d3662da9ce3ef2757fa828b43350 100644 (file)
--- 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();