[Beignet] [PATCH V3 2/2] utests: Add basic arithmetic test case
Zhigang Gong
zhigang.gong at linux.intel.com
Wed Jun 26 01:28:20 PDT 2013
LGTM, pushed. Thanks.
On Wed, Jun 26, 2013 at 03:52:13PM +0800, Ruiling Song wrote:
> test case for + - * / % of data type (u)int8/16/32
> remove duplicated cases.
>
> Signed-off-by: Ruiling Song <ruiling.song at intel.com>
> ---
> kernels/compiler_basic_arithmetic.cl | 53 ++++++++++++++++
> kernels/compiler_sub_bytes.cl | 7 ---
> kernels/compiler_sub_shorts.cl | 7 ---
> utests/CMakeLists.txt | 3 +-
> utests/compiler_basic_arithmetic.cpp | 112 ++++++++++++++++++++++++++++++++++
> utests/compiler_sub_bytes.cpp | 35 -----------
> utests/compiler_sub_shorts.cpp | 36 -----------
> 7 files changed, 166 insertions(+), 87 deletions(-)
> create mode 100644 kernels/compiler_basic_arithmetic.cl
> delete mode 100644 kernels/compiler_sub_bytes.cl
> delete mode 100644 kernels/compiler_sub_shorts.cl
> create mode 100644 utests/compiler_basic_arithmetic.cpp
> delete mode 100644 utests/compiler_sub_bytes.cpp
> delete mode 100644 utests/compiler_sub_shorts.cpp
>
> diff --git a/kernels/compiler_basic_arithmetic.cl b/kernels/compiler_basic_arithmetic.cl
> new file mode 100644
> index 0000000..3e145d8
> --- /dev/null
> +++ b/kernels/compiler_basic_arithmetic.cl
> @@ -0,0 +1,53 @@
> +#define DECL_KERNEL_SUB(type)\
> +__kernel void \
> +compiler_sub_##type(__global type *src0, __global type *src1, __global type *dst) \
> +{ \
> + int id = (int)get_global_id(0); \
> + dst[id] = src0[id] - src1[id]; \
> +}
> +
> +#define DECL_KERNEL_ADD(type)\
> +__kernel void \
> +compiler_add_##type(__global type *src0, __global type *src1, __global type *dst) \
> +{ \
> + int id = (int)get_global_id(0); \
> + dst[id] = src0[id] + src1[id]; \
> +}
> +
> +#define DECL_KERNEL_MUL(type)\
> +__kernel void \
> +compiler_mul_##type(__global type *src0, __global type *src1, __global type *dst) \
> +{ \
> + int id = (int)get_global_id(0); \
> + dst[id] = src0[id] * src1[id]; \
> +}
> +
> +#define DECL_KERNEL_DIV(type)\
> +__kernel void \
> +compiler_div_##type(__global type *src0, __global type *src1, __global type *dst) \
> +{ \
> + int id = (int)get_global_id(0); \
> + dst[id] = src0[id] / src1[id]; \
> +}
> +
> +#define DECL_KERNEL_REM(type)\
> +__kernel void \
> +compiler_rem_##type(__global type *src0, __global type *src1, __global type *dst) \
> +{ \
> + int id = (int)get_global_id(0); \
> + dst[id] = src0[id] % src1[id]; \
> +}
> +
> +#define DECL_KERNEL_FOR_ALL_TYPE(op) \
> +DECL_KERNEL_##op(char) \
> +DECL_KERNEL_##op(uchar) \
> +DECL_KERNEL_##op(short) \
> +DECL_KERNEL_##op(ushort) \
> +DECL_KERNEL_##op(int) \
> +DECL_KERNEL_##op(uint)
> +
> +DECL_KERNEL_FOR_ALL_TYPE(SUB)
> +DECL_KERNEL_FOR_ALL_TYPE(ADD)
> +DECL_KERNEL_FOR_ALL_TYPE(MUL)
> +DECL_KERNEL_FOR_ALL_TYPE(DIV)
> +DECL_KERNEL_FOR_ALL_TYPE(REM)
> diff --git a/kernels/compiler_sub_bytes.cl b/kernels/compiler_sub_bytes.cl
> deleted file mode 100644
> index f058561..0000000
> --- a/kernels/compiler_sub_bytes.cl
> +++ /dev/null
> @@ -1,7 +0,0 @@
> -__kernel void
> -compiler_sub_bytes(__global char *src0, __global char *src1, __global char *dst)
> -{
> - int id = (int)get_global_id(0);
> - dst[id] = src0[id] - src1[id];
> -}
> -
> diff --git a/kernels/compiler_sub_shorts.cl b/kernels/compiler_sub_shorts.cl
> deleted file mode 100644
> index d26de7f..0000000
> --- a/kernels/compiler_sub_shorts.cl
> +++ /dev/null
> @@ -1,7 +0,0 @@
> -__kernel void
> -compiler_sub_shorts(__global short *src0, __global short *src1, __global short *dst)
> -{
> - int id = (int)get_global_id(0);
> - dst[id] = src0[id] - src1[id];
> -}
> -
> diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
> index df59feb..fa36277 100644
> --- a/utests/CMakeLists.txt
> +++ b/utests/CMakeLists.txt
> @@ -5,6 +5,7 @@ link_directories (${LLVM_LIBRARY_DIR})
> set (utests_sources
> cl_create_kernel.cpp
> utest_error.c
> + compiler_basic_arithmetic.cpp
> compiler_displacement_map_element.cpp
> compiler_shader_toy.cpp
> compiler_mandelbrot.cpp
> @@ -57,8 +58,6 @@ set (utests_sources
> compiler_saturate_sub.cpp
> compiler_shift_right.cpp
> compiler_short_scatter.cpp
> - compiler_sub_bytes.cpp
> - compiler_sub_shorts.cpp
> compiler_uint2_copy.cpp
> compiler_uint3_copy.cpp
> compiler_uint8_copy.cpp
> diff --git a/utests/compiler_basic_arithmetic.cpp b/utests/compiler_basic_arithmetic.cpp
> new file mode 100644
> index 0000000..dcdd084
> --- /dev/null
> +++ b/utests/compiler_basic_arithmetic.cpp
> @@ -0,0 +1,112 @@
> +#include "utest_helper.hpp"
> +
> +enum eTestOP {
> + TEST_OP_ADD =0,
> + TEST_OP_SUB,
> + TEST_OP_MUL,
> + TEST_OP_DIV,
> + TEST_OP_REM
> +};
> +
> +template <typename T, eTestOP op>
> +static void test_exec(const char* kernel_name)
> +{
> + const size_t n = 160;
> +
> + // Setup kernel and buffers
> + OCL_CREATE_KERNEL_FROM_FILE("compiler_basic_arithmetic", kernel_name);
> +std::cout <<"kernel name: " << kernel_name << std::endl;
> + buf_data[0] = (T*) malloc(sizeof(T) * n);
> + buf_data[1] = (T*) malloc(sizeof(T) * n);
> + for (uint32_t i = 0; i < n; ++i) ((T*)buf_data[0])[i] = (T) rand();
> + for (uint32_t i = 0; i < n; ++i) ((T*)buf_data[1])[i] = (T) rand();
> + if(op == TEST_OP_DIV || op == TEST_OP_REM) {
> + for (uint32_t i = 0; i < n; ++i) {
> + if(((T*)buf_data[1])[i] == 0)
> + ((T*)buf_data[1])[i] = (T) 1;
> + }
> + }
> + OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(T), buf_data[0]);
> + OCL_CREATE_BUFFER(buf[1], CL_MEM_COPY_HOST_PTR, n * sizeof(T), buf_data[1]);
> + OCL_CREATE_BUFFER(buf[2], 0, n * sizeof(T), NULL);
> +
> + // Run the kernel
> + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
> + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
> + OCL_SET_ARG(2, sizeof(cl_mem), &buf[2]);
> + globals[0] = n;
> + locals[0] = 16;
> + OCL_NDRANGE(1);
> +
> + // Check result
> + OCL_MAP_BUFFER(2);
> + if(op == TEST_OP_SUB) {
> + for (uint32_t i = 0; i < n; ++i)
> + OCL_ASSERT(((T*)buf_data[2])[i] == (T)(((T*)buf_data[0])[i] - ((T*)buf_data[1])[i]));
> + } else if(op == TEST_OP_ADD) {
> + for (uint32_t i = 0; i < n; ++i)
> + OCL_ASSERT(((T*)buf_data[2])[i] == (T)(((T*)buf_data[0])[i] + ((T*)buf_data[1])[i]));
> + } else if(op == TEST_OP_MUL) {
> + for (uint32_t i = 0; i < n; ++i)
> + OCL_ASSERT(((T*)buf_data[2])[i] == (T)(((T*)buf_data[0])[i] * ((T*)buf_data[1])[i]));
> + } else if(op == TEST_OP_DIV) {
> + for (uint32_t i = 0; i < n; ++i)
> + OCL_ASSERT(((T*)buf_data[2])[i] == (T)(((T*)buf_data[0])[i] / ((T*)buf_data[1])[i]));
> + } else {
> + for (uint32_t i = 0; i < n; ++i)
> + OCL_ASSERT(((T*)buf_data[2])[i] == (T)(((T*)buf_data[0])[i] % ((T*)buf_data[1])[i]));
> + }
> + free(buf_data[0]);
> + free(buf_data[1]);
> + buf_data[0] = buf_data[1] = NULL;
> +}
> +
> +#define DECL_TEST_SUB(type, alias) \
> +static void compiler_sub_ ##alias(void)\
> +{\
> + test_exec<type, TEST_OP_SUB>("compiler_sub_" # alias);\
> +}\
> +MAKE_UTEST_FROM_FUNCTION(compiler_sub_ ## alias)
> +
> +#define DECL_TEST_ADD(type, alias) \
> +static void compiler_add_ ##alias(void)\
> +{\
> + test_exec<type, TEST_OP_ADD>("compiler_add_" # alias);\
> +}\
> +MAKE_UTEST_FROM_FUNCTION(compiler_add_ ## alias)
> +
> +#define DECL_TEST_MUL(type, alias) \
> +static void compiler_mul_ ##alias(void)\
> +{\
> + test_exec<type, TEST_OP_MUL>("compiler_mul_" # alias);\
> +}\
> +MAKE_UTEST_FROM_FUNCTION(compiler_mul_ ## alias)
> +
> +#define DECL_TEST_DIV(type, alias) \
> +static void compiler_div_ ##alias(void)\
> +{\
> + test_exec<type, TEST_OP_DIV>("compiler_div_" # alias);\
> +}\
> +MAKE_UTEST_FROM_FUNCTION(compiler_div_ ## alias)
> +
> +#define DECL_TEST_REM(type, alias) \
> +static void compiler_rem_ ##alias(void)\
> +{\
> + test_exec<type, TEST_OP_REM>("compiler_rem_" # alias);\
> +}\
> +MAKE_UTEST_FROM_FUNCTION(compiler_rem_ ## alias)
> +
> +#define DECL_TEST_FOR_ALL_TYPE(op)\
> +DECL_TEST_##op(int8_t, char) \
> +DECL_TEST_##op(uint8_t, uchar) \
> +DECL_TEST_##op(int16_t, short) \
> +DECL_TEST_##op(uint16_t, ushort) \
> +DECL_TEST_##op(int32_t, int) \
> +DECL_TEST_##op(uint32_t, uint)
> +
> +DECL_TEST_FOR_ALL_TYPE(SUB)
> +DECL_TEST_FOR_ALL_TYPE(ADD)
> +DECL_TEST_FOR_ALL_TYPE(MUL)
> +DECL_TEST_FOR_ALL_TYPE(DIV)
> +DECL_TEST_FOR_ALL_TYPE(REM)
> +#undef DECL_TEST_FOR_ALL_TYPE
> diff --git a/utests/compiler_sub_bytes.cpp b/utests/compiler_sub_bytes.cpp
> deleted file mode 100644
> index 740a8fd..0000000
> --- a/utests/compiler_sub_bytes.cpp
> +++ /dev/null
> @@ -1,35 +0,0 @@
> -#include "utest_helper.hpp"
> -
> -static void compiler_sub_bytes(void)
> -{
> - const size_t n = 16;
> -
> - // Setup kernel and buffers
> - OCL_CREATE_KERNEL("compiler_sub_bytes");
> - buf_data[0] = (int8_t*) malloc(sizeof(int8_t) * n);
> - buf_data[1] = (int8_t*) malloc(sizeof(int8_t) * n);
> - for (uint32_t i = 0; i < n; ++i) ((int8_t*)buf_data[0])[i] = (int8_t) rand();
> - for (uint32_t i = 0; i < n; ++i) ((int8_t*)buf_data[1])[i] = (int8_t) rand();
> - OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(int8_t), buf_data[0]);
> - OCL_CREATE_BUFFER(buf[1], CL_MEM_COPY_HOST_PTR, n * sizeof(int8_t), buf_data[1]);
> - OCL_CREATE_BUFFER(buf[2], 0, n * sizeof(int8_t), NULL);
> -
> - // Run the kernel
> - OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
> - OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
> - OCL_SET_ARG(2, sizeof(cl_mem), &buf[2]);
> - globals[0] = n;
> - locals[0] = 16;
> - OCL_NDRANGE(1);
> -
> - // Check result
> - OCL_MAP_BUFFER(2);
> - for (uint32_t i = 0; i < n; ++i)
> - OCL_ASSERT(((int8_t*)buf_data[2])[i] == (int8_t)(((int8_t*)buf_data[0])[i] - ((int8_t*)buf_data[1])[i]));
> - free(buf_data[0]);
> - free(buf_data[1]);
> - buf_data[0] = buf_data[1] = NULL;
> -}
> -
> -MAKE_UTEST_FROM_FUNCTION(compiler_sub_bytes);
> -
> diff --git a/utests/compiler_sub_shorts.cpp b/utests/compiler_sub_shorts.cpp
> deleted file mode 100644
> index 7c24a56..0000000
> --- a/utests/compiler_sub_shorts.cpp
> +++ /dev/null
> @@ -1,36 +0,0 @@
> -#include "utest_helper.hpp"
> -
> -static void compiler_sub_shorts(void)
> -{
> - const size_t n = 16;
> -
> - // Setup kernel and buffers
> - OCL_CREATE_KERNEL("compiler_sub_shorts");
> - buf_data[0] = (int16_t*) malloc(sizeof(int16_t) * n);
> - buf_data[1] = (int16_t*) malloc(sizeof(int16_t) * n);
> - for (uint32_t i = 0; i < n; ++i) ((int16_t*)buf_data[0])[i] = (int16_t) rand();
> - for (uint32_t i = 0; i < n; ++i) ((int16_t*)buf_data[1])[i] = (int16_t) rand();
> - OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(int16_t), buf_data[0]);
> - OCL_CREATE_BUFFER(buf[1], CL_MEM_COPY_HOST_PTR, n * sizeof(int16_t), buf_data[1]);
> - OCL_CREATE_BUFFER(buf[2], 0, n * sizeof(int16_t), NULL);
> -
> - // Run the kernel
> - OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
> - OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
> - OCL_SET_ARG(2, sizeof(cl_mem), &buf[2]);
> - globals[0] = n;
> - locals[0] = 16;
> - OCL_NDRANGE(1);
> -
> - // Check result
> - OCL_MAP_BUFFER(2);
> - for (uint32_t i = 0; i < n; ++i)
> - OCL_ASSERT(((int16_t*)buf_data[2])[i] == (int16_t)(((int16_t*)buf_data[0])[i] - ((int16_t*)buf_data[1])[i]));
> - free(buf_data[0]);
> - free(buf_data[1]);
> - buf_data[0] = buf_data[1] = NULL;
> -}
> -
> -MAKE_UTEST_FROM_FUNCTION(compiler_sub_shorts);
> -
> -
> --
> 1.7.9.5
>
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet
More information about the Beignet
mailing list