From fce9f3eb89eeee3fba2296cad936fac582b1dc9f Mon Sep 17 00:00:00 2001 From: Zhigang Gong Date: Thu, 16 May 2013 10:58:51 +0800 Subject: [PATCH] utests: Refine the fill image0 test case to use map gtt. 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 Tested-by: Simon Richter --- 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) -- 2.7.4