[Beignet] [PATCH 2/2] test cases for 64-bit float
Zhigang Gong
zhigang.gong at linux.intel.com
Thu Jun 20 23:30:26 PDT 2013
Pushed. Thanks for the contribution.
On Fri, Jun 21, 2013 at 12:26:32PM +0800, Homer Hsing wrote:
>
> Signed-off-by: Homer Hsing <homer.xing at intel.com>
> ---
> kernels/compiler_double.cl | 10 ++++++++++
> kernels/compiler_double_2.cl | 10 ++++++++++
> kernels/compiler_double_3.cl | 7 +++++++
> utests/CMakeLists.txt | 3 +++
> utests/compiler_double.cpp | 46 +++++++++++++++++++++++++++++++++++++++++++
> utests/compiler_double_2.cpp | 47 ++++++++++++++++++++++++++++++++++++++++++++
> utests/compiler_double_3.cpp | 46 +++++++++++++++++++++++++++++++++++++++++++
> 7 files changed, 169 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..0d92479
> --- /dev/null
> +++ b/kernels/compiler_double.cl
> @@ -0,0 +1,10 @@
> +#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..706ef69
> --- /dev/null
> +++ b/kernels/compiler_double_2.cl
> @@ -0,0 +1,10 @@
> +#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..562953a
> --- /dev/null
> +++ b/kernels/compiler_double_3.cl
> @@ -0,0 +1,7 @@
> +#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 <cmath>
> +#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 <cmath>
> +#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 <cmath>
> +#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);
> --
> 1.8.1.2
>
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet
More information about the Beignet
mailing list