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

Additional ENABLE_DYNAMIC_CUDA option implemented in cmake. Warning fixes and refactoring.

Browse Source
This commit is contained in:
Alexander Smorkalov
2013-12-19 17:41:04 +04:00
parent 037ffcdf99
commit 5a5c82bb1d
6 changed files with 969 additions and 951 deletions
Download Patch File Download Diff File Expand all files Collapse all files

1
CMakeLists.txt
View File

@@ -201,6 +201,7 @@ OCV_OPTION(INSTALL_TO_MANGLED_PATHS "Enables mangled install paths, that help wi
# OpenCV build options
# ===================================================
OCV_OPTION(ENABLE_DYNAMIC_CUDA "Enabled dynamic CUDA linkage" ON IF ANDROID OR LINUX)
OCV_OPTION(ENABLE_PRECOMPILED_HEADERS "Use precompiled headers" ON IF (NOT IOS) )
OCV_OPTION(ENABLE_SOLUTION_FOLDERS "Solution folder in Visual Studio or in other IDEs" (MSVC_IDE OR CMAKE_GENERATOR MATCHES Xcode) IF (CMAKE_VERSION VERSION_GREATER "2.8.0") )
OCV_OPTION(ENABLE_PROFILING "Enable profiling in the GCC compiler (Add flags: -g -pg)" OFF IF CMAKE_COMPILER_IS_GNUCXX )

14
modules/core/CMakeLists.txt
View File

@@ -1,8 +1,12 @@
set(the_description "The Core Functionality")
if (HAVE_opencv_dynamicuda)
message(STATUS "ENABLE_DYNAMIC_CUDA ${ENABLE_DYNAMIC_CUDA}")
if (ENABLE_DYNAMIC_CUDA)
message(STATUS "Using dynamic cuda approach")
ocv_add_module(core PRIVATE_REQUIRED ${ZLIB_LIBRARIES})
else()
message(STATUS "Link CUDA statically")
ocv_add_module(core PRIVATE_REQUIRED ${ZLIB_LIBRARIES} ${CUDA_LIBRARIES} ${CUDA_npp_LIBRARY})
endif()
@@ -12,7 +16,7 @@ if(HAVE_WINRT)
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /ZW /GS /Gm- /AI\"${WINDOWS_SDK_PATH}/References/CommonConfiguration/Neutral\" /AI\"${VISUAL_STUDIO_PATH}/vcpackages\"")
endif()
if(HAVE_opencv_dynamicuda)
if(ENABLE_DYNAMIC_CUDA)
add_definitions(-DDYNAMIC_CUDA_SUPPORT)
else()
add_definitions(-DUSE_CUDA)
@@ -26,18 +30,18 @@ endif()
file(GLOB lib_cuda_hdrs "include/opencv2/${name}/cuda/*.hpp" "include/opencv2/${name}/cuda/*.h")
file(GLOB lib_cuda_hdrs_detail "include/opencv2/${name}/cuda/detail/*.hpp" "include/opencv2/${name}/cuda/detail/*.h")
if (NOT HAVE_opencv_dynamicuda)
if (NOT ENABLE_DYNAMIC_CUDA)
file(GLOB lib_cuda "../dynamicuda/src/cuda/*.cu*")
endif()
source_group("Cuda Headers" FILES ${lib_cuda_hdrs})
source_group("Cuda Headers\\Detail" FILES ${lib_cuda_hdrs_detail})
if (NOT HAVE_opencv_dynamicuda)
if (NOT ENABLE_DYNAMIC_CUDA)
source_group("Src\\Cuda" FILES ${lib_cuda} ${lib_cuda_hdrs})
endif()
if (HAVE_opencv_dynamicuda)
if (ENABLE_DYNAMIC_CUDA)
ocv_glob_module_sources(SOURCES "${opencv_core_BINARY_DIR}/version_string.inc"
HEADERS ${lib_cuda_hdrs} ${lib_cuda_hdrs_detail})
else()

1
modules/dynamicuda/CMakeLists.txt
View File

@@ -5,6 +5,7 @@ endif()
set(the_description "Dynamic CUDA linkage")
add_definitions(-DUSE_CUDA)
ocv_warnings_disable(CMAKE_CXX_FLAGS -Wundef)
ocv_module_include_directories("${OpenCV_SOURCE_DIR}/modules/gpu/include")
set(OPENCV_MODULE_TYPE SHARED)
if (BUILD_FAT_JAVA_LIB)

