From b3e9efd931ce0aa76cdf1480d1b685a5bcae695f Mon Sep 17 00:00:00 2001
From: Homer Hsing <homer.xing@intel.com>
Date: Wed, 30 Jan 2013 14:05:31 +0800
Subject: [PATCH] Add convert_uchar_sat and test case

'convert_uchar_sat' converts float to uchar saturately.
'convert_uchar_sat' simply calls add_sat.
by 'convert_uchar_sat' function we don't need to clamp(value, 0, 255).
we also add a test case.

Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Lu Guanqun <guanqun.lu@intel.com>
---
 backend/src/ocl_stdlib.h              |  5 ++++
 kernels/compiler_convert_uchar_sat.cl |  4 ++++
 utests/CMakeLists.txt                 |  1 +
 utests/compiler_convert_uchar_sat.cpp | 44 +++++++++++++++++++++++++++++++++++
 4 files changed, 54 insertions(+)
 create mode 100644 kernels/compiler_convert_uchar_sat.cl
 create mode 100644 utests/compiler_convert_uchar_sat.cpp

diff --git a/backend/src/ocl_stdlib.h b/backend/src/ocl_stdlib.h
index d69c046..5c9b766 100644
--- a/backend/src/ocl_stdlib.h
+++ b/backend/src/ocl_stdlib.h
@@ -223,6 +223,11 @@ UDEF(uint);
 UDEF(ulong);
 #undef UDEF
 
+
+uchar INLINE_OVERLOADABLE convert_uchar_sat(float x) {
+    return add_sat((uchar)x, (uchar)0);
+}
+
 #define DEC2(name) INLINE_OVERLOADABLE int2 name(float2 x) { return (name(x.s0), name(x.s1)); }
 #define DEC3(name) INLINE_OVERLOADABLE int3 name(float3 x) { return (name(x.s0), name(x.s1), name(x.s2)); }
 #define DEC4(name) INLINE_OVERLOADABLE int4 name(float4 x) { return (name(x.s0), name(x.s1), name(x.s2), name(x.s3)); }
diff --git a/kernels/compiler_convert_uchar_sat.cl b/kernels/compiler_convert_uchar_sat.cl
new file mode 100644
index 0000000..0c81ecc
--- /dev/null
+++ b/kernels/compiler_convert_uchar_sat.cl
@@ -0,0 +1,4 @@
+kernel void compiler_convert_uchar_sat(global float *src, global uint *dst) {
+  int i = get_global_id(0);
+  dst[i] = convert_uchar_sat(src[i]);
+}
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 3a7176a..a599241 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -18,6 +18,7 @@ ADD_LIBRARY(utests SHARED
   compiler_array2.cpp
   compiler_array3.cpp
   compiler_byte_scatter.cpp
+  compiler_convert_uchar_sat.cpp
   compiler_copy_buffer.cpp
   compiler_copy_image.cpp
   compiler_copy_buffer_row.cpp
diff --git a/utests/compiler_convert_uchar_sat.cpp b/utests/compiler_convert_uchar_sat.cpp
new file mode 100644
index 0000000..da00041
--- /dev/null
+++ b/utests/compiler_convert_uchar_sat.cpp
@@ -0,0 +1,44 @@
+#include "utest_helper.hpp"
+
+static void cpu(int global_id, float *src, int *dst) {
+  float f = src[global_id];
+  dst[global_id] = f > 255 ? 255 : f < 0 ? 0 : f;
+}
+
+void compiler_convert_uchar_sat(void)
+{
+  const size_t n = 16;
+  float cpu_src[16];
+  int cpu_dst[16];
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_convert_uchar_sat");
+  OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(float), NULL);
+  OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(int), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+  globals[0] = 16;
+  locals[0] = 16;
+
+  // Run random tests
+  for (uint32_t pass = 0; pass < 8; ++pass) {
+    OCL_MAP_BUFFER(0);
+    for (int32_t i = 0; i < (int32_t) n; ++i)
+      cpu_src[i] = ((float*)buf_data[0])[i] = (rand() & 1023) / 2;
+    OCL_UNMAP_BUFFER(0);
+
+    // Run the kernel on GPU
+    OCL_NDRANGE(1);
+
+    // Run on CPU
+    for (int32_t i = 0; i < (int32_t) n; ++i) cpu(i, cpu_src, cpu_dst);
+
+    // Compare
+    OCL_MAP_BUFFER(1);
+    for (int32_t i = 0; i < (int32_t) n; ++i)
+      OCL_ASSERT(((int *)buf_data[1])[i] == cpu_dst[i]);
+    OCL_UNMAP_BUFFER(1);
+  }
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_convert_uchar_sat);
-- 
2.7.4