generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

fix GpuMat::setTo implementation

Browse Source 
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
This commit is contained in:
Vladislav Vinogradov 2015-02-27 16:28:06 +03:00
parent 2c9547e314
commit 4f5d30a865
1 changed files with 40 additions and 67 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

107
modules/dynamicuda/src/cuda/matrix_operations.cu
View File

@ -44,6 +44,7 @@
#include "opencv2/gpu/device/transform.hpp" #include "opencv2/gpu/device/transform.hpp"
#include "opencv2/gpu/device/functional.hpp" #include "opencv2/gpu/device/functional.hpp"
#include "opencv2/gpu/device/type_traits.hpp" #include "opencv2/gpu/device/type_traits.hpp"
#include "opencv2/gpu/device/vec_traits.hpp"
namespace cv { namespace gpu { namespace device namespace cv { namespace gpu { namespace device
{ {
@ -105,87 +106,59 @@ namespace cv { namespace gpu { namespace device
////////////////////////////////// SetTo ////////////////////////////////// ////////////////////////////////// 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 <typename T> __device__ __forceinline__ T readScalar(int i);
template <> __device__ __forceinline__ uchar readScalar<uchar>(int i) {return scalar_8u[i];}
template <> __device__ __forceinline__ schar readScalar<schar>(int i) {return scalar_8s[i];}
template <> __device__ __forceinline__ ushort readScalar<ushort>(int i) {return scalar_16u[i];}
template <> __device__ __forceinline__ short readScalar<short>(int i) {return scalar_16s[i];}
template <> __device__ __forceinline__ int readScalar<int>(int i) {return scalar_32s[i];}
template <> __device__ __forceinline__ float readScalar<float>(int i) {return scalar_32f[i];}
template <> __device__ __forceinline__ double readScalar<double>(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<typename T> template<typename T>
__global__ void set_to_without_mask(T* mat, int cols, int rows, size_t step, int channels) __global__ void set_to_without_mask(PtrStepSz<T> mat, typename TypeVec<T, 4>::vec_type val, int channels)
{ {
size_t x = blockIdx.x * blockDim.x + threadIdx.x; const int y = blockIdx.x * blockDim.y + threadIdx.y;
size_t y = blockIdx.y * blockDim.y + threadIdx.y;
if ((x < cols * channels ) && (y < rows)) if (y < mat.rows)
{ {
size_t idx = y * ( step >> shift_and_sizeof<T>::shift ) + x; const T vals[] = {
mat[idx] = readScalar<T>(x % channels); 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<typename T> template<typename T>
__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<T> mat, const PtrStepb mask, typename TypeVec<T, 4>::vec_type val, int channels)
{ {
size_t x = blockIdx.x * blockDim.x + threadIdx.x; const int y = blockIdx.x * blockDim.y + threadIdx.y;
size_t y = blockIdx.y * blockDim.y + threadIdx.y;
if ((x < cols * channels ) && (y < rows)) if (y < mat.rows)
if (mask[y * step_mask + x / channels] != 0) {
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<T>::shift ) + x; if (mask_row[x / channels])
mat[idx] = readScalar<T>(x % channels); {
row[x] = vals[x % channels];
}
} }
}
} }
template <typename T> template <typename T>
void set_to_gpu(PtrStepSzb mat, const T* scalar, PtrStepSzb mask, int channels, cudaStream_t stream) void set_to_gpu(PtrStepSzb mat, const T* scalar, PtrStepSzb mask, int channels, cudaStream_t stream)
{ {
writeScalar(scalar); typedef typename TypeVec<T, 4>::vec_type vec_type;
dim3 threadsPerBlock(32, 8, 1); dim3 block(32, 8);
dim3 numBlocks (mat.cols * channels / threadsPerBlock.x + 1, mat.rows / threadsPerBlock.y + 1, 1); dim3 grid(divUp(mat.rows, block.y));
set_to_with_mask<T><<<numBlocks, threadsPerBlock, 0, stream>>>((T*)mat.data, (uchar*)mask.data, mat.cols, mat.rows, mat.step, channels, mask.step); set_to_with_mask<T><<<grid, block, 0, stream>>>(PtrStepSz<T>(mat), mask, VecTraits<vec_type>::make(scalar), channels);
cudaSafeCall( cudaGetLastError() ); cudaSafeCall( cudaGetLastError() );
if (stream == 0) if (stream == 0)
@ -203,12 +176,12 @@ namespace cv { namespace gpu { namespace device
template <typename T> template <typename T>
void set_to_gpu(PtrStepSzb mat, const T* scalar, int channels, cudaStream_t stream) void set_to_gpu(PtrStepSzb mat, const T* scalar, int channels, cudaStream_t stream)
{ {
writeScalar(scalar); typedef typename TypeVec<T, 4>::vec_type vec_type;
dim3 threadsPerBlock(32, 8, 1); dim3 block(32, 8);
dim3 numBlocks (mat.cols * channels / threadsPerBlock.x + 1, mat.rows / threadsPerBlock.y + 1, 1); dim3 grid(divUp(mat.rows, block.y));
set_to_without_mask<T><<<numBlocks, threadsPerBlock, 0, stream>>>((T*)mat.data, mat.cols, mat.rows, mat.step, channels); set_to_without_mask<T><<<grid, block, 0, stream>>>(PtrStepSz<T>(mat), VecTraits<vec_type>::make(scalar), channels);
cudaSafeCall( cudaGetLastError() ); cudaSafeCall( cudaGetLastError() );
if (stream == 0) if (stream == 0)