[Beignet] [PATCH V2 2/2] utests: Add basic arithmetic test case

Ruiling Song ruiling.song at intel.com
Tue Jun 25 23:45:44 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 |   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



More information about the Beignet mailing list