Added tests for return instruction lowering

11 files changed, 253 insertions, 5 deletions

diff --git a/kernels/compiler_array.cl b/kernels/compiler_array.cl
new file mode 100644
index 00000000..08ec6db4
--- /dev/null
+++ b/kernels/compiler_array.cl
@@ -0,0 +1,12 @@
+__kernel void
+compiler_array(__global int *src, __global int *dst, int x)
+{
+  if (x > 10) {
+    int array[x*256];
+    array[get_local_id(0)] = get_global_id(0);
+    dst[get_global_id(0)] = array[get_local_id(1)];
+  } else
+    dst[get_global_id(0)] = src[get_local_id(1)];
+}
+
+
diff --git a/kernels/compiler_lower_return0.cl b/kernels/compiler_lower_return0.cl
new file mode 100644
index 00000000..fd9846ef
--- /dev/null
+++ b/kernels/compiler_lower_return0.cl
@@ -0,0 +1,8 @@
+__kernel void
+compiler_lower_return0(__global int *src, __global int *dst) {
+  const int id = get_global_id(0);
+  dst[id] = id;
+  if (src[id] > 0) return;
+  dst[id] = src[id];
+}
+
diff --git a/kernels/compiler_lower_return1.cl b/kernels/compiler_lower_return1.cl
new file mode 100644
index 00000000..bcb6b7f0
--- /dev/null
+++ b/kernels/compiler_lower_return1.cl
@@ -0,0 +1,8 @@
+__kernel void
+compiler_lower_return1(__global int *src, __global int *dst) {
+  const int id = get_global_id(0);
+  dst[id] = id;
+  if (id < 11 && (src[id] > 0 || src[id+16] < 2)) return;
+  dst[id] = src[id];
+}
+
diff --git a/kernels/compiler_lower_return2.cl b/kernels/compiler_lower_return2.cl
new file mode 100644
index 00000000..9fa8ad60
--- /dev/null
+++ b/kernels/compiler_lower_return2.cl
@@ -0,0 +1,11 @@
+__kernel void
+compiler_lower_return2(__global int *src, __global int *dst) {
+  const int id = get_global_id(0);
+  dst[id] = id;
+  while (dst[id] > src[id]) {
+    if (dst[id] > 10) return;
+    dst[id]--;
+  }
+  dst[id] += 2;
+}
+
diff --git a/kernels/test_write_only.cl b/kernels/test_write_only.cl
index 738749a9..29fe6e34 100644
--- a/kernels/test_write_only.cl
+++ b/kernels/test_write_only.cl
@@ -1,7 +1,7 @@
 __kernel void
 test_write_only(__global int *dst)
 {
-    int id = (int)get_global_id(0);
-    dst[id] = id;
+  int id = (int)get_global_id(0);
+  dst[id] = id;
 }
 
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 43f8556a..3345a1ad 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -7,7 +7,7 @@ ADD_LIBRARY(utests SHARED
             utest_file_map.cpp
             utest_assert.cpp
             utest.cpp
-            app_mandelbrot.cpp
+#app_mandelbrot.cpp
             compiler_write_only.cpp
             compiler_write_only_shorts.cpp
             compiler_write_only_bytes.cpp
@@ -24,7 +24,10 @@ ADD_LIBRARY(utests SHARED
             compiler_unstructured_branch0.cpp
             compiler_unstructured_branch1.cpp
             compiler_unstructured_branch2.cpp
-            compiler_unstructured_branch3.cpp)
+            compiler_unstructured_branch3.cpp
+            compiler_lower_return0.cpp
+            compiler_lower_return1.cpp
+            compiler_lower_return2.cpp)
 TARGET_LINK_LIBRARIES(utests cl m)
 
 ADD_EXECUTABLE(run utest_run.cpp)
diff --git a/utests/compiler_copy_buffer.cpp b/utests/compiler_copy_buffer.cpp
index ee1dbf1b..b0054241 100644
--- a/utests/compiler_copy_buffer.cpp
+++ b/utests/compiler_copy_buffer.cpp
@@ -25,6 +25,7 @@ static void compiler_copy_buffer(void)
 
   // Setup kernel and buffers
   OCL_CREATE_KERNEL("test_copy_buffer");
+  //OCL_CREATE_KERNEL("compiler_array");
   buf_data[0] = (uint32_t*) malloc(sizeof(uint32_t) * n);
   for (uint32_t i = 0; i < n; ++i) ((uint32_t*)buf_data[0])[i] = i;
   OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(uint32_t), buf_data[0]);
