Added support for some Gen extensions. The idea is mostly to play with Gen

specific hardware and propose simple way to use them. We have three extensions
here:
- Gen register regions. This allows us to perform strided loads in the register
file. To implement that on top of OCL, the idea is to encapsulate them in a
function with a side effect. Not really clean but it works.
- Gen gather from register file. Same idea but here we simply gather data from
a bunch of registers
- Vote any/all. This is basically the same idea as ptx i.e. uniform predicates
for branches.
- block read/write. Just to play with uniform load/store messages

I added a bunch of tests for all that and fix thing here and there to make them
work

diff --git a/kernels/compiler_array1.cl b/kernels/compiler_array1.cl
new file mode 100644 (file)
index 0000000..ad567c2
--- /dev/null
+++ b/kernels/compiler_array1.cl
@@ -0,0 +1,15 @@
+__kernel void
+compiler_array1(__global int *src, __global int *dst)
+{
+  int final[16];
+  for (int i = 0; i < 16; ++i) {
+    int array[16];
+    for (int j = 0; j < src[0]; ++j)
+      array[j] = 1+src[0];
+    for (int j = src[0]; j < 16; ++j)
+      array[j] = get_global_id(0);
+    final[i] = array[i];
+  }
+  dst[get_global_id(0)] = final[get_global_id(0)];
+}
+

diff --git a/kernels/compiler_array2.cl b/kernels/compiler_array2.cl
new file mode 100644 (file)
index 0000000..ae73932
--- /dev/null
+++ b/kernels/compiler_array2.cl
@@ -0,0 +1,13 @@
+__kernel void
+compiler_array2(__global int *src, __global int *dst)
+{
+  int final[16];
+  int array[16];
+  for (int j = 0; j < 16; ++j) array[j] = j;
+  for (int j = 0; j < 16; ++j) final[j] = j+1;
+  if (get_global_id(0) == 15)
+    dst[get_global_id(0)] = final[get_global_id(0)];
+  else
+    dst[get_global_id(0)] = array[15 - get_global_id(0)];
+}
+

diff --git a/kernels/compiler_array3.cl b/kernels/compiler_array3.cl
new file mode 100644 (file)
index 0000000..152c22a
--- /dev/null
+++ b/kernels/compiler_array3.cl
@@ -0,0 +1,14 @@
+__kernel void
+compiler_array3(__global int *src, __global int *dst)
+{
+  int tmp[32];
+  for (int i = 0; i < 16; ++i) {
+    for (int j = 0; j < 16; ++j)
+      tmp[j] = get_global_id(0);
+    for (int j = 0; j < src[0]; ++j)
+      tmp[j] = 1+src[j];
+    tmp[16+i] = tmp[i];
+  }
+  dst[get_global_id(0)] = tmp[16+get_global_id(0)];
+}
+

diff --git a/kernels/compiler_gather_register_file.cl b/kernels/compiler_gather_register_file.cl
new file mode 100644 (file)
index 0000000..773797d
--- /dev/null
+++ b/kernels/compiler_gather_register_file.cl
@@ -0,0 +1,10 @@
+__kernel void
+compiler_gather_register_file(__global uint *src, __global uint *dst)
+{
+  __gen_ocl_force_simd16();
+  int id = (int)get_global_id(0);
+  const int x0 = src[id];
+  const unsigned short index = get_global_id(0);
+  dst[id] = __gen_ocl_rgather(index, x0);
+}
+

diff --git a/kernels/compiler_gather_register_file0.cl b/kernels/compiler_gather_register_file0.cl
new file mode 100644 (file)
index 0000000..0e6d487
--- /dev/null
+++ b/kernels/compiler_gather_register_file0.cl
@@ -0,0 +1,10 @@
+__kernel void
+compiler_gather_register_file0(__global uint *src, __global uint *dst)
+{
+  __gen_ocl_force_simd16();
+  int id = (int)get_global_id(0);
+  const int x0 = src[id];
+  const unsigned short index = 15 - get_global_id(0);
+  dst[id] = __gen_ocl_rgather(index, x0);
+}
+