235
modules/dynamicuda/include/opencv2/dynamicuda/dynamicuda.hpp
View File

@@ -1,9 +1,9 @@
#ifndef __GPUMAT_CUDA_HPP__
#define __GPUMAT_CUDA_HPP__
class DeviceInfoFuncTable
{
public:
class DeviceInfoFuncTable
{
public:
// cv::DeviceInfo
virtual size_t sharedMemPerBlock() const = 0;
virtual void queryMemory(size_t&, size_t&) const = 0;
@@ -37,11 +37,11 @@
virtual void printShortCudaDeviceInfo(int) const = 0;
virtual ~DeviceInfoFuncTable() {};
};
};
class GpuFuncTable
{
public:
class GpuFuncTable
{
public:
virtual ~GpuFuncTable() {}
// GpuMat routines
@@ -60,11 +60,11 @@
virtual void mallocPitch(void** devPtr, size_t* step, size_t width, size_t height) const = 0;
virtual void free(void* devPtr) const = 0;
};
};
class EmptyDeviceInfoFuncTable: public DeviceInfoFuncTable
{
public:
class EmptyDeviceInfoFuncTable: public DeviceInfoFuncTable
{
public:
size_t sharedMemPerBlock() const { throw_nogpu; return 0; }
void queryMemory(size_t&, size_t&) const { throw_nogpu; }
size_t freeMemory() const { throw_nogpu; return 0; }
@@ -98,11 +98,11 @@
void printCudaDeviceInfo(int) const { throw_nogpu; }
void printShortCudaDeviceInfo(int) const { throw_nogpu; }
};
};
class EmptyFuncTable : public GpuFuncTable
{
public:
class EmptyFuncTable : public GpuFuncTable
{
public:
void copy(const Mat&, GpuMat&) const { throw_nogpu; }
void copy(const GpuMat&, Mat&) const { throw_nogpu; }
@@ -117,7 +117,7 @@
void mallocPitch(void**, size_t*, size_t, size_t) const { throw_nogpu; }
void free(void*) const {}
};
};
#if defined(USE_CUDA)
@@ -153,46 +153,46 @@ namespace cv { namespace gpu { namespace device
void convert_gpu(PtrStepSzb src, int sdepth, PtrStepSzb dst, int ddepth, double alpha, double beta, cudaStream_t stream);
}}}
template <typename T> void kernelSetCaller(GpuMat& src, Scalar s, cudaStream_t stream)
{
template <typename T> void kernelSetCaller(GpuMat& src, Scalar s, cudaStream_t stream)
{
Scalar_<T> sf = s;
cv::gpu::device::set_to_gpu(src, sf.val, src.channels(), stream);
}
}
template <typename T> void kernelSetCaller(GpuMat& src, Scalar s, const GpuMat& mask, cudaStream_t stream)
{
template <typename T> void kernelSetCaller(GpuMat& src, Scalar s, const GpuMat& mask, cudaStream_t stream)
{
Scalar_<T> sf = s;
cv::gpu::device::set_to_gpu(src, sf.val, mask, src.channels(), stream);
}
}
template<int n> struct NPPTypeTraits;
template<> struct NPPTypeTraits<CV_8U> { typedef Npp8u npp_type; };
template<> struct NPPTypeTraits<CV_8S> { typedef Npp8s npp_type; };
template<> struct NPPTypeTraits<CV_16U> { typedef Npp16u npp_type; };
template<> struct NPPTypeTraits<CV_16S> { typedef Npp16s npp_type; };
template<> struct NPPTypeTraits<CV_32S> { typedef Npp32s npp_type; };
template<> struct NPPTypeTraits<CV_32F> { typedef Npp32f npp_type; };
template<> struct NPPTypeTraits<CV_64F> { typedef Npp64f npp_type; };
template<int n> struct NPPTypeTraits;
template<> struct NPPTypeTraits<CV_8U> { typedef Npp8u npp_type; };
template<> struct NPPTypeTraits<CV_8S> { typedef Npp8s npp_type; };
template<> struct NPPTypeTraits<CV_16U> { typedef Npp16u npp_type; };
template<> struct NPPTypeTraits<CV_16S> { typedef Npp16s npp_type; };
template<> struct NPPTypeTraits<CV_32S> { typedef Npp32s npp_type; };
template<> struct NPPTypeTraits<CV_32F> { typedef Npp32f npp_type; };
template<> struct NPPTypeTraits<CV_64F> { typedef Npp64f npp_type; };
//////////////////////////////////////////////////////////////////////////
// Convert
//////////////////////////////////////////////////////////////////////////
// Convert
template<int SDEPTH, int DDEPTH> struct NppConvertFunc
{
template<int SDEPTH, int DDEPTH> struct NppConvertFunc
{
typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
typedef typename NPPTypeTraits<DDEPTH>::npp_type dst_t;
typedef NppStatus (*func_ptr)(const src_t* pSrc, int nSrcStep, dst_t* pDst, int nDstStep, NppiSize oSizeROI);
};
template<int DDEPTH> struct NppConvertFunc<CV_32F, DDEPTH>
{
};
template<int DDEPTH> struct NppConvertFunc<CV_32F, DDEPTH>
{
typedef typename NPPTypeTraits<DDEPTH>::npp_type dst_t;
typedef NppStatus (*func_ptr)(const Npp32f* pSrc, int nSrcStep, dst_t* pDst, int nDstStep, NppiSize oSizeROI, NppRoundMode eRoundMode);
};
};
template<int SDEPTH, int DDEPTH, typename NppConvertFunc<SDEPTH, DDEPTH>::func_ptr func> struct NppCvt
{
template<int SDEPTH, int DDEPTH, typename NppConvertFunc<SDEPTH, DDEPTH>::func_ptr func> struct NppCvt
{
typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
typedef typename NPPTypeTraits<DDEPTH>::npp_type dst_t;
@@ -206,10 +206,10 @@ namespace cv { namespace gpu { namespace device
cudaSafeCall( cudaDeviceSynchronize() );
}
};
};
template<int DDEPTH, typename NppConvertFunc<CV_32F, DDEPTH>::func_ptr func> struct NppCvt<CV_32F, DDEPTH, func>
{
template<int DDEPTH, typename NppConvertFunc<CV_32F, DDEPTH>::func_ptr func> struct NppCvt<CV_32F, DDEPTH, func>
{
typedef typename NPPTypeTraits<DDEPTH>::npp_type dst_t;
static void call(const GpuMat& src, GpuMat& dst)
@@ -222,34 +222,34 @@ namespace cv { namespace gpu { namespace device
cudaSafeCall( cudaDeviceSynchronize() );
}
};
};
//////////////////////////////////////////////////////////////////////////
// Set
//////////////////////////////////////////////////////////////////////////
// Set
template<int SDEPTH, int SCN> struct NppSetFunc
{
template<int SDEPTH, int SCN> struct NppSetFunc
{
typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
typedef NppStatus (*func_ptr)(const src_t values[], src_t* pSrc, int nSrcStep, NppiSize oSizeROI);
};
template<int SDEPTH> struct NppSetFunc<SDEPTH, 1>
{
};
template<int SDEPTH> struct NppSetFunc<SDEPTH, 1>
{
typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
typedef NppStatus (*func_ptr)(src_t val, src_t* pSrc, int nSrcStep, NppiSize oSizeROI);
};
template<int SCN> struct NppSetFunc<CV_8S, SCN>
{
};
template<int SCN> struct NppSetFunc<CV_8S, SCN>
{
typedef NppStatus (*func_ptr)(Npp8s values[], Npp8s* pSrc, int nSrcStep, NppiSize oSizeROI);
};
template<> struct NppSetFunc<CV_8S, 1>
{
};
template<> struct NppSetFunc<CV_8S, 1>
{
typedef NppStatus (*func_ptr)(Npp8s val, Npp8s* pSrc, int nSrcStep, NppiSize oSizeROI);
};
};
template<int SDEPTH, int SCN, typename NppSetFunc<SDEPTH, SCN>::func_ptr func> struct NppSet
{
template<int SDEPTH, int SCN, typename NppSetFunc<SDEPTH, SCN>::func_ptr func> struct NppSet
{
typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
static void call(GpuMat& src, Scalar s)
@@ -264,9 +264,9 @@ namespace cv { namespace gpu { namespace device
cudaSafeCall( cudaDeviceSynchronize() );
}
};
template<int SDEPTH, typename NppSetFunc<SDEPTH, 1>::func_ptr func> struct NppSet<SDEPTH, 1, func>
{
};
template<int SDEPTH, typename NppSetFunc<SDEPTH, 1>::func_ptr func> struct NppSet<SDEPTH, 1, func>
{
typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
static void call(GpuMat& src, Scalar s)
@@ -281,23 +281,23 @@ namespace cv { namespace gpu { namespace device
cudaSafeCall( cudaDeviceSynchronize() );
}
};
};
template<int SDEPTH, int SCN> struct NppSetMaskFunc
{
template<int SDEPTH, int SCN> struct NppSetMaskFunc
{
typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
typedef NppStatus (*func_ptr)(const src_t values[], src_t* pSrc, int nSrcStep, NppiSize oSizeROI, const Npp8u* pMask, int nMaskStep);
};
template<int SDEPTH> struct NppSetMaskFunc<SDEPTH, 1>
{
};
template<int SDEPTH> struct NppSetMaskFunc<SDEPTH, 1>
{
typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
typedef NppStatus (*func_ptr)(src_t val, src_t* pSrc, int nSrcStep, NppiSize oSizeROI, const Npp8u* pMask, int nMaskStep);
};
};
template<int SDEPTH, int SCN, typename NppSetMaskFunc<SDEPTH, SCN>::func_ptr func> struct NppSetMask
{
template<int SDEPTH, int SCN, typename NppSetMaskFunc<SDEPTH, SCN>::func_ptr func> struct NppSetMask
{
typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
static void call(GpuMat& src, Scalar s, const GpuMat& mask)
@@ -312,9 +312,9 @@ namespace cv { namespace gpu { namespace device
cudaSafeCall( cudaDeviceSynchronize() );
}
};
template<int SDEPTH, typename NppSetMaskFunc<SDEPTH, 1>::func_ptr func> struct NppSetMask<SDEPTH, 1, func>
{
};
template<int SDEPTH, typename NppSetMaskFunc<SDEPTH, 1>::func_ptr func> struct NppSetMask<SDEPTH, 1, func>
{
typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
static void call(GpuMat& src, Scalar s, const GpuMat& mask)
@@ -329,20 +329,20 @@ namespace cv { namespace gpu { namespace device
cudaSafeCall( cudaDeviceSynchronize() );
}
};
};
//////////////////////////////////////////////////////////////////////////
// CopyMasked
//////////////////////////////////////////////////////////////////////////
// CopyMasked
template<int SDEPTH> struct NppCopyMaskedFunc
{
template<int SDEPTH> struct NppCopyMaskedFunc
{
typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
typedef NppStatus (*func_ptr)(const src_t* pSrc, int nSrcStep, src_t* pDst, int nDstStep, NppiSize oSizeROI, const Npp8u* pMask, int nMaskStep);
};
};
template<int SDEPTH, typename NppCopyMaskedFunc<SDEPTH>::func_ptr func> struct NppCopyMasked
{
template<int SDEPTH, typename NppCopyMaskedFunc<SDEPTH>::func_ptr func> struct NppCopyMasked
{
typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
static void call(const GpuMat& src, GpuMat& dst, const GpuMat& mask, cudaStream_t /*stream*/)
@@ -355,16 +355,24 @@ namespace cv { namespace gpu { namespace device
cudaSafeCall( cudaDeviceSynchronize() );
}
};
};
template <typename T> static inline bool isAligned(const T* ptr, size_t size)
{
template <typename T> static inline bool isAligned(const T* ptr, size_t size)
{
return reinterpret_cast<size_t>(ptr) % size == 0;
}
}
namespace cv { namespace gpu { namespace device
{
void copyWithMask(const GpuMat& src, GpuMat& dst, const GpuMat& mask, cudaStream_t stream = 0)
namespace cv { namespace gpu { namespace device
{
void copyWithMask(const GpuMat& src, GpuMat& dst, const GpuMat& mask, cudaStream_t stream = 0);
void convertTo(const GpuMat& src, GpuMat& dst);
void convertTo(const GpuMat& src, GpuMat& dst, double alpha, double beta, cudaStream_t stream = 0);
void setTo(GpuMat& src, Scalar s, cudaStream_t stream);
void setTo(GpuMat& src, Scalar s, const GpuMat& mask, cudaStream_t stream);
void setTo(GpuMat& src, Scalar s);
void setTo(GpuMat& src, Scalar s, const GpuMat& mask);
void copyWithMask(const GpuMat& src, GpuMat& dst, const GpuMat& mask, cudaStream_t stream)
{
CV_Assert(src.size() == dst.size() && src.type() == dst.type());
CV_Assert(src.size() == mask.size() && mask.depth() == CV_8U && (mask.channels() == 1 || mask.channels() == src.channels()));
@@ -377,7 +385,7 @@ namespace cv { namespace gpu { namespace device
cv::gpu::device::convert_gpu(src.reshape(1), src.depth(), dst.reshape(1), dst.depth(), 1.0, 0.0, 0);
}
void convertTo(const GpuMat& src, GpuMat& dst, double alpha, double beta, cudaStream_t stream = 0)
void convertTo(const GpuMat& src, GpuMat& dst, double alpha, double beta, cudaStream_t stream)
{
cv::gpu::device::convert_gpu(src.reshape(1), src.depth(), dst.reshape(1), dst.depth(), alpha, beta, stream);
}
@@ -417,12 +425,11 @@ namespace cv { namespace gpu { namespace device
{
setTo(src, s, mask, 0);
}
}}}
}}}
class CudaArch
{
public:
class CudaArch
{
public:
CudaArch()
{
fromStr(CUDA_ARCH_BIN, bin);
@@ -461,7 +468,7 @@ namespace cv { namespace gpu { namespace device
}
private:
private:
void fromStr(const string& set_as_str, vector<int>& arr)
{
if (set_as_str.find_first_not_of(" ") == string::npos)
@@ -482,11 +489,11 @@ namespace cv { namespace gpu { namespace device
vector<int> bin;
vector<int> ptx;
vector<int> features;
};
};
class DeviceProps
{
public:
class DeviceProps
{
public:
DeviceProps()
{
props_.resize(10, 0);
@@ -515,15 +522,15 @@ namespace cv { namespace gpu { namespace device
return props_[devID];
}
private:
private:
std::vector<cudaDeviceProp*> props_;
};
};
DeviceProps deviceProps;
DeviceProps deviceProps;
class CudaDeviceInfoFuncTable: DeviceInfoFuncTable
{
public:
class CudaDeviceInfoFuncTable: DeviceInfoFuncTable
{
public:
size_t sharedMemPerBlock() const
{
return deviceProps.get(device_id_)->sharedMemPerBlock;
@@ -817,7 +824,7 @@ namespace cv { namespace gpu { namespace device
fflush(stdout);
}
private:
private:
int device_id_;
std::string name_;
@@ -847,20 +854,22 @@ namespace cv { namespace gpu { namespace device
return -1;
}
};
};
class CudaFuncTable : public GpuFuncTable
{
public:
class CudaFuncTable : public GpuFuncTable
{
public:
void copy(const Mat& src, GpuMat& dst) const
{
cudaSafeCall( cudaMemcpy2D(dst.data, dst.step, src.data, src.step, src.cols * src.elemSize(), src.rows, cudaMemcpyHostToDevice) );
}
void copy(const GpuMat& src, Mat& dst) const
{
cudaSafeCall( cudaMemcpy2D(dst.data, dst.step, src.data, src.step, src.cols * src.elemSize(), src.rows, cudaMemcpyDeviceToHost) );
}
void copy(const GpuMat& src, GpuMat& dst) const
{
cudaSafeCall( cudaMemcpy2D(dst.data, dst.step, src.data, src.step, src.cols * src.elemSize(), src.rows, cudaMemcpyDeviceToDevice) );
@@ -1087,6 +1096,6 @@ namespace cv { namespace gpu { namespace device
{
cudaFree(devPtr);
}
};
};
#endif
#endif

3
modules/dynamicuda/src/main.cpp
View File

@@ -39,6 +39,9 @@ static EmptyFuncTable gpuTable;
extern "C" {
DeviceInfoFuncTable* deviceInfoFactory();
GpuFuncTable* gpuFactory();
DeviceInfoFuncTable* deviceInfoFactory()
{
return (DeviceInfoFuncTable*)&deviceInfoTable;

2
modules/java/CMakeLists.txt
View File

@@ -297,7 +297,7 @@ if(BUILD_FAT_JAVA_LIB)
list(REMOVE_ITEM __deps ${m})
endif()
endforeach()
if (HAVE_opencv_dynamicuda)
if (ENABLE_DYNAMIC_CUDA)
list(REMOVE_ITEM __deps "opencv_dynamicuda")
endif()
if (ANDROID AND HAVE_opencv_gpu)