[Beignet] [PATCH 1/2] utests: Correct box blur

[Zhigang Gong]

Good catch, pushed. Thanks.

On Sat, May 25, 2013 at 10:14:16AM +0200, Dag Lem wrote:
> The box blur test kernel incorrectly calculates the bottom extents of
> the image. This yields visible blocking artifacts in the top of the
> test image (since BMP images are stored bottom to top).
> 
> These calculations are corrected, other extent calculations are
> simplified, and some dead code is removed.
> 
> The reference image is corrected accordingly, and is now identical to the
> reference image for the box blur float buffer test.
> 
> Signed-off-by: Dag Lem <dag at nimrod.no>
> ---
>  kernels/compiler_box_blur.cl      |  51 +++++++-------------------------------
>  kernels/compiler_box_blur_ref.bmp | Bin 49206 -> 49206 bytes
>  2 files changed, 9 insertions(+), 42 deletions(-)
> 
> diff --git a/kernels/compiler_box_blur.cl b/kernels/compiler_box_blur.cl
> index 0c6b657..26936e0 100644
> --- a/kernels/compiler_box_blur.cl
> +++ b/kernels/compiler_box_blur.cl
> @@ -27,7 +27,7 @@ inline uint pack_fp3(float3 u3) {
>      C2 = (from1+from2+from3);\
>      C3 = (from2+from3+r);\
>    } while(0)
> -#if 1
> +
>  __kernel void compiler_box_blur(__global const uint *src,
>                                  __global uint *dst,
>                                  int w,
> @@ -39,27 +39,27 @@ __kernel void compiler_box_blur(__global const uint *src,
>    const int yend = min(y + chunk, h); /* we process a tile in the image */
>  
>    /* Current line (left (1 pixel), center (4 pixels), right (1 pixel)) */
> -  const int left = max(4*x-1 + y*w, y*w);
> -  const int right = min(4*x+4 + y*w, y*w+w-1);
> +  const int left = max(4*x-1, 0) + y*w;
> +  const int right = min(4*x+4, w-1) + y*w;
>    int curr = x + y*(w>>2);
>    HFILTER3(curr0, curr1, curr2, curr3, curr, left, right);
>  
>    /* Top line (left (1 pixel), center (4 pixels), right (1 pixel)) */
>    const int ytop = max(y-1,0);
> -  const int topLeft = max(4*x-1 + ytop*w, ytop*w);
> -  const int topRight = min(4*x+4 + ytop*w, ytop*w+w-1);
> +  const int topLeft = max(4*x-1, 0) + ytop*w;
> +  const int topRight = min(4*x+4, w-1) + ytop*w;
>    const int top = x + ytop*(w>>2);
>    HFILTER3(top0, top1, top2, top3, top, topLeft, topRight);
>  
>    /* To guard bottom line */
>    const int maxBottom = x + (h-1)*(w>>2);
> -  const int maxBottomLeft = max(x-1,0) + (h-1)*w;
> -  const int maxBottomRight = min(x+1,w-1) + (h-1)*w;
> +  const int maxBottomLeft = max(4*x-1,0) + (h-1)*w;
> +  const int maxBottomRight = min(4*x+4,w-1) + (h-1)*w;
>  
>    /* We use a short 3 pixel sliding window */
>    const int ybottom = min(y+1,h-1);
> -  int bottomLeft = max(4*x-1 + ybottom*w, ybottom*w);
> -  int bottomRight = min(4*x+4 + ybottom*w, ybottom*w+w-1);
> +  int bottomLeft = max(4*x-1, 0) + ybottom*w;
> +  int bottomRight = min(4*x+4, w-1) + ybottom*w;
>    int bottom = x + ybottom*(w>>2);
>  
>    /* Top down sliding window */
> @@ -78,36 +78,3 @@ __kernel void compiler_box_blur(__global const uint *src,
>      curr0 = bottom0; curr1 = bottom1; curr2 = bottom2; curr3 = bottom3;
>    }
>  }
> -#else
> -
> -__kernel void compiler_box_blur(__global const uint *src,
> -                                __global uint *dst,
> -                                int w,
> -                                int h,
> -                                int chunk)
> -{
> -  const int x = get_global_id(0);
> -  int y = 0;
> -  const int yend = min(y + 64, h); /* we process a tile in the image */
> -
> -  /* Current line (left (1 pixel), center (4 pixels), right (1 pixel)) */
> -  int curr = x + y*32;
> -
> -  /* Top down sliding window */
> -  for (; y < yend; ++y, curr += (w>>2)) {
> -    float3 d = (float3)(255.f,255.f,255.f);
> -    const uint4 to = (uint4)(pack_fp3(d),pack_fp3(d),pack_fp3(d),pack_fp3(d));
> -#if 0
> -    dst[4*curr+0] = (int)dst;
> -    dst[4*curr+1] = (int)dst;
> -    dst[4*curr+2] = (int)dst;
> -    dst[4*curr+3] = (int)dst;
> -#endif
> -    dst[4*curr+0] = to.x;
> -    dst[4*curr+1] = to.y;
> -    dst[4*curr+2] = to.z;
> -    dst[4*curr+3] = to.w;
> -  }
> -}
> -#endif
> -
> diff --git a/kernels/compiler_box_blur_ref.bmp b/kernels/compiler_box_blur_ref.bmp
> index fd910089f442157a62ceb722978ac152aad7cf5c..149cbba1d8bae1fdd4e4803970e06ec1682a40fc 100644
> GIT binary patch
> delta 771
> zcmYk4&1%9>5QP`MS>Gc>5Ns<36g4qMb7Nzyt87HE^#_cCqMN!fyDnY15fPuqFLyvZ
> z(CNv{y?4%>F%5EQko$Z4S3erR8q?`?u~@8DtIcM!+wCBu(Wu>S7mG#1Fig{gfDesQ
> zsr0E)E|;s-YNOF`9LMuK-}jYLUmCq$kN-gskd%D4+YQ4oilTnMPYhKU%2COZ;tpKm
> zvm}Je08W*YlH_&zflj9*5#qY8ZQJ#F9c7kfA(2n642MG|Wlp9ZXHa7p-vY6L`0C1X
> zxjY;Wr_<?nyIrr><M9~Baivm`x#_{&^lLVo5YTG1 at Ju3p04bpWz$sxfn9XL2Kk14B
> zA!L#y*>1P{{eHb(BNFRyJD<-N3I)7>*MOjaJb)mmR;$qfDV)a at T*M*q3ES}Ibp==m
> zGXJS at n#cvVB6u!B+<x5%oa(m*PjG;D<!OmbPcw92hO7X6Y=O_gV8A?LwqAjoEnza5
> zF#0nKCnGD6hz)6)p3mp|{eHPz_-!Hr7kHNjbiz`Qx$ihQOa{O)A3}+ww%<k^>zmK#
> E4|f=m+W-In
> 
> delta 771
> zcmX|<yN-fj5QPihY~RCZV&P at Iix<2Ain^=@R}(erf`t+_R9Z+!X)JuH{bZ6o#h)`X
> zKmVCILtw-LBX<1=_1$i_-EKL5BkcG4<MDVpoi3LPWVKoi27^o{6OYFeiNt3FLQpQ3
> zi^XC#o6YC*jYgy0Za16FPN$>mx at nqx1;*p?a5yv!L(??L^!t6P-w{Tmk;&b!2xN)L
> zTb2dqGeRnSCeM-KW6<mMq)G|i?RJp_5ah{#R;yL3)f7dkR4Ro+L0aTN+qM~rNtt=E
> zSj^}897eoeuPI;yfGn3waRuTy&h2)4JRZ;Ilk<MR&t@~G=UrY$gXz%&7ywjNr4{^;
> zhUkeBNv_GwdcBU#fR}aGb&*GvlNiEzI2=6BJD<;-bT*w%OQljWnM|cpjD>1Con{3*
> zhkP!VV?G(33`!V-b7I?MGJ#_o4zUhAfK$Q=N^k#(VQ7%$2EzgFVKQHmhftDG)`1u$
> zKvE`S!3Z!k9WXZ>C#uz|tbiW{x5%(ufb+5f+g5E>zS(S;JKD#49q^wBa6qTr4uatI
> ldWB&K at qM3$FaTE&po4$&(VzoB5>7)9PFVrY(;VZm)<5J`ku3lK
> 
> -- 
> 1.8.1.4
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet