[Beignet] [PATCH] utests: Add test case for image box blur.

Zhigang Gong zhigang.gong at linux.intel.com
Mon May 20 22:59:22 PDT 2013


On Mon, May 20, 2013 at 05:21:33PM +0200, Dag Lem wrote:
> This test case can be used to demonstrate two bugs:
> 
> * clCreateImage(ctx, CL_MEM_COPY_HOST_PTR, ...) makes a garbled copy -
>   replace "#if 0" with "#if 1" in compiler_box_blur_image.cpp to demonstrate.
I can't reproduce this bug. I get the same output with the two different code path
one is using CL_MEM_COPY_HOST_PTR and the other is using write image.
> * Kernel float4 arithmetics don't work -
>   replace "#if 0" with "#if 1" in compiler_box_blur_image.cl to demonstrate.
>   For now, int4 arithmetics are used in the kernel.
This bug is a limitation of Gen platform. If you create an image with int type,
you will have to use read/write_imagei to access it. If you create an image with
unorm type, you will have to use read/write_imagef to access it. So if you want
to use float4 in the kernel, you need to modify the cpu side code to set the
channel data type to CL_UNORM_INT8. Then you need to enable the float4 code path
on OCL kernel side.

diff --git a/utests/compiler_box_blur_image.cpp b/utests/compiler_box_blur_image.cpp
index 812a906..29bd466 100644
--- a/utests/compiler_box_blur_image.cpp
+++ b/utests/compiler_box_blur_image.cpp
@@ -16,7 +16,7 @@ static void compiler_box_blur_image()
   src = cl_read_bmp("lenna128x128.bmp", &w, &h);

   format.image_channel_order = CL_RGBA;
-  format.image_channel_data_type = CL_UNSIGNED_INT8;
+  format.image_channel_data_type = CL_UNORM_INT8;
   desc.image_type = CL_MEM_OBJECT_IMAGE2D;
   desc.image_width = w;
   desc.image_height = h;

 
