[Beignet] [PATCH] utest: do not check MV near image border
Yang, Rong R
rong.r.yang at intel.com
Fri Apr 8 06:03:55 UTC 2016
Pushed, thanks.
> -----Original Message-----
> From: Beignet [mailto:beignet-bounces at lists.freedesktop.org] On Behalf Of
> yan.wang at linux.intel.com
> Sent: Thursday, March 17, 2016 17:33
> To: Guo, Yejun <yejun.guo at intel.com>
> Cc: Guo, Yejun <yejun.guo at intel.com>; beignet at 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 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
> >
> > _______________________________________________
> > Beignet mailing list
> > Beignet at lists.freedesktop.org
> > https://lists.freedesktop.org/mailman/listinfo/beignet
> >
>
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> https://lists.freedesktop.org/mailman/listinfo/beignet
More information about the Beignet
mailing list