[Beignet] [PATCH V2 4/9] Utest: Add test case for half type vload\store

[Xiuli Pan]

From: Pan Xiuli <xiuli.pan at intel.com>

V2: Half program is different with normal program, reorder the test case
order.
Signed-off-by: Pan Xiuli <xiuli.pan at intel.com>

Signed-off-by: Pan Xiuli <xiuli.pan at intel.com>
---
 kernels/compiler_vector_load_store.cl |  8 +++++---
 utests/compiler_vector_load_store.cpp | 32 ++++++++++++++++++++++++++------
 2 files changed, 31 insertions(+), 9 deletions(-)

diff --git a/kernels/compiler_vector_load_store.cl b/kernels/compiler_vector_load_store.cl
index aec38b1..d423174 100644
--- a/kernels/compiler_vector_load_store.cl
+++ b/kernels/compiler_vector_load_store.cl
@@ -17,6 +17,7 @@ __kernel void test_##type ##n(__global type *pin, \
   vstore ##n(value, x, pout); \
 }
 
+#ifndef HALF
 #define TEST_ALL_TYPE(n) \
   TEST_TYPE(char,n)  \
   TEST_TYPE(uchar,n) \
@@ -27,11 +28,12 @@ __kernel void test_##type ##n(__global type *pin, \
   TEST_TYPE(float,n) \
   TEST_TYPE(long,n)  \
   TEST_TYPE(ulong,n)
-//  TEST_TYPE(double,n)
-
-#if 0
+#else
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#define TEST_ALL_TYPE(n) \
   TEST_TYPE(half,n)
 #endif
+//  TEST_TYPE(double,n)
 
 TEST_ALL_TYPE(2)
 TEST_ALL_TYPE(3)
diff --git a/utests/compiler_vector_load_store.cpp b/utests/compiler_vector_load_store.cpp
index 5a1a8d1..80e72a9 100644
--- a/utests/compiler_vector_load_store.cpp
+++ b/utests/compiler_vector_load_store.cpp
@@ -1,15 +1,27 @@
 #include "utest_helper.hpp"
 #include <string.h>
+#include <math.h>
 template<typename T>
 static void compiler_vector_load_store(int elemNum, const char *kernelName)
 {
   const size_t n = elemNum * 256;
+  if (strstr(kernelName, "half") != NULL)
+    if (!cl_check_half())
+      return;
 
   // Setup kernel and buffers
-  OCL_CREATE_KERNEL_FROM_FILE("compiler_vector_load_store", kernelName);
+  if (strstr(kernelName, "half") != NULL)
+    OCL_CALL(cl_kernel_init, "compiler_vector_load_store.cl", kernelName,
+                             SOURCE, "-DHALF");
+  else
+    OCL_CREATE_KERNEL_FROM_FILE("compiler_vector_load_store", kernelName);
   buf_data[0] = (T*) malloc(sizeof(T) * n);
-  for (uint32_t i = 0; i < n; ++i)
-    ((T*)buf_data[0])[i] = i;
+  for (uint32_t i = 0; i < n; ++i) {
+    if (strstr(kernelName, "half") != NULL)
+      ((T*)buf_data[0])[i] = __float_to_half(as_uint((float)i/(float)n));
+    else
+      ((T*)buf_data[0])[i] = i;
+  }
   OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(T), buf_data[0]);
   OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(T), NULL);
   free(buf_data[0]);
@@ -28,10 +40,17 @@ static void compiler_vector_load_store(int elemNum, const char *kernelName)
   for (uint32_t i = 0; i < n; ++i)
   {
     int shift = ((i % elemNum) + 1);
-    if (strstr(kernelName, "double") == NULL)
-      OCL_ASSERT(((T*)buf_data[1])[i] == (T)(((T*)buf_data[0])[i] + shift));
-    else
+    if (strstr(kernelName, "double") != NULL)
       OCL_ASSERT((((T*)buf_data[1])[i] - ((T)((T*)buf_data[0])[i] + shift)) < 1e-5);
+    else if (strstr(kernelName, "half") != NULL) {
+      float fdst = as_float(__half_to_float(((T*)buf_data[1])[i]));
+      float fsrc = as_float(__half_to_float((T)(((T*)buf_data[0])[i])));
+      fsrc += shift;
+      //printf("%d (%f, %f)\n",i, fdst, fsrc);
+      OCL_ASSERT((fabs(fsrc - fdst) <= 0.03 * fabs(fdst)));
+    }
+    else
+      OCL_ASSERT(((T*)buf_data[1])[i] == (T)(((T*)buf_data[0])[i] + shift));
   }
   OCL_UNMAP_BUFFER(0);
   OCL_UNMAP_BUFFER(1);
@@ -61,3 +80,4 @@ test_all_vector(float, float, true)
 //test_all_vector(double, double, true)
 test_all_vector(int64_t, long, true)
 test_all_vector(uint64_t, ulong, false)
+test_all_vector(uint16_t, half, false)
-- 
2.7.4