diff --git a/kernels/compiler_gather_register_file1.cl b/kernels/compiler_gather_register_file1.cl
new file mode 100644 (file)
index 0000000..184202c
--- /dev/null
+++ b/kernels/compiler_gather_register_file1.cl
@@ -0,0 +1,11 @@
+__kernel void
+compiler_gather_register_file1(__global uint *src, __global uint *dst)
+{
+  __gen_ocl_force_simd16();
+  int id = (int)get_global_id(0);
+  const int x0 = src[id];
+  const int x1 = src[id+16];
+  const unsigned short index = 2*get_global_id(0);
+  dst[id] = __gen_ocl_rgather(index, x0, x1);
+}
+

diff --git a/kernels/compiler_obread.cl b/kernels/compiler_obread.cl
new file mode 100644 (file)
index 0000000..14658d9
--- /dev/null
+++ b/kernels/compiler_obread.cl
@@ -0,0 +1,8 @@
+__kernel void
+compiler_obread(__global uint *src, __global uint *dst)
+{
+  int id = (int)get_global_id(0);
+  const int to =  __gen_ocl_obread(src+id);
+  dst[id] = to;
+}
+

diff --git a/kernels/compiler_obwrite.cl b/kernels/compiler_obwrite.cl
new file mode 100644 (file)
index 0000000..50e55a1
--- /dev/null
+++ b/kernels/compiler_obwrite.cl
@@ -0,0 +1,8 @@
+__kernel void
+compiler_obwrite(__global uint *src, __global uint *dst)
+{
+  int id = (int)get_global_id(0);
+  const int to =  src[id];
+  __gen_ocl_obwrite(dst+id,to);
+}
+

diff --git a/kernels/compiler_region.cl b/kernels/compiler_region.cl
new file mode 100644 (file)
index 0000000..d74ac7d
--- /dev/null
+++ b/kernels/compiler_region.cl
@@ -0,0 +1,10 @@
+__kernel void
+compiler_region(__global uint *src, __global uint *dst)
+{
+  __gen_ocl_force_simd16();
+  int id = (int)get_global_id(0);
+  const int x0 = src[id];
+  const int x1 = src[id+16];
+  dst[id] = __gen_ocl_region(0, 16, 8, 2, x0, x1);
+}
+

diff --git a/kernels/compiler_region0.cl b/kernels/compiler_region0.cl
new file mode 100644 (file)
index 0000000..5bd57c0
--- /dev/null
+++ b/kernels/compiler_region0.cl
@@ -0,0 +1,11 @@
+__kernel void
+compiler_region0(__global uint *src, __global uint *dst)
+{
+  __gen_ocl_force_simd16();
+  int id = (int)get_global_id(0);
+  const int x0 = src[id];
+  const int x1 = src[id+16];
+  const int x2 = src[id+32];
+  dst[id] = __gen_ocl_region(1, 16, 8, 2, x0, x1, x2);
+}
+

diff --git a/kernels/compiler_region1.cl b/kernels/compiler_region1.cl
new file mode 100644 (file)
index 0000000..9deb63c
--- /dev/null
+++ b/kernels/compiler_region1.cl
@@ -0,0 +1,9 @@
+__kernel void
+compiler_region1(__global uint *src, __global uint *dst)
+{
+  __gen_ocl_force_simd16();
+  int id = (int)get_global_id(0);
+  const int x0 = src[id];
+  dst[id] = __gen_ocl_region(0, 16, 8, 2, x0);
+}
+

diff --git a/kernels/compiler_vote_all.cl b/kernels/compiler_vote_all.cl
new file mode 100644 (file)
index 0000000..1918c1c
--- /dev/null
+++ b/kernels/compiler_vote_all.cl
@@ -0,0 +1,10 @@
+__kernel void
+compiler_vote_all(__global uint *src, __global uint *dst)
+{
+  int id = (int)get_global_id(0);
+  if (__gen_ocl_all(id > 8))
+    dst[id] = src[id];
+  else
+    dst[id] = 0;
+}
+

