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

added enqueueHostCallback method to gpu::Stream

Browse Source
This commit is contained in:
Vladislav Vinogradov
2013-02-13 15:51:27 +04:00
parent ce2fd7fec9
commit a828b60765
4 changed files with 485 additions and 196 deletions
Download Patch File Download Diff File Expand all files Collapse all files

425
modules/gpu/src/cudastream.cpp
View File

@@ -42,51 +42,37 @@
#include "precomp.hpp"
using namespace std;
using namespace cv;
using namespace cv::gpu;
#if defined HAVE_CUDA
struct Stream::Impl
{
static cudaStream_t getStream(const Impl* impl) { return impl ? impl->stream : 0; }
cudaStream_t stream;
int ref_counter;
};
#include "opencv2/gpu/stream_accessor.hpp"
CV_EXPORTS cudaStream_t cv::gpu::StreamAccessor::getStream(const Stream& stream)
{
return Stream::Impl::getStream(stream.impl);
};
#endif /* !defined (HAVE_CUDA) */
#if !defined (HAVE_CUDA)
void cv::gpu::Stream::create() { throw_nogpu(); }
void cv::gpu::Stream::release() { throw_nogpu(); }
cv::gpu::Stream::Stream() : impl(0) { throw_nogpu(); }
cv::gpu::Stream::~Stream() { throw_nogpu(); }
cv::gpu::Stream::Stream(const Stream& /*stream*/) { throw_nogpu(); }
Stream& cv::gpu::Stream::operator=(const Stream& /*stream*/) { throw_nogpu(); return *this; }
bool cv::gpu::Stream::queryIfComplete() { throw_nogpu(); return true; }
cv::gpu::Stream::Stream() { throw_nogpu(); }
cv::gpu::Stream::~Stream() {}
cv::gpu::Stream::Stream(const Stream&) { throw_nogpu(); }
Stream& cv::gpu::Stream::operator=(const Stream&) { throw_nogpu(); return *this; }
bool cv::gpu::Stream::queryIfComplete() { throw_nogpu(); return false; }
void cv::gpu::Stream::waitForCompletion() { throw_nogpu(); }
void cv::gpu::Stream::enqueueDownload(const GpuMat& /*src*/, Mat& /*dst*/) { throw_nogpu(); }
void cv::gpu::Stream::enqueueDownload(const GpuMat& /*src*/, CudaMem& /*dst*/) { throw_nogpu(); }
void cv::gpu::Stream::enqueueUpload(const CudaMem& /*src*/, GpuMat& /*dst*/) { throw_nogpu(); }
void cv::gpu::Stream::enqueueUpload(const Mat& /*src*/, GpuMat& /*dst*/) { throw_nogpu(); }
void cv::gpu::Stream::enqueueCopy(const GpuMat& /*src*/, GpuMat& /*dst*/) { throw_nogpu(); }
void cv::gpu::Stream::enqueueMemSet(GpuMat& /*src*/, Scalar /*val*/) { throw_nogpu(); }
void cv::gpu::Stream::enqueueMemSet(GpuMat& /*src*/, Scalar /*val*/, const GpuMat& /*mask*/) { throw_nogpu(); }
void cv::gpu::Stream::enqueueConvert(const GpuMat& /*src*/, GpuMat& /*dst*/, int /*type*/, double /*a*/, double /*b*/) { throw_nogpu(); }
void cv::gpu::Stream::enqueueDownload(const GpuMat&, Mat&) { throw_nogpu(); }
void cv::gpu::Stream::enqueueDownload(const GpuMat&, CudaMem&) { throw_nogpu(); }
void cv::gpu::Stream::enqueueUpload(const CudaMem&, GpuMat&) { throw_nogpu(); }
void cv::gpu::Stream::enqueueUpload(const Mat&, GpuMat&) { throw_nogpu(); }
void cv::gpu::Stream::enqueueCopy(const GpuMat&, GpuMat&) { throw_nogpu(); }
void cv::gpu::Stream::enqueueMemSet(GpuMat&, Scalar) { throw_nogpu(); }
void cv::gpu::Stream::enqueueMemSet(GpuMat&, Scalar, const GpuMat&) { throw_nogpu(); }
void cv::gpu::Stream::enqueueConvert(const GpuMat&, GpuMat&, int, double, double) { throw_nogpu(); }
void cv::gpu::Stream::enqueueHostCallback(StreamCallback, void*) { throw_nogpu(); }
Stream& cv::gpu::Stream::Null() { throw_nogpu(); static Stream s; return s; }
cv::gpu::Stream::operator bool() const { throw_nogpu(); return false; }
cv::gpu::Stream::Stream(Impl*) { throw_nogpu(); }
void cv::gpu::Stream::create() { throw_nogpu(); }
void cv::gpu::Stream::release() { throw_nogpu(); }
#else /* !defined (HAVE_CUDA) */
#include "opencv2/gpu/stream_accessor.hpp"
namespace cv { namespace gpu
{
void copyWithMask(const GpuMat& src, GpuMat& dst, const GpuMat& mask, cudaStream_t stream);
@@ -95,14 +81,247 @@ namespace cv { namespace gpu
void setTo(GpuMat& src, Scalar s, const GpuMat& mask, cudaStream_t stream);
}}
struct Stream::Impl
{
static cudaStream_t getStream(const Impl* impl)
{
return impl ? impl->stream : 0;
}
cudaStream_t stream;
int ref_counter;
};
cudaStream_t cv::gpu::StreamAccessor::getStream(const Stream& stream)
{
return Stream::Impl::getStream(stream.impl);
}
cv::gpu::Stream::Stream() : impl(0)
{
create();
}
cv::gpu::Stream::~Stream()
{
release();
}
cv::gpu::Stream::Stream(const Stream& stream) : impl(stream.impl)
{
if (impl)
CV_XADD(&impl->ref_counter, 1);
}
Stream& cv::gpu::Stream::operator =(const Stream& stream)
{
if (this != &stream)
{
release();
impl = stream.impl;
if (impl)
CV_XADD(&impl->ref_counter, 1);
}
return *this;
}
bool cv::gpu::Stream::queryIfComplete()
{
cudaStream_t stream = Impl::getStream(impl);
cudaError_t err = cudaStreamQuery(stream);
if (err == cudaErrorNotReady || err == cudaSuccess)
return err == cudaSuccess;
cudaSafeCall(err);
return false;
}
void cv::gpu::Stream::waitForCompletion()
{
cudaStream_t stream = Impl::getStream(impl);
cudaSafeCall( cudaStreamSynchronize(stream) );
}
void cv::gpu::Stream::enqueueDownload(const GpuMat& src, Mat& dst)
{
// if not -> allocation will be done, but after that dst will not point to page locked memory
CV_Assert( src.size() == dst.size() && src.type() == dst.type() );
cudaStream_t stream = Impl::getStream(impl);
size_t bwidth = src.cols * src.elemSize();
cudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyDeviceToHost, stream) );
}
void cv::gpu::Stream::enqueueDownload(const GpuMat& src, CudaMem& dst)
{
dst.create(src.size(), src.type(), CudaMem::ALLOC_PAGE_LOCKED);
cudaStream_t stream = Impl::getStream(impl);
size_t bwidth = src.cols * src.elemSize();
cudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyDeviceToHost, stream) );
}
void cv::gpu::Stream::enqueueUpload(const CudaMem& src, GpuMat& dst)
{
dst.create(src.size(), src.type());
cudaStream_t stream = Impl::getStream(impl);
size_t bwidth = src.cols * src.elemSize();
cudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyHostToDevice, stream) );
}
void cv::gpu::Stream::enqueueUpload(const Mat& src, GpuMat& dst)
{
dst.create(src.size(), src.type());
cudaStream_t stream = Impl::getStream(impl);
size_t bwidth = src.cols * src.elemSize();
cudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyHostToDevice, stream) );
}
void cv::gpu::Stream::enqueueCopy(const GpuMat& src, GpuMat& dst)
{
dst.create(src.size(), src.type());
cudaStream_t stream = Impl::getStream(impl);
size_t bwidth = src.cols * src.elemSize();
cudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyDeviceToDevice, stream) );
}
void cv::gpu::Stream::enqueueMemSet(GpuMat& src, Scalar val)
{
const int sdepth = src.depth();
if (sdepth == CV_64F)
{
if (!deviceSupports(NATIVE_DOUBLE))
CV_Error(CV_StsUnsupportedFormat, "The device doesn't support double");
}
cudaStream_t stream = Impl::getStream(impl);
if (val[0] == 0.0 && val[1] == 0.0 && val[2] == 0.0 && val[3] == 0.0)
{
cudaSafeCall( cudaMemset2DAsync(src.data, src.step, 0, src.cols * src.elemSize(), src.rows, stream) );
return;
}
if (sdepth == CV_8U)
{
int cn = src.channels();
if (cn == 1 || (cn == 2 && val[0] == val[1]) || (cn == 3 && val[0] == val[1] && val[0] == val[2]) || (cn == 4 && val[0] == val[1] && val[0] == val[2] && val[0] == val[3]))
{
int ival = saturate_cast<uchar>(val[0]);
cudaSafeCall( cudaMemset2DAsync(src.data, src.step, ival, src.cols * src.elemSize(), src.rows, stream) );
return;
}
}
setTo(src, val, stream);
}
void cv::gpu::Stream::enqueueMemSet(GpuMat& src, Scalar val, const GpuMat& mask)
{
const int sdepth = src.depth();
if (sdepth == CV_64F)
{
if (!deviceSupports(NATIVE_DOUBLE))
CV_Error(CV_StsUnsupportedFormat, "The device doesn't support double");
}
CV_Assert(mask.type() == CV_8UC1);
cudaStream_t stream = Impl::getStream(impl);
setTo(src, val, mask, stream);
}
void cv::gpu::Stream::enqueueConvert(const GpuMat& src, GpuMat& dst, int dtype, double alpha, double beta)
{
if (dtype < 0)
dtype = src.type();
else
dtype = CV_MAKE_TYPE(CV_MAT_DEPTH(dtype), src.channels());
const int sdepth = src.depth();
const int ddepth = CV_MAT_DEPTH(dtype);
if (sdepth == CV_64F || ddepth == CV_64F)
{
if (!deviceSupports(NATIVE_DOUBLE))
CV_Error(CV_StsUnsupportedFormat, "The device doesn't support double");
}
bool noScale = fabs(alpha - 1) < numeric_limits<double>::epsilon() && fabs(beta) < numeric_limits<double>::epsilon();
if (sdepth == ddepth && noScale)
{
enqueueCopy(src, dst);
return;
}
dst.create(src.size(), dtype);
cudaStream_t stream = Impl::getStream(impl);
convertTo(src, dst, alpha, beta, stream);
}
#if CUDA_VERSION >= 5000
namespace
{
template<class S, class D> void devcopy(const S& src, D& dst, cudaStream_t s, cudaMemcpyKind k)
struct CallbackData
{
dst.create(src.size(), src.type());
size_t bwidth = src.cols * src.elemSize();
cudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, k, s) );
cv::gpu::Stream::StreamCallback callback;
void* userData;
Stream stream;
};
void CUDART_CB cudaStreamCallback(cudaStream_t, cudaError_t status, void* userData)
{
CallbackData* data = reinterpret_cast<CallbackData*>(userData);
data->callback(data->stream, static_cast<int>(status), data->userData);
delete data;
}
}
#endif
void cv::gpu::Stream::enqueueHostCallback(StreamCallback callback, void* userData)
{
#if CUDA_VERSION >= 5000
CallbackData* data = new CallbackData;
data->callback = callback;
data->userData = userData;
data->stream = *this;
cudaStream_t stream = Impl::getStream(impl);
cudaSafeCall( cudaStreamAddCallback(stream, cudaStreamCallback, data, 0) );
#else
(void) callback;
(void) userData;
CV_Error(CV_StsNotImplemented, "This function requires CUDA 5.0");
#endif
}
cv::gpu::Stream& cv::gpu::Stream::Null()
{
static Stream s((Impl*) 0);
return s;
}
cv::gpu::Stream::operator bool() const
{
return impl && impl->stream;
}
cv::gpu::Stream::Stream(Impl* impl_) : impl(impl_)
{
}
void cv::gpu::Stream::create()
@@ -113,7 +332,7 @@ void cv::gpu::Stream::create()
cudaStream_t stream;
cudaSafeCall( cudaStreamCreate( &stream ) );
impl = (Stream::Impl*)fastMalloc(sizeof(Stream::Impl));
impl = (Stream::Impl*) fastMalloc(sizeof(Stream::Impl));
impl->stream = stream;
impl->ref_counter = 1;
@@ -121,133 +340,11 @@ void cv::gpu::Stream::create()
void cv::gpu::Stream::release()
{
if( impl && CV_XADD(&impl->ref_counter, -1) == 1 )
if (impl && CV_XADD(&impl->ref_counter, -1) == 1)
{
cudaSafeCall( cudaStreamDestroy( impl->stream ) );
cv::fastFree( impl );
cudaSafeCall( cudaStreamDestroy(impl->stream) );
cv::fastFree(impl);
}
}
cv::gpu::Stream::Stream() : impl(0) { create(); }
cv::gpu::Stream::~Stream() { release(); }
cv::gpu::Stream::Stream(const Stream& stream) : impl(stream.impl)
{
if( impl )
CV_XADD(&impl->ref_counter, 1);
}
Stream& cv::gpu::Stream::operator=(const Stream& stream)
{
if( this != &stream )
{
if( stream.impl )
CV_XADD(&stream.impl->ref_counter, 1);
release();
impl = stream.impl;
}
return *this;
}
bool cv::gpu::Stream::queryIfComplete()
{
cudaError_t err = cudaStreamQuery( Impl::getStream(impl) );
if (err == cudaErrorNotReady || err == cudaSuccess)
return err == cudaSuccess;
cudaSafeCall(err);
return false;
}
void cv::gpu::Stream::waitForCompletion() { cudaSafeCall( cudaStreamSynchronize( Impl::getStream(impl) ) ); }
void cv::gpu::Stream::enqueueDownload(const GpuMat& src, Mat& dst)
{
// if not -> allocation will be done, but after that dst will not point to page locked memory
CV_Assert(src.cols == dst.cols && src.rows == dst.rows && src.type() == dst.type() );
devcopy(src, dst, Impl::getStream(impl), cudaMemcpyDeviceToHost);
}
void cv::gpu::Stream::enqueueDownload(const GpuMat& src, CudaMem& dst) { devcopy(src, dst, Impl::getStream(impl), cudaMemcpyDeviceToHost); }
void cv::gpu::Stream::enqueueUpload(const CudaMem& src, GpuMat& dst){ devcopy(src, dst, Impl::getStream(impl), cudaMemcpyHostToDevice); }
void cv::gpu::Stream::enqueueUpload(const Mat& src, GpuMat& dst) { devcopy(src, dst, Impl::getStream(impl), cudaMemcpyHostToDevice); }
void cv::gpu::Stream::enqueueCopy(const GpuMat& src, GpuMat& dst) { devcopy(src, dst, Impl::getStream(impl), cudaMemcpyDeviceToDevice); }
void cv::gpu::Stream::enqueueMemSet(GpuMat& src, Scalar s)
{
CV_Assert((src.depth() != CV_64F) ||
(TargetArchs::builtWith(NATIVE_DOUBLE) && DeviceInfo().supports(NATIVE_DOUBLE)));
if (s[0] == 0.0 && s[1] == 0.0 && s[2] == 0.0 && s[3] == 0.0)
{
cudaSafeCall( cudaMemset2DAsync(src.data, src.step, 0, src.cols * src.elemSize(), src.rows, Impl::getStream(impl)) );
return;
}
if (src.depth() == CV_8U)
{
int cn = src.channels();
if (cn == 1 || (cn == 2 && s[0] == s[1]) || (cn == 3 && s[0] == s[1] && s[0] == s[2]) || (cn == 4 && s[0] == s[1] && s[0] == s[2] && s[0] == s[3]))
{
int val = saturate_cast<uchar>(s[0]);
cudaSafeCall( cudaMemset2DAsync(src.data, src.step, val, src.cols * src.elemSize(), src.rows, Impl::getStream(impl)) );
return;
}
}
setTo(src, s, Impl::getStream(impl));
}
void cv::gpu::Stream::enqueueMemSet(GpuMat& src, Scalar val, const GpuMat& mask)
{
CV_Assert((src.depth() != CV_64F) ||
(TargetArchs::builtWith(NATIVE_DOUBLE) && DeviceInfo().supports(NATIVE_DOUBLE)));
CV_Assert(mask.type() == CV_8UC1);
setTo(src, val, mask, Impl::getStream(impl));
}
void cv::gpu::Stream::enqueueConvert(const GpuMat& src, GpuMat& dst, int rtype, double alpha, double beta)
{
CV_Assert((src.depth() != CV_64F && CV_MAT_DEPTH(rtype) != CV_64F) ||
(TargetArchs::builtWith(NATIVE_DOUBLE) && DeviceInfo().supports(NATIVE_DOUBLE)));
bool noScale = fabs(alpha-1) < std::numeric_limits<double>::epsilon() && fabs(beta) < std::numeric_limits<double>::epsilon();
if( rtype < 0 )
rtype = src.type();
else
rtype = CV_MAKETYPE(CV_MAT_DEPTH(rtype), src.channels());
int sdepth = src.depth(), ddepth = CV_MAT_DEPTH(rtype);
if( sdepth == ddepth && noScale )
{
src.copyTo(dst);
return;
}
GpuMat temp;
const GpuMat* psrc = &src;
if( sdepth != ddepth && psrc == &dst )
psrc = &(temp = src);
dst.create( src.size(), rtype );
convertTo(src, dst, alpha, beta, Impl::getStream(impl));
}
cv::gpu::Stream::operator bool() const
{
return impl && impl->stream;
}
cv::gpu::Stream::Stream(Impl* impl_) : impl(impl_) {}
cv::gpu::Stream& cv::gpu::Stream::Null()
{
static Stream s((Impl*)0);
return s;
}
#endif /* !defined (HAVE_CUDA) */