#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
};
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..."
std::cout << "ERROR: Unknown filename extension" << std::endl;
return;
}
+
+ GenerateBlueNoiseOffsets();
}
bool Image::IsValid() const {
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,
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,
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();
- return {};
+ 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 {};
+ }
}
- 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();
+ 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 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, 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 {};
+ }
+ }
+
+ 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;
--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;
}
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) {
- std::cout
- << "ERROR ToColorDitheredWithBlueNoise: Failed to set input buffer"
- << std::endl;
- opencl_handle->CleanupAllKernels();
- return {};
+ 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 {};
+ }
}
- auto output_buffer_id = opencl_handle->CreateKernelBuffer(
- kid, CL_MEM_WRITE_ONLY, this->data_.size(), nullptr);
- if (output_buffer_id == 0) {
+ if (!opencl_handle->SetKernelBufferData(color_kernel_name, kBufferInputName,
+ this->data_.size(),
+ this->data_.data())) {
std::cout
- << "ERROR ToColorDitheredWithBlueNoise: Failed to set output buffer"
+ << "ERROR ToColorDitheredWithBlueNoise: Failed to init input buffer"
<< std::endl;
- opencl_handle->CleanupAllKernels();
+ 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, 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 {};
+ }
+ }
+
+ 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;
--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 {};
}
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"
"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"
"}\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"
"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"
"}\n";
}
- return opencl_color_kernel_;
+ return kOpenCLColorKernel;
}
OpenCLHandle::Ptr Image::GetOpenCLHandle() {
<< 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;
+}
#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();
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();
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,
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({id, kernel_info});
+ kernels_.insert({kernel_name, 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()) {
}
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();
std::cout << "ERROR: OpenCLHandle::CreateKernelBuffer: OpenCLContext is "
"not initialized"
<< std::endl;
- return 0;
+ return false;
}
cl_int err_num;
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_id, {mem_object, buf_size}});
+ buffer_map->insert({buffer_name, {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;
}
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;
}
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;
}
}
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};
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;
}
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;
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;
}
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) {
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;
}
}
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()) {
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;
}
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()) {
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;
}
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;
}
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;
}
#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:
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);
+ bool CreateKernelFromSource(const std::string &kernel_fn,
+ const std::string &kernel_name);
+
/*!
- * \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 char *kernel_fn,
- const char *kernel_name);
+ bool CreateKernelFromSource(const char *kernel_fn,
+ const std::string &kernel_name);
+
/*!
- * \brief Returns the KernelID, to be used with other fns in OpenCLHandle
+ * \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(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 CreateKernelFromFile(const std::string &filename,
- const char *kernel_name);
+ bool CreateKernelFromFile(const std::string &filename,
+ const std::string &kernel_name);
+
/*!
- * \brief Returns the KernelID, to be used with other fns in OpenCLHandle
+ * \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(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 CreateKernelFromFile(const char *filename,
- const char *kernel_name);
+ bool CreateKernelFromFile(const char *filename,
+ const std::string &kernel_name);
/*!
- * \brief Creates a cl_mem buffer and returns its id
+ * \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();
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();
projects / EN605.617.81.FA21_StephenSeo_DitheringProject / commitdiff