diff --git a/kernels/compiler_vote_any.cl b/kernels/compiler_vote_any.cl
new file mode 100644 (file)
index 0000000..0a81e89
--- /dev/null
+++ b/kernels/compiler_vote_any.cl
@@ -0,0 +1,10 @@
+__kernel void
+compiler_vote_any(__global uint *src, __global uint *dst)
+{
+  int id = (int)get_global_id(0);
+  if (__gen_ocl_any(id > 6))
+    dst[id] = src[id];
+  else
+    dst[id] = 0;
+}
+

diff --git a/kernels/compiler_write_only_bytes.cl b/kernels/compiler_write_only_bytes.cl
index d17e54a..0bc0cd8 100644 (file)
--- a/kernels/compiler_write_only_bytes.cl
+++ b/kernels/compiler_write_only_bytes.cl
@@ -1,7 +1,7 @@
-__kernel void
-compiler_write_only_bytes(__global char *dst)
-{
-    int id = (int)get_global_id(0);
-    dst[id] = 2;
-}
-
+__kernel void\r
+compiler_write_only_bytes(__global char *dst)\r
+{\r
+    int id = (int)get_global_id(0);\r
+    dst[id] = 2;\r
+}\r
+\r

diff --git a/kernels/compiler_write_only_shorts.cl b/kernels/compiler_write_only_shorts.cl
index d5e2f26..bfd23cc 100644 (file)
--- a/kernels/compiler_write_only_shorts.cl
+++ b/kernels/compiler_write_only_shorts.cl
@@ -1,7 +1,7 @@
-__kernel void
-compiler_write_only_shorts(__global short *dst)
-{
-    int id = (int)get_global_id(0);
-    dst[id] = 2;
-}
-
+__kernel void\r
+compiler_write_only_shorts(__global short *dst)\r
+{\r
+    int id = (int)get_global_id(0);\r
+    dst[id] = 2;\r
+}\r
+\r

diff --git a/utests/compiler_argument_structure.cpp b/utests/compiler_argument_structure.cpp
index 512e9e7..22464a5 100644 (file)
--- a/utests/compiler_argument_structure.cpp
+++ b/utests/compiler_argument_structure.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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"
 
 struct hop { int x, y; };

diff --git a/utests/compiler_argument_structure_indirect.cpp b/utests/compiler_argument_structure_indirect.cpp
index e572f81..a4584d5 100644 (file)
--- a/utests/compiler_argument_structure_indirect.cpp
+++ b/utests/compiler_argument_structure_indirect.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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"
 
 struct hop { int x[16]; };

diff --git a/utests/compiler_array.cpp b/utests/compiler_array.cpp
index 0cec7d8..8806c99 100644 (file)
--- a/utests/compiler_array.cpp
+++ b/utests/compiler_array.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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"
 
 void compiler_array(void)

