[Beignet] [PATCH v2 2/3] add bswap64 in utest.

xionghu.luo at intel.com xionghu.luo at intel.com
Mon Sep 21 23:51:08 PDT 2015


From: Luo Xionghu <xionghu.luo at intel.com>

Signed-off-by: Luo Xionghu <xionghu.luo at intel.com>
---
 kernels/compiler_bswap.cl | 14 ++++++++++-
 utests/compiler_bswap.cpp | 63 ++++++++++++++++++++++++++++++++++++++++++++---
 2 files changed, 72 insertions(+), 5 deletions(-)

diff --git a/kernels/compiler_bswap.cl b/kernels/compiler_bswap.cl
index 3a0a373..b1432b2 100644
--- a/kernels/compiler_bswap.cl
+++ b/kernels/compiler_bswap.cl
@@ -1,5 +1,15 @@
+#define SWAP64(A)	  \
+((((A) & 0xff00000000000000) >> 56) | \
+    (((A) & 0x00ff000000000000) >> 40) | \
+    (((A) & 0x0000ff0000000000) >> 24) | \
+    (((A) & 0x000000ff00000000) >> 8) |  \
+    (((A) & 0x00000000ff000000) << 8) |  \
+    (((A) & 0x0000000000ff0000) << 24) | \
+    (((A) & 0x000000000000ff00) << 40) | \
+    (((A) & 0x00000000000000ff) << 56) )
+
 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) {
+    int src2, global int * dst2,  short src3, global short * dst3, global ulong* src4, global ulong* dst4, long src5, global long* dst5) {
   if (get_global_id(0) % 2 == 0) {
     dst0[get_global_id(0)] = __builtin_bswap32(src0[get_global_id(0)]);
   } else {
@@ -13,5 +23,7 @@ kernel void compiler_bswap(global uint * src0, global uint * dst0, global ushort
 
   dst2[get_global_id(0)] = __builtin_bswap32(src2);
   dst3[get_global_id(0)] = __builtin_bswap16(src3);
+  dst4[get_global_id(0)] = SWAP64(src4[get_global_id(0)]);
+  dst5[get_global_id(0)] = SWAP64(src5);
 }
 
diff --git a/utests/compiler_bswap.cpp b/utests/compiler_bswap.cpp
index 3af9ef5..ed22750 100644
--- a/utests/compiler_bswap.cpp
+++ b/utests/compiler_bswap.cpp
@@ -7,6 +7,14 @@
     (((uint32_t)(A) & 0x00ff0000) >> 8) | \
     (((uint32_t)(A) & 0x0000ff00) << 8) | \
     (((uint32_t)(A) & 0x000000ff) << 24))
+#define cpu_htonll(A)     ((((uint64_t)(A) & 0xff00000000000000) >> 56) | \
+    (((uint64_t)(A) & 0x00ff000000000000) >> 40) | \
+    (((uint64_t)(A) & 0x0000ff0000000000) >> 24) | \
+    (((uint64_t)(A) & 0x000000ff00000000) >> 8) |  \
+    (((uint64_t)(A) & 0x00000000ff000000) << 8) |  \
+    (((uint64_t)(A) & 0x0000000000ff0000) << 24) | \
+    (((uint64_t)(A) & 0x000000000000ff00) << 40) | \
+    (((uint64_t)(A) & 0x00000000000000ff) << 56) )
 
 
 template <typename T> static void gen_rand_val(T & val)
@@ -22,6 +30,8 @@ template <typename T> static void cpu(int global_id, T *src, T *dst)
     g = cpu_htons(f);
   else if (sizeof(T) == sizeof(int32_t))
     g = cpu_htonl(f);
+  else if (sizeof(T) == sizeof(int64_t))
+    g = cpu_htonll(f);
   dst[global_id] = g;
 }
 
@@ -33,15 +43,19 @@ template <typename T> static void cpu(int global_id, T src, T *dst)
     g = cpu_htons(f);
   else if (sizeof(T) == sizeof(int32_t))
     g = cpu_htonl(f);
+  else if (sizeof(T) == sizeof(int64_t))
+    g = cpu_htonll(f);
   dst[global_id] = g;
 }
 
 template <typename T> inline static void print_data(T& val)
 {
   if(sizeof(T) == sizeof(uint16_t))
-    printf(" 0x%hx", val);
-  else
-    printf(" 0x%x", val);
+    printf(" 0x%hx", (uint16_t)val);
+  else if(sizeof(T) == sizeof(uint32_t))
+    printf(" 0x%x", (uint32_t)val);
+  else if(sizeof(T) == sizeof(uint64_t))
+    printf(" 0x%lx", (uint64_t)val);
 }
 
 template <typename T> static void dump_data(T* raw, T* cpu, T* gpu, int n)
