Stream(const Ptr<Impl>& impl);
friend struct StreamAccessor;
+ friend class BufferPool;
};
class CV_EXPORTS Event
{
class MemoryStack;
- class CV_EXPORTS BufferAllocator : public GpuMat::Allocator
+ class CV_EXPORTS StackAllocator : public GpuMat::Allocator
{
public:
- explicit BufferAllocator(Stream& stream);
- ~BufferAllocator();
+ explicit StackAllocator(cudaStream_t stream);
+ ~StackAllocator();
bool allocate(uchar** devPtr, size_t* step, int** refcount, int rows, int cols, size_t elemSize);
void free(uchar* devPtr, int* refcount);
private:
- BufferAllocator(const BufferAllocator&);
- BufferAllocator& operator =(const BufferAllocator&);
+ StackAllocator(const StackAllocator&);
+ StackAllocator& operator =(const StackAllocator&);
+ cudaStream_t stream_;
MemoryStack* memStack_;
- Stream stream_;
size_t alignment_;
};
+ class CV_EXPORTS BufferPool
+ {
+ public:
+ explicit BufferPool(Stream& stream);
+
+ GpuMat getBuffer(int rows, int cols, int type);
+ GpuMat getBuffer(Size size, int type) { return getBuffer(size.height, size.width, type); }
+
+ private:
+ GpuMat::Allocator* allocator_;
+ };
+
CV_EXPORTS void setBufferAllocatorUsage(bool on);
CV_EXPORTS void allocateMemoryPool(int deviceId, size_t stackSize, int stackCount);
}
/////////////////////////////////////////////////////////////
-/// BufferAllocator
+/// StackAllocator
namespace
{
bool enableMemoryPool = true;
}
-cv::cuda::BufferAllocator::BufferAllocator(Stream& stream) : memStack_(0), stream_(stream)
+cv::cuda::StackAllocator::StackAllocator(cudaStream_t stream) : stream_(stream), memStack_(0)
{
if (enableMemoryPool)
{
}
}
-namespace
+cv::cuda::StackAllocator::~StackAllocator()
{
- void CUDART_CB returnMemStackCallback(cudaStream_t, cudaError_t, void* userData)
- {
- MemoryStack* memStack = static_cast<MemoryStack*>(userData);
- memStack->pool->returnMemStack(memStack);
- }
-}
+ cudaStreamSynchronize(stream_);
-cv::cuda::BufferAllocator::~BufferAllocator()
-{
if (memStack_ != 0)
- CV_CUDEV_SAFE_CALL( cudaStreamAddCallback(StreamAccessor::getStream(stream_), returnMemStackCallback, memStack_, 0) );
+ memStack_->pool->returnMemStack(memStack_);
}
namespace
}
}
-bool cv::cuda::BufferAllocator::allocate(uchar** devPtr, size_t* step, int** refcount, int rows, int cols, size_t elemSize)
+bool cv::cuda::StackAllocator::allocate(uchar** devPtr, size_t* step, int** refcount, int rows, int cols, size_t elemSize)
{
if (memStack_ == 0)
return false;
return true;
}
-void cv::cuda::BufferAllocator::free(uchar* devPtr, int* refcount)
+void cv::cuda::StackAllocator::free(uchar* devPtr, int* refcount)
{
if (memStack_ == 0)
return;
setDevice(currentDevice);
}
+/////////////////////////////////////////////////////////////
+/// BufferPool
+
+GpuMat cv::cuda::BufferPool::getBuffer(int rows, int cols, int type)
+{
+ GpuMat buf(allocator_);
+ buf.create(rows, cols, type);
+ return buf;
+}
+
#endif
{
public:
cudaStream_t stream;
+ Ptr<StackAllocator> stackAllocator_;
Impl();
Impl(cudaStream_t stream);
~Impl();
};
+cv::cuda::BufferPool::BufferPool(Stream& stream) : allocator_(stream.impl_->stackAllocator_.get())
+{
+}
+
cv::cuda::Stream::Impl::Impl() : stream(0)
{
cudaSafeCall( cudaStreamCreate(&stream) );
+
+ stackAllocator_ = makePtr<StackAllocator>(stream);
}
cv::cuda::Stream::Impl::Impl(cudaStream_t stream_) : stream(stream_)
{
+ stackAllocator_ = makePtr<StackAllocator>(stream);
}
cv::cuda::Stream::Impl::~Impl()
{
+ stackAllocator_.release();
+
if (stream)
cudaStreamDestroy(stream);
}
{
void func1(const GpuMat& src, GpuMat& dst, Stream& stream)
{
- BufferAllocator bufAlloc(stream);
+ BufferPool pool(stream);
- GpuMat buf(&bufAlloc);
+ GpuMat buf = pool.getBuffer(src.size(), CV_32FC(src.channels()));
src.convertTo(buf, CV_32F, 1.0 / 255.0, stream);
void func2(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream)
{
- BufferAllocator bufAlloc(stream);
+ BufferPool pool(stream);
- GpuMat buf1(&bufAlloc);
+ GpuMat buf1 = pool.getBuffer(src1.size(), CV_32FC(src1.channels()));
func1(src1, buf1, stream);
- GpuMat buf2(&bufAlloc);
+ GpuMat buf2 = pool.getBuffer(src2.size(), CV_32FC(src2.channels()));
func1(src2, buf2, stream);
using namespace cv;
using namespace cv::cuda;
-struct BufferPool : TestWithParam<DeviceInfo>
+struct BufferPoolTest : TestWithParam<DeviceInfo>
{
};
{
void func1(const GpuMat& src, GpuMat& dst, Stream& stream)
{
- BufferAllocator bufAlloc(stream);
+ BufferPool pool(stream);
- GpuMat buf(&bufAlloc);
+ GpuMat buf = pool.getBuffer(src.size(), CV_32FC(src.channels()));
src.convertTo(buf, CV_32F, 1.0 / 255.0, stream);
void func2(const GpuMat& src, GpuMat& dst, Stream& stream)
{
- BufferAllocator bufAlloc(stream);
+ BufferPool pool(stream);
- GpuMat buf1(&bufAlloc);
+ GpuMat buf1 = pool.getBuffer(saturate_cast<int>(src.rows * 0.5), saturate_cast<int>(src.cols * 0.5), src.type());
cuda::resize(src, buf1, Size(), 0.5, 0.5, cv::INTER_NEAREST, stream);
- GpuMat buf2(&bufAlloc);
+ GpuMat buf2 = pool.getBuffer(buf1.size(), CV_32FC(buf1.channels()));
func1(buf1, buf2, stream);
- GpuMat buf3(&bufAlloc);
+ GpuMat buf3 = pool.getBuffer(src.size(), buf2.type());
cuda::resize(buf2, buf3, src.size(), 0, 0, cv::INTER_NEAREST, stream);
}
}
-CUDA_TEST_P(BufferPool, Test)
+CUDA_TEST_P(BufferPoolTest, SimpleUsage)
{
DeviceInfo devInfo = GetParam();
setDevice(devInfo.deviceID());
ASSERT_MAT_NEAR(dst_gold, dst, 0);
}
-INSTANTIATE_TEST_CASE_P(CUDA_Stream, BufferPool, ALL_DEVICES);
+INSTANTIATE_TEST_CASE_P(CUDA_Stream, BufferPoolTest, ALL_DEVICES);
#endif // HAVE_CUDA