From 1f51e6c0deedea619d9cd6ffb39e692a3af8ae96 Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov <ilya.lavrenov@itseez.com>
Date: Thu, 10 Oct 2013 22:28:09 +0400
Subject: [PATCH] fixed kernel compilation warnings on MacOSX

---
 modules/ocl/src/cl_operations.cpp             | 27 +++++++++++++++++--
 modules/ocl/src/imgproc.cpp                   |  2 +-
 .../src/opencl/arithm_absdiff_nonsaturate.cl  |  4 +--
 modules/ocl/src/opencl/arithm_sum.cl          | 12 ++++-----
 modules/ocl/src/pyrlk.cpp                     |  2 +-
 5 files changed, 35 insertions(+), 12 deletions(-)

diff --git a/modules/ocl/src/cl_operations.cpp b/modules/ocl/src/cl_operations.cpp
index ed13be5dd1..9514cc390b 100644
--- a/modules/ocl/src/cl_operations.cpp
+++ b/modules/ocl/src/cl_operations.cpp
@@ -212,13 +212,35 @@ void openCLVerifyKernel(const Context *ctx, cl_kernel kernel, size_t *localThrea
 static double total_execute_time = 0;
 static double total_kernel_time = 0;
 #endif
+
+static std::string removeDuplicatedWhiteSpaces(const char * buildOptions)
+{
+    if (buildOptions == NULL)
+        return "";
+
+    size_t length = strlen(buildOptions), didx = 0, sidx = 0;
+    while (sidx < length && buildOptions[sidx] == 0)
+        ++sidx;
+
+    std::string opt;
+    opt.resize(length);
+
+    for ( ; sidx < length; ++sidx)
+        if (buildOptions[sidx] != ' ')
+            opt[didx++] = buildOptions[sidx];
+        else if ( !(didx > 0 && opt[didx - 1] == ' ') )
+            opt[didx++] = buildOptions[sidx];
+
+    return opt;
+}
+
 void openCLExecuteKernel_(Context *ctx, const cv::ocl::ProgramEntry* source, string kernelName, size_t globalThreads[3],
                           size_t localThreads[3],  vector< pair<size_t, const void *> > &args, int channels,
                           int depth, const char *build_options)
 {
     //construct kernel name
     //The rule is functionName_Cn_Dn, C represent Channels, D Represent DataType Depth, n represent an integer number
-    //for exmaple split_C2_D2, represent the split kernel with channels =2 and dataType Depth = 2(Data type is char)
+    //for example split_C2_D3, represent the split kernel with channels = 2 and dataType Depth = 3(Data type is short)
     stringstream idxStr;
     if(channels != -1)
         idxStr << "_C" << channels;
@@ -227,7 +249,8 @@ void openCLExecuteKernel_(Context *ctx, const cv::ocl::ProgramEntry* source, str
     kernelName += idxStr.str();
 
     cl_kernel kernel;
-    kernel = openCLGetKernelFromSource(ctx, source, kernelName, build_options);
+    std::string fixedOptions = removeDuplicatedWhiteSpaces(build_options);
+    kernel = openCLGetKernelFromSource(ctx, source, kernelName, fixedOptions.c_str());
 
     if ( localThreads != NULL)
     {
diff --git a/modules/ocl/src/imgproc.cpp b/modules/ocl/src/imgproc.cpp
index 5b00078b72..81ab2fc79f 100644
--- a/modules/ocl/src/imgproc.cpp
+++ b/modules/ocl/src/imgproc.cpp
@@ -1497,7 +1497,7 @@ namespace cv
                     openCLSafeCall(clReleaseKernel(kernel));
 
                     static char opt[20] = {0};
-                    sprintf(opt, " -D WAVE_SIZE=%d", (int)wave_size);
+                    sprintf(opt, "-D WAVE_SIZE=%d", (int)wave_size);
                     openCLExecuteKernel(Context::getContext(), &imgproc_clahe, kernelName, globalThreads, localThreads, args, -1, -1, opt);
                 }
             }
diff --git a/modules/ocl/src/opencl/arithm_absdiff_nonsaturate.cl b/modules/ocl/src/opencl/arithm_absdiff_nonsaturate.cl
index e5d8271394..0208806069 100644
--- a/modules/ocl/src/opencl/arithm_absdiff_nonsaturate.cl
+++ b/modules/ocl/src/opencl/arithm_absdiff_nonsaturate.cl
@@ -70,7 +70,7 @@ __kernel void arithm_absdiff_nonsaturate_binary(__global srcT *src1, int src1_st
         dstT t1 = convertToDstT(src2[src2_index]);
         dstT t2 = t0 - t1;
 
-        dst[dst_index] = t2 >= 0 ? t2 : -t2;
+        dst[dst_index] = t2 >= (dstT)(0) ? t2 : -t2;
     }
 }
 
@@ -88,6 +88,6 @@ __kernel void arithm_absdiff_nonsaturate(__global srcT *src1, int src1_step, int
 
         dstT t0 = convertToDstT(src1[src1_index]);
 
-        dst[dst_index] = t0 >= 0 ? t0 : -t0;
+        dst[dst_index] = t0 >= (dstT)(0) ? t0 : -t0;
     }
 }
diff --git a/modules/ocl/src/opencl/arithm_sum.cl b/modules/ocl/src/opencl/arithm_sum.cl
index 4011f03bea..39bcf949a0 100644
--- a/modules/ocl/src/opencl/arithm_sum.cl
+++ b/modules/ocl/src/opencl/arithm_sum.cl
@@ -51,14 +51,14 @@
 #endif
 #endif
 
-#if defined (FUNC_SUM)
+#if FUNC_SUM
 #define FUNC(a, b) b += a;
-#endif
-#if defined (FUNC_ABS_SUM)
-#define FUNC(a, b) b += a >= 0 ? a : -a;
-#endif
-#if defined (FUNC_SQR_SUM)
+#elif FUNC_ABS_SUM
+#define FUNC(a, b) b += a >= (dstT)(0) ? a : -a;
+#elif FUNC_SQR_SUM
 #define FUNC(a, b) b += a * a;
+#else
+#error No sum function
 #endif
 
 /**************************************Array buffer SUM**************************************/
diff --git a/modules/ocl/src/pyrlk.cpp b/modules/ocl/src/pyrlk.cpp
index 8e8692e77f..bd9b18e8c3 100644
--- a/modules/ocl/src/pyrlk.cpp
+++ b/modules/ocl/src/pyrlk.cpp
@@ -134,7 +134,7 @@ static void lkSparse_run(oclMat &I, oclMat &J,
             openCLSafeCall(clReleaseKernel(kernel));
 
             static char opt[32] = {0};
-            sprintf(opt, " -D WAVE_SIZE=%d", wave_size);
+            sprintf(opt, "-D WAVE_SIZE=%d", wave_size);
 
             openCLExecuteKernel(clCxt, &pyrlk, kernelName, globalThreads, localThreads,
                                 args, I.oclchannels(), I.depth(), opt);