diff --git a/utests/compiler_array0.cpp b/utests/compiler_array0.cpp
index 9e3535d..7cf2bbb 100644 (file)
--- a/utests/compiler_array0.cpp
+++ b/utests/compiler_array0.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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) {

diff --git a/utests/compiler_array1.cpp b/utests/compiler_array1.cpp
new file mode 100644 (file)
index 0000000..fe1ecec
--- /dev/null
+++ b/utests/compiler_array1.cpp
@@ -0,0 +1,52 @@
+#include "utest_helper.hpp"
+
+static void cpu(int global_id, int *src, int *dst) {
+  int final[16];
+  for (int i = 0; i < 16; ++i) {
+    int array[16];
+    for (int j = 0; j < src[0]; ++j)
+      array[j] = 1+src[0];
+    for (int j = src[0]; j < 16; ++j)
+      array[j] = global_id;
+    final[i] = array[i];
+  }
+  dst[global_id] = final[global_id];
+}
+
+void compiler_array1(void)
+{
+  const size_t n = 16;
+  int cpu_dst[16], cpu_src[16];
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_array1");
+  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;
+
+  // Run random tests
+  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_array1);
+

diff --git a/utests/compiler_array2.cpp b/utests/compiler_array2.cpp
new file mode 100644 (file)
index 0000000..61ca9da
--- /dev/null
+++ b/utests/compiler_array2.cpp
@@ -0,0 +1,50 @@
+#include "utest_helper.hpp"
+
+static void cpu(int global_id, int *src, int *dst) {
+  int final[16];
+  int array[16];
+  for (int j = 0; j < 16; ++j) array[j] = j;
+  for (int j = 0; j < 16; ++j) final[j] = j+1;
+  if (global_id == 15)
+    dst[global_id] = final[global_id];
+  else
+    dst[global_id] = array[15 - global_id];
+}
+
+void compiler_array2(void)
+{
+  const size_t n = 16;
+  int cpu_dst[16], cpu_src[16];
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_array2");
+  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;
+
+  // Run random tests
+  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_array2);
+

diff --git a/utests/compiler_array3.cpp b/utests/compiler_array3.cpp
new file mode 100644 (file)
index 0000000..865b1e5
--- /dev/null
+++ b/utests/compiler_array3.cpp
@@ -0,0 +1,51 @@
+#include "utest_helper.hpp"
+
+static void cpu(int global_id, int *src, int *dst) {
+  int tmp[32];
+  for (int i = 0; i < 16; ++i) {
+    for (int j = 0; j < 16; ++j)
+      tmp[j] = global_id;
+    for (int j = 0; j < src[0]; ++j)
+      tmp[j] = 1+src[j];
+    tmp[16+i] = tmp[i];
+  }
+  dst[global_id] = tmp[16+global_id];
+}
+
+void compiler_array3(void)
+{
+  const size_t n = 16;
+  int cpu_dst[16], cpu_src[16];
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_array3");
+  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;
+
+  // Run random tests
+  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_array3);
+

diff --git a/utests/compiler_byte_scatter.cpp b/utests/compiler_byte_scatter.cpp
index 115f5df..11300da 100644 (file)
--- a/utests/compiler_byte_scatter.cpp
+++ b/utests/compiler_byte_scatter.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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_byte_scatter(void)

diff --git a/utests/compiler_copy_buffer.cpp b/utests/compiler_copy_buffer.cpp
index b005424..8066efe 100644 (file)
--- a/utests/compiler_copy_buffer.cpp
+++ b/utests/compiler_copy_buffer.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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_copy_buffer(void)

diff --git a/utests/compiler_copy_buffer_row.cpp b/utests/compiler_copy_buffer_row.cpp
index cfaea0e..12c0592 100644 (file)
--- a/utests/compiler_copy_buffer_row.cpp
+++ b/utests/compiler_copy_buffer_row.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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_copy_buffer_row(void)

diff --git a/utests/compiler_function_argument.cpp b/utests/compiler_function_argument.cpp
index 29775ce..a39523b 100644 (file)
--- a/utests/compiler_function_argument.cpp
+++ b/utests/compiler_function_argument.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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"
 
 void compiler_function_argument(void)

diff --git a/utests/compiler_function_argument0.cpp b/utests/compiler_function_argument0.cpp
index 3aac9cb..2e4227e 100644 (file)
--- a/utests/compiler_function_argument0.cpp
+++ b/utests/compiler_function_argument0.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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"
 
 void compiler_function_argument0(void)

diff --git a/utests/compiler_function_argument1.cpp b/utests/compiler_function_argument1.cpp
index dff46a8..48a7677 100644 (file)
--- a/utests/compiler_function_argument1.cpp
+++ b/utests/compiler_function_argument1.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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"
 
 void compiler_function_argument1(void)

diff --git a/utests/compiler_gather_register_file.cpp b/utests/compiler_gather_register_file.cpp
new file mode 100644 (file)
index 0000000..a877b58
--- /dev/null
+++ b/utests/compiler_gather_register_file.cpp
@@ -0,0 +1,31 @@
+#include "utest_helper.hpp"
+
+static void compiler_gather_register_file(void)
+{
+  const size_t n = 48;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_gather_register_file");
+  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]);
+  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);
+
+  // Check result
+  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER(1);
+  for (uint32_t i = 0; i < 16; ++i)
+    OCL_ASSERT(((uint32_t*)buf_data[1])[i] == i);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_gather_register_file);
+

