From d852c88b9f5f21733f205438fcda798b8b68d17f Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Jos=C3=A9=20Carlos=20Cuevas?= Date: Sat, 31 Aug 2013 17:48:20 +0200 Subject: [PATCH] Made the OpenCL version work after finding out what was wrong --- .mandel_classic.c.swp | Bin 32768 -> 16384 bytes Makefile | 15 ++- main.c | 211 ++++++++++++++++++++++++------------------ mandelbrot_kernel.cl | 35 ++++--- mandelclassic | Bin 12095 -> 12111 bytes test | Bin 10180 -> 10172 bytes 6 files changed, 153 insertions(+), 108 deletions(-) diff --git a/.mandel_classic.c.swp b/.mandel_classic.c.swp index 59426f6f67eaa5f869f836d332a078a4f582f632..4e914e81974d6e773b0c3fb928c08ab6cc8dddc0 100644 GIT binary patch delta 321 zcmZo@U}|V!oFKtuq`6V*Jo7{YRz}{<@0o>p7+EGW3#fAbhk`ki4F$vv|3P_03=m++ z2%(`A12a@y8%nRrPf8qv*O$sam^)tB`7!tXGQUVMNFcn+)85kA;aSRYk0r5RP z28KC691q04Kx_fTYC!yemw{m!5a$4~5fFdnVPJR!#8-g$6cC>T;`Knh5{Op-aS;#~ z0&zSL`vb8b4}&v<9gv{|#A-my1;j6aj=2ZKyMTBK5O)D_3lKL0aWN1V0dWctCjl|Y a;R)QE8Cm9VZDzCi#XQ+3S!8lu(F6cYML?YZ literal 32768 zcmeI536NaHd4Pu-#aLiMFfNWXQY6jnIlDq2v|1TywL+_`q>WZf46+!FX5MO_G&{4J zV>OGw5DEi6FgU4-!H77NU4X$M6n1P5mn}?@lrbS(G3Kxd1_LoRhn+Y@3?|>--S65% z$AXhkys7$j-s}GRcz^fn?!S-O&FQmxR;$xm<|p_Zl}LQ}0k!UpnVSxHvo(<@I~AuY z?tA%EO10|H4!4-!i`*x4gspw$(pD!|t+Wnji+QK8xm>E+%S^72tyJ7xOKvYQ#on<7 zVh!w<25Qy8=7lp8^V{3f5^naalhq#{J+NO<#KN%#VhzL^h&2#vAl5*vfmj2v24W4o zw>3~5IW%!0k8q%Uj9vEofN|f4*!vCk`?_)d2iyCV_WO--{)gIoFP=p<0WY1ytiNl& zmyPp3$hs%(_l@KHt#kYpYarG@tbtequ?Aud#2Sb-5NjaTK&*jS1F;5T4NRheEQ|IZ z#QhqX0AT&!v;Tj2WFqkd+y}S8-@tY7S@;YL!4?>Vvtc=$0)GN0!U=E)91K61nMgbU z_ru-rHMkZ&4d+1)dZ7na!YnunX2RQK?lt%~_&(eY*TU8C1^8GNZA{-CD zI)d=;^s6!8&3A5qtKS(4Vg6rXL;X*hU zHbOHrLIb=tBa!$mJPu!nOJD%j!+bazJ_Iiuo=7|o&%rn08*mjApc7`o2jLe?lKcSf zhcCf(unT%&70iVL;ZY`Az5-u{iy;e}pb-v&r^fVHmc-N8u=VmdTt4;cmDA zu7Hcch0XA%unHEz>2M;v#01TY@G#s3yWnCt7dAp4oCb%%q3|m8?cd=Zkh*s@L_W%X z##E=;;#4!aY;MTOXU38vrAnsWq~vQX)sgVz>QTA+WO6&Tv^%JhZq+Ggt8S^NmZ)?p zVcdPUPGwW`#U_ibwN>?&O50RzM3rh~6C^R)9n3p}s%vT2%5EDi$!#Z@WV)q&fof1= zW2uxXJJnjbsL~xuC%v7_Z^sY8Mpe&{kzxC2;$s}}wik`WVIJH%AflR;#BZF*Sy#54 zcS~+wp|#w$AScOWog~^)l9;`c=u&Di;nSzKF9_vLXN|b}oU&WWH@B@wkXRuJpR#m;~439La zItA@LtMS;5qt=gkQsXGOAopP>sSe|ubdAAzI6Fc$a)+~9oXl7gML|VsQsZ&xTZ*sMA+2M(uf-}r_VrRnTF4**x87bkA z6h|AWcS&~S(qxguf*X8H?j(KVF`TVue3es`t-deN8u#9s4b^gm^vL=~|8i_2PiHHV zm8<)nA}CC{9a63*D<pw z=1{9smR5l(*p^OH;P$l%iA#Lr#Osrbhob8=?TfrB(R$^}NVKMkG^+L%Uc8-&QF-5V z{SFEwX>y_^y|R-VuZX=i$Md7KO^7jy(v+cXEB0KJ?=MXo6-By*sq{RB8mFk~?{hWQ zr>w41sBupX!TU>>(uCPp)NlUwpv_3X7k$x_NfRzLrdpM~9CF+(Lsi{tx3{EezI6YS zqSsH-K=Em_mwXv7X^_l2Vf#DsXB!<8nb**0acOI8$041tZF2ma)PtPtaP!q6oss!1 zGMou#(ObIOory}dT+3C}i0kB> zOt!qGlBkZ2IHK$cJ|Q7q^}T0!A9H30WX0}EG9#wwqmH$rTdC%wZa5U&%@f?*k~oR| zFShL=*uG-_C#@0o9qjgt;S4wkUdI-H8pMXb1Xh9A=s(29z6LIZIut?d@qSnhy>Kl2 z5gY`+CLOWKe+m!5{ctOM7A}T!VIj0Y66V2&;aCv+{7{g*9t`m()y&^^j5Z*s0PnhIn| zVI_x$SwwKex8kK`;xOS&j!YOocdR(*WW?e#j>fE%#B@SLXhoRTA1Na@n4GmbTUR~a zOsyKt7HZBSH7Eb^;;bg-^Ud`8L?ua_D+=yNvXU!1PSFz(bwu?Sz0>|m=Ta{B^V7ay zft>>1&wwJHZc0oe)bAyvw!t^Z3YE<5m$SufvRWZ3U7a0ScjmICR;$fR21{j(XY|O{!EnTAm0;}ANJkD^*hqEKiM!b-+M?h?A3P4!fA&Mr7rgW%d8pCn?bvQ{lC&S zAF}rUINSrf;96J*b72n1+Wy#$jRSb-5NjaTK&*lH zlm=wzB}Ri@dtxMdYT9l$@*HJ9$U@Ii#;N_a@?bXSs0K6qH+BX_MHj1iRYw_DRWVY= zD)HZexie+$1MD$zmvqd^EG@a`>=Y)i;ferm~m-T3EJuNo{0@7Z#i+EH@&YkR(j_KoDE6GMp_G zGNJWdn}!#e*n@$L-IHdLK_u{zLNrU*{suG2YD7jhff*SaPdK=@Wz25B2s9kplTjHd zxK&#}gRqi_*;gcBh*(=d|Oln>&`wU3umq13|rA7o0Y z-BiMfsa3X@t6mnvfo$1%k1?4D`*w*|XI2_|zsYDadSDx}-YCJVZ;Q@kSY+!>v<9Ji z5x&7k*IK<%MP8C7?``rfsJA**(JTr}n&|aiLrNN%n(QgHO={RvM5I~Nn-Lpq=@qiI zA{&9rN#p7VPwDl47yJKp_LqzOzrps(zrenK0d~Wka0mP)$o~I2^uath3SPyYe+8a` ze}daV_W0idm%(2^30#o5fOB9qw81=h89V=3coKHQU2q3{3qA&CK?A&ro&O-*315NF zz$ZZV@^1o}6IclyFbjUge8B7QBs>ORgG*opw!(7ggvD?y90hMO2k;WeT){PPF`NTu z!wP5vZ-4)xTqoca<^*1XXW$8V0B(aDU@L5bv*Dv~I-CZl!U7PzpTze8@G{r`2=~GH z&+ZHTVY9VLhbb2zZHo z|1ZLBxC^d?Pr?rPIs5ychmXSm$o~Is=z?W1AKKs;c!s_G_rg7}8@>eB!4KA|7zbTOSX&ypb7AW<20BT~|Ednq;`6NADAdIl5B~ zh6uLT4#rl}>=;KG$k|btQSkcd4V;vEd}mKseHXNRd0&QwIXl5ue8D3w-zWKVe@K1sc`*^N|hiP6u$ zS|aD9q6dp4#^5o0(Z)tI<`r*wXY|o6J9^b_a%e}~<%}8!#u_=7HGjc;zJndfA}5`^ z@oitYFo;r|M@l=)xwWxW6XS8>v5^I0lKW)6o+8?$9B1Ra-6@;~N~z6Ew`>dK!Y9@w zRrwH^oS@@;osrv_5Ldm599GHV2^tCQn2&0Y!*nTA$o8b0)pVe~$?_qO8adMvt`?zM zAlVJc{XaVD7&ei}Bu^by5x$`+JDt;irZRAP$`q8@*e>eOnhHDWR|${F5BqUu!s^pR zp@UoLiD&%+X&Mja(f0MIUwGQc8_H^n&rTzT5B6Di|5T#MmzhSAu;X7wQ{MolgDO7`WL^_^EU9oI!^PKg$S(+Gof|jdy&>sC?w>`4{pT>sY1Y-YV z?&;5uxxW!^fGmjr3VYqf_Y8Oh+g;A_pAVyO2DCsD#I8RMeumxt4D5mmxX=n8fg|Bx zu+y)E^{@(#haZDO zA!QcQnwMw~J0U89vS$prz`L3oYOS!IRvNTnLqjSPnQG9H1P`QpZQ=2|m3PXD4KHr5Z=zgM=hxpa{8>!b1b0r8sewkGL}8u`zAsAY0O0 zb0FUw71XhBY*K3mdV52O1PO?y!?H%{5tqf@dYg^A!#)K|d1T#R#E}d|mEdKJO^>W| zCe>UU!DN|h@*?#nsGF3YqGpPFu0Kf~4H$YJNqY*Z7W*t_&5*|nKU-AoQ^pW6f;-2n zBO$4KpQ+-AjM~UWAwxhyMr|cYcd9haYtw07pM6V1g2*UIe+NN!CcNi+ck7G1 zRnhYB6N*-)h`PvM+iBL0l6(cZ^9@{a-b)PL3Vk#N<}{L8(Ti0xKO%vNDSelAS)!U{ z{2%mIl6${J4EF!()m?8+5=91Tbji7IX@~3~lxox9$@rBCUB(!$&2r}1-Zc}nCC$W2 zuxLD0Fy!PIxD9A(3}e)E%jpXV9$DZ`$kPWFhHc%v-3ROV!4`h;rMRV~r6siWAhgI7 zR6V;UZ(niS6?iRxSIU(QlZyFDwA~77lI7o0PTD={BinxYmVU;z+GyV2ydZ^ORE$m5 zS9V6SWrs@TNSIQgrhnQ;CUYdHoz^K8O67{z;#7DzMwwa>(ft&m?G*g`D>droost?Z zsa-cB*qZi6SB;o{36#CIBLRw*!_OVoH9RP=>5Z8*B#CiVLk>*`HAB~B zuWm>tZ20*sO7Fr$NOLHP8XPMAOfXKRLK4-!(>Z_ zIg#g!KECymQDLY8NSiQ81yBuM6|ljiPn$+CQ_J@jRD3z^MX%?#dM?5F*pDMADX$hJH9g8G5F%cc5*^>O1-u2f!3|J>GvG*g z06YJ3SOGu6zJCz@87_iRI34~NJ_L_q*WU+sgRH6F1fPdIbig7I+y4mo8TS3-@I^Qm zj)R|I<9`9<_wy9Ygzr&Kw}F(GD`ob-^pR@(?tk97jT`hfUhZTx7|%$^04Sc3&~+kO zMd%0P83}8C#WNBz0yZ-mcDlr~4>IwLL}2BGjrEN!@r;D`+p+PCL_%8Fct!#x$T&Nm zk 0.00001; zoom = zoom * 0.98) { - // Set the arguments of the kernel - ret = clEnqueueWriteBuffer(command_queue, kernel_current_line, CL_TRUE, 0, - sizeof(cl_int), ¤t_line, 0, NULL, NULL); - - ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), &kernel_res_x); - ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), &kernel_res_y); - ret = clSetKernelArg(kernel, 2, sizeof(cl_mem), &kernel_current_line); - ret = clSetKernelArg(kernel, 3, sizeof(cl_mem), graph_mem_obj); - - // Execute the OpenCL kernel on the list - size_t global_item_size = res_x; // Process the entire screen - size_t local_item_size = 64; // Process in groups of 64 - ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, - &global_item_size, &local_item_size, 0, NULL, NULL); - - // Read the memory buffer graph_mem_obj on the device to the local variable graph_dots - ret = clEnqueueReadBuffer(command_queue, graph_mem_obj, CL_TRUE, 0, - res_x * sizeof(cl_int), graph_line, 0, NULL, NULL); - - for (i = 0; i < 800; i++) + for (current_line = 0; current_line < res_y; current_line++) { - graph_dots[(current_line * 800) + i] = graph_line[i]; - } - } + // Set the arguments of the kernel + ret = clEnqueueWriteBuffer(command_queue, kernel_current_line, CL_TRUE, 0, + sizeof(int), ¤t_line, 0, NULL, NULL); - // Display the result to the screen - /* for(i = 0; i < 3078; i++) - printf("Linear: %d -> %d\n", i, graph_dots[i]); */ + ret = clEnqueueWriteBuffer(command_queue, kernel_zoom_level, CL_TRUE, 0, + sizeof(float), &zoom, 0, NULL, NULL); - printf("Rendering...\n"); + clFinish(command_queue); - int iteration; - Uint32 *pixel; - // Lock surface - SDL_LockSurface(screen); - // rank = screen->pitch/sizeof(Uint32); - pixel = (Uint32*)screen->pixels; - /* Draw all dots */ - for(i = 0;i < total_res;i++) - { - // Get the iterations for the point - // printf("Point %d\n", i); - iteration = graph_dots[i]; - if ((iteration < 1000) && (iteration >= 0)) { - pixel[i] = SDL_MapRGBA(screen->format, - red_scale[iteration], - 0, - blue_scale[iteration], - 255); - } - else - { - pixel[i] = SDL_MapRGBA(screen->format, - 0, - 0, - 0, - 255); - } + ret = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &kernel_current_line); + ret = clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &kernel_zoom_level); - } - // Unlock surface - SDL_UnlockSurface(screen); + // Execute the OpenCL kernel on the list + size_t global_item_size = res_x; // Process the entire line + size_t local_item_size = 32; // Process in groups of 64 + ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, + &global_item_size, &local_item_size, 0, NULL, NULL); - // Draw to the scree - SDL_Flip(screen); + if (ret != CL_SUCCESS) + { + printf("Error while executing kernel\n"); + printf("Error code %d\n", ret); + } + + // Wait for the computation to finish + clFinish(command_queue); + + // Read the memory buffer graph_mem_obj on the device to the local variable graph_dots + ret = clEnqueueReadBuffer(command_queue, graph_mem_obj, CL_TRUE, 0, + res_x * sizeof(int), graph_line, 0, NULL, NULL); + + if (ret != CL_SUCCESS) + printf("Error while reading results buffer\n"); + + clFinish(command_queue); + + int line_count; + for (line_count = 0; line_count < res_x; line_count++) + { + int temp_val = graph_line[line_count]; + graph_dots[(current_line * res_x) + line_count] = temp_val; + } + } + + int iteration; + Uint32 *pixel; + // Lock surface + // SDL_LockSurface(screen); + // rank = screen->pitch/sizeof(Uint32); + pixel = (Uint32*)screen->pixels; + /* Draw all dots */ + for(i = 0;i < total_res;i++) + { + // Get the iterations for the point + // printf("Point %d\n", i); + iteration = graph_dots[i]; + if ((iteration < 128) && (iteration > 0)) { + pixel[i] = SDL_MapRGBA(screen->format, + 0, + 20 + iteration, + 0, + 255); + } + else if ((iteration >= 128) && (iteration < ITERATIONS)) + { + pixel[i] = SDL_MapRGBA(screen->format, + iteration, + 148, + iteration, + 255); + } + else + { + pixel[i] = SDL_MapRGBA(screen->format, + 0, + 0, + 0, + 255); + } + } + // Unlock surface + // SDL_UnlockSurface(screen); + + // Draw to the scree + SDL_Flip(screen); + } // Clean up ret = clFlush(command_queue); ret = clFinish(command_queue); ret = clReleaseKernel(kernel); ret = clReleaseProgram(program); - // ret = clReleaseMemObject(a_mem_obj); - // ret = clReleaseMemObject(b_mem_obj); + ret = clReleaseMemObject(kernel_res_x); + ret = clReleaseMemObject(kernel_res_y); + ret = clReleaseMemObject(kernel_current_line); ret = clReleaseMemObject(graph_mem_obj); ret = clReleaseCommandQueue(command_queue); ret = clReleaseContext(context); // free(A); // free(B); - free(graph_dots); + // free(graph_dots); + // free(graph_line); SDL_Event ev; int active; diff --git a/mandelbrot_kernel.cl b/mandelbrot_kernel.cl index 029599b..d42f85d 100644 --- a/mandelbrot_kernel.cl +++ b/mandelbrot_kernel.cl @@ -1,39 +1,45 @@ -float map_x(int x, int width) +float map_x(int x, int width, float zoom) { - return (((float)x / (float)width) * 3.5) - 2.5; + return (((float)x / (float)width) * (3.5 * zoom)) - (2.5 - (1.0 - zoom)); } -float map_y(int y, int height) +float map_y(int y, int height, float zoom) { - return (((float)y / (float)height) * 2.0) - 1.0; + return (((float)y / (float)height) * (2.0 * zoom)) - (1.00001 - (1.0 - zoom)); } -__kernel void mandelbrot_point(__global const int res_x, __global const int res_y, __global const int line, __global int *graph_line) +__kernel void mandelbrot_point(__global const int *res_x, + __global const int *res_y, + __global const int *line, + __global const float *zoom, + __global int *graph_line) { // Get the index of the current element int image_x = get_global_id(0); - int image_y = line; - float pos_x = map_x(image_x, res_x); - float pos_y = map_y(image_y, res_y); + int image_y = *line; + float pos_x = map_x(image_x, *res_x, *zoom); + float pos_y = map_y(image_y, *res_y, *zoom); float x = 0.0; float y = 0.0; int iteration = 0; - int max_iteration = 255; - float xtemp; + int max_iteration = 256; + float xtemp, xx, yy; while (iteration < max_iteration) { - xtemp = x * x - y * y + pos_x; + xx = x * x; + yy = y * y; + if ((xx) + (yy) > (4.0)) break; + + xtemp = xx - yy + pos_x; y = 2.0 * x * y + pos_y; x = xtemp; iteration++; - - if ((x * x) + (y * y) >= (4.0)) break; } - if (iteration >= max_iteration) + if (iteration > max_iteration) { graph_line[image_x] = 0; } @@ -41,5 +47,4 @@ __kernel void mandelbrot_point(__global const int res_x, __global const int res_ { graph_line[image_x] = iteration; } - } diff --git a/mandelclassic b/mandelclassic index 3cebd922ee21317274cbe794bb8927361dc4d45b..4c61ca940241435362b07d52600dfa4257e3ab69 100755 GIT binary patch delta 2432 zcmZvedu&rx9LMjuTV3zg_V#Y=ZaabAF6|x#);6L*h*?WExRpT$f({*Q;5J7L2+P#{ zp;>i7c2(k)p9xcZOidKfgzElqKH@5<8~h_v6H$B&qNYWNb3+|Ixc<)BZCuoo+;i^d z`}>{O?|1HFWb4OUb&oV@k8!lj)*P37qWYC+x6Cvy*K(Y~1cLQ7c=ed}9=Vn2JwQmM zq?@X4syZwuEt2A9E}Y_yu^R=Bke6Kh?#h4u(wg6Dk3Lg>Z|2IYou@q;<)K!7H&1rM zc)sjyDr=qXX4bdj)K1RJH36~6qzSy1FbNDLEI!_=duLtg!t&;HYJvv~WEd;3+3B)# zWf)#Zbq3YBzGH9>)hesaPDja=Gay-IXhPMRld#P|9MyhQ3+Jdaum{y)tIg|}qgu^v z6rd);*H+o?oMZLCIQH=fHhb|LwF-zJgNmwaj=Bt@f=ptN5Ek?XoeHP!O?O)|E6_vj z`@ni)*3-DY3{6IG9@e5ShIf=ZsWCfYU_r<`WYQ>Ky0rWpS&BIi>c zWITo3+i9VK&V&opfd(y`WPF}7-EMX_<1Z=G@MT*We?nQH+{pM8Wg5tAjPY^GMU*3q z-==J%tT28?JJ&Fui@$di zN6vI$Csnij=DAwsEvz*bsiKa)AbTNL=YzbcU!RNLGuF*dYBd!%B`bP5`mjBTmQy;b5 z4}yIW*#|e+{hNc>8)vh8h{@>N`52iJH`nVOqAT4Y-j>Ejylsrl+P$8x=^%Q|F1|yo zS*wc=ufxTPHI_rD54Yi->crKP+2ttekg}Kd{`v=-JR8ED*GIp5haM?){{yMaVqDeU z^gX`4XYaWn`sy$khfq=Xx1WioZ;v;}x5T%`Teh4?K9I7-#k#M4JQcl^(8jl_+Ih9_ zyUE6liP5xzTUAHj*Gp(Knl8kRkK%NfJJ@)`E%;0=*`zL}uL!ko#-VE8-^Xt`F^nr< zPsQaeBh+L}h`uD%hAB~hNH{|H>gaZ?B>nKFV+DC+Xxi~Cft2LuZ_bBLrE0$M3S5=S z_$61MtT^PiUNH>J(8hpsT3nRI>wF|XFw{E zBY_#i7_AFb(V1zl>^zJR4r2cnpRyr7`y_@ib2mD5h*e`h9_WC6=XIs)(KaPEot+Y! z-J95b^ZZRbj5{Oq*C2}fp2E>EXVEm~>P8Rqa9;1>VS!Z6+jv+h?J83HE?P%6qF?4A zxT~azXz)SFYf61r*WF#pj$KNyP3i9JYTwpY$Zcrf)z;p%W7nNEH8r5SR`K~UIO9?Z znjb&@OI7nvV7TT>t(OfWje`yTVi*&c+|#hwz0?xHGnIlJ?pm@8j<~~Q0KRY^ENno( zck%N44hKqaa@x?n2lM5jX95d+RvNGjPa8%A{wnnphP`DI}E$L=>lxG(T5KNGMB%ev|q^L;B}ac9N3$rO`p~<7g%gKmtj{ zLq>}V8WKHnv^2ph8V=c>X)1*tmB4?1CG^(rz6T6Oq8e;60>sgigPMrgbm zSBEbjS~OO3oPgg%I&%t0=)Zp_6{fWnq9GY4PJ~v2b)gdmL(4GvUxdP>6NH6f;)LZ3 NeIzopdEpOve*^heu_OQh delta 2421 zcmZvee@qld6vub=Snm$**j>0|Zwbgj_<1PDrCK${a;Kn6gOyTj!O9O?z%#MtZ2Unr zgtKVMX-(8;Q?WmSHt|=CV)>)AHF9cOXw#%%tkr4^SlitBqd~Nhigx>Exr?NBlAYPl z_q{hW`(}1$f2!@R-DMN6ncWO66E(y5cIO>-HAtGqL|cZjYXZS~3SQob-7D2=x)%t^ z5&aR>A5r}P zv>bs`dI|RFP3FllJur-FKdQ91Xz2&hAi)@_hNy(?W$>VyYcQE(C#fM=i)xj@B*#xu z$0p5!hZcim7A9E@Fo-cd*i*tJ)ds{UK?qgLBsCXYsM?Gss5BvsaUdD$}}V40gWH=$Smcc#-o&JZo*v}-=Z8t zxl`jC$gN(E3cEDnSL#617WQj=fig{d*r)LilxYIP^%|d}Y@}SS@mG{-D#LD#Kc^f= zxlrR{lueZF8Xx5CFEL&-6$DM#PaWbZGaB!q%u^ou2ZgNdlxg0>LmIzJS)@Fm@h0ew zo#CiKRSBFslWB1?u3+d8M)uCf`_O50LSZRBLf~BByWV{u#y#(N5PxnIHLh5H032;dZMbh1ww)Au3&EqooNiL zHkEQsd|ELBUzny^4nI~^Pe5PV))!MYXMCMo*};`T*ffo-gd}r%HHRJGm?iCPZ|xx9K52%&__<_%w~qgWKzTxb62Uz^%2!>*2e)$C{qj^U z26iVDrjFzG=ilqV?lBpzAm(UZKT6on?AO8dgbeny4%oys&?%&|&*|X2kjJ{m;7=i$ z%^ZVR(UCs(P*riAD>H@r{InomxA>DG5-Qd;WagIbn0`+WOT`NI=O=Jfe3hO45J-Zw zz=AP-#hJPEv1n8{1v7`8n0s*ry4mnvT#ry4IyLdR{kQkvGEbyOX}e?St;C?v9@v~% zMLOVO;&D?&o;3YEE?Xl6?ki^=zE~5K@9Ht zOI3$(pdVqLbr#o&tJ@6kSc}LE_|obmM<8tdFoS-NPUB^Jq^dT$Eh421+lzSF$6HMDa@s7TJ<^38s*#=^H62Lpk%Ollt{^=9oN@RT4O1N7o zSs6c&EQ$ERm1WgQ=;=b54J(l7k+>qO4n5CiNwN@0tTrjEitu6?GvkJPX!ZZc>P0KT zw3kDAwj@V$5sm#c659zcXVZDD%67{2k~?bF+F!6HjD=~2Q`t^jtl|IZCsx}b3An6c z4$dKx@n|>MXtgso%%CU7iMd#tBPB=wNuoZ>(I)|)VwV61b2iA)rH`6VM$7}?%B{oy zhuv7);a;v2Q th>mA}O0px?a_~4L9Jt!C2v_}pR)-xK^(A$A=?^*ei{Kyrp^!qw z7Su+|pc92x4r2bA3H!tOhnV^PX+UXYV(||l3POYlRS+6b=%%pNsn2un5*xg5@ArPr zIrrXk?|tunIrrI|b%f8aKk(s+MTLpdQ^s_`99tjOrLKQTDGK`%a6YYAU8~i#T5yVA zR7znH!b)LNbV~3MzdC&SQhszHx!Q<)dhJl)%9W7)*2%B!qHB+9=5X=efNq<_MP)Y1 z3|R1}n(;p`L$AU#X4o~)9j;l?dxI%SSKl-ST=%J=de=LE9uA-TRQS~w3^n_#9{3Hd zV*|lZtIz6#`vz2)MeBB-)en!O^)Y9a*PLyfQUT2FkZn5&9}i)!Io5TRa|g?ARqQ# z!>z2oG#TW#!}Fd=cssC_N#kkebI0lH$gS=T_6Kf7`$#Z-MvdXlV@^Qhp zJPKu{%<&C~BzPH;kriHnTEyX&U5kF;<#otvDe2(%U)Iu6yZaclp@DSpjO=S^*wqXu z#$+&J22^9Rbnu)sVri-Ohg+~8!&r|;9%r^e3U zSGA07)ozS^!g&eaJUrr7SQ}Nv{I4TLSs7J@`r$+F5oWz+jSS(`byJgV6~YI^iot(Vt{ugc1 Bhywrs delta 1511 zcmZ9Me`s4(6vyw)kLD#!oA=tpmzUP&^_NX2qpf51RUJu&(C65!Zi|~@X)06thr$XS zMxFLa721NBa*$=9kTJoCY&HZWe+Z3CT>la5e?y!^L0Lf2s?gM`&vV{hT=2p>=YGED z+SFBTiWCBW@)gwb()I zCsr-~(}zrFRTGhur#4jVC!a{Bk|$FqbVzuEyb3ArVg4TEy!#z99Ikt(9bqTj<4Xlx z153{j`1x;vr`fYBF}*<{vl@sR zq?L&%A68oTJ#2-(w?oFEmjsoub5aT?x%C>BH~oeCIh6SC*XTC%r9%ZbCMEDrH9JTA)0I?fx#d6?cV9&Z;|5W5Tbl;L%ZRxx^zqIT;>|2{qf6vZf?gDp|! z38;xvcqJ3k5Z7T=y1)(Cm7;tZl#WH7hl?F1&%s(xCsetswBEVw#-%5e5$~Oke_k2) zR_$Lw6^hC@#3G{*j)*Xa{xtdmKAFg<_s(ym7{RPA5$dSs?Qbm)V_oCk7~UVX7|bKl z1yaMWy(`W&*hYI2#=7HJ@;ur?so4E7=NXvokx`oM2{}8Mf%VF`gEgce%fe?!w1<%M zE|!CBq&(V?GmMM$MrE9~HyVORqeplaE=FgRguN&pywzSn%ULH|21y;k^GTx3z`QDB zKAp_FYPZtZggu1)Vs$!j%$6!EjlXIm+KE$h7Q$+Hugqg5rS}5oF(~#;N{t;iQcu8k bpA2fhhGjn4KZ2LE(m#p)EwUwL;Ii{y>bs6}