[Beignet] [PATCH] test case for DW multiplication

[Zhigang Gong]

Pushed, thanks.
On Tue, Jun 04, 2013 at 03:52:13PM +0800, Homer Hsing wrote:
> This case tests whether a predication bug of DW multiplication
> had been fixed.
> 
> Signed-off-by: Homer Hsing <homer.xing at intel.com>
> ---
>  kernels/compiler_displacement_map_element.cl | 12 ++++++
>  utests/CMakeLists.txt                        |  1 +
>  utests/compiler_displacement_map_element.cpp | 64 ++++++++++++++++++++++++++++
>  3 files changed, 77 insertions(+)
>  create mode 100644 kernels/compiler_displacement_map_element.cl
>  create mode 100644 utests/compiler_displacement_map_element.cpp
> 
> diff --git a/kernels/compiler_displacement_map_element.cl b/kernels/compiler_displacement_map_element.cl
> new file mode 100644
> index 0000000..b94d9a0
> --- /dev/null
> +++ b/kernels/compiler_displacement_map_element.cl
> @@ -0,0 +1,12 @@
> +kernel void compiler_displacement_map_element(const global uint *in, const global uint *offset, int w, int h, global uint *out) {
> +    const int cx = get_global_id(0);
> +    const int cy = get_global_id(1);
> +    uint c = offset[cy * w + cx];
> +    int x_pos = cx + c;
> +    int y_pos = cy + c;
> +    if(0 <= x_pos && x_pos < w && 0 <= y_pos && y_pos < h)
> +        out[cy * w + cx] = in[y_pos * w + x_pos];
> +    else
> +        out[cy * w + cx] = 0;
> +}
> +
> diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
> index 322b727..618bd5d 100644
> --- a/utests/CMakeLists.txt
> +++ b/utests/CMakeLists.txt
> @@ -5,6 +5,7 @@ link_directories (${LLVM_LIBRARY_DIR})
>  set (utests_sources
>    cl_create_kernel.cpp
>    utest_error.c
> +  compiler_displacement_map_element.cpp
>    compiler_shader_toy.cpp
>    compiler_mandelbrot.cpp
>    compiler_mandelbrot_alternate.cpp
> diff --git a/utests/compiler_displacement_map_element.cpp b/utests/compiler_displacement_map_element.cpp
> new file mode 100644
> index 0000000..98041ec
> --- /dev/null
> +++ b/utests/compiler_displacement_map_element.cpp
> @@ -0,0 +1,64 @@
> +#include "utest_helper.hpp"
> +
> +typedef unsigned int uint;
> +constexpr int W = 16, H = 16;
> +constexpr int SIZE = W * H;
> +uint in_1[SIZE];
> +uint disp_map[SIZE];
> +uint out_1[SIZE];
> +
> +uint cpu(const int cx, const int cy, const uint *in, const uint *disp_map, int w, int h) {
> +  uint c = disp_map[cy * w + cx];
> +  int x_pos = cx + c;
> +  int y_pos = cy + c;
> +  if(0 <= x_pos && x_pos < w && 0 <= y_pos && y_pos < h)
> +    return in[y_pos * w + x_pos];
> +  else
> +    return 0;
> +}
> +
> +void test() {
> +  OCL_MAP_BUFFER(2);
> +  for(int y=0; y<H; y++)
> +    for(int x=0; x<W; x++) {
> +      uint out = ((uint*)buf_data[2]) [y * W + x];
> +      uint wish = cpu(x, y, in_1, disp_map, W, H);
> +      if(out != wish)
> +        printf("XXX %d %d %x %x\n", x, y, out, wish);
> +      OCL_ASSERT(out == wish);
> +    }
> +  OCL_UNMAP_BUFFER(2);
> +}
> +
> +void displacement_map_element(void) {
> +  int i, pass;
> +
> +  OCL_CREATE_KERNEL("compiler_displacement_map_element");
> +  OCL_CREATE_BUFFER(buf[0], 0, SIZE * sizeof(uint), NULL);
> +  OCL_CREATE_BUFFER(buf[1], 0, SIZE * sizeof(uint), NULL);
> +  OCL_CREATE_BUFFER(buf[2], 0, SIZE * sizeof(uint), NULL);
> +  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
> +  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
> +  OCL_SET_ARG(2, sizeof(W), &W);
> +  OCL_SET_ARG(3, sizeof(H), &H);
> +  OCL_SET_ARG(4, sizeof(cl_mem), &buf[2]);
> +  globals[0] = W;
> +  globals[1] = H;
> +  locals[0] = 16;
> +  locals[1] = 16;
> +
> +  for (pass = 0; pass < 8; pass ++) {
> +    OCL_MAP_BUFFER(0);
> +    OCL_MAP_BUFFER(1);
> +    for (i = 0; i < SIZE; i ++) {
> +      in_1[i] = ((uint*)buf_data[0])[i] = ((rand() & 0xFFFF) << 16) | (rand() & 0xFFFF);
> +      disp_map[i] = ((uint*)buf_data[1])[i] = rand() & 3;
> +    }
> +    OCL_UNMAP_BUFFER(0);
> +    OCL_UNMAP_BUFFER(1);
> +    OCL_NDRANGE(2);
> +    test();
> +  }
> +}
> +
> +MAKE_UTEST_FROM_FUNCTION(displacement_map_element);
> -- 
> 1.8.1.2
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet