Skip to content

Commit

Permalink
fix GpuMat::setTo implementation
Browse files Browse the repository at this point in the history 
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
  • Loading branch information
Vladislav Vinogradov committed Feb 27, 2015
1 parent 2c9547e commit 4f5d30a
Showing 1 changed file with 40 additions and 67 deletions.
107 changes: 40 additions & 67 deletions modules/dynamicuda/src/cuda/matrix_operations.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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
{
Expand Down Expand Up @@ -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 <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>
__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;
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<T>::shift ) + x;
mat[idx] = readScalar<T>(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<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;
size_t y = blockIdx.y * blockDim.y + threadIdx.y;
const int y = blockIdx.x * blockDim.y + threadIdx.y;

if (y < mat.rows)
{
const T vals[] = {
val.x, val.y, val.z, val.w
};

if ((x < cols * channels ) && (y < rows))
if (mask[y * step_mask + x / channels] != 0)
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;
mat[idx] = readScalar<T>(x % channels);
if (mask_row[x / channels])
{
row[x] = vals[x % channels];
}
}
}
}

template <typename T>
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 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><<<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() );

if (stream == 0)
Expand All @@ -203,12 +176,12 @@ namespace cv { namespace gpu { namespace device
template <typename T>
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 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><<<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() );

if (stream == 0)
Expand Down

0 comments on commit 4f5d30a

Please sign in to comment.