stephenseo/blue_noise_generation
1
0
Fork 0
Code Issues Pull requests Projects Releases Wiki Activity

Make usage of OpenCL conditional (configuring)

Browse source
This commit is contained in:
Stephen Seo 2023-03-31 18:46:10 +09:00
parent 2f768f43bd
commit 62d15771ad
2 changed files with 269 additions and 187 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

25
CMakeLists.txt
View file

@ -18,16 +18,35 @@ if(NOT DEFINED CMAKE_BUILD_TYPE OR NOT CMAKE_BUILD_TYPE)
endif() endif()
find_package(Threads REQUIRED) find_package(Threads REQUIRED)
find_package(OpenCL REQUIRED) if(NOT DEFINED DISABLE_OPENCL)
find_package(OpenCL QUIET)
if(NOT DEFINED OpenCL_FOUND)
set(DISABLE_OPENCL True)
message(WARNING "OpenCL not found, OpenCL usage is disabled.")
endif()
endif()
find_package(PNG REQUIRED) find_package(PNG REQUIRED)
add_executable(Dithering ${Dithering_SOURCES}) add_executable(Dithering ${Dithering_SOURCES})
target_compile_features(Dithering PUBLIC cxx_std_17) target_compile_features(Dithering PUBLIC cxx_std_17)
target_include_directories(Dithering PUBLIC if(DEFINED DISABLE_OPENCL AND DISABLE_OPENCL)
message(STATUS "OpenCL usage is disabled.")
target_include_directories(Dithering PUBLIC
Threads::Threads
${PNG_INCLUDE_DIRS})
target_link_libraries(Dithering PUBLIC
Threads::Threads
${PNG_LIBRARIES})
target_compile_definitions(Dithering PRIVATE DITHERING_OPENCL_ENABLED=0)
else()
message(STATUS "OpenCL usage is enabled.")
target_include_directories(Dithering PUBLIC
Threads::Threads Threads::Threads
${OpenCL_INCLUDE_DIRS} ${OpenCL_INCLUDE_DIRS}
${PNG_INCLUDE_DIRS}) ${PNG_INCLUDE_DIRS})
target_link_libraries(Dithering PUBLIC target_link_libraries(Dithering PUBLIC
Threads::Threads Threads::Threads
${OpenCL_LIBRARIES} ${OpenCL_LIBRARIES}
${PNG_LIBRARIES}) ${PNG_LIBRARIES})
target_compile_definitions(Dithering PRIVATE DITHERING_OPENCL_ENABLED=1)
endif()

283
src/blue_noise.cpp
View file

