Added several tests for unstructured branches

10 files changed, 289 insertions, 4 deletions

diff --git a/kernels/compiler_unstructured_branch0.cl b/kernels/compiler_unstructured_branch0.cl
new file mode 100644
index 00000000..66da6e02
--- /dev/null
+++ b/kernels/compiler_unstructured_branch0.cl
@@ -0,0 +1,14 @@
+__kernel void
+compiler_unstructured_branch0(__global int *src, __global int *dst)
+{
+  int id = (int)get_global_id(0);
+  dst[id] = src[id];
+  if (dst[id] >= 0) goto label;
+
+  do {
+    dst[id] = 1;
+  label:
+    id += get_local_size(0);
+  } while (id < 32);
+}
+
diff --git a/kernels/compiler_unstructured_branch1.cl b/kernels/compiler_unstructured_branch1.cl
new file mode 100644
index 00000000..fb937e0a
--- /dev/null
+++ b/kernels/compiler_unstructured_branch1.cl
@@ -0,0 +1,14 @@
+__kernel void
+compiler_unstructured_branch1(__global int *src, __global int *dst)
+{
+  int id = (int)get_global_id(0);
+  dst[id] = src[id];
+  if (dst[id] >= 0) goto label1;
+  dst[id] = 1;
+  if (src[id] <= 2) goto label2;
+  label1:
+  dst[id] -= 2;
+  label2:
+  dst[id] += 2;
+}
+
diff --git a/kernels/compiler_unstructured_branch2.cl b/kernels/compiler_unstructured_branch2.cl
new file mode 100644
index 00000000..546f2532
--- /dev/null
+++ b/kernels/compiler_unstructured_branch2.cl
@@ -0,0 +1,18 @@
+__kernel void
+compiler_unstructured_branch2(__global int *src, __global int *dst)
+{
+  int id = (int)get_global_id(0);
+  dst[id] = src[id];
+  if (dst[id] < 0) goto label1;
+  dst[id] = 1;
+  if (dst[id] > src[id]) goto label3;
+  dst[id]++;
+  if (src[id] <= 2) goto label2;
+  label1:
+  dst[id] -= 2;
+  label2:
+  dst[id] += 2;
+  label3:
+  dst[id] *= 3;
+}
+
diff --git a/kernels/test_copy_buffer.cl b/kernels/test_copy_buffer.cl
index 927db27a..d6e23b2e 100644
--- a/kernels/test_copy_buffer.cl
+++ b/kernels/test_copy_buffer.cl
@@ -1,7 +1,7 @@
 __kernel void
 test_copy_buffer(__global float* src, __global float* dst)
 {
-    int id = (int)get_global_id(0);
-    dst[id] = src[id];
+  int id = (int)get_global_id(0);
+  dst[id] = src[id];
 }
 
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 4a54d2ab..644ebb19 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -9,7 +9,10 @@ ADD_LIBRARY(utests SHARED
             utest.cpp
             compiler_write_only.cpp
             compiler_copy_buffer.cpp
-            compiler_copy_buffer_row.cpp)
+            compiler_copy_buffer_row.cpp
+            compiler_unstructured_branch0.cpp
+            compiler_unstructured_branch1.cpp
+            compiler_unstructured_branch2.cpp)
 TARGET_LINK_LIBRARIES(utests cl m)
 
 ADD_EXECUTABLE(run utest_run.cpp)
