From d2e1318341a4361378bae92e9898bcec852febe8 Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov <ilya.lavrenov@itseez.com>
Date: Wed, 27 Nov 2013 23:42:10 +0400
Subject: [PATCH] RGBA <-> mRGBA

---
 modules/imgproc/src/color.cpp           |  9 +++++++
 modules/imgproc/src/opencl/cvtcolor.cl  | 43 +++++++++++++++++++++++++++++++++
 modules/imgproc/test/ocl/test_color.cpp |  4 +--
 3 files changed, 54 insertions(+), 2 deletions(-)

diff --git a/modules/imgproc/src/color.cpp b/modules/imgproc/src/color.cpp
index 5ce2811..264490c 100644
--- a/modules/imgproc/src/color.cpp
+++ b/modules/imgproc/src/color.cpp
@@ -3010,6 +3010,15 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
                         depth, dcn, bidx, hrange, 6.f/hrange));
         break;
     }
+    case COLOR_RGBA2mRGBA: case COLOR_mRGBA2RGBA:
+    {
+        CV_Assert(scn == 4 && depth == CV_8U);
+        dcn = 4;
+
+        k.create(code == COLOR_RGBA2mRGBA ? "RGBA2mRGBA" : "mRGBA2RGBA", ocl::imgproc::cvtcolor_oclsrc,
+                 format("-D depth=%d -D dcn=4 -D scn=4 -D bidx=3", depth));
+        break;
+    }
     default:
         ;
     }
diff --git a/modules/imgproc/src/opencl/cvtcolor.cl b/modules/imgproc/src/opencl/cvtcolor.cl
index 53bb5a9..ca69629 100644
--- a/modules/imgproc/src/opencl/cvtcolor.cl
+++ b/modules/imgproc/src/opencl/cvtcolor.cl
@@ -988,12 +988,55 @@ __kernel void HLS2RGB(__global const uchar* srcptr, int src_step, int src_offset
 
 #endif
 
+/////////////////////////// RGBA <-> mRGBA (alpha premultiplied) //////////////
 
+#ifdef DEPTH_0
 
+__kernel void RGBA2mRGBA(__global const uchar* src, int src_step, int src_offset,
+                         __global uchar* dst, int dst_step, int dst_offset,
+                         int rows, int cols)
+{
+    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(__global const uchar* src, int src_step, int src_offset,
+                         __global uchar* dst, int dst_step, int dst_offset,
+                         int rows, int cols)
+{
+    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/imgproc/test/ocl/test_color.cpp b/modules/imgproc/test/ocl/test_color.cpp
index aad04b6..22b6b5c 100644
--- a/modules/imgproc/test/ocl/test_color.cpp
+++ b/modules/imgproc/test/ocl/test_color.cpp
@@ -247,8 +247,8 @@ OCL_TEST_P(CvtColor8u, GRAY2BGR555) { performTest(1, 2, CVTCODE(GRAY2BGR555)); }
 
 // RGBA <-> mRGBA
 
-//OCL_TEST_P(CvtColor8u, RGBA2mRGBA) { performTest(4, 4, CVTCODE(RGBA2mRGBA)); }
-//OCL_TEST_P(CvtColor8u, mRGBA2RGBA) { performTest(4, 4, CVTCODE(mRGBA2RGBA)); }
+OCL_TEST_P(CvtColor8u, RGBA2mRGBA) { performTest(4, 4, CVTCODE(RGBA2mRGBA)); }
+OCL_TEST_P(CvtColor8u, mRGBA2RGBA) { performTest(4, 4, CVTCODE(mRGBA2RGBA)); }
 
 // YUV -> RGBA_NV12
 
-- 
2.7.4