Impl OpenCL Dithering with Blue-Noise

May need some cleanup (like adding comments).
This commit is contained in:
Stephen Seo 2021-11-26 23:29:48 +09:00
parent 165fa65cc1
commit 732c99403a
5 changed files with 374 additions and 1 deletions

View file

@ -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_;
}

View file

@ -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_;

View file

@ -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;
}

View file

@ -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,

View file

@ -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();