From 82a3267d61de233e31aede5f044fee38de6a8f65 Mon Sep 17 00:00:00 2001 From: Homer Hsing Date: Fri, 1 Feb 2013 15:52:44 +0800 Subject: [PATCH] Also make "arithmetic shift right" work Only add a line of code ... Looks like Mr. Ben has forgotten that line ... Also add a test case. Signed-off-by: Homer Hsing Reviewed-by: Zhigang Gong --- kernels/compiler_arith_shift_right.cl | 4 ++++ utests/CMakeLists.txt | 1 + utests/compiler_arith_shift_right.cpp | 43 +++++++++++++++++++++++++++++++++++ 3 files changed, 48 insertions(+) create mode 100644 kernels/compiler_arith_shift_right.cl create mode 100644 utests/compiler_arith_shift_right.cpp diff --git a/kernels/compiler_arith_shift_right.cl b/kernels/compiler_arith_shift_right.cl new file mode 100644 index 0000000..03a4d8d --- /dev/null +++ b/kernels/compiler_arith_shift_right.cl @@ -0,0 +1,4 @@ +kernel void compiler_arith_shift_right(global int *src, global int *dst) { + int i = get_global_id(0); + dst[i] = src[i] >> 24; +} diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index c67cf34..50e0baa 100644 --- a/utests/CMakeLists.txt +++ b/utests/CMakeLists.txt @@ -12,6 +12,7 @@ ADD_LIBRARY(utests SHARED compiler_box_blur.cpp compiler_insert_to_constant.cpp compiler_argument_structure.cpp + compiler_arith_shift_right.cpp compiler_array0.cpp compiler_array.cpp compiler_array1.cpp diff --git a/utests/compiler_arith_shift_right.cpp b/utests/compiler_arith_shift_right.cpp new file mode 100644 index 0000000..6485571 --- /dev/null +++ b/utests/compiler_arith_shift_right.cpp @@ -0,0 +1,43 @@ +#include "utest_helper.hpp" + +static void cpu(int global_id, int *src, int *dst) { + dst[global_id] = src[global_id] >> 24; +} + +void compiler_arith_shift_right(void) +{ + const size_t n = 16; + int cpu_src[16]; + int cpu_dst[16]; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_arith_shift_right"); + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(int), NULL); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(int), 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] = ((int*)buf_data[0])[i] = 0x80000000 | rand(); + 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 < (int32_t) n; ++i) + OCL_ASSERT(((int *)buf_data[1])[i] == cpu_dst[i]); + OCL_UNMAP_BUFFER(1); + } +} + +MAKE_UTEST_FROM_FUNCTION(compiler_arith_shift_right); -- 2.7.4