From: Junyan He <junyan.he at linux.intel.com>
Signed-off-by: Junyan He <junyan.he at linux.intel.com>
---
kernels/test_fill_image_2d_array.cl | 13 +++++
utests/CMakeLists.txt | 1 +
utests/compiler_fill_image_2d_array.cpp | 84 +++++++++++++++++++++++++++++++++
3 files changed, 98 insertions(+)
create mode 100644 kernels/test_fill_image_2d_array.cl
create mode 100644 utests/compiler_fill_image_2d_array.cpp
diff --git a/kernels/test_fill_image_2d_array.cl b/kernels/test_fill_image_2d_array.cl
new file mode 100644
index 0000000..e756010
--- /dev/null
+++ b/kernels/test_fill_image_2d_array.cl
@@ -0,0 +1,13 @@
+__kernel void
+test_fill_image_2d_array(__write_only image2d_array_t dst)
+{
+ int coordx;
+ int coordy;
+ int coordz;
+ coordx = (int)get_global_id(0);
+ coordy = (int)get_global_id(1);
+ coordz = (int)get_global_id(2);
+ uint4 color4 = {0, 1, 2 ,3};
+ if (coordz < 7)
+ write_imageui(dst, (int3)(coordx, coordy, coordz), color4);
+}
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 57d5909..cf3addc 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -188,6 +188,7 @@ set (utests_sources
image_1D_buffer.cpp
compare_image_2d_and_1d_array.cpp
compiler_fill_image_1d_array.cpp
+ compiler_fill_image_2d_array.cpp
compiler_constant_expr.cpp
vload_bench.cpp
utest_assert.cpp
diff --git a/utests/compiler_fill_image_2d_array.cpp b/utests/compiler_fill_image_2d_array.cpp
new file mode 100644
index 0000000..649b416
--- /dev/null
+++ b/utests/compiler_fill_image_2d_array.cpp
@@ -0,0 +1,84 @@
+#include <string.h>
+#include "utest_helper.hpp"
+
+static void compiler_fill_image_2d_array(void)
+{
+ const size_t w = 64;
+ const size_t h = 16;
+ const size_t array = 8;
+ cl_image_format format;
+ cl_image_desc desc;
+ size_t origin[3] = { };
+ size_t region[3];
+ uint32_t* dst;
+
+ memset(&desc, 0x0, sizeof(cl_image_desc));
+ memset(&format, 0x0, sizeof(cl_image_format));
+
+ format.image_channel_order = CL_RGBA;
+ format.image_channel_data_type = CL_UNSIGNED_INT8;
+ desc.image_type = CL_MEM_OBJECT_IMAGE2D_ARRAY;
+ desc.image_width = w;
+ desc.image_height = h;
+ desc.image_row_pitch = 0;//w * sizeof(uint32_t);
+ desc.image_array_size = array;
+
+ // Setup kernel and images
+ OCL_CREATE_KERNEL("test_fill_image_2d_array");
+
+ OCL_CREATE_IMAGE(buf[0], 0, &format, &desc, NULL);
+
+ OCL_MAP_BUFFER_GTT(0);
+ memset(buf_data[0], 0, sizeof(uint32_t) * w * h * array);
+ OCL_UNMAP_BUFFER_GTT(0);
+
+ // Run the kernel
+ OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+ globals[0] = w/2;
+ locals[0] = 16;
+ globals[1] = h;
+ locals[1] = 8;
+ globals[2] = array;
+ locals[2] = 8;
+ OCL_NDRANGE(3);
+
+ // Check result
+ region[0] = w;
+ region[1] = h;
+ region[2] = array;
+ dst = (uint32_t*)malloc(w*h*array*sizeof(uint32_t));
+ OCL_READ_IMAGE(buf[0], origin, region, dst);
+
+#if 0
+ printf("------ The image result is: -------\n");
+ for (uint32_t k = 0; k < array; k++) {
+ for (uint32_t j = 0; j < h; j++) {
+ for (uint32_t i = 0; i < w; i++) {
+ printf(" %2x", dst[k*h*w + j*w + i]);
+ }
+ printf("\n");
+ }
+ printf("\n");
+ }
+#endif
+
+ for (uint32_t k = 0; k < array - 1; k++) {
+ for (uint32_t j = 0; j < h; j++) {
+ for (uint32_t i = 0; i < w/2; i++) {
+ OCL_ASSERT(dst[k*w*h + j*w + i] == 0x03020100);
+ }
+ for (uint32_t i = w/2; i < w; i++) {
+ OCL_ASSERT(dst[k*w*h + j*w + i] == 0);
+ }
+ }
+ }
+
+ for (uint32_t j = 0; j < h; j++) {
+ for (uint32_t i = 0; i < w; i++) {
+ OCL_ASSERT(dst[(array - 1)*w*h + j*w + i] == 0x0);
+ }
+ }
+ free(dst);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_fill_image_2d_array);
--
1.9.1