From cc34a8ac3c291f107bc226a09aebe86969aaf085 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Tue, 16 Apr 2013 13:17:51 +0400 Subject: [PATCH] new implementation for GpuMat::setTo (without constant memory) --- modules/core/src/cuda/matrix_operations.cu | 130 +++++---------------- 1 file changed, 31 insertions(+), 99 deletions(-) diff --git a/modules/core/src/cuda/matrix_operations.cu b/modules/core/src/cuda/matrix_operations.cu index d16a88df1..7de5205ec 100644 --- a/modules/core/src/cuda/matrix_operations.cu +++ b/modules/core/src/cuda/matrix_operations.cu @@ -44,42 +44,33 @@ #include "opencv2/core/cuda/transform.hpp" #include "opencv2/core/cuda/functional.hpp" #include "opencv2/core/cuda/type_traits.hpp" +#include "opencv2/core/cuda/vec_traits.hpp" #include "matrix_operations.hpp" namespace cv { namespace gpu { namespace cudev { - template struct shift_and_sizeof; - template <> struct shift_and_sizeof { enum { shift = 0 }; }; - template <> struct shift_and_sizeof { enum { shift = 0 }; }; - template <> struct shift_and_sizeof { enum { shift = 1 }; }; - template <> struct shift_and_sizeof { enum { shift = 1 }; }; - template <> struct shift_and_sizeof { enum { shift = 2 }; }; - template <> struct shift_and_sizeof { enum { shift = 2 }; }; - template <> struct shift_and_sizeof { enum { shift = 3 }; }; - - /////////////////////////////////////////////////////////////////////////// - ////////////////////////////////// CopyTo ///////////////////////////////// /////////////////////////////////////////////////////////////////////////// + // copyWithMask template void copyWithMask(PtrStepSzb src, PtrStepSzb dst, int cn, PtrStepSzb mask, bool multiChannelMask, cudaStream_t stream) { if (multiChannelMask) - cv::gpu::cudev::transform((PtrStepSz)src, (PtrStepSz)dst, identity(), SingleMask(mask), stream); + cv::gpu::cudev::transform((PtrStepSz) src, (PtrStepSz) dst, identity(), SingleMask(mask), stream); else - cv::gpu::cudev::transform((PtrStepSz)src, (PtrStepSz)dst, identity(), SingleMaskChannels(mask, cn), stream); + cv::gpu::cudev::transform((PtrStepSz) src, (PtrStepSz) dst, identity(), SingleMaskChannels(mask, cn), stream); } void copyWithMask(PtrStepSzb src, PtrStepSzb dst, size_t elemSize1, int cn, PtrStepSzb mask, bool multiChannelMask, cudaStream_t stream) { typedef void (*func_t)(PtrStepSzb src, PtrStepSzb dst, int cn, PtrStepSzb mask, bool multiChannelMask, cudaStream_t stream); - static func_t tab[] = + static const func_t tab[] = { 0, - copyWithMask, - copyWithMask, + copyWithMask, + copyWithMask, 0, copyWithMask, 0, @@ -88,81 +79,39 @@ namespace cv { namespace gpu { namespace cudev copyWithMask }; - tab[elemSize1](src, dst, cn, mask, multiChannelMask, stream); + const func_t func = tab[elemSize1]; + CV_DbgAssert( func != 0 ); + + func(src, dst, cn, mask, multiChannelMask, stream); } /////////////////////////////////////////////////////////////////////////// - ////////////////////////////////// SetTo ////////////////////////////////// - /////////////////////////////////////////////////////////////////////////// + // set - __constant__ uchar scalar_8u[4]; - __constant__ schar scalar_8s[4]; - __constant__ ushort scalar_16u[4]; - __constant__ short scalar_16s[4]; - __constant__ int scalar_32s[4]; - __constant__ float scalar_32f[4]; - __constant__ double scalar_64f[4]; + template + __global__ void set(PtrStepSz mat, const Mask mask, const int channels, const typename TypeVec::vec_type value) + { + const int x = blockIdx.x * blockDim.x + threadIdx.x; + const int y = blockIdx.y * blockDim.y + threadIdx.y; - template __device__ __forceinline__ T readScalar(int i); - template <> __device__ __forceinline__ uchar readScalar(int i) {return scalar_8u[i];} - template <> __device__ __forceinline__ schar readScalar(int i) {return scalar_8s[i];} - template <> __device__ __forceinline__ ushort readScalar(int i) {return scalar_16u[i];} - template <> __device__ __forceinline__ short readScalar(int i) {return scalar_16s[i];} - template <> __device__ __forceinline__ int readScalar(int i) {return scalar_32s[i];} - template <> __device__ __forceinline__ float readScalar(int i) {return scalar_32f[i];} - template <> __device__ __forceinline__ double readScalar(int i) {return scalar_64f[i];} + if (x >= mat.cols * channels || y >= mat.rows) + return; - static inline void writeScalar(const uchar* vals) - { - cudaSafeCall( cudaMemcpyToSymbol(scalar_8u, vals, sizeof(uchar) * 4) ); - } - static inline void writeScalar(const schar* vals) - { - cudaSafeCall( cudaMemcpyToSymbol(scalar_8s, vals, sizeof(schar) * 4) ); - } - static inline void writeScalar(const ushort* vals) - { - cudaSafeCall( cudaMemcpyToSymbol(scalar_16u, vals, sizeof(ushort) * 4) ); - } - static inline void writeScalar(const short* vals) - { - cudaSafeCall( cudaMemcpyToSymbol(scalar_16s, vals, sizeof(short) * 4) ); - } - static inline void writeScalar(const int* vals) - { - cudaSafeCall( cudaMemcpyToSymbol(scalar_32s, vals, sizeof(int) * 4) ); - } - static inline void writeScalar(const float* vals) - { - cudaSafeCall( cudaMemcpyToSymbol(scalar_32f, vals, sizeof(float) * 4) ); - } - static inline void writeScalar(const double* vals) - { - cudaSafeCall( cudaMemcpyToSymbol(scalar_64f, vals, sizeof(double) * 4) ); - } + const T scalar[4] = {value.x, value.y, value.z, value.w}; - template - __global__ void set(T* mat, int cols, int rows, size_t step, int channels) - { - size_t x = blockIdx.x * blockDim.x + threadIdx.x; - size_t y = blockIdx.y * blockDim.y + threadIdx.y; - - if ((x < cols * channels ) && (y < rows)) - { - size_t idx = y * ( step >> shift_and_sizeof::shift ) + x; - mat[idx] = readScalar(x % channels); - } + if (mask(y, x / channels)) + mat(y, x) = scalar[x % channels]; } template void set(PtrStepSz mat, const T* scalar, int channels, cudaStream_t stream) { - writeScalar(scalar); + typedef typename TypeVec::vec_type scalar_t; - dim3 threadsPerBlock(32, 8, 1); - dim3 numBlocks(mat.cols * channels / threadsPerBlock.x + 1, mat.rows / threadsPerBlock.y + 1, 1); + dim3 block(32, 8); + dim3 grid(divUp(mat.cols * channels, block.x), divUp(mat.rows, block.y)); - set<<>>(mat.data, mat.cols, mat.rows, mat.step, channels); + set<<>>(mat, WithOutMask(), channels, VecTraits::make(scalar)); cudaSafeCall( cudaGetLastError() ); if (stream == 0) @@ -177,29 +126,15 @@ namespace cv { namespace gpu { namespace cudev template void set(PtrStepSz mat, const float* scalar, int channels, cudaStream_t stream); template void set(PtrStepSz mat, const double* scalar, int channels, cudaStream_t stream); - template - __global__ void set(T* mat, const uchar* mask, int cols, int rows, size_t step, int channels, size_t step_mask) - { - size_t x = blockIdx.x * blockDim.x + threadIdx.x; - size_t y = blockIdx.y * blockDim.y + threadIdx.y; - - if ((x < cols * channels ) && (y < rows)) - if (mask[y * step_mask + x / channels] != 0) - { - size_t idx = y * ( step >> shift_and_sizeof::shift ) + x; - mat[idx] = readScalar(x % channels); - } - } - template void set(PtrStepSz mat, const T* scalar, PtrStepSzb mask, int channels, cudaStream_t stream) { - writeScalar(scalar); + typedef typename TypeVec::vec_type scalar_t; - dim3 threadsPerBlock(32, 8, 1); - dim3 numBlocks(mat.cols * channels / threadsPerBlock.x + 1, mat.rows / threadsPerBlock.y + 1, 1); + dim3 block(32, 8); + dim3 grid(divUp(mat.cols * channels, block.x), divUp(mat.rows, block.y)); - set<<>>(mat.data, mask.data, mat.cols, mat.rows, mat.step, channels, mask.step); + set<<>>(mat, SingleMask(mask), channels, VecTraits::make(scalar)); cudaSafeCall( cudaGetLastError() ); if (stream == 0) @@ -215,8 +150,7 @@ namespace cv { namespace gpu { namespace cudev template void set(PtrStepSz mat, const double* scalar, PtrStepSzb mask, int channels, cudaStream_t stream); /////////////////////////////////////////////////////////////////////////// - //////////////////////////////// ConvertTo //////////////////////////////// - /////////////////////////////////////////////////////////////////////////// + // convert template struct Convertor : unary_function { @@ -281,8 +215,6 @@ namespace cv { namespace gpu { namespace cudev template void cvt_(PtrStepSzb src, PtrStepSzb dst, double alpha, double beta, cudaStream_t stream) { - cudaSafeCall( cudaSetDoubleForDevice(&alpha) ); - cudaSafeCall( cudaSetDoubleForDevice(&beta) ); Convertor op(static_cast(alpha), static_cast(beta)); cv::gpu::cudev::transform((PtrStepSz)src, (PtrStepSz)dst, op, WithOutMask(), stream); }