[Beignet] [PATCH 2/2] test 64bit-integer shifting
Zhigang Gong
zhigang.gong at linux.intel.com
Mon Aug 12 02:10:10 PDT 2013
Could you modify your test case to add some branches?
We definitely need the unit test case to cover the
branch predication for the those 64bit data type.
On Mon, Aug 12, 2013 at 04:45:13PM +0800, Homer Hsing wrote:
> Signed-off-by: Homer Hsing <homer.xing at intel.com>
> ---
> kernels/compiler_long_asr.cl | 4 ++++
> kernels/compiler_long_shl.cl | 4 ++++
> kernels/compiler_long_shr.cl | 4 ++++
> utests/CMakeLists.txt | 3 +++
> utests/compiler_long_asr.cpp | 38 ++++++++++++++++++++++++++++++++++++++
> utests/compiler_long_shl.cpp | 38 ++++++++++++++++++++++++++++++++++++++
> utests/compiler_long_shr.cpp | 38 ++++++++++++++++++++++++++++++++++++++
> 7 files changed, 129 insertions(+)
> create mode 100644 kernels/compiler_long_asr.cl
> create mode 100644 kernels/compiler_long_shl.cl
> create mode 100644 kernels/compiler_long_shr.cl
> create mode 100644 utests/compiler_long_asr.cpp
> create mode 100644 utests/compiler_long_shl.cpp
> create mode 100644 utests/compiler_long_shr.cpp
>
> diff --git a/kernels/compiler_long_asr.cl b/kernels/compiler_long_asr.cl
> new file mode 100644
> index 0000000..ba3bdf1
> --- /dev/null
> +++ b/kernels/compiler_long_asr.cl
> @@ -0,0 +1,4 @@
> +kernel void compiler_long_asr(global long *src, global long *dst) {
> + int i = get_global_id(0);
> + dst[i] = src[i] >> i;
> +}
> diff --git a/kernels/compiler_long_shl.cl b/kernels/compiler_long_shl.cl
> new file mode 100644
> index 0000000..c197a62
> --- /dev/null
> +++ b/kernels/compiler_long_shl.cl
> @@ -0,0 +1,4 @@
> +kernel void compiler_long_shl(global long *src, global long *dst) {
> + int i = get_global_id(0);
> + dst[i] = src[i] << i;
> +}
> diff --git a/kernels/compiler_long_shr.cl b/kernels/compiler_long_shr.cl
> new file mode 100644
> index 0000000..ee989d6
> --- /dev/null
> +++ b/kernels/compiler_long_shr.cl
> @@ -0,0 +1,4 @@
> +kernel void compiler_long_shr(global ulong *src, global ulong *dst) {
> + int i = get_global_id(0);
> + dst[i] = src[i] >> i;
> +}
> diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
> index 8740cda..b3d039e 100644
> --- a/utests/CMakeLists.txt
> +++ b/utests/CMakeLists.txt
> @@ -126,6 +126,9 @@ set (utests_sources
> compiler_long.cpp
> compiler_long_2.cpp
> compiler_long_convert.cpp
> + compiler_long_shl.cpp
> + compiler_long_shr.cpp
> + compiler_long_asr.cpp
> utest_assert.cpp
> utest.cpp
> utest_file_map.cpp
> diff --git a/utests/compiler_long_asr.cpp b/utests/compiler_long_asr.cpp
> new file mode 100644
> index 0000000..a950aa7
> --- /dev/null
> +++ b/utests/compiler_long_asr.cpp
> @@ -0,0 +1,38 @@
> +#include <cstdint>
> +#include <cstring>
> +#include <iostream>
> +#include "utest_helper.hpp"
> +
> +void compiler_long_asr(void)
> +{
> + const size_t n = 64;
> + int64_t src[n];
> +
> + // Setup kernel and buffers
> + OCL_CREATE_KERNEL("compiler_long_asr");
> + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(int64_t), NULL);
> + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(int64_t), 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 (int32_t i = 0; i < (int32_t) n; ++i)
> + src[i] = (int64_t)1 << 63;
> + OCL_MAP_BUFFER(0);
> + memcpy(buf_data[0], src, sizeof(src));
> + OCL_UNMAP_BUFFER(0);
> +
> + // Run the kernel on GPU
> + OCL_NDRANGE(1);
> +
> + // Compare
> + OCL_MAP_BUFFER(1);
> + int64_t *dest = ((int64_t *)buf_data[1]);
> + for (int32_t i = 0; i < (int32_t) n; ++i)
> + OCL_ASSERT(dest[i] == src[i] >> i);
> + OCL_UNMAP_BUFFER(2);
> +}
> +
> +MAKE_UTEST_FROM_FUNCTION(compiler_long_asr);
> diff --git a/utests/compiler_long_shl.cpp b/utests/compiler_long_shl.cpp
> new file mode 100644
> index 0000000..d0d12da
> --- /dev/null
> +++ b/utests/compiler_long_shl.cpp
> @@ -0,0 +1,38 @@
> +#include <cstdint>
> +#include <cstring>
> +#include <iostream>
> +#include "utest_helper.hpp"
> +
> +void compiler_long_shl(void)
> +{
> + const size_t n = 64;
> + int64_t src[n];
> +
> + // Setup kernel and buffers
> + OCL_CREATE_KERNEL("compiler_long_shl");
> + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(int64_t), NULL);
> + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(int64_t), 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 (int32_t i = 0; i < (int32_t) n; ++i)
> + src[i] = 1;
> + OCL_MAP_BUFFER(0);
> + memcpy(buf_data[0], src, sizeof(src));
> + OCL_UNMAP_BUFFER(0);
> +
> + // Run the kernel on GPU
> + OCL_NDRANGE(1);
> +
> + // Compare
> + OCL_MAP_BUFFER(1);
> + int64_t *dest = ((int64_t *)buf_data[1]);
> + for (int32_t i = 0; i < (int32_t) n; ++i)
> + OCL_ASSERT(dest[i] == ((int64_t)1) << i);
> + OCL_UNMAP_BUFFER(2);
> +}
> +
> +MAKE_UTEST_FROM_FUNCTION(compiler_long_shl);
> diff --git a/utests/compiler_long_shr.cpp b/utests/compiler_long_shr.cpp
> new file mode 100644
> index 0000000..91333ed
> --- /dev/null
> +++ b/utests/compiler_long_shr.cpp
> @@ -0,0 +1,38 @@
> +#include <cstdint>
> +#include <cstring>
> +#include <iostream>
> +#include "utest_helper.hpp"
> +
> +void compiler_long_shr(void)
> +{
> + const size_t n = 64;
> + uint64_t src[n];
> +
> + // Setup kernel and buffers
> + OCL_CREATE_KERNEL("compiler_long_shr");
> + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint64_t), NULL);
> + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint64_t), 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 (int32_t i = 0; i < (int32_t) n; ++i)
> + src[i] = (uint64_t)1 << 63;
> + OCL_MAP_BUFFER(0);
> + memcpy(buf_data[0], src, sizeof(src));
> + OCL_UNMAP_BUFFER(0);
> +
> + // Run the kernel on GPU
> + OCL_NDRANGE(1);
> +
> + // Compare
> + OCL_MAP_BUFFER(1);
> + uint64_t *dest = ((uint64_t *)buf_data[1]);
> + for (int32_t i = 0; i < (int32_t) n; ++i)
> + OCL_ASSERT(dest[i] == src[i] >> i);
> + OCL_UNMAP_BUFFER(2);
> +}
> +
> +MAKE_UTEST_FROM_FUNCTION(compiler_long_shr);
> --
> 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