From 6c00747f00077af235ca082e45c00c1dbfe0bca1 Mon Sep 17 00:00:00 2001 From: Benjamin Segovia Date: Mon, 14 May 2012 18:24:39 +0000 Subject: [PATCH] Added tests for return instruction lowering --- kernels/compiler_array.cl | 12 ++++++ kernels/compiler_lower_return0.cl | 8 ++++ kernels/compiler_lower_return1.cl | 8 ++++ kernels/compiler_lower_return2.cl | 11 +++++ kernels/test_write_only.cl | 4 +- utests/CMakeLists.txt | 7 ++- utests/compiler_copy_buffer.cpp | 1 + utests/compiler_lower_return0.cpp | 73 ++++++++++++++++++++++++++++++++ utests/compiler_lower_return1.cpp | 66 +++++++++++++++++++++++++++++ utests/compiler_lower_return2.cpp | 67 +++++++++++++++++++++++++++++ utests/compiler_unstructured_branch0.cpp | 1 - 11 files changed, 253 insertions(+), 5 deletions(-) create mode 100644 kernels/compiler_array.cl create mode 100644 kernels/compiler_lower_return0.cl create mode 100644 kernels/compiler_lower_return1.cl create mode 100644 kernels/compiler_lower_return2.cl create mode 100644 utests/compiler_lower_return0.cpp create mode 100644 utests/compiler_lower_return1.cpp create mode 100644 utests/compiler_lower_return2.cpp diff --git a/kernels/compiler_array.cl b/kernels/compiler_array.cl new file mode 100644 index 0000000..08ec6db --- /dev/null +++ b/kernels/compiler_array.cl @@ -0,0 +1,12 @@ +__kernel void +compiler_array(__global int *src, __global int *dst, int x) +{ + if (x > 10) { + int array[x*256]; + array[get_local_id(0)] = get_global_id(0); + dst[get_global_id(0)] = array[get_local_id(1)]; + } else + dst[get_global_id(0)] = src[get_local_id(1)]; +} + + diff --git a/kernels/compiler_lower_return0.cl b/kernels/compiler_lower_return0.cl new file mode 100644 index 0000000..fd9846e --- /dev/null +++ b/kernels/compiler_lower_return0.cl @@ -0,0 +1,8 @@ +__kernel void +compiler_lower_return0(__global int *src, __global int *dst) { + const int id = get_global_id(0); + dst[id] = id; + if (src[id] > 0) return; + dst[id] = src[id]; +} + diff --git a/kernels/compiler_lower_return1.cl b/kernels/compiler_lower_return1.cl new file mode 100644 index 0000000..bcb6b7f --- /dev/null +++ b/kernels/compiler_lower_return1.cl @@ -0,0 +1,8 @@ +__kernel void +compiler_lower_return1(__global int *src, __global int *dst) { + const int id = get_global_id(0); + dst[id] = id; + if (id < 11 && (src[id] > 0 || src[id+16] < 2)) return; + dst[id] = src[id]; +} + diff --git a/kernels/compiler_lower_return2.cl b/kernels/compiler_lower_return2.cl new file mode 100644 index 0000000..9fa8ad6 --- /dev/null +++ b/kernels/compiler_lower_return2.cl @@ -0,0 +1,11 @@ +__kernel void +compiler_lower_return2(__global int *src, __global int *dst) { + const int id = get_global_id(0); + dst[id] = id; + while (dst[id] > src[id]) { + if (dst[id] > 10) return; + dst[id]--; + } + dst[id] += 2; +} + diff --git a/kernels/test_write_only.cl b/kernels/test_write_only.cl index 738749a..29fe6e3 100644 --- a/kernels/test_write_only.cl +++ b/kernels/test_write_only.cl @@ -1,7 +1,7 @@ __kernel void test_write_only(__global int *dst) { - int id = (int)get_global_id(0); - dst[id] = id; + int id = (int)get_global_id(0); + dst[id] = id; } diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index 43f8556..3345a1a 100644 --- a/utests/CMakeLists.txt +++ b/utests/CMakeLists.txt @@ -7,7 +7,7 @@ ADD_LIBRARY(utests SHARED utest_file_map.cpp utest_assert.cpp utest.cpp - app_mandelbrot.cpp +#app_mandelbrot.cpp compiler_write_only.cpp compiler_write_only_shorts.cpp compiler_write_only_bytes.cpp @@ -24,7 +24,10 @@ ADD_LIBRARY(utests SHARED compiler_unstructured_branch0.cpp compiler_unstructured_branch1.cpp compiler_unstructured_branch2.cpp - compiler_unstructured_branch3.cpp) + compiler_unstructured_branch3.cpp + compiler_lower_return0.cpp + compiler_lower_return1.cpp + compiler_lower_return2.cpp) TARGET_LINK_LIBRARIES(utests cl m) ADD_EXECUTABLE(run utest_run.cpp) diff --git a/utests/compiler_copy_buffer.cpp b/utests/compiler_copy_buffer.cpp index ee1dbf1..b005424 100644 --- a/utests/compiler_copy_buffer.cpp +++ b/utests/compiler_copy_buffer.cpp @@ -25,6 +25,7 @@ static void compiler_copy_buffer(void) // Setup kernel and buffers OCL_CREATE_KERNEL("test_copy_buffer"); + //OCL_CREATE_KERNEL("compiler_array"); buf_data[0] = (uint32_t*) malloc(sizeof(uint32_t) * n); for (uint32_t i = 0; i < n; ++i) ((uint32_t*)buf_data[0])[i] = i; OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(uint32_t), buf_data[0]); diff --git a/utests/compiler_lower_return0.cpp b/utests/compiler_lower_return0.cpp new file mode 100644 index 0000000..47c7f68 --- /dev/null +++ b/utests/compiler_lower_return0.cpp @@ -0,0 +1,73 @@ +/* + * Copyright © 2012 Intel Corporation + * + * This library is free software; you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public + * License as published by the Free Software Foundation; either + * version 2 of the License, or (at your option) any later version. + * + * This library is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with this library. If not, see . + * + * Author: Benjamin Segovia + */ + +#include "utest_helper.hpp" + +static void compiler_lower_return0(void) +{ + const size_t n = 32; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_lower_return0"); + buf_data[0] = (uint32_t*) malloc(sizeof(uint32_t) * n); + for (uint32_t i = 0; i < n; ++i) ((uint32_t*)buf_data[0])[i] = 2; + OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(uint32_t), buf_data[0]); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL); + free(buf_data[0]); + buf_data[0] = NULL; + + // Run the kernel + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = n; + locals[0] = 16; + OCL_NDRANGE(1); + + // First control flow + OCL_MAP_BUFFER(0); + OCL_MAP_BUFFER(1); + for (int32_t i = 0; i < 32; ++i) + OCL_ASSERT(((int32_t*)buf_data[1])[i] == i); + + // Second control flow + for (uint32_t i = 0; i < n; ++i) ((int32_t*)buf_data[0])[i] = -2; + OCL_UNMAP_BUFFER(0); + OCL_UNMAP_BUFFER(1); + OCL_NDRANGE(1); + OCL_MAP_BUFFER(0); + OCL_MAP_BUFFER(1); + for (uint32_t i = 0; i < 32; ++i) + OCL_ASSERT(((int32_t*)buf_data[1])[i] == -2); + + // Third control flow + for (uint32_t i = 0; i < 8; ++i) ((int32_t*)buf_data[0])[i] = 2; + OCL_UNMAP_BUFFER(0); + OCL_UNMAP_BUFFER(1); + OCL_NDRANGE(1); + OCL_MAP_BUFFER(0); + OCL_MAP_BUFFER(1); + for (int32_t i = 0; i < 8; ++i) + OCL_ASSERT(((int32_t*)buf_data[1])[i] == i); + for (int32_t i = 8; i < 32; ++i) + OCL_ASSERT(((int32_t*)buf_data[1])[i] == -2); +} + +MAKE_UTEST_FROM_FUNCTION(compiler_lower_return0); + + diff --git a/utests/compiler_lower_return1.cpp b/utests/compiler_lower_return1.cpp new file mode 100644 index 0000000..0550cce --- /dev/null +++ b/utests/compiler_lower_return1.cpp @@ -0,0 +1,66 @@ +/* + * Copyright © 2012 Intel Corporation + * + * This library is free software; you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public + * License as published by the Free Software Foundation; either + * version 2 of the License, or (at your option) any later version. + * + * This library is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with this library. If not, see . + * + * Author: Benjamin Segovia + */ + +#include "utest_helper.hpp" + +static void compiler_lower_return1(void) +{ + const size_t n = 32; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_lower_return1"); + buf_data[0] = (uint32_t*) malloc(sizeof(uint32_t) * n); + for (uint32_t i = 0; i < n; ++i) ((uint32_t*)buf_data[0])[i] = 2; + OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(uint32_t), buf_data[0]); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL); + free(buf_data[0]); + buf_data[0] = NULL; + + // Run the kernel + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = 16; + locals[0] = 16; + OCL_NDRANGE(1); + + // First control flow + OCL_MAP_BUFFER(0); + OCL_MAP_BUFFER(1); + for (int32_t i = 0; i < 11; ++i) + OCL_ASSERT(((int32_t*)buf_data[1])[i] == i); + for (int32_t i = 11; i < 16; ++i) + OCL_ASSERT(((int32_t*)buf_data[1])[i] == 2); + + // Second control flow + for (uint32_t i = 0; i < 4; ++i) ((int32_t*)buf_data[0])[i] = -2; + OCL_UNMAP_BUFFER(0); + OCL_UNMAP_BUFFER(1); + OCL_NDRANGE(1); + OCL_MAP_BUFFER(0); + OCL_MAP_BUFFER(1); + for (int32_t i = 0; i < 4; ++i) + OCL_ASSERT(((int32_t*)buf_data[1])[i] == -2); + for (int32_t i = 4; i < 11; ++i) + OCL_ASSERT(((int32_t*)buf_data[1])[i] == i); + for (int32_t i = 11; i < 16; ++i) + OCL_ASSERT(((int32_t*)buf_data[1])[i] == 2); +} + +MAKE_UTEST_FROM_FUNCTION(compiler_lower_return1); + diff --git a/utests/compiler_lower_return2.cpp b/utests/compiler_lower_return2.cpp new file mode 100644 index 0000000..bdc9ef6 --- /dev/null +++ b/utests/compiler_lower_return2.cpp @@ -0,0 +1,67 @@ +/* + * Copyright © 2012 Intel Corporation + * + * This library is free software; you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public + * License as published by the Free Software Foundation; either + * version 2 of the License, or (at your option) any later version. + * + * This library is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with this library. If not, see . + * + * Author: Benjamin Segovia + */ + +#include "utest_helper.hpp" + +static void cpu(int global_id, int *src, int *dst) { + const int id = global_id; + dst[id] = id; + while (dst[id] > src[id]) { + if (dst[id] > 10) return; + dst[id]--; + } + dst[id] += 2; +} + +static void compiler_lower_return2(void) +{ + const size_t n = 16; + int cpu_dst[16], cpu_src[16]; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_lower_return2"); + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL); + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = 16; + locals[0] = 16; + + for (uint32_t pass = 0; pass < 8; ++pass) { + OCL_MAP_BUFFER(0); + for (int32_t i = 0; i < (int32_t) n; ++i) + cpu_src[i] = ((int32_t*)buf_data[0])[i] = rand() % 16; + OCL_UNMAP_BUFFER(0); + + // Run the kernel on GPU + OCL_NDRANGE(1); + + // Run on CPU + for (int32_t i = 0; i <(int32_t) n; ++i) cpu(i, cpu_src, cpu_dst); + + // Compare + OCL_MAP_BUFFER(1); + for (int32_t i = 0; i < 11; ++i) + OCL_ASSERT(((int32_t*)buf_data[1])[i] == cpu_dst[i]); + OCL_UNMAP_BUFFER(1); + } +} + +MAKE_UTEST_FROM_FUNCTION(compiler_lower_return2); + diff --git a/utests/compiler_unstructured_branch0.cpp b/utests/compiler_unstructured_branch0.cpp index 12ebade..a314a51 100644 --- a/utests/compiler_unstructured_branch0.cpp +++ b/utests/compiler_unstructured_branch0.cpp @@ -72,4 +72,3 @@ static void compiler_unstructured_branch0(void) MAKE_UTEST_FROM_FUNCTION(compiler_unstructured_branch0); - -- 2.7.4