[Beignet] [PATCH 2/2] utests: Refine the fill image0 test case to use map gtt.

[Zhigang Gong]

Now we don't fill the whole image to a const color. we
fill it according to the coords. Then we can use map gtt
to get the result and verify the result easily on cpu
side.

Signed-off-by: Zhigang Gong <zhigang.gong at linux.intel.com>
---
 kernels/test_fill_image0.cl     |  2 +-
 utests/compiler_fill_image0.cpp |  7 +++----
 utests/utest_helper.hpp         | 11 +++++++++++
 3 files changed, 15 insertions(+), 5 deletions(-)

diff --git a/kernels/test_fill_image0.cl b/kernels/test_fill_image0.cl
index ad1339f..9428092 100644
--- a/kernels/test_fill_image0.cl
+++ b/kernels/test_fill_image0.cl
@@ -2,8 +2,8 @@ __kernel void
 test_fill_image0(__write_only image2d_t dst)
 {
   int2 coord;
-  int4 color4 = {0x12, 0x34, 0x56, 0x78};
   coord.x = (int)get_global_id(0);
   coord.y = (int)get_global_id(1);
+  int4 color4 = {coord.y & 0xFF, (coord.y & 0xFF00) >> 8, coord.x & 0xFF, (coord.x & 0xFF00) >> 8};
   write_imagei(dst, coord, color4);
 }
diff --git a/utests/compiler_fill_image0.cpp b/utests/compiler_fill_image0.cpp
index cf76be3..7c8f40e 100644
--- a/utests/compiler_fill_image0.cpp
+++ b/utests/compiler_fill_image0.cpp
@@ -14,7 +14,6 @@ static void compiler_fill_image0(void)
   desc.image_height = h;
   desc.image_row_pitch = 0;
 
-
   // Setup kernel and images
   OCL_CREATE_KERNEL("test_fill_image0");
 
@@ -29,11 +28,11 @@ static void compiler_fill_image0(void)
   OCL_NDRANGE(2);
 
   // Check result
-  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER_GTT(0);
   for (uint32_t j = 0; j < h; ++j)
     for (uint32_t i = 0; i < w; i++)
-      OCL_ASSERT(((uint32_t*)buf_data[0])[j * w + i] == 0x78563412);
-  OCL_UNMAP_BUFFER(0);
+      OCL_ASSERT(((uint32_t*)buf_data[0])[j * w + i] == (i << 16 | j));
+  OCL_UNMAP_BUFFER_GTT(0);
 }
 
 MAKE_UTEST_FROM_FUNCTION(compiler_fill_image0);
diff --git a/utests/utest_helper.hpp b/utests/utest_helper.hpp
index 5258416..d882fc7 100644
--- a/utests/utest_helper.hpp
+++ b/utests/utest_helper.hpp
@@ -111,6 +111,17 @@ extern EGLSurface  eglSurface;
     } \
   } while (0)
 
+#define OCL_MAP_BUFFER_GTT(ID) \
+    OCL_CALL2(clMapBufferGTTIntel, buf_data[ID], buf[ID])
+
+#define OCL_UNMAP_BUFFER_GTT(ID) \
+  do { \
+    if (buf[ID] != NULL) { \
+      OCL_CALL (clUnmapBufferGTTIntel, buf[ID]); \
+      buf_data[ID] = NULL; \
+    } \
+  } while (0)
+
 #define OCL_NDRANGE(DIM_N) \
     OCL_CALL (clEnqueueNDRangeKernel, queue, kernel, DIM_N, NULL, globals, locals, 0, NULL, NULL)
 
-- 
1.7.11.7