From 62d15771ad2fd5e64040d6ad5d00d5025f5770d0 Mon Sep 17 00:00:00 2001 From: Stephen Seo Date: Fri, 31 Mar 2023 18:46:10 +0900 Subject: [PATCH] Make usage of OpenCL conditional (configuring) --- CMakeLists.txt | 37 +++- src/blue_noise.cpp | 421 ++++++++++++++++++++++++++------------------- 2 files changed, 270 insertions(+), 188 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index ac66099..63df9d1 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -18,16 +18,35 @@ if(NOT DEFINED CMAKE_BUILD_TYPE OR NOT CMAKE_BUILD_TYPE) endif() 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) add_executable(Dithering ${Dithering_SOURCES}) target_compile_features(Dithering PUBLIC cxx_std_17) -target_include_directories(Dithering PUBLIC - Threads::Threads - ${OpenCL_INCLUDE_DIRS} - ${PNG_INCLUDE_DIRS}) -target_link_libraries(Dithering PUBLIC - Threads::Threads - ${OpenCL_LIBRARIES} - ${PNG_LIBRARIES}) +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 + ${OpenCL_INCLUDE_DIRS} + ${PNG_INCLUDE_DIRS}) + target_link_libraries(Dithering PUBLIC + Threads::Threads + ${OpenCL_LIBRARIES} + ${PNG_LIBRARIES}) + target_compile_definitions(Dithering PRIVATE DITHERING_OPENCL_ENABLED=1) +endif() diff --git a/src/blue_noise.cpp b/src/blue_noise.cpp index 95fca22..f372fd5 100644 --- a/src/blue_noise.cpp +++ b/src/blue_noise.cpp @@ -1,24 +1,26 @@ #include "blue_noise.hpp" -#include #include -#include +#include #include +#include #include +#include #include #include -#include +#if DITHERING_OPENCL_ENABLED == 1 #include +#endif #include "image.hpp" - -image::Bl dither::blue_noise(int width, int height, int threads, bool use_opencl) { - +image::Bl dither::blue_noise(int width, int height, int threads, + bool use_opencl) { bool using_opencl = false; - if(use_opencl) { +#if DITHERING_OPENCL_ENABLED == 1 + if (use_opencl) { // try to use OpenCL do { 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; err = clGetPlatformIDs(1, &platform, nullptr); - if(err != CL_SUCCESS) { + if (err != CL_SUCCESS) { std::cerr << "OpenCL: Failed to identify a platform\n"; break; } - err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, nullptr); - if(err != CL_SUCCESS) { + err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, + nullptr); + if (err != CL_SUCCESS) { std::cerr << "OpenCL: Failed to get a device\n"; break; } - context = clCreateContext(nullptr, 1, &device, nullptr, nullptr, &err); + context = clCreateContext(nullptr, 1, &device, nullptr, nullptr, + &err); { char buf[1024]; std::ifstream program_file("src/blue_noise.cl"); 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; } std::string program_string; - while(program_file.good()) { + while (program_file.good()) { 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); } } const char *string_ptr = program_string.c_str(); std::size_t program_size = program_string.size(); - program = clCreateProgramWithSource(context, 1, (const char**)&string_ptr, &program_size, &err); - if(err != CL_SUCCESS) { + program = clCreateProgramWithSource(context, 1, + (const char **)&string_ptr, + &program_size, &err); + if (err != CL_SUCCESS) { std::cerr << "OpenCL: Failed to create the program\n"; clReleaseContext(context); break; } - err = clBuildProgram(program, 1, &device, nullptr, nullptr, nullptr); - if(err != CL_SUCCESS) { + err = clBuildProgram(program, 1, &device, nullptr, nullptr, + nullptr); + if (err != CL_SUCCESS) { std::cerr << "OpenCL: Failed to build the program\n"; std::size_t log_size; - clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, nullptr, &log_size); - std::unique_ptr log = std::make_unique(log_size + 1); + clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, + 0, nullptr, &log_size); + std::unique_ptr log = + std::make_unique(log_size + 1); 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; clReleaseProgram(program); @@ -87,29 +99,36 @@ image::Bl dither::blue_noise(int width, int height, int threads, bool use_opencl std::cout << "OpenCL: Initialized, trying cl_impl..." << std::endl; std::vector result = internal::blue_noise_cl_impl( - width, height, filter_size, context, device, program); + width, height, filter_size, context, device, program); clReleaseProgram(program); clReleaseContext(context); - if(!result.empty()) { + if (!result.empty()) { return internal::rangeToBl(result, width); } std::cout << "ERROR: Empty result\n"; } while (false); } +#else + std::clog << "WARNING: Not compiled with OpenCL support!\n"; +#endif - if(!using_opencl) { - std::cout << "OpenCL: Failed to setup/use or is not enabled, using regular impl..." - << std::endl; - return internal::rangeToBl(internal::blue_noise_impl(width, height, threads), width); + if (!using_opencl) { + std::cout << "OpenCL: Failed to setup/use or is not enabled, using " + "regular impl..." + << std::endl; + return internal::rangeToBl( + internal::blue_noise_impl(width, height, threads), width); } std::cout << "ERROR: Invalid state (end of blue_noise fn)\n"; return {}; } -std::vector dither::internal::blue_noise_impl(int width, int height, int threads) { +std::vector dither::internal::blue_noise_impl(int width, + int height, + int threads) { int count = width * height; std::vector filter_out; filter_out.resize(count); @@ -119,50 +138,54 @@ std::vector dither::internal::blue_noise_impl(int width, int heigh pbp.resize(count); #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 FILE *random_noise_image = fopen("random_noise.pbm", "w"); fprintf(random_noise_image, "P1\n%d %d\n", width, height); - for(int y = 0; y < height; ++y) { - for(int x = 0; x < width; ++x) { - fprintf(random_noise_image, "%d ", pbp[utility::twoToOne(x, y, width, height)] ? 1 : 0); + for (int y = 0; y < height; ++y) { + for (int x = 0; x < width; ++x) { + fprintf(random_noise_image, "%d ", + pbp[utility::twoToOne(x, y, width, height)] ? 1 : 0); } fputc('\n', random_noise_image); } fclose(random_noise_image); #endif -//#ifndef NDEBUG + // #ifndef NDEBUG int iterations = 0; -//#endif + // #endif int filter_size = (width + height) / 2; - std::unique_ptr> precomputed = std::make_unique>(internal::precompute_gaussian(filter_size)); + std::unique_ptr> precomputed = + std::make_unique>( + internal::precompute_gaussian(filter_size)); - internal::compute_filter(pbp, width, height, count, filter_size, - filter_out, precomputed.get(), threads); + internal::compute_filter(pbp, width, height, count, filter_size, filter_out, + precomputed.get(), threads); #ifndef NDEBUG internal::write_filter(filter_out, width, "filter_out_start.pgm"); #endif std::cout << "Begin BinaryArray generation loop\n"; - while(true) { + while (true) { #ifndef NDEBUG -// if(++iterations % 10 == 0) { - printf("Iteration %d\n", ++iterations); + // if(++iterations % 10 == 0) { + printf("Iteration %d\n", ++iterations); // } #endif // get filter values internal::compute_filter(pbp, width, height, count, filter_size, - filter_out, precomputed.get(), threads); + filter_out, precomputed.get(), threads); -//#ifndef NDEBUG -// for(int i = 0; i < count; ++i) { -// int x, y; -// std::tie(x, y) = internal::oneToTwo(i, width); -// printf("%d (%d, %d): %f\n", i, x, y, filter_out[i]); -// } -//#endif + // #ifndef NDEBUG + // for(int i = 0; i < count; ++i) { + // int x, y; + // std::tie(x, y) = internal::oneToTwo(i, width); + // printf("%d (%d, %d): %f\n", i, x, y, filter_out[i]); + // } + // #endif int min, max; std::tie(min, max) = internal::filter_minmax(filter_out, pbp); @@ -172,27 +195,30 @@ std::vector dither::internal::blue_noise_impl(int width, int heigh // get filter values again internal::compute_filter(pbp, width, height, count, filter_size, - filter_out, precomputed.get(), threads); + filter_out, precomputed.get(), threads); // get second buffer's 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; break; } else { pbp[second_min] = true; } - if(iterations % 100 == 0) { + if (iterations % 100 == 0) { // generate blue_noise image from pbp #ifndef NDEBUG FILE *blue_noise_image = fopen("blue_noise.pbm", "w"); fprintf(blue_noise_image, "P1\n%d %d\n", width, height); - for(int y = 0; y < height; ++y) { - for(int x = 0; x < width; ++x) { - fprintf(blue_noise_image, "%d ", pbp[utility::twoToOne(x, y, width, height)] ? 1 : 0); + for (int y = 0; y < height; ++y) { + for (int x = 0; x < width; ++x) { + fprintf(blue_noise_image, "%d ", + pbp[utility::twoToOne(x, y, width, height)] ? 1 + : 0); } fputc('\n', blue_noise_image); } @@ -200,8 +226,8 @@ std::vector dither::internal::blue_noise_impl(int width, int heigh #endif } } - internal::compute_filter(pbp, width, height, count, filter_size, - filter_out, precomputed.get(), threads); + internal::compute_filter(pbp, width, height, count, filter_size, filter_out, + precomputed.get(), threads); #ifndef NDEBUG internal::write_filter(filter_out, width, "filter_out_final.pgm"); #endif @@ -210,9 +236,10 @@ std::vector dither::internal::blue_noise_impl(int width, int heigh // generate blue_noise image from pbp FILE *blue_noise_image = fopen("blue_noise.pbm", "w"); fprintf(blue_noise_image, "P1\n%d %d\n", width, height); - for(int y = 0; y < height; ++y) { - for(int x = 0; x < width; ++x) { - fprintf(blue_noise_image, "%d ", pbp[utility::twoToOne(x, y, width, height)] ? 1 : 0); + for (int y = 0; y < height; ++y) { + for (int x = 0; x < width; ++x) { + fprintf(blue_noise_image, "%d ", + pbp[utility::twoToOne(x, y, width, height)] ? 1 : 0); } fputc('\n', blue_noise_image); } @@ -230,20 +257,22 @@ std::vector dither::internal::blue_noise_impl(int width, int heigh std::cout << i << ' '; #endif internal::compute_filter(pbp, width, height, count, filter_size, - filter_out, precomputed.get(), threads); - std::tie(std::ignore, max) = internal::filter_minmax(filter_out, pbp); + filter_out, precomputed.get(), threads); + std::tie(std::ignore, max) = + internal::filter_minmax(filter_out, pbp); pbp[max] = false; dither_array[max] = i; } pbp = pbp_copy; } 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 std::cout << i << ' '; #endif internal::compute_filter(pbp, width, height, count, filter_size, - filter_out, precomputed.get(), threads); + filter_out, precomputed.get(), threads); std::tie(min, std::ignore) = internal::filter_minmax(filter_out, pbp); pbp[min] = true; dither_array[min] = i; @@ -254,11 +283,12 @@ std::vector dither::internal::blue_noise_impl(int width, int heigh #ifndef NDEBUG std::cout << i << ' '; #endif - for(unsigned int i = 0; i < pbp.size(); ++i) { + for (unsigned int i = 0; i < pbp.size(); ++i) { reversed_pbp[i] = !pbp[i]; } - internal::compute_filter(reversed_pbp, width, height, count, filter_size, - filter_out, precomputed.get(), threads); + internal::compute_filter(reversed_pbp, width, height, count, + filter_size, filter_out, precomputed.get(), + threads); std::tie(std::ignore, max) = internal::filter_minmax(filter_out, pbp); pbp[max] = true; dither_array[max] = i; @@ -267,8 +297,10 @@ std::vector dither::internal::blue_noise_impl(int width, int heigh return dither_array; } +#if DITHERING_OPENCL_ENABLED == 1 std::vector 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_kernel kernel; cl_command_queue queue; @@ -284,12 +316,18 @@ std::vector dither::internal::blue_noise_cl_impl( queue = clCreateCommandQueueWithProperties(context, device, nullptr, &err); - d_filter_out = clCreateBuffer(context, CL_MEM_WRITE_ONLY, count * sizeof(float), 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); - if(err != CL_SUCCESS) { + d_filter_out = clCreateBuffer(context, CL_MEM_WRITE_ONLY, + count * sizeof(float), 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); + if (err != CL_SUCCESS) { std::cerr << "OpenCL: Failed to write to d_precomputed buffer\n"; clReleaseMemObject(d_pbp); clReleaseMemObject(d_precomputed); @@ -299,33 +337,33 @@ std::vector dither::internal::blue_noise_cl_impl( } kernel = clCreateKernel(program, "do_filter", &err); - if(err != CL_SUCCESS) { + if (err != CL_SUCCESS) { std::cerr << "OpenCL: Failed to create kernel: "; - switch(err) { - case CL_INVALID_PROGRAM: - std::cerr << "invalid program\n"; - break; - case CL_INVALID_PROGRAM_EXECUTABLE: - std::cerr << "invalid program executable\n"; - break; - case CL_INVALID_KERNEL_NAME: - std::cerr << "invalid kernel name\n"; - break; - case CL_INVALID_KERNEL_DEFINITION: - std::cerr << "invalid kernel definition\n"; - break; - case CL_INVALID_VALUE: - std::cerr << "invalid value\n"; - break; - case CL_OUT_OF_RESOURCES: - std::cerr << "out of resources\n"; - break; - case CL_OUT_OF_HOST_MEMORY: - std::cerr << "out of host memory\n"; - break; - default: - std::cerr << "unknown error\n"; - break; + switch (err) { + case CL_INVALID_PROGRAM: + std::cerr << "invalid program\n"; + break; + case CL_INVALID_PROGRAM_EXECUTABLE: + std::cerr << "invalid program executable\n"; + break; + case CL_INVALID_KERNEL_NAME: + std::cerr << "invalid kernel name\n"; + break; + case CL_INVALID_KERNEL_DEFINITION: + std::cerr << "invalid kernel definition\n"; + break; + case CL_INVALID_VALUE: + std::cerr << "invalid value\n"; + break; + case CL_OUT_OF_RESOURCES: + std::cerr << "out of resources\n"; + break; + case CL_OUT_OF_HOST_MEMORY: + std::cerr << "out of host memory\n"; + break; + default: + std::cerr << "unknown error\n"; + break; } clReleaseMemObject(d_pbp); clReleaseMemObject(d_precomputed); @@ -334,7 +372,8 @@ std::vector dither::internal::blue_noise_cl_impl( 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"; clReleaseKernel(kernel); clReleaseMemObject(d_pbp); @@ -343,7 +382,8 @@ std::vector dither::internal::blue_noise_cl_impl( clReleaseCommandQueue(queue); 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"; clReleaseKernel(kernel); clReleaseMemObject(d_pbp); @@ -352,7 +392,7 @@ std::vector dither::internal::blue_noise_cl_impl( clReleaseCommandQueue(queue); 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"; clReleaseKernel(kernel); clReleaseMemObject(d_pbp); @@ -361,7 +401,7 @@ std::vector dither::internal::blue_noise_cl_impl( clReleaseCommandQueue(queue); 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"; clReleaseKernel(kernel); clReleaseMemObject(d_pbp); @@ -370,7 +410,7 @@ std::vector dither::internal::blue_noise_cl_impl( clReleaseCommandQueue(queue); 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"; clReleaseKernel(kernel); clReleaseMemObject(d_pbp); @@ -381,7 +421,8 @@ std::vector dither::internal::blue_noise_cl_impl( } if (filter_size % 2 == 0) { 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"; clReleaseKernel(kernel); clReleaseMemObject(d_pbp); @@ -391,7 +432,8 @@ std::vector dither::internal::blue_noise_cl_impl( return {}; } } 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"; clReleaseKernel(kernel); clReleaseMemObject(d_pbp); @@ -402,7 +444,9 @@ std::vector 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"; clReleaseKernel(kernel); clReleaseMemObject(d_pbp); @@ -411,90 +455,99 @@ std::vector dither::internal::blue_noise_cl_impl( clReleaseCommandQueue(queue); 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::endl; + std::cout << "OpenCL: global = " << global_size + << ", local = " << local_size << std::endl; std::vector filter(count); bool reversed_pbp = false; const auto get_filter = [&queue, &kernel, &global_size, &local_size, - &d_filter_out, &d_pbp, &pbp, &pbp_i, &count, &filter, &err, &reversed_pbp] () -> bool { - for(unsigned int i = 0; i < pbp.size(); ++i) { + &d_filter_out, &d_pbp, &pbp, &pbp_i, &count, + &filter, &err, &reversed_pbp]() -> bool { + for (unsigned int i = 0; i < pbp.size(); ++i) { if (reversed_pbp) { pbp_i[i] = pbp[i] ? 0 : 1; } else { 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"; return false; } - if(err = clEnqueueNDRangeKernel( - queue, kernel, 1, nullptr, &global_size, &local_size, - 0, nullptr, nullptr); err != CL_SUCCESS) { + if (err = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, + &global_size, &local_size, 0, nullptr, + nullptr); + err != CL_SUCCESS) { std::cerr << "OpenCL: Failed to enqueue task: "; - switch(err) { - case CL_INVALID_PROGRAM_EXECUTABLE: - std::cerr << "invalid program executable\n"; - break; - case CL_INVALID_COMMAND_QUEUE: - std::cerr << "invalid command queue\n"; - break; - case CL_INVALID_KERNEL: - std::cerr << "invalid kernel\n"; - break; - case CL_INVALID_CONTEXT: - std::cerr << "invalid context\n"; - break; - case CL_INVALID_KERNEL_ARGS: - std::cerr << "invalid kernel args\n"; - break; - case CL_INVALID_WORK_DIMENSION: - std::cerr << "invalid work dimension\n"; - break; - case CL_INVALID_GLOBAL_WORK_SIZE: - std::cerr << "invalid global work size\n"; - break; - case CL_INVALID_GLOBAL_OFFSET: - std::cerr << "invalid global offset\n"; - break; - case CL_INVALID_WORK_GROUP_SIZE: - std::cerr << "invalid work group size\n"; - break; - case CL_INVALID_WORK_ITEM_SIZE: - std::cerr << "invalid work item size\n"; - break; - case CL_MISALIGNED_SUB_BUFFER_OFFSET: - std::cerr << "misaligned sub buffer offset\n"; - break; - default: - std::cerr << "Unknown\n"; - break; + switch (err) { + case CL_INVALID_PROGRAM_EXECUTABLE: + std::cerr << "invalid program executable\n"; + break; + case CL_INVALID_COMMAND_QUEUE: + std::cerr << "invalid command queue\n"; + break; + case CL_INVALID_KERNEL: + std::cerr << "invalid kernel\n"; + break; + case CL_INVALID_CONTEXT: + std::cerr << "invalid context\n"; + break; + case CL_INVALID_KERNEL_ARGS: + std::cerr << "invalid kernel args\n"; + break; + case CL_INVALID_WORK_DIMENSION: + std::cerr << "invalid work dimension\n"; + break; + case CL_INVALID_GLOBAL_WORK_SIZE: + std::cerr << "invalid global work size\n"; + break; + case CL_INVALID_GLOBAL_OFFSET: + std::cerr << "invalid global offset\n"; + break; + case CL_INVALID_WORK_GROUP_SIZE: + std::cerr << "invalid work group size\n"; + break; + case CL_INVALID_WORK_ITEM_SIZE: + std::cerr << "invalid work item size\n"; + break; + case CL_MISALIGNED_SUB_BUFFER_OFFSET: + std::cerr << "misaligned sub buffer offset\n"; + break; + default: + std::cerr << "Unknown\n"; + break; } return false; } 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; }; { #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 FILE *random_noise_image = fopen("random_noise.pbm", "w"); fprintf(random_noise_image, "P1\n%d %d\n", width, height); - for(int y = 0; y < height; ++y) { - for(int x = 0; x < width; ++x) { - fprintf(random_noise_image, "%d ", pbp[utility::twoToOne(x, y, width, height)] ? 1 : 0); + for (int y = 0; y < height; ++y) { + for (int x = 0; x < width; ++x) { + fprintf(random_noise_image, "%d ", + pbp[utility::twoToOne(x, y, width, height)] ? 1 : 0); } fputc('\n', random_noise_image); } @@ -502,7 +555,7 @@ std::vector dither::internal::blue_noise_cl_impl( #endif } - if(!get_filter()) { + if (!get_filter()) { std::cerr << "OpenCL: Failed to execute do_filter (at start)\n"; clReleaseKernel(kernel); clReleaseMemObject(d_pbp); @@ -519,12 +572,12 @@ std::vector dither::internal::blue_noise_cl_impl( int iterations = 0; std::cout << "Begin BinaryArray generation loop\n"; - while(true) { + while (true) { #ifndef NDEBUG printf("Iteration %d\n", ++iterations); #endif - if(!get_filter()) { + if (!get_filter()) { std::cerr << "OpenCL: Failed to execute do_filter\n"; break; } @@ -534,32 +587,35 @@ std::vector dither::internal::blue_noise_cl_impl( pbp[max] = false; - if(!get_filter()) { + if (!get_filter()) { std::cerr << "OpenCL: Failed to execute do_filter\n"; break; } // get second buffer's 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; break; } else { pbp[second_min] = true; } - if(iterations % 100 == 0) { + if (iterations % 100 == 0) { #ifndef NDEBUG std::cout << "max was " << max << ", second_min is " << second_min - << std::endl; + << std::endl; // generate blue_noise image from pbp FILE *blue_noise_image = fopen("blue_noise.pbm", "w"); fprintf(blue_noise_image, "P1\n%d %d\n", width, height); - for(int y = 0; y < height; ++y) { - for(int x = 0; x < width; ++x) { - fprintf(blue_noise_image, "%d ", pbp[utility::twoToOne(x, y, width, height)] ? 1 : 0); + for (int y = 0; y < height; ++y) { + for (int x = 0; x < width; ++x) { + fprintf(blue_noise_image, "%d ", + pbp[utility::twoToOne(x, y, width, height)] ? 1 + : 0); } fputc('\n', blue_noise_image); } @@ -568,16 +624,17 @@ std::vector dither::internal::blue_noise_cl_impl( } } - if(!get_filter()) { + if (!get_filter()) { std::cerr << "OpenCL: Failed to execute do_filter (at end)\n"; } else { #ifndef NDEBUG internal::write_filter(filter, width, "filter_out_final.pgm"); FILE *blue_noise_image = fopen("blue_noise.pbm", "w"); fprintf(blue_noise_image, "P1\n%d %d\n", width, height); - for(int y = 0; y < height; ++y) { - for(int x = 0; x < width; ++x) { - fprintf(blue_noise_image, "%d ", pbp[utility::twoToOne(x, y, width, height)] ? 1 : 0); + for (int y = 0; y < height; ++y) { + for (int x = 0; x < width; ++x) { + fprintf(blue_noise_image, "%d ", + pbp[utility::twoToOne(x, y, width, height)] ? 1 : 0); } fputc('\n', blue_noise_image); } @@ -588,7 +645,8 @@ std::vector dither::internal::blue_noise_cl_impl( #ifndef NDEBUG { 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 @@ -620,11 +678,13 @@ std::vector dither::internal::blue_noise_cl_impl( pbp = pbp_copy; #ifndef NDEBUG 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 } 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 std::cout << i << ' '; #endif @@ -643,7 +703,8 @@ std::vector dither::internal::blue_noise_cl_impl( #ifndef NDEBUG { 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(); internal::write_filter(filter, width, "filter_mid.pgm"); image::Bl pbp_image = toBl(pbp, width); @@ -675,7 +736,8 @@ std::vector dither::internal::blue_noise_cl_impl( get_filter(); internal::write_filter(filter, width, "filter_after.pgm"); 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 @@ -686,3 +748,4 @@ std::vector dither::internal::blue_noise_cl_impl( clReleaseCommandQueue(queue); return dither_array; } +#endif -- 2.49.0