utest: do not check MV near image border

Submitted by Guo Yejun on March 17, 2016, 1:42 a.m.

Details

Message ID 1458178944-2865-1-git-send-email-yejun.guo@intel.com
State New
Headers show
Series "utest: do not check MV near image border" ( rev: 1 ) in Beignet

Not browsing as part of any series.

Commit Message

Guo Yejun March 17, 2016, 1:42 a.m.
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@intel.com>
---
 utests/builtin_kernel_block_motion_estimate_intel.cpp | 18 +++++++++++-------
 1 file changed, 11 insertions(+), 7 deletions(-)

Patch hide | download patch | download mbox

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);
 

Comments

Now this case could passed when previous test_printf case has multiply tests.
VME engine seems to read data out of specified image buffer which is based
on drm bo.
If this drm bo of src/ref image object reuse from previous bo with garbage
by coincidence, it will cause different MV results.

Yan Wang

> 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@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
>
> _______________________________________________
> Beignet mailing list
> Beignet@lists.freedesktop.org
> https://lists.freedesktop.org/mailman/listinfo/beignet
>
Pushed, thanks.

> -----Original Message-----

> From: Beignet [mailto:beignet-bounces@lists.freedesktop.org] On Behalf Of

> yan.wang@linux.intel.com

> Sent: Thursday, March 17, 2016 17:33

> To: Guo, Yejun <yejun.guo@intel.com>

> Cc: Guo, Yejun <yejun.guo@intel.com>; beignet@lists.freedesktop.org

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

> 

> Now this case could passed when previous test_printf case has multiply tests.

> VME engine seems to read data out of specified image buffer which is based

> on drm bo.

> If this drm bo of src/ref image object reuse from previous bo with garbage by

> coincidence, it will cause different MV results.

> 

> Yan Wang

> 

> > 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@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

> >

> > _______________________________________________

> > Beignet mailing list

> > Beignet@lists.freedesktop.org

> > https://lists.freedesktop.org/mailman/listinfo/beignet

> >

> 

> _______________________________________________

> Beignet mailing list

> Beignet@lists.freedesktop.org

> https://lists.freedesktop.org/mailman/listinfo/beignet