diff --git a/utests/compiler_gather_register_file0.cpp b/utests/compiler_gather_register_file0.cpp
new file mode 100644 (file)
index 0000000..f9e293e
--- /dev/null
+++ b/utests/compiler_gather_register_file0.cpp
@@ -0,0 +1,32 @@
+#include "utest_helper.hpp"
+
+static void compiler_gather_register_file0(void)
+{
+  const size_t n = 48;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_gather_register_file0");
+  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]);
+  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);
+
+  // Check result
+  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER(1);
+  for (uint32_t i = 0; i < 16; ++i)
+    OCL_ASSERT(((uint32_t*)buf_data[1])[i] == 15-i);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_gather_register_file0);
+
+

diff --git a/utests/compiler_gather_register_file1.cpp b/utests/compiler_gather_register_file1.cpp
new file mode 100644 (file)
index 0000000..4956b7e
--- /dev/null
+++ b/utests/compiler_gather_register_file1.cpp
@@ -0,0 +1,31 @@
+#include "utest_helper.hpp"
+
+static void compiler_gather_register_file1(void)
+{
+  const size_t n = 48;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_gather_register_file1");
+  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]);
+  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);
+
+  // Check result
+  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*i);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_gather_register_file1);
+

diff --git a/utests/compiler_if_else.cpp b/utests/compiler_if_else.cpp
index 44b1a87..e38b23f 100644 (file)
--- a/utests/compiler_if_else.cpp
+++ b/utests/compiler_if_else.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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)

diff --git a/utests/compiler_local_slm.cpp b/utests/compiler_local_slm.cpp
index f9c9e0c..aa9a2fe 100644 (file)
--- a/utests/compiler_local_slm.cpp
+++ b/utests/compiler_local_slm.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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"
 
 void compiler_local_slm(void)

diff --git a/utests/compiler_lower_return0.cpp b/utests/compiler_lower_return0.cpp
index 47c7f68..0e9dbd0 100644 (file)
--- a/utests/compiler_lower_return0.cpp
+++ b/utests/compiler_lower_return0.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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)

diff --git a/utests/compiler_lower_return1.cpp b/utests/compiler_lower_return1.cpp
index 0550cce..b4f1fe3 100644 (file)
--- a/utests/compiler_lower_return1.cpp
+++ b/utests/compiler_lower_return1.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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)

diff --git a/utests/compiler_lower_return2.cpp b/utests/compiler_lower_return2.cpp
index bdc9ef6..1e34036 100644 (file)
--- a/utests/compiler_lower_return2.cpp
+++ b/utests/compiler_lower_return2.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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) {

diff --git a/utests/compiler_obread.cpp b/utests/compiler_obread.cpp
new file mode 100644 (file)
index 0000000..99831e2
--- /dev/null
+++ b/utests/compiler_obread.cpp
@@ -0,0 +1,31 @@
+#include "utest_helper.hpp"
+
+static void compiler_obread(void)
+{
+  const size_t n = 48;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_obread");
+  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]);
+  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);
+
+  // Check result
+  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER(1);
+  for (uint32_t i = 0; i < 16; ++i)
+    OCL_ASSERT(((uint32_t*)buf_data[1])[i] == i);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_obread);
+

