summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
-rw-r--r--kernels/compiler_array1.cl15
-rw-r--r--kernels/compiler_array2.cl13
-rw-r--r--kernels/compiler_array3.cl14
-rw-r--r--kernels/compiler_gather_register_file.cl10
-rw-r--r--kernels/compiler_gather_register_file0.cl10
-rw-r--r--kernels/compiler_gather_register_file1.cl11
-rw-r--r--kernels/compiler_obread.cl8
-rw-r--r--kernels/compiler_obwrite.cl8
-rw-r--r--kernels/compiler_region.cl10
-rw-r--r--kernels/compiler_region0.cl11
-rw-r--r--kernels/compiler_region1.cl9
-rw-r--r--kernels/compiler_vote_all.cl10
-rw-r--r--kernels/compiler_vote_any.cl10
-rw-r--r--kernels/compiler_write_only_bytes.cl14
-rw-r--r--kernels/compiler_write_only_shorts.cl14
-rw-r--r--utests/compiler_argument_structure.cpp19
-rw-r--r--utests/compiler_argument_structure_indirect.cpp19
-rw-r--r--utests/compiler_array.cpp19
-rw-r--r--utests/compiler_array0.cpp19
-rw-r--r--utests/compiler_array1.cpp52
-rw-r--r--utests/compiler_array2.cpp50
-rw-r--r--utests/compiler_array3.cpp51
-rw-r--r--utests/compiler_byte_scatter.cpp19
-rw-r--r--utests/compiler_copy_buffer.cpp19
-rw-r--r--utests/compiler_copy_buffer_row.cpp19
-rw-r--r--utests/compiler_function_argument.cpp19
-rw-r--r--utests/compiler_function_argument0.cpp19
-rw-r--r--utests/compiler_function_argument1.cpp19
-rw-r--r--utests/compiler_gather_register_file.cpp31
-rw-r--r--utests/compiler_gather_register_file0.cpp32
-rw-r--r--utests/compiler_gather_register_file1.cpp31
-rw-r--r--utests/compiler_if_else.cpp19
-rw-r--r--utests/compiler_local_slm.cpp19
-rw-r--r--utests/compiler_lower_return0.cpp19
-rw-r--r--utests/compiler_lower_return1.cpp19
-rw-r--r--utests/compiler_lower_return2.cpp19
-rw-r--r--utests/compiler_obread.cpp31
-rw-r--r--utests/compiler_obwrite.cpp31
-rw-r--r--utests/compiler_region.cpp31
-rw-r--r--utests/compiler_region0.cpp31
-rw-r--r--utests/compiler_region1.cpp33
-rw-r--r--utests/compiler_short_scatter.cpp19
-rw-r--r--utests/compiler_sub_bytes.cpp19
-rw-r--r--utests/compiler_sub_shorts.cpp19
-rw-r--r--utests/compiler_uint2_copy.cpp19
-rw-r--r--utests/compiler_uint3_copy.cpp19
-rw-r--r--utests/compiler_uint3_unaligned_copy.cpp19
-rw-r--r--utests/compiler_unstructured_branch0.cpp19
-rw-r--r--utests/compiler_unstructured_branch1.cpp19
-rw-r--r--utests/compiler_unstructured_branch2.cpp19
-rw-r--r--utests/compiler_unstructured_branch3.cpp19
-rw-r--r--utests/compiler_vote_all.cpp33
-rw-r--r--utests/compiler_vote_any.cpp31
-rw-r--r--utests/compiler_write_only.cpp19
-rw-r--r--utests/compiler_write_only_bytes.cpp19
-rw-r--r--utests/compiler_write_only_shorts.cpp19
56 files changed, 621 insertions, 546 deletions
diff --git a/kernels/compiler_array1.cl b/kernels/compiler_array1.cl
new file mode 100644
index 00000000..ad567c2e
--- /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
index 00000000..ae73932a
--- /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
index 00000000..152c22a0
--- /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
index 00000000..773797d0
--- /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
index 00000000..0e6d4875
--- /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
index 00000000..184202c6
--- /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
index 00000000..14658d96
--- /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
index 00000000..50e55a17
--- /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
index 00000000..d74ac7d3
--- /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
index 00000000..5bd57c05
--- /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
index 00000000..9deb63cd
--- /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
index 00000000..1918c1cd
--- /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
index 00000000..0a81e893
--- /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 d17e54a8..0bc0cd8b 100644
--- 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
+compiler_write_only_bytes(__global char *dst)
+{
+ int id = (int)get_global_id(0);
+ dst[id] = 2;
+}
+
diff --git a/kernels/compiler_write_only_shorts.cl b/kernels/compiler_write_only_shorts.cl
index d5e2f265..bfd23ccf 100644
--- 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
+compiler_write_only_shorts(__global short *dst)
+{
+ int id = (int)get_global_id(0);
+ dst[id] = 2;
+}
+
diff --git a/utests/compiler_argument_structure.cpp b/utests/compiler_argument_structure.cpp
index 512e9e73..22464a5f 100644
--- 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 e572f81a..a4584d52 100644
--- 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 0cec7d8d..8806c993 100644
--- 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 9e3535dc..7cf2bbb2 100644
--- 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
index 00000000..fe1ecec3
--- /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
index 00000000..61ca9da8
--- /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
index 00000000..865b1e54
--- /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 115f5df3..11300daa 100644
--- 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 b0054241..8066efee 100644
--- 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 cfaea0e9..12c05929 100644
--- 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 29775ceb..a39523be 100644
--- 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 3aac9cb7..2e4227e4 100644
--- 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 dff46a81..48a76776 100644
--- 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
index 00000000..a877b585
--- /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
index 00000000..f9e293ed
--- /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
index 00000000..4956b7e1
--- /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 44b1a87a..e38b23ff 100644
--- 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 f9c9e0c5..aa9a2fec 100644
--- 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 47c7f683..0e9dbd0d 100644
--- 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 0550cce3..b4f1fe3e 100644
--- 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 bdc9ef6c..1e34036e 100644
--- 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
index 00000000..99831e23
--- /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
index 00000000..5f6b63b5
--- /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
index 00000000..395c227e
--- /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
index 00000000..56ddb68b
--- /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
index 00000000..9b03b8e6
--- /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 8f5f4c2f..17467441 100644
--- 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 f2262f16..49a52619 100644
--- 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 bf4271f1..4aeeca35 100644
--- 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 386cfc09..8eb43142 100644
--- 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 277d6efb..320f4058 100644
--- 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 9ccdb420..d42b4c38 100644
--- 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 a314a519..128a53ed 100644
--- 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 5c3614c9..6021f5b6 100644
--- 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 1808ef13..d61c6b56 100644
--- 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 732bcf2a..0c6992a8 100644
--- 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
index 00000000..9fcffd92
--- /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
index 00000000..91bcb03d
--- /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 b42bfd22..e04f855b 100644
--- 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 19a8ee19..1a13cdb4 100644
--- 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 4db5e119..19988fe1 100644
--- 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)