@@ -78,7 +92,7 @@ template <typename T> static void dump_data(T raw, T* cpu, T* gpu, int n)
 
 void compiler_bswap(void)
 {
-  const size_t n = 32;
+  const size_t n = 16;
   uint32_t src0[n];
   uint16_t src1[n];
   uint32_t dst0[n];
@@ -87,6 +101,10 @@ void compiler_bswap(void)
   int32_t dst2[n];
   int16_t src3 = static_cast<int16_t>(rand());
   int16_t dst3[n];
+  uint64_t src4[n];
+  uint64_t dst4[n];
+  int64_t src5 = static_cast<int64_t>(rand()) << 32| static_cast<int64_t>(rand());
+  int64_t dst5[n];
 
   // Setup kernel and buffers
   OCL_CREATE_KERNEL_FROM_FILE("compiler_bswap", "compiler_bswap");
@@ -108,6 +126,15 @@ void compiler_bswap(void)
   OCL_CREATE_BUFFER(buf[5], 0, sizeof(dst3), NULL);
   OCL_SET_ARG(7, sizeof(cl_mem), &buf[5]);
 
+  OCL_CREATE_BUFFER(buf[6], 0, sizeof(src4), NULL);
+  OCL_SET_ARG(8, sizeof(cl_mem), &buf[6]);
+  OCL_CREATE_BUFFER(buf[7], 0, sizeof(dst4), NULL);
+  OCL_SET_ARG(9, sizeof(cl_mem), &buf[7]);
+
+  OCL_SET_ARG(10, sizeof(int64_t), &src5);
+  OCL_CREATE_BUFFER(buf[8], 0, sizeof(dst5), NULL);
+  OCL_SET_ARG(11, sizeof(cl_mem), &buf[8]);
+
   OCL_MAP_BUFFER(0);
   for (int32_t i = 0; i < (int32_t) n; ++i) {
     gen_rand_val(src0[i]);
@@ -142,6 +169,16 @@ void compiler_bswap(void)
   memset(buf_data[5], 0, sizeof(dst3));
   OCL_UNMAP_BUFFER(5);
 
+  OCL_MAP_BUFFER(6);
+  for (int32_t i = 0; i < (int32_t) n; ++i) {
+    uint64_t x, y;
+    gen_rand_val(x);
+    gen_rand_val(y);
+    src4[i] = (x << 32)| y;
+  }
+  memcpy(buf_data[6], src4, sizeof(src4));
+  OCL_UNMAP_BUFFER(6);
+
   globals[0] = n;
   locals[0] = 16;
   OCL_NDRANGE(1);
@@ -173,6 +210,14 @@ void compiler_bswap(void)
   for (int32_t i = 0; i < (int32_t) n; ++i)
     cpu(i, src3, dst3);
 
+  // Run on CPU
+  for (int32_t i = 0; i < (int32_t) n; ++i)
+    cpu(i, src4, dst4);
+
+  // Run on CPU
+  for (int32_t i = 0; i < (int32_t) n; ++i)
+    cpu(i, src5, dst5);
+
   OCL_MAP_BUFFER(1);
   //dump_data(src0, dst0, (uint32_t *)buf_data[1], n);
   OCL_ASSERT(!memcmp(buf_data[1], dst0, sizeof(dst0)));
@@ -192,6 +237,16 @@ void compiler_bswap(void)
   //dump_data(src3, dst3, (int16_t *)buf_data[5], n);
   OCL_ASSERT(!memcmp(buf_data[5], dst3, sizeof(dst3)));
   OCL_UNMAP_BUFFER(5);
+
+  OCL_MAP_BUFFER(7);
+  //dump_data(src4, dst4, (uint64_t *)buf_data[7], n);
+  OCL_ASSERT(!memcmp(buf_data[7], dst4, sizeof(dst4)));
+  OCL_UNMAP_BUFFER(7);
+
+  OCL_MAP_BUFFER(8);
+  //dump_data(src5, dst5, (int64_t *)buf_data[8], n);
+  OCL_ASSERT(!memcmp(buf_data[8], dst5, sizeof(dst5)));
+  OCL_UNMAP_BUFFER(8);
 }
 
 MAKE_UTEST_FROM_FUNCTION(compiler_bswap);
-- 
1.9.1



More information about the Beignet mailing list