utest: Add test cases for half.

Signed-off-by: Junyan He <junyan.he@linux.intel.com>
Reviewed-by: "Yang, Rong R" <rong.r.yang@intel.com>

4 files changed, 105 insertions, 0 deletions

diff --git a/kernels/compiler_half.cl b/kernels/compiler_half.cl
new file mode 100644
index 00000000..dc227664
--- /dev/null
+++ b/kernels/compiler_half.cl
@@ -0,0 +1,11 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+kernel void compiler_half_basic(global half *src, global half *dst) {
+  int i = get_global_id(0);
+  half hf = 2.5;
+  half val = src[i];
+  val = val + hf;
+  val = val*val;
+  val = val/(half)1.8;
+  dst[i] = val;
+}
+
diff --git a/kernels/compiler_half_convert.cl b/kernels/compiler_half_convert.cl
new file mode 100644
index 00000000..c28921e0
--- /dev/null
+++ b/kernels/compiler_half_convert.cl
@@ -0,0 +1,56 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+kernel void compiler_half_to_long_sat(global half *src, global long *dst) {
+  int i = get_global_id(0);
+  dst[i] = convert_long_sat(src[i]);
+}
+
+kernel void compiler_ulong_to_half(global ulong *src, global half *dst) {
+  int i = get_global_id(0);
+  dst[i] = convert_half(src[i]);
+}
+
+kernel void compiler_half_to_long(global half *src, global long *dst) {
+  int i = get_global_id(0);
+  dst[i] = convert_long(src[i]);
+}
+
+kernel void compiler_int_to_half(global int *src, global half *dst) {
+  int i = get_global_id(0);
+  dst[i] = convert_half(src[i]);
+}
+
+kernel void compiler_uchar_to_half(global uchar *src, global half *dst) {
+  int i = get_global_id(0);
+  dst[i] = convert_half(src[i]);
+}
+
+kernel void compiler_half_to_uint_sat(global half *src, global uint *dst) {
+  int i = get_global_id(0);
+  dst[i] = convert_uint(src[i]);
+}
+
+kernel void compiler_half_to_ushort_sat(global half *src, global ushort *dst) {
+  int i = get_global_id(0);
+  dst[i] = convert_ushort(src[i]);
+}
+
+kernel void compiler_half_to_char_sat(global half *src, global char *dst) {
+  int i = get_global_id(0);
+  dst[i] = convert_char_sat(src[i]);
+}
+
+kernel void compiler_half2_as_int(global half2 *src, global int *dst) {
+  int i = get_global_id(0);
+  dst[i] = as_int(src[i]);
+}
+
+kernel void compiler_half_as_char2(global half *src, global char2 *dst) {
+  int i = get_global_id(0);
+  dst[i] = as_char2(src[i]);
+}
+
+kernel void compiler_half_to_float(global half4 *src, global float4 *dst) {
+  int i = get_global_id(0);
+  dst[i] = convert_float4(src[i]);
+}
diff --git a/kernels/compiler_half_math.cl b/kernels/compiler_half_math.cl
new file mode 100644
index 00000000..a11a9566
--- /dev/null
+++ b/kernels/compiler_half_math.cl
@@ -0,0 +1,28 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+#define MATH_KERNEL_ARG1(NAME) \
+    kernel void compiler_half_math_##NAME(global half *src, global half *dst) { \
+    int i = get_global_id(0); \
+    dst[i] = NAME(src[i]); \
+  }
+
+MATH_KERNEL_ARG1(sin);
+MATH_KERNEL_ARG1(cos);
+MATH_KERNEL_ARG1(sinh);
+MATH_KERNEL_ARG1(cosh);
+MATH_KERNEL_ARG1(tan);
+MATH_KERNEL_ARG1(log10);
+MATH_KERNEL_ARG1(log);
+MATH_KERNEL_ARG1(trunc);
+MATH_KERNEL_ARG1(exp);
+MATH_KERNEL_ARG1(sqrt);
+MATH_KERNEL_ARG1(ceil);
+
+#define MATH_KERNEL_ARG2(NAME) \
+    kernel void compiler_half_math_##NAME(global half4 *src0, global half4 *src1, global half4 *dst) { \
+    int i = get_global_id(0); \
+    dst[i] = NAME(src0[i], src1[i]); \
+  }
+MATH_KERNEL_ARG2(fmod);
+MATH_KERNEL_ARG2(fmax);
+MATH_KERNEL_ARG2(fmin);
diff --git a/kernels/compiler_half_relation.cl b/kernels/compiler_half_relation.cl
new file mode 100644
index 00000000..dfb01e60
--- /dev/null
+++ b/kernels/compiler_half_relation.cl
@@ -0,0 +1,10 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+kernel void compiler_half_isnan(global half2 *src, global short2 *dst) {
+  int i = get_global_id(0);
+  dst[i] = isnan(src[i]);
+}
+
+kernel void compiler_half_isinf(global half *src, global int *dst) {
+  int i = get_global_id(0);
+  dst[i] = isinf(src[i]);
+}