diff --git a/modules/core/src/matrix_transform.cpp b/modules/core/src/matrix_transform.cpp
index ce859e8eb8..245cd89843 100644
--- a/modules/core/src/matrix_transform.cpp
+++ b/modules/core/src/matrix_transform.cpp
@@ -973,7 +973,12 @@ static bool ocl_flip(InputArray _src, OutputArray _dst, int flipCode )
     if (cn > 4)
         return false;
 
-    const char * kernelName;
+    Size size = _src.size();
+    _dst.create(size, type);
+    UMat src = _src.getUMat(), dst = _dst.getUMat();
+    bool inplace = (dst.u == src.u);
+
+    String kernelName;
     if (flipCode == 0)
         kernelName = "arithm_flip_rows", flipType = FLIP_ROWS;
     else if (flipCode > 0)
@@ -981,33 +986,65 @@ static bool ocl_flip(InputArray _src, OutputArray _dst, int flipCode )
     else
         kernelName = "arithm_flip_rows_cols", flipType = FLIP_BOTH;
 
+    if(inplace)
+        kernelName += "_inplace";
+
     int pxPerWIy = (dev.isIntel() && (dev.type() & ocl::Device::TYPE_GPU)) ? 4 : 1;
     kercn = (cn!=3 || flipType == FLIP_ROWS) ? std::max(kercn, cn) : cn;
+    const int TILE_SIZE = 32, BLOCK_ROWS = 8;
 
-    ocl::Kernel k(kernelName, ocl::core::flip_oclsrc,
-        format( "-D T=%s -D T1=%s -D DEPTH=%d -D cn=%d -D PIX_PER_WI_Y=%d -D kercn=%d",
+    ocl::Kernel k(kernelName.c_str(), ocl::core::flip_oclsrc,
+                format( "-D T=%s -D T1=%s -D DEPTH=%d -D cn=%d -D PIX_PER_WI_Y=%d -D kercn=%d -D TILE_SIZE=%d -D BLOCK_ROWS=%d%s",
                 kercn != cn ? ocl::typeToStr(CV_MAKE_TYPE(depth, kercn)) : ocl::vecopTypeToStr(CV_MAKE_TYPE(depth, kercn)),
-                kercn != cn ? ocl::typeToStr(depth) : ocl::vecopTypeToStr(depth), depth, cn, pxPerWIy, kercn));
+                kercn != cn ? ocl::typeToStr(depth) : ocl::vecopTypeToStr(depth), depth, cn, pxPerWIy, kercn, TILE_SIZE, BLOCK_ROWS,
+                inplace ? " -D INPLACE" : ""));
     if (k.empty())
         return false;
 
-    Size size = _src.size();
-    _dst.create(size, type);
-    UMat src = _src.getUMat(), dst = _dst.getUMat();
-
     int cols = size.width * cn / kercn, rows = size.height;
-    cols = flipType == FLIP_COLS ? (cols + 1) >> 1 : cols;
-    rows = flipType & FLIP_ROWS ? (rows + 1) >> 1 : rows;
+    int work_cols = flipType == FLIP_COLS ? (cols + 1) >> 1 : cols;
+    int work_rows = flipType & FLIP_ROWS ? (rows + 1) >> 1 : rows;
 
-    k.args(ocl::KernelArg::ReadOnlyNoSize(src),
-           ocl::KernelArg::WriteOnly(dst, cn, kercn), rows, cols);
+    if (inplace)
+    {
+        k.args(ocl::KernelArg::ReadWriteNoSize(dst), rows, cols);
 
-    size_t maxWorkGroupSize = dev.maxWorkGroupSize();
-    CV_Assert(maxWorkGroupSize % 4 == 0);
+        int gs_cols, gs_rows;
+        if (flipType == FLIP_COLS)
+        {
+            gs_cols = work_cols;
+            gs_rows = rows;
+        }
+        else if (flipType == FLIP_ROWS)
+        {
+            gs_cols = cols;
+            gs_rows = work_rows;
+        }
+        else // FLIP_BOTH
+        {
+            gs_cols = cols;
+            gs_rows = rows;
+        }
 
-    size_t globalsize[2] = { (size_t)cols, ((size_t)rows + pxPerWIy - 1) / pxPerWIy },
-            localsize[2] = { maxWorkGroupSize / 4, 4 };
-    return k.run(2, globalsize, (flipType == FLIP_COLS) && !dev.isIntel() ? localsize : NULL, false);
+        size_t globalsize[2] = {
+            (size_t)divUp(gs_cols, TILE_SIZE) * TILE_SIZE,
+            (size_t)divUp(gs_rows, TILE_SIZE) * BLOCK_ROWS
+        };
+        size_t localsize[2] = { TILE_SIZE, BLOCK_ROWS };
+        return k.run(2, globalsize, localsize, false);
+    }
+    else
+    {
+        k.args(ocl::KernelArg::ReadOnlyNoSize(src),
+           ocl::KernelArg::WriteOnly(dst, cn, kercn), work_rows, work_cols);
+
+        size_t maxWorkGroupSize = dev.maxWorkGroupSize();
+        CV_Assert(maxWorkGroupSize % 4 == 0);
+
+        size_t globalsize[2] = { (size_t)work_cols, ((size_t)work_rows + pxPerWIy - 1) / pxPerWIy };
+        size_t localsize[2]  = { maxWorkGroupSize / 4, 4 };
+        return k.run(2, globalsize, (flipType == FLIP_COLS) && !dev.isIntel() ? localsize : NULL, false);
+    }
 }
 
 #endif
