[Beignet] [PATCH] utest: do not check MV near image border

Guo Yejun yejun.guo at intel.com
Thu Mar 17 01:42:24 UTC 2016


if the image width and height is not aligned, the VME hardware block
could use the data out of the image, there is no clear rule defines
the behavior of this case, so do not check the MVs near the border.

Signed-off-by: Guo Yejun <yejun.guo at intel.com>
---
 utests/builtin_kernel_block_motion_estimate_intel.cpp | 18 +++++++++++-------
 1 file changed, 11 insertions(+), 7 deletions(-)

diff --git a/utests/builtin_kernel_block_motion_estimate_intel.cpp b/utests/builtin_kernel_block_motion_estimate_intel.cpp
index 12bcb7d..008b27c 100644
--- a/utests/builtin_kernel_block_motion_estimate_intel.cpp
+++ b/utests/builtin_kernel_block_motion_estimate_intel.cpp
@@ -48,7 +48,7 @@ void builtin_kernel_block_motion_estimate_intel(void)
       if (i >= 32 && i <= 47 && j >= 16 && j <= 31)
         image_data2[w * j + i] = image_data1[w * j + i] = 100;
       else
-        image_data2[w * j + i] = image_data1[w * j + i] = 0;
+        image_data2[w * j + i] = image_data1[w * j + i] = 17;
     }
   }
 
@@ -61,8 +61,9 @@ void builtin_kernel_block_motion_estimate_intel(void)
   OCL_CREATE_IMAGE(buf[0], CL_MEM_COPY_HOST_PTR, &format, &desc, image_data1);        //src
   OCL_CREATE_IMAGE(buf[1], CL_MEM_COPY_HOST_PTR, &format, &desc, image_data2);        //ref
 
-  const size_t mv = (80/16) * (48/16);
-  OCL_CREATE_BUFFER(buf[2], 0, mv * sizeof(int) * 4, NULL);
+  const size_t mv_w = (w + 15) / 16;
+  const size_t mv_h = (h + 15) / 16;
+  OCL_CREATE_BUFFER(buf[2], 0, mv_w * mv_h * sizeof(short) * 2, NULL);
 
   OCL_SET_ARG(0, sizeof(cl_accelerator_intel), &accel);
   OCL_SET_ARG(1, sizeof(cl_mem), &buf[0]);
@@ -76,7 +77,7 @@ void builtin_kernel_block_motion_estimate_intel(void)
   OCL_CALL(clEnqueueNDRangeKernel, queue, kernel, 2, NULL, globals, NULL, 0, NULL, NULL);
 
   OCL_MAP_BUFFER(2);
-  short expected[] = {-64, -48,
+  short expected[] = {-64, -48,   //S13.2 fixed point value
                     -64, -48,
                     -64, -48,
                     -64, -48,
@@ -92,9 +93,12 @@ void builtin_kernel_block_motion_estimate_intel(void)
                     0, -48,
                     -64, -48};
   short* res = (short*)buf_data[2];
-  for (uint32_t j = 0; j < mv; ++j) {
-    OCL_ASSERT(res[j * 2 + 0] == expected[j * 2 + 0]);
-    OCL_ASSERT(res[j * 2 + 1] == expected[j * 2 + 1]);
+  for (uint32_t j = 0; j < mv_h - 1; ++j) {
+    for (uint32_t i = 0; i < mv_w - 1; ++i) {
+        uint32_t index = j * mv_w * 2 + i * 2;
+        OCL_ASSERT(res[index + 0] == expected[index + 0]);
+        OCL_ASSERT(res[index + 1] == expected[index + 1]);
+    }
   }
   OCL_UNMAP_BUFFER(2);
 
-- 
1.9.1



More information about the Beignet mailing list