diff --git a/utests/compiler_obwrite.cpp b/utests/compiler_obwrite.cpp
new file mode 100644 (file)
index 0000000..5f6b63b
--- /dev/null
+++ b/utests/compiler_obwrite.cpp
@@ -0,0 +1,31 @@
+#include "utest_helper.hpp"
+
+static void compiler_obwrite(void)
+{
+  const size_t n = 48;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_obwrite");
+  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]);
+  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);
+
+  // Check result
+  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER(1);
+  for (uint32_t i = 0; i < 16; ++i)
+    OCL_ASSERT(((uint32_t*)buf_data[1])[i] == i);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_obwrite);
+

diff --git a/utests/compiler_region.cpp b/utests/compiler_region.cpp
new file mode 100644 (file)
index 0000000..395c227
--- /dev/null
+++ b/utests/compiler_region.cpp
@@ -0,0 +1,31 @@
+#include "utest_helper.hpp"
+
+static void compiler_region(void)
+{
+  const size_t n = 32;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_region");
+  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]);
+  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);
+
+  // Check result
+  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*i);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_region);
+

diff --git a/utests/compiler_region0.cpp b/utests/compiler_region0.cpp
new file mode 100644 (file)
index 0000000..56ddb68
--- /dev/null
+++ b/utests/compiler_region0.cpp
@@ -0,0 +1,31 @@
+#include "utest_helper.hpp"
+
+static void compiler_region0(void)
+{
+  const size_t n = 48;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_region0");
+  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]);
+  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);
+
+  // Check result
+  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*i+1);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_region0);
+

diff --git a/utests/compiler_region1.cpp b/utests/compiler_region1.cpp
new file mode 100644 (file)
index 0000000..9b03b8e
--- /dev/null
+++ b/utests/compiler_region1.cpp
@@ -0,0 +1,33 @@
+#include "utest_helper.hpp"
+
+static void compiler_region1(void)
+{
+  const size_t n = 32;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_region1");
+  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]);
+  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);
+
+  // Check result
+  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER(1);
+  for (uint32_t i = 0; i < 8; ++i)
+    OCL_ASSERT(((uint32_t*)buf_data[1])[i] == 2*i);
+  for (uint32_t i = 8; i < 16; ++i)
+    OCL_ASSERT(((uint32_t*)buf_data[1])[i] == 0);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_region1);
+

diff --git a/utests/compiler_short_scatter.cpp b/utests/compiler_short_scatter.cpp
index 8f5f4c2..1746744 100644 (file)
--- a/utests/compiler_short_scatter.cpp
+++ b/utests/compiler_short_scatter.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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_short_scatter(void)

diff --git a/utests/compiler_sub_bytes.cpp b/utests/compiler_sub_bytes.cpp
index f2262f1..49a5261 100644 (file)
--- a/utests/compiler_sub_bytes.cpp
+++ b/utests/compiler_sub_bytes.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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_sub_bytes(void)

diff --git a/utests/compiler_sub_shorts.cpp b/utests/compiler_sub_shorts.cpp
index bf4271f..4aeeca3 100644 (file)
--- a/utests/compiler_sub_shorts.cpp
+++ b/utests/compiler_sub_shorts.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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_sub_shorts(void)

diff --git a/utests/compiler_uint2_copy.cpp b/utests/compiler_uint2_copy.cpp
index 386cfc0..8eb4314 100644 (file)
--- a/utests/compiler_uint2_copy.cpp
+++ b/utests/compiler_uint2_copy.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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_uint2_copy(void)

diff --git a/utests/compiler_uint3_copy.cpp b/utests/compiler_uint3_copy.cpp
index 277d6ef..320f405 100644 (file)
--- a/utests/compiler_uint3_copy.cpp
+++ b/utests/compiler_uint3_copy.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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_uint3_copy(void)

diff --git a/utests/compiler_uint3_unaligned_copy.cpp b/utests/compiler_uint3_unaligned_copy.cpp
index 9ccdb42..d42b4c3 100644 (file)
--- a/utests/compiler_uint3_unaligned_copy.cpp
+++ b/utests/compiler_uint3_unaligned_copy.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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_uint3_unaligned_copy(void)

