]> git.seodisparate.com - blue_noise_generation/commitdiff
Make usage of OpenCL conditional (configuring)
authorStephen Seo <seo.disparate@gmail.com>
Fri, 31 Mar 2023 09:46:10 +0000 (18:46 +0900)
committerStephen Seo <seo.disparate@gmail.com>
Fri, 31 Mar 2023 09:46:10 +0000 (18:46 +0900)
CMakeLists.txt
src/blue_noise.cpp

index ac66099ba684956d894d16f6565774f97962e51e..63df9d13a413b52ac35465d0b17471ca092c6be7 100644 (file)
@@ -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()
index 95fca220f46464c9c55e03b54084896e36150d34..f372fd54df04a196ae2ad6fe83d699b088aa0be5 100644 (file)
@@ -1,24 +1,26 @@
 #include "blue_noise.hpp"
 
-#include <random>
 #include <cassert>
-#include <iostream>
+#include <cstdio>
 #include <fstream>
+#include <iostream>
 #include <memory>
+#include <random>
 #include <string>
 #include <unordered_set>
-#include <cstdio>
 
+#if DITHERING_OPENCL_ENABLED == 1
 #include <CL/opencl.h>
+#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<char[]> log = std::make_unique<char[]>(log_size + 1);
+                    clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG,
+                                          0, nullptr, &log_size);
+                    std::unique_ptr<char[]> log =
+                            std::make_unique<char[]>(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<unsigned int> 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<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;
     std::vector<float> filter_out;
     filter_out.resize(count);
@@ -119,50 +138,54 @@ std::vector<unsigned int> 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<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,
-            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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> dither::internal::blue_noise_impl(int width, int heigh
     return dither_array;
 }
 
+#if DITHERING_OPENCL_ENABLED == 1
 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_kernel kernel;
     cl_command_queue queue;
@@ -284,12 +316,18 @@ std::vector<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<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";
         clReleaseKernel(kernel);
         clReleaseMemObject(d_pbp);
@@ -411,90 +455,99 @@ std::vector<unsigned int> 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<float> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<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";
     } 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> 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<unsigned int> dither::internal::blue_noise_cl_impl(
     clReleaseCommandQueue(queue);
     return dither_array;
 }
+#endif