[Beignet] [PATCH V2 7/7] Utests: refine the previous fake 3D test cases.
Zhigang Gong
zhigang.gong at linux.intel.com
Thu Sep 12 23:01:48 PDT 2013
All the previous 3D test cases are only using depth 1, and not
really touch the 3D read/write code path. Now fix them.
Signed-off-by: Zhigang Gong <zhigang.gong at linux.intel.com>
---
kernels/test_copy_image_3d.cl | 27 +++++++++++++++++++++-----
kernels/test_fill_image_3d.cl | 2 +-
kernels/test_fill_image_3d_2.cl | 2 +-
utests/compiler_copy_image_3d.cpp | 36 ++++++++++++++++++++++-------------
utests/compiler_fill_image_3d.cpp | 6 ++++--
utests/compiler_fill_image_3d_2.cpp | 10 ++++++----
6 files changed, 57 insertions(+), 26 deletions(-)
diff --git a/kernels/test_copy_image_3d.cl b/kernels/test_copy_image_3d.cl
index 766227a..103fb69 100644
--- a/kernels/test_copy_image_3d.cl
+++ b/kernels/test_copy_image_3d.cl
@@ -1,11 +1,28 @@
__kernel void
-test_copy_image_3d(__read_only image3d_t src, __write_only image3d_t dst, sampler_t sampler)
+test_copy_image_3d(__read_only image3d_t src,
+ __write_only image3d_t dst,
+ sampler_t sampler,
+ __write_only image2d_t buf0,
+ __write_only image2d_t buf1,
+ __write_only image2d_t buf2,
+ __write_only image2d_t buf3)
{
int4 coord;
- int4 color;
+ int2 coord2;
+ float4 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);
+ coord.z = (int)get_global_id(2);
+ coord2.x = coord.x;
+ coord2.y = coord.y;
+ color = read_imagef(src, sampler, coord);
+ write_imagef(dst, coord, color);
+ if (coord.z == 0)
+ write_imagef(buf0, coord2, color);
+ else if (coord.z == 1)
+ write_imagef(buf1, coord2, color);
+ else if (coord.z == 2)
+ write_imagef(buf2, coord2, color);
+ else if (coord.z == 3)
+ write_imagef(buf3, coord2, color);
}
diff --git a/kernels/test_fill_image_3d.cl b/kernels/test_fill_image_3d.cl
index 0f0c6fd..4988f69 100644
--- a/kernels/test_fill_image_3d.cl
+++ b/kernels/test_fill_image_3d.cl
@@ -9,6 +9,6 @@ test_fill_image_3d(__write_only image3d_t dst, uint color)
color4.s3 = color & 0xFF;
coord.x = (int)get_global_id(0);
coord.y = (int)get_global_id(1);
- coord.z = 0;
+ coord.z = (int)get_global_id(2);
write_imagei(dst, coord, color4);
}
diff --git a/kernels/test_fill_image_3d_2.cl b/kernels/test_fill_image_3d_2.cl
index 22b6452..1f9eaa1 100644
--- a/kernels/test_fill_image_3d_2.cl
+++ b/kernels/test_fill_image_3d_2.cl
@@ -5,6 +5,6 @@ test_fill_image_3d_2(__write_only image3d_t dst)
int4 color4 = {0x12, 0x34, 0x56, 0x78};
coord.x = (int)get_global_id(0);
coord.y = (int)get_global_id(1);
- coord.z = 0;
+ coord.z = (int)get_global_id(2);
write_imagei(dst, coord, color4);
}
diff --git a/utests/compiler_copy_image_3d.cpp b/utests/compiler_copy_image_3d.cpp
index 5290090..e7567dc 100644
--- a/utests/compiler_copy_image_3d.cpp
+++ b/utests/compiler_copy_image_3d.cpp
@@ -1,10 +1,11 @@
#include "utest_helper.hpp"
+#include "string.h"
static void compiler_copy_image_3d(void)
{
const size_t w = 512;
const size_t h = 512;
- const size_t depth = 1;
+ const size_t depth = 4;
cl_image_format format;
cl_sampler sampler;
@@ -14,12 +15,14 @@ static void compiler_copy_image_3d(void)
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;
+ ((float*)buf_data[0])[k*w*h + j*w + i] = (k << 10) + (j << 10) + i;
format.image_channel_order = CL_RGBA;
- format.image_channel_data_type = CL_UNSIGNED_INT8;
- OCL_CREATE_IMAGE3D(buf[0], CL_MEM_COPY_HOST_PTR, &format, w, h, depth, 0, 0, buf_data[0]);
+ format.image_channel_data_type = CL_UNORM_INT8;
+ OCL_CREATE_IMAGE3D(buf[0], CL_MEM_COPY_HOST_PTR, &format, w, h, depth, w*4, w*h*4, buf_data[0]);
OCL_CREATE_IMAGE3D(buf[1], 0, &format, w, h, depth, 0, 0, NULL);
+ for(int i = 0; i < depth; i++)
+ OCL_CREATE_IMAGE2D(buf[2 + i], 0, &format, w, h, 0, NULL);
OCL_CREATE_SAMPLER(sampler, CL_ADDRESS_REPEAT, CL_FILTER_NEAREST);
free(buf_data[0]);
buf_data[0] = NULL;
@@ -28,21 +31,28 @@ static void compiler_copy_image_3d(void)
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);
+ for(int i = 0; i < depth; i++)
+ OCL_SET_ARG(3 + i, sizeof(cl_mem), &buf[2 + i]);
globals[0] = w;
globals[1] = h;
- locals[0] = 16;
- locals[1] = 16;
- OCL_NDRANGE(2);
+ globals[2] = depth;
+ locals[0] = 64;
+ locals[1] = 1;
+ locals[2] = 1;
+ OCL_NDRANGE(3);
// Check result
- OCL_MAP_BUFFER(0);
- OCL_MAP_BUFFER(1);
+ for(int i = 0; i < depth + 2; i++)
+ OCL_MAP_BUFFER_GTT(i);
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);
+ for (uint32_t i = 0; i < w; i++) {
+ OCL_ASSERT(((float*)buf_data[0])[k*w*((h+1)&-2LL) + j*w + i] == ((float*)buf_data[1])[k*w*((h+1)&-2LL) + j*w + i]);
+ OCL_ASSERT(((float*)buf_data[0])[k*w*((h+1)&-2LL) + j*w + i] == ((float*)buf_data[k + 2])[j * w + i]);
+ }
+
+ for(int i = 0; i < depth + 2; i++)
+ OCL_UNMAP_BUFFER_GTT(i);
}
MAKE_UTEST_FROM_FUNCTION(compiler_copy_image_3d);
diff --git a/utests/compiler_fill_image_3d.cpp b/utests/compiler_fill_image_3d.cpp
index 4b3d4e3..6a679fb 100644
--- a/utests/compiler_fill_image_3d.cpp
+++ b/utests/compiler_fill_image_3d.cpp
@@ -4,7 +4,7 @@ static void compiler_fill_image_3d(void)
{
const size_t w = 512;
const size_t h = 512;
- const size_t depth = 1;
+ const size_t depth = 5;
uint32_t color = 0x12345678;
cl_image_format format;
@@ -21,9 +21,11 @@ static void compiler_fill_image_3d(void)
OCL_SET_ARG(1, sizeof(color), &color);
globals[0] = w;
globals[1] = h;
+ globals[2] = depth;
locals[0] = 16;
locals[1] = 16;
- OCL_NDRANGE(2);
+ locals[2] = 1;
+ OCL_NDRANGE(3);
// Check result
OCL_MAP_BUFFER(0);
diff --git a/utests/compiler_fill_image_3d_2.cpp b/utests/compiler_fill_image_3d_2.cpp
index 8ecc3e3..f5ff792 100644
--- a/utests/compiler_fill_image_3d_2.cpp
+++ b/utests/compiler_fill_image_3d_2.cpp
@@ -4,7 +4,7 @@ static void compiler_fill_image_3d_2(void)
{
const size_t w = 512;
const size_t h = 512;
- const size_t depth = 1;
+ const size_t depth = 5;
cl_image_format format;
format.image_channel_order = CL_RGBA;
@@ -19,17 +19,19 @@ static void compiler_fill_image_3d_2(void)
OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
globals[0] = w;
globals[1] = h;
+ globals[2] = depth;
locals[0] = 16;
locals[1] = 16;
- OCL_NDRANGE(2);
+ locals[2] = 1;
+ OCL_NDRANGE(3);
// Check result
- OCL_MAP_BUFFER(0);
+ OCL_MAP_BUFFER_GTT(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);
+ OCL_UNMAP_BUFFER_GTT(0);
}
MAKE_UTEST_FROM_FUNCTION(compiler_fill_image_3d_2);
--
1.7.9.5
More information about the Beignet
mailing list