@ -1,24 +1,26 @@
#include "blue_noise.hpp" #include "blue_noise.hpp"
#include <random>
#include <cassert> #include <cassert>
#include <iostream> #include <cstdio>
#include <fstream> #include <fstream>
#include <iostream>
#include <memory> #include <memory>
#include <random>
#include <string> #include <string>
#include <unordered_set> #include <unordered_set>
#include <cstdio>
#if DITHERING_OPENCL_ENABLED == 1
#include <CL/opencl.h> #include <CL/opencl.h>
#endif
#include "image.hpp" #include "image.hpp"
image::Bl dither::blue_noise(int width, int height, int threads,
image::Bl dither::blue_noise(int width, int height, int threads, bool use_opencl) { bool use_opencl) {
bool using_opencl = false; bool using_opencl = false;
if(use_opencl) { #if DITHERING_OPENCL_ENABLED == 1
if (use_opencl) {
// try to use OpenCL // try to use OpenCL
do { do {
cl_device_id device; cl_device_id device;
@ -31,52 +33,62 @@ image::Bl dither::blue_noise(int width, int height, int threads, bool use_opencl
int filter_size = (width + height) / 2; int filter_size = (width + height) / 2;
err = clGetPlatformIDs(1, &platform, nullptr); err = clGetPlatformIDs(1, &platform, nullptr);
if(err != CL_SUCCESS) { if (err != CL_SUCCESS) {
std::cerr << "OpenCL: Failed to identify a platform\n"; std::cerr << "OpenCL: Failed to identify a platform\n";
break; break;
} }
err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, nullptr); err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device,
if(err != CL_SUCCESS) { nullptr);
if (err != CL_SUCCESS) {
std::cerr << "OpenCL: Failed to get a device\n"; std::cerr << "OpenCL: Failed to get a device\n";
break; break;
} }
context = clCreateContext(nullptr, 1, &device, nullptr, nullptr, &err); context = clCreateContext(nullptr, 1, &device, nullptr, nullptr,
&err);
{ {
char buf[1024]; char buf[1024];
std::ifstream program_file("src/blue_noise.cl"); std::ifstream program_file("src/blue_noise.cl");
if (!program_file.good()) { if (!program_file.good()) {
std::cerr << "ERROR: Failed to read \"src/blue_noise.cl\" (not found?)\n"; std::cerr << "ERROR: Failed to read \"src/blue_noise.cl\" "
"(not found?)\n";
break; break;
} }
std::string program_string; std::string program_string;
while(program_file.good()) { while (program_file.good()) {
program_file.read(buf, 1024); program_file.read(buf, 1024);
if(int read_count = program_file.gcount(); read_count > 0) { if (int read_count = program_file.gcount();
read_count > 0) {
program_string.append(buf, read_count); program_string.append(buf, read_count);
} }
} }
const char *string_ptr = program_string.c_str(); const char *string_ptr = program_string.c_str();
std::size_t program_size = program_string.size(); std::size_t program_size = program_string.size();
program = clCreateProgramWithSource(context, 1, (const char**)&string_ptr, &program_size, &err); program = clCreateProgramWithSource(context, 1,
if(err != CL_SUCCESS) { (const char **)&string_ptr,
&program_size, &err);
if (err != CL_SUCCESS) {
std::cerr << "OpenCL: Failed to create the program\n"; std::cerr << "OpenCL: Failed to create the program\n";
clReleaseContext(context); clReleaseContext(context);
break; break;
} }
err = clBuildProgram(program, 1, &device, nullptr, nullptr, nullptr); err = clBuildProgram(program, 1, &device, nullptr, nullptr,
if(err != CL_SUCCESS) { nullptr);
if (err != CL_SUCCESS) {
std::cerr << "OpenCL: Failed to build the program\n"; std::cerr << "OpenCL: Failed to build the program\n";
std::size_t log_size; std::size_t log_size;
clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, nullptr, &log_size); clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG,
std::unique_ptr<char[]> log = std::make_unique<char[]>(log_size + 1); 0, nullptr, &log_size);
std::unique_ptr<char[]> log =
std::make_unique<char[]>(log_size + 1);
log[log_size] = 0; log[log_size] = 0;
clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, log_size, log.get(), nullptr); clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG,
log_size, log.get(), nullptr);
std::cerr << log.get() << std::endl; std::cerr << log.get() << std::endl;
clReleaseProgram(program); clReleaseProgram(program);
@ -92,24 +104,31 @@ image::Bl dither::blue_noise(int width, int height, int threads, bool use_opencl
clReleaseProgram(program); clReleaseProgram(program);
clReleaseContext(context); clReleaseContext(context);
if(!result.empty()) { if (!result.empty()) {
return internal::rangeToBl(result, width); return internal::rangeToBl(result, width);
} }
std::cout << "ERROR: Empty result\n"; std::cout << "ERROR: Empty result\n";
} while (false); } while (false);
} }
#else
std::clog << "WARNING: Not compiled with OpenCL support!\n";
#endif
if(!using_opencl) { if (!using_opencl) {
std::cout << "OpenCL: Failed to setup/use or is not enabled, using regular impl..." std::cout << "OpenCL: Failed to setup/use or is not enabled, using "
"regular impl..."
<< std::endl; << std::endl;
return internal::rangeToBl(internal::blue_noise_impl(width, height, threads), width); return internal::rangeToBl(
internal::blue_noise_impl(width, height, threads), width);
} }
std::cout << "ERROR: Invalid state (end of blue_noise fn)\n"; std::cout << "ERROR: Invalid state (end of blue_noise fn)\n";
return {}; return {};
} }
std::vector<unsigned int> dither::internal::blue_noise_impl(int width, int height, int threads) { std::vector<unsigned int> dither::internal::blue_noise_impl(int width,
int height,
int threads) {
int count = width * height; int count = width * height;
std::vector<float> filter_out; std::vector<float> filter_out;
filter_out.resize(count); filter_out.resize(count);
@ -119,36 +138,40 @@ std::vector<unsigned int> dither::internal::blue_noise_impl(int width, int heigh
pbp.resize(count); pbp.resize(count);
#ifndef NDEBUG #ifndef NDEBUG
printf("Inserting %d pixels into image of max count %d\n", pixel_count, count); printf("Inserting %d pixels into image of max count %d\n", pixel_count,
count);
// generate image from randomized pbp // generate image from randomized pbp
FILE *random_noise_image = fopen("random_noise.pbm", "w"); FILE *random_noise_image = fopen("random_noise.pbm", "w");
fprintf(random_noise_image, "P1\n%d %d\n", width, height); fprintf(random_noise_image, "P1\n%d %d\n", width, height);
for(int y = 0; y < height; ++y) { for (int y = 0; y < height; ++y) {
for(int x = 0; x < width; ++x) { for (int x = 0; x < width; ++x) {
fprintf(random_noise_image, "%d ", pbp[utility::twoToOne(x, y, width, height)] ? 1 : 0); fprintf(random_noise_image, "%d ",
pbp[utility::twoToOne(x, y, width, height)] ? 1 : 0);
} }
fputc('\n', random_noise_image); fputc('\n', random_noise_image);
} }
fclose(random_noise_image); fclose(random_noise_image);
#endif #endif
//#ifndef NDEBUG // #ifndef NDEBUG
int iterations = 0; int iterations = 0;
//#endif // #endif
int filter_size = (width + height) / 2; int filter_size = (width + height) / 2;
std::unique_ptr<std::vector<float>> precomputed = std::make_unique<std::vector<float>>(internal::precompute_gaussian(filter_size)); std::unique_ptr<std::vector<float>> precomputed =
std::make_unique<std::vector<float>>(
internal::precompute_gaussian(filter_size));
internal::compute_filter(pbp, width, height, count, filter_size, internal::compute_filter(pbp, width, height, count, filter_size, filter_out,
filter_out, precomputed.get(), threads); precomputed.get(), threads);
#ifndef NDEBUG #ifndef NDEBUG
internal::write_filter(filter_out, width, "filter_out_start.pgm"); internal::write_filter(filter_out, width, "filter_out_start.pgm");
#endif #endif
std::cout << "Begin BinaryArray generation loop\n"; std::cout << "Begin BinaryArray generation loop\n";
while(true) { while (true) {
#ifndef NDEBUG #ifndef NDEBUG
// if(++iterations % 10 == 0) { // if(++iterations % 10 == 0) {
printf("Iteration %d\n", ++iterations); printf("Iteration %d\n", ++iterations);
// } // }
#endif #endif
@ -156,13 +179,13 @@ std::vector<unsigned int> dither::internal::blue_noise_impl(int width, int heigh
internal::compute_filter(pbp, width, height, count, filter_size, internal::compute_filter(pbp, width, height, count, filter_size,
filter_out, precomputed.get(), threads); filter_out, precomputed.get(), threads);
//#ifndef NDEBUG // #ifndef NDEBUG
// for(int i = 0; i < count; ++i) { // for(int i = 0; i < count; ++i) {
// int x, y; // int x, y;
// std::tie(x, y) = internal::oneToTwo(i, width); // std::tie(x, y) = internal::oneToTwo(i, width);
// printf("%d (%d, %d): %f\n", i, x, y, filter_out[i]); // printf("%d (%d, %d): %f\n", i, x, y, filter_out[i]);
// } // }
//#endif // #endif
int min, max; int min, max;
std::tie(min, max) = internal::filter_minmax(filter_out, pbp); std::tie(min, max) = internal::filter_minmax(filter_out, pbp);
@ -176,23 +199,26 @@ std::vector<unsigned int> dither::internal::blue_noise_impl(int width, int heigh
// get second buffer's min // get second buffer's min
int second_min; int second_min;
std::tie(second_min, std::ignore) = internal::filter_minmax(filter_out, pbp); std::tie(second_min, std::ignore) =
internal::filter_minmax(filter_out, pbp);
if(second_min == max) { if (second_min == max) {
pbp[max] = true; pbp[max] = true;
break; break;
} else { } else {
pbp[second_min] = true; pbp[second_min] = true;
} }
if(iterations % 100 == 0) { if (iterations % 100 == 0) {
// generate blue_noise image from pbp // generate blue_noise image from pbp
#ifndef NDEBUG #ifndef NDEBUG
FILE *blue_noise_image = fopen("blue_noise.pbm", "w"); FILE *blue_noise_image = fopen("blue_noise.pbm", "w");
fprintf(blue_noise_image, "P1\n%d %d\n", width, height); fprintf(blue_noise_image, "P1\n%d %d\n", width, height);
for(int y = 0; y < height; ++y) { for (int y = 0; y < height; ++y) {
for(int x = 0; x < width; ++x) { for (int x = 0; x < width; ++x) {
fprintf(blue_noise_image, "%d ", pbp[utility::twoToOne(x, y, width, height)] ? 1 : 0); fprintf(blue_noise_image, "%d ",
pbp[utility::twoToOne(x, y, width, height)] ? 1
: 0);
} }
fputc('\n', blue_noise_image); fputc('\n', blue_noise_image);
} }
@ -200,8 +226,8 @@ std::vector<unsigned int> dither::internal::blue_noise_impl(int width, int heigh
#endif #endif
} }
} }
internal::compute_filter(pbp, width, height, count, filter_size, internal::compute_filter(pbp, width, height, count, filter_size, filter_out,
filter_out, precomputed.get(), threads); precomputed.get(), threads);
#ifndef NDEBUG #ifndef NDEBUG
internal::write_filter(filter_out, width, "filter_out_final.pgm"); internal::write_filter(filter_out, width, "filter_out_final.pgm");
#endif #endif
@ -210,9 +236,10 @@ std::vector<unsigned int> dither::internal::blue_noise_impl(int width, int heigh
// generate blue_noise image from pbp // generate blue_noise image from pbp
FILE *blue_noise_image = fopen("blue_noise.pbm", "w"); FILE *blue_noise_image = fopen("blue_noise.pbm", "w");
fprintf(blue_noise_image, "P1\n%d %d\n", width, height); fprintf(blue_noise_image, "P1\n%d %d\n", width, height);
for(int y = 0; y < height; ++y) { for (int y = 0; y < height; ++y) {
for(int x = 0; x < width; ++x) { for (int x = 0; x < width; ++x) {
fprintf(blue_noise_image, "%d ", pbp[utility::twoToOne(x, y, width, height)] ? 1 : 0); fprintf(blue_noise_image, "%d ",
pbp[utility::twoToOne(x, y, width, height)] ? 1 : 0);
} }
fputc('\n', blue_noise_image); fputc('\n', blue_noise_image);
} }
@ -231,14 +258,16 @@ std::vector<unsigned int> dither::internal::blue_noise_impl(int width, int heigh
#endif #endif
internal::compute_filter(pbp, width, height, count, filter_size, internal::compute_filter(pbp, width, height, count, filter_size,
filter_out, precomputed.get(), threads); filter_out, precomputed.get(), threads);
std::tie(std::ignore, max) = internal::filter_minmax(filter_out, pbp); std::tie(std::ignore, max) =
internal::filter_minmax(filter_out, pbp);
pbp[max] = false; pbp[max] = false;
dither_array[max] = i; dither_array[max] = i;
} }
pbp = pbp_copy; pbp = pbp_copy;
} }
std::cout << "\nRanking remainder of first half of pixels...\n"; std::cout << "\nRanking remainder of first half of pixels...\n";
for (unsigned int i = pixel_count; i < (unsigned int)((count + 1) / 2); ++i) { for (unsigned int i = pixel_count; i < (unsigned int)((count + 1) / 2);
++i) {
#ifndef NDEBUG #ifndef NDEBUG
std::cout << i << ' '; std::cout << i << ' ';
#endif #endif
@ -254,11 +283,12 @@ std::vector<unsigned int> dither::internal::blue_noise_impl(int width, int heigh
#ifndef NDEBUG #ifndef NDEBUG
std::cout << i << ' '; std::cout << i << ' ';
#endif #endif
for(unsigned int i = 0; i < pbp.size(); ++i) { for (unsigned int i = 0; i < pbp.size(); ++i) {
reversed_pbp[i] = !pbp[i]; reversed_pbp[i] = !pbp[i];
} }
internal::compute_filter(reversed_pbp, width, height, count, filter_size, internal::compute_filter(reversed_pbp, width, height, count,
filter_out, precomputed.get(), threads); filter_size, filter_out, precomputed.get(),
threads);
std::tie(std::ignore, max) = internal::filter_minmax(filter_out, pbp); std::tie(std::ignore, max) = internal::filter_minmax(filter_out, pbp);
pbp[max] = true; pbp[max] = true;
dither_array[max] = i; dither_array[max] = i;
@ -267,8 +297,10 @@ std::vector<unsigned int> dither::internal::blue_noise_impl(int width, int heigh
return dither_array; return dither_array;
} }
#if DITHERING_OPENCL_ENABLED == 1
std::vector<unsigned int> dither::internal::blue_noise_cl_impl( std::vector<unsigned int> dither::internal::blue_noise_cl_impl(
const int width, const int height, const int filter_size, cl_context context, cl_device_id device, cl_program program) { const int width, const int height, const int filter_size,
cl_context context, cl_device_id device, cl_program program) {
cl_int err; cl_int err;
cl_kernel kernel; cl_kernel kernel;
cl_command_queue queue; cl_command_queue queue;
@ -284,12 +316,18 @@ std::vector<unsigned int> dither::internal::blue_noise_cl_impl(
queue = clCreateCommandQueueWithProperties(context, device, nullptr, &err); queue = clCreateCommandQueueWithProperties(context, device, nullptr, &err);
d_filter_out = clCreateBuffer(context, CL_MEM_WRITE_ONLY, count * sizeof(float), nullptr, nullptr); d_filter_out = clCreateBuffer(context, CL_MEM_WRITE_ONLY,
d_precomputed = clCreateBuffer(context, CL_MEM_READ_ONLY, precomputed.size() * sizeof(float), nullptr, nullptr); count * sizeof(float), nullptr, nullptr);
d_pbp = clCreateBuffer(context, CL_MEM_READ_ONLY, count * sizeof(int), nullptr, nullptr); d_precomputed = clCreateBuffer(context, CL_MEM_READ_ONLY,
precomputed.size() * sizeof(float), nullptr,
nullptr);
d_pbp = clCreateBuffer(context, CL_MEM_READ_ONLY, count * sizeof(int),
nullptr, nullptr);
err = clEnqueueWriteBuffer(queue, d_precomputed, CL_TRUE, 0, precomputed.size() * sizeof(float), &precomputed[0], 0, nullptr, nullptr); err = clEnqueueWriteBuffer(queue, d_precomputed, CL_TRUE, 0,
if(err != CL_SUCCESS) { precomputed.size() * sizeof(float),
&precomputed[0], 0, nullptr, nullptr);
if (err != CL_SUCCESS) {
std::cerr << "OpenCL: Failed to write to d_precomputed buffer\n"; std::cerr << "OpenCL: Failed to write to d_precomputed buffer\n";
clReleaseMemObject(d_pbp); clReleaseMemObject(d_pbp);
clReleaseMemObject(d_precomputed); clReleaseMemObject(d_precomputed);
@ -299,9 +337,9 @@ std::vector<unsigned int> dither::internal::blue_noise_cl_impl(
} }
kernel = clCreateKernel(program, "do_filter", &err); kernel = clCreateKernel(program, "do_filter", &err);
if(err != CL_SUCCESS) { if (err != CL_SUCCESS) {
std::cerr << "OpenCL: Failed to create kernel: "; std::cerr << "OpenCL: Failed to create kernel: ";
switch(err) { switch (err) {
case CL_INVALID_PROGRAM: case CL_INVALID_PROGRAM:
std::cerr << "invalid program\n"; std::cerr << "invalid program\n";
break; break;
@ -334,7 +372,8 @@ std::vector<unsigned int> dither::internal::blue_noise_cl_impl(
return {}; return {};
} }
if(clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_filter_out) != CL_SUCCESS) { if (clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_filter_out) !=
CL_SUCCESS) {
std::cerr << "OpenCL: Failed to set kernel arg 0\n"; std::cerr << "OpenCL: Failed to set kernel arg 0\n";
clReleaseKernel(kernel); clReleaseKernel(kernel);
clReleaseMemObject(d_pbp); clReleaseMemObject(d_pbp);
@ -343,7 +382,8 @@ std::vector<unsigned int> dither::internal::blue_noise_cl_impl(
clReleaseCommandQueue(queue); clReleaseCommandQueue(queue);
return {}; return {};
} }
if(clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_precomputed) != CL_SUCCESS) { if (clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_precomputed) !=
CL_SUCCESS) {
std::cerr << "OpenCL: Failed to set kernel arg 1\n"; std::cerr << "OpenCL: Failed to set kernel arg 1\n";
clReleaseKernel(kernel); clReleaseKernel(kernel);
clReleaseMemObject(d_pbp); clReleaseMemObject(d_pbp);
@ -352,7 +392,7 @@ std::vector<unsigned int> dither::internal::blue_noise_cl_impl(
clReleaseCommandQueue(queue); clReleaseCommandQueue(queue);
return {}; return {};
} }
if(clSetKernelArg(kernel, 2, sizeof(cl_mem), &d_pbp) != CL_SUCCESS) { if (clSetKernelArg(kernel, 2, sizeof(cl_mem), &d_pbp) != CL_SUCCESS) {
std::cerr << "OpenCL: Failed to set kernel arg 2\n"; std::cerr << "OpenCL: Failed to set kernel arg 2\n";
clReleaseKernel(kernel); clReleaseKernel(kernel);
clReleaseMemObject(d_pbp); clReleaseMemObject(d_pbp);
@ -361,7 +401,7 @@ std::vector<unsigned int> dither::internal::blue_noise_cl_impl(
clReleaseCommandQueue(queue); clReleaseCommandQueue(queue);
return {}; return {};
} }
if(clSetKernelArg(kernel, 3, sizeof(int), &width) != CL_SUCCESS) { if (clSetKernelArg(kernel, 3, sizeof(int), &width) != CL_SUCCESS) {
std::cerr << "OpenCL: Failed to set kernel arg 3\n"; std::cerr << "OpenCL: Failed to set kernel arg 3\n";
clReleaseKernel(kernel); clReleaseKernel(kernel);
clReleaseMemObject(d_pbp); clReleaseMemObject(d_pbp);
@ -370,7 +410,7 @@ std::vector<unsigned int> dither::internal::blue_noise_cl_impl(
clReleaseCommandQueue(queue); clReleaseCommandQueue(queue);
return {}; return {};
} }
if(clSetKernelArg(kernel, 4, sizeof(int), &height) != CL_SUCCESS) { if (clSetKernelArg(kernel, 4, sizeof(int), &height) != CL_SUCCESS) {
std::cerr << "OpenCL: Failed to set kernel arg 4\n"; std::cerr << "OpenCL: Failed to set kernel arg 4\n";
clReleaseKernel(kernel); clReleaseKernel(kernel);
clReleaseMemObject(d_pbp); clReleaseMemObject(d_pbp);
@ -381,7 +421,8 @@ std::vector<unsigned int> dither::internal::blue_noise_cl_impl(
} }
if (filter_size % 2 == 0) { if (filter_size % 2 == 0) {
int filter_size_odd = filter_size + 1; int filter_size_odd = filter_size + 1;
if(clSetKernelArg(kernel, 5, sizeof(int), &filter_size_odd) != CL_SUCCESS) { if (clSetKernelArg(kernel, 5, sizeof(int), &filter_size_odd) !=
CL_SUCCESS) {
std::cerr << "OpenCL: Failed to set kernel arg 4\n"; std::cerr << "OpenCL: Failed to set kernel arg 4\n";
clReleaseKernel(kernel); clReleaseKernel(kernel);
clReleaseMemObject(d_pbp); clReleaseMemObject(d_pbp);
@ -391,7 +432,8 @@ std::vector<unsigned int> dither::internal::blue_noise_cl_impl(
return {}; return {};
} }
} else { } else {
if(clSetKernelArg(kernel, 5, sizeof(int), &filter_size) != CL_SUCCESS) { if (clSetKernelArg(kernel, 5, sizeof(int), &filter_size) !=
CL_SUCCESS) {
std::cerr << "OpenCL: Failed to set kernel arg 4\n"; std::cerr << "OpenCL: Failed to set kernel arg 4\n";
clReleaseKernel(kernel); clReleaseKernel(kernel);
clReleaseMemObject(d_pbp); clReleaseMemObject(d_pbp);
@ -402,7 +444,9 @@ std::vector<unsigned int> dither::internal::blue_noise_cl_impl(
} }
} }
if(clGetKernelWorkGroupInfo(kernel, device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(std::size_t), &local_size, nullptr) != CL_SUCCESS) { if (clGetKernelWorkGroupInfo(kernel, device, CL_KERNEL_WORK_GROUP_SIZE,
sizeof(std::size_t), &local_size,
nullptr) != CL_SUCCESS) {
std::cerr << "OpenCL: Failed to get work group size\n"; std::cerr << "OpenCL: Failed to get work group size\n";
clReleaseKernel(kernel); clReleaseKernel(kernel);
clReleaseMemObject(d_pbp); clReleaseMemObject(d_pbp);
@ -411,34 +455,39 @@ std::vector<unsigned int> dither::internal::blue_noise_cl_impl(
clReleaseCommandQueue(queue); clReleaseCommandQueue(queue);
return {}; return {};
} }
global_size = (std::size_t)std::ceil(count / (float)local_size) * local_size; global_size =
(std::size_t)std::ceil(count / (float)local_size) * local_size;
std::cout << "OpenCL: global = " << global_size << ", local = " << local_size std::cout << "OpenCL: global = " << global_size
<< std::endl; << ", local = " << local_size << std::endl;
std::vector<float> filter(count); std::vector<float> filter(count);
bool reversed_pbp = false; bool reversed_pbp = false;
const auto get_filter = [&queue, &kernel, &global_size, &local_size, const auto get_filter = [&queue, &kernel, &global_size, &local_size,
&d_filter_out, &d_pbp, &pbp, &pbp_i, &count, &filter, &err, &reversed_pbp] () -> bool { &d_filter_out, &d_pbp, &pbp, &pbp_i, &count,
for(unsigned int i = 0; i < pbp.size(); ++i) { &filter, &err, &reversed_pbp]() -> bool {
for (unsigned int i = 0; i < pbp.size(); ++i) {
if (reversed_pbp) { if (reversed_pbp) {
pbp_i[i] = pbp[i] ? 0 : 1; pbp_i[i] = pbp[i] ? 0 : 1;
} else { } else {
pbp_i[i] = pbp[i] ? 1 : 0; pbp_i[i] = pbp[i] ? 1 : 0;
} }
} }
if(clEnqueueWriteBuffer(queue, d_pbp, CL_TRUE, 0, count * sizeof(int), &pbp_i[0], 0, nullptr, nullptr) != CL_SUCCESS) { if (clEnqueueWriteBuffer(queue, d_pbp, CL_TRUE, 0, count * sizeof(int),
&pbp_i[0], 0, nullptr,
nullptr) != CL_SUCCESS) {
std::cerr << "OpenCL: Failed to write to d_pbp buffer\n"; std::cerr << "OpenCL: Failed to write to d_pbp buffer\n";
return false; return false;
} }
if(err = clEnqueueNDRangeKernel( if (err = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr,
queue, kernel, 1, nullptr, &global_size, &local_size, &global_size, &local_size, 0, nullptr,
0, nullptr, nullptr); err != CL_SUCCESS) { nullptr);
err != CL_SUCCESS) {
std::cerr << "OpenCL: Failed to enqueue task: "; std::cerr << "OpenCL: Failed to enqueue task: ";
switch(err) { switch (err) {
case CL_INVALID_PROGRAM_EXECUTABLE: case CL_INVALID_PROGRAM_EXECUTABLE:
std::cerr << "invalid program executable\n"; std::cerr << "invalid program executable\n";
break; break;
@ -481,20 +530,24 @@ std::vector<unsigned int> dither::internal::blue_noise_cl_impl(
clFinish(queue); clFinish(queue);
clEnqueueReadBuffer(queue, d_filter_out, CL_TRUE, 0, count * sizeof(float), &filter[0], 0, nullptr, nullptr); clEnqueueReadBuffer(queue, d_filter_out, CL_TRUE, 0,
count * sizeof(float), &filter[0], 0, nullptr,
nullptr);
return true; return true;
}; };
{ {
#ifndef NDEBUG #ifndef NDEBUG
printf("Inserting %d pixels into image of max count %d\n", pixel_count, count); printf("Inserting %d pixels into image of max count %d\n", pixel_count,
count);
// generate image from randomized pbp // generate image from randomized pbp
FILE *random_noise_image = fopen("random_noise.pbm", "w"); FILE *random_noise_image = fopen("random_noise.pbm", "w");
fprintf(random_noise_image, "P1\n%d %d\n", width, height); fprintf(random_noise_image, "P1\n%d %d\n", width, height);
for(int y = 0; y < height; ++y) { for (int y = 0; y < height; ++y) {
for(int x = 0; x < width; ++x) { for (int x = 0; x < width; ++x) {
fprintf(random_noise_image, "%d ", pbp[utility::twoToOne(x, y, width, height)] ? 1 : 0); fprintf(random_noise_image, "%d ",
pbp[utility::twoToOne(x, y, width, height)] ? 1 : 0);
} }
fputc('\n', random_noise_image); fputc('\n', random_noise_image);
} }
@ -502,7 +555,7 @@ std::vector<unsigned int> dither::internal::blue_noise_cl_impl(
#endif #endif
} }
if(!get_filter()) { if (!get_filter()) {
std::cerr << "OpenCL: Failed to execute do_filter (at start)\n"; std::cerr << "OpenCL: Failed to execute do_filter (at start)\n";
clReleaseKernel(kernel); clReleaseKernel(kernel);
clReleaseMemObject(d_pbp); clReleaseMemObject(d_pbp);
@ -519,12 +572,12 @@ std::vector<unsigned int> dither::internal::blue_noise_cl_impl(
int iterations = 0; int iterations = 0;
std::cout << "Begin BinaryArray generation loop\n"; std::cout << "Begin BinaryArray generation loop\n";
while(true) { while (true) {
#ifndef NDEBUG #ifndef NDEBUG
printf("Iteration %d\n", ++iterations); printf("Iteration %d\n", ++iterations);
#endif #endif
if(!get_filter()) { if (!get_filter()) {
std::cerr << "OpenCL: Failed to execute do_filter\n"; std::cerr << "OpenCL: Failed to execute do_filter\n";
break; break;
} }
@ -534,32 +587,35 @@ std::vector<unsigned int> dither::internal::blue_noise_cl_impl(
pbp[max] = false; pbp[max] = false;
if(!get_filter()) { if (!get_filter()) {
std::cerr << "OpenCL: Failed to execute do_filter\n"; std::cerr << "OpenCL: Failed to execute do_filter\n";
break; break;
} }
// get second buffer's min // get second buffer's min
int second_min; int second_min;
std::tie(second_min, std::ignore) = internal::filter_minmax(filter, pbp); std::tie(second_min, std::ignore) =
internal::filter_minmax(filter, pbp);
if(second_min == max) { if (second_min == max) {
pbp[max] = true; pbp[max] = true;
break; break;
} else { } else {
pbp[second_min] = true; pbp[second_min] = true;
} }
if(iterations % 100 == 0) { if (iterations % 100 == 0) {
#ifndef NDEBUG #ifndef NDEBUG
std::cout << "max was " << max << ", second_min is " << second_min std::cout << "max was " << max << ", second_min is " << second_min
<< std::endl; << std::endl;
// generate blue_noise image from pbp // generate blue_noise image from pbp
FILE *blue_noise_image = fopen("blue_noise.pbm", "w"); FILE *blue_noise_image = fopen("blue_noise.pbm", "w");
fprintf(blue_noise_image, "P1\n%d %d\n", width, height); fprintf(blue_noise_image, "P1\n%d %d\n", width, height);
for(int y = 0; y < height; ++y) { for (int y = 0; y < height; ++y) {
for(int x = 0; x < width; ++x) { for (int x = 0; x < width; ++x) {
fprintf(blue_noise_image, "%d ", pbp[utility::twoToOne(x, y, width, height)] ? 1 : 0); fprintf(blue_noise_image, "%d ",
pbp[utility::twoToOne(x, y, width, height)] ? 1
: 0);
} }
fputc('\n', blue_noise_image); fputc('\n', blue_noise_image);
} }
@ -568,16 +624,17 @@ std::vector<unsigned int> dither::internal::blue_noise_cl_impl(
} }
} }
if(!get_filter()) { if (!get_filter()) {
std::cerr << "OpenCL: Failed to execute do_filter (at end)\n"; std::cerr << "OpenCL: Failed to execute do_filter (at end)\n";
} else { } else {
#ifndef NDEBUG #ifndef NDEBUG
internal::write_filter(filter, width, "filter_out_final.pgm"); internal::write_filter(filter, width, "filter_out_final.pgm");
FILE *blue_noise_image = fopen("blue_noise.pbm", "w"); FILE *blue_noise_image = fopen("blue_noise.pbm", "w");
fprintf(blue_noise_image, "P1\n%d %d\n", width, height); fprintf(blue_noise_image, "P1\n%d %d\n", width, height);
for(int y = 0; y < height; ++y) { for (int y = 0; y < height; ++y) {
for(int x = 0; x < width; ++x) { for (int x = 0; x < width; ++x) {
fprintf(blue_noise_image, "%d ", pbp[utility::twoToOne(x, y, width, height)] ? 1 : 0); fprintf(blue_noise_image, "%d ",
pbp[utility::twoToOne(x, y, width, height)] ? 1 : 0);
} }
fputc('\n', blue_noise_image); fputc('\n', blue_noise_image);
} }
@ -588,7 +645,8 @@ std::vector<unsigned int> dither::internal::blue_noise_cl_impl(
#ifndef NDEBUG #ifndef NDEBUG
{ {
image::Bl pbp_image = toBl(pbp, width); image::Bl pbp_image = toBl(pbp, width);
pbp_image.writeToFile(image::file_type::PNG, true, "debug_pbp_before.png"); pbp_image.writeToFile(image::file_type::PNG, true,
"debug_pbp_before.png");
} }
#endif #endif
@ -620,11 +678,13 @@ std::vector<unsigned int> dither::internal::blue_noise_cl_impl(
pbp = pbp_copy; pbp = pbp_copy;
#ifndef NDEBUG #ifndef NDEBUG
image::Bl min_pixels = internal::rangeToBl(dither_array, width); image::Bl min_pixels = internal::rangeToBl(dither_array, width);
min_pixels.writeToFile(image::file_type::PNG, true, "da_min_pixels.png"); min_pixels.writeToFile(image::file_type::PNG, true,
"da_min_pixels.png");
#endif #endif
} }
std::cout << "\nRanking remainder of first half of pixels...\n"; std::cout << "\nRanking remainder of first half of pixels...\n";
for (unsigned int i = pixel_count; i < (unsigned int)((count + 1) / 2); ++i) { for (unsigned int i = pixel_count; i < (unsigned int)((count + 1) / 2);
++i) {
#ifndef NDEBUG #ifndef NDEBUG
std::cout << i << ' '; std::cout << i << ' ';
#endif #endif
@ -643,7 +703,8 @@ std::vector<unsigned int> dither::internal::blue_noise_cl_impl(
#ifndef NDEBUG #ifndef NDEBUG
{ {
image::Bl min_pixels = internal::rangeToBl(dither_array, width); image::Bl min_pixels = internal::rangeToBl(dither_array, width);
min_pixels.writeToFile(image::file_type::PNG, true, "da_mid_pixels.png"); min_pixels.writeToFile(image::file_type::PNG, true,
"da_mid_pixels.png");
get_filter(); get_filter();
internal::write_filter(filter, width, "filter_mid.pgm"); internal::write_filter(filter, width, "filter_mid.pgm");
image::Bl pbp_image = toBl(pbp, width); image::Bl pbp_image = toBl(pbp, width);
@ -675,7 +736,8 @@ std::vector<unsigned int> dither::internal::blue_noise_cl_impl(
get_filter(); get_filter();
internal::write_filter(filter, width, "filter_after.pgm"); internal::write_filter(filter, width, "filter_after.pgm");
image::Bl pbp_image = toBl(pbp, width); image::Bl pbp_image = toBl(pbp, width);
pbp_image.writeToFile(image::file_type::PNG, true, "debug_pbp_after.png"); pbp_image.writeToFile(image::file_type::PNG, true,
"debug_pbp_after.png");
} }
#endif #endif
@ -686,3 +748,4 @@ std::vector<unsigned int> dither::internal::blue_noise_cl_impl(
clReleaseCommandQueue(queue); clReleaseCommandQueue(queue);
return dither_array; return dither_array;
} }
#endif