Pushed, thanks.
> -----Original Message----- > From: Beignet [mailto:beignet-boun...@lists.freedesktop.org] On Behalf Of > yan.w...@linux.intel.com > Sent: Thursday, March 17, 2016 17:33 > To: Guo, Yejun <yejun....@intel.com> > Cc: Guo, Yejun <yejun....@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....@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 _______________________________________________ Beignet mailing list Beignet@lists.freedesktop.org https://lists.freedesktop.org/mailman/listinfo/beignet