diff --git a/utests/compiler_unstructured_branch0.cpp b/utests/compiler_unstructured_branch0.cpp
index a314a51..128a53e 100644 (file)
--- a/utests/compiler_unstructured_branch0.cpp
+++ b/utests/compiler_unstructured_branch0.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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)

diff --git a/utests/compiler_unstructured_branch1.cpp b/utests/compiler_unstructured_branch1.cpp
index 5c3614c..6021f5b 100644 (file)
--- a/utests/compiler_unstructured_branch1.cpp
+++ b/utests/compiler_unstructured_branch1.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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)

diff --git a/utests/compiler_unstructured_branch2.cpp b/utests/compiler_unstructured_branch2.cpp
index 1808ef1..d61c6b5 100644 (file)
--- a/utests/compiler_unstructured_branch2.cpp
+++ b/utests/compiler_unstructured_branch2.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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)

diff --git a/utests/compiler_unstructured_branch3.cpp b/utests/compiler_unstructured_branch3.cpp
index 732bcf2..0c6992a 100644 (file)
--- a/utests/compiler_unstructured_branch3.cpp
+++ b/utests/compiler_unstructured_branch3.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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)

diff --git a/utests/compiler_vote_all.cpp b/utests/compiler_vote_all.cpp
new file mode 100644 (file)
index 0000000..9fcffd9
--- /dev/null
+++ b/utests/compiler_vote_all.cpp
@@ -0,0 +1,33 @@
+#include "utest_helper.hpp"
+
+static void compiler_vote_all(void)
+{
+  const size_t n = 32;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_vote_all");
+  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]);
+  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] = 32;
+  locals[0] = 16;
+  OCL_NDRANGE(1);
+
+  // Check result
+  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER(1);
+  for (uint32_t i = 0; i < 16; ++i)
+    OCL_ASSERT(((uint32_t*)buf_data[1])[i] == 0);
+  for (uint32_t i = 16; i < 32; ++i)
+    OCL_ASSERT(((uint32_t*)buf_data[1])[i] == i);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_vote_all);
+

diff --git a/utests/compiler_vote_any.cpp b/utests/compiler_vote_any.cpp
new file mode 100644 (file)
index 0000000..91bcb03
--- /dev/null
+++ b/utests/compiler_vote_any.cpp
@@ -0,0 +1,31 @@
+#include "utest_helper.hpp"
+
+static void compiler_vote_any(void)
+{
+  const size_t n = 32;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_vote_any");
+  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]);
+  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] = 32;
+  locals[0] = 16;
+  OCL_NDRANGE(1);
+
+  // Check result
+  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER(1);
+  for (uint32_t i = 0; i < 32; ++i)
+    OCL_ASSERT(((uint32_t*)buf_data[1])[i] == i);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_vote_any);
+

diff --git a/utests/compiler_write_only.cpp b/utests/compiler_write_only.cpp
index b42bfd2..e04f855 100644 (file)
--- a/utests/compiler_write_only.cpp
+++ b/utests/compiler_write_only.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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"
 
 void compiler_write_only(void)

diff --git a/utests/compiler_write_only_bytes.cpp b/utests/compiler_write_only_bytes.cpp
index 19a8ee1..1a13cdb 100644 (file)
--- a/utests/compiler_write_only_bytes.cpp
+++ b/utests/compiler_write_only_bytes.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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"
 
 void compiler_write_only_bytes(void)

diff --git a/utests/compiler_write_only_shorts.cpp b/utests/compiler_write_only_shorts.cpp
index 4db5e11..19988fe 100644 (file)
--- a/utests/compiler_write_only_shorts.cpp
+++ b/utests/compiler_write_only_shorts.cpp
@@ -1,22 +1,3 @@
-/* 
- * 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"
 
 void compiler_write_only_shorts(void)