diff --git a/modules/core/src/opencl/flip.cl b/modules/core/src/opencl/flip.cl
index afd14e4e1f..dbe23c73a4 100644
--- a/modules/core/src/opencl/flip.cl
+++ b/modules/core/src/opencl/flip.cl
@@ -63,6 +63,9 @@
 #endif
 #define TSIZE ((int)sizeof(T1)*3)
 #endif
+#define LDS_STEP (TILE_SIZE + 1)
+
+#ifndef INPLACE
 
 __kernel void arithm_flip_rows(__global const uchar * srcptr, int src_step, int src_offset,
                                __global uchar * dstptr, int dst_step, int dst_offset,
@@ -183,3 +186,196 @@ __kernel void arithm_flip_cols(__global const uchar * srcptr, int src_step, int
         }
     }
 }
+
+#else
+
+__kernel void arithm_flip_rows_inplace(__global uchar * srcptr, int src_step, int src_offset,
+                                       int rows, int cols)
+{
+    int gp_x = get_group_id(0);
+    int gp_y = get_group_id(1);
+    int lx = get_local_id(0);
+    int ly = get_local_id(1);
+
+    __local T tile_top[TILE_SIZE * LDS_STEP];
+    __local T tile_bottom[TILE_SIZE * LDS_STEP];
+
+    int half_rows = (rows + 1) / 2;
+
+    int x = gp_x * TILE_SIZE + lx;
+    int y_top = gp_y * TILE_SIZE + ly;
+    int y_bottom = rows - 1 - y_top;
+
+    #pragma unroll
+    for (int i = 0; i < TILE_SIZE; i += BLOCK_ROWS)
+    {
+        if (x < cols && y_top + i < half_rows)
+        {
+            int curr_y_top = y_top + i;
+            int curr_y_bottom = rows - 1 - curr_y_top;
+
+            T val_top = loadpix(srcptr + mad24(curr_y_top, src_step, mad24(x, TSIZE, src_offset)));
+            T val_bottom = loadpix(srcptr + mad24(curr_y_bottom, src_step, mad24(x, TSIZE, src_offset)));
+
+            tile_top[mad24(ly + i, LDS_STEP, lx)] = val_top;
+            tile_bottom[mad24(ly + i, LDS_STEP, lx)] = val_bottom;
+        }
+    }
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    #pragma unroll
+    for (int i = 0; i < TILE_SIZE; i += BLOCK_ROWS)
+    {
+        if (x < cols && y_top + i < half_rows)
+        {
+            int curr_y_top = y_top + i;
+            int curr_y_bottom = rows - 1 - curr_y_top;
+
+            storepix(tile_bottom[mad24(ly + i, LDS_STEP, lx)],
+                    srcptr + mad24(curr_y_top, src_step, mad24(x, TSIZE, src_offset)));
+
+            if (curr_y_top != curr_y_bottom)
+            {
+                storepix(tile_top[mad24(ly + i, LDS_STEP, lx)],
+                        srcptr + mad24(curr_y_bottom, src_step, mad24(x, TSIZE, src_offset)));
+            }
+        }
+    }
+}
+
+__kernel void arithm_flip_rows_cols_inplace(__global uchar * srcptr, int src_step, int src_offset,
+                                            int rows, int cols)
+{
+    int gp_x = get_group_id(0);
+    int gp_y = get_group_id(1);
+    int lx = get_local_id(0);
+    int ly = get_local_id(1);
+
+    __local T tile_first[TILE_SIZE * LDS_STEP];
+    __local T tile_second[TILE_SIZE * LDS_STEP];
+
+    int total_pixels = rows * cols;
+    int half_pixels = (total_pixels + 1) / 2;
+
+    int x_first = gp_x * TILE_SIZE + lx;
+    int y_first = gp_y * TILE_SIZE + ly;
+
+    int x_second = cols - 1 - x_first;
+    int y_second = rows - 1 - y_first;
+
+    #pragma unroll
+    for (int i = 0; i < TILE_SIZE; i += BLOCK_ROWS)
+    {
+        int curr_y_first = y_first + i;
+        int curr_y_second = rows - 1 - curr_y_first;
+        int linear_idx = curr_y_first * cols + x_first;
+
+        if (x_first < cols && curr_y_first < rows && linear_idx < half_pixels)
+        {
+            T val_first = loadpix(srcptr + mad24(curr_y_first, src_step, mad24(x_first, TSIZE, src_offset)));
+            T val_second = loadpix(srcptr + mad24(curr_y_second, src_step, mad24(x_second, TSIZE, src_offset)));
+
+#if kercn == 2
+#if cn == 1
+            val_first = val_first.s10;
+            val_second = val_second.s10;
+#endif
+#elif kercn == 4
+#if cn == 1
+            val_first = val_first.s3210;
+            val_second = val_second.s3210;
+#elif cn == 2
+            val_first = val_first.s2301;
+            val_second = val_second.s2301;
+#endif
+#endif
+            tile_first[mad24(ly + i, LDS_STEP, lx)] = val_first;
+            tile_second[mad24(ly + i, LDS_STEP, lx)] = val_second;
+        }
+    }
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    #pragma unroll
+    for (int i = 0; i < TILE_SIZE; i += BLOCK_ROWS)
+    {
+        int curr_y_first = y_first + i;
+        int curr_y_second = rows - 1 - curr_y_first;
+        int linear_idx = curr_y_first * cols + x_first;
+
+        if (x_first < cols && curr_y_first < rows && linear_idx < half_pixels)
+        {
+            storepix(tile_second[mad24(ly + i, LDS_STEP, lx)],
+                    srcptr + mad24(curr_y_first, src_step, mad24(x_first, TSIZE, src_offset)));
+
+            if (linear_idx != total_pixels - 1 - linear_idx)
+            {
+                storepix(tile_first[mad24(ly + i, LDS_STEP, lx)],
+                        srcptr + mad24(curr_y_second, src_step, mad24(x_second, TSIZE, src_offset)));
+            }
+        }
+    }
+}
+
+__kernel void arithm_flip_cols_inplace(__global uchar * srcptr, int src_step, int src_offset,
+                                       int rows, int cols)
+{
+    int gp_x = get_group_id(0);
+    int gp_y = get_group_id(1);
+    int lx = get_local_id(0);
+    int ly = get_local_id(1);
+
+    __local T tile_left[TILE_SIZE * LDS_STEP];
+    __local T tile_right[TILE_SIZE * LDS_STEP];
+
+    int half_cols = (cols + 1) / 2;
+
+    int x_left = gp_x * TILE_SIZE + lx;
+    int y = gp_y * TILE_SIZE + ly;
+    int x_right = cols - 1 - x_left;
+
+    #pragma unroll
+    for (int i = 0; i < TILE_SIZE; i += BLOCK_ROWS)
+    {
+        if (y + i < rows && x_left < half_cols)
+        {
+            T val_left = loadpix(srcptr + mad24(y + i, src_step, mad24(x_left, TSIZE, src_offset)));
+            T val_right = loadpix(srcptr + mad24(y + i, src_step, mad24(x_right, TSIZE, src_offset)));
+
+#if kercn == 2
+#if cn == 1
+            val_left = val_left.s10;
+            val_right = val_right.s10;
+#endif
+#elif kercn == 4
+#if cn == 1
+            val_left = val_left.s3210;
+            val_right = val_right.s3210;
+#elif cn == 2
+            val_left = val_left.s2301;
+            val_right = val_right.s2301;
+#endif
+#endif
+            tile_left[mad24(ly + i, LDS_STEP, lx)] = val_left;
+            tile_right[mad24(ly + i, LDS_STEP, lx)] = val_right;
+        }
+    }
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    #pragma unroll
+    for (int i = 0; i < TILE_SIZE; i += BLOCK_ROWS)
+    {
+        if (y + i < rows && x_left < half_cols)
+        {
+            storepix(tile_right[mad24(ly + i, LDS_STEP, lx)],
+                    srcptr + mad24(y + i, src_step, mad24(x_left, TSIZE, src_offset)));
+
+            if (x_left != x_right)
+            {
+                storepix(tile_left[mad24(ly + i, LDS_STEP, lx)],
+                        srcptr + mad24(y + i, src_step, mad24(x_right, TSIZE, src_offset)));
+            }
+        }
+    }
+}
+
+#endif // INPLACE
\ No newline at end of file