All the previous 3D test cases are only using depth 1, and not really touch the 3D read/write code path. Now fix them.
v3: fix compilation warnings. Signed-off-by: Zhigang Gong <[email protected]> --- 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..ff493e7 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(uint32_t 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(uint32_t 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(uint32_t 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(uint32_t 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 _______________________________________________ Beignet mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/beignet
