From 4f5d30a865bd45f144cc70aed1057ac500a72e0c Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Fri, 27 Feb 2015 16:28:06 +0300 Subject: [PATCH] fix GpuMat::setTo implementation previous implementation was not thread/stream safe, since it used constant memory new implementation doesn't use any global objects, so it is thread/stream safe --- .../dynamicuda/src/cuda/matrix_operations.cu | 107 +++++++----------- 1 file changed, 40 insertions(+), 67 deletions(-) diff --git a/modules/dynamicuda/src/cuda/matrix_operations.cu b/modules/dynamicuda/src/cuda/matrix_operations.cu index c9b12de104..35bca63200 100644 --- a/modules/dynamicuda/src/cuda/matrix_operations.cu +++ b/modules/dynamicuda/src/cuda/matrix_operations.cu @@ -44,6 +44,7 @@ #include "opencv2/gpu/device/transform.hpp" #include "opencv2/gpu/device/functional.hpp" #include "opencv2/gpu/device/type_traits.hpp" +#include "opencv2/gpu/device/vec_traits.hpp" namespace cv { namespace gpu { namespace device { @@ -105,87 +106,59 @@ namespace cv { namespace gpu { namespace device ////////////////////////////////// SetTo ////////////////////////////////// /////////////////////////////////////////////////////////////////////////// - __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 __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];} - - void writeScalar(const uchar* vals) - { - cudaSafeCall( cudaMemcpyToSymbol(scalar_8u, vals, sizeof(uchar) * 4) ); - } - void writeScalar(const schar* vals) - { - cudaSafeCall( cudaMemcpyToSymbol(scalar_8s, vals, sizeof(schar) * 4) ); - } - void writeScalar(const ushort* vals) - { - cudaSafeCall( cudaMemcpyToSymbol(scalar_16u, vals, sizeof(ushort) * 4) ); - } - void writeScalar(const short* vals) - { - cudaSafeCall( cudaMemcpyToSymbol(scalar_16s, vals, sizeof(short) * 4) ); - } - void writeScalar(const int* vals) - { - cudaSafeCall( cudaMemcpyToSymbol(scalar_32s, vals, sizeof(int) * 4) ); - } - void writeScalar(const float* vals) - { - cudaSafeCall( cudaMemcpyToSymbol(scalar_32f, vals, sizeof(float) * 4) ); - } - void writeScalar(const double* vals) - { - cudaSafeCall( cudaMemcpyToSymbol(scalar_64f, vals, sizeof(double) * 4) ); - } - template - __global__ void set_to_without_mask(T* mat, int cols, int rows, size_t step, int channels) + __global__ void set_to_without_mask(PtrStepSz mat, typename TypeVec::vec_type val, int channels) { - size_t x = blockIdx.x * blockDim.x + threadIdx.x; - size_t y = blockIdx.y * blockDim.y + threadIdx.y; + const int y = blockIdx.x * blockDim.y + threadIdx.y; - if ((x < cols * channels ) && (y < rows)) + if (y < mat.rows) { - size_t idx = y * ( step >> shift_and_sizeof::shift ) + x; - mat[idx] = readScalar(x % channels); + const T vals[] = { + val.x, val.y, val.z, val.w + }; + + T* row = mat.ptr(y); + + for (int x = threadIdx.x; x < mat.cols * channels; x += 32) + { + row[x] = vals[x % channels]; + } } } template - __global__ void set_to_with_mask(T* mat, const uchar* mask, int cols, int rows, size_t step, int channels, size_t step_mask) + __global__ void set_to_with_mask(PtrStepSz mat, const PtrStepb mask, typename TypeVec::vec_type val, int channels) { - size_t x = blockIdx.x * blockDim.x + threadIdx.x; - size_t y = blockIdx.y * blockDim.y + threadIdx.y; + const int y = blockIdx.x * blockDim.y + threadIdx.y; - if ((x < cols * channels ) && (y < rows)) - if (mask[y * step_mask + x / channels] != 0) + if (y < mat.rows) + { + const T vals[] = { + val.x, val.y, val.z, val.w + }; + + T* row = mat.ptr(y); + const uchar* mask_row = mask.ptr(y); + + for (int x = threadIdx.x; x < mat.cols * channels; x += 32) { - size_t idx = y * ( step >> shift_and_sizeof::shift ) + x; - mat[idx] = readScalar(x % channels); + if (mask_row[x / channels]) + { + row[x] = vals[x % channels]; + } } + } } + template void set_to_gpu(PtrStepSzb mat, const T* scalar, PtrStepSzb mask, int channels, cudaStream_t stream) { - writeScalar(scalar); + typedef typename TypeVec::vec_type vec_type; - 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.rows, block.y)); - set_to_with_mask<<>>((T*)mat.data, (uchar*)mask.data, mat.cols, mat.rows, mat.step, channels, mask.step); + set_to_with_mask<<>>(PtrStepSz(mat), mask, VecTraits::make(scalar), channels); cudaSafeCall( cudaGetLastError() ); if (stream == 0) @@ -203,12 +176,12 @@ namespace cv { namespace gpu { namespace device template void set_to_gpu(PtrStepSzb mat, const T* scalar, int channels, cudaStream_t stream) { - writeScalar(scalar); + typedef typename TypeVec::vec_type vec_type; - 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.rows, block.y)); - set_to_without_mask<<>>((T*)mat.data, mat.cols, mat.rows, mat.step, channels); + set_to_without_mask<<>>(PtrStepSz(mat), VecTraits::make(scalar), channels); cudaSafeCall( cudaGetLastError() ); if (stream == 0)