diff --git a/modules/cudawarping/src/cuda/warp.cu b/modules/cudawarping/src/cuda/warp.cu index 2412f6ee88d..51da3d478f1 100644 --- a/modules/cudawarping/src/cuda/warp.cu +++ b/modules/cudawarping/src/cuda/warp.cu @@ -53,55 +53,75 @@ namespace cv { namespace cuda { namespace device { namespace imgproc { - __constant__ float c_warpMat[3 * 3]; - struct AffineTransform { - static __device__ __forceinline__ float2 calcCoord(int x, int y) + static const int rows = 2; + static __device__ __forceinline__ float2 calcCoord(const float warpMat[AffineTransform::rows * 3], int x, int y) { - const float xcoo = c_warpMat[0] * x + c_warpMat[1] * y + c_warpMat[2]; - const float ycoo = c_warpMat[3] * x + c_warpMat[4] * y + c_warpMat[5]; + const float xcoo = warpMat[0] * x + warpMat[1] * y + warpMat[2]; + const float ycoo = warpMat[3] * x + warpMat[4] * y + warpMat[5]; return make_float2(xcoo, ycoo); } + + struct Coefficients + { + Coefficients(const float* c_) + { + for(int i = 0; i < AffineTransform::rows * 3; i++) + c[i] = c_[i]; + } + float c[AffineTransform::rows * 3]; + }; }; struct PerspectiveTransform { - static __device__ __forceinline__ float2 calcCoord(int x, int y) + static const int rows = 3; + static __device__ __forceinline__ float2 calcCoord(const float warpMat[PerspectiveTransform::rows * 3], int x, int y) { - const float coeff = 1.0f / (c_warpMat[6] * x + c_warpMat[7] * y + c_warpMat[8]); + const float coeff = 1.0f / (warpMat[6] * x + warpMat[7] * y + warpMat[8]); - const float xcoo = coeff * (c_warpMat[0] * x + c_warpMat[1] * y + c_warpMat[2]); - const float ycoo = coeff * (c_warpMat[3] * x + c_warpMat[4] * y + c_warpMat[5]); + const float xcoo = coeff * (warpMat[0] * x + warpMat[1] * y + warpMat[2]); + const float ycoo = coeff * (warpMat[3] * x + warpMat[4] * y + warpMat[5]); return make_float2(xcoo, ycoo); } + struct Coefficients + { + Coefficients(const float* c_) + { + for(int i = 0; i < PerspectiveTransform::rows * 3; i++) + c[i] = c_[i]; + } + + float c[PerspectiveTransform::rows * 3]; + }; }; /////////////////////////////////////////////////////////////////// // Build Maps - template __global__ void buildWarpMaps(PtrStepSzf xmap, PtrStepf ymap) + template __global__ void buildWarpMaps(PtrStepSzf xmap, PtrStepf ymap, const typename Transform::Coefficients warpMat) { const int x = blockDim.x * blockIdx.x + threadIdx.x; const int y = blockDim.y * blockIdx.y + threadIdx.y; if (x < xmap.cols && y < xmap.rows) { - const float2 coord = Transform::calcCoord(x, y); + const float2 coord = Transform::calcCoord(warpMat.c, x, y); xmap(y, x) = coord.x; ymap(y, x) = coord.y; } } - template void buildWarpMaps_caller(PtrStepSzf xmap, PtrStepSzf ymap, cudaStream_t stream) + template void buildWarpMaps_caller(PtrStepSzf xmap, PtrStepSzf ymap, const float warpMat[Transform::rows * 3], cudaStream_t stream) { dim3 block(32, 8); dim3 grid(divUp(xmap.cols, block.x), divUp(xmap.rows, block.y)); - buildWarpMaps<<>>(xmap, ymap); + buildWarpMaps<<>>(xmap, ymap, warpMat); cudaSafeCall( cudaGetLastError() ); if (stream == 0) @@ -110,29 +130,25 @@ namespace cv { namespace cuda { namespace device void buildWarpAffineMaps_gpu(float coeffs[2 * 3], PtrStepSzf xmap, PtrStepSzf ymap, cudaStream_t stream) { - cudaSafeCall( cudaMemcpyToSymbol(c_warpMat, coeffs, 2 * 3 * sizeof(float)) ); - - buildWarpMaps_caller(xmap, ymap, stream); + buildWarpMaps_caller(xmap, ymap, coeffs, stream); } void buildWarpPerspectiveMaps_gpu(float coeffs[3 * 3], PtrStepSzf xmap, PtrStepSzf ymap, cudaStream_t stream) { - cudaSafeCall( cudaMemcpyToSymbol(c_warpMat, coeffs, 3 * 3 * sizeof(float)) ); - - buildWarpMaps_caller(xmap, ymap, stream); + buildWarpMaps_caller(xmap, ymap, coeffs, stream); } /////////////////////////////////////////////////////////////////// // Warp - template __global__ void warp(const Ptr2D src, PtrStepSz dst) + template __global__ void warp(const Ptr2D src, PtrStepSz dst, const typename Transform::Coefficients warpMat) { const int x = blockDim.x * blockIdx.x + threadIdx.x; const int y = blockDim.y * blockIdx.y + threadIdx.y; if (x < dst.cols && y < dst.rows) { - const float2 coord = Transform::calcCoord(x, y); + const float2 coord = Transform::calcCoord(warpMat.c, x, y); dst.ptr(y)[x] = saturate_cast(src(coord.y, coord.x)); } @@ -140,7 +156,7 @@ namespace cv { namespace cuda { namespace device template class Filter, template class B, typename T> struct WarpDispatcherStream { - static void call(PtrStepSz src, PtrStepSz dst, const float* borderValue, cudaStream_t stream, bool) + static void call(PtrStepSz src, PtrStepSz dst, const float* borderValue, const float warpMat[Transform::rows*3], cudaStream_t stream, bool) { typedef typename TypeVec::cn>::vec_type work_type; @@ -151,14 +167,14 @@ namespace cv { namespace cuda { namespace device BorderReader< PtrStep, B > brdSrc(src, brd); Filter< BorderReader< PtrStep, B > > filter_src(brdSrc); - warp<<>>(filter_src, dst); + warp<<>>(filter_src, dst, warpMat); cudaSafeCall( cudaGetLastError() ); } }; template class Filter, template class B, typename T> struct WarpDispatcherNonStream { - static void call(PtrStepSz src, PtrStepSz srcWhole, int xoff, int yoff, PtrStepSz dst, const float* borderValue, bool) + static void call(PtrStepSz src, PtrStepSz srcWhole, int xoff, int yoff, PtrStepSz dst, const float* borderValue, const float warpMat[Transform::rows*3], bool) { CV_UNUSED(xoff); CV_UNUSED(yoff); @@ -173,7 +189,7 @@ namespace cv { namespace cuda { namespace device BorderReader< PtrStep, B > brdSrc(src, brd); Filter< BorderReader< PtrStep, B > > filter_src(brdSrc); - warp<<>>(filter_src, dst); + warp<<>>(filter_src, dst, warpMat); cudaSafeCall( cudaGetLastError() ); cudaSafeCall( cudaDeviceSynchronize() ); @@ -195,7 +211,7 @@ namespace cv { namespace cuda { namespace device }; \ template class Filter, template class B> struct WarpDispatcherNonStream \ { \ - static void call(PtrStepSz< type > src, PtrStepSz< type > srcWhole, int xoff, int yoff, PtrStepSz< type > dst, const float* borderValue, bool cc20) \ + static void call(PtrStepSz< type > src, PtrStepSz< type > srcWhole, int xoff, int yoff, PtrStepSz< type > dst, const float* borderValue, const float warpMat[Transform::rows*3], bool cc20) \ { \ typedef typename TypeVec::cn>::vec_type work_type; \ dim3 block(32, cc20 ? 8 : 4); \ @@ -205,14 +221,14 @@ namespace cv { namespace cuda { namespace device B brd(src.rows, src.cols, VecTraits::make(borderValue)); \ BorderReader< tex_warp_ ## type ##_reader, B > brdSrc(texSrc, brd); \ Filter< BorderReader< tex_warp_ ## type ##_reader, B > > filter_src(brdSrc); \ - warp<<>>(filter_src, dst); \ + warp<<>>(filter_src, dst, warpMat); \ cudaSafeCall( cudaGetLastError() ); \ cudaSafeCall( cudaDeviceSynchronize() ); \ } \ }; \ template class Filter> struct WarpDispatcherNonStream \ { \ - static void call(PtrStepSz< type > src, PtrStepSz< type > srcWhole, int xoff, int yoff, PtrStepSz< type > dst, const float*, bool) \ + static void call(PtrStepSz< type > src, PtrStepSz< type > srcWhole, int xoff, int yoff, PtrStepSz< type > dst, const float*, const float warpMat[Transform::rows*3], bool) \ { \ dim3 block(32, 8); \ dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y)); \ @@ -221,14 +237,14 @@ namespace cv { namespace cuda { namespace device if (srcWhole.cols == src.cols && srcWhole.rows == src.rows) \ { \ Filter< tex_warp_ ## type ##_reader > filter_src(texSrc); \ - warp<<>>(filter_src, dst); \ + warp<<>>(filter_src, dst, warpMat); \ } \ else \ { \ BrdReplicate brd(src.rows, src.cols); \ BorderReader< tex_warp_ ## type ##_reader, BrdReplicate > brdSrc(texSrc, brd); \ Filter< BorderReader< tex_warp_ ## type ##_reader, BrdReplicate > > filter_src(brdSrc); \ - warp<<>>(filter_src, dst); \ + warp<<>>(filter_src, dst, warpMat); \ } \ cudaSafeCall( cudaGetLastError() ); \ cudaSafeCall( cudaDeviceSynchronize() ); \ @@ -263,20 +279,20 @@ namespace cv { namespace cuda { namespace device template class Filter, template class B, typename T> struct WarpDispatcher { - static void call(PtrStepSz src, PtrStepSz srcWhole, int xoff, int yoff, PtrStepSz dst, const float* borderValue, cudaStream_t stream, bool cc20) + static void call(PtrStepSz src, PtrStepSz srcWhole, int xoff, int yoff, PtrStepSz dst, const float* borderValue, const float warpMat[Transform::rows*3], cudaStream_t stream, bool cc20) { if (stream == 0) - WarpDispatcherNonStream::call(src, srcWhole, xoff, yoff, dst, borderValue, cc20); + WarpDispatcherNonStream::call(src, srcWhole, xoff, yoff, dst, borderValue, warpMat, cc20); else - WarpDispatcherStream::call(src, dst, borderValue, stream, cc20); + WarpDispatcherStream::call(src, dst, borderValue, warpMat, stream, cc20); } }; template void warp_caller(PtrStepSzb src, PtrStepSzb srcWhole, int xoff, int yoff, PtrStepSzb dst, int interpolation, - int borderMode, const float* borderValue, cudaStream_t stream, bool cc20) + int borderMode, const float* borderValue, const float warpMat[Transform::rows*3], cudaStream_t stream, bool cc20) { - typedef void (*func_t)(PtrStepSz src, PtrStepSz srcWhole, int xoff, int yoff, PtrStepSz dst, const float* borderValue, cudaStream_t stream, bool cc20); + typedef void (*func_t)(PtrStepSz src, PtrStepSz srcWhole, int xoff, int yoff, PtrStepSz dst, const float* borderValue, const float warpMat[Transform::rows*3], cudaStream_t stream, bool cc20); static const func_t funcs[3][5] = { @@ -304,15 +320,13 @@ namespace cv { namespace cuda { namespace device }; funcs[interpolation][borderMode](static_cast< PtrStepSz >(src), static_cast< PtrStepSz >(srcWhole), xoff, yoff, - static_cast< PtrStepSz >(dst), borderValue, stream, cc20); + static_cast< PtrStepSz >(dst), borderValue, warpMat, stream, cc20); } template void warpAffine_gpu(PtrStepSzb src, PtrStepSzb srcWhole, int xoff, int yoff, float coeffs[2 * 3], PtrStepSzb dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, bool cc20) { - cudaSafeCall( cudaMemcpyToSymbol(c_warpMat, coeffs, 2 * 3 * sizeof(float)) ); - - warp_caller(src, srcWhole, xoff, yoff, dst, interpolation, borderMode, borderValue, stream, cc20); + warp_caller(src, srcWhole, xoff, yoff, dst, interpolation, borderMode, borderValue, coeffs, stream, cc20); } template void warpAffine_gpu(PtrStepSzb src, PtrStepSzb srcWhole, int xoff, int yoff, float coeffs[2 * 3], PtrStepSzb dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, bool cc20); @@ -348,9 +362,7 @@ namespace cv { namespace cuda { namespace device template void warpPerspective_gpu(PtrStepSzb src, PtrStepSzb srcWhole, int xoff, int yoff, float coeffs[3 * 3], PtrStepSzb dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, bool cc20) { - cudaSafeCall( cudaMemcpyToSymbol(c_warpMat, coeffs, 3 * 3 * sizeof(float)) ); - - warp_caller(src, srcWhole, xoff, yoff, dst, interpolation, borderMode, borderValue, stream, cc20); + warp_caller(src, srcWhole, xoff, yoff, dst, interpolation, borderMode, borderValue, coeffs, stream, cc20); } template void warpPerspective_gpu(PtrStepSzb src, PtrStepSzb srcWhole, int xoff, int yoff, float coeffs[3 * 3], PtrStepSzb dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, bool cc20);