Merge pull request #3172 from ilya-lavrenov:tapi_transpose_amd

modules/core/src/matrix.cpp

@@ -3011,9 +3011,9 @@ static bool ocl_transpose( InputArray _src, OutputArray _dst )
     }
 
     ocl::Kernel k(kernelName.c_str(), ocl::core::transpose_oclsrc,
-                  format("-D T=%s -D T1=%s -D cn=%d -D TILE_DIM=%d -D BLOCK_ROWS=%d -D rowsPerWI=%d",
+                  format("-D T=%s -D T1=%s -D cn=%d -D TILE_DIM=%d -D BLOCK_ROWS=%d -D rowsPerWI=%d%s",
                          ocl::memopTypeToStr(type), ocl::memopTypeToStr(depth),
-                         cn, TILE_DIM, BLOCK_ROWS, rowsPerWI));
+                         cn, TILE_DIM, BLOCK_ROWS, rowsPerWI, inplace ? " -D INPLACE" : ""));
     if (k.empty())
         return false;
 

modules/core/src/opencl/transpose.cl

@@ -53,6 +53,8 @@
 #define TSIZE ((int)sizeof(T1)*3)
 #endif
 
+#ifndef INPLACE
+
 #define LDS_STEP      (TILE_DIM + 1)
 
 __kernel void transpose(__global const uchar * srcptr, int src_step, int src_offset, int src_rows, int src_cols,
@@ -114,6 +116,8 @@ __kernel void transpose(__global const uchar * srcptr, int src_step, int src_off
     }
 }
 
+#else
+
 __kernel void transpose_inplace(__global uchar * srcptr, int src_step, int src_offset, int src_rows)
 {
     int x = get_global_id(0);
@@ -138,3 +142,5 @@ __kernel void transpose_inplace(__global uchar * srcptr, int src_step, int src_o
             }
     }
 }
+
+#endif // INPLACE