diff --git a/utests/compiler_unstructured_branch0.cpp b/utests/compiler_unstructured_branch0.cpp
new file mode 100644
index 00000000..12ebade3
--- /dev/null
+++ b/utests/compiler_unstructured_branch0.cpp
@@ -0,0 +1,75 @@
+/* 
+ * Copyright © 2012 Intel Corporation
+ *
+ * This library is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 2 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library. If not, see <http://www.gnu.org/licenses/>.
+ *
+ * Author: Benjamin Segovia <benjamin.segovia@intel.com>
+ */
+
+#include "utest_helper.hpp"
+
+static void compiler_unstructured_branch0(void)
+{
+  const size_t n = 32;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_unstructured_branch0");
+  buf_data[0] = (uint32_t*) malloc(sizeof(uint32_t) * n);
+  for (uint32_t i = 0; i < n; ++i) ((uint32_t*)buf_data[0])[i] = 2;
+  OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(uint32_t), buf_data[0]);
+  OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL);
+  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]);
+  globals[0] = 16;
+  locals[0] = 16;
+  OCL_NDRANGE(1);
+
+  // First control flow
+  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER(1);
+  for (uint32_t i = 0; i < 16; ++i)
+    OCL_ASSERT(((int32_t*)buf_data[1])[i] == 2);
+  for (uint32_t i = 16; i < 32; ++i)
+    OCL_ASSERT(((int32_t*)buf_data[1])[i] == 1);
+
+  // Second control flow
+  for (uint32_t i = 0; i < n; ++i) ((int32_t*)buf_data[0])[i] = -2;
+  OCL_UNMAP_BUFFER(0);
+  OCL_UNMAP_BUFFER(1);
+  OCL_NDRANGE(1);
+  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER(1);
+  for (uint32_t i = 0; i < 32; ++i)
+    OCL_ASSERT(((int32_t*)buf_data[1])[i] == 1);
+
+  // Third control flow
+  for (uint32_t i = 0; i < 8; ++i) ((int32_t*)buf_data[0])[i] = 2;
+  OCL_UNMAP_BUFFER(0);
+  OCL_UNMAP_BUFFER(1);
+  OCL_NDRANGE(1);
+  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER(1);
+  for (uint32_t i = 0; i < 8; ++i)
+    OCL_ASSERT(((int32_t*)buf_data[1])[i] == 2);
+  for (uint32_t i = 8; i < 32; ++i)
+    OCL_ASSERT(((int32_t*)buf_data[1])[i] == 1);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_unstructured_branch0);
+
+
diff --git a/utests/compiler_unstructured_branch1.cpp b/utests/compiler_unstructured_branch1.cpp
new file mode 100644
index 00000000..5c3614c9
--- /dev/null
+++ b/utests/compiler_unstructured_branch1.cpp
@@ -0,0 +1,73 @@
+/* 
+ * Copyright © 2012 Intel Corporation
+ *
+ * This library is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 2 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library. If not, see <http://www.gnu.org/licenses/>.
+ *
+ * Author: Benjamin Segovia <benjamin.segovia@intel.com>
+ */
+
+#include "utest_helper.hpp"
+
+static void compiler_unstructured_branch1(void)
+{
+  const size_t n = 16;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_unstructured_branch1");
+  buf_data[0] = (uint32_t*) malloc(sizeof(uint32_t) * n);
+  for (uint32_t i = 0; i < n; ++i) ((uint32_t*)buf_data[0])[i] = 2;
+  OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(uint32_t), buf_data[0]);
+  OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL);
+  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]);
+  globals[0] = 16;
+  locals[0] = 16;
+  OCL_NDRANGE(1);
+
+  // First control flow
+  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER(1);
+  for (uint32_t i = 0; i < n; ++i)
+    OCL_ASSERT(((int32_t*)buf_data[1])[i] == 2);
+
+  // Second control flow
+  for (uint32_t i = 0; i < n; ++i) ((int32_t*)buf_data[0])[i] = -2;
+  OCL_UNMAP_BUFFER(0);
+  OCL_UNMAP_BUFFER(1);
+  OCL_NDRANGE(1);
+  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER(1);
+  for (uint32_t i = 0; i < n; ++i)
+    OCL_ASSERT(((uint32_t*)buf_data[1])[i] == 3);
+
+  // Third control flow
+  for (uint32_t i = 0; i < 8; ++i) ((int32_t*)buf_data[0])[i] = 2;
+  for (uint32_t i = 8; i < n; ++i) ((int32_t*)buf_data[0])[i] = -2;
+  OCL_UNMAP_BUFFER(0);
+  OCL_UNMAP_BUFFER(1);
+  OCL_NDRANGE(1);
+  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER(1);
+  for (uint32_t i = 0; i < 8; ++i)
+    OCL_ASSERT(((int32_t*)buf_data[1])[i] == 2);
+  for (uint32_t i = 8; i < n; ++i)
+    OCL_ASSERT(((int32_t*)buf_data[1])[i] == 3);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_unstructured_branch1);
+
diff --git a/utests/compiler_unstructured_branch2.cpp b/utests/compiler_unstructured_branch2.cpp
new file mode 100644
index 00000000..1808ef13
--- /dev/null
+++ b/utests/compiler_unstructured_branch2.cpp
@@ -0,0 +1,87 @@
+/* 
+ * Copyright © 2012 Intel Corporation
+ *
+ * This library is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 2 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library. If not, see <http://www.gnu.org/licenses/>.
+ *
+ * Author: Benjamin Segovia <benjamin.segovia@intel.com>
+ */
+
+#include "utest_helper.hpp"
+
+static void compiler_unstructured_branch2(void)
+{
+  const size_t n = 16;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_unstructured_branch2");
+  buf_data[0] = (uint32_t*) malloc(sizeof(uint32_t) * n);
+  for (uint32_t i = 0; i < n; ++i) ((uint32_t*)buf_data[0])[i] = 2;
+  OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(uint32_t), buf_data[0]);
+  OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL);
+  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]);
+  globals[0] = 16;
+  locals[0] = 16;
+  OCL_NDRANGE(1);
+
+  // First control flow
+  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER(1);
+  for (uint32_t i = 0; i < n; ++i)
+    OCL_ASSERT(((int32_t*)buf_data[1])[i] == 12);
+
+  // Second control flow
+  for (uint32_t i = 0; i < n; ++i) ((int32_t*)buf_data[0])[i] = -2;
+  OCL_UNMAP_BUFFER(0);
+  OCL_UNMAP_BUFFER(1);
+  OCL_NDRANGE(1);
+  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER(1);
+  for (uint32_t i = 0; i < n; ++i)
+    OCL_ASSERT(((int32_t*)buf_data[1])[i] == -6);
+
+  // Third control flow
+  for (uint32_t i = 0; i < 8; ++i) ((int32_t*)buf_data[0])[i] = 2;
+  for (uint32_t i = 8; i < n; ++i) ((int32_t*)buf_data[0])[i] = -2;
+  OCL_UNMAP_BUFFER(0);
+  OCL_UNMAP_BUFFER(1);
+  OCL_NDRANGE(1);
+  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER(1);
+  for (uint32_t i = 0; i < 8; ++i)
+    OCL_ASSERT(((int32_t*)buf_data[1])[i] == 12);
+  for (uint32_t i = 8; i < n; ++i)
+    OCL_ASSERT(((int32_t*)buf_data[1])[i] == -6);
+
+  // Fourth control flow
+  for (uint32_t i = 0; i < 4; ++i) ((int32_t*)buf_data[0])[i] = 1;
+  for (uint32_t i = 4; i < 8; ++i) ((int32_t*)buf_data[0])[i] = 2;
+  for (uint32_t i = 8; i < n; ++i) ((int32_t*)buf_data[0])[i] = -2;
+  OCL_UNMAP_BUFFER(0);
+  OCL_UNMAP_BUFFER(1);
+  OCL_NDRANGE(1);
+  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER(1);
+  for (uint32_t i = 0; i < 8; ++i)
+    OCL_ASSERT(((int32_t*)buf_data[1])[i] == 12);
+  for (uint32_t i = 8; i < n; ++i)
+    OCL_ASSERT(((int32_t*)buf_data[1])[i] == -6);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_unstructured_branch2);
+
diff --git a/utests/utest_assert.cpp b/utests/utest_assert.cpp
index f3b9a006..fe3a99c1 100644
--- a/utests/utest_assert.cpp
+++ b/utests/utest_assert.cpp
@@ -36,6 +36,7 @@ void onFailedAssertion(const char *msg, const char *file, const char *fn, int li
                         + std::string(file)
                         + ", function " + std::string(fn)
                         + ", line " + std::string(lineString);
+  assert(0);
   throw Exception(str);
 }
 
diff --git a/utests/utest_helper.cpp b/utests/utest_helper.cpp
index d9979ae1..e3cc27f1 100644
--- a/utests/utest_helper.cpp
+++ b/utests/utest_helper.cpp
@@ -26,7 +26,7 @@
 #include <cstdio>
 #include <cstdint>
 #include <cstring>
-#include <casserth>
+#include <cassert>
 
 #define FATAL(...) \
 do { \