diff --git a/utests/compiler_lower_return0.cpp b/utests/compiler_lower_return0.cpp
new file mode 100644
index 00000000..47c7f683
--- /dev/null
+++ b/utests/compiler_lower_return0.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_lower_return0(void)
+{
+  const size_t n = 32;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_lower_return0");
+  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] = n;
+  locals[0] = 16;
+  OCL_NDRANGE(1);
+
+  // First control flow
+  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER(1);
+  for (int32_t i = 0; i < 32; ++i)
+    OCL_ASSERT(((int32_t*)buf_data[1])[i] == i);
+
+  // 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] == -2);
+
+  // 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 (int32_t i = 0; i < 8; ++i)
+    OCL_ASSERT(((int32_t*)buf_data[1])[i] == i);
+  for (int32_t i = 8; i < 32; ++i)
+    OCL_ASSERT(((int32_t*)buf_data[1])[i] == -2);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_lower_return0);
+
+
diff --git a/utests/compiler_lower_return1.cpp b/utests/compiler_lower_return1.cpp
new file mode 100644
index 00000000..0550cce3
--- /dev/null
+++ b/utests/compiler_lower_return1.cpp
@@ -0,0 +1,66 @@
+/* 
+ * 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_lower_return1(void)
+{
+  const size_t n = 32;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_lower_return1");
+  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 (int32_t i = 0; i < 11; ++i)
+    OCL_ASSERT(((int32_t*)buf_data[1])[i] == i);
+  for (int32_t i = 11; i < 16; ++i)
+    OCL_ASSERT(((int32_t*)buf_data[1])[i] == 2);
+
+  // Second control flow
+  for (uint32_t i = 0; i < 4; ++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 (int32_t i = 0; i < 4; ++i)
+    OCL_ASSERT(((int32_t*)buf_data[1])[i] == -2);
+  for (int32_t i = 4; i < 11; ++i)
+    OCL_ASSERT(((int32_t*)buf_data[1])[i] == i);
+  for (int32_t i = 11; i < 16; ++i)
+    OCL_ASSERT(((int32_t*)buf_data[1])[i] == 2);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_lower_return1);
+
diff --git a/utests/compiler_lower_return2.cpp b/utests/compiler_lower_return2.cpp
new file mode 100644
index 00000000..bdc9ef6c
--- /dev/null
+++ b/utests/compiler_lower_return2.cpp
@@ -0,0 +1,67 @@
+/* 
+ * 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 cpu(int global_id, int *src, int *dst) {
+  const int id = global_id;
+  dst[id] = id;
+  while (dst[id] > src[id]) {
+    if (dst[id] > 10) return;
+    dst[id]--;
+  }
+  dst[id] += 2;
+}
+
+static void compiler_lower_return2(void)
+{
+  const size_t n = 16;
+  int cpu_dst[16], cpu_src[16];
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_lower_return2");
+  OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL);
+  OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+  globals[0] = 16;
+  locals[0] = 16;
+
+  for (uint32_t pass = 0; pass < 8; ++pass) {
+    OCL_MAP_BUFFER(0);
+    for (int32_t i = 0; i < (int32_t) n; ++i)
+      cpu_src[i] = ((int32_t*)buf_data[0])[i] = rand() % 16;
+    OCL_UNMAP_BUFFER(0);
+
+    // Run the kernel on GPU
+    OCL_NDRANGE(1);
+
+    // Run on CPU
+    for (int32_t i = 0; i <(int32_t) n; ++i) cpu(i, cpu_src, cpu_dst);
+
+    // Compare
+    OCL_MAP_BUFFER(1);
+    for (int32_t i = 0; i < 11; ++i)
+      OCL_ASSERT(((int32_t*)buf_data[1])[i] == cpu_dst[i]);
+    OCL_UNMAP_BUFFER(1);
+  }
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_lower_return2);
+
diff --git a/utests/compiler_unstructured_branch0.cpp b/utests/compiler_unstructured_branch0.cpp
index 12ebade3..a314a519 100644
--- a/utests/compiler_unstructured_branch0.cpp
+++ b/utests/compiler_unstructured_branch0.cpp
@@ -72,4 +72,3 @@ static void compiler_unstructured_branch0(void)
 
 MAKE_UTEST_FROM_FUNCTION(compiler_unstructured_branch0);
 
-