From 6e5eaac63d86eec95b7f5fca3d8b75c5cfa62a52 Mon Sep 17 00:00:00 2001 From: Stephen Seo Date: Wed, 1 Dec 2021 16:02:34 +0900 Subject: [PATCH] Refactoring of opencl_handle and image for reuse Refactoring was done so that the OpenCL kernels wouldn't need to be recompiled every time an image was dithered, among other efficiency fixes. --- src/image.cc | 510 +++++++++++++++++++++++++++++-------------- src/image.h | 29 ++- src/main.cc | 2 +- src/opencl_handle.cc | 292 +++++++++++++++---------- src/opencl_handle.h | 198 +++++++++-------- 5 files changed, 667 insertions(+), 364 deletions(-) diff --git a/src/image.cc b/src/image.cc index 23c1c37..90712d2 100644 --- a/src/image.cc +++ b/src/image.cc @@ -8,15 +8,29 @@ #include #include -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 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 Image::kDitherBWPalette = { png_color{0, 0, 0}, // black png_color{255, 255, 255} // white }; -const std::array Image::dither_color_palette_ = { +const std::array 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 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::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::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::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 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::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::ToColorDitheredWithBlueNoise(Image *blue_noise) { std::unique_ptr(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 &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; +} diff --git a/src/image.h b/src/image.h index 2ccc8e5..797f19b 100644 --- a/src/image.h +++ b/src/image.h @@ -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 dither_bw_palette_; - static const std::array dither_color_palette_; + static constexpr unsigned int kBlueNoiseOffsetMax = 128; + static const char *kOpenCLGrayscaleKernel; + static const char *kOpenCLColorKernel; + static const std::array kDitherBWPalette; + static const std::array 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 blue_noise_offsets_; /// Internally holds rgba or grayscale (1 channel) std::vector 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 &buffer_names, + const Image *input_image, + const Image *blue_noise_image) const; }; #endif diff --git a/src/main.cc b/src/main.cc index 90e7ba5..9309377 100644 --- a/src/main.cc +++ b/src/main.cc @@ -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; } diff --git a/src/opencl_handle.cc b/src/opencl_handle.cc index 786e64b..addfba8 100644 --- a/src/opencl_handle.cc +++ b/src/opencl_handle.cc @@ -1,13 +1,14 @@ #include "opencl_handle.h" +#include + #include #include #include 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 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 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 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; } diff --git a/src/opencl_handle.h b/src/opencl_handle.h index b93d914..8618189 100644 --- a/src/opencl_handle.h +++ b/src/opencl_handle.h @@ -14,21 +14,18 @@ #include #endif -typedef unsigned int KernelID; -typedef unsigned int BufferID; - class OpenCLContext { public: typedef std::shared_ptr Ptr; typedef std::weak_ptr 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 GetGlobalWorkSize(KernelID kernel_id); + std::array 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 mem_objects_; - BufferID buffer_id_counter_; + std::unordered_map mem_objects_; }; OpenCLHandle(); OpenCLContext::WeakPtr opencl_ptr_; - std::unordered_map kernels_; - KernelID kernel_counter_; + std::unordered_map kernels_; }; ~OpenCLContext();