From fe76b2116a76b6bc1dab0fac4dec35366d45da1e Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov <ilya.lavrenov@itseez.com>
Date: Thu, 14 Nov 2013 12:02:32 +0400
Subject: [PATCH] added RGBA <-> mRGBA

---
 modules/ocl/src/color.cpp           |  9 +++++
 modules/ocl/src/opencl/cvt_color.cl | 53 +++++++++++++++++++++++++++++
 modules/ocl/test/test_color.cpp     |  5 +++
 3 files changed, 67 insertions(+)

diff --git a/modules/ocl/src/color.cpp b/modules/ocl/src/color.cpp
index b807afd41c..f41d29782b 100644
--- a/modules/ocl/src/color.cpp
+++ b/modules/ocl/src/color.cpp
@@ -461,6 +461,15 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn)
         toRGB_caller(src, dst, bidx, kernelName, format(" -D hrange=%d -D hscale=%f", hrange, 6.f/hrange));
         break;
     }
+    case CV_RGBA2mRGBA: case CV_mRGBA2RGBA:
+        {
+            CV_Assert(scn == 4 && depth == CV_8U);
+            dst.create(sz, CV_MAKETYPE(depth, 4));
+            std::string kernelName = code == CV_RGBA2mRGBA ? "RGBA2mRGBA" : "mRGBA2RGBA";
+
+            fromRGB_caller(src, dst, 0, kernelName);
+            break;
+        }
     default:
         CV_Error( CV_StsBadFlag, "Unknown/unsupported color conversion code" );
     }
diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl
index 70922c3ff9..bf3b6cfa76 100644
--- a/modules/ocl/src/opencl/cvt_color.cl
+++ b/modules/ocl/src/opencl/cvt_color.cl
@@ -962,3 +962,56 @@ __kernel void HLS2RGB(int cols, int rows, int src_step, int dst_step, int bidx,
 }
 
 #endif
+
+/////////////////////////// RGBA <-> mRGBA (alpha premultiplied) //////////////
+
+#ifdef DEPTH_0
+
+__kernel void RGBA2mRGBA(int cols, int rows, int src_step, int dst_step,
+                        int bidx, __global const uchar * src, __global uchar * dst,
+                        int src_offset, int dst_offset)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (y < rows && x < cols)
+    {
+        x <<= 2;
+        int src_idx = mad24(y, src_step, src_offset + x);
+        int dst_idx = mad24(y, dst_step, dst_offset + x);
+
+        uchar v0 = src[src_idx], v1 = src[src_idx + 1];
+        uchar v2 = src[src_idx + 2], v3 = src[src_idx + 3];
+
+        dst[dst_idx] = (v0 * v3 + HALF_MAX) / MAX_NUM;
+        dst[dst_idx + 1] = (v1 * v3 + HALF_MAX) / MAX_NUM;
+        dst[dst_idx + 2] = (v2 * v3 + HALF_MAX) / MAX_NUM;
+        dst[dst_idx + 3] = v3;
+    }
+}
+
+__kernel void mRGBA2RGBA(int cols, int rows, int src_step, int dst_step, int bidx,
+                        __global const uchar * src, __global uchar * dst,
+                        int src_offset, int dst_offset)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (y < rows && x < cols)
+    {
+        x <<= 2;
+        int src_idx = mad24(y, src_step, src_offset + x);
+        int dst_idx = mad24(y, dst_step, dst_offset + x);
+
+        uchar v0 = src[src_idx], v1 = src[src_idx + 1];
+        uchar v2 = src[src_idx + 2], v3 = src[src_idx + 3];
+        uchar v3_half = v3 / 2;
+
+        dst[dst_idx] = v3 == 0 ? 0 : (v0 * MAX_NUM + v3_half) / v3;
+        dst[dst_idx + 1] = v3 == 0 ? 0 : (v1 * MAX_NUM + v3_half) / v3;
+        dst[dst_idx + 2] = v3 == 0 ? 0 : (v2 * MAX_NUM + v3_half) / v3;
+        dst[dst_idx + 3] = v3;
+    }
+}
+
+#endif
diff --git a/modules/ocl/test/test_color.cpp b/modules/ocl/test/test_color.cpp
index 935e974628..7c4fe39420 100644
--- a/modules/ocl/test/test_color.cpp
+++ b/modules/ocl/test/test_color.cpp
@@ -259,6 +259,11 @@ OCL_TEST_P(CvtColor8u, BGR5552GRAY) { doTest(2, 1, CVTCODE(BGR5552GRAY)); }
 OCL_TEST_P(CvtColor8u, GRAY2BGR565) { doTest(1, 2, CVTCODE(GRAY2BGR565)); }
 OCL_TEST_P(CvtColor8u, GRAY2BGR555) { doTest(1, 2, CVTCODE(GRAY2BGR555)); }
 
+// RGBA <-> mRGBA
+
+OCL_TEST_P(CvtColor8u, RGBA2mRGBA) { doTest(4, 4, CVTCODE(RGBA2mRGBA)); }
+OCL_TEST_P(CvtColor8u, mRGBA2RGBA) { doTest(4, 4, CVTCODE(mRGBA2RGBA)); }
+
 // YUV -> RGBA_NV12
 
 struct CvtColor_YUV420 :