[Beignet] [V2 PATCH 7/7] utest: Update the test case for bswap.
junyan.he at inbox.com
junyan.he at inbox.com
Thu Mar 5 23:24:13 PST 2015
From: Junyan He <junyan.he at linux.intel.com>
Signed-off-by: Junyan He <junyan.he at linux.intel.com>
---
kernels/compiler_bswap.cl | 24 +++---
utests/compiler_bswap.cpp | 203 +++++++++++++++++++++++++++++++--------------
2 files changed, 156 insertions(+), 71 deletions(-)
diff --git a/kernels/compiler_bswap.cl b/kernels/compiler_bswap.cl
index 97313b1..3a0a373 100644
--- a/kernels/compiler_bswap.cl
+++ b/kernels/compiler_bswap.cl
@@ -1,13 +1,17 @@
-#define TEST_TYPE(TYPE, LENGTH) \
-kernel void compiler_bswap_##TYPE(global TYPE * src, global TYPE * dst){ \
- dst[get_global_id(0)]= __builtin_bswap##LENGTH(src[get_global_id(0)]); \
- dst[get_global_id(0)]= __builtin_bswap##LENGTH(dst[get_global_id(0)] -1 ); \
-}
+kernel void compiler_bswap(global uint * src0, global uint * dst0, global ushort * src1, global ushort * dst1,
+ int src2, global int * dst2, short src3, global short * dst3) {
+ if (get_global_id(0) % 2 == 0) {
+ dst0[get_global_id(0)] = __builtin_bswap32(src0[get_global_id(0)]);
+ } else {
+ dst0[get_global_id(0)] = src0[get_global_id(0)];
+ }
+ dst1[get_global_id(0)] = __builtin_bswap16(src1[get_global_id(0)]);
+ if (get_global_id(0) % 2 == 1) {
+ dst1[get_global_id(0)] = __builtin_bswap16(dst1[get_global_id(0)] + 1);
+ }
-TEST_TYPE(short, 16)
-TEST_TYPE(ushort, 16)
-TEST_TYPE(int, 32)
-TEST_TYPE(uint, 32)
+ dst2[get_global_id(0)] = __builtin_bswap32(src2);
+ dst3[get_global_id(0)] = __builtin_bswap16(src3);
+}
-#undef TEST_TYPE
diff --git a/utests/compiler_bswap.cpp b/utests/compiler_bswap.cpp
index 9475b99..3af9ef5 100644
--- a/utests/compiler_bswap.cpp
+++ b/utests/compiler_bswap.cpp
@@ -1,7 +1,6 @@
#include "utest_helper.hpp"
#include "string.h"
-namespace {
#define cpu_htons(A) ((((uint16_t)(A) & 0xff00) >> 8) | \
(((uint16_t)(A) & 0x00ff) << 8))
#define cpu_htonl(A) ((((uint32_t)(A) & 0xff000000) >> 24) | \
@@ -9,108 +8,190 @@ namespace {
(((uint32_t)(A) & 0x0000ff00) << 8) | \
(((uint32_t)(A) & 0x000000ff) << 24))
+
+template <typename T> static void gen_rand_val(T & val)
+{
+ val = static_cast<T>(rand());//(0xAABBCCDD);//
+}
+
template <typename T> static void cpu(int global_id, T *src, T *dst)
{
- T f = src[global_id];
- T g = 0;
- if(sizeof(T) == sizeof(int16_t))
- g = cpu_htons(f);
- else if(sizeof(T) == sizeof(int32_t))
- g = cpu_htonl(f);
- dst[global_id] = g;
+ T f = src[global_id];
+ T g = 0;
+ if (sizeof(T) == sizeof(int16_t))
+ g = cpu_htons(f);
+ else if (sizeof(T) == sizeof(int32_t))
+ g = cpu_htonl(f);
+ dst[global_id] = g;
}
-template <typename T> static void gen_rand_val (T & val)
+template <typename T> static void cpu(int global_id, T src, T *dst)
{
- val = static_cast<T>(rand() );
+ T f = src;
+ T g = 0;
+ if (sizeof(T) == sizeof(int16_t))
+ g = cpu_htons(f);
+ else if (sizeof(T) == sizeof(int32_t))
+ g = cpu_htonl(f);
+ dst[global_id] = g;
}
-template <typename T>
-inline static void print_data (T& val)
+template <typename T> inline static void print_data(T& val)
{
- if(sizeof(T) == sizeof(uint16_t))
- printf(" %hx", val);
- else
- printf(" %x", val);
+ if(sizeof(T) == sizeof(uint16_t))
+ printf(" 0x%hx", val);
+ else
+ printf(" 0x%x", val);
}
-template <typename T> static void dump_data (T* src, T* dst, int n)
+template <typename T> static void dump_data(T* raw, T* cpu, T* gpu, int n)
{
- printf("\nRaw: \n");
- for (int32_t i = 0; i < (int32_t) n; ++i) {
- print_data(((T *)buf_data[0])[i]);
- }
+ printf("\nRaw: \n");
+ for (int32_t i = 0; i < (int32_t) n; ++i) {
+ print_data(raw[i]);
+ }
- printf("\nCPU: \n");
- for (int32_t i = 0; i < (int32_t) n; ++i) {
- print_data(dst[i]);
- }
- printf("\nGPU: \n");
- for (int32_t i = 0; i < (int32_t) n; ++i) {
- print_data(((T *)buf_data[1])[i]);
- }
+ printf("\nCPU: \n");
+ for (int32_t i = 0; i < (int32_t) n; ++i) {
+ print_data(cpu[i]);
+ }
+ printf("\nGPU: \n");
+ for (int32_t i = 0; i < (int32_t) n; ++i) {
+ print_data(gpu[i]);
+ }
}
-template<typename T>
-void test(const char *kernel_name)
+template <typename T> static void dump_data(T raw, T* cpu, T* gpu, int n)
{
- const size_t n = 64;
- T cpu_dst[n];
- T cpu_src[n];
+ printf("\nRaw: \n");
+ print_data(raw);
+
+ printf("\nCPU: \n");
+ for (int32_t i = 0; i < (int32_t) n; ++i) {
+ print_data(cpu[i]);
+ }
+ printf("\nGPU: \n");
+ for (int32_t i = 0; i < (int32_t) n; ++i) {
+ print_data(gpu[i]);
+ }
+}
+
+void compiler_bswap(void)
+{
+ const size_t n = 32;
+ uint32_t src0[n];
+ uint16_t src1[n];
+ uint32_t dst0[n];
+ uint16_t dst1[n];
+ int32_t src2 = static_cast<int32_t>(rand());
+ int32_t dst2[n];
+ int16_t src3 = static_cast<int16_t>(rand());
+ int16_t dst3[n];
// Setup kernel and buffers
- OCL_CREATE_KERNEL_FROM_FILE("compiler_bswap", kernel_name);
- OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(T), NULL);
- OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(T), NULL);
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_bswap", "compiler_bswap");
+ OCL_CREATE_BUFFER(buf[0], 0, sizeof(src0), NULL);
OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+ OCL_CREATE_BUFFER(buf[1], 0, sizeof(dst0), NULL);
OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+ OCL_CREATE_BUFFER(buf[2], 0, sizeof(src1), NULL);
+ OCL_SET_ARG(2, sizeof(cl_mem), &buf[2]);
+ OCL_CREATE_BUFFER(buf[3], 0, sizeof(dst1), NULL);
+ OCL_SET_ARG(3, sizeof(cl_mem), &buf[3]);
+
+ OCL_SET_ARG(4, sizeof(int32_t), &src2);
+ OCL_CREATE_BUFFER(buf[4], 0, sizeof(dst2), NULL);
+ OCL_SET_ARG(5, sizeof(cl_mem), &buf[4]);
+
+ OCL_SET_ARG(6, sizeof(int16_t), &src3);
+ OCL_CREATE_BUFFER(buf[5], 0, sizeof(dst3), NULL);
+ OCL_SET_ARG(7, sizeof(cl_mem), &buf[5]);
+
OCL_MAP_BUFFER(0);
for (int32_t i = 0; i < (int32_t) n; ++i) {
- gen_rand_val(cpu_src[i]);
+ gen_rand_val(src0[i]);
}
-
- memcpy(buf_data[0], cpu_src, sizeof(T) * n);
+ memcpy(buf_data[0], src0, sizeof(src0));
+ OCL_UNMAP_BUFFER(0);
/* Clear the dst buffer to avoid random data. */
OCL_MAP_BUFFER(1);
- memset(buf_data[1], 0, sizeof(T) * n);
+ memset(buf_data[1], 0, sizeof(dst0));
OCL_UNMAP_BUFFER(1);
+ OCL_MAP_BUFFER(2);
+ for (int32_t i = 0; i < (int32_t) n; ++i) {
+ gen_rand_val(src1[i]);
+ }
+ memcpy(buf_data[2], src1, sizeof(src1));
+ OCL_UNMAP_BUFFER(2);
+
+ /* Clear the dst buffer to avoid random data. */
+ OCL_MAP_BUFFER(3);
+ memset(buf_data[3], 0, sizeof(dst1));
+ OCL_UNMAP_BUFFER(3);
+
+ /* Clear the dst buffer to avoid random data. */
+ OCL_MAP_BUFFER(4);
+ memset(buf_data[4], 0, sizeof(dst2));
+ OCL_UNMAP_BUFFER(4);
+
+ /* Clear the dst buffer to avoid random data. */
+ OCL_MAP_BUFFER(5);
+ memset(buf_data[5], 0, sizeof(dst3));
+ OCL_UNMAP_BUFFER(5);
+
globals[0] = n;
locals[0] = 16;
OCL_NDRANGE(1);
// Run on CPU
- for (int32_t i = 0; i < (int32_t) n; ++i)
- cpu(i, cpu_src, cpu_dst);
+ for (int32_t i = 0; i < (int32_t) n; ++i) {
+ if (i%2) {
+ dst0[i] = src0[i];
+ continue;
+ }
+ cpu(i, src0, dst0);
+ }
+
+ // Run on CPU
+ for (int32_t i = 0; i < (int32_t) n; ++i) {
+ cpu(i, src1, dst1);
+
+ if (i%2) {
+ dst1[i] = dst1[i] + 1;
+ cpu(i, dst1, dst1);
+ }
+ }
+ // Run on CPU
for (int32_t i = 0; i < (int32_t) n; ++i)
- cpu_dst[i] = cpu_dst[i] -1;
+ cpu(i, src2, dst2);
// Run on CPU
for (int32_t i = 0; i < (int32_t) n; ++i)
- cpu(i, cpu_dst, cpu_dst);
+ cpu(i, src3, dst3);
OCL_MAP_BUFFER(1);
- // dump_data(cpu_src, cpu_dst, n);
+ //dump_data(src0, dst0, (uint32_t *)buf_data[1], n);
+ OCL_ASSERT(!memcmp(buf_data[1], dst0, sizeof(dst0)));
+ OCL_UNMAP_BUFFER(1);
- OCL_ASSERT(!memcmp(buf_data[1], cpu_dst, sizeof(T) * n));
+ OCL_MAP_BUFFER(3);
+ //dump_data(src1, dst1, (uint16_t *)buf_data[3], n);
+ OCL_ASSERT(!memcmp(buf_data[3], dst1, sizeof(dst1)));
+ OCL_UNMAP_BUFFER(3);
- OCL_UNMAP_BUFFER(1);
- OCL_UNMAP_BUFFER(0);
-}
+ OCL_MAP_BUFFER(4);
+ //dump_data(src2, dst2, (int32_t *)buf_data[4], n);
+ OCL_ASSERT(!memcmp(buf_data[4], dst2, sizeof(dst2)));
+ OCL_UNMAP_BUFFER(4);
+ OCL_MAP_BUFFER(5);
+ //dump_data(src3, dst3, (int16_t *)buf_data[5], n);
+ OCL_ASSERT(!memcmp(buf_data[5], dst3, sizeof(dst3)));
+ OCL_UNMAP_BUFFER(5);
}
-#define compiler_bswap(type, kernel) \
-static void compiler_bswap_ ##type(void)\
-{\
- test<type>(# kernel);\
-}\
-MAKE_UTEST_FROM_FUNCTION(compiler_bswap_ ## type);
-
-compiler_bswap(int16_t, compiler_bswap_short)
-compiler_bswap(uint16_t, compiler_bswap_ushort)
-compiler_bswap(int32_t, compiler_bswap_int)
-compiler_bswap(uint32_t, compiler_bswap_uint)
+MAKE_UTEST_FROM_FUNCTION(compiler_bswap);
--
1.7.9.5
More information about the Beignet
mailing list