From db199abf207daad52b34fac600f3388472a9966c Mon Sep 17 00:00:00 2001 From: Homer Hsing Date: Fri, 21 Jun 2013 12:26:32 +0800 Subject: [PATCH] test cases for 64-bit float Signed-off-by: Homer Hsing Reviewed-by: Zhigang Gong Tested-by: Yang Rong --- kernels/compiler_double.cl | 9 +++++++++ kernels/compiler_double_2.cl | 9 +++++++++ kernels/compiler_double_3.cl | 6 ++++++ utests/CMakeLists.txt | 3 +++ utests/compiler_double.cpp | 46 +++++++++++++++++++++++++++++++++++++++++++ utests/compiler_double_2.cpp | 47 ++++++++++++++++++++++++++++++++++++++++++++ utests/compiler_double_3.cpp | 46 +++++++++++++++++++++++++++++++++++++++++++ 7 files changed, 166 insertions(+) create mode 100644 kernels/compiler_double.cl create mode 100644 kernels/compiler_double_2.cl create mode 100644 kernels/compiler_double_3.cl create mode 100644 utests/compiler_double.cpp create mode 100644 utests/compiler_double_2.cpp create mode 100644 utests/compiler_double_3.cpp diff --git a/kernels/compiler_double.cl b/kernels/compiler_double.cl new file mode 100644 index 0000000..a84f142 --- /dev/null +++ b/kernels/compiler_double.cl @@ -0,0 +1,9 @@ +#pragma OPENCL EXTENSION cl_khr_fp64 : enable +kernel void compiler_double(global double *src, global double *dst) { + int i = get_global_id(0); + double d = 1.234567890123456789; + if (i < 14) + dst[i] = d * (src[i] + d); + else + dst[i] = 14; +} diff --git a/kernels/compiler_double_2.cl b/kernels/compiler_double_2.cl new file mode 100644 index 0000000..20ee614 --- /dev/null +++ b/kernels/compiler_double_2.cl @@ -0,0 +1,9 @@ +#pragma OPENCL EXTENSION cl_khr_fp64 : enable +kernel void compiler_double_2(global float *src, global double *dst) { + int i = get_global_id(0); + float d = 1.234567890123456789f; + if (i < 14) + dst[i] = d * (d + src[i]); + else + dst[i] = 14; +} diff --git a/kernels/compiler_double_3.cl b/kernels/compiler_double_3.cl new file mode 100644 index 0000000..8b32404 --- /dev/null +++ b/kernels/compiler_double_3.cl @@ -0,0 +1,6 @@ +#pragma OPENCL EXTENSION cl_khr_fp64 : enable +kernel void compiler_double_3(global float *src, global double *dst) { + int i = get_global_id(0); + float d = 1.234567890123456789f; + dst[i] = i < 14 ? d : 14; +} diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index 108fa06..a913c2b 100644 --- a/utests/CMakeLists.txt +++ b/utests/CMakeLists.txt @@ -27,6 +27,9 @@ set (utests_sources compiler_copy_image.cpp compiler_copy_image_3d.cpp compiler_copy_buffer_row.cpp + compiler_double.cpp + compiler_double_2.cpp + compiler_double_3.cpp compiler_fabs.cpp compiler_fill_image.cpp compiler_fill_image0.cpp diff --git a/utests/compiler_double.cpp b/utests/compiler_double.cpp new file mode 100644 index 0000000..7c54ddf --- /dev/null +++ b/utests/compiler_double.cpp @@ -0,0 +1,46 @@ +#include +#include "utest_helper.hpp" + +static void cpu(int global_id, double *src, double *dst) { + double f = src[global_id]; + double d = 1.234567890123456789; + dst[global_id] = global_id < 14 ? (d * (f + d)) : 14; +} + +void compiler_double(void) +{ + const size_t n = 16; + double cpu_dst[n], cpu_src[n]; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_double"); + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(double), NULL); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(double), NULL); + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = n; + locals[0] = 16; + + // Run random tests + for (uint32_t pass = 0; pass < 1; ++pass) { + OCL_MAP_BUFFER(0); + for (int32_t i = 0; i < (int32_t) n; ++i) + cpu_src[i] = ((double*)buf_data[0])[i] = .1f * (rand() & 15) - .75f; + 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(fabs(((double*)buf_data[1])[i] - cpu_dst[i]) < 1e-4); + OCL_UNMAP_BUFFER(1); + } +} + +MAKE_UTEST_FROM_FUNCTION(compiler_double); diff --git a/utests/compiler_double_2.cpp b/utests/compiler_double_2.cpp new file mode 100644 index 0000000..7e3ae4b --- /dev/null +++ b/utests/compiler_double_2.cpp @@ -0,0 +1,47 @@ +#include +#include "utest_helper.hpp" + +static void cpu(int global_id, float *src, double *dst) { + float f = src[global_id]; + float d = 1.234567890123456789; + dst[global_id] = global_id < 14 ? d * (d + f) : 14; +} + +void compiler_double_2(void) +{ + const size_t n = 16; + float cpu_src[n]; + double cpu_dst[n]; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_double_2"); + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(float), NULL); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(double), NULL); + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = n; + locals[0] = 16; + + // Run random tests + for (uint32_t pass = 0; pass < 1; ++pass) { + OCL_MAP_BUFFER(0); + for (int32_t i = 0; i < (int32_t) n; ++i) + cpu_src[i] = ((float*)buf_data[0])[i] = .1f * (rand() & 15) - .75f; + 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(fabs(((double*)buf_data[1])[i] - cpu_dst[i]) < 1e-4); + OCL_UNMAP_BUFFER(1); + } +} + +MAKE_UTEST_FROM_FUNCTION(compiler_double_2); diff --git a/utests/compiler_double_3.cpp b/utests/compiler_double_3.cpp new file mode 100644 index 0000000..294950d --- /dev/null +++ b/utests/compiler_double_3.cpp @@ -0,0 +1,46 @@ +#include +#include "utest_helper.hpp" + +static void cpu(int global_id, float *src, double *dst) { + float d = 1.234567890123456789; + dst[global_id] = global_id < 14 ? d : 14; +} + +void compiler_double_3(void) +{ + const size_t n = 16; + float cpu_src[n]; + double cpu_dst[n]; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_double_3"); + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(float), NULL); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(double), NULL); + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = n; + locals[0] = 16; + + // Run random tests + for (uint32_t pass = 0; pass < 1; ++pass) { + OCL_MAP_BUFFER(0); + for (int32_t i = 0; i < (int32_t) n; ++i) + cpu_src[i] = ((float*)buf_data[0])[i] = .1f * (rand() & 15) - .75f; + 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(fabs(((double*)buf_data[1])[i] - cpu_dst[i]) < 1e-4); + OCL_UNMAP_BUFFER(1); + } +} + +MAKE_UTEST_FROM_FUNCTION(compiler_double_3); -- 2.7.4