[Beignet] [PATCH 5/5] test cases for image3d_t
Homer Hsing
homer.xing at intel.com
Sun May 5 17:45:52 PDT 2013
test cases for image3d_t
Signed-off-by: Homer Hsing <homer.xing at intel.com>
---
kernels/test_copy_image_3d.cl | 11 ++++++++
kernels/test_fill_image_3d.cl | 14 ++++++++++
kernels/test_fill_image_3d_2.cl | 10 +++++++
utests/CMakeLists.txt | 3 ++
utests/compiler_copy_image_3d.cpp | 55 +++++++++++++++++++++++++++++++++++++
utests/compiler_fill_image_3d.cpp | 44 +++++++++++++++++++++++++++++
utests/compiler_fill_image_3d_2.cpp | 42 ++++++++++++++++++++++++++++
7 files changed, 179 insertions(+)
create mode 100644 kernels/test_copy_image_3d.cl
create mode 100644 kernels/test_fill_image_3d.cl
create mode 100644 kernels/test_fill_image_3d_2.cl
create mode 100644 utests/compiler_copy_image_3d.cpp
create mode 100644 utests/compiler_fill_image_3d.cpp
create mode 100644 utests/compiler_fill_image_3d_2.cpp
diff --git a/kernels/test_copy_image_3d.cl b/kernels/test_copy_image_3d.cl
new file mode 100644
index 0000000..766227a
--- /dev/null
+++ b/kernels/test_copy_image_3d.cl
@@ -0,0 +1,11 @@
+__kernel void
+test_copy_image_3d(__read_only image3d_t src, __write_only image3d_t dst, sampler_t sampler)
+{
+ int4 coord;
+ int4 color;
+ coord.x = (int)get_global_id(0);
+ coord.y = (int)get_global_id(1);
+ coord.z = 0;
+ color = read_imagei(src, sampler, coord);
+ write_imagei(dst, coord, color);
+}
diff --git a/kernels/test_fill_image_3d.cl b/kernels/test_fill_image_3d.cl
new file mode 100644
index 0000000..0f0c6fd
--- /dev/null
+++ b/kernels/test_fill_image_3d.cl
@@ -0,0 +1,14 @@
+__kernel void
+test_fill_image_3d(__write_only image3d_t dst, uint color)
+{
+ int4 coord;
+ int4 color4;
+ color4.s0 = (color >> 24) & 0xFF;
+ color4.s1 = (color >> 16) & 0xFF;
+ color4.s2 = (color >> 8) & 0xFF;
+ color4.s3 = color & 0xFF;
+ coord.x = (int)get_global_id(0);
+ coord.y = (int)get_global_id(1);
+ coord.z = 0;
+ write_imagei(dst, coord, color4);
+}
diff --git a/kernels/test_fill_image_3d_2.cl b/kernels/test_fill_image_3d_2.cl
new file mode 100644
index 0000000..22b6452
--- /dev/null
+++ b/kernels/test_fill_image_3d_2.cl
@@ -0,0 +1,10 @@
+__kernel void
+test_fill_image_3d_2(__write_only image3d_t dst)
+{
+ int4 coord;
+ int4 color4 = {0x12, 0x34, 0x56, 0x78};
+ coord.x = (int)get_global_id(0);
+ coord.y = (int)get_global_id(1);
+ coord.z = 0;
+ write_imagei(dst, coord, color4);
+}
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index a02d91f..1cecec8 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -23,10 +23,13 @@ set (utests_sources
compiler_convert_uchar_sat.cpp
compiler_copy_buffer.cpp
compiler_copy_image.cpp
+ compiler_copy_image_3d.cpp
compiler_copy_buffer_row.cpp
compiler_fabs.cpp
compiler_fill_image.cpp
compiler_fill_image0.cpp
+ compiler_fill_image_3d.cpp
+ compiler_fill_image_3d_2.cpp
compiler_function_argument0.cpp
compiler_function_argument1.cpp
compiler_function_argument.cpp
diff --git a/utests/compiler_copy_image_3d.cpp b/utests/compiler_copy_image_3d.cpp
new file mode 100644
index 0000000..9816fe4
--- /dev/null
+++ b/utests/compiler_copy_image_3d.cpp
@@ -0,0 +1,55 @@
+#include "utest_helper.hpp"
+
+static void compiler_copy_image_3d(void)
+{
+ const size_t w = 512;
+ const size_t h = 512;
+ const size_t depth = 1;
+ cl_image_format format;
+ cl_image_desc desc;
+ cl_sampler sampler;
+
+ // Setup kernel and images
+ OCL_CREATE_KERNEL("test_copy_image_3d");
+ buf_data[0] = (uint32_t*) malloc(sizeof(uint32_t) * w * h * depth);
+ for (uint32_t k = 0; k < depth; k++)
+ for (uint32_t j = 0; j < h; j++)
+ for (uint32_t i = 0; i < w; i++)
+ ((uint32_t*)buf_data[0])[k*w*h + j*w + i] = k*w*h + j*w + i;
+
+ format.image_channel_order = CL_RGBA;
+ format.image_channel_data_type = CL_UNSIGNED_INT8;
+ desc.image_type = CL_MEM_OBJECT_IMAGE3D;
+ desc.image_width = w;
+ desc.image_height = h;
+ desc.image_depth = depth;
+ desc.image_row_pitch = 0;
+ desc.image_slice_pitch = 0;
+ OCL_CREATE_IMAGE(buf[0], CL_MEM_COPY_HOST_PTR, &format, &desc, buf_data[0]);
+ OCL_CREATE_IMAGE(buf[1], 0, &format, &desc, NULL);
+ OCL_CREATE_SAMPLER(sampler, CL_ADDRESS_REPEAT, CL_FILTER_NEAREST);
+ free(buf_data[0]);
+ buf_data[0] = 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(sampler), &sampler);
+ globals[0] = w;
+ globals[1] = h;
+ locals[0] = 16;
+ locals[1] = 16;
+ OCL_NDRANGE(2);
+
+ // Check result
+ OCL_MAP_BUFFER(0);
+ OCL_MAP_BUFFER(1);
+ for (uint32_t k = 0; k < depth; k++)
+ for (uint32_t j = 0; j < h; ++j)
+ for (uint32_t i = 0; i < w; i++)
+ OCL_ASSERT(((uint32_t*)buf_data[0])[k*w*h + j*w + i] == ((uint32_t*)buf_data[1])[k*w*h + j*w + i]);
+ OCL_UNMAP_BUFFER(0);
+ OCL_UNMAP_BUFFER(1);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_copy_image_3d);
diff --git a/utests/compiler_fill_image_3d.cpp b/utests/compiler_fill_image_3d.cpp
new file mode 100644
index 0000000..5d98b0b
--- /dev/null
+++ b/utests/compiler_fill_image_3d.cpp
@@ -0,0 +1,44 @@
+#include "utest_helper.hpp"
+
+static void compiler_fill_image_3d(void)
+{
+ const size_t w = 512;
+ const size_t h = 512;
+ const size_t depth = 1;
+ uint32_t color = 0x12345678;
+ cl_image_format format;
+ cl_image_desc desc;
+
+ format.image_channel_order = CL_RGBA;
+ format.image_channel_data_type = CL_UNSIGNED_INT8;
+ desc.image_type = CL_MEM_OBJECT_IMAGE3D;
+ desc.image_width = w;
+ desc.image_height = h;
+ desc.image_depth = depth;
+ desc.image_row_pitch = 0;
+ desc.image_slice_pitch = 0;
+
+ // Setup kernel and images
+ OCL_CREATE_KERNEL("test_fill_image_3d");
+
+ OCL_CREATE_IMAGE(buf[0], 0, &format, &desc, NULL);
+
+ // Run the kernel
+ OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+ OCL_SET_ARG(1, sizeof(color), &color);
+ globals[0] = w;
+ globals[1] = h;
+ locals[0] = 16;
+ locals[1] = 16;
+ OCL_NDRANGE(2);
+
+ // Check result
+ OCL_MAP_BUFFER(0);
+ for (uint32_t k = 0; k < depth; k++)
+ for (uint32_t j = 0; j < h; ++j)
+ for (uint32_t i = 0; i < w; i++)
+ OCL_ASSERT(((uint32_t*)buf_data[0])[k*w*h + j*w + i] == 0x78563412);
+ OCL_UNMAP_BUFFER(0);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_fill_image_3d);
diff --git a/utests/compiler_fill_image_3d_2.cpp b/utests/compiler_fill_image_3d_2.cpp
new file mode 100644
index 0000000..3b4edb9
--- /dev/null
+++ b/utests/compiler_fill_image_3d_2.cpp
@@ -0,0 +1,42 @@
+#include "utest_helper.hpp"
+
+static void compiler_fill_image_3d_2(void)
+{
+ const size_t w = 512;
+ const size_t h = 512;
+ const size_t depth = 1;
+ cl_image_format format;
+ cl_image_desc desc;
+
+ format.image_channel_order = CL_RGBA;
+ format.image_channel_data_type = CL_UNSIGNED_INT8;
+ desc.image_type = CL_MEM_OBJECT_IMAGE3D;
+ desc.image_width = w;
+ desc.image_height = h;
+ desc.image_depth = depth;
+ desc.image_row_pitch = 0;
+ desc.image_slice_pitch = 0;
+
+ // Setup kernel and images
+ OCL_CREATE_KERNEL("test_fill_image_3d_2");
+
+ OCL_CREATE_IMAGE(buf[0], 0, &format, &desc, NULL);
+
+ // Run the kernel
+ OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+ globals[0] = w;
+ globals[1] = h;
+ locals[0] = 16;
+ locals[1] = 16;
+ OCL_NDRANGE(2);
+
+ // Check result
+ OCL_MAP_BUFFER(0);
+ for (uint32_t k = 0; k < depth; k++)
+ for (uint32_t j = 0; j < h; ++j)
+ for (uint32_t i = 0; i < w; i++)
+ OCL_ASSERT(((uint32_t*)buf_data[0])[k*w*h + j*w + i] == 0x78563412);
+ OCL_UNMAP_BUFFER(0);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_fill_image_3d_2);
--
1.8.1.2
More information about the Beignet
mailing list