Added more tests for structured and unstructured branches Debugged the unstructured...

diff --git a/kernels/compiler_if_else.cl b/kernels/compiler_if_else.cl
new file mode 100644 (file)
index 0000000..7ae8f99
--- /dev/null
+++ b/kernels/compiler_if_else.cl
@@ -0,0 +1,14 @@
+__kernel void
+compiler_if_else(__global int *src, __global int *dst)
+{
+  int id = (int)get_global_id(0);
+  dst[id] = src[id];
+  if (dst[id] >= 0) {
+    dst[id] = src[id+1];
+    src[id] = 1;
+  } else {
+    dst[id]--;
+    src[id] = 2;
+  }
+}
+

diff --git a/kernels/compiler_unstructured_branch3.cl b/kernels/compiler_unstructured_branch3.cl
new file mode 100644 (file)
index 0000000..67b4761
--- /dev/null
+++ b/kernels/compiler_unstructured_branch3.cl
@@ -0,0 +1,16 @@
+__kernel void
+compiler_unstructured_branch3(__global int *src, __global int *dst)
+{
+  int id = (int)get_global_id(0);
+  dst[id] = src[id];
+  if (dst[id] >= 2) goto label1;
+  dst[id] = 1;
+  if (src[id] < 2) goto label2;
+  dst[id]--;
+  label1:
+  dst[id] -= 2;
+  label2:
+  dst[id] += 2;
+}
+
+

diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 644ebb1..aa0dd0d 100644 (file)
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -10,9 +10,11 @@ ADD_LIBRARY(utests SHARED
             compiler_write_only.cpp
             compiler_copy_buffer.cpp
             compiler_copy_buffer_row.cpp
+            compiler_if_else.cpp
             compiler_unstructured_branch0.cpp
             compiler_unstructured_branch1.cpp
-            compiler_unstructured_branch2.cpp)
+            compiler_unstructured_branch2.cpp
+            compiler_unstructured_branch3.cpp)
 TARGET_LINK_LIBRARIES(utests cl m)
 
 ADD_EXECUTABLE(run utest_run.cpp)

diff --git a/utests/compiler_if_else.cpp b/utests/compiler_if_else.cpp
new file mode 100644 (file)
index 0000000..62a7917
--- /dev/null
+++ b/utests/compiler_if_else.cpp
@@ -0,0 +1,83 @@
+/* 
+ * 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_if_else(void)
+{
+  const size_t n = 17;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_if_else");
+  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);
+    OCL_ASSERT(((int32_t*)buf_data[0])[i] == 1);
+  }
+
+  // Second control flow
+  for (uint32_t i = 0; i < n; ++i) ((int32_t*)buf_data[0])[i] = -1;
+  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 < 16; ++i) {
+    OCL_ASSERT(((uint32_t*)buf_data[1])[i] == -2);
+    OCL_ASSERT(((uint32_t*)buf_data[0])[i] == 2);
+  }
+
+  // Third control flow
+  for (uint32_t i = 0; i < 4; ++i) ((int32_t*)buf_data[0])[i] = 2;
+  for (uint32_t i = 4; i < n; ++i) ((int32_t*)buf_data[0])[i] = -1;
+  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 < 3; ++i) {
+    OCL_ASSERT(((int32_t*)buf_data[1])[i] == 2);
+    OCL_ASSERT(((int32_t*)buf_data[0])[i] == 1);
+  }
+  OCL_ASSERT(((int32_t*)buf_data[1])[3] == -1);
+  OCL_ASSERT(((int32_t*)buf_data[0])[3] == 1);
+  for (uint32_t i = 4; i < 16; ++i) {
+    OCL_ASSERT(((int32_t*)buf_data[1])[i] == -2);
+    OCL_ASSERT(((int32_t*)buf_data[0])[i] == 2);
+  }
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_if_else);
+

diff --git a/utests/compiler_unstructured_branch3.cpp b/utests/compiler_unstructured_branch3.cpp
new file mode 100644 (file)
index 0000000..be143d2
--- /dev/null
+++ b/utests/compiler_unstructured_branch3.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_branch3(void)
+{
+  const size_t n = 16;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_unstructured_branch3");
+  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] = 0;
+  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] = 0;
+  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_branch3);
+