diff options
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) |