[Beignet] [PATCH 2/2] utests: Add basic arithmetic test case
Ruiling Song
ruiling.song at intel.com
Tue Jun 25 00:38:49 PDT 2013
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 | 73 +++++++++++++++++++
kernels/compiler_sub_bytes.cl | 7 --
kernels/compiler_sub_shorts.cl | 7 --
utests/CMakeLists.txt | 3 +-
utests/compiler_basic_arithmetic.cpp | 132 ++++++++++++++++++++++++++++++++++
utests/compiler_sub_bytes.cpp | 35 ---------
utests/compiler_sub_shorts.cpp | 36 ----------
7 files changed, 206 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..2bc2c27
--- /dev/null
+++ b/kernels/compiler_basic_arithmetic.cl
@@ -0,0 +1,73 @@
+#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]; \
+}
+DECL_KERNEL_SUB(char)
+DECL_KERNEL_SUB(uchar)
+DECL_KERNEL_SUB(short)
+DECL_KERNEL_SUB(ushort)
+DECL_KERNEL_SUB(int)
+DECL_KERNEL_SUB(uint)
+
+DECL_KERNEL_ADD(char)
+DECL_KERNEL_ADD(uchar)
+DECL_KERNEL_ADD(short)
+DECL_KERNEL_ADD(ushort)
+DECL_KERNEL_ADD(int)
+DECL_KERNEL_ADD(uint)
+
+DECL_KERNEL_MUL(char)
+DECL_KERNEL_MUL(uchar)
+DECL_KERNEL_MUL(short)
+DECL_KERNEL_MUL(ushort)
+DECL_KERNEL_MUL(int)
+DECL_KERNEL_MUL(uint)
+
+DECL_KERNEL_DIV(char)
+DECL_KERNEL_DIV(uchar)
+DECL_KERNEL_DIV(short)
+DECL_KERNEL_DIV(ushort)
+DECL_KERNEL_DIV(int)
+DECL_KERNEL_DIV(uint)
+
+DECL_KERNEL_REM(char)
+DECL_KERNEL_REM(uchar)
+DECL_KERNEL_REM(short)
+DECL_KERNEL_REM(ushort)
+DECL_KERNEL_REM(int)
+DECL_KERNEL_REM(uint)
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 c009d99..d6bbebe 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
@@ -55,8 +56,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..5ab5f44
--- /dev/null
+++ b/utests/compiler_basic_arithmetic.cpp
@@ -0,0 +1,132 @@
+#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)
+
+DECL_TEST_SUB(int8_t, char);
+DECL_TEST_SUB(uint8_t, uchar);
+DECL_TEST_SUB(int16_t, short);
+DECL_TEST_SUB(uint16_t, ushort);
+DECL_TEST_SUB(int32_t, int);
+DECL_TEST_SUB(uint32_t, uint);
+
+DECL_TEST_ADD(int8_t, char);
+DECL_TEST_ADD(uint8_t, uchar);
+DECL_TEST_ADD(int16_t, short);
+DECL_TEST_ADD(uint16_t, ushort);
+DECL_TEST_ADD(int32_t, int);
+DECL_TEST_ADD(uint32_t, uint);
+
+DECL_TEST_MUL(int8_t, char);
+DECL_TEST_MUL(uint8_t, uchar);
+DECL_TEST_MUL(int16_t, short);
+DECL_TEST_MUL(uint16_t, ushort);
+DECL_TEST_MUL(int32_t, int);
+DECL_TEST_MUL(uint32_t, uint);
+
+DECL_TEST_DIV(int8_t, char);
+DECL_TEST_DIV(uint8_t, uchar);
+DECL_TEST_DIV(int16_t, short);
+DECL_TEST_DIV(uint16_t, ushort);
+DECL_TEST_DIV(int32_t, int);
+DECL_TEST_DIV(uint32_t, uint);
+
+DECL_TEST_REM(int8_t, char);
+DECL_TEST_REM(uint8_t, uchar);
+DECL_TEST_REM(int16_t, short);
+DECL_TEST_REM(uint16_t, ushort);
+DECL_TEST_REM(int32_t, int);
+DECL_TEST_REM(uint32_t, uint);
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
More information about the Beignet
mailing list