diff options
author | Guo Yejun <yejun.guo@intel.com> | 2016-03-17 09:42:24 +0800 |
---|---|---|
committer | Yang Rong <rong.r.yang@intel.com> | 2016-04-08 14:34:00 +0800 |
commit | 4f2d8c79aeb9d6c2a85e0cf7e6a85707c5c3cbc5 (patch) | |
tree | 73fb32b537346448cad1b197b815dcd255f6e598 /utests/builtin_kernel_block_motion_estimate_intel.cpp | |
parent | 14bd8855dddcf683df8138c1062bc65b05d46f94 (diff) | |
download | beignet-4f2d8c79aeb9d6c2a85e0cf7e6a85707c5c3cbc5.tar.gz |
utest: do not check MV near image border
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>
Reviewed-by: Yan Wang <yan.wang@linux.intel.com>
Diffstat (limited to 'utests/builtin_kernel_block_motion_estimate_intel.cpp')
-rw-r--r-- | utests/builtin_kernel_block_motion_estimate_intel.cpp | 18 |
1 files 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 12bcb7d8..008b27c8 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); |