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

[Dag Lem]

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