From: bsegovia Date: Mon, 17 Sep 2012 12:46:31 +0000 (+0000) Subject: Added support for some Gen extensions. The idea is mostly to play with Gen X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=c27f1b7668c068621e5ebb4282704afbdbf907be;p=contrib%2Fbeignet.git Added support for some Gen extensions. The idea is mostly to play with Gen specific hardware and propose simple way to use them. We have three extensions here: - Gen register regions. This allows us to perform strided loads in the register file. To implement that on top of OCL, the idea is to encapsulate them in a function with a side effect. Not really clean but it works. - Gen gather from register file. Same idea but here we simply gather data from a bunch of registers - Vote any/all. This is basically the same idea as ptx i.e. uniform predicates for branches. - block read/write. Just to play with uniform load/store messages I added a bunch of tests for all that and fix thing here and there to make them work --- diff --git a/kernels/compiler_array1.cl b/kernels/compiler_array1.cl new file mode 100644 index 0000000..ad567c2 --- /dev/null +++ b/kernels/compiler_array1.cl @@ -0,0 +1,15 @@ +__kernel void +compiler_array1(__global int *src, __global int *dst) +{ + int final[16]; + for (int i = 0; i < 16; ++i) { + int array[16]; + for (int j = 0; j < src[0]; ++j) + array[j] = 1+src[0]; + for (int j = src[0]; j < 16; ++j) + array[j] = get_global_id(0); + final[i] = array[i]; + } + dst[get_global_id(0)] = final[get_global_id(0)]; +} + diff --git a/kernels/compiler_array2.cl b/kernels/compiler_array2.cl new file mode 100644 index 0000000..ae73932 --- /dev/null +++ b/kernels/compiler_array2.cl @@ -0,0 +1,13 @@ +__kernel void +compiler_array2(__global int *src, __global int *dst) +{ + int final[16]; + int array[16]; + for (int j = 0; j < 16; ++j) array[j] = j; + for (int j = 0; j < 16; ++j) final[j] = j+1; + if (get_global_id(0) == 15) + dst[get_global_id(0)] = final[get_global_id(0)]; + else + dst[get_global_id(0)] = array[15 - get_global_id(0)]; +} + diff --git a/kernels/compiler_array3.cl b/kernels/compiler_array3.cl new file mode 100644 index 0000000..152c22a --- /dev/null +++ b/kernels/compiler_array3.cl @@ -0,0 +1,14 @@ +__kernel void +compiler_array3(__global int *src, __global int *dst) +{ + int tmp[32]; + for (int i = 0; i < 16; ++i) { + for (int j = 0; j < 16; ++j) + tmp[j] = get_global_id(0); + for (int j = 0; j < src[0]; ++j) + tmp[j] = 1+src[j]; + tmp[16+i] = tmp[i]; + } + dst[get_global_id(0)] = tmp[16+get_global_id(0)]; +} + diff --git a/kernels/compiler_gather_register_file.cl b/kernels/compiler_gather_register_file.cl new file mode 100644 index 0000000..773797d --- /dev/null +++ b/kernels/compiler_gather_register_file.cl @@ -0,0 +1,10 @@ +__kernel void +compiler_gather_register_file(__global uint *src, __global uint *dst) +{ + __gen_ocl_force_simd16(); + int id = (int)get_global_id(0); + const int x0 = src[id]; + const unsigned short index = get_global_id(0); + dst[id] = __gen_ocl_rgather(index, x0); +} + diff --git a/kernels/compiler_gather_register_file0.cl b/kernels/compiler_gather_register_file0.cl new file mode 100644 index 0000000..0e6d487 --- /dev/null +++ b/kernels/compiler_gather_register_file0.cl @@ -0,0 +1,10 @@ +__kernel void +compiler_gather_register_file0(__global uint *src, __global uint *dst) +{ + __gen_ocl_force_simd16(); + int id = (int)get_global_id(0); + const int x0 = src[id]; + const unsigned short index = 15 - get_global_id(0); + dst[id] = __gen_ocl_rgather(index, x0); +} + diff --git a/kernels/compiler_gather_register_file1.cl b/kernels/compiler_gather_register_file1.cl new file mode 100644 index 0000000..184202c --- /dev/null +++ b/kernels/compiler_gather_register_file1.cl @@ -0,0 +1,11 @@ +__kernel void +compiler_gather_register_file1(__global uint *src, __global uint *dst) +{ + __gen_ocl_force_simd16(); + int id = (int)get_global_id(0); + const int x0 = src[id]; + const int x1 = src[id+16]; + const unsigned short index = 2*get_global_id(0); + dst[id] = __gen_ocl_rgather(index, x0, x1); +} + diff --git a/kernels/compiler_obread.cl b/kernels/compiler_obread.cl new file mode 100644 index 0000000..14658d9 --- /dev/null +++ b/kernels/compiler_obread.cl @@ -0,0 +1,8 @@ +__kernel void +compiler_obread(__global uint *src, __global uint *dst) +{ + int id = (int)get_global_id(0); + const int to = __gen_ocl_obread(src+id); + dst[id] = to; +} + diff --git a/kernels/compiler_obwrite.cl b/kernels/compiler_obwrite.cl new file mode 100644 index 0000000..50e55a1 --- /dev/null +++ b/kernels/compiler_obwrite.cl @@ -0,0 +1,8 @@ +__kernel void +compiler_obwrite(__global uint *src, __global uint *dst) +{ + int id = (int)get_global_id(0); + const int to = src[id]; + __gen_ocl_obwrite(dst+id,to); +} + diff --git a/kernels/compiler_region.cl b/kernels/compiler_region.cl new file mode 100644 index 0000000..d74ac7d --- /dev/null +++ b/kernels/compiler_region.cl @@ -0,0 +1,10 @@ +__kernel void +compiler_region(__global uint *src, __global uint *dst) +{ + __gen_ocl_force_simd16(); + int id = (int)get_global_id(0); + const int x0 = src[id]; + const int x1 = src[id+16]; + dst[id] = __gen_ocl_region(0, 16, 8, 2, x0, x1); +} + diff --git a/kernels/compiler_region0.cl b/kernels/compiler_region0.cl new file mode 100644 index 0000000..5bd57c0 --- /dev/null +++ b/kernels/compiler_region0.cl @@ -0,0 +1,11 @@ +__kernel void +compiler_region0(__global uint *src, __global uint *dst) +{ + __gen_ocl_force_simd16(); + int id = (int)get_global_id(0); + const int x0 = src[id]; + const int x1 = src[id+16]; + const int x2 = src[id+32]; + dst[id] = __gen_ocl_region(1, 16, 8, 2, x0, x1, x2); +} + diff --git a/kernels/compiler_region1.cl b/kernels/compiler_region1.cl new file mode 100644 index 0000000..9deb63c --- /dev/null +++ b/kernels/compiler_region1.cl @@ -0,0 +1,9 @@ +__kernel void +compiler_region1(__global uint *src, __global uint *dst) +{ + __gen_ocl_force_simd16(); + int id = (int)get_global_id(0); + const int x0 = src[id]; + dst[id] = __gen_ocl_region(0, 16, 8, 2, x0); +} + diff --git a/kernels/compiler_vote_all.cl b/kernels/compiler_vote_all.cl new file mode 100644 index 0000000..1918c1c --- /dev/null +++ b/kernels/compiler_vote_all.cl @@ -0,0 +1,10 @@ +__kernel void +compiler_vote_all(__global uint *src, __global uint *dst) +{ + int id = (int)get_global_id(0); + if (__gen_ocl_all(id > 8)) + dst[id] = src[id]; + else + dst[id] = 0; +} + diff --git a/kernels/compiler_vote_any.cl b/kernels/compiler_vote_any.cl new file mode 100644 index 0000000..0a81e89 --- /dev/null +++ b/kernels/compiler_vote_any.cl @@ -0,0 +1,10 @@ +__kernel void +compiler_vote_any(__global uint *src, __global uint *dst) +{ + int id = (int)get_global_id(0); + if (__gen_ocl_any(id > 6)) + dst[id] = src[id]; + else + dst[id] = 0; +} + diff --git a/kernels/compiler_write_only_bytes.cl b/kernels/compiler_write_only_bytes.cl index d17e54a..0bc0cd8 100644 --- 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 d5e2f26..bfd23cc 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 512e9e7..22464a5 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 . - * - * Author: Benjamin Segovia - */ - #include "utest_helper.hpp" struct hop { int x, y; }; diff --git a/utests/compiler_argument_structure_indirect.cpp b/utests/compiler_argument_structure_indirect.cpp index e572f81..a4584d5 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #include "utest_helper.hpp" struct hop { int x[16]; }; diff --git a/utests/compiler_array.cpp b/utests/compiler_array.cpp index 0cec7d8..8806c99 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #include "utest_helper.hpp" void compiler_array(void) diff --git a/utests/compiler_array0.cpp b/utests/compiler_array0.cpp index 9e3535d..7cf2bbb 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 . - * - * Author: Benjamin Segovia - */ - #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 0000000..fe1ecec --- /dev/null +++ b/utests/compiler_array1.cpp @@ -0,0 +1,52 @@ +#include "utest_helper.hpp" + +static void cpu(int global_id, int *src, int *dst) { + int final[16]; + for (int i = 0; i < 16; ++i) { + int array[16]; + for (int j = 0; j < src[0]; ++j) + array[j] = 1+src[0]; + for (int j = src[0]; j < 16; ++j) + array[j] = global_id; + final[i] = array[i]; + } + dst[global_id] = final[global_id]; +} + +void compiler_array1(void) +{ + const size_t n = 16; + int cpu_dst[16], cpu_src[16]; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_array1"); + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL); + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = 16; + locals[0] = 16; + + // Run random tests + for (uint32_t pass = 0; pass < 8; ++pass) { + OCL_MAP_BUFFER(0); + for (int32_t i = 0; i < (int32_t) n; ++i) + cpu_src[i] = ((int32_t*)buf_data[0])[i] = rand() % 16; + OCL_UNMAP_BUFFER(0); + + // Run the kernel on GPU + OCL_NDRANGE(1); + + // Run on CPU + for (int32_t i = 0; i <(int32_t) n; ++i) cpu(i, cpu_src, cpu_dst); + + // Compare + OCL_MAP_BUFFER(1); + for (int32_t i = 0; i < 11; ++i) + OCL_ASSERT(((int32_t*)buf_data[1])[i] == cpu_dst[i]); + OCL_UNMAP_BUFFER(1); + } +} + +MAKE_UTEST_FROM_FUNCTION(compiler_array1); + diff --git a/utests/compiler_array2.cpp b/utests/compiler_array2.cpp new file mode 100644 index 0000000..61ca9da --- /dev/null +++ b/utests/compiler_array2.cpp @@ -0,0 +1,50 @@ +#include "utest_helper.hpp" + +static void cpu(int global_id, int *src, int *dst) { + int final[16]; + int array[16]; + for (int j = 0; j < 16; ++j) array[j] = j; + for (int j = 0; j < 16; ++j) final[j] = j+1; + if (global_id == 15) + dst[global_id] = final[global_id]; + else + dst[global_id] = array[15 - global_id]; +} + +void compiler_array2(void) +{ + const size_t n = 16; + int cpu_dst[16], cpu_src[16]; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_array2"); + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL); + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = 16; + locals[0] = 16; + + // Run random tests + for (uint32_t pass = 0; pass < 8; ++pass) { + OCL_MAP_BUFFER(0); + for (int32_t i = 0; i < (int32_t) n; ++i) + cpu_src[i] = ((int32_t*)buf_data[0])[i] = rand() % 16; + OCL_UNMAP_BUFFER(0); + + // Run the kernel on GPU + OCL_NDRANGE(1); + + // Run on CPU + for (int32_t i = 0; i <(int32_t) n; ++i) cpu(i, cpu_src, cpu_dst); + + // Compare + OCL_MAP_BUFFER(1); + for (int32_t i = 0; i < 11; ++i) + OCL_ASSERT(((int32_t*)buf_data[1])[i] == cpu_dst[i]); + OCL_UNMAP_BUFFER(1); + } +} + +MAKE_UTEST_FROM_FUNCTION(compiler_array2); + diff --git a/utests/compiler_array3.cpp b/utests/compiler_array3.cpp new file mode 100644 index 0000000..865b1e5 --- /dev/null +++ b/utests/compiler_array3.cpp @@ -0,0 +1,51 @@ +#include "utest_helper.hpp" + +static void cpu(int global_id, int *src, int *dst) { + int tmp[32]; + for (int i = 0; i < 16; ++i) { + for (int j = 0; j < 16; ++j) + tmp[j] = global_id; + for (int j = 0; j < src[0]; ++j) + tmp[j] = 1+src[j]; + tmp[16+i] = tmp[i]; + } + dst[global_id] = tmp[16+global_id]; +} + +void compiler_array3(void) +{ + const size_t n = 16; + int cpu_dst[16], cpu_src[16]; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_array3"); + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL); + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = 16; + locals[0] = 16; + + // Run random tests + for (uint32_t pass = 0; pass < 8; ++pass) { + OCL_MAP_BUFFER(0); + for (int32_t i = 0; i < (int32_t) n; ++i) + cpu_src[i] = ((int32_t*)buf_data[0])[i] = rand() % 16; + OCL_UNMAP_BUFFER(0); + + // Run the kernel on GPU + OCL_NDRANGE(1); + + // Run on CPU + for (int32_t i = 0; i <(int32_t) n; ++i) cpu(i, cpu_src, cpu_dst); + + // Compare + OCL_MAP_BUFFER(1); + for (int32_t i = 0; i < 11; ++i) + OCL_ASSERT(((int32_t*)buf_data[1])[i] == cpu_dst[i]); + OCL_UNMAP_BUFFER(1); + } +} + +MAKE_UTEST_FROM_FUNCTION(compiler_array3); + diff --git a/utests/compiler_byte_scatter.cpp b/utests/compiler_byte_scatter.cpp index 115f5df..11300da 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #include "utest_helper.hpp" static void compiler_byte_scatter(void) diff --git a/utests/compiler_copy_buffer.cpp b/utests/compiler_copy_buffer.cpp index b005424..8066efe 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #include "utest_helper.hpp" static void compiler_copy_buffer(void) diff --git a/utests/compiler_copy_buffer_row.cpp b/utests/compiler_copy_buffer_row.cpp index cfaea0e..12c0592 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #include "utest_helper.hpp" static void compiler_copy_buffer_row(void) diff --git a/utests/compiler_function_argument.cpp b/utests/compiler_function_argument.cpp index 29775ce..a39523b 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #include "utest_helper.hpp" void compiler_function_argument(void) diff --git a/utests/compiler_function_argument0.cpp b/utests/compiler_function_argument0.cpp index 3aac9cb..2e4227e 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #include "utest_helper.hpp" void compiler_function_argument0(void) diff --git a/utests/compiler_function_argument1.cpp b/utests/compiler_function_argument1.cpp index dff46a8..48a7677 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #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 0000000..a877b58 --- /dev/null +++ b/utests/compiler_gather_register_file.cpp @@ -0,0 +1,31 @@ +#include "utest_helper.hpp" + +static void compiler_gather_register_file(void) +{ + const size_t n = 48; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_gather_register_file"); + buf_data[0] = (uint32_t*) malloc(sizeof(uint32_t) * n); + for (uint32_t i = 0; i < n; ++i) ((uint32_t*)buf_data[0])[i] = i; + OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(uint32_t), buf_data[0]); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL); + free(buf_data[0]); + buf_data[0] = NULL; + + // Run the kernel + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = 16; + locals[0] = 16; + OCL_NDRANGE(1); + + // Check result + OCL_MAP_BUFFER(0); + OCL_MAP_BUFFER(1); + for (uint32_t i = 0; i < 16; ++i) + OCL_ASSERT(((uint32_t*)buf_data[1])[i] == i); +} + +MAKE_UTEST_FROM_FUNCTION(compiler_gather_register_file); + diff --git a/utests/compiler_gather_register_file0.cpp b/utests/compiler_gather_register_file0.cpp new file mode 100644 index 0000000..f9e293e --- /dev/null +++ b/utests/compiler_gather_register_file0.cpp @@ -0,0 +1,32 @@ +#include "utest_helper.hpp" + +static void compiler_gather_register_file0(void) +{ + const size_t n = 48; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_gather_register_file0"); + buf_data[0] = (uint32_t*) malloc(sizeof(uint32_t) * n); + for (uint32_t i = 0; i < n; ++i) ((uint32_t*)buf_data[0])[i] = i; + OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(uint32_t), buf_data[0]); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL); + free(buf_data[0]); + buf_data[0] = NULL; + + // Run the kernel + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = 16; + locals[0] = 16; + OCL_NDRANGE(1); + + // Check result + OCL_MAP_BUFFER(0); + OCL_MAP_BUFFER(1); + for (uint32_t i = 0; i < 16; ++i) + OCL_ASSERT(((uint32_t*)buf_data[1])[i] == 15-i); +} + +MAKE_UTEST_FROM_FUNCTION(compiler_gather_register_file0); + + diff --git a/utests/compiler_gather_register_file1.cpp b/utests/compiler_gather_register_file1.cpp new file mode 100644 index 0000000..4956b7e --- /dev/null +++ b/utests/compiler_gather_register_file1.cpp @@ -0,0 +1,31 @@ +#include "utest_helper.hpp" + +static void compiler_gather_register_file1(void) +{ + const size_t n = 48; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_gather_register_file1"); + buf_data[0] = (uint32_t*) malloc(sizeof(uint32_t) * n); + for (uint32_t i = 0; i < n; ++i) ((uint32_t*)buf_data[0])[i] = i; + OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(uint32_t), buf_data[0]); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL); + free(buf_data[0]); + buf_data[0] = NULL; + + // Run the kernel + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = 16; + locals[0] = 16; + OCL_NDRANGE(1); + + // Check result + OCL_MAP_BUFFER(0); + OCL_MAP_BUFFER(1); + for (uint32_t i = 0; i < 16; ++i) + OCL_ASSERT(((uint32_t*)buf_data[1])[i] == 2*i); +} + +MAKE_UTEST_FROM_FUNCTION(compiler_gather_register_file1); + diff --git a/utests/compiler_if_else.cpp b/utests/compiler_if_else.cpp index 44b1a87..e38b23f 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #include "utest_helper.hpp" static void compiler_if_else(void) diff --git a/utests/compiler_local_slm.cpp b/utests/compiler_local_slm.cpp index f9c9e0c..aa9a2fe 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #include "utest_helper.hpp" void compiler_local_slm(void) diff --git a/utests/compiler_lower_return0.cpp b/utests/compiler_lower_return0.cpp index 47c7f68..0e9dbd0 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #include "utest_helper.hpp" static void compiler_lower_return0(void) diff --git a/utests/compiler_lower_return1.cpp b/utests/compiler_lower_return1.cpp index 0550cce..b4f1fe3 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #include "utest_helper.hpp" static void compiler_lower_return1(void) diff --git a/utests/compiler_lower_return2.cpp b/utests/compiler_lower_return2.cpp index bdc9ef6..1e34036 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #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 0000000..99831e2 --- /dev/null +++ b/utests/compiler_obread.cpp @@ -0,0 +1,31 @@ +#include "utest_helper.hpp" + +static void compiler_obread(void) +{ + const size_t n = 48; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_obread"); + buf_data[0] = (uint32_t*) malloc(sizeof(uint32_t) * n); + for (uint32_t i = 0; i < n; ++i) ((uint32_t*)buf_data[0])[i] = i; + OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(uint32_t), buf_data[0]); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL); + free(buf_data[0]); + buf_data[0] = NULL; + + // Run the kernel + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = 16; + locals[0] = 16; + OCL_NDRANGE(1); + + // Check result + OCL_MAP_BUFFER(0); + OCL_MAP_BUFFER(1); + for (uint32_t i = 0; i < 16; ++i) + OCL_ASSERT(((uint32_t*)buf_data[1])[i] == i); +} + +MAKE_UTEST_FROM_FUNCTION(compiler_obread); + diff --git a/utests/compiler_obwrite.cpp b/utests/compiler_obwrite.cpp new file mode 100644 index 0000000..5f6b63b --- /dev/null +++ b/utests/compiler_obwrite.cpp @@ -0,0 +1,31 @@ +#include "utest_helper.hpp" + +static void compiler_obwrite(void) +{ + const size_t n = 48; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_obwrite"); + buf_data[0] = (uint32_t*) malloc(sizeof(uint32_t) * n); + for (uint32_t i = 0; i < n; ++i) ((uint32_t*)buf_data[0])[i] = i; + OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(uint32_t), buf_data[0]); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL); + free(buf_data[0]); + buf_data[0] = NULL; + + // Run the kernel + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = 16; + locals[0] = 16; + OCL_NDRANGE(1); + + // Check result + OCL_MAP_BUFFER(0); + OCL_MAP_BUFFER(1); + for (uint32_t i = 0; i < 16; ++i) + OCL_ASSERT(((uint32_t*)buf_data[1])[i] == i); +} + +MAKE_UTEST_FROM_FUNCTION(compiler_obwrite); + diff --git a/utests/compiler_region.cpp b/utests/compiler_region.cpp new file mode 100644 index 0000000..395c227 --- /dev/null +++ b/utests/compiler_region.cpp @@ -0,0 +1,31 @@ +#include "utest_helper.hpp" + +static void compiler_region(void) +{ + const size_t n = 32; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_region"); + buf_data[0] = (uint32_t*) malloc(sizeof(uint32_t) * n); + for (uint32_t i = 0; i < n; ++i) ((uint32_t*)buf_data[0])[i] = i; + OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(uint32_t), buf_data[0]); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL); + free(buf_data[0]); + buf_data[0] = NULL; + + // Run the kernel + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = 16; + locals[0] = 16; + OCL_NDRANGE(1); + + // Check result + OCL_MAP_BUFFER(0); + OCL_MAP_BUFFER(1); + for (uint32_t i = 0; i < 16; ++i) + OCL_ASSERT(((uint32_t*)buf_data[1])[i] == 2*i); +} + +MAKE_UTEST_FROM_FUNCTION(compiler_region); + diff --git a/utests/compiler_region0.cpp b/utests/compiler_region0.cpp new file mode 100644 index 0000000..56ddb68 --- /dev/null +++ b/utests/compiler_region0.cpp @@ -0,0 +1,31 @@ +#include "utest_helper.hpp" + +static void compiler_region0(void) +{ + const size_t n = 48; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_region0"); + buf_data[0] = (uint32_t*) malloc(sizeof(uint32_t) * n); + for (uint32_t i = 0; i < n; ++i) ((uint32_t*)buf_data[0])[i] = i; + OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(uint32_t), buf_data[0]); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL); + free(buf_data[0]); + buf_data[0] = NULL; + + // Run the kernel + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = 16; + locals[0] = 16; + OCL_NDRANGE(1); + + // Check result + OCL_MAP_BUFFER(0); + OCL_MAP_BUFFER(1); + for (uint32_t i = 0; i < 16; ++i) + OCL_ASSERT(((uint32_t*)buf_data[1])[i] == 2*i+1); +} + +MAKE_UTEST_FROM_FUNCTION(compiler_region0); + diff --git a/utests/compiler_region1.cpp b/utests/compiler_region1.cpp new file mode 100644 index 0000000..9b03b8e --- /dev/null +++ b/utests/compiler_region1.cpp @@ -0,0 +1,33 @@ +#include "utest_helper.hpp" + +static void compiler_region1(void) +{ + const size_t n = 32; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_region1"); + buf_data[0] = (uint32_t*) malloc(sizeof(uint32_t) * n); + for (uint32_t i = 0; i < n; ++i) ((uint32_t*)buf_data[0])[i] = i; + OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(uint32_t), buf_data[0]); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL); + free(buf_data[0]); + buf_data[0] = NULL; + + // Run the kernel + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = 16; + locals[0] = 16; + OCL_NDRANGE(1); + + // Check result + OCL_MAP_BUFFER(0); + OCL_MAP_BUFFER(1); + for (uint32_t i = 0; i < 8; ++i) + OCL_ASSERT(((uint32_t*)buf_data[1])[i] == 2*i); + for (uint32_t i = 8; i < 16; ++i) + OCL_ASSERT(((uint32_t*)buf_data[1])[i] == 0); +} + +MAKE_UTEST_FROM_FUNCTION(compiler_region1); + diff --git a/utests/compiler_short_scatter.cpp b/utests/compiler_short_scatter.cpp index 8f5f4c2..1746744 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #include "utest_helper.hpp" static void compiler_short_scatter(void) diff --git a/utests/compiler_sub_bytes.cpp b/utests/compiler_sub_bytes.cpp index f2262f1..49a5261 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #include "utest_helper.hpp" static void compiler_sub_bytes(void) diff --git a/utests/compiler_sub_shorts.cpp b/utests/compiler_sub_shorts.cpp index bf4271f..4aeeca3 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #include "utest_helper.hpp" static void compiler_sub_shorts(void) diff --git a/utests/compiler_uint2_copy.cpp b/utests/compiler_uint2_copy.cpp index 386cfc0..8eb4314 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #include "utest_helper.hpp" static void compiler_uint2_copy(void) diff --git a/utests/compiler_uint3_copy.cpp b/utests/compiler_uint3_copy.cpp index 277d6ef..320f405 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #include "utest_helper.hpp" static void compiler_uint3_copy(void) diff --git a/utests/compiler_uint3_unaligned_copy.cpp b/utests/compiler_uint3_unaligned_copy.cpp index 9ccdb42..d42b4c3 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #include "utest_helper.hpp" static void compiler_uint3_unaligned_copy(void) diff --git a/utests/compiler_unstructured_branch0.cpp b/utests/compiler_unstructured_branch0.cpp index a314a51..128a53e 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #include "utest_helper.hpp" static void compiler_unstructured_branch0(void) diff --git a/utests/compiler_unstructured_branch1.cpp b/utests/compiler_unstructured_branch1.cpp index 5c3614c..6021f5b 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #include "utest_helper.hpp" static void compiler_unstructured_branch1(void) diff --git a/utests/compiler_unstructured_branch2.cpp b/utests/compiler_unstructured_branch2.cpp index 1808ef1..d61c6b5 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #include "utest_helper.hpp" static void compiler_unstructured_branch2(void) diff --git a/utests/compiler_unstructured_branch3.cpp b/utests/compiler_unstructured_branch3.cpp index 732bcf2..0c6992a 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #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 0000000..9fcffd9 --- /dev/null +++ b/utests/compiler_vote_all.cpp @@ -0,0 +1,33 @@ +#include "utest_helper.hpp" + +static void compiler_vote_all(void) +{ + const size_t n = 32; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_vote_all"); + buf_data[0] = (uint32_t*) malloc(sizeof(uint32_t) * n); + for (uint32_t i = 0; i < n; ++i) ((uint32_t*)buf_data[0])[i] = i; + OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(uint32_t), buf_data[0]); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL); + free(buf_data[0]); + buf_data[0] = NULL; + + // Run the kernel + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = 32; + locals[0] = 16; + OCL_NDRANGE(1); + + // Check result + OCL_MAP_BUFFER(0); + OCL_MAP_BUFFER(1); + for (uint32_t i = 0; i < 16; ++i) + OCL_ASSERT(((uint32_t*)buf_data[1])[i] == 0); + for (uint32_t i = 16; i < 32; ++i) + OCL_ASSERT(((uint32_t*)buf_data[1])[i] == i); +} + +MAKE_UTEST_FROM_FUNCTION(compiler_vote_all); + diff --git a/utests/compiler_vote_any.cpp b/utests/compiler_vote_any.cpp new file mode 100644 index 0000000..91bcb03 --- /dev/null +++ b/utests/compiler_vote_any.cpp @@ -0,0 +1,31 @@ +#include "utest_helper.hpp" + +static void compiler_vote_any(void) +{ + const size_t n = 32; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_vote_any"); + buf_data[0] = (uint32_t*) malloc(sizeof(uint32_t) * n); + for (uint32_t i = 0; i < n; ++i) ((uint32_t*)buf_data[0])[i] = i; + OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(uint32_t), buf_data[0]); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL); + free(buf_data[0]); + buf_data[0] = NULL; + + // Run the kernel + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = 32; + locals[0] = 16; + OCL_NDRANGE(1); + + // Check result + OCL_MAP_BUFFER(0); + OCL_MAP_BUFFER(1); + for (uint32_t i = 0; i < 32; ++i) + OCL_ASSERT(((uint32_t*)buf_data[1])[i] == i); +} + +MAKE_UTEST_FROM_FUNCTION(compiler_vote_any); + diff --git a/utests/compiler_write_only.cpp b/utests/compiler_write_only.cpp index b42bfd2..e04f855 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #include "utest_helper.hpp" void compiler_write_only(void) diff --git a/utests/compiler_write_only_bytes.cpp b/utests/compiler_write_only_bytes.cpp index 19a8ee1..1a13cdb 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #include "utest_helper.hpp" void compiler_write_only_bytes(void) diff --git a/utests/compiler_write_only_shorts.cpp b/utests/compiler_write_only_shorts.cpp index 4db5e11..19988fe 100644 --- 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 . - * - * Author: Benjamin Segovia - */ - #include "utest_helper.hpp" void compiler_write_only_shorts(void)