mirror/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2024-11-29 05:29:54 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

Thrust allocator for temporary storage used with thrust algorithms.

Browse Source
This commit is contained in:
Dan 2016-02-18 00:23:20 -05:00
parent e04f372083
commit da93a5527d
2 changed files with 50 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

9
modules/core/include/opencv2/core/cuda/utility.hpp
View File

@ -54,6 +54,15 @@
namespace cv { namespace cuda { namespace device namespace cv { namespace cuda { namespace device
{ {
struct CV_EXPORTS ThrustAllocator
{
typedef uchar value_type;
virtual __device__ __host__ uchar* allocate(size_t numBytes) = 0;
virtual __device__ __host__ void free(uchar* ptr) = 0;
static ThrustAllocator& getAllocator();
static void setAllocator(ThrustAllocator* allocator);
};
#define OPENCV_CUDA_LOG_WARP_SIZE (5) #define OPENCV_CUDA_LOG_WARP_SIZE (5)
#define OPENCV_CUDA_WARP_SIZE (1 << OPENCV_CUDA_LOG_WARP_SIZE) #define OPENCV_CUDA_WARP_SIZE (1 << OPENCV_CUDA_LOG_WARP_SIZE)
#define OPENCV_CUDA_LOG_MEM_BANKS ((__CUDA_ARCH__ >= 200) ? 5 : 4) // 32 banks on fermi, 16 on tesla #define OPENCV_CUDA_LOG_MEM_BANKS ((__CUDA_ARCH__ >= 200) ? 5 : 4) // 32 banks on fermi, 16 on tesla

41
modules/core/src/cuda/gpu_mat.cu
View File

@ -50,11 +50,52 @@
#include "opencv2/core/cuda.hpp" #include "opencv2/core/cuda.hpp"
#include "opencv2/cudev.hpp" #include "opencv2/cudev.hpp"
#include "opencv2/core/cuda/utility.hpp"
using namespace cv; using namespace cv;
using namespace cv::cuda; using namespace cv::cuda;
using namespace cv::cudev; using namespace cv::cudev;
namespace
{
class DefaultThrustAllocator: public cv::cuda::device::ThrustAllocator
{
public:
__device__ __host__ uchar* allocate(size_t numBytes)
{
#ifndef __CUDA_ARCH__
uchar* ptr;
CV_CUDEV_SAFE_CALL(cudaMalloc(&ptr, numBytes));
return ptr;
#else
return NULL;
#endif
}
__device__ __host__ void free(uchar* ptr)
{
#ifndef __CUDA_ARCH__
CV_CUDEV_SAFE_CALL(cudaFree(ptr));
#endif
}
};
DefaultThrustAllocator defaultThrustAllocator;
cv::cuda::device::ThrustAllocator* g_thrustAllocator = &defaultThrustAllocator;
}
cv::cuda::device::ThrustAllocator& cv::cuda::device::ThrustAllocator::getAllocator()
{
return *g_thrustAllocator;
}
void cv::cuda::device::ThrustAllocator::setAllocator(cv::cuda::device::ThrustAllocator* allocator)
{
if(allocator == NULL)
g_thrustAllocator = &defaultThrustAllocator;
else
g_thrustAllocator = allocator;
}
namespace namespace
{ {
class DefaultAllocator : public GpuMat::Allocator class DefaultAllocator : public GpuMat::Allocator