[Beignet] [PATCH] A bug in branch instructions

[Homer Hsing]

Hi,

This test case shows a branch bug.

Following code runs wrong:

  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;

Our code may generates wrong prediction flags in Gen assembly ...

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 63c873d..c322535 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