From 3cece9ae1da91a57b7100c9a40be3d2881bc02ee Mon Sep 17 00:00:00 2001 From: Stephen Seo Date: Sat, 27 Nov 2021 13:13:27 +0900 Subject: [PATCH] Impl dithering on red/green/blue channels Also add res/blue_noise_64x64.png for use in dithering. --- res/blue_noise_64x64.png | Bin 0 -> 4228 bytes src/image.cc | 230 ++++++++++++++++++++++++++++++++++++++- src/image.h | 17 ++- src/main.cc | 3 +- 4 files changed, 247 insertions(+), 3 deletions(-) create mode 100644 res/blue_noise_64x64.png diff --git a/res/blue_noise_64x64.png b/res/blue_noise_64x64.png new file mode 100644 index 0000000000000000000000000000000000000000..3c1bdb54f2c36c479e8422df06e06edb4c9b360e GIT binary patch literal 4228 zcmV-~5PR>5P)QSxqe=Zf8kP-4X5^1GBywPVGsDnE_yJ$}PP+nxZ{3q#+COC@U+ z-;4vzZxcM;A@_kw1^BU!DbGVJqrl^JheM^{ww2QBBFJnnUHO_lfLP3o{jE-m_b|3e zaE(m>v4L602mPN_yE(^MF`WKUrE4$Dr*}Qv5pVjGdg)sQIlAl{Y%xrKXEu=u>`4%8 z$fh2(J!&N648K%;=LR`9Vke)vJSck@=Qr6^m_f&f547B7_OjNUyy-8`9JmDXhiUi@ zM}(>XULezKnFNeX#0z4$b=^{(Zuw#F2#r`+iPr~|V+ffkgEaGS z&uEfAs|kTI5^#krt6LQ~tpdA|qYCLnA)YiEfCCkbYUrThLwtN}>U{3~Uh`s-Drb?p z)ISIVge@Zq<%HY!uM~-Hsdw9FGW(WSt%5(Q<3-wZrdZn|4H-t@w(kHbmjTA2^bXyc zJFqFC3-l=7!7cC!tKkqv^JX^Uzb89^zEauoTF7vW9_2%iwrb@DwMWYcm_++30*-s;y zOx$f(>a$-E{ixGgIh9Gn03Ls?;#b)!X!rw*9)E4vwRGrR7C*xInYVo07|CIZQP*CV zvloTnm%2U1e=el&e8nB8G0TfJmq%uXBfa?&jQ~H=5R-KY%b_^HsYd&q7dDC?x1if! z1#TWXRsMFKF1dUJ#ytGXUMzYKRP=7p3SB*u_;&_R*tX^Yq2N+`rQ`s;Rx$9yNQLNC zB;aEax=ZoNQu}u#uSDR;iK!|E+Uy)d^l2k{Ob4m%)R`TwHHxy;4^g<|r5A2K$uw^m z(J*KL`iiI;W3@9Am2ZQ}Fx!1;E1Cy5kMew<6z^%jlS^i$(~qv-n!o05Mkj>8Vom!S zX6%SKVI_e0j$g9!mPfV@07`JrLz@0~)VuZ!pIDIyrsHM7-b@#}FkMO4JbcYHh6+?P z5n4GRk;Px@6Wf-(Eq0W0*{eEAErGIK)%e&@(Y-v@I;R+zU)pjX6WAIJdmuvxy^QRY00 zvC?N6c_~(n&kpY>q6~P(^lh|Mon0)dLK$YH0oqkShM>pxtTqZ~~)3W7!a$(7+kWFmdpVFY?H zV%Z4lN~?VdpF6b$ePXYJHTJcCo4I+xM6u!_N`2Y_(1Pcy*K13NA#(tbM#o>)@TlXs z7i%{Wxr^$>I-&T$O6`C%bk-xiW91%*>dHOv2+&6(>;_x@lWiwhkHGP;S-mqEcariI z!C(3SwiIuNE0c34W0KW@@^hvPM~@v=iV~Z+Ql<}>?JUY~OQsZ(P@-(VR;SfK8cy7z z_zYt!K_74UT*MUiERd=PlE1CC~Y0g+5b|uiYqi z(<|hF4w5f)imu85gFncqdY$EqqX*nnx^4=ZdjMmVt3T{&Sn}K4 zxh1PTU-pZ{gF1k}W^{$wX5cTrmwQKR2Ri00P>YCfb&`Urx68Yc;Qa!TD5uq*ok~p; z7^fZpiNEE|Dz8Ee!8P$f39LqEvOn{b5Gt77TI!xDR+%%kOyG>c+3ZuJf?nuQ)_|%= z+n>2-A2i{+5X?H?OX@Vr0RARwh6MJL=3LQ&w%1`J>JH9yyhiG97(xny{Sn(9^Kum` zSv3oC)bj(Byg>-hCspzX&6#q8OP*VRxCfDUQ~+f=vY${jZ?h_)4p1zf$9s}8SRkoj z%Z9YlWIC*9i=;lf?suw>yd6o%b|LtUX2N$hg-q(U1k^0+9nxy2*bV@d(Gu~-fZP#D z{&eVx-~mbZq}ql41)w_hB$31AMa>A~Wt_|>{Xl7pEU;RlHs=kLtl1T0LiLhvsYHt_ z^B%bX?MQemT^yatW|17iYByWF9A^;5HBsXnb*EPZcO!vVFo;J3Vb>C*@6kWv5PY** z`591UG@^zbU-b^nJMrT z)0+&0vuQU`4{DOfaV+VUBj2_mgZ?iVkZ+@b0bTwKQYN0{6Y-ik zYx8+a>>ZrYZArFb2uir=0b}!6H+&wHXs$;h3*|HS23w{wQ`Ur@b;$s$);64EM!t8` z3g2xu*-vu+g=R{7-cXVxuZP=EF#3}m=!TtbGo{cYm(``k~=$aYsns~PX^=GPxOo%X#IFNp(ysaaHfGp zs^@GsWj=Wf@+;c8|=mLB7yEE{DOll~c7f|l0B5q9vz=!>zEx!#x z%&?drh~jB~OfTaJ%rj=409nnAXWlLJA&kdZqd?RGiB1jtqgT_KB3Tn}u?)QYdfmhX zJh_YJqe3N9w{eBmrFt(}>q)cA5!Rb!y^h@w`T!cSH&2_wT&8Xk{~hCSH0{b_yF(>v zw?W{{l`Gqh9aJ-ipJMwLGIa3Fl|2CCCSdmu)V~B{JdYy(KNwMUwLt)V<_PZ}fB@Jw zvqxpQf~_T=8;{w7>ky1~r&R7?K+vRYxI#o303AsYw|hn)_Jh+!2XZOdVZ{SY0KzCo)HA5}RVWX1F80D2lfQY+VE7@SSiWi5opG-%lS5th zGn3&u`31QxMq<7>s?~oC&KNOlo$k<=4nu_YZUcgk(weXfYDZgRNXP+(I$%QuCaC!w zTjGFe!=9&kpt8(K@ByAW?(qRn>YrSA`F!NZm!fXvWM&12>ct=fSN8y~b)Zq;#GB=4 zL6jG5&4x0h>>!%gK?u=0@D|5i3%j7_%L3b0?h(u>n1Ze+oooNgn(IK_d-_PZWvF-- zTH1*KGFQAfjUh@g(*y5=+OrG%Z&G}~6I8X5WKNrT;5&U2NOwGtr!k3jc z@<;jM~M84O2nfm^J5pom|>$0Ds&a_$bK}oJ)eZm)#Lw#-(m9 zU4zj8vI&O6>RFQyI%$>HifqKIDwf50|CqpH?-!sSGqUW)bqQs-AUL%e%%@eqKA37k zi}G#G1^rWpF_rf~;=u1v0GaqJX%#QLP2R&3UNq|h+gS~=GG_oIheo$ldgPF7DWCPw zMw0q-h)6Z*7KhK~R4M>c|G}X*iYuK z*KEoQ`DD#Y0JuI^fDp(`2fOf@a;;khry|CG-d^gEZ`iIZKbt3d#|>7VE=;&6 a|GNNfP562#=s`gM0000 const char *Image::opencl_grayscale_kernel_ = nullptr; +const char *Image::opencl_color_kernel_ = nullptr; Image::Image() : data_(), width_(0), height_(0), is_grayscale_(true) {} @@ -750,7 +751,7 @@ std::unique_ptr Image::ToGrayscaleDitheredWithBlueNoise( auto blue_noise_buffer_id = opencl_handle->CreateKernelBuffer( kid, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, blue_noise->data_.size(), - (void *)blue_noise->data_.data()); + blue_noise->data_.data()); if (blue_noise_buffer_id == 0) { std::cout << "ERROR ToGrayscaleDitheredWithBlueNoise: Failed to set " "blue-noise buffer" @@ -874,6 +875,184 @@ std::unique_ptr Image::ToGrayscaleDitheredWithBlueNoise( return grayscale_image; } +std::unique_ptr Image::ToColorDitheredWithBlueNoise(Image *blue_noise) { + if (!blue_noise->IsGrayscale()) { + std::cout + << "ERROR ToColorDitheredWithBlueNoise: blue_noise is not grayscale" + << std::endl; + return {}; + } + + auto opencl_handle = GetOpenCLHandle(); + if (!opencl_handle) { + std::cout + << "ERROR ToColorDitheredWithBlueNoise: Failed to get OpenCLHandle" + << std::endl; + return {}; + } + + // set up kernel and buffers + auto kid = opencl_handle->CreateKernelFromSource(GetColorDitheringKernel(), + "ColorDither"); + if (kid == 0) { + std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to create " + "OpenCL Kernel" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + + auto input_buffer_id = opencl_handle->CreateKernelBuffer( + kid, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, this->data_.size(), + this->data_.data()); + if (input_buffer_id == 0) { + std::cout + << "ERROR ToColorDitheredWithBlueNoise: Failed to set input buffer" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + + auto output_buffer_id = opencl_handle->CreateKernelBuffer( + kid, CL_MEM_WRITE_ONLY, this->data_.size(), nullptr); + if (output_buffer_id == 0) { + std::cout + << "ERROR ToColorDitheredWithBlueNoise: Failed to set output buffer" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + + auto blue_noise_buffer_id = opencl_handle->CreateKernelBuffer( + kid, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, blue_noise->data_.size(), + blue_noise->data_.data()); + if (blue_noise_buffer_id == 0) { + std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to set " + "blue-noise buffer" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + + std::srand(std::time(nullptr)); + std::array blue_noise_offsets = { + std::rand() % blue_noise->GetSize(), std::rand() % blue_noise->GetSize(), + std::rand() % blue_noise->GetSize()}; + + while (blue_noise_offsets[0] == blue_noise_offsets[1] || + blue_noise_offsets[1] == blue_noise_offsets[2] || + blue_noise_offsets[0] == blue_noise_offsets[2]) { + blue_noise_offsets = {std::rand() % blue_noise->GetSize(), + std::rand() % blue_noise->GetSize(), + std::rand() % blue_noise->GetSize()}; + } + + auto blue_noise_offsets_buffer_id = opencl_handle->CreateKernelBuffer( + kid, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(unsigned int) * 3, + blue_noise_offsets.data()); + + // assign buffers/data to kernel parameters + if (!opencl_handle->AssignKernelBuffer(kid, 0, input_buffer_id)) { + std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to set parameter 0" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + if (!opencl_handle->AssignKernelBuffer(kid, 1, blue_noise_buffer_id)) { + std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to set parameter 1" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + if (!opencl_handle->AssignKernelBuffer(kid, 2, output_buffer_id)) { + std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to set parameter 2" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + unsigned int input_width = this->GetWidth(); + if (!opencl_handle->AssignKernelArgument(kid, 3, sizeof(unsigned int), + &input_width)) { + std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to set parameter 3" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + unsigned int input_height = this->GetHeight(); + if (!opencl_handle->AssignKernelArgument(kid, 4, sizeof(unsigned int), + &input_height)) { + std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to set parameter 4" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + unsigned int blue_noise_width = blue_noise->GetWidth(); + if (!opencl_handle->AssignKernelArgument(kid, 5, sizeof(unsigned int), + &blue_noise_width)) { + std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to set parameter 5" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + unsigned int blue_noise_height = blue_noise->GetHeight(); + if (!opencl_handle->AssignKernelArgument(kid, 6, sizeof(unsigned int), + &blue_noise_height)) { + std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to set parameter 6" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + if (!opencl_handle->AssignKernelBuffer(kid, 7, + blue_noise_offsets_buffer_id)) { + std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to set parameter 7" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + + auto work_group_size = opencl_handle->GetWorkGroupSize(kid); + std::cout << "Got work_group_size == " << work_group_size << std::endl; + + std::size_t work_group_size_0 = std::sqrt(work_group_size); + std::size_t work_group_size_1 = work_group_size_0; + + while (work_group_size_0 > 1 && input_width % work_group_size_0 != 0) { + --work_group_size_0; + } + while (work_group_size_1 > 1 && input_height % work_group_size_1 != 0) { + --work_group_size_1; + } + + std::cout << "Using WIDTHxHEIGHT: " << input_width << "x" << input_height + << " with work_group_sizes: " << work_group_size_0 << "x" + << work_group_size_1 << std::endl; + + if (!opencl_handle->ExecuteKernel2D(kid, input_width, input_height, + work_group_size_0, work_group_size_1, + true)) { + std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to execute Kernel" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + + std::unique_ptr result_image = + std::unique_ptr(new Image(*this)); + + if (!opencl_handle->GetBufferData(kid, output_buffer_id, + result_image->GetSize(), + result_image->data_.data())) { + std::cout << "ERROR ToColorDitheredWithBlueNoise: Failed to get output " + "buffer data" + << std::endl; + opencl_handle->CleanupAllKernels(); + return {}; + } + + opencl_handle->CleanupAllKernels(); + return result_image; +} + const char *Image::GetGrayscaleDitheringKernel() { if (opencl_grayscale_kernel_ == nullptr) { opencl_grayscale_kernel_ = @@ -910,6 +1089,55 @@ const char *Image::GetGrayscaleDitheringKernel() { return opencl_grayscale_kernel_; } +const char *Image::GetColorDitheringKernel() { + if (opencl_color_kernel_ == nullptr) { + opencl_color_kernel_ = + "unsigned int BN_INDEX(\n" + "unsigned int x,\n" + "unsigned int y,\n" + "unsigned int o,\n" + "unsigned int bn_width,\n" + "unsigned int bn_height) {\n" + "unsigned int offset_x = (o % bn_width + x) % bn_width;\n" + "unsigned int offset_y = (o / bn_width + y) % bn_height;\n" + "return offset_x + offset_y * bn_width;\n" + "}\n" + "\n" + "__kernel void ColorDither(\n" + "__global const unsigned char *input,\n" + "__global const unsigned char *blue_noise,\n" + "__global unsigned char *output,\n" + "const unsigned int input_width,\n" + "const unsigned int input_height,\n" + "const unsigned int blue_noise_width,\n" + "const unsigned int blue_noise_height,\n" + "__global const unsigned int *blue_noise_offsets) {\n" + "unsigned int idx = get_global_id(0);\n" + "unsigned int idy = get_global_id(1);\n" + " unsigned int b_i[3] = {\n" + " BN_INDEX(idx, idy, blue_noise_offsets[0], blue_noise_width,\n" + " blue_noise_height),\n" + " BN_INDEX(idx, idy, blue_noise_offsets[1], blue_noise_width,\n" + " blue_noise_height),\n" + " BN_INDEX(idx, idy, blue_noise_offsets[2], blue_noise_width,\n" + " blue_noise_height)\n" + " };\n" + // input is 4 bytes per pixel, alpha channel is merely copied + " for (unsigned int i = 0; i < 4; ++i) {\n" + " unsigned int input_index = idx * 4 + idy * input_width * 4 + i;\n" + " if (i < 3) {\n" + " output[input_index] = input[input_index] > blue_noise[b_i[i]] ? " + " 255 : 0;\n" + " } else {\n" + " output[input_index] = input[input_index];\n" + " }\n" + " }\n" + "}\n"; + } + + return opencl_color_kernel_; +} + OpenCLHandle::Ptr Image::GetOpenCLHandle() { if (!opencl_handle_) { opencl_handle_ = OpenCLContext::GetHandle(); diff --git a/src/image.h b/src/image.h index d7051be..3e3ef4b 100644 --- a/src/image.h +++ b/src/image.h @@ -105,14 +105,29 @@ class Image { */ std::unique_ptr ToGrayscaleDitheredWithBlueNoise(Image *blue_noise); - /// Returns the Dithering Kernel function as a C string + /*! + * \brief Returns a colored dithered version of the current Image. + * + * Unlike the grayscaled version, this dithers the red, green, and blue + * channels. There may be mixed pixels as a result, such as yellow, cyan, + * magenta, or white. + * + * \return A std::unique_ptr holding an Image on success, empty otherwise. + */ + std::unique_ptr ToColorDitheredWithBlueNoise(Image *blue_noise); + + /// Returns the grayscale Dithering Kernel function as a C string static const char *GetGrayscaleDitheringKernel(); + /// Returns the color Dithering Kernel function as a C string + static const char *GetColorDitheringKernel(); + /// Returns the OpenCLHandle::Ptr instance OpenCLHandle::Ptr GetOpenCLHandle(); private: static const char *opencl_grayscale_kernel_; + static const char *opencl_color_kernel_; OpenCLHandle::Ptr opencl_handle_; /// Internally holds rgba std::vector data_; diff --git a/src/main.cc b/src/main.cc index bbdddd0..3172e32 100644 --- a/src/main.cc +++ b/src/main.cc @@ -18,7 +18,8 @@ int main(int argc, char **argv) { return 1; } - auto output = input.ToGrayscaleDitheredWithBlueNoise(&bluenoise); + // auto output = input.ToGrayscaleDitheredWithBlueNoise(&bluenoise); + auto output = input.ToColorDitheredWithBlueNoise(&bluenoise); if (!output || !output->IsValid()) { std::cout << "ERROR: output Image is invalid" << std::endl; return 1;