summaryrefslogtreecommitdiff
path: root/utests/builtin_kernel_block_motion_estimate_intel.cpp
diff options
context:
space:
mode:
authorGuo Yejun <yejun.guo@intel.com>2016-03-17 09:42:24 +0800
committerYang Rong <rong.r.yang@intel.com>2016-04-08 14:34:00 +0800
commit4f2d8c79aeb9d6c2a85e0cf7e6a85707c5c3cbc5 (patch)
tree73fb32b537346448cad1b197b815dcd255f6e598 /utests/builtin_kernel_block_motion_estimate_intel.cpp
parent14bd8855dddcf683df8138c1062bc65b05d46f94 (diff)
downloadbeignet-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.cpp18
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);