> 
> Currently, the "golden image" is a copy of compiler_box_blur_float_ref.bmp -
> this must possibly be changed when the second bug above is fixed.
> 
> Signed-off-by: Dag Lem <dag at nimrod.no>
> ---
>  kernels/compiler_box_blur_image.cl      |  31 +++++++++++++++
>  kernels/compiler_box_blur_image_ref.bmp | Bin 0 -> 49206 bytes
>  utests/CMakeLists.txt                   |   1 +
>  utests/compiler_box_blur_image.cpp      |  65 ++++++++++++++++++++++++++++++++
>  utests/utest_helper.hpp                 |   6 +++
>  5 files changed, 103 insertions(+)
>  create mode 100644 kernels/compiler_box_blur_image.cl
>  create mode 100644 kernels/compiler_box_blur_image_ref.bmp
>  create mode 100644 utests/compiler_box_blur_image.cpp
> 
> diff --git a/kernels/compiler_box_blur_image.cl b/kernels/compiler_box_blur_image.cl
> new file mode 100644
> index 0000000..c8fd6cc
> --- /dev/null
> +++ b/kernels/compiler_box_blur_image.cl
> @@ -0,0 +1,31 @@
> +__kernel void compiler_box_blur_image(__read_only image2d_t src,
> +                                      __write_only image2d_t dst)
> +{
> +  const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
> +                            CLK_ADDRESS_CLAMP_TO_EDGE |
> +                            CLK_FILTER_NEAREST;
> +  const int2 coord = (int2)(get_global_id(0), get_global_id(1));
> +  int x, y;
> +#if 0
> +  // FIXME: float4 arithmetics don't work (yield zero results).
> +  float4 sum = 0;
> +
> +  for (y = 0; y < 3; y++) {
> +    for (x = 0; x < 3; x++) {
> +      sum += read_imagef(src, sampler, (int2)(coord.x + x, coord.y + y));
> +    }
> +  }
> +
> +  write_imagef(dst, coord, (1.0f/9.0f)*sum);
> +#else
> +  int4 sum = 0;
> +
> +  for (y = 0; y < 3; y++) {
> +    for (x = 0; x < 3; x++) {
> +      sum += read_imagei(src, sampler, (int2)(coord.x + x, coord.y + y));
> +    }
> +  }
> +
> +  write_imagei(dst, coord, sum/9);
> +#endif
> +}
> diff --git a/kernels/compiler_box_blur_image_ref.bmp b/kernels/compiler_box_blur_image_ref.bmp
> new file mode 100644
> index 0000000000000000000000000000000000000000..149cbba1d8bae1fdd4e4803970e06ec1682a40fc
> GIT binary patch
> literal 49206
> zcmZwQXWJ%KmA7%{Tg+Pn(kLhjf(Rlu2xc*%h!M;=D`LPL#w;qJB7%wt=GGQ4Ac_H6
> zZAlWo%>1fO?ltAzvybDc*;ie=_Fn5;|F!n6>u!AZKI6C>AF#ILT=)2Y7vBH>{_6q%
> zt*70%UjO6r%x6CHS<iabi(mZWcfIRfU;EnEzW at F2f9XqKdfVIH_WIYq{v#jx$alW;
> zoolYS=ErNh`s%B{@P#kD^{sDx%2S^5xW_&2=%bH*)T18dTIbkfkA3u`AN`oeJm#^F
> zeeB~O|M(M5IN_LMj`7ta9`T6o`^r<#d*1Ut^q~)Z{No at 0;0HgbfBNaCA9d7G#ytGt
> z4?p6FBMv|Oa2?M+{pnAC=}TYwhBv&yayrjg+pAyw>JNP21FZ7+CqMbg&#di#4r!nM
> z^rt`dsZSZ8W60+}|M|~-?sH%M@|VBzm9JcU at x|Z#<~Ogr@=Ct^<R?G*#V>ww{q@)X
> z{`bHC%U}L-;J|^~ZoBRFwdojOm2=NM*P^d{<ty26)>&t<;?z at 5<p`Tu$dVuZ=tsZ&
> z<u6Tg^UXK!- at o6wFMQz(pY)_B8Ft1QX8`-67rp4^FMqk4XP$ZHv!DHJ{nx(swXb^B
> ztDgJZ=NiM)lTJDbE>U>PTi){ZuYWzrZ+`QetsDPque~<@Km6ejIU+O$Kk<oAL=`t}
> zjlj|IKmPdR&3W?4CqvCnUp?}Xk2J<rIIPGc{9$~@JKn*h(@s0B8C)Osu!rs0vquz6
> zB`#;5eKu3^cYXfzpU+=+-v0KtfB3^6jz7Fw_=Eh}&wdtg<ZJrw_=;@6r>;vcz4Wrn
> zF2kRwe)hAU{rcCx{@w3>_s2i}@o#_o+u#5G_kaECU&!Bi=bd1F;~U=~NVpj63gdI0
> z^PDF?`N>@47YiAG;e{7ob=6hP^QS-k>86`*`qsC;^`<wy33efP>QkSp^MV(=K<v&r
> z=N#d~{3S1WiD;g8-g!JS1|?pT+!LPg1a5xyt6#;R*@S1z$3FHk`WIqS(b)XOFMjbo
> z?|BanETMDIrZEX&@Xg6c-}(r<5LiO6++>ECktYC7J?&{v!{0E^8p(`8U{0zKWn(A-
> zbK{9eWBmKy|9*yj;uD_$IsR0Eu&a-5`ic(w3V)5+7hQDGC6`=6{5<;J_r6C_`}Xbo
> z&2N75hd=y*_;0=S)_?rtANZ61zyJO3U;XM=o_fVAUO~6W^A|2KB|d0-_wL13T*#nk
> zLL`ndND6q6kp{D0jN~}4g#h_xGp!~Z at qh4xAN-JqJOsD+fA_oJB`}j4!%b3QYWyVv
> zFD|&?f|tGQWehm+#1q2{7s!-<+;PX52fFkRtou?GSelW%V*0w*y-uxTp7aEt4GlJx
> zkw+eRB>fw2Dk;En<LOY+d5Qv|c;`FcDKUdQ3P+gE<1Ym9heBL at Mf_45e7LT-;tDmE
> zQ|bo(M4R~k`OkmifAHWz`j_d-BY>Jk!6)PBF at yVx^ZVZSzRH@)M)pD`h!V&AzA_B?
> znrC}T2t?E4nvQ7l{Xq|UkPv{zQ+Q3~aRY>aBlxF(2D6HU;!gum$G-xOHfW+?DxuM6
> z-hlQ9`Naz|W&`Hfszs7WAQTpN3ZlUt5e0#hN8D$L5+v!k%x06&iSjk#AIRxH^5UW{
> zxg#ivKmOnT_O}U4j{p4UKi?J~;>W+@?<&(3N%WX!Me1d_%$14neeZh_AW=&QtH}l3
> z)gyoY^Pi=FJ;kEZ7IPs3Ut>APA?1#wt|G|}9Ee|2&lN7)lL|?1gh&5;VQuwa6x2Er
> z5*^=ygB577Yfx+P_m$9S!K8eB{8bw(GD+ at mO7uy~RI=5_;#Fw{UsIGM2;({a?svbN
> z>{(9#3_Jh)^W{FrgUo$r23}+0CbM-k{DZIZPw^pstfu&YMz at r&-Eqeq;M1@%^#7XI
> zyvB-TNdU0B^U;rfbdt;D_*<6^mIlW at B8b2J?QgAx62F9i!Cc^VmRDHBLUSL1!`4QD
> za@?jP1fncn at u&aUe=TwO<(Ko_gJ_ezrUTj<gkyB*U#sIQA<&l*<_BB}5(3*08EA|l
> zW(|L50F(4=krmuhz)654tP6;EX at F3m5`;F++*fBLLhxPnHTUCB{B)K0!Qbw at fBoxU
> zXXm|ZZR!%45xEqA-6jirX$7_A%wP4hgmp#1<m!VhSosoA`wMx{GlogMD6m{WIqs$=
> z1n$G{)TbKrk%gAVU%U*n5dE9^r$7CvE#L<~_<`bvgS~@3!B78W53**@U4?%F=^}~0
> zW(@xF%MuEXXXP?csV5Xd+l*9ajDa0+9aYg(KDv6u0NMb=2~Eg|%CqglRUsUx5ge4d
> zQ{^E}a7t=*mGLAbk(G$+uDj08i`Xr<+!BAvSNWZPNY3L938qCr5OuUAqdxN)OK9+F
> z)-sZ4#ne_}ACW{#kdMa0NLq!UL_iyMU?&P9qWIhL6k$b`h4`yoC>drsnOOagzx_*?
> zoI~0!u at wch!3P|@weX)v$%3-tU(&?QQ%*TWF%bpg2cK<9I{tQKOUsM0qwPSX3Wu3J
> zLJYzr?v{?ftB>u%q{z;c2}nr+3+3Axo#^m~h~>0o3*(5r(uw`X8*fx<@UQqIF9g!T
> zQ(_GQZ6|NnqF`DJ35t`61e6+<2JMC$ZUEmrCefUNR6G_+0X_;rnqN})oOATqF9c$t
> z&OGpe547JYL?$$qRHZH-4YQoBiox_R3y4Q$BSn<3Ed=oDgHQP&D<ew1kEAOo#)yTu
> zXa?A|8Kao1rNOME5Qg1BAgrtDn56Itgp#EMc}mDUE*vsfpTR`9<xiPHfC*eMI!<s0
> zeBzhdgq=J1Bj)5YC(jN)&`AxY+bSRg-#kJIg|_Cwk_gsSloY5+gjAA5ijNu8FdKoj
> z$U~iExe!QsSIi{>u=+wU{88c*%kj6IZ at FwR1Edo8P7CQ@SQP<srA#GA;G7k8sGU$#
> zuMXx00BC?t{AGa at u)5S{Jn`64Sy(RcmpwAeBsL3!)gLogf;`nGsKZq85)Lzf&ukDd
> z(R74>MDbUAI`ap=W4~PP<fB~7YZv$$boO7|5aXcIQ4r&=Un|k at hgS{ch{!tpgD-u|
> zaL+yWh=o<ei&m4mG%p0zA0G<`A<;iiB?8MC=_x)^APwLzVeMA7Bope(Wpk2f<sDvc
> zOsWvnv~6E at W<w8uE&e7K9eF|FzGc80-}pwgP6~5Dl5)XEi6H1IzjG<gSH5~yF4&*}
> zkuo0F*8~xN at R=R|i2_hne#J+^;=?b+NB+uUCm$zIzd0yAPCkl8x>fuEBrqj{4eV(i
> zQ{kY+zd(z7`j-OoOE#(e9 at p-^`)*hBTbf9~WHZ`ID2<_tIc>Y6353HZi5g~p)!x;3
> z0uVTPCt>lIu*{YVHhQdBhzU5=e^(jgj+y0`&9t&(fR1k!BGI8+t(}Eoa+&d{$^!Kk
> zPhmw~`c%7EC4%B87Ip`dOFc`QM4&z8aUAfkl?e2SRpJM;V~@NBpZ)ow_*3lg at 8pS_
> zibU9UAIEZrnWs&}RYQh2MOjP#P$*W`5?1wp>W@(R2quL<-C+{^wgZk75lh$#oqyB^
> z*&$m!VhP`37=LFT^#^=~k^Rc4(rrR|23CClInE%f9vY19_z0aefzh>-JpB{77 at -H3
> zR06ZyWTzOh#LS*GgU8KZ6BN at D4CU+C#KofYmq(_m2?YLFIrf0h+AaTb+2+{K4i1VH
> zZMeX{_uhNeSSGR7p$nM=53dQu0+~4lSxyLSKxkOh=5E?1Bz_Z$1<{&;?D6Dd{1rFu
> zQ;|f_mt_iwbQ2o2#lh7hasegIDt+l6fAvRSuE{$QRG&Ra-atX5qj|AnxdQHH8g55+
> zrKm#VYH9qL%}0afnyY829r+%tZyp`<TTc1n2>xutpTYR6U^onX%7ke6s{o3>;-jjk
> ze_jCP6^Yaxbvm(tf8TxgQ7qq?=QPM#f;27WX48DN<<EGjXLbBFMF|<nWh1tg;4cdD
> zOX+h?09+FSz%@!M1-ipQYB!8jyk42vxA7PH_!Ga$JO5kMY at YI4rNdA4A;l6fp{f<Q
> zk+B00XwunkA}$Eytp32R?BhlfVolGyR{3!*Sri=8mO)h^K{1teJuY6x)0Om(e^qEL
> z3=4l+rT at fFnH7J=dy$O4;^X8$Cq_TG*);8Oiis*JuL%;Rz?!XQ=0f0VEz#)|N(K_a
> z9Y=SdgcW}qu|P;DWB90Gx&lt{B1?fjYKjaj)-pp!Ws5pph=@qhe>ux+B~EtuXmy1P
> z7W}0EP~fN=?qJA%Q4k%K+<oN3E-3oykF^}Fd!46YA#ev(&lXKjsQ?q!V&RT?hCg`-
> zqFqCQ6-SNxDvuT={H1f^xATIpc%s1(W#Q1`?<~YG$(05qthH#;O(MiWlLZ3djwLK-
> zQPAu`%%B8KL<m at NfiDCORQ6?-AZccmmZz!;dZHi}_;<(NCxa|ua&?uuI)HhW{+VqH
> zaB|~_t6`3UfYS+Lu+w?8U13z}MZuL9WM7j7<|G5FJI1iwRk1Y{{g8oSe6Rk*-^|kl
> zMFe?47)Ht?E=ZzEkdg?9#tSWI!{4SB{5D at qI{AB1Rex;09oMPC5^X at TScz67F$bY0
> z5b?K!DuF-v<jM^2MS(HkU`6>tpaglpAX*1(mZPe$1Y8OTgmNm%o{}SCDtr>qO=d`T
> z$}}6}r7i4SkjSbi8yE&JIQT33Y<A}Kh}N^d1)LY|z(u;QL~e4x8O&M*q-5kcZ5VGt
> zYKPjjzK)3K3vqSKLh_6LJ?ju<Dg;y^#^YZKxH8iIOaJrEOZj#&KRa~t>}LnNq*yo|
> z at kkjY8=5h{Rj at KY0<de#9<c5LCI?>w4R&LA<Z&IGDZy+i8~B?2-C37yW<$od#Z!^>
> zk@$58(mZBULe&IuAuu)|=9F$lJD22DcJSA?p>WM3NHWs--}@J0m3_}@=CqcM0i3d;
> zrW1dfFJbW~du2p=dPL%oefkf(t5zWJ7gJ;mV;)D!BQ}E|o2ao65GO>oOUF5TBAbav
> zSypzSP#jc#CHmG|Zv}_);R28Wt9+K-%0l*E4yr%+^Ve+jsRiFc6(7;}EF~zU;JdO%
> z8TZSE6p&f2ma_+`DxEQ4kqGKA`0Bj85HeL%1f_;CG at y(Ch`*(2UJ3{SFLZ2b2Fp9p
> z4DS7l{*hsMGqc=BeL?JXL%1!3za`W<{E1daxP?GySmlIc-Dr>a at oyu>A9nB!6Khim
> zp9%(%oNX6aO3w;E{?Iv5;IDpy at 5~84;A#SW3V~dZAkv(jq{0Rs at j~B~ZH$yeVV7L6
> zV^z&pG^{RjaiC<yq;wVq$zm<!C6}hkdEHRKgh_}6S1GU+2E<j=Od=MT^ASuTK(DFD
> zL(Q=Z*ndUAiH=pgW;RtDY^wza7Sd(8CBSihP$GzmVJ3HFq*~|7X5=}-1wnU**PgPd
> zhP-StRpd?HhUlZ{SjDl53zlX80^rlX*%CkUaH;<kt4RUcqnuuNl^REM4S8i5Qc1*S
> zUSu{|>foUHN?Egsi0KHP%@nTG$YR_iuVKK_szeY1UqNk)fn9!yf+p7_ATb+>JUE)&
> zuyRCx79}d_&rL9yBpjtuXCB3e7y9&2vRI3?Xn2JE0_Sm8^AHb2G;cY(yj4WGqZhM5
> z!@pyLC71y$VT14Sjy-7}{F)_9ZWW;s0+Wj@{$v2YsnkRDv-%G^1rooe6i~S0>CCSt
> zIGo!8WHg at C#VU)6ugRIgWz*__FC&08tq3YLuxpkh7%nI#8sPeT!Q5AxIf>IL)D~lB
> zM_viC3pfZ#H at iREq?P&k0-ZzvA8>keV(<6|1?d`wGqn at H{a51!%k2onrtVpb2-y!v
> zOu#Xx*|;MFN)3gOH>W5+YvE$B`*|TKW<XmAJgZfC*yAQLtkN-IDP{@g;V<+GmpLmw
> zOu}EIRX$RnI!U at vbY@V3q_A87)6!*VURy$Z<1Yk~5qynpB5O8TKr<Sb87~BYBfuA5
> zS=yW;rOz1E$)f!n2mA0B3$J6&0MZi~UXYe3NCXBr!+>um<-R+JNhougs>;KD`sWLc
> zQ9}EUVZnDL`|1Oya1&X(fU-mUSn(^tc$`xnr(Ovn?-~zpja7hatTK;QMQR~qSZ*!z
> z6N#|W2~#m)6#_zl7}2uaQC#s<|Cz0E95?oG5OC|7*>WvmLU!`VE>QmZjfI_N!$C)?
> zx>}{ZUZ7RE?TN^+l2P0heDHe<Ue7_x at 1in72w-H1${@t_Mp2IO8vd?6$`0rR+M!7D
> zs`)BE{;tk)3M}=~24Pc?icxc5Fq;r?4Lx#_9bDj!Ex1~3PR4UYEPTZP!sb2*IvTuN
> z{R at Exl>4O8wN*p`Mw^rnC=_Onzrv!-Dw57=$U88onfjQJJ;Kc6qd6_7ud&ugUUx!c
> zIbVcWbrLT{4CfRsxb%ab`X<E`H-T;dEXF84juyW1DviI_>VbdPU3a<DiOUNDdl+o{
> zAVTD8Q1A{hc%>WuRD?hH9?`hNJC&bKz~@f86@#_(zk)no7K399y%}HTDi?~RnT0 at z
> zxg22xZ~y at r{55w3#kYM0zRB@#30JKnmQ!5(r4<5X2}s9I<ZHtp;b>dJY?gDiG at EH&
> zLa~-rW|nR~3M-f()48VK$BG+<z7h)o!e4aMB2?W$Oy$9!>@^)fF%<%Z$g78AnC(D@
> zns5jIQT!EmM)C+KvF1qn2cQ1AV=ZDv9$QfWS-d2#Et at -uit+f<Kgg=I0}t{-K!H3e
> z7b;CS#Cd^!?DF!3kLG4AI2E6wpkMuAgGr<k$w at 29XRvb7o-0*5Hru}>0!I`F!SZ<X
> z7jAQAGyQNR{xXXWabPW&Xqyvlt-k!y^tb#Opae+)h-{an9)BsI5b5KuFAE$>)gn=_
> zDLahu6nR?!{R=(6RMahh(X?^6#$VP1U+(Kb&3FZYBb;)zrAl6vIqaq~p(wzm<36<z
> zs{5j5QA=ZjzsFfpW0jj?Dg=?o1bJ|P5?o0Nz7oM@!jr929FS2tsS9~qfR2NN0Uq(V
> zGHXTiGd2BNh~sjZocLl4cW|H?vNxONRYEh5eu2;tm(f3`gn%*35PEDeM_xmI%U^0}
> zLQvx%-+9fX(Z3R5LSEQ^xBT6d9kAO{6?PnyamicGE+St4(GH~nZd<fjs6VzFviG2m
> zI-KGQ`(=miLkK*TEcLbEGps66p}<7Ez&E*fir}{)%Ki8YD+?)5A}E#MSgWg!(#LX9
> zZoji)O*oiQ3BunhNr-+tMe2YfuPr&S>2pB=v5-wcLp~+QJ{1P`_$xuvF9q;dh~iIO
> z$&A0&wE at 9GS0+?Bhrbdem#aU#ApTzc at W+JNTrdxnR`W&Bfq~k=hg~isrUVHJM_A}U
> zOD!Vk5ln;t_QVg1kKj+u0bhu{ycloJCbUH2=Q#2j1-4hZ3W|@e?(+pnWr_yiHNOy8
> z&MKOS6CVC*j>E0KrRkr+c9Zx=KKQCxx!|keFXC2d50ZEBPxhsNC}`>r_(~A|HEU^-
> zk+a;ylfdvdp<^3L$g4l`XAg$Np9rClvko(=5c^eEl-faXrZ at Y$kt$DRpIj*<**J~k
> z!$*k^2pV*VKN|5jn|+o^T70ZU8{nH<`BM39y9iT-hO-I`vNsp^U3~;ZM)0MBO1prF
> zSAq(GMd^nbbfrUtq)0_<gIAFs_;pHz>d$Ne46u+i)KRFtYlq#nEudV`--=%#z(=05
> ztg^HN0KE7U43&EG9sbH1FL<r;XoM|8ZL;_?8+o{NWRSQh^)wH at 0S#7&Jc|jN`DK3m
> z<%pv!ydZ3u$O!hBP+FOpRh4~l1fTM~>=3^*KkO<1!e(>FeNmpcFnB2(ms~=tR>X&4
> zG*-l2OFiO|T=BQA`&2|}D!kSVU11jmH}irc>RSe|K_=pF)A0&S>X at hbWZyOZ$ZPtF
> z2L2XJW;_LrgJu>-Te*|c?7!CHpe-0uSA<;wENul*1i>VG)ye*)(>6`I!fI21h_%aR
> zcIl-IHk-I`mI3&OB_`lAtO;2~kQOqhkdd)@q!RJVdh_VC(N~(Jo?*t&zi&0<1xxg~
> zrid|BqP3#s;?E2n926ye>RRzH{2j$BJNRn at f3#!?<nga=dj(V8sXf6*JK&h>h((Q?
> zd?z#9Z2b<C`16Hw)G_tP{;T!=WeHzVzLJ8xK6=2CsK-kIvk{U2`N|p9S%@!{sg9#6
> z1BFN{`#Xm8Y_P<s@~8=#K*Zm)Dt3=cEA+_Q4C)N#sLFcHJdA0JF{aXlzhP-oni!A2
> zC@>!9GR{J3Z~Lzj#0=~D%K0CEcu7taXr5pMBK}o=64ii`S(=d<BT`<${Okrb$Y{6K
> zD){aclu|=>SenKHNX87K*z~XdJO6uS#vd;Ahja;ys#TTv)4B01*Qm5;c1rO>2l(<D
> ze;s?D2i+7GaTHng4HiOGRTMFnr2hh8D!Iup<$_{8Xa<Ah4><T9hgU;YDv`Y_=00jV
> z#^3{5#E8>iX-grcK>TT5l)-m(p8!N$H2G@@Y(2iF?4<&suaz43SAUF at 0{9zmq at v-6
> zZot9c#a~2(0Qmwz^RU<S<4;Sbb#mVFcm81s{!~r#Xc&e+yny3RGo_M1xXDL*0sh8I
> z8R+QWxx!6F7z#;^43d;4G&lpEEL8x-CuKt6K??~Syi|>CCk_7M>1wtb2rS1O8Iwyi
> znd|reef#!NkxicqW-uYt>W}5Dg?Su=+oJfhQ?&%sTIMf^=|A#a-8Z2{U9DRQvM$m3
> zwa2f}9dZhRzW1_{V$MQ>j=9Djs|y6NsR?FgJe8(@#t(l3LMly2*GK3ysA;&SjtL4c
> z2wTW(=!&*FBSt(mQuyqcN)YdaMnr^w8B6uq`I~}ps6RrfVi*sIk>*MKoaH#|uB-)I
> zgM~ON%e3nJynp|GzdYLzue<I#aso%lMBIkroKvphU- at M`{_b;-N%nCOH{P^nJCwf_
> zQ$8DX7$f5>ftz2v{qf6tt|;Iie0{aZ$BMsL(CS(|mOs=YBOLJ;(DZLA at HPA`0fG&J
> zwpC+^%t(Dtd6rNCG-2nciZ*WG3#g6<!8DS#MaNX&n^Q?qI!tAUr2(2m-oeby=%ab8
> zrQF$k(Ww5k^LkbTA9HYOK*3c+yuvWspFjKw27kx7_I*JVZe>Rb2)E^sXBf*Zs`Rl$
> zjF{~dB-}GUw)&bSX7kMq*ZBKNp#8Z`-2K8J1mMF?{D<;a4WQ6Of%uWc6IcPopYixh
> zVf at v9H$fw5_1qRf*s!QYuzQxOYbMe0QlMxXK(y))y=fkFpT>BqUr7-lV}Ru>+6saJ
> zy_J~C9r`x|8iv_h?Z5FS5(-2EbGx2`rXh?BQ!Cv9;VWl3H5q>yr%*^ej_|q^h&lf0
> zUnx-TR9BkO1ReGX?uvimzvSC`kY<Kt1fQHFi`8vyie7*5^jie at g`f`EPh*3cz(>S5
> zn~X`(JWRzr{z72R;g2N!V5^XNmcc?`9<#yQgvR*RuP%m}oRXzRs~~J4!xSjtlNVO(
> ztZ3H*p~|c{2&)NAA_VlW;<W`h$?z9R&VmzvoHZXwQ5_H|@uDU1f{2zJrxjI?2#v2S
> zEd~4nDdVJM{KY~oWive}m?T6M7HSm>>5TaC_m$<4(R75sS~ePocKWvq;O{I%{K}3t
> z<m)I1HDqA4fg)*LNk>a;R{tG-hCc{sIP)08NPP?$EDCs<M2BBm+=NLpt)!-Tji>kr
> zR?;cF%CGUZMJwWFX2Zbu_kG`*+;R$n(BKb?J<+fc|L`C26A$|dSXpcL5nRVPzVOI+
> zu at gbsHn~^}fsDi7a`;mbJ+P-tSN$wEf%f>+A2;O}OOV%e$^u~$YkvsyA}j>{;;-^l
> z^LrzhD1eWc#u5fp8EgS6q>S6eUzWumjSeJiGguLGwP#^&9#Ie!eY-Os(W{uRTkG;4
> z1vvEy7TJ4BIQ%<0OLLh!8p3Lzntb at -htp@TB+8wN0V!4Ek1*|33S5~gTQwaz1mD(a
> zuvXLnRFmxjWJ#P at xx*KBDnSf at 9ek;$!#0|?kP1nY^kX)LA#Z+cxBT4~0<Bzts9-^d
> zzaN6|*JObbq#?g8N5}&OAU9i`8d)Vht)l*j<LF;<*_Fyl1zDd9D#NPxv)LJTbw=^w
> zg#=OI at u%lxMrk^tOk?<aL?TFkcRZ_BNN4NfW)l1vV}LLjj6bXTZGp`As}2|pi<)n%
> z1y%<E0W!?_zd-1iR{d8H?84qFJ!=LE=ZnI~Dk;f*Nh%z$OEhaSl5%`x4{cjZeG>vJ
> zT6&t<d}|4zH?6%%eC?i&Ui9x2)U~sa5U536B}MzME8P+R4s5Vun@;>~zM{Zvs>72w
> zA(ht at B26fq7gE*LSHusUm`aDv?ZBEDcFOlxW)5;$$c!g5!?;OxQb00toP*pkZKojo
> z4W at reW)+xrPj2!T0Vps*USa35zWW9UfsvZxuhf{+Zvsvu{&7fKfEkRDp$wLf_!Eg_
> zv?wpQX1iA&<t!mtT3D?m#Z);dFz^LM3RqMMFk1*vvcDLkXgl{x0cTMmuzGs|?D*Rh
> z)FrFe=s*5ap5B0^rQz at 7z-)fuFFI)JFj8G5OC({mRZ*!>GFvo_0gb|K>td>&`_<RK
> z?xrXU873|!v4cR}*6%81tZOYQ;y57ekxe3QA*G3=IfzF50f*h$K=TwnfcVzq;8Qq9
> z at CVIl1O&$tV-UmNefi5B<`Dq>OE=32H`gr33%Tj4%vzVu&Es(viiK7bB2V+0rNM_q
> zM at m{w2=LbgqT{IgV<VOY_+zDkZ`*~1zs!J<Y^<e-Re#dI%9j4I1<LNmI}%cY;_up4
> zLbUv{O$Z8MV@^nXq5qb!ibXrsn8#E0XxS$P$P!6Z!I4VF%O?DlyEfnS&ubdvV}~^5
> zT?jk{j;}zd_$WmBl2`G;- at Aa2;fB8fj=pA=k5T}6{Bdx{LR5jgw1veTsb^YwB-hMf
> z>G%su;`gmq2 at +WXkPfVpamcGeuA1K^C|?L%r33g at u_)~^gA!bcmU!ad`Nx%a)gO?t
> zrETODMun)#BWk`PY_pjp{s{YsA^i&hcVxO8r#JbIHkp|Le|IDZdvGrGNJs_azMD*H
> zPh=SW+~htff-g5^x+}Hd>?XVhOcjFKFz_|I0D9 at 4N>OFJO?LQ8I(P7Qr4Tk_A_0{k
> zb4o7ro1tBx&9~w&3Ysy at ZfPMf*iRqU9scdb_5$J;3(G0LLO^NE;H)dQ^zT{W;~JZ7
> z0r(RV{l{P3z)g{KbrwQfGfDh0XR~?mcL#sug`i4H>UIGsK-lt#y7ZMMO~?F{Ny%QC
> za9|#7;BPG-eI*TfDn|1QVJ8Ghin9la(}3!L at ydd?7zDJ=Kf_=CGM?E~!AQbT6`3S1
> zs*sG3vxGqa(t-RuXkE&5wTeK~KiRwb%GFvz;GhD&__|~1Qh at jw&~-Tytqn9HBtLFi
> zMcrVGx<O|6`~8G#Y<9)pv`&c(g9zC8JG9V-TwtM%18os>l at 8{iA|WPqj*vG8Ib{hE
> zx14#vXMp<-ftnJ;7!`?kR(G|E0V*E-<IgX8CjNd~;GNtg5>3Y=YFGSOivUNA-^Cx6
> zz9loXS<3)Lvcr!AbodbnY9Skz!@?1TQCKaF|D=G7Xcs6K3}_c75~f<kV77G at C3qnK
> ze$B<3ltg}kW-U_0B>swa_5c6n-y_bSN{14SKU}UH!Q5`H_{dz5RXym|<oH)!XbgNa
> zzz)6~;XBRq(dkHY6f{_fl~77%0FMB%16qz;r8Wnpzi$;B15_T3_+5Qef*>-6{t>`m
> zLGZ1;PgDN|4u9YJ%7kPi;wrz;t2V5aXpUiQBN+C$`W7W_3MI~vYNdd+3PJjp3k;S8
> zMru}Z$Cjan$buyU*Je`|5k~@ipa8_WP{?!uaT9C$=aC25Jo=|b<MB~{0Qb=rAjbuo
> zFM>ktikLVmLHeQ~C>$)8Iiha{JY6j+S!hpb;;CkX7Aay7P`V;%r4LMjwj3;w`p9?1
> zM}d-+ZPq>_?~W at V)qjJ<5fEO0Pk3g74j=bBWD+))2?KwEXE6Tsk3UsZHdQCS$bBWq
> zY@*KqSBVz?$pSV5&cgW10wM5`k;(|{G_Qg0pQ($8)is5h8hK5(3X$3(b488&p7Ph6
> z_?t at o$5Xrj)W at GJog$p{1Og&DQCLV?E(?J<&Bhm;xg*Ev-waBV`<!w*qB{I at RtW3@
> zwmkEo$81Kb9yF}}2nu0vjhiX}g at 8cJ9MmYV;!t at pk}x=8uyY{z4uJw>7%T|0pNary
> zCk+Szo2 at I>LZg!~7^aC4mqktNU?1 at he<3IeR at X9#EQPEMOWr9&-obf9`S@!Lpmrgk
> z#zH`oQqm(7MJ>e7c>JMoGyVpnr?0U}pEmHPPggfNMez1nrP~GvIR3OmV>(>3u3_Tn
> z<CNL5V9%aC{1rqx<WV*-21)bFGI0Tc%rxGC3=3`W8WytAVXBj-!M=rv*WmL-U-&#_
> zE3s!WK-Fa`VF<ePl-kN2{=Q=W%6|FBQIx4Hn~)MjaKyihKZAX%u}ugnL24%RB#zW0
> z08#L2EDHF8Z(2<uvNS=)-y`CsY^HyYsJ9#k4b1qr=k_3^0zo+`gx9cx&pTg%sUq=(
> zgE}5D4?SDaX2ClRd;wB_=>M?84x>E;fO4eqqk&VL!s27EROW{lB{kVRf^J<IVT`0u
> zeAIT|I>s0&C|-W~YrhagH+_X9Ygs}7)}nk at QQ)1VHh|gLMjo0NxkD0=!fQ^;Nd#@t
> z{rWSASWI)LED$g7S*{g37P16VdZ2uzKvRPFt2mRcGD1IqZ}|%cA2DZ!CS(X}W?-|Z
> zp^d6BYBJl*AR_9I?ABDIGD6>Y^VlqW%j at z&*_R!*VEiBY(1*H7a->K|c2$B=2Azkw
> zK)MQ%BqbhZGk_VY4X1EsHXE>26``gXmY466$Q?^-!l&tgOc*M>#s-owGtO*t5?I+U
> z7Xeb4ji*c|!A{I&!9ss+!Czyt1f?t>dx^m8YLO`D2n|F^3XELPD8cU at mVnVJLdFsM
> zugW6 at Yy>{?g>3?)8Lx(@JbXt|1jWv`agb$fBSJAHuw;*CTf9Q&aY>~3lRf6;zG3=K
> zp^hBg6e){RH8Zdk1v4iDrrJtGQ;L~Z7N|S~!#e_IlGH$$Tv_f}A;4c;fVCeX?2N}#
> z!VqgAaJ4RJk&U>smS99g!(WCvobZ=}@h4+hz)jfEw#&E%-_=cbgg`41WM$=uN5rWw
> zG!=jJ0I at WNqM(@jmbI{R+~I*=R5bkAkDh1Q!&)6PpbCQc@|RWC6?c`vgoMNwxXf=O
> z6JkveJ9!ek#JByqa+9+<;zE!ffry>3gYVd2Gr}20{Kd=2VrO+o0cVe>(S{vX at v&2|
> z$5W2nQl2EFs*o{Qj6CJ>rTne$=5~J&sL1 at 5jz2S)<m#iDdBiVl>0jHX0Omew6PF2%
> z5QrK_Bm(#vg_DZpEd&z5<d{$_{lp)3HroN6K=7A(N)wy8u#0~?hzE_a7V{WM59Yza
> z%yt$Yx#F)8p%6$7?h{_*)gR>uRgJ&y3krJ_ZL1T5ZPcE{9*&?S1Vz>`TMXmTGdT-I
> zO)R*{82mlr?B{XH^tQ-a{0){q*0Ly(TTarcKdNZ_?Y~T7iR=&peZq6zgAPobrGPSm
> zHqM&KgYRbd-O>8Cbf5<z(4>Pa0c)LtScR&Hkg?EkS^Z(T6wrx3zj(*fsXyu8vx=Ch
> z%6gBRKmO7 at O=6CcC{&2V0{_mR+^0PpxLnm$3fz$e_G=|hG?jNLppw{*Z3IGql11g0
> zv;tE=@u6)Y6CK+G{@k}8N#V}kYN^DODfoMe4QdyAM6gO|9{fR8ZHkN3C>a?kxO|aT
> z>_HX(vOoy5Mf$7w-R!IGutBFlSY7NqE)c|74PS$>{EuSApI^QLob2g8{!~PV{33pC
> zl63V49FvHe5U@>v6wgkH5OruYfK$d`1vuA4K`it&ex-Kg at h3>9Ag#T?#%D$QCV$x@
> z1kOS*3I`v}FTbqlXe#?SVle*9P-=L}Y~t~V+T{pp7Z3{-OrCm*VdNwP at wX@o`Q<)i
> zWTpGeOaClkIX9h-NEb;4m at 3rxBdqD9|59M$;wyq=k`{k?gjo8wqIFX>dEjqmG8_H-
> zDDK9CuT_6M<&FnyV!?6Z!9 at C7$Zmx^+L)-f%2!)}JH*c&MU~8`osazGqgdF2p->Ps
> z_iYmvf*MREkse%?BiaxJLb6t;P0FDV`<G*?)s1&{0$kS6zYu8pg%1#UD%qto8>E}%
> z=wE8!(6&-_Vv;5h<Y9o*mu=4c*ivcyqwNZmb=}lY_9~ACzAH^?_^S^l;ny?_rL&J}
> zi1?Zr|MX9Iatch}N^=Z{f8i4a at Zlvwy0S?L0gZviS~G(W5tSN4V`>rd_;Zs}sw<`G
> z;GhqMCh^<&6kf?<ezP%yRkmYS7#&LWdE{*1Tm-waLno3&XfodBYu7-&668hL5AS}C
> z5D3pIkC0lziUdI3e3aeDvqvihSi+u45RE|rAf9Lsf42GP+ at k4-PRwPBW+!5&r)p%Q
> zQ|0MuT_plZCujT<4;M7fil#+DpcWyZ5LA26&9t5P at fSf&6~d(YK$FP0qmo!i2skJO
> za+AoBm-4~^1s%$WksQZ;CFl~97mm%i5jOtbs^kJsC9K9l<XJw-F~+v5?D*BC-{Jcs
> zpud3l3zVQh#2%2X$Q~{*RZ(J(CL)@wWCJ4^=GQGT;)~^^z?6C&ie;PaMq>l<hh3PA
> zO!Jr{&rWw-hkWqcL%A9M>OcOpWPUKImu+h7 at CRSY*L<t^$Y?WYs)0lhO{Exl;bS(v
> z86d}1e(yd6CImJ(*tzD85O5ZJUi(V%!9icXaGbU&ft#?a>hzDW3a`1D6v+>N1xn>H
> zfX;FD!#mCUtA{^4A at AQI96WeXMyNu3X9k7yi*5MZXN!IzU>lPdsR;)Hyi>aXWDM+O
> zfe!3iW>7-dX-21h<ky6_NL+9Ouo_PMs4_;yFu6Htj29aI7UKIZ{wd!i^k1Lx4tU6O
> zM3W234)`jA;zRuUf?`E_cKi!(b-rA1<$FymL`0-IIN3zRcis(y&pdY86sSdBq(iT4
> zZc>ghiVyKCJK+1<H~7R4zN^$ANM}xj at mEy&E(NSe?QFoGJ(#OJLZG>6zcxn2BWdA7
> zTf;{hP9uJqmHu_aLh}*x`hA;jq3ZyHzl<oj=7hbSuKI&N&Pqr7uSmgI^OTRSn5#TE
> z3xOsUa7k4GT~cEGX<bBMlmd2vwg3+DR1u(ao(5m1j8GOR)0XYH=-QleJo1cWo+3}}
> zLXZgMf*+0~g&a|%@b|m>F8;UNa*NZ2d}Kx&jr1oNK3X at _>|vD<6d*y70yHC1eB>8x
> z#~f8`HDEO#g+|j63#y=JG152Q)mMfguVJW|YZ6=^CD8EC1tH+0r|?G at eA2Z;30cg8
> zUpVmhxSAh)#m8Ev5_+y_?n?wR6MA~)JB9O?o7CPm=c;j;ry4a%dGG}b7OQ)l<2=Hd
> z)XCm&S@;vYzMD!=A&?hD-akK3`7u|2D*h^xF%0&KVOsSKdFCNtyK#pNG>~Bw;-fsZ
> zO$!Pe$iM+Z3Pd}phaD at ZHQ4!{BwY2o=~@V^75^B<pMXVH9iuDo<L?m#-1zuIN5IC5
> znn1YXZ>sns47DV6GVIqIhMCq`v;+|sDWIBpHn23t(|qKhbO1qIRAyupWXjj^kFpGt
> zgF?(m?(-3C{A~nilZq8J2|_<|L<v&o>EECKgrF=C1$n_LW>@^ZxUr4pd=zy2D?!M!
> zjXi9$M|&#~1u5x4wbU8XeoFS~9N`*m%yz&$%||!229yQW?GZ4f`qh8qck#E4`qtq`
> zQ#0kDgauH;-&Z8v`GGd at hu0gZ@bR4iGTnVq(9p9Bln8<-1hN2dn&&lVp(cAb(bln3
> z;17lO0u?}9tWK*obVmibO!@e$HkOM#<HbTn5)@m2!;dXMYZrDEGS3o{b)}x#Yqm7z
> z4%9ef&LsRTO%?9oF9n##0GW<SzjB~Tv8A5Qp$*Sz2B&|5L`f7xQ^TKz at tN}{{b+<F
> zUc`(u__8nlg}_sVfO3SuoV0Bt_K`dQ at z1?*u+xz}=B~lF5olf+Oe?L3obWgcZ}2J8
> zM`Z_pS11q=0$7Ab7I at H9tc4nXyMRFG+kBTJ(b^RG{p)Xc?E)o&a2SKOLTYm-01$li
> zilM^*tKclwl_317!pB0u5|xv)gydML7!sJ$2b9#C0yt4f|HEI-$`?<eM6u?fUcS(?
> z)$x~8o&^nj0BI at yv<kkZ6=asvj7PLd0gsyye^D?ecRcIVVIBuCdlukq)&#;8J?Tct
> zWFVS~0P=p20UQJ_3qkr98W3#Fbg20#B5dG-Rpd4DIIA``0<q9^7WJ%uMfIzZEU*z2
> z1&2})EKUF3XqYNt9chTfO-jID>ovr>8s_mA7iL(U*H~GdYjA<D$-Y9cD;I)K`I=%L
> zf2 at GgaASiB-9g4NRmkuMU!y4bLjML#{fDb9ApIADNdX;EP at h4N^~5IJ3TFHT!YUeP
> z+pl%VKv)G=69@{&Srjx8Q6j_w2Ng+Ef$f{<k&QVSYzdHob+zXTM6qCi=8ms~z`K~&
> z0l~6;Fj55RhX7=yY(>)9gXQp2KGZe>9hSK23xS~MC_!GFZ5fUwy$I017XLONR~Tz+
> z?Y~UWHDbo!mQ8i>*I|qjgg<wfA$&U-e2uWCe*@GX*=lM0r338SD8wI0BgtE!X_D7w
> zQzaJNpTGP!ESeAyKgZkkeAM?iHA*qS$sT7>Aaz%T$Q{h}1wxFd6nvGC2$c>FD)k`S
> zZLKH at mR985L|ezlLf~6u*Z{sF#sxJHfBNSpp(5iWO9WlPL0iSshQBDtO+SFcWnUvR
> zJa<ShM=$ehI#Ys7%R=xmv5GxQgMFIHud$>buRK0gR7~lgdgCtyBYwu$L_w28#nC*h
> zvLVWM+7m%FO$Om?e!dF?{?O6Cj*mKufJ`*GcO0DUgTzm1a-Rb6SF43UT)^i(_!2=B
> z1W|tRk<=Y*h&=uba0h?o$nJo at ukcqDsoh#^_Bd<VfIl}SJK^D`uqd!%Eek{ee`QAq
> z^zrwRFqrB=(sBe+SPB8`(<(d2>lm=*FJ!_;s~Std*KtI#r}G8>>JMY^7p(XbC;k#&
> z+A>e$f<##8um3B at V&p6O at d!_;P8<c at D;TNZZyrquDnUka!F?ZHePjc_5C&Wn at CTp2
> z`c~HpgcMM(#K_y3d^Djfm=Hj1X(fpAZNajjcQM&v7jPdUrN(j;$Z<CG_X0DBj*uxP
> zPWVEs=m~*?St+3YXrjO>#?TDcqzqMynrRp5?^K>N at 2kk;AjI0D{);bGd<H)LqQG~Q
> z at HY=Lz!wGT^$6Xje_KFh6KAa|R0v!(H=z(Um7k}jk943Ylt7jUR?+WcS1I7Wa%v$?
> znFn5ppcD%M%|q>@iUGJ9NCw>RXsTWbf$T$!%u;(PK=zs_xTb$0kd?0D0toDkv1qIN
> zsQkiZ<Ac#hGgpY>k39ZTKrTo|OgQTiJjKo<iC-+FgFEcOzp|sV;=lhQ^eMl(Ar^{o
> z`u8mU;0u9{2N`KCj`NyHAg_{N#IFR&Vok?=(@IH_bJ*jfMB at cA>?AXFhW at R>Wz01P
> zZ+-I{AV!+Lk<!zgQp|GFK`rtYP{m#|Z}~SfvOs`Hv<<>|kMnv2uN-M&p~*+}r&Eyk
> zUrlW!b;J}9t2+kRd}Sg7oQ0qxV3ME&HL%Jf7w`w)jw~*|l0>Sj7h!`J>968vgAll*
> ziKeLv!InSGBaf=q*2w^7$fo#%PX-kbh at hrGF2~<TZQ*|I^8RHK9b_Q2Vy`^PFJYuH
> z*v&7*ImkzIN{8T!h{jnbY7BLZmk6R at X0f&|&5N&_q+)3w`)Npmq)9c|0Q2xy77%=D
> zhh1~>@ll%)$WZUaKGF#?mbR6oNgeYL5BqHmrloc%APQ=cy#Rk7r5<BAI3+0A!!86}
> z8KWJ_f8oAD3m{KB&<dp>R{x!!@yAnof?(I9f5KxAwZ~rvf6Y8flL*?y{LDiCAh2Jh
> zC310rj{dEp=D^ND)9SbihmWNIcZ5J+2n2{GEurq&chpSdWrt0Z4J=_2GEx8;O)NO&
> zc*$&<$DA3&pyq@<{`&OqMHqiSN*H6XM<jyKw<3=$q_)7wLa|2P60Y>G>8py&V7XJM
> zU>S-3F6EE at x8jEdE)7H8>4+nJT0Q>6FE=&y$AeJBzw;-Zd$8ud#|vfrZC-w3;VJ&?
> zVHgPs0r7+1p6jEv+^0J5IT(K*9ZobMfLhc at 5el2K%4Vd>;OfoGP2q480mnt at z&5+H
> zk)p;1s<sOY0on*x_F1Kg2>$f16$QV}`GH89a2!lCs1M45PQ_f>nq2)+{A-d<U9v8Z
> zSjA0+S_rh!e6l}O{E+VOTmCPA@(v$NXh2D%f8yWu{>4Ep{r9-o=@dbMAW_hy9tTxd
> z1f&wHqyWu}Zv|gp1Khqq?V_Mx)2CiLh^wazKwIMA%sk%;jTq4bQ`=TbE;B<!MV;b6
> zpl1fyc~KPtCLyn*f$t;o_)7tuew`y5=^z(8P7<aPeJgUoihS=livN;q>wnSX3}8_v
> z;YQWe9)BHAY2a`9cN2fm5Yy-fLvqS;=N}O#d+>?D0Y*mf(HKQo!{1XLueskx9sJp8
> zIh!5dot^Yeh#NlmSNx$)^Qj$b0~8$hg#&yFq$sB*Tbn)tXlq8+3{x3+L1{eF;jd=A
> z0cLQY7qWm0j$w9T at Z|#R5<wFMeR)xv5ct7FF0kL@*5$Y;Lkfbrq`twf_>=-_kvqIL
> ztP;eS#P7Or-)St2d^hyrLO%YkQb1A*x_Jn=HxXGtz#>bo#(+-)iVX5ob5tHv`RJyv
> z%wTEu3!iyV3O?gC6%}WGaSBhsb)a^NHq0Gg`6y)k&1Ry~*X(hw%+qlb6Sk>u6e3}v
> zb_LD-im;Cw4M=21P$7^A$WI7#BtlIVNC6Fee<l+SkC>nRWMDZ{Np01p^H2Q6LJ9I>
> zXg^XgnXL+u0UNYc5la?e?&_mXR~^J4)rta!wqQ$ekQZD2{Kek at tJ?>ShZ=URK)Cki
> z>noCGHl;~_i!vTn+pZ+-6~hvkkEEGRt~EDx at aMIlr)ohQcYrlH)UJ%g#1(%ZEhlH0
> z?IY&y^sT{e5>kg!A*lZE2!GiZ1VSN-e<4taxZn{hrdV=P54ntd?*ckvq3>)AqZH^M
> z%4`Z`bBRzCT<2pGMjq9gkNDt>m_ at 5U%;rw~L3RevY`*&H6=MWZ(`mQT$B at 5#9R5V!
> z9_o=^_T3RC<TbMA7yTp8evJY6x1Hq(UvOp>uy~r;5(sOeQ#kOUb0=O`+>{-S7hI^2
> zI!pMfAA*8U^O}BzX!!SM8WC}VwU%{-(i^n3>_N^v9f61+e^>9q-c9UKQhQQh8-nqo
> zSi~Pg<Xvmf(Z3+Z- at 2@2Iq#}J2E<>H at t}_$6t();YVN4a9+B)l2rR)fTY!W~7Knwt
> zz)jW~EEkN?DlFClUsGzB9d_<RO$o3=3M{Hc!B_5E%XZ_q)ejD$py|7#`Kl0<1-_zf
> z^*{KoP-{990#!(p8nAl`3bD{97!`o>JBxy^Nda%@-o>Oj{jj{uWq+;w)vx#z1pom7
> zEg8Ey&Xq*G=3RmbBVstKKSE%;3H$KZ at pvT&Pprs{Y;1-0Syrk1<N>}dMl8(cT<0s|
> zZ$HJK{!>2tHAYrj7+^wIjaAwz?d5{1kN9W-hY0X~ZLkM{rtfC4kQbbCbvE;Dl^=70
> zuL(iT&CWvVl}QZt(Mt`mxH$!-d{<2 at w0;Z}Amu{x5;^QrW<n7BwgY{{rkPV9z;Sx1
> z{_{xCS<7oQ-~u`R&Ob_!QZFJC2#;5S2;0&mhaO|VmsZHvgcyIPAR!QITZ7^Qd&NiP
> zXS~gh%`AtV;1Q<zh5OZ41J0 at p(`Q~+ at L}%<H?%RI1CPU`JHB!gcJK{$rx2)<>W|%)
> z{tE%}Lf{VlI}7O;;M;=-5BeVyW`k4W{2b91;3^BGfYk6Vz?Zzhffu4M$_IZMp0G#0
> zD>(S0;VF&_0q^h^P5XjidCk${U)fnFST6PsoHcJ)R#&K126f)iy}WQWzeE`R=Cqs=
> zBxeQPx?S;SxyMaH<O4bKg&-yS%1!+ULmlVzPr8~wK&|O0JH)IRDMLLa7CLGX!cZeG
> z3dolfG_S7-#BH}N|HJ+sl(zmt<Or(l#J?=?qr~35d+AVJ-SRI4iJxEb=OchZpaC0y
> z^e_=V at lsU9S_pLF51^mFH2gaPSxfV{D|y2!gItxA0Lld)O>T?;`4)fmUyP(cO%@0N
> z_)M~p`p<Yq;&1hejBElb<^iX6C8er at IV`9WvqGWysOePa-A7n4ui0&V^pyEaF>x#k
> z_*Z#W>dQq4z;`0W|L(i*mIaKI_4otss`B8kBI($O{Zz<JNv&w`2=++<Otko`W81AW
> zL|73dHR7NAWC8xw9}K~_%e4PuUcu238vMg6i1=Gt{ZV9;eIYhLM=p2;)A08qfP5Q)
> zj#CrtQq_ZO<0AvC9&May6>}jFi^5SoN!w6Es!&8&EHr}eU`gBepYqq=94Op<0Yyf*
> zqhz;|S)w+TC+x0usy{-&FV3RvCpb+A+78^&{1Qu_Fggeog4ujm{QtpU2*3f~RXU`9
> zMZhD<4)--2oJMHXm{tB~ZTS}hc$2r<C<@BBxbp~J_Y=R#eG7ZP|LDjLFKghNTwl~^
> z!&4kYMo0aD8g}pt0hMaVyBUAl#yq?gY5?7&e8tBdAE%^~H%Yrm{QMGVh=3wF0$_kU
> za>3(b1iOE{VbAr58C>tU<Bogpz1M at Zr1DU^rV3SmbTlD2{q)m0>qH^Agakz)s05Av
> z)sUJf=-5(Civ}GcfV^E#lcDi9%>Ju6|A<AhN4}0gNFU{;?6V>}Q6-hc?_2!ApJpxm
> zHAT<<uA>S`J&Q^@@HI3D743khYMK`gfzYY8?BXvMC|tv>hNq4)A{RUh1ve8vMEEl`
> z at sO_QTS(=>pX^=NBsk!Izzu)GASxTKU06&1Qb0#V5(0m<^b<YrC|P{P)N<P<5hNYp
> zek#=JP_=0(HvRK`mH$;@;SP!7 at 5cvC2@<bL7Wf+Y=^uY at SV)>*2o!c7g_RzvPL4wr
> z|0)k>Ic~#pwXb<86Kn2UU2(PnsXs!XwaLc63h&w>C)z?l*c7gTtm$Bmyhgp0jHIjq
> z+(Oi at Lyb(80#0Id;@4~ioWabHbaJQq-?Q|8&pr15VkFI at t^N?aSokO-!0+a1>wh_A
> z6(JS^tKeT25E%Y~GT4zPX;lkGScuSIvL`~#Q|{o8`SR=A`iED~3P-g|dGTQNA7srA
> z>k$HzgJeHj1A#gGXi~sJWDmPwIi_l&pf3u_7Yp%Md at A_(>(ecI69U+YzorCHJLXU$
> z=88P#1V4>>1&C{usror)#J?)!wT0ylHW;emoaKTQJz^eN;P7+befK$n8c($Tfl$DW
> zr+oeNzwQakf6h#(IEb=PLg6Xo3jrHAXp+c-FJClag+(%w1Q80fCU)-NZ-ZFnf2Li)
> zv-p<+Q{G9v#t{bqS>r*>@fQVOxnF|Vf-RKh36Jl}uVm2|6j9J;qy)rYs{|?9Y7YJi
> zJM3+>BmM-Byjp>nCJ+NZfOuL<suBSH6$j4hFql=UtAAT^#1Tg*?6A;}rR9;tL0<8(
> ziW&Ui=D(Bb%@=UA!=C;-7x|fjq{Z4oLLd>4#2?sFV9URaRVCp*Zbl;CRwz<i{{H@@
> z{?LEoZ|@d36$W9nYfk=*hf%?aKlnwVBParFGmJZ$kzK)epBWx!C*{*T?42j8Ij;Cm
> z&BrIfyTXotMMFp3K%VlQTlD$j3`X0cKm0i>5qxBlw;ST8Z6VO%7jy)`04qu{^^g=#
> zJ at r%r)O^5^7c%jhD7e#o!-$j2OfH*5nE{2sReY!6pZ at tCfADRft{zbk*jD|g5d0Ir
> zc;PQEbmGrDNvFbaQ0h at No_&?_Cku#;BdU#`Ai%E?p^hxDFDOSsV0EtmHs2{bYCir_
> zfbyk4vIji!T_X>;=G>liJ!pQi(QNn3OOuv}zoNu7*pFNrhJI|J1Ts*K9SAjL#4!3-
> z_NN3Pznp~Dvyi8>Dw0GnlKXtbP$(^AwzaNYE<56{6MtnKf5VKa_zMAf2%kmChWI=G
> zU at o#+`q!z*6F>Oax?&5yEAl0U&}d>*3fPDZut`yVQ2<=?Q3&wYVYWi#L_~q}G6lX6
> zRC(HD at z+v*HAlq&zq+AryNR<L<S92%#ea6j_)As5{V5q(c<rIwBoZn at fF&9 at wklrx
> zvx^cm#~zwTTSMNnn(Aaz at S}tkO~Nn2w$bB>s!*;af4Sh_Rbfp3TrT3SB2G5xpLv?%
> z!+_wkP+y*UR-ieEe21LMJHqZn-hH9bFHB-e=SmRTnh@}mol*dKtxchPREV55V@~#}
> zLHf5pI;G4}hyGX9FJ(vbpL!9$xLkYfwU$e1?6fYV)qg;AgoB;bMaJYGHwH69=|J89
> zj{wLX*^L9*s*sZp!7D+P9n7^g0?VK2jp=teG|V7;wj25<Z#BgJi at c;0BQ5<~!X6<&
> zwm<xthsp3)Au%^nQ+AddIB)so)5rKDOxT>&q&ZW0LGzgVN?+rtJFZ3=Lyf{D1ju6^
> zc8jixTO>yhjze}&HKdx;Q4 at Lfr`A!k%0umx55D5#KKRnG^P7}c0|gyMD5xv`G)5Rs
> zcZvy~e$S^U!!T!n9kAAUl7;p%?^#4%3gGWSAwWL<U4=kFb8a!d*-8QRA1);0F9bRm
> z(trE~$RQkmz6%`w1`E0{3EV>eKl!t^?G;soWt at hstB>HTIqePw$W0xC&8!KN5R?K%
> zf#VSI9U^*Af?#wCl?C=eTdjhFE#=qRe22U}_s%=-6dDhTubS`2Nydu;W5l|dO|GFT
> z+>F7W^3 at h$q79b7&Wzyefa5UfY6cum+4=vu7<^6c#2<XE5+nuufsAgsuh^QzawdUV
> zxo`?%m2dGU<M>OE^e+ndL&2EpKQBBonimC6Y51FsYw;(3<WY at 3X|)Sf`Hf+WX1qXX
> zhRM+QOL>k1!qiUAWYCX+?Zt|J$8xkqLCw*0G{r~RY0q`FadVmF1O}<f$#TZHiGTWc
> z)wd#J at Q1WbGr1~0%y6xvZ!rEoVuG7^wc)z91;ehXOLg$~M(oYP>K>6ho|W1<yk-yb
> z=3yR~q+))0OaEK`9+yWdq>)qhgYViQ6W-wuevQ1XTiKWUus6m{&9?=sy_i+tXtufP
> zt@=O@>W{M at o~n?uXs-b3581cXBCq&}2>xDFy)KE7ja62HPcx!Gv1%p*s22p}HN~IH
> zg3fpc43)<=!6XrhcH~O|CY2FFz(O^|c!Z;^qfx$(-Eox>QbSmuvHt%sh_G5`5;M at 7
> zMHv=<8m9mF)4$9e{wU#Z1_}``j at WkjGHv+dY_P~$sM^jKfpAm7v78W;ChqKnX!*bN
> z4+{>i$~F11Q>rSJ^bdZ;$6J^Z<TxY?6e5tTMWUdAuktwU%Lw}ycL<EN>JNi8zb!CB
> ze22dng at RScQ&+XR(~hr1K{H-+b3)*w`Hk<$6lz&e<?*o&{(@o`aG*kuJ)AX(L#D}j
> z4IQLJLdIMD3&EWKF&zG*|2#6G#Nn6v6I-e_o^9bR96kz%uY^F80>-$bS=6`bAAhQF
> z_ at plT4t?V9z^3?n74;V2b<8d-1icYB2nm5iuoqx13-Ff+UYZD)NsLS+ywe%}kgEKW
> z%QgPWAQWuttYU!s3Lru1r+ at reg`SR^6e60a+IelnF^3)VijQkMf&r?KEbvc8>4D3%
> zEd<J}b)|q1u+YKGiXD)s4uADuppnGiLQ0TR6FI@|TEm~^oMpW1208u5A7>xo1zD#U
> zZPUasUseCX*GdHII<&KasTCj0>EBgz{siAkfhG|OfxXyG+k`qM#5`h{2D_^S*WoYy
> zqZ)r9))y=`u!NnMSUT+S_YwcD_-kc{Z!M$>opHt)usgUGg_sM6rtcIaFQkcMlIFcw
> z!BJ5WR{bP9!=xUmh>M_enfQGx5C~WrLnRb{6;dbc;%Fn#-|`Qv at xf2wuv-OTKEgsX
> zAQTrXY89Ns4}}oW7)z{fGphdhK at xV#S9!GHyK4ArHU%4me33hr=sYIeH6OVt3e_KY
> zk=K+s4#Ftu#NFyX1|M)D4?E!QYYeNoZ<xnDOM8XDyBp?Wq3NXm*`NhNIDC{2MB;?W
> zJF`)2c?wtjeN_nL2pa$#{i{Fn1^M9PFRz7wgB5nORol%5xF)TFkGXFfF#HW`z_RbI
> z|9hSw&=~L32k=P+*nSAYU;UwZ*j;6T2ENenh#)N_4Qa9}C)^|<E18#0 at P`^ngVEzJ
> zwwAM~dV&UjH8=HxZ}9N%zH9vPk)Cv76(OL1A#kn!6bRVsbQPC9d-m8~?NA`ofVhhw
> zZ4j-c1;zc&|K`WReax#rDwPx%@fSGKA~Viz at e(%pfYd-XkA7<R74pWoBkSE{M#ml%
> zfP*Yyo8s^Asrb{sX6NnvL+$u`9dlKb4D*O3;3Wez7$eJEHT-k1`h&m0$a4+b>PfUw
> z)qLAc-%fM0W~7d*8Ng?WJqSz$RCvEjXdox`YEj{EQwa2RJP5L<B&-#YVL2w2bBF5U
> zug^&Lad!B3{#2>t0@=IMmHtO>_*Yv at xKx%+2u$K5+RipmgP`LkXgbI%JIGs3V-kJR
> z6*K4`e<|SbDGTCX2_k#vA|g=)%wy?kS;8ss<pTXPQZqoT(X)yW^EdMPMa2DK?j}oo
> z1YAUz=PFI=+d&xSy~E#S?Ec2!1HKdhne3f>Dn9K3>EBmA5=@YF$c*oCFiA&l0rC~s
> z*g&XgdsA~v3Vw$QyRas&sf*;emj2b2w7o{){{I!4L9NJGt`L+C?f|H at _u*h;VuqcF
> z_#J+nXKcZ)QUL#sK*B*n!elm+TTz|Vs%)-|!N2+=1Z at OrqtM5{K*QUL7B!nb=DgO5
> zQEG=h^1(L)fW+ at b+fRI6_pm|<b{lb0kG~KAULbrMf3CsHev^~}3ZQ01s_O`3z+?yD
> zl<)AP>|^d|p;h at 6F0M`1^BRiOUKFY at u9_(ON(ThoNKx_Wv=2TMyvBhIR#(z7r+n?l
> z>)(Qx8nkbpKC=bLBW5tCW)jx|VO<Hh4FU5q5`QBxv`~T%z83-O+7Vo9ViA9+wZ!ki
> z8qu<i?1|s&o=Rv2b=Wx{YW%fiFAA_11 at PT>9sWYVQ?pr%`wWP`n}vp*L<O~-m-vM^
> z_++njsGxrWlj-!v>ybSg8Z41t0<aST_mS5C1Vk*Hf6 at SF5g;eFDLyKU6krUS34<h9
> zYmAa*W~iGs{@r0fAt)V4!Y+`7_*+q~C3}#q1)XAIT5b9|Lc at I{wXy(w9h7XKn0wd8
> z-#eg6&IKG$@(-~9pR*A1T2UDO94Q2C1Zcy~5 at tB__pPhGrD+Ls;;;DV!>-{^{5Cb!
> zwGb#lyZ8e#E!sY=VY#3G8ieKdPdtU7t4vn}*aHZYLg5bVzB0H;aBgxTz#mw%xoQ4F
> zfIr|=K^ye%>W<ViiK$Y)JgQl!4it58E(Dg=kWc)u*z~>bf$t}7r!ryUyZLcYD1?Cc
> z9fVw+HvKf at TJfQO8^Q2Lp3FqS9gD`F#ze<Kh$R-qKm9xT*mJl1mmJu6|Ft@$anN}3
> z8#1wS1qY>40H1UVTCMxg=@{v1FjgW0cn^9?YpTX{$5SSmJ66}2G5pmZ`nMOjPpIJI
> z2`_&o8lEv9@)HrECxfEkDg2=(eurM<6F-JbvbmdIxUEI2&O!l4-cLZ at zeGVeN&$r7
> z^`M6Tgus1`NnM$NKQFfY^+*5lmj&W3$BCcU8aJ6VB%>Nq*V-tyk*=DNn#V1=W$0t%
> zJqxvG(Kg;q<Gbm8cih}QOaJj_p#h2y_>f`?c%BY_o%)5qRSt at j`dkY5*4BUEzya{p
> zA4g*MiC=szls!HQfh<sYM4=Q=g7DYqKf*+=2?6mdur>V&fwfGI=hUALJb-JW;J$F^
> z<8O7*=eQP=uq4a+^bFm!ly5*|;_q=Q3gtS}a+!ecdeEE}@{~@K%nY6qGL1=9hWHq-
> z_<)Hj+4GbRUG+<ZeAk~S7%2tp6qK*yNanBqjy<+r9V0C&z811v_P8S#>;<x*$}a at V
> zHn4_62uyB8sI?3z5uCS0gbfCCjX#|${$9KAhkeDr at Yj4L_0oU**^pM5fpBm_?Fy7D
> zl6^IC5gK<qJLq)Gqt!9e7_DcEMHBj}W<{+ at W*$~izVav#hT#t;2RY7J1YGqsFsD(z
> z>IA+tX(R9yN~*oe?^jm8xH<yKCaCcr!B9(PSkX^CetE0rkUjF!L9 at DrT2vsST?jJ9
> zx5FO+ko(zz;K83?v<d<Kb~k2l(CXaRc<PEzZJ_zMMK>QCz;fI)^DO8uSjtOH;O0S{
> z0?{gldBioxo1~lOu`96+4H at by{PE-+@?7Sx#CP>^;Ku~{;KyG?9DV)Pr}ERkS8oBL
> zd6_k}w~P}D3o)B4y|?=P8Tocz9a#{6HV8p$xw3&hV~lUjZ`0u#&iK>6Y9kiOzK(D>
> z{}2H0+JD2v00i*a3Lb5DV(vO$*<s`9jt7mIYmT`3YEl1A1=Hy%_l=LXr5oAf&4vm1
> zS!kp^6xH~L-L;OJ$Y at Y&EB<Bwt>ct({q@(Ay~^XY*V)*Cnb+XwJM5ZtON3U}0mtu_
> zC!KT>@=}2K at i#f(r9k}M5exj4pmt<>aMS+FHG&7<%UAkGzLx%dTXl`UcohOykXL5^
> zS2Tj(&4KA=_sfELB(|Vw1^})Z<9 at l&488)KomT9t?zox$nOB8050B^y>S@&hS9j3h
> zwE-IJ;McHHZEn8#X3D?)_S*$QQt%XfS<i9UwL0<7W+s&dC!c&W@}i&#ar~`lW_`=)
> zlL7vkBbHdOgj)C|i1F|6L-Ub$P5)Kb_*>m08uxL4C}{eF*PTHZN!La$h}Sj)Jcz1~
> z&Y~ShgY1sqE*j<zPpzoqX*apg27)9`1T^OP%I3?jah&h<<Bx`#U)dl2KGG!SD!=06
> z4-GzA0({N<+YneW1u}*W>;Z&C{W at eXpnvZL@h=3W8~#4hyqUet;jhCm?_bp)`>zrd
> z|LXrlq58uH0<hNTKNQixFk^Jw!41iwS7><Tt7(BPE(Z9h`L-asGdJC_*7)mwUIX6d
> zn1u|pUD7R=aTWzlpPQj@^%08UzeM8V<v&UD8t!*a2TnirICBajf7P_M6$>q;XKhvy
> zcVb8n^^u`fzaiTTRH0eU%)>vcq<}()zbFWYLzf9vn{EBaU$gBBflj4$R<~}bVSx*@
> zfNc4%LcX8>%5?!N99oZbGx#xgr;o1F=_cAHZ{6`)*@p!mX9N5Rl^Kq6cm}zuE(BqB
> zQ`2{+5Ez*((YWD;8~lA=<#FbO#fqj?);e`8rDtsx1)MR7L0_Xu%}5x%9or6uzm6re
> zMG8A&xbn*7zdf|S=uiufJV^gaP-P!>4gWO`%O8JDo(3!z=K$iV!*Mse4;`NPdk?Lc
> zG_so|h5j}WJ~n0rzmzH88rjW)oBq4<oxgaHD8qnap^qN%+LfDJuCVJw-gRZLToxib
> zObF8N%qoB4_sfGFL|F6M5`<?t<hKeWTHIFr%m3b*uO>jgl5r(n;m)L1cjt;L?*E at 4
> zlL(B&4eEaRguTkI!vOa+ozl<}EHNIFphXF{B<dP}05xNLG`O20Qky7*dK(UX-}<FU
> zGlKz3T-P$&mDg5b2AAW{JR$JWafLh4hBy91p8M{{JlQV+SQ`lb?%=?q^7|(iWXT2)
> bOly7%;i&~5YIigrQ@*d<3_jb$PR;pWU;lcu
> 
> literal 0
> HcmV?d00001
> 
> diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
> index 2ba01c4..5f1694b 100644
> --- a/utests/CMakeLists.txt
> +++ b/utests/CMakeLists.txt
> @@ -9,6 +9,7 @@ set (utests_sources
>    compiler_mandelbrot.cpp
>    compiler_mandelbrot_alternate.cpp
>    compiler_box_blur_float.cpp
> +  compiler_box_blur_image.cpp
>    compiler_box_blur.cpp
>    compiler_insert_to_constant.cpp
>    compiler_argument_structure.cpp
> diff --git a/utests/compiler_box_blur_image.cpp b/utests/compiler_box_blur_image.cpp
> new file mode 100644
> index 0000000..812a906
> --- /dev/null
> +++ b/utests/compiler_box_blur_image.cpp
> @@ -0,0 +1,65 @@
> +#include "utest_helper.hpp"
> +#include <cmath>
> +
> +static void compiler_box_blur_image()
> +{
> +  int w, h;
> +  cl_image_format format;
> +  cl_image_desc desc;
> +  size_t origin[3] = { 0, 0, 0 };
> +  size_t region[3];
> +  int *src, *dst;
> +
> +  OCL_CREATE_KERNEL("compiler_box_blur_image");
> +
> +  /* Load the picture */
> +  src = cl_read_bmp("lenna128x128.bmp", &w, &h);
> +
> +  format.image_channel_order = CL_RGBA;
> +  format.image_channel_data_type = CL_UNSIGNED_INT8;
> +  desc.image_type = CL_MEM_OBJECT_IMAGE2D;
> +  desc.image_width = w;
> +  desc.image_height = h;
> +  desc.image_depth = 1;
> +  desc.image_row_pitch = w*sizeof(uint32_t);
> +  desc.image_slice_pitch = 0;
> +  desc.num_mip_levels = 0;
> +  desc.num_samples = 0;
> +  desc.buffer = NULL;
> +  region[0] = w;
> +  region[1] = h;
> +  region[2] = 1;
> +
> +  /* Run the kernel */
> +#if 0
> +  // FIXME: CL_MEM_COPY_HOST_PTR yields a garbled copy.
> +  OCL_CREATE_IMAGE(buf[0], CL_MEM_COPY_HOST_PTR, &format, &desc, src);
> +#else
> +  desc.image_row_pitch = 0;
> +  OCL_CREATE_IMAGE(buf[0], 0, &format, &desc, NULL);
> +  OCL_WRITE_IMAGE(buf[0], origin, region, src);
> +#endif
> +  free(src);
> +  desc.image_row_pitch = 0;
> +  OCL_CREATE_IMAGE(buf[1], 0, &format, &desc, NULL);
> +  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
> +  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
> +  globals[0] = w;
> +  globals[1] = h;
> +  locals[0] = 16;
> +  locals[1] = 16;
> +  OCL_NDRANGE(2);
> +  dst = (int*)malloc(w*h*sizeof(uint32_t));
> +  OCL_READ_IMAGE(buf[1], origin, region, dst);
> +
> +  /* Save the image (for debug purpose) */
> +  cl_write_bmp(dst, w, h, "compiler_box_blur_image.bmp");
> +
> +  /* Compare with the golden image */
> +  OCL_CHECK_IMAGE(dst, w, h, "compiler_box_blur_image_ref.bmp");
> +
> +  free(dst);
> +}
> +
> +MAKE_UTEST_FROM_FUNCTION(compiler_box_blur_image);
> +
> diff --git a/utests/utest_helper.hpp b/utests/utest_helper.hpp
> index d882fc7..c28615e 100644
> --- a/utests/utest_helper.hpp
> +++ b/utests/utest_helper.hpp
> @@ -88,6 +88,12 @@ extern EGLSurface  eglSurface;
>  #define OCL_CREATE_IMAGE(IMAGE, FLAGS, FORMAT, DESC, DATA) \
>      OCL_CALL2(clCreateImage, IMAGE, ctx, FLAGS, FORMAT, DESC, DATA)
>  
> +#define OCL_READ_IMAGE(IMAGE, ORIGIN, REGION, DATA) \
> +    OCL_CALL(clEnqueueReadImage, queue, IMAGE, CL_TRUE, ORIGIN, REGION, 0, 0, DATA, 0, NULL, NULL)
> +
> +#define OCL_WRITE_IMAGE(IMAGE, ORIGIN, REGION, DATA) \
> +    OCL_CALL(clEnqueueWriteImage, queue, IMAGE, CL_TRUE, ORIGIN, REGION, 0, 0, DATA, 0, NULL, NULL)
> +
>  #define OCL_CREATE_GL_IMAGE(IMAGE, FLAGS, TARGET, LEVEL, TEXTURE) \
>      OCL_CALL2(clCreateFromGLTexture, IMAGE, ctx, FLAGS, TARGET, LEVEL, TEXTURE)
>  
> -- 
> 1.8.1.4
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet


More information about the Beignet mailing list