[Beignet] [PATCH] utests: Add basic arithmetic test case
Song, Ruiling
ruiling.song at intel.com
Tue Jun 25 01:13:05 PDT 2013
I refined some macro definition in this version. Please ignore previous one.
-----Original Message-----
From: beignet-bounces+ruiling.song=intel.com at lists.freedesktop.org [mailto:beignet-bounces+ruiling.song=intel.com at lists.freedesktop.org] On Behalf Of Ruiling Song
Sent: Tuesday, June 25, 2013 4:11 PM
To: beignet at lists.freedesktop.org
Cc: Song, Ruiling
Subject: [Beignet] [PATCH] utests: Add basic arithmetic test case
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 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..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