[Beignet] [PATCH] Add the test cases for 1D Image Array
junyan.he at inbox.com
junyan.he at inbox.com
Fri Jun 20 03:07:40 PDT 2014
From: Junyan He <junyan.he at linux.intel.com>
Signed-off-by: Junyan He <junyan.he at linux.intel.com>
---
kernels/compare_image_2d_and_1d_array.cl | 13 ++++++
kernels/test_get_image_info_array.cl | 25 ++++++++++
utests/CMakeLists.txt | 2 +
utests/compare_image_2d_and_1d_array.cpp | 79 ++++++++++++++++++++++++++++++++
utests/compiler_get_image_info_array.cpp | 64 ++++++++++++++++++++++++++
5 files changed, 183 insertions(+)
create mode 100644 kernels/compare_image_2d_and_1d_array.cl
create mode 100644 kernels/test_get_image_info_array.cl
create mode 100644 utests/compare_image_2d_and_1d_array.cpp
create mode 100644 utests/compiler_get_image_info_array.cpp
diff --git a/kernels/compare_image_2d_and_1d_array.cl b/kernels/compare_image_2d_and_1d_array.cl
new file mode 100644
index 0000000..5cd077e
--- /dev/null
+++ b/kernels/compare_image_2d_and_1d_array.cl
@@ -0,0 +1,13 @@
+__kernel void
+compare_image_2d_and_1d_array(image2d_t a1, image1d_array_t a2, sampler_t sampler)
+{
+ float2 coord;
+ int4 color1;
+ int4 color2;
+ coord.x = (float)get_global_id(0) + 0.3;
+ coord.y = (float)get_global_id(1) + 0.3;
+ color1 = read_imagei(a1, sampler, coord);
+ color2 = read_imagei(a2, sampler, coord);
+// printf("########## x y is (%f, %f), color1 is (%d %d %d %d), color2 is (%d %d %d %d)\n",
+// coord.x, coord.y, color1.x, color1.y, color1.z, color1.w, color2.x, color2.y, color2.z, color2.w);
+}
diff --git a/kernels/test_get_image_info_array.cl b/kernels/test_get_image_info_array.cl
new file mode 100644
index 0000000..333da77
--- /dev/null
+++ b/kernels/test_get_image_info_array.cl
@@ -0,0 +1,25 @@
+__kernel void
+test_get_image_info_array(__write_only image1d_array_t a1, __write_only image2d_array_t a2, __global int *result)
+{
+ int w, h, array_sz;
+
+ w = get_image_width(a1);
+ array_sz = (int)get_image_array_size(a1);
+ int channel_data_type = get_image_channel_data_type(a1);
+ int channel_order = get_image_channel_order(a1);
+ result[0] = w;
+ result[1] = array_sz;
+ result[2] = channel_data_type;
+ result[3] = channel_order;
+
+ w = get_image_width(a2);
+ h = get_image_height(a2);
+ array_sz = (int)get_image_array_size(a2);
+ channel_data_type = get_image_channel_data_type(a2);
+ channel_order = get_image_channel_order(a2);
+ result[4] = w;
+ result[5] = h;
+ result[6] = array_sz;
+ result[7] = channel_data_type;
+ result[8] = channel_order;
+}
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index f0e62e2..641a73b 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -122,6 +122,7 @@ set (utests_sources
compiler_volatile.cpp
compiler_copy_image1.cpp
compiler_get_image_info.cpp
+ compiler_get_image_info_array.cpp
compiler_vect_compare.cpp
compiler_vector_load_store.cpp
compiler_vector_inc.cpp
@@ -182,6 +183,7 @@ set (utests_sources
enqueue_fill_buf.cpp
enqueue_built_in_kernels.cpp
image_1D_buffer.cpp
+ compare_image_2d_and_1d_array.cpp
utest_assert.cpp
utest.cpp
utest_file_map.cpp
diff --git a/utests/compare_image_2d_and_1d_array.cpp b/utests/compare_image_2d_and_1d_array.cpp
new file mode 100644
index 0000000..f2c828e
--- /dev/null
+++ b/utests/compare_image_2d_and_1d_array.cpp
@@ -0,0 +1,79 @@
+#include <string.h>
+#include "utest_helper.hpp"
+
+static void compare_image_2d_and_1d_array(void)
+{
+ const int w = 64;
+ const int h = 32;
+ cl_image_format format;
+ cl_image_desc desc;
+ cl_sampler sampler;
+
+ // Create the 1D array buffer.
+ memset(&desc, 0x0, sizeof(cl_image_desc));
+ memset(&format, 0x0, sizeof(cl_image_format));
+
+ uint32_t* image_data1 = (uint32_t *)malloc(w * h * sizeof(uint32_t));
+ uint32_t* image_data2 = (uint32_t *)malloc(w * h * sizeof(uint32_t));
+ for (int j = 0; j < h; j++) {
+ for (int i = 0; i < w; i++) {
+ char a = 0;
+ if (j % 2 == 0)
+ a = (j + 3) & 0x3f;
+
+ image_data2[w * j + i] = image_data1[w * j + i] = a << 24 | a << 16 | a << 8 | a;
+ }
+ }
+
+ format.image_channel_order = CL_RGBA;
+ format.image_channel_data_type = CL_UNSIGNED_INT8;
+ desc.image_type = CL_MEM_OBJECT_IMAGE2D;
+ desc.image_width = w;
+ desc.image_height = h;
+ desc.image_row_pitch = w * sizeof(uint32_t);
+ OCL_CREATE_IMAGE(buf[0], CL_MEM_COPY_HOST_PTR, &format, &desc, image_data1);
+
+ // Create the 2D array buffer.
+ 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_IMAGE1D_ARRAY;
+ desc.image_width = w;
+ desc.image_array_size = h;
+ desc.image_row_pitch = w * sizeof(uint32_t);
+ OCL_CREATE_IMAGE(buf[1], CL_MEM_COPY_HOST_PTR, &format, &desc, image_data2);
+
+ OCL_CREATE_SAMPLER(sampler, CL_ADDRESS_REPEAT, CL_FILTER_LINEAR);
+
+ // Setup kernel and images
+ OCL_CREATE_KERNEL("compare_image_2d_and_1d_array");
+
+ // Run the kernel
+ OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+ OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+ OCL_SET_ARG(2, sizeof(cl_sampler), &sampler);
+ globals[0] = 32;
+ globals[1] = 16;
+ locals[0] = 32;
+ locals[1] = 16;
+ OCL_NDRANGE(2);
+
+ OCL_MAP_BUFFER_GTT(0);
+ OCL_MAP_BUFFER_GTT(1);
+ for (int j = 0; j < h; ++j) {
+ for (int i = 0; i < w; i++) {
+ // Because the array index will not join the sample caculation, the result should
+ // be different between the 2D and 1D_array.
+ if (j % 2 == 0)
+ OCL_ASSERT(((uint32_t*)buf_data[0])[j * w + i] == ((uint32_t*)buf_data[1])[j * w + i]);
+ }
+ }
+ OCL_UNMAP_BUFFER_GTT(0);
+ OCL_UNMAP_BUFFER_GTT(1);
+
+ OCL_CALL(clReleaseSampler, sampler);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compare_image_2d_and_1d_array);
diff --git a/utests/compiler_get_image_info_array.cpp b/utests/compiler_get_image_info_array.cpp
new file mode 100644
index 0000000..970877d
--- /dev/null
+++ b/utests/compiler_get_image_info_array.cpp
@@ -0,0 +1,64 @@
+#include <string.h>
+#include "utest_helper.hpp"
+
+static void compiler_get_image_info_array(void)
+{
+ const int w = 256;
+ const int h = 512;
+ const int array_size1 = 10;
+ const int array_size2 = 3;
+ cl_image_format format;
+ cl_image_desc desc;
+
+ // Create the 1D array buffer.
+ 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_IMAGE1D_ARRAY;
+ desc.image_width = w;
+ desc.image_array_size = array_size1;
+ OCL_CREATE_IMAGE(buf[0], 0, &format, &desc, NULL);
+
+ // Create the 2D array buffer.
+ 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_array_size = array_size2;
+ OCL_CREATE_IMAGE(buf[1], 0, &format, &desc, NULL);
+
+ // Setup kernel and images
+ OCL_CREATE_KERNEL("test_get_image_info_array");
+
+ OCL_CREATE_BUFFER(buf[2], 0, 32 * sizeof(int), NULL);
+
+ // Run the kernel
+ OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+ OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+ OCL_SET_ARG(2, sizeof(cl_mem), &buf[2]);
+ globals[0] = 32;
+ locals[0] = 16;
+ OCL_NDRANGE(1);
+
+ // Check result
+ OCL_MAP_BUFFER(2);
+ OCL_ASSERT(((int*)buf_data[2])[0] == w);
+ OCL_ASSERT(((int*)buf_data[2])[1] == array_size1);
+ OCL_ASSERT(((int*)buf_data[2])[2] == CL_UNSIGNED_INT8);
+ OCL_ASSERT(((int*)buf_data[2])[3] == CL_RGBA);
+
+ OCL_ASSERT(((int*)buf_data[2])[4] == w);
+ OCL_ASSERT(((int*)buf_data[2])[5] == h);
+ OCL_ASSERT(((int*)buf_data[2])[6] == array_size2);
+ OCL_ASSERT(((int*)buf_data[2])[7] == CL_UNSIGNED_INT8);
+ OCL_ASSERT(((int*)buf_data[2])[8] == CL_RGBA);
+ OCL_UNMAP_BUFFER(2);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_get_image_info_array);
--
1.8.3.2
More information about the Beignet
mailing list