stephenseo/EN605.617.81.FA21_StephenSeo_DitheringProject
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

Refactoring of opencl_handle and image for reuse

Browse source 
Refactoring was done so that the OpenCL kernels wouldn't need to be
recompiled every time an image was dithered, among other efficiency
fixes.
This commit is contained in:
Stephen Seo 2021-12-01 16:02:34 +09:00
parent 11f48592bf
commit 6e5eaac63d
5 changed files with 667 additions and 364 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

510
src/image.cc
View file

@ -8,15 +8,29 @@
#include <fstream>
#include <iostream>
const char *Image::opencl_grayscale_kernel_ = nullptr;
const char *Image::opencl_color_kernel_ = nullptr;
#define IGPUP_PROJECT_GRAYSCALE_KERNEL_NAME_ "GrayscaleDither"
#define IGPUP_PROJECT_COLOR_KERNEL_NAME_ "ColorDither"
const std::array<png_color, 2> Image::dither_bw_palette_ = {
const char *Image::kOpenCLGrayscaleKernel = nullptr;
const char *Image::kOpenCLColorKernel = nullptr;
const std::string Image::kBufferInputName = "DitherBufferInput";
const std::string Image::kBufferOutputName = "DitherBufferOutput";
const std::string Image::kBufferBlueNoiseName = "DitherBufferBlueNoise";
const std::string Image::kBufferBlueNoiseOffsetsName =
"DitherBufferBlueNoiseOffsets";
const std::string Image::kGrayscaleKernelName =
IGPUP_PROJECT_GRAYSCALE_KERNEL_NAME_;
const std::string Image::kColorKernelName = IGPUP_PROJECT_COLOR_KERNEL_NAME_;
const std::string Image::kEmptyString = {};
const std::array<png_color, 2> Image::kDitherBWPalette = {
png_color{0, 0, 0}, // black
png_color{255, 255, 255} // white
};
const std::array<png_color, 8> Image::dither_color_palette_ = {
const std::array<png_color, 8> Image::kDitherColorPalette = {
png_color{0, 0, 0}, // black
png_color{255, 255, 255}, // white
png_color{255, 0, 0}, // red
@ -28,22 +42,28 @@ const std::array<png_color, 8> Image::dither_color_palette_ = {
};
Image::Image()
: data_(),
: blue_noise_offsets_{0, 0, 0},
data_(),
width_(0),
height_(0),
is_grayscale_(true),
is_dithered_grayscale_(false),
is_dithered_color_(false) {}
is_dithered_color_(false),
is_preserving_blue_noise_offsets_(true) {
GenerateBlueNoiseOffsets();
}
Image::Image(const char *filename) : Image(std::string(filename)) {}
Image::Image(const std::string &filename)
: data_(),
: blue_noise_offsets_{0, 0, 0},
data_(),
width_(0),
height_(0),
is_grayscale_(true),
is_dithered_grayscale_(false),
is_dithered_color_(false) {
is_dithered_color_(false),
is_preserving_blue_noise_offsets_(true) {
if (filename.compare(filename.size() - 4, filename.size(), ".png") == 0) {
// filename expected to be .png
std::cout << "INFO: PNG filename extension detected, decoding..."
@ -66,6 +86,8 @@ Image::Image(const std::string &filename)
std::cout << "ERROR: Unknown filename extension" << std::endl;
return;
}
GenerateBlueNoiseOffsets();
}
bool Image::IsValid() const {
@ -145,8 +167,8 @@ bool Image::SaveAsPNG(const std::string &filename, bool overwrite) {
png_set_IHDR(png_ptr, png_info_ptr, width_, height_, 1,
PNG_COLOR_TYPE_PALETTE, PNG_INTERLACE_NONE,
PNG_COMPRESSION_TYPE_DEFAULT, PNG_FILTER_TYPE_DEFAULT);
png_set_PLTE(png_ptr, png_info_ptr, dither_bw_palette_.data(),
dither_bw_palette_.size());
png_set_PLTE(png_ptr, png_info_ptr, kDitherBWPalette.data(),
kDitherBWPalette.size());
} else {
png_set_IHDR(png_ptr, png_info_ptr, width_, height_, 8,
PNG_COLOR_TYPE_GRAY, PNG_INTERLACE_NONE,
@ -157,8 +179,8 @@ bool Image::SaveAsPNG(const std::string &filename, bool overwrite) {
png_set_IHDR(png_ptr, png_info_ptr, width_, height_, 4,
PNG_COLOR_TYPE_PALETTE, PNG_INTERLACE_NONE,
PNG_COMPRESSION_TYPE_DEFAULT, PNG_FILTER_TYPE_DEFAULT);
png_set_PLTE(png_ptr, png_info_ptr, dither_color_palette_.data(),
dither_color_palette_.size());
png_set_PLTE(png_ptr, png_info_ptr, kDitherColorPalette.data(),
kDitherColorPalette.size());
} else {
png_set_IHDR(png_ptr, png_info_ptr, width_, height_, 8,
PNG_COLOR_TYPE_RGB_ALPHA, PNG_INTERLACE_NONE,
@ -378,126 +400,158 @@ std::unique_ptr<Image> Image::ToGrayscaleDitheredWithBlueNoise(
return {};
}
// first check if existing kernel/buffers can be used
if (opencl_handle->HasKernel(kGrayscaleKernelName) &&
!VerifyOpenCLBuffers(
kGrayscaleKernelName,
{kBufferInputName, kBufferOutputName, kBufferBlueNoiseName},
grayscale_image.get(), blue_noise)) {
opencl_handle->CleanupKernel(kGrayscaleKernelName);
}
// set up kernel and buffers
auto kid = opencl_handle->CreateKernelFromSource(
GetGrayscaleDitheringKernel(), "Dither");
if (kid == 0) {
std::cout << "ERROR ToGrayscaleDitheredWithBlueNoise: Failed to create "
"OpenCL Kernel"
const std::string &grayscale_kernel_name = GetGrayscaleKernelName();
if (grayscale_kernel_name.empty() ||
!opencl_handle->HasKernel(grayscale_kernel_name)) {
std::cout << "ERROR ToGrayscaleDitheredWithBlueNoise: Failed to init 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 ToGrayscaleDitheredWithBlueNoise: Failed to set input buffer"
<< std::endl;
opencl_handle->CleanupAllKernels();
if (!opencl_handle->HasBuffer(grayscale_kernel_name, kBufferInputName)) {
if (!opencl_handle->CreateKernelBuffer(
grayscale_kernel_name, CL_MEM_READ_ONLY,
grayscale_image->data_.size(), nullptr, kBufferInputName)) {
std::cout << "ERROR ToGrayscaleDitheredWithBlueNoise: Failed to alloc "
"input buffer"
<< std::endl;
opencl_handle->CleanupKernel(grayscale_kernel_name);
return {};
}
}
if (!opencl_handle->SetKernelBufferData(
grayscale_kernel_name, kBufferInputName,
grayscale_image->data_.size(), grayscale_image->data_.data())) {
std::cout << "ERROR ToGrayscaleDitheredWithBlueNoise: Failed to init "
"input buffer"
<< std::endl;
opencl_handle->CleanupKernel(grayscale_kernel_name);
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 ToGrayscaleDitheredWithBlueNoise: Failed to set output buffer"
<< std::endl;
opencl_handle->CleanupAllKernels();
return {};
if (!opencl_handle->HasBuffer(grayscale_kernel_name, kBufferOutputName)) {
if (!opencl_handle->CreateKernelBuffer(
grayscale_kernel_name, CL_MEM_WRITE_ONLY,
grayscale_image->data_.size(), nullptr, kBufferOutputName)) {
std::cout << "ERROR ToGrayscaleDitheredWithBlueNoise: Failed to set "
"output buffer"
<< std::endl;
opencl_handle->CleanupKernel(grayscale_kernel_name);
return {};
}
}
auto blue_noise_buffer_id = opencl_handle->CreateKernelBuffer(
kid, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, blue_noise->data_.size(),
blue_noise->data_.data());
if (blue_noise_buffer_id == 0) {
std::cout << "ERROR ToGrayscaleDitheredWithBlueNoise: Failed to set "
if (!opencl_handle->HasBuffer(grayscale_kernel_name, kBufferBlueNoiseName)) {
if (!opencl_handle->CreateKernelBuffer(
grayscale_kernel_name, CL_MEM_READ_ONLY, blue_noise->data_.size(),
nullptr, kBufferBlueNoiseName)) {
std::cout << "ERROR ToGrayscaleDitheredWithBlueNoise: Failed to alloc "
"blue-noise buffer"
<< std::endl;
opencl_handle->CleanupKernel(grayscale_kernel_name);
return {};
}
}
if (!opencl_handle->SetKernelBufferData(
grayscale_kernel_name, kBufferBlueNoiseName, blue_noise->data_.size(),
blue_noise->data_.data())) {
std::cout << "ERROR ToGrayscaleDitheredWithBlueNoise: Failed to init "
"blue-noise buffer"
<< std::endl;
opencl_handle->CleanupAllKernels();
opencl_handle->CleanupKernel(grayscale_kernel_name);
return {};
}
// assign buffers/data to kernel parameters
if (!opencl_handle->AssignKernelBuffer(kid, 0, input_buffer_id)) {
if (!opencl_handle->AssignKernelBuffer(grayscale_kernel_name, 0,
kBufferInputName)) {
std::cout
<< "ERROR ToGrayscaleDitheredWithBlueNoise: Failed to set parameter 0"
<< std::endl;
opencl_handle->CleanupAllKernels();
opencl_handle->CleanupKernel(grayscale_kernel_name);
return {};
}
if (!opencl_handle->AssignKernelBuffer(kid, 1, blue_noise_buffer_id)) {
if (!opencl_handle->AssignKernelBuffer(grayscale_kernel_name, 1,
kBufferBlueNoiseName)) {
std::cout
<< "ERROR ToGrayscaleDitheredWithBlueNoise: Failed to set parameter 1"
<< std::endl;
opencl_handle->CleanupAllKernels();
opencl_handle->CleanupKernel(grayscale_kernel_name);
return {};
}
if (!opencl_handle->AssignKernelBuffer(kid, 2, output_buffer_id)) {
if (!opencl_handle->AssignKernelBuffer(grayscale_kernel_name, 2,
kBufferOutputName)) {
std::cout
<< "ERROR ToGrayscaleDitheredWithBlueNoise: Failed to set parameter 2"
<< std::endl;
opencl_handle->CleanupAllKernels();
opencl_handle->CleanupKernel(grayscale_kernel_name);
return {};
}
unsigned int width = grayscale_image->GetWidth();
if (!opencl_handle->AssignKernelArgument(kid, 3, sizeof(unsigned int),
&width)) {
if (!opencl_handle->AssignKernelArgument(grayscale_kernel_name, 3,
sizeof(unsigned int), &width)) {
std::cout
<< "ERROR ToGrayscaleDitheredWithBlueNoise: Failed to set parameter 3"
<< std::endl;
opencl_handle->CleanupAllKernels();
opencl_handle->CleanupKernel(grayscale_kernel_name);
return {};
}
unsigned int height = grayscale_image->GetHeight();
if (!opencl_handle->AssignKernelArgument(kid, 4, sizeof(unsigned int),
&height)) {
if (!opencl_handle->AssignKernelArgument(grayscale_kernel_name, 4,
sizeof(unsigned int), &height)) {
std::cout
<< "ERROR ToGrayscaleDitheredWithBlueNoise: Failed to set parameter 4"
<< std::endl;
opencl_handle->CleanupAllKernels();
opencl_handle->CleanupKernel(grayscale_kernel_name);
return {};
}
unsigned int blue_noise_width = blue_noise->GetWidth();
if (!opencl_handle->AssignKernelArgument(kid, 5, sizeof(unsigned int),
&blue_noise_width)) {
if (!opencl_handle->AssignKernelArgument(
grayscale_kernel_name, 5, sizeof(unsigned int), &blue_noise_width)) {
std::cout
<< "ERROR ToGrayscaleDitheredWithBlueNoise: Failed to set parameter 5"
<< std::endl;
opencl_handle->CleanupAllKernels();
opencl_handle->CleanupKernel(grayscale_kernel_name);
return {};
}
unsigned int blue_noise_height = blue_noise->GetHeight();
if (!opencl_handle->AssignKernelArgument(kid, 6, sizeof(unsigned int),
&blue_noise_height)) {
if (!opencl_handle->AssignKernelArgument(
grayscale_kernel_name, 6, sizeof(unsigned int), &blue_noise_height)) {
std::cout
<< "ERROR ToGrayscaleDitheredWithBlueNoise: Failed to set parameter 6"
<< std::endl;
opencl_handle->CleanupAllKernels();
opencl_handle->CleanupKernel(grayscale_kernel_name);
return {};
}
std::srand(std::time(nullptr));
unsigned int blue_noise_offset =
std::rand() % (blue_noise_width * blue_noise_height);
if (!opencl_handle->AssignKernelArgument(kid, 7, sizeof(unsigned int),
&blue_noise_offset)) {
if (!is_preserving_blue_noise_offsets_) {
GenerateBlueNoiseOffsets();
}
if (!opencl_handle->AssignKernelArgument(grayscale_kernel_name, 7,
sizeof(unsigned int),
&blue_noise_offsets_.at(0))) {
std::cout
<< "ERROR ToGrayscaleDitheredWithBlueNoise: Failed to set parameter 7"
<< std::endl;
opencl_handle->CleanupAllKernels();
opencl_handle->CleanupKernel(grayscale_kernel_name);
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;
auto work_group_size = opencl_handle->GetWorkGroupSize(grayscale_kernel_name);
// DEBUG
// std::cout << "Got work_group_size == " << 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;
@ -509,30 +563,31 @@ std::unique_ptr<Image> Image::ToGrayscaleDitheredWithBlueNoise(
--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;
// DEBUG
// 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)) {
if (!opencl_handle->ExecuteKernel2D(grayscale_kernel_name, width, height,
work_group_size_0, work_group_size_1,
true)) {
std::cout
<< "ERROR ToGrayscaleDitheredWithBlueNoise: Failed to execute Kernel"
<< std::endl;
opencl_handle->CleanupAllKernels();
opencl_handle->CleanupKernel(grayscale_kernel_name);
return {};
}
if (!opencl_handle->GetBufferData(kid, output_buffer_id,
if (!opencl_handle->GetBufferData(grayscale_kernel_name, kBufferOutputName,
grayscale_image->GetSize(),
grayscale_image->data_.data())) {
std::cout << "ERROR ToGrayscaleDitheredWithBlueNoise: Failed to get output "
"buffer data"
<< std::endl;
opencl_handle->CleanupAllKernels();
opencl_handle->CleanupKernel(grayscale_kernel_name);
return {};
}
opencl_handle->CleanupAllKernels();
return grayscale_image;
}
@ -559,127 +614,178 @@ std::unique_ptr<Image> Image::ToColorDitheredWithBlueNoise(Image *blue_noise) {
return {};
}
// first check if existing kernel/buffers can be used
if (opencl_handle->HasKernel(kColorKernelName) &&
!VerifyOpenCLBuffers(
kColorKernelName,
{kBufferInputName, kBufferOutputName, kBufferBlueNoiseName}, this,
blue_noise)) {
opencl_handle->CleanupKernel(kColorKernelName);
}
// set up kernel and buffers
auto kid = opencl_handle->CreateKernelFromSource(GetColorDitheringKernel(),
"ColorDither");
if (kid == 0) {
std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to create "
const std::string &color_kernel_name = GetColorKernelName();
if (color_kernel_name.empty() ||
!opencl_handle->HasKernel(color_kernel_name)) {
std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to init "
"OpenCL Kernel"
<< std::endl;
opencl_handle->CleanupAllKernels();
opencl_handle->CleanupKernel(color_kernel_name);
return {};
}
auto input_buffer_id = opencl_handle->CreateKernelBuffer(
kid, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, this->data_.size(),
this->data_.data());
if (input_buffer_id == 0) {
if (!opencl_handle->HasBuffer(color_kernel_name, kBufferInputName)) {
if (!opencl_handle->CreateKernelBuffer(color_kernel_name, CL_MEM_READ_ONLY,
this->data_.size(), nullptr,
kBufferInputName)) {
std::cout
<< "ERROR ToColorDitheredWithBlueNoise: Failed to alloc input buffer"
<< std::endl;
opencl_handle->CleanupKernel(color_kernel_name);
return {};
}
}
if (!opencl_handle->SetKernelBufferData(color_kernel_name, kBufferInputName,
this->data_.size(),
this->data_.data())) {
std::cout
<< "ERROR ToColorDitheredWithBlueNoise: Failed to set input buffer"
<< "ERROR ToColorDitheredWithBlueNoise: Failed to init input buffer"
<< std::endl;
opencl_handle->CleanupAllKernels();
opencl_handle->CleanupKernel(color_kernel_name);
return {};
}
auto output_buffer_id = opencl_handle->CreateKernelBuffer(
kid, CL_MEM_WRITE_ONLY, this->data_.size(), nullptr);
if (output_buffer_id == 0) {
std::cout
<< "ERROR ToColorDitheredWithBlueNoise: Failed to set output buffer"
<< std::endl;
opencl_handle->CleanupAllKernels();
return {};
if (!opencl_handle->HasBuffer(color_kernel_name, kBufferOutputName)) {
if (!opencl_handle->CreateKernelBuffer(color_kernel_name, CL_MEM_WRITE_ONLY,
this->data_.size(), nullptr,
kBufferOutputName)) {
std::cout
<< "ERROR ToColorDitheredWithBlueNoise: Failed to set output buffer"
<< std::endl;
opencl_handle->CleanupKernel(color_kernel_name);
return {};
}
}
auto blue_noise_buffer_id = opencl_handle->CreateKernelBuffer(
kid, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, blue_noise->data_.size(),
blue_noise->data_.data());
if (blue_noise_buffer_id == 0) {
std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to set "
if (!opencl_handle->HasBuffer(color_kernel_name, kBufferBlueNoiseName)) {
if (!opencl_handle->CreateKernelBuffer(color_kernel_name, CL_MEM_READ_ONLY,
blue_noise->data_.size(), nullptr,
kBufferBlueNoiseName)) {
std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to alloc "
"blue-noise buffer"
<< std::endl;
opencl_handle->CleanupKernel(color_kernel_name);
return {};
}
}
if (!opencl_handle->SetKernelBufferData(
color_kernel_name, kBufferBlueNoiseName, blue_noise->data_.size(),
blue_noise->data_.data())) {
std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to init "
"blue-noise buffer"
<< std::endl;
opencl_handle->CleanupAllKernels();
opencl_handle->CleanupKernel(color_kernel_name);
return {};
}
std::srand(std::time(nullptr));
std::array<unsigned int, 3> blue_noise_offsets = {
std::rand() % blue_noise->GetSize(), std::rand() % blue_noise->GetSize(),
std::rand() % blue_noise->GetSize()};
if (!opencl_handle->HasBuffer(color_kernel_name,
kBufferBlueNoiseOffsetsName)) {
if (!is_preserving_blue_noise_offsets_) {
GenerateBlueNoiseOffsets();
}
while (blue_noise_offsets[0] == blue_noise_offsets[1] ||
blue_noise_offsets[1] == blue_noise_offsets[2] ||
blue_noise_offsets[0] == blue_noise_offsets[2]) {
blue_noise_offsets = {std::rand() % blue_noise->GetSize(),
std::rand() % blue_noise->GetSize(),
std::rand() % blue_noise->GetSize()};
if (!opencl_handle->CreateKernelBuffer(
color_kernel_name, CL_MEM_READ_ONLY,
sizeof(unsigned int) * blue_noise_offsets_.size(), nullptr,
kBufferBlueNoiseOffsetsName)) {
std::cout
<< "ERROR ToColorDitheredWithBlueNoise: Failed to alloc blue-noise "
"offsets buffer"
<< std::endl;
opencl_handle->CleanupKernel(color_kernel_name);
return {};
}
}
auto blue_noise_offsets_buffer_id = opencl_handle->CreateKernelBuffer(
kid, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(unsigned int) * 3,
blue_noise_offsets.data());
if (!opencl_handle->SetKernelBufferData(
color_kernel_name, kBufferBlueNoiseOffsetsName,
blue_noise_offsets_.size() * sizeof(unsigned int),
blue_noise_offsets_.data())) {
std::cout
<< "ERROR ToColorDitheredWithBlueNoise: Failed to init blue-noise "
"offsets buffer"
<< std::endl;
opencl_handle->CleanupKernel(color_kernel_name);
return {};
}
// assign buffers/data to kernel parameters
if (!opencl_handle->AssignKernelBuffer(kid, 0, input_buffer_id)) {
if (!opencl_handle->AssignKernelBuffer(color_kernel_name, 0,
kBufferInputName)) {
std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to set parameter 0"
<< std::endl;
opencl_handle->CleanupAllKernels();
opencl_handle->CleanupKernel(color_kernel_name);
return {};
}
if (!opencl_handle->AssignKernelBuffer(kid, 1, blue_noise_buffer_id)) {
if (!opencl_handle->AssignKernelBuffer(color_kernel_name, 1,
kBufferBlueNoiseName)) {
std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to set parameter 1"
<< std::endl;
opencl_handle->CleanupAllKernels();
opencl_handle->CleanupKernel(color_kernel_name);
return {};
}
if (!opencl_handle->AssignKernelBuffer(kid, 2, output_buffer_id)) {
if (!opencl_handle->AssignKernelBuffer(color_kernel_name, 2,
kBufferOutputName)) {
std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to set parameter 2"
<< std::endl;
opencl_handle->CleanupAllKernels();
opencl_handle->CleanupKernel(color_kernel_name);
return {};
}
unsigned int input_width = this->GetWidth();
if (!opencl_handle->AssignKernelArgument(kid, 3, sizeof(unsigned int),
&input_width)) {
if (!opencl_handle->AssignKernelArgument(
color_kernel_name, 3, sizeof(unsigned int), &input_width)) {
std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to set parameter 3"
<< std::endl;
opencl_handle->CleanupAllKernels();
opencl_handle->CleanupKernel(color_kernel_name);
return {};
}
unsigned int input_height = this->GetHeight();
if (!opencl_handle->AssignKernelArgument(kid, 4, sizeof(unsigned int),
&input_height)) {
if (!opencl_handle->AssignKernelArgument(
color_kernel_name, 4, sizeof(unsigned int), &input_height)) {
std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to set parameter 4"
<< std::endl;
opencl_handle->CleanupAllKernels();
opencl_handle->CleanupKernel(color_kernel_name);
return {};
}
unsigned int blue_noise_width = blue_noise->GetWidth();
if (!opencl_handle->AssignKernelArgument(kid, 5, sizeof(unsigned int),
&blue_noise_width)) {
if (!opencl_handle->AssignKernelArgument(
color_kernel_name, 5, sizeof(unsigned int), &blue_noise_width)) {
std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to set parameter 5"
<< std::endl;
opencl_handle->CleanupAllKernels();
opencl_handle->CleanupKernel(color_kernel_name);
return {};
}
unsigned int blue_noise_height = blue_noise->GetHeight();
if (!opencl_handle->AssignKernelArgument(kid, 6, sizeof(unsigned int),
&blue_noise_height)) {
if (!opencl_handle->AssignKernelArgument(
color_kernel_name, 6, sizeof(unsigned int), &blue_noise_height)) {
std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to set parameter 6"
<< std::endl;
opencl_handle->CleanupAllKernels();
opencl_handle->CleanupKernel(color_kernel_name);
return {};
}
if (!opencl_handle->AssignKernelBuffer(kid, 7,
blue_noise_offsets_buffer_id)) {
if (!opencl_handle->AssignKernelBuffer(color_kernel_name, 7,
kBufferBlueNoiseOffsetsName)) {
std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to set parameter 7"
<< std::endl;
opencl_handle->CleanupAllKernels();
opencl_handle->CleanupKernel(color_kernel_name);
return {};
}
auto work_group_size = opencl_handle->GetWorkGroupSize(kid);
std::cout << "Got work_group_size == " << work_group_size << std::endl;
auto work_group_size = opencl_handle->GetWorkGroupSize(color_kernel_name);
// DEBUG
// std::cout << "Got work_group_size == " << 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;
@ -691,16 +797,17 @@ std::unique_ptr<Image> Image::ToColorDitheredWithBlueNoise(Image *blue_noise) {
--work_group_size_1;
}
std::cout << "Using WIDTHxHEIGHT: " << input_width << "x" << input_height
<< " with work_group_sizes: " << work_group_size_0 << "x"
<< work_group_size_1 << std::endl;
// DEBUG
// std::cout << "Using WIDTHxHEIGHT: " << input_width << "x" << input_height
// << " with work_group_sizes: " << work_group_size_0 << "x"
// << work_group_size_1 << std::endl;
if (!opencl_handle->ExecuteKernel2D(kid, input_width, input_height,
work_group_size_0, work_group_size_1,
true)) {
if (!opencl_handle->ExecuteKernel2D(color_kernel_name, input_width,
input_height, work_group_size_0,
work_group_size_1, true)) {
std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to execute Kernel"
<< std::endl;
opencl_handle->CleanupAllKernels();
opencl_handle->CleanupKernel(color_kernel_name);
return {};
}
@ -708,23 +815,22 @@ std::unique_ptr<Image> Image::ToColorDitheredWithBlueNoise(Image *blue_noise) {
std::unique_ptr<Image>(new Image(*this));
result_image->is_dithered_color_ = true;
if (!opencl_handle->GetBufferData(kid, output_buffer_id,
if (!opencl_handle->GetBufferData(color_kernel_name, kBufferOutputName,
result_image->GetSize(),
result_image->data_.data())) {
std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to get output "
"buffer data"
<< std::endl;
opencl_handle->CleanupAllKernels();
opencl_handle->CleanupKernel(color_kernel_name);
return {};
}
opencl_handle->CleanupAllKernels();
return result_image;
}
const char *Image::GetGrayscaleDitheringKernel() {
if (opencl_grayscale_kernel_ == nullptr) {
opencl_grayscale_kernel_ =
if (kOpenCLGrayscaleKernel == nullptr) {
kOpenCLGrayscaleKernel =
"unsigned int BN_INDEX(\n"
"unsigned int x,\n"
"unsigned int y,\n"
@ -736,7 +842,8 @@ const char *Image::GetGrayscaleDitheringKernel() {
"return offset_x + offset_y * bn_width;\n"
"}\n"
"\n"
"__kernel void Dither(\n"
"__kernel void " IGPUP_PROJECT_GRAYSCALE_KERNEL_NAME_
"(\n"
"__global const unsigned char *input,\n"
"__global const unsigned char *blue_noise,\n"
"__global unsigned char *output,\n"
@ -755,12 +862,12 @@ const char *Image::GetGrayscaleDitheringKernel() {
"}\n";
}
return opencl_grayscale_kernel_;
return kOpenCLGrayscaleKernel;
}
const char *Image::GetColorDitheringKernel() {
if (opencl_color_kernel_ == nullptr) {
opencl_color_kernel_ =
if (kOpenCLColorKernel == nullptr) {
kOpenCLColorKernel =
"unsigned int BN_INDEX(\n"
"unsigned int x,\n"
"unsigned int y,\n"
@ -772,7 +879,8 @@ const char *Image::GetColorDitheringKernel() {
"return offset_x + offset_y * bn_width;\n"
"}\n"
"\n"
"__kernel void ColorDither(\n"
"__kernel void " IGPUP_PROJECT_COLOR_KERNEL_NAME_
"(\n"
"__global const unsigned char *input,\n"
"__global const unsigned char *blue_noise,\n"
"__global unsigned char *output,\n"
@ -804,7 +912,7 @@ const char *Image::GetColorDitheringKernel() {
"}\n";
}
return opencl_color_kernel_;
return kOpenCLColorKernel;
}
OpenCLHandle::Ptr Image::GetOpenCLHandle() {
@ -1269,3 +1377,87 @@ void Image::DecodePPM(const std::string &filename) {
<< filename << '"' << std::endl;
}
}
const std::string &Image::GetGrayscaleKernelName() {
if (!GetOpenCLHandle()) {
return kEmptyString;
} else if (!opencl_handle_->HasKernel(kGrayscaleKernelName)) {
if (!opencl_handle_->CreateKernelFromSource(GetGrayscaleDitheringKernel(),
kGrayscaleKernelName)) {
std::cout << "ERROR: Failed to create " << kGrayscaleKernelName
<< " OpenCL Kernel" << std::endl;
return kEmptyString;
}
}
return kGrayscaleKernelName;
}
const std::string &Image::GetColorKernelName() {
if (!GetOpenCLHandle()) {
return kEmptyString;
} else if (!opencl_handle_->HasKernel(kColorKernelName)) {
if (!opencl_handle_->CreateKernelFromSource(GetColorDitheringKernel(),
kColorKernelName)) {
std::cout << "ERROR: Failed to create " << kColorKernelName
<< " OpenCL Kernel" << std::endl;
return kEmptyString;
}
}
return kColorKernelName;
}
void Image::GenerateBlueNoiseOffsets() {
std::srand(std::time(nullptr));
while (DuplicateBlueNoiseOffsetExists()) {
for (unsigned int i = 0; i < blue_noise_offsets_.size(); ++i) {
blue_noise_offsets_.at(i) = rand() % kBlueNoiseOffsetMax;
}
}
}
bool Image::DuplicateBlueNoiseOffsetExists() const {
for (unsigned int i = 1; i < blue_noise_offsets_.size(); ++i) {
if (blue_noise_offsets_.at(i - 1) == blue_noise_offsets_.at(i)) {
return true;
}
}
if (blue_noise_offsets_.size() > 1 &&
blue_noise_offsets_.at(0) ==
blue_noise_offsets_.at(blue_noise_offsets_.size() - 1)) {
return true;
}
return false;
}
bool Image::VerifyOpenCLBuffers(const std::string &kernel_name,
const std::vector<std::string> &buffer_names,
const Image *input_image,
const Image *blue_noise_image) const {
std::size_t size;
for (auto &buffer_name : buffer_names) {
size = opencl_handle_->GetBufferSize(kernel_name, buffer_name);
if (size == 0) {
return false;
}
if (buffer_name == kBufferInputName || buffer_name == kBufferOutputName) {
if (input_image->is_grayscale_) {
if (size != input_image->width_ * input_image->height_) {
return false;
}
} else {
if (size != input_image->width_ * input_image->height_ * 4) {
return false;
}
}
} else if (buffer_name == kBufferBlueNoiseName) {
if (size != blue_noise_image->width_ * blue_noise_image->height_) {
return false;
}
}
}
return true;
}

29
src/image.h
View file

@ -130,11 +130,20 @@ class Image {
private:
friend class Video;
static const char *opencl_grayscale_kernel_;
static const char *opencl_color_kernel_;
static const std::array<png_color, 2> dither_bw_palette_;
static const std::array<png_color, 8> dither_color_palette_;
static constexpr unsigned int kBlueNoiseOffsetMax = 128;
static const char *kOpenCLGrayscaleKernel;
static const char *kOpenCLColorKernel;
static const std::array<png_color, 2> kDitherBWPalette;
static const std::array<png_color, 8> kDitherColorPalette;
static const std::string kBufferInputName;
static const std::string kBufferOutputName;
static const std::string kBufferBlueNoiseName;
static const std::string kBufferBlueNoiseOffsetsName;
static const std::string kGrayscaleKernelName;
static const std::string kColorKernelName;
static const std::string kEmptyString;
OpenCLHandle::Ptr opencl_handle_;
std::array<unsigned int, 3> blue_noise_offsets_;
/// Internally holds rgba or grayscale (1 channel)
std::vector<uint8_t> data_;
unsigned int width_;
@ -142,10 +151,22 @@ class Image {
bool is_grayscale_;
bool is_dithered_grayscale_;
bool is_dithered_color_;
bool is_preserving_blue_noise_offsets_;
void DecodePNG(const std::string &filename);
void DecodePGM(const std::string &filename);
void DecodePPM(const std::string &filename);
const std::string &GetGrayscaleKernelName();
const std::string &GetColorKernelName();
void GenerateBlueNoiseOffsets();
bool DuplicateBlueNoiseOffsetExists() const;
bool VerifyOpenCLBuffers(const std::string &kernel_name,
const std::vector<std::string> &buffer_names,
const Image *input_image,
const Image *blue_noise_image) const;
};
#endif

2
src/main.cc
View file

@ -10,7 +10,7 @@ int main(int argc, char **argv) {
return 1;
}
Video video("input.mp4");
if (!video.DitherVideo("output.mp4", &blue_noise)) {
if (!video.DitherVideo("output.mp4", &blue_noise, false)) {
std::cout << "ERROR: Failed to dither video" << std::endl;
return 1;
}

292
src/opencl_handle.cc
View file

@ -1,13 +1,14 @@
#include "opencl_handle.h"
#include <CL/cl.h>
#include <fstream>
#include <iostream>
#include <vector>
OpenCLContext::Ptr OpenCLContext::instance_ = {};
OpenCLContext::OpenCLHandle::OpenCLHandle()
: opencl_ptr_(), kernels_(), kernel_counter_(0) {}
OpenCLContext::OpenCLHandle::OpenCLHandle() : opencl_ptr_(), kernels_() {}
OpenCLContext::OpenCLHandle::~OpenCLHandle() {
CleanupAllKernels();
@ -23,21 +24,26 @@ bool OpenCLContext::OpenCLHandle::IsValid() const {
return context->IsValid();
}
KernelID OpenCLContext::OpenCLHandle::CreateKernelFromSource(
const std::string &kernel_fn, const char *kernel_name) {
bool OpenCLContext::OpenCLHandle::CreateKernelFromSource(
const std::string &kernel_fn, const std::string &kernel_name) {
if (!IsValid()) {
std::cout << "ERROR: OpenCLContext is not initialized" << std::endl;
return 0;
return false;
} else if (HasKernel(kernel_name)) {
std::cout
<< "ERROR: OpenCLContext already has kernel with given kernel_name \""
<< kernel_name << '"' << std::endl;
return false;
}
cl_int err_num;
KernelInfo kernel_info = {nullptr, nullptr, {}, 0};
KernelInfo kernel_info = {nullptr, nullptr, {}};
OpenCLContext::Ptr context_ptr = opencl_ptr_.lock();
if (!context_ptr) {
std::cout << "ERROR: OpenCLHandle: OpenCLContext is not initialized"
<< std::endl;
return 0;
return false;
}
const char *source_c_str = kernel_fn.c_str();
@ -46,7 +52,7 @@ KernelID OpenCLContext::OpenCLHandle::CreateKernelFromSource(
if (err_num != CL_SUCCESS) {
std::cout << "ERROR: OpenCLHandle: Failed to create program from source"
<< std::endl;
return 0;
return false;
}
err_num = clBuildProgram(kernel_info.program_, 0, nullptr, nullptr, nullptr,
@ -61,49 +67,54 @@ KernelID OpenCLContext::OpenCLHandle::CreateKernelFromSource(
build_log.data(), nullptr);
std::cout << build_log.data();
clReleaseProgram(kernel_info.program_);
return 0;
return false;
}
kernel_info.kernel_ =
clCreateKernel(kernel_info.program_, kernel_name, &err_num);
clCreateKernel(kernel_info.program_, kernel_name.c_str(), &err_num);
if (err_num != CL_SUCCESS) {
std::cout << "ERROR: OpenCLHandle: Failed to create kernel object from "
<< "source" << std::endl;
clReleaseProgram(kernel_info.program_);
return 0;
return false;
}
KernelID id;
do {
id = ++kernel_counter_;
} while (id == 0 || kernels_.find(id) != kernels_.end());
kernels_.insert({kernel_name, kernel_info});
kernels_.insert({id, kernel_info});
return id;
return true;
}
KernelID OpenCLContext::OpenCLHandle::CreateKernelFromSource(
const char *kernel_fn, const char *kernel_name) {
bool OpenCLContext::OpenCLHandle::CreateKernelFromSource(
const char *kernel_fn, const std::string &kernel_name) {
if (!IsValid()) {
std::cout << "ERROR: OpenCLContext is not initialized" << std::endl;
return 0;
return false;
} else if (HasKernel(kernel_name)) {
std::cout
<< "ERROR: OpenCLContext already has kernel with given kernel_name \""
<< kernel_name << '"' << std::endl;
return false;
}
return CreateKernelFromSource(std::string(kernel_fn), kernel_name);
}
KernelID OpenCLContext::OpenCLHandle::CreateKernelFromFile(
const std::string &filename, const char *kernel_name) {
bool OpenCLContext::OpenCLHandle::CreateKernelFromFile(
const std::string &filename, const std::string &kernel_name) {
if (!IsValid()) {
std::cout << "ERROR: OpenCLContext is not initialized" << std::endl;
return 0;
return false;
} else if (HasKernel(kernel_name)) {
std::cout
<< "ERROR: OpenCLContext already has kernel with given kernel_name \""
<< kernel_name << '"' << std::endl;
return false;
}
std::string source;
{
char buf[1024];
std::ifstream ifs(filename);
if (!ifs.is_open()) {
return 0;
return false;
}
while (ifs.good()) {
@ -112,38 +123,52 @@ KernelID OpenCLContext::OpenCLHandle::CreateKernelFromFile(
}
if (source.empty()) {
return 0;
return false;
}
}
return CreateKernelFromSource(source, kernel_name);
}
KernelID OpenCLContext::OpenCLHandle::CreateKernelFromFile(
const char *filename, const char *kernel_name) {
bool OpenCLContext::OpenCLHandle::CreateKernelFromFile(
const char *filename, const std::string &kernel_name) {
if (!IsValid()) {
std::cout << "ERROR: OpenCLContext is not initialized" << std::endl;
return 0;
return false;
} else if (HasKernel(kernel_name)) {
std::cout
<< "ERROR: OpenCLContext already has kernel with given kernel_name \""
<< kernel_name << '"' << std::endl;
return false;
}
return CreateKernelFromFile(std::string(filename), kernel_name);
}
BufferID OpenCLContext::OpenCLHandle::CreateKernelBuffer(KernelID kernel_id,
cl_mem_flags flags,
std::size_t buf_size,
void *host_ptr) {
bool OpenCLContext::OpenCLHandle::CreateKernelBuffer(
const std::string &kernel_name, cl_mem_flags flags, std::size_t buf_size,
void *host_ptr, const std::string &buffer_name) {
if (!IsValid()) {
std::cout << "ERROR: OpenCLHandle::CreateKernelBuffer: OpenCLContext is "
"not initialized"
<< std::endl;
return 0;
return false;
}
auto kernel_info_iter = kernels_.find(kernel_id);
auto kernel_info_iter = kernels_.find(kernel_name);
if (kernel_info_iter == kernels_.end()) {
std::cout
<< "ERROR: OpenCLHandle::CreateKernelBuffer: Got Invalid kernel_id"
<< std::endl;
return 0;
std::cout << "ERROR: OpenCLHandle::CreateKernelBuffer: Given kernel \""
<< kernel_name << "\" doesn't exist" << std::endl;
return false;
}
auto *buffer_map = &kernel_info_iter->second.mem_objects_;
{
auto buffer_info_iter = buffer_map->find(buffer_name);
if (buffer_info_iter != buffer_map->end()) {
std::cout
<< "ERROR: OpenCLHandle::CreateKernelBuffer: Buffer with name \""
<< buffer_name << "\" already exists" << std::endl;
return false;
}
}
auto opencl_context = opencl_ptr_.lock();
@ -151,7 +176,7 @@ BufferID OpenCLContext::OpenCLHandle::CreateKernelBuffer(KernelID kernel_id,
std::cout << "ERROR: OpenCLHandle::CreateKernelBuffer: OpenCLContext is "
"not initialized"
<< std::endl;
return 0;
return false;
}
cl_int err_num;
@ -163,39 +188,33 @@ BufferID OpenCLContext::OpenCLHandle::CreateKernelBuffer(KernelID kernel_id,
std::cout
<< "ERROR: OpenCLHandle::CreateKernelBuffer: Failed to create buffer"
<< std::endl;
return 0;
return false;
}
BufferID buffer_id;
auto *buffer_map = &kernel_info_iter->second.mem_objects_;
do {
buffer_id = ++kernel_info_iter->second.buffer_id_counter_;
} while (buffer_id == 0 || buffer_map->find(buffer_id) != buffer_map->end());
buffer_map->insert({buffer_name, {mem_object, buf_size}});
buffer_map->insert({buffer_id, {mem_object, buf_size}});
return buffer_id;
return true;
}
bool OpenCLContext::OpenCLHandle::SetKernelBufferData(KernelID kernel_id,
BufferID buffer_id,
std::size_t data_size,
void *data_ptr) {
bool OpenCLContext::OpenCLHandle::SetKernelBufferData(
const std::string &kernel_name, const std::string &buffer_name,
std::size_t data_size, void *data_ptr) {
if (!IsValid()) {
std::cout << "ERROR: OpenCLContext is not initialized" << std::endl;
return false;
}
auto kernel_info_iter = kernels_.find(kernel_id);
auto kernel_info_iter = kernels_.find(kernel_name);
if (kernel_info_iter == kernels_.end()) {
std::cout << "ERROR: OpenCLHandle::SetKernelBufferData: Invalid KernelID"
<< std::endl;
std::cout << "ERROR: OpenCLHandle::SetKernelBufferData: Kernel with name \""
<< kernel_name << "\" doesn't exist" << std::endl;
return false;
}
auto buffer_info_iter = kernel_info_iter->second.mem_objects_.find(buffer_id);
if (buffer_info_iter == kernel_info_iter->second.mem_objects_.end()) {
std::cout << "ERROR: OpenCLHandle::SetKernelBufferData: Invalid BufferID"
<< std::endl;
auto *buffer_map = &kernel_info_iter->second.mem_objects_;
auto buffer_info_iter = buffer_map->find(buffer_name);
if (buffer_info_iter == buffer_map->end()) {
std::cout << "ERROR: OpenCLHandle::SetKernelBufferData: Buffer with name \""
<< buffer_name << "\" doesn't exist" << std::endl;
return false;
}
@ -233,27 +252,25 @@ bool OpenCLContext::OpenCLHandle::SetKernelBufferData(KernelID kernel_id,
return true;
}
bool OpenCLContext::OpenCLHandle::AssignKernelBuffer(KernelID kernel_id,
unsigned int idx,
BufferID buffer_id) {
bool OpenCLContext::OpenCLHandle::AssignKernelBuffer(
const std::string &kernel_name, unsigned int idx,
const std::string &buffer_name) {
if (!IsValid()) {
std::cout << "ERROR: OpenCLContext is not initialized" << std::endl;
return false;
}
auto kernel_iter = kernels_.find(kernel_id);
auto kernel_iter = kernels_.find(kernel_name);
if (kernel_iter == kernels_.end()) {
std::cout
<< "ERROR: OpenCLHandle::AssignKernelBuffer: no kernel with given id"
<< std::endl;
std::cout << "ERROR: OpenCLHandle::AssignKernelBuffer: Kernel with name \""
<< kernel_name << "\" doesn't exist" << std::endl;
return false;
}
auto *buffer_map = &kernel_iter->second.mem_objects_;
auto buffer_info_iter = buffer_map->find(buffer_id);
auto buffer_info_iter = buffer_map->find(buffer_name);
if (buffer_info_iter == buffer_map->end()) {
std::cout << "ERROR: OpenCLHandle::AssignKernelBuffer: no buffer in "
"kernel_info with given id"
<< std::endl;
std::cout << "ERROR: OpenCLHandle::AssignKernelBuffer: buffer with name \""
<< buffer_name << "\" doesn't exist" << std::endl;
return false;
}
@ -271,19 +288,18 @@ bool OpenCLContext::OpenCLHandle::AssignKernelBuffer(KernelID kernel_id,
return true;
}
bool OpenCLContext::OpenCLHandle::AssignKernelArgument(KernelID kernel_id,
unsigned int idx,
std::size_t data_size,
const void *data_ptr) {
bool OpenCLContext::OpenCLHandle::AssignKernelArgument(
const std::string &kernel_name, unsigned int idx, std::size_t data_size,
const void *data_ptr) {
if (!IsValid()) {
std::cout << "ERROR: OpenCLContext is not initialized" << std::endl;
return false;
}
auto iter = kernels_.find(kernel_id);
auto iter = kernels_.find(kernel_name);
if (iter == kernels_.end()) {
std::cout
<< "ERROR: OpenCLHandle::AssignKernelArgument: no kernel with given id"
<< std::endl;
<< "ERROR: OpenCLHandle::AssignKernelArgument: Kernel with name \""
<< kernel_name << "\" doesn't exist" << std::endl;
return false;
}
@ -309,7 +325,7 @@ bool OpenCLContext::OpenCLHandle::AssignKernelArgument(KernelID kernel_id,
}
std::array<std::size_t, 3> OpenCLContext::OpenCLHandle::GetGlobalWorkSize(
KernelID kernel_id) {
const std::string &kernel_name) {
if (!IsValid()) {
std::cout << "ERROR: OpenCLContext is not initialized" << std::endl;
return {0, 0, 0};
@ -324,10 +340,10 @@ std::array<std::size_t, 3> OpenCLContext::OpenCLHandle::GetGlobalWorkSize(
return sizes;
}
auto kernel_iter = kernels_.find(kernel_id);
auto kernel_iter = kernels_.find(kernel_name);
if (kernel_iter == kernels_.end()) {
std::cout << "ERROR: OpenCLHandle::GetGlobalWorkSize: Invalid kernel_id"
<< std::endl;
std::cout << "ERROR: OpenCLHandle::GetGlobalWorkSize: Kernel with name \""
<< kernel_name << "\" doesn't exist" << std::endl;
return sizes;
}
@ -345,7 +361,8 @@ std::array<std::size_t, 3> OpenCLContext::OpenCLHandle::GetGlobalWorkSize(
return sizes;
}
std::size_t OpenCLContext::OpenCLHandle::GetWorkGroupSize(KernelID kernel_id) {
std::size_t OpenCLContext::OpenCLHandle::GetWorkGroupSize(
const std::string &kernel_name) {
if (!IsValid()) {
std::cout << "ERROR: OpenCLContext is not initialized" << std::endl;
return 0;
@ -358,10 +375,10 @@ std::size_t OpenCLContext::OpenCLHandle::GetWorkGroupSize(KernelID kernel_id) {
return 0;
}
auto kernel_iter = kernels_.find(kernel_id);
auto kernel_iter = kernels_.find(kernel_name);
if (kernel_iter == kernels_.end()) {
std::cout << "ERROR: OpenCLHandle::GetWorkGroupSize: Invalid kernel_id"
<< std::endl;
std::cout << "ERROR: OpenCLHandle::GetWorkGroupSize: Kernel with name \""
<< kernel_name << "\" doesn't exist" << std::endl;
return 0;
}
@ -404,7 +421,7 @@ std::size_t OpenCLContext::OpenCLHandle::GetDeviceMaxWorkGroupSize() {
return value;
}
bool OpenCLContext::OpenCLHandle::ExecuteKernel(KernelID kernel_id,
bool OpenCLContext::OpenCLHandle::ExecuteKernel(const std::string &kernel_name,
std::size_t global_work_size,
std::size_t local_work_size,
bool is_blocking) {
@ -420,10 +437,10 @@ bool OpenCLContext::OpenCLHandle::ExecuteKernel(KernelID kernel_id,
return false;
}
auto kernel_iter = kernels_.find(kernel_id);
auto kernel_iter = kernels_.find(kernel_name);
if (kernel_iter == kernels_.end()) {
std::cout << "ERROR: OpenCLHandle::ExecuteKernel: Invalid kernel_id"
<< std::endl;
std::cout << "ERROR: OpenCLHandle::ExecuteKernel: Kernel with name \""
<< kernel_name << "\" doesn't exist" << std::endl;
return false;
}
@ -452,7 +469,7 @@ bool OpenCLContext::OpenCLHandle::ExecuteKernel(KernelID kernel_id,
}
bool OpenCLContext::OpenCLHandle::ExecuteKernel2D(
KernelID kernel_id, std::size_t global_work_size_0,
const std::string &kernel_name, 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()) {
@ -467,10 +484,10 @@ bool OpenCLContext::OpenCLHandle::ExecuteKernel2D(
return false;
}
auto kernel_iter = kernels_.find(kernel_id);
auto kernel_iter = kernels_.find(kernel_name);
if (kernel_iter == kernels_.end()) {
std::cout << "ERROR: OpenCLHandle::ExecuteKernel2D: Invalid kernel_id"
<< std::endl;
std::cout << "ERROR: OpenCLHandle::ExecuteKernel2D: Kernel with name \""
<< kernel_name << "\" doesn't exist" << std::endl;
return false;
}
@ -503,8 +520,8 @@ bool OpenCLContext::OpenCLHandle::ExecuteKernel2D(
return true;
}
bool OpenCLContext::OpenCLHandle::GetBufferData(KernelID kernel_id,
BufferID buffer_id,
bool OpenCLContext::OpenCLHandle::GetBufferData(const std::string &kernel_name,
const std::string &buffer_name,
std::size_t out_size,
void *data_out) {
if (!IsValid()) {
@ -517,17 +534,17 @@ bool OpenCLContext::OpenCLHandle::GetBufferData(KernelID kernel_id,
std::cout << "ERROR: OpenCLContext is not initialized" << std::endl;
}
auto kernel_iter = kernels_.find(kernel_id);
auto kernel_iter = kernels_.find(kernel_name);
if (kernel_iter == kernels_.end()) {
std::cout << "ERROR: OpenCLHandle::GetBufferData: Invalid kernel_id"
<< std::endl;
std::cout << "ERROR: OpenCLHandle::GetBufferData: Kernel with name \""
<< kernel_name << "\" doesn't exist" << std::endl;
return false;
}
auto buffer_iter = kernel_iter->second.mem_objects_.find(buffer_id);
auto buffer_iter = kernel_iter->second.mem_objects_.find(buffer_name);
if (buffer_iter == kernel_iter->second.mem_objects_.end()) {
std::cout << "ERROR: OpenCLHandle::GetBufferData: Invalid buffer_id"
<< std::endl;
std::cout << "ERROR: OpenCLHandle::GetBufferData: Buffer with name \""
<< buffer_name << "\" doesn't exist" << std::endl;
return false;
}
@ -558,23 +575,71 @@ bool OpenCLContext::OpenCLHandle::GetBufferData(KernelID kernel_id,
return true;
}
bool OpenCLContext::OpenCLHandle::CleanupBuffer(KernelID kernel_id,
BufferID buffer_id) {
bool OpenCLContext::OpenCLHandle::HasKernel(
const std::string &kernel_name) const {
return kernels_.find(kernel_name) != kernels_.end();
}
bool OpenCLContext::OpenCLHandle::HasBuffer(
const std::string &kernel_name, const std::string &buffer_name) const {
auto kernel_iter = kernels_.find(kernel_name);
if (kernel_iter == kernels_.end()) {
return false;
}
auto *buffer_map = &kernel_iter->second.mem_objects_;
auto buffer_iter = buffer_map->find(buffer_name);
if (buffer_iter == buffer_map->end()) {
return false;
}
return true;
}
std::size_t OpenCLContext::OpenCLHandle::GetBufferSize(
const std::string &kernel_name, const std::string &buffer_name) const {
auto kernel_iter = kernels_.find(kernel_name);
if (kernel_iter == kernels_.end()) {
return 0;
}
auto *buffer_map = &kernel_iter->second.mem_objects_;
auto buffer_iter = buffer_map->find(buffer_name);
if (buffer_iter == buffer_map->end()) {
return 0;
}
std::size_t size = 0;
cl_int err = clGetMemObjectInfo(buffer_iter->second.mem, CL_MEM_SIZE,
sizeof(std::size_t), &size, nullptr);
if (err != CL_SUCCESS) {
std::cout << "ERROR: Failed to query size of device buffer \""
<< buffer_name << "\" with kernel \"" << kernel_name << '"'
<< std::endl;
return 0;
}
return size;
}
bool OpenCLContext::OpenCLHandle::CleanupBuffer(
const std::string &kernel_name, const std::string &buffer_name) {
if (!IsValid()) {
std::cout << "ERROR: OpenCLContext is not initialized" << std::endl;
return false;
}
auto kernel_iter = kernels_.find(kernel_id);
auto kernel_iter = kernels_.find(kernel_name);
if (kernel_iter == kernels_.end()) {
std::cout << "ERROR: OpenCLHandle::CleanupBuffer: Invalid kernel_id"
<< std::endl;
std::cout << "ERROR: OpenCLHandle::CleanupBuffer: Kernel with name \""
<< kernel_name << "\" doesn't exist" << std::endl;
return false;
}
auto buffer_iter = kernel_iter->second.mem_objects_.find(buffer_id);
auto buffer_iter = kernel_iter->second.mem_objects_.find(buffer_name);
if (buffer_iter == kernel_iter->second.mem_objects_.end()) {
std::cout << "ERROR: OpenCLHandle::CleanupBuffer: Invalid buffer_id"
<< std::endl;
std::cout << "ERROR: OpenCLHandle::CleanupBuffer: Buffer with name \""
<< buffer_name << "\" doesn't exist" << std::endl;
return false;
}
@ -584,13 +649,16 @@ bool OpenCLContext::OpenCLHandle::CleanupBuffer(KernelID kernel_id,
return true;
}
bool OpenCLContext::OpenCLHandle::CleanupKernel(KernelID id) {
bool OpenCLContext::OpenCLHandle::CleanupKernel(
const std::string &kernel_name) {
if (!IsValid()) {
std::cout << "ERROR: OpenCLContext is not initialized" << std::endl;
return false;
}
auto iter = kernels_.find(id);
auto iter = kernels_.find(kernel_name);
if (iter == kernels_.end()) {
std::cout << "WARNING CleanupKernel: Kernel with name \"" << kernel_name
<< "\" doesn't exist" << std::endl;
return false;
}

198
src/opencl_handle.h
View file

@ -14,21 +14,18 @@
#include <CL/cl.h>
#endif
typedef unsigned int KernelID;
typedef unsigned int BufferID;
class OpenCLContext {
public:
typedef std::shared_ptr<OpenCLContext> Ptr;
typedef std::weak_ptr<OpenCLContext> WeakPtr;
/*!
* \brief A simplified handle to OpenCL
* \brief A simplified handle to OpenCL.
*
* This class can only be obtained by a call to OpenCLContext::GetHandle()
* This class can only be obtained by a call to OpenCLContext::GetHandle().
*
* OpenCL is automatically cleaned up when all shared ptrs of OpenCLHandle are
* destructed
* destructed.
*/
class OpenCLHandle {
public:
@ -48,148 +45,175 @@ class OpenCLContext {
bool IsValid() const;
/*!
* \brief Returns the KernelID, to be used with other fns in OpenCLHandle
* \brief Compiles a kernel from source that can be referenced with the
* given kernel name.
*
* The created kernel can be free'd with a call to CleanupKernel(KernelID)
* The created kernel can be free'd with a call to CleanupKernel().
*
* \return KernelID with value 0 on failure, non-zero otherwise
* \return True on success.
*/
KernelID CreateKernelFromSource(const std::string &kernel_fn,
const char *kernel_name);
/*!
* \brief Returns the KernelID, to be used with other fns in OpenCLHandle
*
* The created kernel can be free'd with a call to CleanupKernel(KernelID)
*
* \return KernelID with value 0 on failure, non-zero otherwise
*/
KernelID CreateKernelFromSource(const char *kernel_fn,
const char *kernel_name);
/*!
* \brief Returns the KernelID, to be used with other fns in OpenCLHandle
*
* The created kernel can be free'd with a call to CleanupKernel(KernelID)
*
* \return KernelID with value 0 on failure, non-zero otherwise
*/
KernelID CreateKernelFromFile(const std::string &filename,
const char *kernel_name);
/*!
* \brief Returns the KernelID, to be used with other fns in OpenCLHandle
*
* The created kernel can be free'd with a call to CleanupKernel(KernelID)
*
* \return KernelID with value 0 on failure, non-zero otherwise
*/
KernelID CreateKernelFromFile(const char *filename,
const char *kernel_name);
bool CreateKernelFromSource(const std::string &kernel_fn,
const std::string &kernel_name);
/*!
* \brief Creates a cl_mem buffer and returns its id
* \brief Compiles a kernel from source that can be referenced with the
* given kernel name.
*
* The created kernel can be free'd with a call to CleanupKernel().
*
* \return True on success.
*/
bool CreateKernelFromSource(const char *kernel_fn,
const std::string &kernel_name);
/*!
* \brief Compiles a kernel from a file that can be referenced with the
* given kernel name.
*
* The created kernel can be free'd with a call to CleanupKernel().
*
* \return True on success.
*/
bool CreateKernelFromFile(const std::string &filename,
const std::string &kernel_name);
/*!
* \brief Compiles a kernel from a file that can be referenced with the
* given kernel name.
*
* The created kernel can be free'd with a call to CleanupKernel().
*
* \return True on success.
*/
bool CreateKernelFromFile(const char *filename,
const std::string &kernel_name);
/*!
* \brief Creates a cl_mem buffer that can be referenced with the given
* buffer_name.
*
* Note that the created buffer is stored with the specified kernel's data.
* This means that the created buffer can only be accessed with the
* KernelID that was used to create it.
* This means that the created buffer can only be accessed with the kernel
* that was used to create it.
*
* If buf_size is set to 0 and host_ptr set to nullptr, then the created
* buffer will be uninitialized.
* If host_ptr set to nullptr, then the created buffer will be
* uninitialized.
*
* \return non-zero BufferID on success
* \return True on success.
*/
BufferID CreateKernelBuffer(KernelID kernel_id, cl_mem_flags flags,
std::size_t buf_size, void *host_ptr);
bool CreateKernelBuffer(const std::string &kernel_name, cl_mem_flags flags,
std::size_t buf_size, void *host_ptr,
const std::string &buffer_name);
/*!
* \brief Assign host data to existing device buffer
*
* \return true on success
* The kernel referenced by kernel_name must exist, and the buffer
* referenced by buffer_name must also exist.
*
* \return True on success.
*/
bool SetKernelBufferData(KernelID kernel_id, BufferID buffer_id,
bool SetKernelBufferData(const std::string &kernel_name,
const std::string &buffer_name,
std::size_t data_size, void *data_ptr);
/*!
* \brief Assign a previously created buffer to a kernel function's
* parameter
* parameter.
*
* \return true on success
* \return true on success.
*/
bool AssignKernelBuffer(KernelID kernel_id, unsigned int idx,
BufferID buffer_id);
bool AssignKernelBuffer(const std::string &kernel_name, unsigned int idx,
const std::string &buffer_name);
/*!
* \brief Assign data to a kernel function's parameter
* \brief Assign data to a kernel function's parameter.
*
* id refers to the kernel's id, and idx refers to the parameter index for
* the kernel function.
* idx refers to the parameter index for the kernel function.
*
* \return true on success
* \return true on success.
*/
bool AssignKernelArgument(KernelID kernel_id, unsigned int idx,
bool AssignKernelArgument(const std::string &kernel_name, unsigned int idx,
std::size_t data_size, const void *data_ptr);
/*!
* \brief Gets the sizes associated with CL_KERNEL_GLOBAL_WORK_SIZE
* \brief Gets the sizes associated with CL_KERNEL_GLOBAL_WORK_SIZE.
*
* \return {0, 0, 0} on failure
* \return {0, 0, 0} on failure.
*/
std::array<std::size_t, 3> GetGlobalWorkSize(KernelID kernel_id);
std::array<std::size_t, 3> GetGlobalWorkSize(
const std::string &kernel_name);
/*!
* \brief Gets the size associated with CL_KERNEL_WORK_GROUP_SIZE
* \brief Gets the size associated with CL_KERNEL_WORK_GROUP_SIZE.
*
* \return 0 on failure
* \return 0 on failure.
*/
std::size_t GetWorkGroupSize(KernelID kernel_id);
std::size_t GetWorkGroupSize(const std::string &kernel_name);
std::size_t GetDeviceMaxWorkGroupSize();
/*!
* \brief Executes the kernel with the given kernel_id
* \brief Executes the kernel with the given kernel_name.
*
* \return true on success
* \return true on success.
*/
bool ExecuteKernel(KernelID kernel_id, std::size_t global_work_size,
bool ExecuteKernel(const std::string &kernel_name,
std::size_t global_work_size,
std::size_t local_work_size, bool is_blocking);
/*!
* \brief Executes the kernel with the given kernel_id
* \brief Executes the kernel with the given kernel_name.
*
* \return true on success
* \return true on success.
*/
bool ExecuteKernel2D(KernelID kernel_id, std::size_t global_work_size_0,
bool ExecuteKernel2D(const std::string &kernel_name,
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
* \brief Copies device memory to data_out.
*
* \return true on success
* \return True on success.
*/
bool GetBufferData(KernelID kernel_id, BufferID buffer_id,
std::size_t out_size, void *data_out);
bool GetBufferData(const std::string &kernel_name,
const std::string &buffer_name, std::size_t out_size,
void *data_out);
/// Returns true if the kernel exists
bool HasKernel(const std::string &kernel_name) const;
/// Returns true if the buffer exists with the kernel
bool HasBuffer(const std::string &kernel_name,
const std::string &buffer_name) const;
/// Returns the buffer size in bytes, or 0 if error
std::size_t GetBufferSize(const std::string &kernel_name,
const std::string &buffer_name) const;
/*!
* \brief Cleans up a mem buffer
* \brief Cleans up a mem buffer.
*
* If using CleanupKernel(KernelID id), there is no need to call this
* function with the same kernel_id as it will cleanup the associated mem
* buffers.
* If using CleanupKernel(), there is no need to call this function with the
* same kernel_id as it will cleanup the associated mem buffers.
*
* \return true if clean has occurred
* \return true if clean has occurred.
*/
bool CleanupBuffer(KernelID kernel_id, BufferID buffer_id);
bool CleanupBuffer(const std::string &kernel_name,
const std::string &buffer_name);
/*!
* \brief Cleans up a kernel object and its associated data (like mem
* buffers)
* \brief Cleans up a kernel object and its associated data (including mem
* buffers).
*
* \return true if cleanup has occurred
* \return true if cleanup has occurred.
*/
bool CleanupKernel(KernelID id);
bool CleanupKernel(const std::string &kernel_name);
/*!
* \brief Cleans up all Kernel data (including mem buffers)
* \brief Cleans up all Kernel data (including mem buffers).
*/
void CleanupAllKernels();
@ -204,16 +228,14 @@ class OpenCLContext {
struct KernelInfo {
cl_kernel kernel_;
cl_program program_;
std::unordered_map<BufferID, BufferInfo> mem_objects_;
BufferID buffer_id_counter_;
std::unordered_map<std::string, BufferInfo> mem_objects_;
};
OpenCLHandle();
OpenCLContext::WeakPtr opencl_ptr_;
std::unordered_map<KernelID, KernelInfo> kernels_;
KernelID kernel_counter_;
std::unordered_map<std::string, KernelInfo> kernels_;
};
~OpenCLContext();