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

Merge pull request #1299 from jet47:gpu-cuda-rename

Browse Source
This commit is contained in:
Alexander Smorkalov
2013-09-23 10:31:46 +04:00
committed by OpenCV Buildbot
parent c0c575d68e 6bbac2a7d9
commit 298a1d50d2
601 changed files with 8873 additions and 9711 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
modules/core/include/opencv2/core/base.hpp
View File

@@ -490,7 +490,7 @@ namespace ogl
class CV_EXPORTS Arrays;
}
namespace gpu
namespace cuda
{
class CV_EXPORTS GpuMat;
class CV_EXPORTS CudaMem;

20
modules/core/include/opencv2/core/gpu.hpp → modules/core/include/opencv2/core/cuda.hpp
View File

@@ -41,17 +41,17 @@
//
//M*/
#ifndef __OPENCV_CORE_GPU_HPP__
#define __OPENCV_CORE_GPU_HPP__
#ifndef __OPENCV_CORE_CUDA_HPP__
#define __OPENCV_CORE_CUDA_HPP__
#ifndef __cplusplus
# error gpu.hpp header must be compiled as C++
# error cuda.hpp header must be compiled as C++
#endif
#include "opencv2/core.hpp"
#include "opencv2/core/gpu_types.hpp"
#include "opencv2/core/cuda_types.hpp"
namespace cv { namespace gpu {
namespace cv { namespace cuda {
//////////////////////////////// GpuMat ///////////////////////////////
@@ -453,7 +453,7 @@ enum FeatureSet
//! checks whether current device supports the given feature
CV_EXPORTS bool deviceSupports(FeatureSet feature_set);
//! information about what GPU archs this OpenCV GPU module was compiled for
//! information about what GPU archs this OpenCV CUDA module was compiled for
class CV_EXPORTS TargetArchs
{
public:
@@ -654,7 +654,7 @@ public:
//! checks whether device supports the given feature
bool supports(FeatureSet feature_set) const;
//! checks whether the GPU module can be run on the given device
//! checks whether the CUDA module can be run on the given device
bool isCompatible() const;
private:
@@ -664,9 +664,9 @@ private:
CV_EXPORTS void printCudaDeviceInfo(int device);
CV_EXPORTS void printShortCudaDeviceInfo(int device);
}} // namespace cv { namespace gpu {
}} // namespace cv { namespace cuda {
#include "opencv2/core/gpu.inl.hpp"
#include "opencv2/core/cuda.inl.hpp"
#endif /* __OPENCV_CORE_GPU_HPP__ */
#endif /* __OPENCV_CORE_CUDA_HPP__ */

14
modules/core/include/opencv2/core/gpu.inl.hpp → modules/core/include/opencv2/core/cuda.inl.hpp
View File

@@ -41,12 +41,12 @@
//
//M*/
#ifndef __OPENCV_CORE_GPUINL_HPP__
#define __OPENCV_CORE_GPUINL_HPP__
#ifndef __OPENCV_CORE_CUDAINL_HPP__
#define __OPENCV_CORE_CUDAINL_HPP__
#include "opencv2/core/gpu.hpp"
#include "opencv2/core/cuda.hpp"
namespace cv { namespace gpu {
namespace cv { namespace cuda {
//////////////////////////////// GpuMat ///////////////////////////////
@@ -587,14 +587,14 @@ bool DeviceInfo::supports(FeatureSet feature_set) const
return version >= feature_set;
}
}} // namespace cv { namespace gpu {
}} // namespace cv { namespace cuda {
//////////////////////////////// Mat ////////////////////////////////
namespace cv {
inline
Mat::Mat(const gpu::GpuMat& m)
Mat::Mat(const cuda::GpuMat& m)
: flags(0), dims(0), rows(0), cols(0), data(0), refcount(0), datastart(0), dataend(0), datalimit(0), allocator(0), size(&rows)
{
m.download(*this);
@@ -602,4 +602,4 @@ Mat::Mat(const gpu::GpuMat& m)
}
#endif // __OPENCV_CORE_GPUINL_HPP__
#endif // __OPENCV_CORE_CUDAINL_HPP__

8
modules/core/include/opencv2/core/cuda/block.hpp
View File

@@ -40,10 +40,10 @@
//
//M*/
#ifndef __OPENCV_GPU_DEVICE_BLOCK_HPP__
#define __OPENCV_GPU_DEVICE_BLOCK_HPP__
#ifndef __OPENCV_CUDA_DEVICE_BLOCK_HPP__
#define __OPENCV_CUDA_DEVICE_BLOCK_HPP__
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
struct Block
{
@@ -200,4 +200,4 @@ namespace cv { namespace gpu { namespace cudev
};
}}}
#endif /* __OPENCV_GPU_DEVICE_BLOCK_HPP__ */
#endif /* __OPENCV_CUDA_DEVICE_BLOCK_HPP__ */

10
modules/core/include/opencv2/core/cuda/border_interpolate.hpp
View File

@@ -40,14 +40,14 @@
//
//M*/
#ifndef __OPENCV_GPU_BORDER_INTERPOLATE_HPP__
#define __OPENCV_GPU_BORDER_INTERPOLATE_HPP__
#ifndef __OPENCV_CUDA_BORDER_INTERPOLATE_HPP__
#define __OPENCV_CUDA_BORDER_INTERPOLATE_HPP__
#include "saturate_cast.hpp"
#include "vec_traits.hpp"
#include "vec_math.hpp"
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
//////////////////////////////////////////////////////////////
// BrdConstant
@@ -709,6 +709,6 @@ namespace cv { namespace gpu { namespace cudev
int width;
D val;
};
}}} // namespace cv { namespace gpu { namespace cudev
}}} // namespace cv { namespace cuda { namespace cudev
#endif // __OPENCV_GPU_BORDER_INTERPOLATE_HPP__
#endif // __OPENCV_CUDA_BORDER_INTERPOLATE_HPP__

406
modules/core/include/opencv2/core/cuda/color.hpp
View File

@@ -40,262 +40,262 @@
//
//M*/
#ifndef __OPENCV_GPU_COLOR_HPP__
#define __OPENCV_GPU_COLOR_HPP__
#ifndef __OPENCV_CUDA_COLOR_HPP__
#define __OPENCV_CUDA_COLOR_HPP__
#include "detail/color_detail.hpp"
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
// All OPENCV_GPU_IMPLEMENT_*_TRAITS(ColorSpace1_to_ColorSpace2, ...) macros implements
// All OPENCV_CUDA_IMPLEMENT_*_TRAITS(ColorSpace1_to_ColorSpace2, ...) macros implements
// template <typename T> class ColorSpace1_to_ColorSpace2_traits
// {
// typedef ... functor_type;
// static __host__ __device__ functor_type create_functor();
// };
OPENCV_GPU_IMPLEMENT_RGB2RGB_TRAITS(bgr_to_rgb, 3, 3, 2)
OPENCV_GPU_IMPLEMENT_RGB2RGB_TRAITS(bgr_to_bgra, 3, 4, 0)
OPENCV_GPU_IMPLEMENT_RGB2RGB_TRAITS(bgr_to_rgba, 3, 4, 2)
OPENCV_GPU_IMPLEMENT_RGB2RGB_TRAITS(bgra_to_bgr, 4, 3, 0)
OPENCV_GPU_IMPLEMENT_RGB2RGB_TRAITS(bgra_to_rgb, 4, 3, 2)
OPENCV_GPU_IMPLEMENT_RGB2RGB_TRAITS(bgra_to_rgba, 4, 4, 2)
OPENCV_CUDA_IMPLEMENT_RGB2RGB_TRAITS(bgr_to_rgb, 3, 3, 2)
OPENCV_CUDA_IMPLEMENT_RGB2RGB_TRAITS(bgr_to_bgra, 3, 4, 0)
OPENCV_CUDA_IMPLEMENT_RGB2RGB_TRAITS(bgr_to_rgba, 3, 4, 2)
OPENCV_CUDA_IMPLEMENT_RGB2RGB_TRAITS(bgra_to_bgr, 4, 3, 0)
OPENCV_CUDA_IMPLEMENT_RGB2RGB_TRAITS(bgra_to_rgb, 4, 3, 2)
OPENCV_CUDA_IMPLEMENT_RGB2RGB_TRAITS(bgra_to_rgba, 4, 4, 2)
#undef OPENCV_GPU_IMPLEMENT_RGB2RGB_TRAITS
#undef OPENCV_CUDA_IMPLEMENT_RGB2RGB_TRAITS
OPENCV_GPU_IMPLEMENT_RGB2RGB5x5_TRAITS(bgr_to_bgr555, 3, 0, 5)
OPENCV_GPU_IMPLEMENT_RGB2RGB5x5_TRAITS(bgr_to_bgr565, 3, 0, 6)
OPENCV_GPU_IMPLEMENT_RGB2RGB5x5_TRAITS(rgb_to_bgr555, 3, 2, 5)
OPENCV_GPU_IMPLEMENT_RGB2RGB5x5_TRAITS(rgb_to_bgr565, 3, 2, 6)
OPENCV_GPU_IMPLEMENT_RGB2RGB5x5_TRAITS(bgra_to_bgr555, 4, 0, 5)
OPENCV_GPU_IMPLEMENT_RGB2RGB5x5_TRAITS(bgra_to_bgr565, 4, 0, 6)
OPENCV_GPU_IMPLEMENT_RGB2RGB5x5_TRAITS(rgba_to_bgr555, 4, 2, 5)
OPENCV_GPU_IMPLEMENT_RGB2RGB5x5_TRAITS(rgba_to_bgr565, 4, 2, 6)
OPENCV_CUDA_IMPLEMENT_RGB2RGB5x5_TRAITS(bgr_to_bgr555, 3, 0, 5)
OPENCV_CUDA_IMPLEMENT_RGB2RGB5x5_TRAITS(bgr_to_bgr565, 3, 0, 6)
OPENCV_CUDA_IMPLEMENT_RGB2RGB5x5_TRAITS(rgb_to_bgr555, 3, 2, 5)
OPENCV_CUDA_IMPLEMENT_RGB2RGB5x5_TRAITS(rgb_to_bgr565, 3, 2, 6)
OPENCV_CUDA_IMPLEMENT_RGB2RGB5x5_TRAITS(bgra_to_bgr555, 4, 0, 5)
OPENCV_CUDA_IMPLEMENT_RGB2RGB5x5_TRAITS(bgra_to_bgr565, 4, 0, 6)
OPENCV_CUDA_IMPLEMENT_RGB2RGB5x5_TRAITS(rgba_to_bgr555, 4, 2, 5)
OPENCV_CUDA_IMPLEMENT_RGB2RGB5x5_TRAITS(rgba_to_bgr565, 4, 2, 6)
#undef OPENCV_GPU_IMPLEMENT_RGB2RGB5x5_TRAITS
#undef OPENCV_CUDA_IMPLEMENT_RGB2RGB5x5_TRAITS
OPENCV_GPU_IMPLEMENT_RGB5x52RGB_TRAITS(bgr555_to_rgb, 3, 2, 5)
OPENCV_GPU_IMPLEMENT_RGB5x52RGB_TRAITS(bgr565_to_rgb, 3, 2, 6)
OPENCV_GPU_IMPLEMENT_RGB5x52RGB_TRAITS(bgr555_to_bgr, 3, 0, 5)
OPENCV_GPU_IMPLEMENT_RGB5x52RGB_TRAITS(bgr565_to_bgr, 3, 0, 6)
OPENCV_GPU_IMPLEMENT_RGB5x52RGB_TRAITS(bgr555_to_rgba, 4, 2, 5)
OPENCV_GPU_IMPLEMENT_RGB5x52RGB_TRAITS(bgr565_to_rgba, 4, 2, 6)
OPENCV_GPU_IMPLEMENT_RGB5x52RGB_TRAITS(bgr555_to_bgra, 4, 0, 5)
OPENCV_GPU_IMPLEMENT_RGB5x52RGB_TRAITS(bgr565_to_bgra, 4, 0, 6)
OPENCV_CUDA_IMPLEMENT_RGB5x52RGB_TRAITS(bgr555_to_rgb, 3, 2, 5)
OPENCV_CUDA_IMPLEMENT_RGB5x52RGB_TRAITS(bgr565_to_rgb, 3, 2, 6)
OPENCV_CUDA_IMPLEMENT_RGB5x52RGB_TRAITS(bgr555_to_bgr, 3, 0, 5)
OPENCV_CUDA_IMPLEMENT_RGB5x52RGB_TRAITS(bgr565_to_bgr, 3, 0, 6)
OPENCV_CUDA_IMPLEMENT_RGB5x52RGB_TRAITS(bgr555_to_rgba, 4, 2, 5)
OPENCV_CUDA_IMPLEMENT_RGB5x52RGB_TRAITS(bgr565_to_rgba, 4, 2, 6)
OPENCV_CUDA_IMPLEMENT_RGB5x52RGB_TRAITS(bgr555_to_bgra, 4, 0, 5)
OPENCV_CUDA_IMPLEMENT_RGB5x52RGB_TRAITS(bgr565_to_bgra, 4, 0, 6)
#undef OPENCV_GPU_IMPLEMENT_RGB5x52RGB_TRAITS
#undef OPENCV_CUDA_IMPLEMENT_RGB5x52RGB_TRAITS
OPENCV_GPU_IMPLEMENT_GRAY2RGB_TRAITS(gray_to_bgr, 3)
OPENCV_GPU_IMPLEMENT_GRAY2RGB_TRAITS(gray_to_bgra, 4)
OPENCV_CUDA_IMPLEMENT_GRAY2RGB_TRAITS(gray_to_bgr, 3)
OPENCV_CUDA_IMPLEMENT_GRAY2RGB_TRAITS(gray_to_bgra, 4)
#undef OPENCV_GPU_IMPLEMENT_GRAY2RGB_TRAITS
#undef OPENCV_CUDA_IMPLEMENT_GRAY2RGB_TRAITS
OPENCV_GPU_IMPLEMENT_GRAY2RGB5x5_TRAITS(gray_to_bgr555, 5)
OPENCV_GPU_IMPLEMENT_GRAY2RGB5x5_TRAITS(gray_to_bgr565, 6)
OPENCV_CUDA_IMPLEMENT_GRAY2RGB5x5_TRAITS(gray_to_bgr555, 5)
OPENCV_CUDA_IMPLEMENT_GRAY2RGB5x5_TRAITS(gray_to_bgr565, 6)
#undef OPENCV_GPU_IMPLEMENT_GRAY2RGB5x5_TRAITS
#undef OPENCV_CUDA_IMPLEMENT_GRAY2RGB5x5_TRAITS
OPENCV_GPU_IMPLEMENT_RGB5x52GRAY_TRAITS(bgr555_to_gray, 5)
OPENCV_GPU_IMPLEMENT_RGB5x52GRAY_TRAITS(bgr565_to_gray, 6)
OPENCV_CUDA_IMPLEMENT_RGB5x52GRAY_TRAITS(bgr555_to_gray, 5)
OPENCV_CUDA_IMPLEMENT_RGB5x52GRAY_TRAITS(bgr565_to_gray, 6)
#undef OPENCV_GPU_IMPLEMENT_RGB5x52GRAY_TRAITS
#undef OPENCV_CUDA_IMPLEMENT_RGB5x52GRAY_TRAITS
OPENCV_GPU_IMPLEMENT_RGB2GRAY_TRAITS(rgb_to_gray, 3, 2)
OPENCV_GPU_IMPLEMENT_RGB2GRAY_TRAITS(bgr_to_gray, 3, 0)
OPENCV_GPU_IMPLEMENT_RGB2GRAY_TRAITS(rgba_to_gray, 4, 2)
OPENCV_GPU_IMPLEMENT_RGB2GRAY_TRAITS(bgra_to_gray, 4, 0)
OPENCV_CUDA_IMPLEMENT_RGB2GRAY_TRAITS(rgb_to_gray, 3, 2)
OPENCV_CUDA_IMPLEMENT_RGB2GRAY_TRAITS(bgr_to_gray, 3, 0)
OPENCV_CUDA_IMPLEMENT_RGB2GRAY_TRAITS(rgba_to_gray, 4, 2)
OPENCV_CUDA_IMPLEMENT_RGB2GRAY_TRAITS(bgra_to_gray, 4, 0)
#undef OPENCV_GPU_IMPLEMENT_RGB2GRAY_TRAITS
#undef OPENCV_CUDA_IMPLEMENT_RGB2GRAY_TRAITS
OPENCV_GPU_IMPLEMENT_RGB2YUV_TRAITS(rgb_to_yuv, 3, 3, 2)
OPENCV_GPU_IMPLEMENT_RGB2YUV_TRAITS(rgba_to_yuv, 4, 3, 2)
OPENCV_GPU_IMPLEMENT_RGB2YUV_TRAITS(rgb_to_yuv4, 3, 4, 2)
OPENCV_GPU_IMPLEMENT_RGB2YUV_TRAITS(rgba_to_yuv4, 4, 4, 2)
OPENCV_GPU_IMPLEMENT_RGB2YUV_TRAITS(bgr_to_yuv, 3, 3, 0)
OPENCV_GPU_IMPLEMENT_RGB2YUV_TRAITS(bgra_to_yuv, 4, 3, 0)
OPENCV_GPU_IMPLEMENT_RGB2YUV_TRAITS(bgr_to_yuv4, 3, 4, 0)
OPENCV_GPU_IMPLEMENT_RGB2YUV_TRAITS(bgra_to_yuv4, 4, 4, 0)
OPENCV_CUDA_IMPLEMENT_RGB2YUV_TRAITS(rgb_to_yuv, 3, 3, 2)
OPENCV_CUDA_IMPLEMENT_RGB2YUV_TRAITS(rgba_to_yuv, 4, 3, 2)
OPENCV_CUDA_IMPLEMENT_RGB2YUV_TRAITS(rgb_to_yuv4, 3, 4, 2)
OPENCV_CUDA_IMPLEMENT_RGB2YUV_TRAITS(rgba_to_yuv4, 4, 4, 2)
OPENCV_CUDA_IMPLEMENT_RGB2YUV_TRAITS(bgr_to_yuv, 3, 3, 0)
OPENCV_CUDA_IMPLEMENT_RGB2YUV_TRAITS(bgra_to_yuv, 4, 3, 0)
OPENCV_CUDA_IMPLEMENT_RGB2YUV_TRAITS(bgr_to_yuv4, 3, 4, 0)
OPENCV_CUDA_IMPLEMENT_RGB2YUV_TRAITS(bgra_to_yuv4, 4, 4, 0)
#undef OPENCV_GPU_IMPLEMENT_RGB2YUV_TRAITS
#undef OPENCV_CUDA_IMPLEMENT_RGB2YUV_TRAITS
OPENCV_GPU_IMPLEMENT_YUV2RGB_TRAITS(yuv_to_rgb, 3, 3, 2)
OPENCV_GPU_IMPLEMENT_YUV2RGB_TRAITS(yuv_to_rgba, 3, 4, 2)
OPENCV_GPU_IMPLEMENT_YUV2RGB_TRAITS(yuv4_to_rgb, 4, 3, 2)
OPENCV_GPU_IMPLEMENT_YUV2RGB_TRAITS(yuv4_to_rgba, 4, 4, 2)
OPENCV_GPU_IMPLEMENT_YUV2RGB_TRAITS(yuv_to_bgr, 3, 3, 0)
OPENCV_GPU_IMPLEMENT_YUV2RGB_TRAITS(yuv_to_bgra, 3, 4, 0)
OPENCV_GPU_IMPLEMENT_YUV2RGB_TRAITS(yuv4_to_bgr, 4, 3, 0)
OPENCV_GPU_IMPLEMENT_YUV2RGB_TRAITS(yuv4_to_bgra, 4, 4, 0)
OPENCV_CUDA_IMPLEMENT_YUV2RGB_TRAITS(yuv_to_rgb, 3, 3, 2)
OPENCV_CUDA_IMPLEMENT_YUV2RGB_TRAITS(yuv_to_rgba, 3, 4, 2)
OPENCV_CUDA_IMPLEMENT_YUV2RGB_TRAITS(yuv4_to_rgb, 4, 3, 2)
OPENCV_CUDA_IMPLEMENT_YUV2RGB_TRAITS(yuv4_to_rgba, 4, 4, 2)
OPENCV_CUDA_IMPLEMENT_YUV2RGB_TRAITS(yuv_to_bgr, 3, 3, 0)
OPENCV_CUDA_IMPLEMENT_YUV2RGB_TRAITS(yuv_to_bgra, 3, 4, 0)
OPENCV_CUDA_IMPLEMENT_YUV2RGB_TRAITS(yuv4_to_bgr, 4, 3, 0)
OPENCV_CUDA_IMPLEMENT_YUV2RGB_TRAITS(yuv4_to_bgra, 4, 4, 0)
#undef OPENCV_GPU_IMPLEMENT_YUV2RGB_TRAITS
#undef OPENCV_CUDA_IMPLEMENT_YUV2RGB_TRAITS
OPENCV_GPU_IMPLEMENT_RGB2YCrCb_TRAITS(rgb_to_YCrCb, 3, 3, 2)
OPENCV_GPU_IMPLEMENT_RGB2YCrCb_TRAITS(rgba_to_YCrCb, 4, 3, 2)
OPENCV_GPU_IMPLEMENT_RGB2YCrCb_TRAITS(rgb_to_YCrCb4, 3, 4, 2)
OPENCV_GPU_IMPLEMENT_RGB2YCrCb_TRAITS(rgba_to_YCrCb4, 4, 4, 2)
OPENCV_GPU_IMPLEMENT_RGB2YCrCb_TRAITS(bgr_to_YCrCb, 3, 3, 0)
OPENCV_GPU_IMPLEMENT_RGB2YCrCb_TRAITS(bgra_to_YCrCb, 4, 3, 0)
OPENCV_GPU_IMPLEMENT_RGB2YCrCb_TRAITS(bgr_to_YCrCb4, 3, 4, 0)
OPENCV_GPU_IMPLEMENT_RGB2YCrCb_TRAITS(bgra_to_YCrCb4, 4, 4, 0)
OPENCV_CUDA_IMPLEMENT_RGB2YCrCb_TRAITS(rgb_to_YCrCb, 3, 3, 2)
OPENCV_CUDA_IMPLEMENT_RGB2YCrCb_TRAITS(rgba_to_YCrCb, 4, 3, 2)
OPENCV_CUDA_IMPLEMENT_RGB2YCrCb_TRAITS(rgb_to_YCrCb4, 3, 4, 2)
OPENCV_CUDA_IMPLEMENT_RGB2YCrCb_TRAITS(rgba_to_YCrCb4, 4, 4, 2)
OPENCV_CUDA_IMPLEMENT_RGB2YCrCb_TRAITS(bgr_to_YCrCb, 3, 3, 0)
OPENCV_CUDA_IMPLEMENT_RGB2YCrCb_TRAITS(bgra_to_YCrCb, 4, 3, 0)
OPENCV_CUDA_IMPLEMENT_RGB2YCrCb_TRAITS(bgr_to_YCrCb4, 3, 4, 0)
OPENCV_CUDA_IMPLEMENT_RGB2YCrCb_TRAITS(bgra_to_YCrCb4, 4, 4, 0)
#undef OPENCV_GPU_IMPLEMENT_RGB2YCrCb_TRAITS
#undef OPENCV_CUDA_IMPLEMENT_RGB2YCrCb_TRAITS
OPENCV_GPU_IMPLEMENT_YCrCb2RGB_TRAITS(YCrCb_to_rgb, 3, 3, 2)
OPENCV_GPU_IMPLEMENT_YCrCb2RGB_TRAITS(YCrCb_to_rgba, 3, 4, 2)
OPENCV_GPU_IMPLEMENT_YCrCb2RGB_TRAITS(YCrCb4_to_rgb, 4, 3, 2)
OPENCV_GPU_IMPLEMENT_YCrCb2RGB_TRAITS(YCrCb4_to_rgba, 4, 4, 2)
OPENCV_GPU_IMPLEMENT_YCrCb2RGB_TRAITS(YCrCb_to_bgr, 3, 3, 0)
OPENCV_GPU_IMPLEMENT_YCrCb2RGB_TRAITS(YCrCb_to_bgra, 3, 4, 0)
OPENCV_GPU_IMPLEMENT_YCrCb2RGB_TRAITS(YCrCb4_to_bgr, 4, 3, 0)
OPENCV_GPU_IMPLEMENT_YCrCb2RGB_TRAITS(YCrCb4_to_bgra, 4, 4, 0)
OPENCV_CUDA_IMPLEMENT_YCrCb2RGB_TRAITS(YCrCb_to_rgb, 3, 3, 2)
OPENCV_CUDA_IMPLEMENT_YCrCb2RGB_TRAITS(YCrCb_to_rgba, 3, 4, 2)
OPENCV_CUDA_IMPLEMENT_YCrCb2RGB_TRAITS(YCrCb4_to_rgb, 4, 3, 2)
OPENCV_CUDA_IMPLEMENT_YCrCb2RGB_TRAITS(YCrCb4_to_rgba, 4, 4, 2)
OPENCV_CUDA_IMPLEMENT_YCrCb2RGB_TRAITS(YCrCb_to_bgr, 3, 3, 0)
OPENCV_CUDA_IMPLEMENT_YCrCb2RGB_TRAITS(YCrCb_to_bgra, 3, 4, 0)
OPENCV_CUDA_IMPLEMENT_YCrCb2RGB_TRAITS(YCrCb4_to_bgr, 4, 3, 0)
OPENCV_CUDA_IMPLEMENT_YCrCb2RGB_TRAITS(YCrCb4_to_bgra, 4, 4, 0)
#undef OPENCV_GPU_IMPLEMENT_YCrCb2RGB_TRAITS
#undef OPENCV_CUDA_IMPLEMENT_YCrCb2RGB_TRAITS
OPENCV_GPU_IMPLEMENT_RGB2XYZ_TRAITS(rgb_to_xyz, 3, 3, 2)
OPENCV_GPU_IMPLEMENT_RGB2XYZ_TRAITS(rgba_to_xyz, 4, 3, 2)
OPENCV_GPU_IMPLEMENT_RGB2XYZ_TRAITS(rgb_to_xyz4, 3, 4, 2)
OPENCV_GPU_IMPLEMENT_RGB2XYZ_TRAITS(rgba_to_xyz4, 4, 4, 2)
OPENCV_GPU_IMPLEMENT_RGB2XYZ_TRAITS(bgr_to_xyz, 3, 3, 0)
OPENCV_GPU_IMPLEMENT_RGB2XYZ_TRAITS(bgra_to_xyz, 4, 3, 0)
OPENCV_GPU_IMPLEMENT_RGB2XYZ_TRAITS(bgr_to_xyz4, 3, 4, 0)
OPENCV_GPU_IMPLEMENT_RGB2XYZ_TRAITS(bgra_to_xyz4, 4, 4, 0)
OPENCV_CUDA_IMPLEMENT_RGB2XYZ_TRAITS(rgb_to_xyz, 3, 3, 2)
OPENCV_CUDA_IMPLEMENT_RGB2XYZ_TRAITS(rgba_to_xyz, 4, 3, 2)
OPENCV_CUDA_IMPLEMENT_RGB2XYZ_TRAITS(rgb_to_xyz4, 3, 4, 2)
OPENCV_CUDA_IMPLEMENT_RGB2XYZ_TRAITS(rgba_to_xyz4, 4, 4, 2)
OPENCV_CUDA_IMPLEMENT_RGB2XYZ_TRAITS(bgr_to_xyz, 3, 3, 0)
OPENCV_CUDA_IMPLEMENT_RGB2XYZ_TRAITS(bgra_to_xyz, 4, 3, 0)
OPENCV_CUDA_IMPLEMENT_RGB2XYZ_TRAITS(bgr_to_xyz4, 3, 4, 0)
OPENCV_CUDA_IMPLEMENT_RGB2XYZ_TRAITS(bgra_to_xyz4, 4, 4, 0)
#undef OPENCV_GPU_IMPLEMENT_RGB2XYZ_TRAITS
#undef OPENCV_CUDA_IMPLEMENT_RGB2XYZ_TRAITS
OPENCV_GPU_IMPLEMENT_XYZ2RGB_TRAITS(xyz_to_rgb, 3, 3, 2)
OPENCV_GPU_IMPLEMENT_XYZ2RGB_TRAITS(xyz4_to_rgb, 4, 3, 2)
OPENCV_GPU_IMPLEMENT_XYZ2RGB_TRAITS(xyz_to_rgba, 3, 4, 2)
OPENCV_GPU_IMPLEMENT_XYZ2RGB_TRAITS(xyz4_to_rgba, 4, 4, 2)
OPENCV_GPU_IMPLEMENT_XYZ2RGB_TRAITS(xyz_to_bgr, 3, 3, 0)
OPENCV_GPU_IMPLEMENT_XYZ2RGB_TRAITS(xyz4_to_bgr, 4, 3, 0)
OPENCV_GPU_IMPLEMENT_XYZ2RGB_TRAITS(xyz_to_bgra, 3, 4, 0)
OPENCV_GPU_IMPLEMENT_XYZ2RGB_TRAITS(xyz4_to_bgra, 4, 4, 0)
OPENCV_CUDA_IMPLEMENT_XYZ2RGB_TRAITS(xyz_to_rgb, 3, 3, 2)
OPENCV_CUDA_IMPLEMENT_XYZ2RGB_TRAITS(xyz4_to_rgb, 4, 3, 2)
OPENCV_CUDA_IMPLEMENT_XYZ2RGB_TRAITS(xyz_to_rgba, 3, 4, 2)
OPENCV_CUDA_IMPLEMENT_XYZ2RGB_TRAITS(xyz4_to_rgba, 4, 4, 2)
OPENCV_CUDA_IMPLEMENT_XYZ2RGB_TRAITS(xyz_to_bgr, 3, 3, 0)
OPENCV_CUDA_IMPLEMENT_XYZ2RGB_TRAITS(xyz4_to_bgr, 4, 3, 0)
OPENCV_CUDA_IMPLEMENT_XYZ2RGB_TRAITS(xyz_to_bgra, 3, 4, 0)
OPENCV_CUDA_IMPLEMENT_XYZ2RGB_TRAITS(xyz4_to_bgra, 4, 4, 0)
#undef OPENCV_GPU_IMPLEMENT_XYZ2RGB_TRAITS
#undef OPENCV_CUDA_IMPLEMENT_XYZ2RGB_TRAITS
OPENCV_GPU_IMPLEMENT_RGB2HSV_TRAITS(rgb_to_hsv, 3, 3, 2)
OPENCV_GPU_IMPLEMENT_RGB2HSV_TRAITS(rgba_to_hsv, 4, 3, 2)
OPENCV_GPU_IMPLEMENT_RGB2HSV_TRAITS(rgb_to_hsv4, 3, 4, 2)
OPENCV_GPU_IMPLEMENT_RGB2HSV_TRAITS(rgba_to_hsv4, 4, 4, 2)
OPENCV_GPU_IMPLEMENT_RGB2HSV_TRAITS(bgr_to_hsv, 3, 3, 0)
OPENCV_GPU_IMPLEMENT_RGB2HSV_TRAITS(bgra_to_hsv, 4, 3, 0)
OPENCV_GPU_IMPLEMENT_RGB2HSV_TRAITS(bgr_to_hsv4, 3, 4, 0)
OPENCV_GPU_IMPLEMENT_RGB2HSV_TRAITS(bgra_to_hsv4, 4, 4, 0)
OPENCV_CUDA_IMPLEMENT_RGB2HSV_TRAITS(rgb_to_hsv, 3, 3, 2)
OPENCV_CUDA_IMPLEMENT_RGB2HSV_TRAITS(rgba_to_hsv, 4, 3, 2)
OPENCV_CUDA_IMPLEMENT_RGB2HSV_TRAITS(rgb_to_hsv4, 3, 4, 2)
OPENCV_CUDA_IMPLEMENT_RGB2HSV_TRAITS(rgba_to_hsv4, 4, 4, 2)
OPENCV_CUDA_IMPLEMENT_RGB2HSV_TRAITS(bgr_to_hsv, 3, 3, 0)
OPENCV_CUDA_IMPLEMENT_RGB2HSV_TRAITS(bgra_to_hsv, 4, 3, 0)
OPENCV_CUDA_IMPLEMENT_RGB2HSV_TRAITS(bgr_to_hsv4, 3, 4, 0)
OPENCV_CUDA_IMPLEMENT_RGB2HSV_TRAITS(bgra_to_hsv4, 4, 4, 0)
#undef OPENCV_GPU_IMPLEMENT_RGB2HSV_TRAITS
#undef OPENCV_CUDA_IMPLEMENT_RGB2HSV_TRAITS
OPENCV_GPU_IMPLEMENT_HSV2RGB_TRAITS(hsv_to_rgb, 3, 3, 2)
OPENCV_GPU_IMPLEMENT_HSV2RGB_TRAITS(hsv_to_rgba, 3, 4, 2)
OPENCV_GPU_IMPLEMENT_HSV2RGB_TRAITS(hsv4_to_rgb, 4, 3, 2)
OPENCV_GPU_IMPLEMENT_HSV2RGB_TRAITS(hsv4_to_rgba, 4, 4, 2)
OPENCV_GPU_IMPLEMENT_HSV2RGB_TRAITS(hsv_to_bgr, 3, 3, 0)
OPENCV_GPU_IMPLEMENT_HSV2RGB_TRAITS(hsv_to_bgra, 3, 4, 0)
OPENCV_GPU_IMPLEMENT_HSV2RGB_TRAITS(hsv4_to_bgr, 4, 3, 0)
OPENCV_GPU_IMPLEMENT_HSV2RGB_TRAITS(hsv4_to_bgra, 4, 4, 0)
OPENCV_CUDA_IMPLEMENT_HSV2RGB_TRAITS(hsv_to_rgb, 3, 3, 2)
OPENCV_CUDA_IMPLEMENT_HSV2RGB_TRAITS(hsv_to_rgba, 3, 4, 2)
OPENCV_CUDA_IMPLEMENT_HSV2RGB_TRAITS(hsv4_to_rgb, 4, 3, 2)
OPENCV_CUDA_IMPLEMENT_HSV2RGB_TRAITS(hsv4_to_rgba, 4, 4, 2)
OPENCV_CUDA_IMPLEMENT_HSV2RGB_TRAITS(hsv_to_bgr, 3, 3, 0)
OPENCV_CUDA_IMPLEMENT_HSV2RGB_TRAITS(hsv_to_bgra, 3, 4, 0)
OPENCV_CUDA_IMPLEMENT_HSV2RGB_TRAITS(hsv4_to_bgr, 4, 3, 0)
OPENCV_CUDA_IMPLEMENT_HSV2RGB_TRAITS(hsv4_to_bgra, 4, 4, 0)
#undef OPENCV_GPU_IMPLEMENT_HSV2RGB_TRAITS
#undef OPENCV_CUDA_IMPLEMENT_HSV2RGB_TRAITS
OPENCV_GPU_IMPLEMENT_RGB2HLS_TRAITS(rgb_to_hls, 3, 3, 2)
OPENCV_GPU_IMPLEMENT_RGB2HLS_TRAITS(rgba_to_hls, 4, 3, 2)
OPENCV_GPU_IMPLEMENT_RGB2HLS_TRAITS(rgb_to_hls4, 3, 4, 2)
OPENCV_GPU_IMPLEMENT_RGB2HLS_TRAITS(rgba_to_hls4, 4, 4, 2)
OPENCV_GPU_IMPLEMENT_RGB2HLS_TRAITS(bgr_to_hls, 3, 3, 0)
OPENCV_GPU_IMPLEMENT_RGB2HLS_TRAITS(bgra_to_hls, 4, 3, 0)
OPENCV_GPU_IMPLEMENT_RGB2HLS_TRAITS(bgr_to_hls4, 3, 4, 0)
OPENCV_GPU_IMPLEMENT_RGB2HLS_TRAITS(bgra_to_hls4, 4, 4, 0)
OPENCV_CUDA_IMPLEMENT_RGB2HLS_TRAITS(rgb_to_hls, 3, 3, 2)
OPENCV_CUDA_IMPLEMENT_RGB2HLS_TRAITS(rgba_to_hls, 4, 3, 2)
OPENCV_CUDA_IMPLEMENT_RGB2HLS_TRAITS(rgb_to_hls4, 3, 4, 2)
OPENCV_CUDA_IMPLEMENT_RGB2HLS_TRAITS(rgba_to_hls4, 4, 4, 2)
OPENCV_CUDA_IMPLEMENT_RGB2HLS_TRAITS(bgr_to_hls, 3, 3, 0)
OPENCV_CUDA_IMPLEMENT_RGB2HLS_TRAITS(bgra_to_hls, 4, 3, 0)
OPENCV_CUDA_IMPLEMENT_RGB2HLS_TRAITS(bgr_to_hls4, 3, 4, 0)
OPENCV_CUDA_IMPLEMENT_RGB2HLS_TRAITS(bgra_to_hls4, 4, 4, 0)
#undef OPENCV_GPU_IMPLEMENT_RGB2HLS_TRAITS
#undef OPENCV_CUDA_IMPLEMENT_RGB2HLS_TRAITS
OPENCV_GPU_IMPLEMENT_HLS2RGB_TRAITS(hls_to_rgb, 3, 3, 2)
OPENCV_GPU_IMPLEMENT_HLS2RGB_TRAITS(hls_to_rgba, 3, 4, 2)
OPENCV_GPU_IMPLEMENT_HLS2RGB_TRAITS(hls4_to_rgb, 4, 3, 2)
OPENCV_GPU_IMPLEMENT_HLS2RGB_TRAITS(hls4_to_rgba, 4, 4, 2)
OPENCV_GPU_IMPLEMENT_HLS2RGB_TRAITS(hls_to_bgr, 3, 3, 0)
OPENCV_GPU_IMPLEMENT_HLS2RGB_TRAITS(hls_to_bgra, 3, 4, 0)
OPENCV_GPU_IMPLEMENT_HLS2RGB_TRAITS(hls4_to_bgr, 4, 3, 0)
OPENCV_GPU_IMPLEMENT_HLS2RGB_TRAITS(hls4_to_bgra, 4, 4, 0)
OPENCV_CUDA_IMPLEMENT_HLS2RGB_TRAITS(hls_to_rgb, 3, 3, 2)
OPENCV_CUDA_IMPLEMENT_HLS2RGB_TRAITS(hls_to_rgba, 3, 4, 2)
OPENCV_CUDA_IMPLEMENT_HLS2RGB_TRAITS(hls4_to_rgb, 4, 3, 2)
OPENCV_CUDA_IMPLEMENT_HLS2RGB_TRAITS(hls4_to_rgba, 4, 4, 2)
OPENCV_CUDA_IMPLEMENT_HLS2RGB_TRAITS(hls_to_bgr, 3, 3, 0)
OPENCV_CUDA_IMPLEMENT_HLS2RGB_TRAITS(hls_to_bgra, 3, 4, 0)
OPENCV_CUDA_IMPLEMENT_HLS2RGB_TRAITS(hls4_to_bgr, 4, 3, 0)
OPENCV_CUDA_IMPLEMENT_HLS2RGB_TRAITS(hls4_to_bgra, 4, 4, 0)
#undef OPENCV_GPU_IMPLEMENT_HLS2RGB_TRAITS
#undef OPENCV_CUDA_IMPLEMENT_HLS2RGB_TRAITS
OPENCV_GPU_IMPLEMENT_RGB2Lab_TRAITS(rgb_to_lab, 3, 3, true, 2)
OPENCV_GPU_IMPLEMENT_RGB2Lab_TRAITS(rgba_to_lab, 4, 3, true, 2)
OPENCV_GPU_IMPLEMENT_RGB2Lab_TRAITS(rgb_to_lab4, 3, 4, true, 2)
OPENCV_GPU_IMPLEMENT_RGB2Lab_TRAITS(rgba_to_lab4, 4, 4, true, 2)
OPENCV_GPU_IMPLEMENT_RGB2Lab_TRAITS(bgr_to_lab, 3, 3, true, 0)
OPENCV_GPU_IMPLEMENT_RGB2Lab_TRAITS(bgra_to_lab, 4, 3, true, 0)
OPENCV_GPU_IMPLEMENT_RGB2Lab_TRAITS(bgr_to_lab4, 3, 4, true, 0)
OPENCV_GPU_IMPLEMENT_RGB2Lab_TRAITS(bgra_to_lab4, 4, 4, true, 0)
OPENCV_CUDA_IMPLEMENT_RGB2Lab_TRAITS(rgb_to_lab, 3, 3, true, 2)
OPENCV_CUDA_IMPLEMENT_RGB2Lab_TRAITS(rgba_to_lab, 4, 3, true, 2)
OPENCV_CUDA_IMPLEMENT_RGB2Lab_TRAITS(rgb_to_lab4, 3, 4, true, 2)
OPENCV_CUDA_IMPLEMENT_RGB2Lab_TRAITS(rgba_to_lab4, 4, 4, true, 2)
OPENCV_CUDA_IMPLEMENT_RGB2Lab_TRAITS(bgr_to_lab, 3, 3, true, 0)
OPENCV_CUDA_IMPLEMENT_RGB2Lab_TRAITS(bgra_to_lab, 4, 3, true, 0)
OPENCV_CUDA_IMPLEMENT_RGB2Lab_TRAITS(bgr_to_lab4, 3, 4, true, 0)
OPENCV_CUDA_IMPLEMENT_RGB2Lab_TRAITS(bgra_to_lab4, 4, 4, true, 0)
OPENCV_GPU_IMPLEMENT_RGB2Lab_TRAITS(lrgb_to_lab, 3, 3, false, 2)
OPENCV_GPU_IMPLEMENT_RGB2Lab_TRAITS(lrgba_to_lab, 4, 3, false, 2)
OPENCV_GPU_IMPLEMENT_RGB2Lab_TRAITS(lrgb_to_lab4, 3, 4, false, 2)
OPENCV_GPU_IMPLEMENT_RGB2Lab_TRAITS(lrgba_to_lab4, 4, 4, false, 2)
OPENCV_GPU_IMPLEMENT_RGB2Lab_TRAITS(lbgr_to_lab, 3, 3, false, 0)
OPENCV_GPU_IMPLEMENT_RGB2Lab_TRAITS(lbgra_to_lab, 4, 3, false, 0)
OPENCV_GPU_IMPLEMENT_RGB2Lab_TRAITS(lbgr_to_lab4, 3, 4, false, 0)
OPENCV_GPU_IMPLEMENT_RGB2Lab_TRAITS(lbgra_to_lab4, 4, 4, false, 0)
OPENCV_CUDA_IMPLEMENT_RGB2Lab_TRAITS(lrgb_to_lab, 3, 3, false, 2)
OPENCV_CUDA_IMPLEMENT_RGB2Lab_TRAITS(lrgba_to_lab, 4, 3, false, 2)
OPENCV_CUDA_IMPLEMENT_RGB2Lab_TRAITS(lrgb_to_lab4, 3, 4, false, 2)
OPENCV_CUDA_IMPLEMENT_RGB2Lab_TRAITS(lrgba_to_lab4, 4, 4, false, 2)
OPENCV_CUDA_IMPLEMENT_RGB2Lab_TRAITS(lbgr_to_lab, 3, 3, false, 0)
OPENCV_CUDA_IMPLEMENT_RGB2Lab_TRAITS(lbgra_to_lab, 4, 3, false, 0)
OPENCV_CUDA_IMPLEMENT_RGB2Lab_TRAITS(lbgr_to_lab4, 3, 4, false, 0)
OPENCV_CUDA_IMPLEMENT_RGB2Lab_TRAITS(lbgra_to_lab4, 4, 4, false, 0)
#undef OPENCV_GPU_IMPLEMENT_RGB2Lab_TRAITS
#undef OPENCV_CUDA_IMPLEMENT_RGB2Lab_TRAITS
OPENCV_GPU_IMPLEMENT_Lab2RGB_TRAITS(lab_to_rgb, 3, 3, true, 2)
OPENCV_GPU_IMPLEMENT_Lab2RGB_TRAITS(lab4_to_rgb, 4, 3, true, 2)
OPENCV_GPU_IMPLEMENT_Lab2RGB_TRAITS(lab_to_rgba, 3, 4, true, 2)
OPENCV_GPU_IMPLEMENT_Lab2RGB_TRAITS(lab4_to_rgba, 4, 4, true, 2)
OPENCV_GPU_IMPLEMENT_Lab2RGB_TRAITS(lab_to_bgr, 3, 3, true, 0)
OPENCV_GPU_IMPLEMENT_Lab2RGB_TRAITS(lab4_to_bgr, 4, 3, true, 0)
OPENCV_GPU_IMPLEMENT_Lab2RGB_TRAITS(lab_to_bgra, 3, 4, true, 0)
OPENCV_GPU_IMPLEMENT_Lab2RGB_TRAITS(lab4_to_bgra, 4, 4, true, 0)
OPENCV_CUDA_IMPLEMENT_Lab2RGB_TRAITS(lab_to_rgb, 3, 3, true, 2)
OPENCV_CUDA_IMPLEMENT_Lab2RGB_TRAITS(lab4_to_rgb, 4, 3, true, 2)
OPENCV_CUDA_IMPLEMENT_Lab2RGB_TRAITS(lab_to_rgba, 3, 4, true, 2)
OPENCV_CUDA_IMPLEMENT_Lab2RGB_TRAITS(lab4_to_rgba, 4, 4, true, 2)
OPENCV_CUDA_IMPLEMENT_Lab2RGB_TRAITS(lab_to_bgr, 3, 3, true, 0)
OPENCV_CUDA_IMPLEMENT_Lab2RGB_TRAITS(lab4_to_bgr, 4, 3, true, 0)
OPENCV_CUDA_IMPLEMENT_Lab2RGB_TRAITS(lab_to_bgra, 3, 4, true, 0)
OPENCV_CUDA_IMPLEMENT_Lab2RGB_TRAITS(lab4_to_bgra, 4, 4, true, 0)
OPENCV_GPU_IMPLEMENT_Lab2RGB_TRAITS(lab_to_lrgb, 3, 3, false, 2)
OPENCV_GPU_IMPLEMENT_Lab2RGB_TRAITS(lab4_to_lrgb, 4, 3, false, 2)
OPENCV_GPU_IMPLEMENT_Lab2RGB_TRAITS(lab_to_lrgba, 3, 4, false, 2)
OPENCV_GPU_IMPLEMENT_Lab2RGB_TRAITS(lab4_to_lrgba, 4, 4, false, 2)
OPENCV_GPU_IMPLEMENT_Lab2RGB_TRAITS(lab_to_lbgr, 3, 3, false, 0)
OPENCV_GPU_IMPLEMENT_Lab2RGB_TRAITS(lab4_to_lbgr, 4, 3, false, 0)
OPENCV_GPU_IMPLEMENT_Lab2RGB_TRAITS(lab_to_lbgra, 3, 4, false, 0)
OPENCV_GPU_IMPLEMENT_Lab2RGB_TRAITS(lab4_to_lbgra, 4, 4, false, 0)
OPENCV_CUDA_IMPLEMENT_Lab2RGB_TRAITS(lab_to_lrgb, 3, 3, false, 2)
OPENCV_CUDA_IMPLEMENT_Lab2RGB_TRAITS(lab4_to_lrgb, 4, 3, false, 2)
OPENCV_CUDA_IMPLEMENT_Lab2RGB_TRAITS(lab_to_lrgba, 3, 4, false, 2)
OPENCV_CUDA_IMPLEMENT_Lab2RGB_TRAITS(lab4_to_lrgba, 4, 4, false, 2)
OPENCV_CUDA_IMPLEMENT_Lab2RGB_TRAITS(lab_to_lbgr, 3, 3, false, 0)
OPENCV_CUDA_IMPLEMENT_Lab2RGB_TRAITS(lab4_to_lbgr, 4, 3, false, 0)
OPENCV_CUDA_IMPLEMENT_Lab2RGB_TRAITS(lab_to_lbgra, 3, 4, false, 0)
OPENCV_CUDA_IMPLEMENT_Lab2RGB_TRAITS(lab4_to_lbgra, 4, 4, false, 0)
#undef OPENCV_GPU_IMPLEMENT_Lab2RGB_TRAITS
#undef OPENCV_CUDA_IMPLEMENT_Lab2RGB_TRAITS
OPENCV_GPU_IMPLEMENT_RGB2Luv_TRAITS(rgb_to_luv, 3, 3, true, 2)
OPENCV_GPU_IMPLEMENT_RGB2Luv_TRAITS(rgba_to_luv, 4, 3, true, 2)
OPENCV_GPU_IMPLEMENT_RGB2Luv_TRAITS(rgb_to_luv4, 3, 4, true, 2)
OPENCV_GPU_IMPLEMENT_RGB2Luv_TRAITS(rgba_to_luv4, 4, 4, true, 2)
OPENCV_GPU_IMPLEMENT_RGB2Luv_TRAITS(bgr_to_luv, 3, 3, true, 0)
OPENCV_GPU_IMPLEMENT_RGB2Luv_TRAITS(bgra_to_luv, 4, 3, true, 0)
OPENCV_GPU_IMPLEMENT_RGB2Luv_TRAITS(bgr_to_luv4, 3, 4, true, 0)
OPENCV_GPU_IMPLEMENT_RGB2Luv_TRAITS(bgra_to_luv4, 4, 4, true, 0)
OPENCV_CUDA_IMPLEMENT_RGB2Luv_TRAITS(rgb_to_luv, 3, 3, true, 2)
OPENCV_CUDA_IMPLEMENT_RGB2Luv_TRAITS(rgba_to_luv, 4, 3, true, 2)
OPENCV_CUDA_IMPLEMENT_RGB2Luv_TRAITS(rgb_to_luv4, 3, 4, true, 2)
OPENCV_CUDA_IMPLEMENT_RGB2Luv_TRAITS(rgba_to_luv4, 4, 4, true, 2)
OPENCV_CUDA_IMPLEMENT_RGB2Luv_TRAITS(bgr_to_luv, 3, 3, true, 0)
OPENCV_CUDA_IMPLEMENT_RGB2Luv_TRAITS(bgra_to_luv, 4, 3, true, 0)
OPENCV_CUDA_IMPLEMENT_RGB2Luv_TRAITS(bgr_to_luv4, 3, 4, true, 0)
OPENCV_CUDA_IMPLEMENT_RGB2Luv_TRAITS(bgra_to_luv4, 4, 4, true, 0)
OPENCV_GPU_IMPLEMENT_RGB2Luv_TRAITS(lrgb_to_luv, 3, 3, false, 2)
OPENCV_GPU_IMPLEMENT_RGB2Luv_TRAITS(lrgba_to_luv, 4, 3, false, 2)
OPENCV_GPU_IMPLEMENT_RGB2Luv_TRAITS(lrgb_to_luv4, 3, 4, false, 2)
OPENCV_GPU_IMPLEMENT_RGB2Luv_TRAITS(lrgba_to_luv4, 4, 4, false, 2)
OPENCV_GPU_IMPLEMENT_RGB2Luv_TRAITS(lbgr_to_luv, 3, 3, false, 0)
OPENCV_GPU_IMPLEMENT_RGB2Luv_TRAITS(lbgra_to_luv, 4, 3, false, 0)
OPENCV_GPU_IMPLEMENT_RGB2Luv_TRAITS(lbgr_to_luv4, 3, 4, false, 0)
OPENCV_GPU_IMPLEMENT_RGB2Luv_TRAITS(lbgra_to_luv4, 4, 4, false, 0)
OPENCV_CUDA_IMPLEMENT_RGB2Luv_TRAITS(lrgb_to_luv, 3, 3, false, 2)
OPENCV_CUDA_IMPLEMENT_RGB2Luv_TRAITS(lrgba_to_luv, 4, 3, false, 2)
OPENCV_CUDA_IMPLEMENT_RGB2Luv_TRAITS(lrgb_to_luv4, 3, 4, false, 2)
OPENCV_CUDA_IMPLEMENT_RGB2Luv_TRAITS(lrgba_to_luv4, 4, 4, false, 2)
OPENCV_CUDA_IMPLEMENT_RGB2Luv_TRAITS(lbgr_to_luv, 3, 3, false, 0)
OPENCV_CUDA_IMPLEMENT_RGB2Luv_TRAITS(lbgra_to_luv, 4, 3, false, 0)
OPENCV_CUDA_IMPLEMENT_RGB2Luv_TRAITS(lbgr_to_luv4, 3, 4, false, 0)
OPENCV_CUDA_IMPLEMENT_RGB2Luv_TRAITS(lbgra_to_luv4, 4, 4, false, 0)
#undef OPENCV_GPU_IMPLEMENT_RGB2Luv_TRAITS
#undef OPENCV_CUDA_IMPLEMENT_RGB2Luv_TRAITS
OPENCV_GPU_IMPLEMENT_Luv2RGB_TRAITS(luv_to_rgb, 3, 3, true, 2)
OPENCV_GPU_IMPLEMENT_Luv2RGB_TRAITS(luv4_to_rgb, 4, 3, true, 2)
OPENCV_GPU_IMPLEMENT_Luv2RGB_TRAITS(luv_to_rgba, 3, 4, true, 2)
OPENCV_GPU_IMPLEMENT_Luv2RGB_TRAITS(luv4_to_rgba, 4, 4, true, 2)
OPENCV_GPU_IMPLEMENT_Luv2RGB_TRAITS(luv_to_bgr, 3, 3, true, 0)
OPENCV_GPU_IMPLEMENT_Luv2RGB_TRAITS(luv4_to_bgr, 4, 3, true, 0)
OPENCV_GPU_IMPLEMENT_Luv2RGB_TRAITS(luv_to_bgra, 3, 4, true, 0)
OPENCV_GPU_IMPLEMENT_Luv2RGB_TRAITS(luv4_to_bgra, 4, 4, true, 0)
OPENCV_CUDA_IMPLEMENT_Luv2RGB_TRAITS(luv_to_rgb, 3, 3, true, 2)
OPENCV_CUDA_IMPLEMENT_Luv2RGB_TRAITS(luv4_to_rgb, 4, 3, true, 2)
OPENCV_CUDA_IMPLEMENT_Luv2RGB_TRAITS(luv_to_rgba, 3, 4, true, 2)
OPENCV_CUDA_IMPLEMENT_Luv2RGB_TRAITS(luv4_to_rgba, 4, 4, true, 2)
OPENCV_CUDA_IMPLEMENT_Luv2RGB_TRAITS(luv_to_bgr, 3, 3, true, 0)
OPENCV_CUDA_IMPLEMENT_Luv2RGB_TRAITS(luv4_to_bgr, 4, 3, true, 0)
OPENCV_CUDA_IMPLEMENT_Luv2RGB_TRAITS(luv_to_bgra, 3, 4, true, 0)
OPENCV_CUDA_IMPLEMENT_Luv2RGB_TRAITS(luv4_to_bgra, 4, 4, true, 0)
OPENCV_GPU_IMPLEMENT_Luv2RGB_TRAITS(luv_to_lrgb, 3, 3, false, 2)
OPENCV_GPU_IMPLEMENT_Luv2RGB_TRAITS(luv4_to_lrgb, 4, 3, false, 2)
OPENCV_GPU_IMPLEMENT_Luv2RGB_TRAITS(luv_to_lrgba, 3, 4, false, 2)
OPENCV_GPU_IMPLEMENT_Luv2RGB_TRAITS(luv4_to_lrgba, 4, 4, false, 2)
OPENCV_GPU_IMPLEMENT_Luv2RGB_TRAITS(luv_to_lbgr, 3, 3, false, 0)
OPENCV_GPU_IMPLEMENT_Luv2RGB_TRAITS(luv4_to_lbgr, 4, 3, false, 0)
OPENCV_GPU_IMPLEMENT_Luv2RGB_TRAITS(luv_to_lbgra, 3, 4, false, 0)
OPENCV_GPU_IMPLEMENT_Luv2RGB_TRAITS(luv4_to_lbgra, 4, 4, false, 0)
OPENCV_CUDA_IMPLEMENT_Luv2RGB_TRAITS(luv_to_lrgb, 3, 3, false, 2)
OPENCV_CUDA_IMPLEMENT_Luv2RGB_TRAITS(luv4_to_lrgb, 4, 3, false, 2)
OPENCV_CUDA_IMPLEMENT_Luv2RGB_TRAITS(luv_to_lrgba, 3, 4, false, 2)
OPENCV_CUDA_IMPLEMENT_Luv2RGB_TRAITS(luv4_to_lrgba, 4, 4, false, 2)
OPENCV_CUDA_IMPLEMENT_Luv2RGB_TRAITS(luv_to_lbgr, 3, 3, false, 0)
OPENCV_CUDA_IMPLEMENT_Luv2RGB_TRAITS(luv4_to_lbgr, 4, 3, false, 0)
OPENCV_CUDA_IMPLEMENT_Luv2RGB_TRAITS(luv_to_lbgra, 3, 4, false, 0)
OPENCV_CUDA_IMPLEMENT_Luv2RGB_TRAITS(luv4_to_lbgra, 4, 4, false, 0)
#undef OPENCV_GPU_IMPLEMENT_Luv2RGB_TRAITS
}}} // namespace cv { namespace gpu { namespace cudev
#undef OPENCV_CUDA_IMPLEMENT_Luv2RGB_TRAITS
}}} // namespace cv { namespace cuda { namespace cudev
#endif // __OPENCV_GPU_BORDER_INTERPOLATE_HPP__
#endif // __OPENCV_CUDA_BORDER_INTERPOLATE_HPP__

20
modules/core/include/opencv2/core/cuda/common.hpp
View File

@@ -40,11 +40,11 @@
//
//M*/
#ifndef __OPENCV_GPU_COMMON_HPP__
#define __OPENCV_GPU_COMMON_HPP__
#ifndef __OPENCV_CUDA_COMMON_HPP__
#define __OPENCV_CUDA_COMMON_HPP__
#include <cuda_runtime.h>
#include "opencv2/core/gpu_types.hpp"
#include "opencv2/core/cuda_types.hpp"
#include "opencv2/core/cvdef.h"
#include "opencv2/core/base.hpp"
@@ -56,7 +56,7 @@
#endif
#endif
namespace cv { namespace gpu {
namespace cv { namespace cuda {
static inline void checkCudaError(cudaError_t err, const char* file, const int line, const char* func)
{
if (cudaSuccess != err)
@@ -66,13 +66,13 @@ namespace cv { namespace gpu {
#ifndef cudaSafeCall
#if defined(__GNUC__)
#define cudaSafeCall(expr) cv::gpu::checkCudaError(expr, __FILE__, __LINE__, __func__)
#define cudaSafeCall(expr) cv::cuda::checkCudaError(expr, __FILE__, __LINE__, __func__)
#else /* defined(__CUDACC__) || defined(__MSVC__) */
#define cudaSafeCall(expr) cv::gpu::checkCudaError(expr, __FILE__, __LINE__, "")
#define cudaSafeCall(expr) cv::cuda::checkCudaError(expr, __FILE__, __LINE__, "")
#endif
#endif
namespace cv { namespace gpu
namespace cv { namespace cuda
{
template <typename T> static inline bool isAligned(const T* ptr, size_t size)
{
@@ -85,9 +85,9 @@ namespace cv { namespace gpu
}
}}
namespace cv { namespace gpu
namespace cv { namespace cuda
{
namespace cudev
namespace device
{
__host__ __device__ __forceinline__ int divUp(int total, int grain)
{
@@ -104,4 +104,4 @@ namespace cv { namespace gpu
#endif // __OPENCV_GPU_COMMON_HPP__
#endif // __OPENCV_CUDA_COMMON_HPP__

46
modules/core/include/opencv2/core/cuda/datamov_utils.hpp
View File

@@ -40,12 +40,12 @@
//
//M*/
#ifndef __OPENCV_GPU_DATAMOV_UTILS_HPP__
#define __OPENCV_GPU_DATAMOV_UTILS_HPP__
#ifndef __OPENCV_CUDA_DATAMOV_UTILS_HPP__
#define __OPENCV_CUDA_DATAMOV_UTILS_HPP__
#include "common.hpp"
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 200
@@ -59,47 +59,47 @@ namespace cv { namespace gpu { namespace cudev
#if defined(_WIN64) || defined(__LP64__)
// 64-bit register modifier for inlined asm
#define OPENCV_GPU_ASM_PTR "l"
#define OPENCV_CUDA_ASM_PTR "l"
#else
// 32-bit register modifier for inlined asm
#define OPENCV_GPU_ASM_PTR "r"
#define OPENCV_CUDA_ASM_PTR "r"
#endif
template<class T> struct ForceGlob;
#define OPENCV_GPU_DEFINE_FORCE_GLOB(base_type, ptx_type, reg_mod) \
#define OPENCV_CUDA_DEFINE_FORCE_GLOB(base_type, ptx_type, reg_mod) \
template <> struct ForceGlob<base_type> \
{ \
__device__ __forceinline__ static void Load(const base_type* ptr, int offset, base_type& val) \
{ \
asm("ld.global."#ptx_type" %0, [%1];" : "="#reg_mod(val) : OPENCV_GPU_ASM_PTR(ptr + offset)); \
asm("ld.global."#ptx_type" %0, [%1];" : "="#reg_mod(val) : OPENCV_CUDA_ASM_PTR(ptr + offset)); \
} \
};
#define OPENCV_GPU_DEFINE_FORCE_GLOB_B(base_type, ptx_type) \
#define OPENCV_CUDA_DEFINE_FORCE_GLOB_B(base_type, ptx_type) \
template <> struct ForceGlob<base_type> \
{ \
__device__ __forceinline__ static void Load(const base_type* ptr, int offset, base_type& val) \
{ \
asm("ld.global."#ptx_type" %0, [%1];" : "=r"(*reinterpret_cast<uint*>(&val)) : OPENCV_GPU_ASM_PTR(ptr + offset)); \
asm("ld.global."#ptx_type" %0, [%1];" : "=r"(*reinterpret_cast<uint*>(&val)) : OPENCV_CUDA_ASM_PTR(ptr + offset)); \
} \
};
OPENCV_GPU_DEFINE_FORCE_GLOB_B(uchar, u8)
OPENCV_GPU_DEFINE_FORCE_GLOB_B(schar, s8)
OPENCV_GPU_DEFINE_FORCE_GLOB_B(char, b8)
OPENCV_GPU_DEFINE_FORCE_GLOB (ushort, u16, h)
OPENCV_GPU_DEFINE_FORCE_GLOB (short, s16, h)
OPENCV_GPU_DEFINE_FORCE_GLOB (uint, u32, r)
OPENCV_GPU_DEFINE_FORCE_GLOB (int, s32, r)
OPENCV_GPU_DEFINE_FORCE_GLOB (float, f32, f)
OPENCV_GPU_DEFINE_FORCE_GLOB (double, f64, d)
OPENCV_CUDA_DEFINE_FORCE_GLOB_B(uchar, u8)
OPENCV_CUDA_DEFINE_FORCE_GLOB_B(schar, s8)
OPENCV_CUDA_DEFINE_FORCE_GLOB_B(char, b8)
OPENCV_CUDA_DEFINE_FORCE_GLOB (ushort, u16, h)
OPENCV_CUDA_DEFINE_FORCE_GLOB (short, s16, h)
OPENCV_CUDA_DEFINE_FORCE_GLOB (uint, u32, r)
OPENCV_CUDA_DEFINE_FORCE_GLOB (int, s32, r)
OPENCV_CUDA_DEFINE_FORCE_GLOB (float, f32, f)
OPENCV_CUDA_DEFINE_FORCE_GLOB (double, f64, d)
#undef OPENCV_GPU_DEFINE_FORCE_GLOB
#undef OPENCV_GPU_DEFINE_FORCE_GLOB_B
#undef OPENCV_GPU_ASM_PTR
#undef OPENCV_CUDA_DEFINE_FORCE_GLOB
#undef OPENCV_CUDA_DEFINE_FORCE_GLOB_B
#undef OPENCV_CUDA_ASM_PTR
#endif // __CUDA_ARCH__ >= 200
}}} // namespace cv { namespace gpu { namespace cudev
}}} // namespace cv { namespace cuda { namespace cudev
#endif // __OPENCV_GPU_DATAMOV_UTILS_HPP__
#endif // __OPENCV_CUDA_DATAMOV_UTILS_HPP__

118
modules/core/include/opencv2/core/cuda/detail/color_detail.hpp
View File

@@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_GPU_COLOR_DETAIL_HPP__
#define __OPENCV_GPU_COLOR_DETAIL_HPP__
#ifndef __OPENCV_CUDA_COLOR_DETAIL_HPP__
#define __OPENCV_CUDA_COLOR_DETAIL_HPP__
#include "../common.hpp"
#include "../vec_traits.hpp"
@@ -49,7 +49,7 @@
#include "../limits.hpp"
#include "../functional.hpp"
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
#ifndef CV_DESCALE
#define CV_DESCALE(x, n) (((x) + (1 << ((n)-1))) >> (n))
@@ -143,10 +143,10 @@ namespace cv { namespace gpu { namespace cudev
};
}
#define OPENCV_GPU_IMPLEMENT_RGB2RGB_TRAITS(name, scn, dcn, bidx) \
#define OPENCV_CUDA_IMPLEMENT_RGB2RGB_TRAITS(name, scn, dcn, bidx) \
template <typename T> struct name ## _traits \
{ \
typedef ::cv::gpu::cudev::color_detail::RGB2RGB<T, scn, dcn, bidx> functor_type; \
typedef ::cv::cuda::device::color_detail::RGB2RGB<T, scn, dcn, bidx> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -216,10 +216,10 @@ namespace cv { namespace gpu { namespace cudev
};
}
#define OPENCV_GPU_IMPLEMENT_RGB2RGB5x5_TRAITS(name, scn, bidx, green_bits) \
#define OPENCV_CUDA_IMPLEMENT_RGB2RGB5x5_TRAITS(name, scn, bidx, green_bits) \
struct name ## _traits \
{ \
typedef ::cv::gpu::cudev::color_detail::RGB2RGB5x5<scn, bidx, green_bits> functor_type; \
typedef ::cv::cuda::device::color_detail::RGB2RGB5x5<scn, bidx, green_bits> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -297,10 +297,10 @@ namespace cv { namespace gpu { namespace cudev
};
}
#define OPENCV_GPU_IMPLEMENT_RGB5x52RGB_TRAITS(name, dcn, bidx, green_bits) \
#define OPENCV_CUDA_IMPLEMENT_RGB5x52RGB_TRAITS(name, dcn, bidx, green_bits) \
struct name ## _traits \
{ \
typedef ::cv::gpu::cudev::color_detail::RGB5x52RGB<dcn, bidx, green_bits> functor_type; \
typedef ::cv::cuda::device::color_detail::RGB5x52RGB<dcn, bidx, green_bits> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -343,10 +343,10 @@ namespace cv { namespace gpu { namespace cudev
};
}
#define OPENCV_GPU_IMPLEMENT_GRAY2RGB_TRAITS(name, dcn) \
#define OPENCV_CUDA_IMPLEMENT_GRAY2RGB_TRAITS(name, dcn) \
template <typename T> struct name ## _traits \
{ \
typedef ::cv::gpu::cudev::color_detail::Gray2RGB<T, dcn> functor_type; \
typedef ::cv::cuda::device::color_detail::Gray2RGB<T, dcn> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -385,10 +385,10 @@ namespace cv { namespace gpu { namespace cudev
};
}
#define OPENCV_GPU_IMPLEMENT_GRAY2RGB5x5_TRAITS(name, green_bits) \
#define OPENCV_CUDA_IMPLEMENT_GRAY2RGB5x5_TRAITS(name, green_bits) \
struct name ## _traits \
{ \
typedef ::cv::gpu::cudev::color_detail::Gray2RGB5x5<green_bits> functor_type; \
typedef ::cv::cuda::device::color_detail::Gray2RGB5x5<green_bits> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -427,10 +427,10 @@ namespace cv { namespace gpu { namespace cudev
};
}
#define OPENCV_GPU_IMPLEMENT_RGB5x52GRAY_TRAITS(name, green_bits) \
#define OPENCV_CUDA_IMPLEMENT_RGB5x52GRAY_TRAITS(name, green_bits) \
struct name ## _traits \
{ \
typedef ::cv::gpu::cudev::color_detail::RGB5x52Gray<green_bits> functor_type; \
typedef ::cv::cuda::device::color_detail::RGB5x52Gray<green_bits> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -478,10 +478,10 @@ namespace cv { namespace gpu { namespace cudev
};
}
#define OPENCV_GPU_IMPLEMENT_RGB2GRAY_TRAITS(name, scn, bidx) \
#define OPENCV_CUDA_IMPLEMENT_RGB2GRAY_TRAITS(name, scn, bidx) \
template <typename T> struct name ## _traits \
{ \
typedef ::cv::gpu::cudev::color_detail::RGB2Gray<T, scn, bidx> functor_type; \
typedef ::cv::cuda::device::color_detail::RGB2Gray<T, scn, bidx> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -529,10 +529,10 @@ namespace cv { namespace gpu { namespace cudev
};
}
#define OPENCV_GPU_IMPLEMENT_RGB2YUV_TRAITS(name, scn, dcn, bidx) \
#define OPENCV_CUDA_IMPLEMENT_RGB2YUV_TRAITS(name, scn, dcn, bidx) \
template <typename T> struct name ## _traits \
{ \
typedef ::cv::gpu::cudev::color_detail::RGB2YUV<T, scn, dcn, bidx> functor_type; \
typedef ::cv::cuda::device::color_detail::RGB2YUV<T, scn, dcn, bidx> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -617,10 +617,10 @@ namespace cv { namespace gpu { namespace cudev
};
}
#define OPENCV_GPU_IMPLEMENT_YUV2RGB_TRAITS(name, scn, dcn, bidx) \
#define OPENCV_CUDA_IMPLEMENT_YUV2RGB_TRAITS(name, scn, dcn, bidx) \
template <typename T> struct name ## _traits \
{ \
typedef ::cv::gpu::cudev::color_detail::YUV2RGB<T, scn, dcn, bidx> functor_type; \
typedef ::cv::cuda::device::color_detail::YUV2RGB<T, scn, dcn, bidx> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -696,10 +696,10 @@ namespace cv { namespace gpu { namespace cudev
};
}
#define OPENCV_GPU_IMPLEMENT_RGB2YCrCb_TRAITS(name, scn, dcn, bidx) \
#define OPENCV_CUDA_IMPLEMENT_RGB2YCrCb_TRAITS(name, scn, dcn, bidx) \
template <typename T> struct name ## _traits \
{ \
typedef ::cv::gpu::cudev::color_detail::RGB2YCrCb<T, scn, dcn, bidx> functor_type; \
typedef ::cv::cuda::device::color_detail::RGB2YCrCb<T, scn, dcn, bidx> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -775,10 +775,10 @@ namespace cv { namespace gpu { namespace cudev
};
}
#define OPENCV_GPU_IMPLEMENT_YCrCb2RGB_TRAITS(name, scn, dcn, bidx) \
#define OPENCV_CUDA_IMPLEMENT_YCrCb2RGB_TRAITS(name, scn, dcn, bidx) \
template <typename T> struct name ## _traits \
{ \
typedef ::cv::gpu::cudev::color_detail::YCrCb2RGB<T, scn, dcn, bidx> functor_type; \
typedef ::cv::cuda::device::color_detail::YCrCb2RGB<T, scn, dcn, bidx> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -851,10 +851,10 @@ namespace cv { namespace gpu { namespace cudev
};
}
#define OPENCV_GPU_IMPLEMENT_RGB2XYZ_TRAITS(name, scn, dcn, bidx) \
#define OPENCV_CUDA_IMPLEMENT_RGB2XYZ_TRAITS(name, scn, dcn, bidx) \
template <typename T> struct name ## _traits \
{ \
typedef ::cv::gpu::cudev::color_detail::RGB2XYZ<T, scn, dcn, bidx> functor_type; \
typedef ::cv::cuda::device::color_detail::RGB2XYZ<T, scn, dcn, bidx> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -926,10 +926,10 @@ namespace cv { namespace gpu { namespace cudev
};
}
#define OPENCV_GPU_IMPLEMENT_XYZ2RGB_TRAITS(name, scn, dcn, bidx) \
#define OPENCV_CUDA_IMPLEMENT_XYZ2RGB_TRAITS(name, scn, dcn, bidx) \
template <typename T> struct name ## _traits \
{ \
typedef ::cv::gpu::cudev::color_detail::XYZ2RGB<T, scn, dcn, bidx> functor_type; \
typedef ::cv::cuda::device::color_detail::XYZ2RGB<T, scn, dcn, bidx> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -1064,10 +1064,10 @@ namespace cv { namespace gpu { namespace cudev
};
}
#define OPENCV_GPU_IMPLEMENT_RGB2HSV_TRAITS(name, scn, dcn, bidx) \
#define OPENCV_CUDA_IMPLEMENT_RGB2HSV_TRAITS(name, scn, dcn, bidx) \
template <typename T> struct name ## _traits \
{ \
typedef ::cv::gpu::cudev::color_detail::RGB2HSV<T, scn, dcn, bidx, 180> functor_type; \
typedef ::cv::cuda::device::color_detail::RGB2HSV<T, scn, dcn, bidx, 180> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -1075,7 +1075,7 @@ namespace cv { namespace gpu { namespace cudev
}; \
template <typename T> struct name ## _full_traits \
{ \
typedef ::cv::gpu::cudev::color_detail::RGB2HSV<T, scn, dcn, bidx, 256> functor_type; \
typedef ::cv::cuda::device::color_detail::RGB2HSV<T, scn, dcn, bidx, 256> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -1083,7 +1083,7 @@ namespace cv { namespace gpu { namespace cudev
}; \
template <> struct name ## _traits<float> \
{ \
typedef ::cv::gpu::cudev::color_detail::RGB2HSV<float, scn, dcn, bidx, 360> functor_type; \
typedef ::cv::cuda::device::color_detail::RGB2HSV<float, scn, dcn, bidx, 360> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -1091,7 +1091,7 @@ namespace cv { namespace gpu { namespace cudev
}; \
template <> struct name ## _full_traits<float> \
{ \
typedef ::cv::gpu::cudev::color_detail::RGB2HSV<float, scn, dcn, bidx, 360> functor_type; \
typedef ::cv::cuda::device::color_detail::RGB2HSV<float, scn, dcn, bidx, 360> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -1204,10 +1204,10 @@ namespace cv { namespace gpu { namespace cudev
};
}
#define OPENCV_GPU_IMPLEMENT_HSV2RGB_TRAITS(name, scn, dcn, bidx) \
#define OPENCV_CUDA_IMPLEMENT_HSV2RGB_TRAITS(name, scn, dcn, bidx) \
template <typename T> struct name ## _traits \
{ \
typedef ::cv::gpu::cudev::color_detail::HSV2RGB<T, scn, dcn, bidx, 180> functor_type; \
typedef ::cv::cuda::device::color_detail::HSV2RGB<T, scn, dcn, bidx, 180> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -1215,7 +1215,7 @@ namespace cv { namespace gpu { namespace cudev
}; \
template <typename T> struct name ## _full_traits \
{ \
typedef ::cv::gpu::cudev::color_detail::HSV2RGB<T, scn, dcn, bidx, 255> functor_type; \
typedef ::cv::cuda::device::color_detail::HSV2RGB<T, scn, dcn, bidx, 255> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -1223,7 +1223,7 @@ namespace cv { namespace gpu { namespace cudev
}; \
template <> struct name ## _traits<float> \
{ \
typedef ::cv::gpu::cudev::color_detail::HSV2RGB<float, scn, dcn, bidx, 360> functor_type; \
typedef ::cv::cuda::device::color_detail::HSV2RGB<float, scn, dcn, bidx, 360> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -1231,7 +1231,7 @@ namespace cv { namespace gpu { namespace cudev
}; \
template <> struct name ## _full_traits<float> \
{ \
typedef ::cv::gpu::cudev::color_detail::HSV2RGB<float, scn, dcn, bidx, 360> functor_type; \
typedef ::cv::cuda::device::color_detail::HSV2RGB<float, scn, dcn, bidx, 360> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -1337,10 +1337,10 @@ namespace cv { namespace gpu { namespace cudev
};
}
#define OPENCV_GPU_IMPLEMENT_RGB2HLS_TRAITS(name, scn, dcn, bidx) \
#define OPENCV_CUDA_IMPLEMENT_RGB2HLS_TRAITS(name, scn, dcn, bidx) \
template <typename T> struct name ## _traits \
{ \
typedef ::cv::gpu::cudev::color_detail::RGB2HLS<T, scn, dcn, bidx, 180> functor_type; \
typedef ::cv::cuda::device::color_detail::RGB2HLS<T, scn, dcn, bidx, 180> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -1348,7 +1348,7 @@ namespace cv { namespace gpu { namespace cudev
}; \
template <typename T> struct name ## _full_traits \
{ \
typedef ::cv::gpu::cudev::color_detail::RGB2HLS<T, scn, dcn, bidx, 256> functor_type; \
typedef ::cv::cuda::device::color_detail::RGB2HLS<T, scn, dcn, bidx, 256> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -1356,7 +1356,7 @@ namespace cv { namespace gpu { namespace cudev
}; \
template <> struct name ## _traits<float> \
{ \
typedef ::cv::gpu::cudev::color_detail::RGB2HLS<float, scn, dcn, bidx, 360> functor_type; \
typedef ::cv::cuda::device::color_detail::RGB2HLS<float, scn, dcn, bidx, 360> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -1364,7 +1364,7 @@ namespace cv { namespace gpu { namespace cudev
}; \
template <> struct name ## _full_traits<float> \
{ \
typedef ::cv::gpu::cudev::color_detail::RGB2HLS<float, scn, dcn, bidx, 360> functor_type; \
typedef ::cv::cuda::device::color_detail::RGB2HLS<float, scn, dcn, bidx, 360> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -1477,10 +1477,10 @@ namespace cv { namespace gpu { namespace cudev
};
}
#define OPENCV_GPU_IMPLEMENT_HLS2RGB_TRAITS(name, scn, dcn, bidx) \
#define OPENCV_CUDA_IMPLEMENT_HLS2RGB_TRAITS(name, scn, dcn, bidx) \
template <typename T> struct name ## _traits \
{ \
typedef ::cv::gpu::cudev::color_detail::HLS2RGB<T, scn, dcn, bidx, 180> functor_type; \
typedef ::cv::cuda::device::color_detail::HLS2RGB<T, scn, dcn, bidx, 180> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -1488,7 +1488,7 @@ namespace cv { namespace gpu { namespace cudev
}; \
template <typename T> struct name ## _full_traits \
{ \
typedef ::cv::gpu::cudev::color_detail::HLS2RGB<T, scn, dcn, bidx, 255> functor_type; \
typedef ::cv::cuda::device::color_detail::HLS2RGB<T, scn, dcn, bidx, 255> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -1496,7 +1496,7 @@ namespace cv { namespace gpu { namespace cudev
}; \
template <> struct name ## _traits<float> \
{ \
typedef ::cv::gpu::cudev::color_detail::HLS2RGB<float, scn, dcn, bidx, 360> functor_type; \
typedef ::cv::cuda::device::color_detail::HLS2RGB<float, scn, dcn, bidx, 360> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -1504,7 +1504,7 @@ namespace cv { namespace gpu { namespace cudev
}; \
template <> struct name ## _full_traits<float> \
{ \
typedef ::cv::gpu::cudev::color_detail::HLS2RGB<float, scn, dcn, bidx, 360> functor_type; \
typedef ::cv::cuda::device::color_detail::HLS2RGB<float, scn, dcn, bidx, 360> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -1646,10 +1646,10 @@ namespace cv { namespace gpu { namespace cudev
};
}
#define OPENCV_GPU_IMPLEMENT_RGB2Lab_TRAITS(name, scn, dcn, srgb, blueIdx) \
#define OPENCV_CUDA_IMPLEMENT_RGB2Lab_TRAITS(name, scn, dcn, srgb, blueIdx) \
template <typename T> struct name ## _traits \
{ \
typedef ::cv::gpu::cudev::color_detail::RGB2Lab<T, scn, dcn, srgb, blueIdx> functor_type; \
typedef ::cv::cuda::device::color_detail::RGB2Lab<T, scn, dcn, srgb, blueIdx> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -1759,10 +1759,10 @@ namespace cv { namespace gpu { namespace cudev
};
}
#define OPENCV_GPU_IMPLEMENT_Lab2RGB_TRAITS(name, scn, dcn, srgb, blueIdx) \
#define OPENCV_CUDA_IMPLEMENT_Lab2RGB_TRAITS(name, scn, dcn, srgb, blueIdx) \
template <typename T> struct name ## _traits \
{ \
typedef ::cv::gpu::cudev::color_detail::Lab2RGB<T, scn, dcn, srgb, blueIdx> functor_type; \
typedef ::cv::cuda::device::color_detail::Lab2RGB<T, scn, dcn, srgb, blueIdx> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -1858,10 +1858,10 @@ namespace cv { namespace gpu { namespace cudev
};
}
#define OPENCV_GPU_IMPLEMENT_RGB2Luv_TRAITS(name, scn, dcn, srgb, blueIdx) \
#define OPENCV_CUDA_IMPLEMENT_RGB2Luv_TRAITS(name, scn, dcn, srgb, blueIdx) \
template <typename T> struct name ## _traits \
{ \
typedef ::cv::gpu::cudev::color_detail::RGB2Luv<T, scn, dcn, srgb, blueIdx> functor_type; \
typedef ::cv::cuda::device::color_detail::RGB2Luv<T, scn, dcn, srgb, blueIdx> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -1959,10 +1959,10 @@ namespace cv { namespace gpu { namespace cudev
};
}
#define OPENCV_GPU_IMPLEMENT_Luv2RGB_TRAITS(name, scn, dcn, srgb, blueIdx) \
#define OPENCV_CUDA_IMPLEMENT_Luv2RGB_TRAITS(name, scn, dcn, srgb, blueIdx) \
template <typename T> struct name ## _traits \
{ \
typedef ::cv::gpu::cudev::color_detail::Luv2RGB<T, scn, dcn, srgb, blueIdx> functor_type; \
typedef ::cv::cuda::device::color_detail::Luv2RGB<T, scn, dcn, srgb, blueIdx> functor_type; \
static __host__ __device__ __forceinline__ functor_type create_functor() \
{ \
return functor_type(); \
@@ -1971,6 +1971,6 @@ namespace cv { namespace gpu { namespace cudev
#undef CV_DESCALE
}}} // namespace cv { namespace gpu { namespace cudev
}}} // namespace cv { namespace cuda { namespace cudev
#endif // __OPENCV_GPU_COLOR_DETAIL_HPP__
#endif // __OPENCV_CUDA_COLOR_DETAIL_HPP__

8
modules/core/include/opencv2/core/cuda/detail/reduce.hpp
View File

@@ -40,14 +40,14 @@
//
//M*/
#ifndef __OPENCV_GPU_REDUCE_DETAIL_HPP__
#define __OPENCV_GPU_REDUCE_DETAIL_HPP__
#ifndef __OPENCV_CUDA_REDUCE_DETAIL_HPP__
#define __OPENCV_CUDA_REDUCE_DETAIL_HPP__
#include <thrust/tuple.h>
#include "../warp.hpp"
#include "../warp_shuffle.hpp"
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
namespace reduce_detail
{
@@ -358,4 +358,4 @@ namespace cv { namespace gpu { namespace cudev
}
}}}
#endif // __OPENCV_GPU_REDUCE_DETAIL_HPP__
#endif // __OPENCV_CUDA_REDUCE_DETAIL_HPP__

8
modules/core/include/opencv2/core/cuda/detail/reduce_key_val.hpp
View File

@@ -40,14 +40,14 @@
//
//M*/
#ifndef __OPENCV_GPU_PRED_VAL_REDUCE_DETAIL_HPP__
#define __OPENCV_GPU_PRED_VAL_REDUCE_DETAIL_HPP__
#ifndef __OPENCV_CUDA_PRED_VAL_REDUCE_DETAIL_HPP__
#define __OPENCV_CUDA_PRED_VAL_REDUCE_DETAIL_HPP__
#include <thrust/tuple.h>
#include "../warp.hpp"
#include "../warp_shuffle.hpp"
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
namespace reduce_key_val_detail
{
@@ -495,4 +495,4 @@ namespace cv { namespace gpu { namespace cudev
}
}}}
#endif // __OPENCV_GPU_PRED_VAL_REDUCE_DETAIL_HPP__
#endif // __OPENCV_CUDA_PRED_VAL_REDUCE_DETAIL_HPP__

10
modules/core/include/opencv2/core/cuda/detail/transform_detail.hpp
View File

@@ -40,14 +40,14 @@
//
//M*/
#ifndef __OPENCV_GPU_TRANSFORM_DETAIL_HPP__
#define __OPENCV_GPU_TRANSFORM_DETAIL_HPP__
#ifndef __OPENCV_CUDA_TRANSFORM_DETAIL_HPP__
#define __OPENCV_CUDA_TRANSFORM_DETAIL_HPP__
#include "../common.hpp"
#include "../vec_traits.hpp"
#include "../functional.hpp"
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
namespace transform_detail
{
@@ -390,6 +390,6 @@ namespace cv { namespace gpu { namespace cudev
}
};
} // namespace transform_detail
}}} // namespace cv { namespace gpu { namespace cudev
}}} // namespace cv { namespace cuda { namespace cudev
#endif // __OPENCV_GPU_TRANSFORM_DETAIL_HPP__
#endif // __OPENCV_CUDA_TRANSFORM_DETAIL_HPP__

10
modules/core/include/opencv2/core/cuda/detail/type_traits_detail.hpp
View File

@@ -40,13 +40,13 @@
//
//M*/
#ifndef __OPENCV_GPU_TYPE_TRAITS_DETAIL_HPP__
#define __OPENCV_GPU_TYPE_TRAITS_DETAIL_HPP__
#ifndef __OPENCV_CUDA_TYPE_TRAITS_DETAIL_HPP__
#define __OPENCV_CUDA_TYPE_TRAITS_DETAIL_HPP__
#include "../common.hpp"
#include "../vec_traits.hpp"
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
namespace type_traits_detail
{
@@ -182,6 +182,6 @@ namespace cv { namespace gpu { namespace cudev
enum { value = 1 };
};
} // namespace type_traits_detail
}}} // namespace cv { namespace gpu { namespace cudev
}}} // namespace cv { namespace cuda { namespace cudev
#endif // __OPENCV_GPU_TYPE_TRAITS_DETAIL_HPP__
#endif // __OPENCV_CUDA_TYPE_TRAITS_DETAIL_HPP__

10
modules/core/include/opencv2/core/cuda/detail/vec_distance_detail.hpp
View File

@@ -40,12 +40,12 @@
//
//M*/
#ifndef __OPENCV_GPU_VEC_DISTANCE_DETAIL_HPP__
#define __OPENCV_GPU_VEC_DISTANCE_DETAIL_HPP__
#ifndef __OPENCV_CUDA_VEC_DISTANCE_DETAIL_HPP__
#define __OPENCV_CUDA_VEC_DISTANCE_DETAIL_HPP__
#include "../datamov_utils.hpp"
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
namespace vec_distance_detail
{
@@ -112,6 +112,6 @@ namespace cv { namespace gpu { namespace cudev
}
};
} // namespace vec_distance_detail
}}} // namespace cv { namespace gpu { namespace cudev
}}} // namespace cv { namespace cuda { namespace cudev
#endif // __OPENCV_GPU_VEC_DISTANCE_DETAIL_HPP__
#endif // __OPENCV_CUDA_VEC_DISTANCE_DETAIL_HPP__

8
modules/core/include/opencv2/core/cuda/dynamic_smem.hpp
View File

@@ -40,10 +40,10 @@
//
//M*/
#ifndef __OPENCV_GPU_DYNAMIC_SMEM_HPP__
#define __OPENCV_GPU_DYNAMIC_SMEM_HPP__
#ifndef __OPENCV_CUDA_DYNAMIC_SMEM_HPP__
#define __OPENCV_CUDA_DYNAMIC_SMEM_HPP__
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
template<class T> struct DynamicSharedMem
{
@@ -77,4 +77,4 @@ namespace cv { namespace gpu { namespace cudev
};
}}}
#endif // __OPENCV_GPU_DYNAMIC_SMEM_HPP__
#endif // __OPENCV_CUDA_DYNAMIC_SMEM_HPP__

10
modules/core/include/opencv2/core/cuda/emulation.hpp
View File

@@ -40,13 +40,13 @@
//
//M*/
#ifndef OPENCV_GPU_EMULATION_HPP_
#define OPENCV_GPU_EMULATION_HPP_
#ifndef OPENCV_CUDA_EMULATION_HPP_
#define OPENCV_CUDA_EMULATION_HPP_
#include "common.hpp"
#include "warp_reduce.hpp"
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
struct Emulation
{
@@ -256,6 +256,6 @@ namespace cv { namespace gpu { namespace cudev
}
};
}; //struct Emulation
}}} // namespace cv { namespace gpu { namespace cudev
}}} // namespace cv { namespace cuda { namespace cudev
#endif /* OPENCV_GPU_EMULATION_HPP_ */
#endif /* OPENCV_CUDA_EMULATION_HPP_ */

10
modules/core/include/opencv2/core/cuda/filters.hpp
View File

@@ -40,15 +40,15 @@
//
//M*/
#ifndef __OPENCV_GPU_FILTERS_HPP__
#define __OPENCV_GPU_FILTERS_HPP__
#ifndef __OPENCV_CUDA_FILTERS_HPP__
#define __OPENCV_CUDA_FILTERS_HPP__
#include "saturate_cast.hpp"
#include "vec_traits.hpp"
#include "vec_math.hpp"
#include "type_traits.hpp"
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
template <typename Ptr2D> struct PointFilter
{
@@ -273,6 +273,6 @@ namespace cv { namespace gpu { namespace cudev
float scale_x, scale_y;
int width, haight;
};
}}} // namespace cv { namespace gpu { namespace cudev
}}} // namespace cv { namespace cuda { namespace cudev
#endif // __OPENCV_GPU_FILTERS_HPP__
#endif // __OPENCV_CUDA_FILTERS_HPP__

10
modules/core/include/opencv2/core/cuda/funcattrib.hpp
View File

@@ -40,12 +40,12 @@
//
//M*/
#ifndef __OPENCV_GPU_DEVICE_FUNCATTRIB_HPP_
#define __OPENCV_GPU_DEVICE_FUNCATTRIB_HPP_
#ifndef __OPENCV_CUDA_DEVICE_FUNCATTRIB_HPP_
#define __OPENCV_CUDA_DEVICE_FUNCATTRIB_HPP_
#include <cstdio>
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
template<class Func>
void printFuncAttrib(Func& func)
@@ -66,6 +66,6 @@ namespace cv { namespace gpu { namespace cudev
printf("\n");
fflush(stdout);
}
}}} // namespace cv { namespace gpu { namespace cudev
}}} // namespace cv { namespace cuda { namespace cudev
#endif /* __OPENCV_GPU_DEVICE_FUNCATTRIB_HPP_ */
#endif /* __OPENCV_CUDA_DEVICE_FUNCATTRIB_HPP_ */

106
modules/core/include/opencv2/core/cuda/functional.hpp
View File

@@ -40,8 +40,8 @@
//
//M*/
#ifndef __OPENCV_GPU_FUNCTIONAL_HPP__
#define __OPENCV_GPU_FUNCTIONAL_HPP__
#ifndef __OPENCV_CUDA_FUNCTIONAL_HPP__
#define __OPENCV_CUDA_FUNCTIONAL_HPP__
#include <functional>
#include "saturate_cast.hpp"
@@ -49,7 +49,7 @@
#include "type_traits.hpp"
#include "device_functions.h"
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
// Function Objects
template<typename Argument, typename Result> struct unary_function : public std::unary_function<Argument, Result> {};
@@ -298,7 +298,7 @@ namespace cv { namespace gpu { namespace cudev
// Min/Max Operations
#define OPENCV_GPU_IMPLEMENT_MINMAX(name, type, op) \
#define OPENCV_CUDA_IMPLEMENT_MINMAX(name, type, op) \
template <> struct name<type> : binary_function<type, type, type> \
{ \
__device__ __forceinline__ type operator()(type lhs, type rhs) const {return op(lhs, rhs);} \
@@ -316,15 +316,15 @@ namespace cv { namespace gpu { namespace cudev
__host__ __device__ __forceinline__ maximum(const maximum&) {}
};
OPENCV_GPU_IMPLEMENT_MINMAX(maximum, uchar, ::max)
OPENCV_GPU_IMPLEMENT_MINMAX(maximum, schar, ::max)
OPENCV_GPU_IMPLEMENT_MINMAX(maximum, char, ::max)
OPENCV_GPU_IMPLEMENT_MINMAX(maximum, ushort, ::max)
OPENCV_GPU_IMPLEMENT_MINMAX(maximum, short, ::max)
OPENCV_GPU_IMPLEMENT_MINMAX(maximum, int, ::max)
OPENCV_GPU_IMPLEMENT_MINMAX(maximum, uint, ::max)
OPENCV_GPU_IMPLEMENT_MINMAX(maximum, float, ::fmax)
OPENCV_GPU_IMPLEMENT_MINMAX(maximum, double, ::fmax)
OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, uchar, ::max)
OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, schar, ::max)
OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, char, ::max)
OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, ushort, ::max)
OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, short, ::max)
OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, int, ::max)
OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, uint, ::max)
OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, float, ::fmax)
OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, double, ::fmax)
template <typename T> struct minimum : binary_function<T, T, T>
{
@@ -336,17 +336,17 @@ namespace cv { namespace gpu { namespace cudev
__host__ __device__ __forceinline__ minimum(const minimum&) {}
};
OPENCV_GPU_IMPLEMENT_MINMAX(minimum, uchar, ::min)
OPENCV_GPU_IMPLEMENT_MINMAX(minimum, schar, ::min)
OPENCV_GPU_IMPLEMENT_MINMAX(minimum, char, ::min)
OPENCV_GPU_IMPLEMENT_MINMAX(minimum, ushort, ::min)
OPENCV_GPU_IMPLEMENT_MINMAX(minimum, short, ::min)
OPENCV_GPU_IMPLEMENT_MINMAX(minimum, int, ::min)
OPENCV_GPU_IMPLEMENT_MINMAX(minimum, uint, ::min)
OPENCV_GPU_IMPLEMENT_MINMAX(minimum, float, ::fmin)
OPENCV_GPU_IMPLEMENT_MINMAX(minimum, double, ::fmin)
OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, uchar, ::min)
OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, schar, ::min)
OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, char, ::min)
OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, ushort, ::min)
OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, short, ::min)
OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, int, ::min)
OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, uint, ::min)
OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, float, ::fmin)
OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, double, ::fmin)
#undef OPENCV_GPU_IMPLEMENT_MINMAX
#undef OPENCV_CUDA_IMPLEMENT_MINMAX
// Math functions
@@ -451,7 +451,7 @@ namespace cv { namespace gpu { namespace cudev
__host__ __device__ __forceinline__ abs_func(const abs_func&) {}
};
#define OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(name, func) \
#define OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(name, func) \
template <typename T> struct name ## _func : unary_function<T, float> \
{ \
__device__ __forceinline__ float operator ()(typename TypeTraits<T>::ParameterType v) const \
@@ -471,7 +471,7 @@ namespace cv { namespace gpu { namespace cudev
__host__ __device__ __forceinline__ name ## _func(const name ## _func&) {} \
};
#define OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR(name, func) \
#define OPENCV_CUDA_IMPLEMENT_BIN_FUNCTOR(name, func) \
template <typename T> struct name ## _func : binary_function<T, T, float> \
{ \
__device__ __forceinline__ float operator ()(typename TypeTraits<T>::ParameterType v1, typename TypeTraits<T>::ParameterType v2) const \
@@ -491,33 +491,33 @@ namespace cv { namespace gpu { namespace cudev
__host__ __device__ __forceinline__ name ## _func(const name ## _func&) {} \
};
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(sqrt, ::sqrt)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(exp, ::exp)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(exp2, ::exp2)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(exp10, ::exp10)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(log, ::log)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(log2, ::log2)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(log10, ::log10)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(sin, ::sin)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(cos, ::cos)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(tan, ::tan)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(asin, ::asin)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(acos, ::acos)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(atan, ::atan)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(sinh, ::sinh)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(cosh, ::cosh)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(tanh, ::tanh)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(asinh, ::asinh)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(acosh, ::acosh)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(atanh, ::atanh)
OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(sqrt, ::sqrt)
OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(exp, ::exp)
OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(exp2, ::exp2)
OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(exp10, ::exp10)
OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(log, ::log)
OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(log2, ::log2)
OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(log10, ::log10)
OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(sin, ::sin)
OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(cos, ::cos)
OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(tan, ::tan)
OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(asin, ::asin)
OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(acos, ::acos)
OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(atan, ::atan)
OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(sinh, ::sinh)
OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(cosh, ::cosh)
OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(tanh, ::tanh)
OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(asinh, ::asinh)
OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(acosh, ::acosh)
OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(atanh, ::atanh)
OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR(hypot, ::hypot)
OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR(atan2, ::atan2)
OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR(pow, ::pow)
OPENCV_CUDA_IMPLEMENT_BIN_FUNCTOR(hypot, ::hypot)
OPENCV_CUDA_IMPLEMENT_BIN_FUNCTOR(atan2, ::atan2)
OPENCV_CUDA_IMPLEMENT_BIN_FUNCTOR(pow, ::pow)
#undef OPENCV_GPU_IMPLEMENT_UN_FUNCTOR
#undef OPENCV_GPU_IMPLEMENT_UN_FUNCTOR_NO_DOUBLE
#undef OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR
#undef OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR
#undef OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR_NO_DOUBLE
#undef OPENCV_CUDA_IMPLEMENT_BIN_FUNCTOR
template<typename T> struct hypot_sqr_func : binary_function<T, T, float>
{
@@ -782,8 +782,8 @@ namespace cv { namespace gpu { namespace cudev
template <typename Func> struct TransformFunctorTraits : DefaultTransformFunctorTraits<Func> {};
#define OPENCV_GPU_TRANSFORM_FUNCTOR_TRAITS(type) \
#define OPENCV_CUDA_TRANSFORM_FUNCTOR_TRAITS(type) \
template <> struct TransformFunctorTraits< type > : DefaultTransformFunctorTraits< type >
}}} // namespace cv { namespace gpu { namespace cudev
}}} // namespace cv { namespace cuda { namespace cudev
#endif // __OPENCV_GPU_FUNCTIONAL_HPP__
#endif // __OPENCV_CUDA_FUNCTIONAL_HPP__

10
modules/core/include/opencv2/core/cuda/limits.hpp
View File

@@ -40,14 +40,14 @@
//
//M*/
#ifndef __OPENCV_GPU_LIMITS_GPU_HPP__
#define __OPENCV_GPU_LIMITS_GPU_HPP__
#ifndef __OPENCV_CUDA_LIMITS_HPP__
#define __OPENCV_CUDA_LIMITS_HPP__
#include <limits.h>
#include <float.h>
#include "common.hpp"
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
template <class T> struct numeric_limits;
@@ -117,6 +117,6 @@ template <> struct numeric_limits<double>
static const bool is_signed = true;
};
}}} // namespace cv { namespace gpu { namespace cudev {
}}} // namespace cv { namespace cuda { namespace cudev {
#endif // __OPENCV_GPU_LIMITS_GPU_HPP__
#endif // __OPENCV_CUDA_LIMITS_HPP__

8
modules/core/include/opencv2/core/cuda/reduce.hpp
View File

@@ -40,14 +40,14 @@
//
//M*/
#ifndef __OPENCV_GPU_REDUCE_HPP__
#define __OPENCV_GPU_REDUCE_HPP__
#ifndef __OPENCV_CUDA_REDUCE_HPP__
#define __OPENCV_CUDA_REDUCE_HPP__
#include <thrust/tuple.h>
#include "detail/reduce.hpp"
#include "detail/reduce_key_val.hpp"
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
template <int N, typename T, class Op>
__device__ __forceinline__ void reduce(volatile T* smem, T& val, unsigned int tid, const Op& op)
@@ -194,4 +194,4 @@ namespace cv { namespace gpu { namespace cudev
}
}}}
#endif // __OPENCV_GPU_UTILITY_HPP__
#endif // __OPENCV_CUDA_UTILITY_HPP__

8
modules/core/include/opencv2/core/cuda/saturate_cast.hpp
View File

@@ -40,12 +40,12 @@
//
//M*/
#ifndef __OPENCV_GPU_SATURATE_CAST_HPP__
#define __OPENCV_GPU_SATURATE_CAST_HPP__
#ifndef __OPENCV_CUDA_SATURATE_CAST_HPP__
#define __OPENCV_CUDA_SATURATE_CAST_HPP__
#include "common.hpp"
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(uchar v) { return _Tp(v); }
template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(schar v) { return _Tp(v); }
@@ -281,4 +281,4 @@ namespace cv { namespace gpu { namespace cudev
}
}}}
#endif /* __OPENCV_GPU_SATURATE_CAST_HPP__ */
#endif /* __OPENCV_CUDA_SATURATE_CAST_HPP__ */

28
modules/core/include/opencv2/core/cuda/scan.hpp
View File

@@ -40,15 +40,15 @@
//
//M*/
#ifndef __OPENCV_GPU_SCAN_HPP__
#define __OPENCV_GPU_SCAN_HPP__
#ifndef __OPENCV_CUDA_SCAN_HPP__
#define __OPENCV_CUDA_SCAN_HPP__
#include "opencv2/core/cuda/common.hpp"
#include "opencv2/core/cuda/utility.hpp"
#include "opencv2/core/cuda/warp.hpp"
#include "opencv2/core/cuda/warp_shuffle.hpp"
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
enum ScanKind { EXCLUSIVE = 0, INCLUSIVE = 1 };
@@ -174,22 +174,22 @@ namespace cv { namespace gpu { namespace cudev
__device__ T warpScanInclusive(T idata, volatile T* s_Data, unsigned int tid)
{
#if __CUDA_ARCH__ >= 300
const unsigned int laneId = cv::gpu::cudev::Warp::laneId();
const unsigned int laneId = cv::cuda::device::Warp::laneId();
// scan on shuffl functions
#pragma unroll
for (int i = 1; i <= (OPENCV_GPU_WARP_SIZE / 2); i *= 2)
for (int i = 1; i <= (OPENCV_CUDA_WARP_SIZE / 2); i *= 2)
{
const T n = cv::gpu::cudev::shfl_up(idata, i);
const T n = cv::cuda::device::shfl_up(idata, i);
if (laneId >= i)
idata += n;
}
return idata;
#else
unsigned int pos = 2 * tid - (tid & (OPENCV_GPU_WARP_SIZE - 1));
unsigned int pos = 2 * tid - (tid & (OPENCV_CUDA_WARP_SIZE - 1));
s_Data[pos] = 0;
pos += OPENCV_GPU_WARP_SIZE;
pos += OPENCV_CUDA_WARP_SIZE;
s_Data[pos] = idata;
s_Data[pos] += s_Data[pos - 1];
@@ -211,7 +211,7 @@ namespace cv { namespace gpu { namespace cudev
template <int tiNumScanThreads, typename T>
__device__ T blockScanInclusive(T idata, volatile T* s_Data, unsigned int tid)
{
if (tiNumScanThreads > OPENCV_GPU_WARP_SIZE)
if (tiNumScanThreads > OPENCV_CUDA_WARP_SIZE)
{
//Bottom-level inclusive warp scan
T warpResult = warpScanInclusive(idata, s_Data, tid);
@@ -219,15 +219,15 @@ namespace cv { namespace gpu { namespace cudev
//Save top elements of each warp for exclusive warp scan
//sync to wait for warp scans to complete (because s_Data is being overwritten)
__syncthreads();
if ((tid & (OPENCV_GPU_WARP_SIZE - 1)) == (OPENCV_GPU_WARP_SIZE - 1))
if ((tid & (OPENCV_CUDA_WARP_SIZE - 1)) == (OPENCV_CUDA_WARP_SIZE - 1))
{
s_Data[tid >> OPENCV_GPU_LOG_WARP_SIZE] = warpResult;
s_Data[tid >> OPENCV_CUDA_LOG_WARP_SIZE] = warpResult;
}
//wait for warp scans to complete
__syncthreads();
if (tid < (tiNumScanThreads / OPENCV_GPU_WARP_SIZE) )
if (tid < (tiNumScanThreads / OPENCV_CUDA_WARP_SIZE) )
{
//grab top warp elements
T val = s_Data[tid];
@@ -238,7 +238,7 @@ namespace cv { namespace gpu { namespace cudev
//return updated warp scans with exclusive scan results
__syncthreads();
return warpResult + s_Data[tid >> OPENCV_GPU_LOG_WARP_SIZE];
return warpResult + s_Data[tid >> OPENCV_CUDA_LOG_WARP_SIZE];
}
else
{
@@ -247,4 +247,4 @@ namespace cv { namespace gpu { namespace cudev
}
}}}
#endif // __OPENCV_GPU_SCAN_HPP__
#endif // __OPENCV_CUDA_SCAN_HPP__

8
modules/core/include/opencv2/core/cuda/simd_functions.hpp
View File

@@ -70,8 +70,8 @@
* POSSIBILITY OF SUCH DAMAGE.
*/
#ifndef __OPENCV_GPU_SIMD_FUNCTIONS_HPP__
#define __OPENCV_GPU_SIMD_FUNCTIONS_HPP__
#ifndef __OPENCV_CUDA_SIMD_FUNCTIONS_HPP__
#define __OPENCV_CUDA_SIMD_FUNCTIONS_HPP__
#include "common.hpp"
@@ -123,7 +123,7 @@
vmin4(a,b) per-byte unsigned minimum: min(a, b)
*/
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
// 2
@@ -906,4 +906,4 @@ namespace cv { namespace gpu { namespace cudev
}
}}}
#endif // __OPENCV_GPU_SIMD_FUNCTIONS_HPP__
#endif // __OPENCV_CUDA_SIMD_FUNCTIONS_HPP__

8
modules/core/include/opencv2/core/cuda/transform.hpp
View File

@@ -40,14 +40,14 @@
//
//M*/
#ifndef __OPENCV_GPU_TRANSFORM_HPP__
#define __OPENCV_GPU_TRANSFORM_HPP__
#ifndef __OPENCV_CUDA_TRANSFORM_HPP__
#define __OPENCV_CUDA_TRANSFORM_HPP__
#include "common.hpp"
#include "utility.hpp"
#include "detail/transform_detail.hpp"
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
template <typename T, typename D, typename UnOp, typename Mask>
static inline void transform(PtrStepSz<T> src, PtrStepSz<D> dst, UnOp op, const Mask& mask, cudaStream_t stream)
@@ -64,4 +64,4 @@ namespace cv { namespace gpu { namespace cudev
}
}}}
#endif // __OPENCV_GPU_TRANSFORM_HPP__
#endif // __OPENCV_CUDA_TRANSFORM_HPP__

8
modules/core/include/opencv2/core/cuda/type_traits.hpp
View File

@@ -40,12 +40,12 @@
//
//M*/
#ifndef __OPENCV_GPU_TYPE_TRAITS_HPP__
#define __OPENCV_GPU_TYPE_TRAITS_HPP__
#ifndef __OPENCV_CUDA_TYPE_TRAITS_HPP__
#define __OPENCV_CUDA_TYPE_TRAITS_HPP__
#include "detail/type_traits_detail.hpp"
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
template <typename T> struct IsSimpleParameter
{
@@ -79,4 +79,4 @@ namespace cv { namespace gpu { namespace cudev
};
}}}
#endif // __OPENCV_GPU_TYPE_TRAITS_HPP__
#endif // __OPENCV_CUDA_TYPE_TRAITS_HPP__

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

@@ -40,18 +40,18 @@
//
//M*/
#ifndef __OPENCV_GPU_UTILITY_HPP__
#define __OPENCV_GPU_UTILITY_HPP__
#ifndef __OPENCV_CUDA_UTILITY_HPP__
#define __OPENCV_CUDA_UTILITY_HPP__
#include "saturate_cast.hpp"
#include "datamov_utils.hpp"
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
#define OPENCV_GPU_LOG_WARP_SIZE (5)
#define OPENCV_GPU_WARP_SIZE (1 << OPENCV_GPU_LOG_WARP_SIZE)
#define OPENCV_GPU_LOG_MEM_BANKS ((__CUDA_ARCH__ >= 200) ? 5 : 4) // 32 banks on fermi, 16 on tesla
#define OPENCV_GPU_MEM_BANKS (1 << OPENCV_GPU_LOG_MEM_BANKS)
#define OPENCV_CUDA_LOG_WARP_SIZE (5)
#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_MEM_BANKS (1 << OPENCV_CUDA_LOG_MEM_BANKS)
///////////////////////////////////////////////////////////////////////////////
// swap
@@ -208,6 +208,6 @@ namespace cv { namespace gpu { namespace cudev
return false;
}
}}} // namespace cv { namespace gpu { namespace cudev
}}} // namespace cv { namespace cuda { namespace cudev
#endif // __OPENCV_GPU_UTILITY_HPP__
#endif // __OPENCV_CUDA_UTILITY_HPP__

10
modules/core/include/opencv2/core/cuda/vec_distance.hpp
View File

@@ -40,14 +40,14 @@
//
//M*/
#ifndef __OPENCV_GPU_VEC_DISTANCE_HPP__
#define __OPENCV_GPU_VEC_DISTANCE_HPP__
#ifndef __OPENCV_CUDA_VEC_DISTANCE_HPP__
#define __OPENCV_CUDA_VEC_DISTANCE_HPP__
#include "reduce.hpp"
#include "functional.hpp"
#include "detail/vec_distance_detail.hpp"
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
template <typename T> struct L1Dist
{
@@ -219,6 +219,6 @@ namespace cv { namespace gpu { namespace cudev
U vec1Vals[MAX_LEN / THREAD_DIM];
};
}}} // namespace cv { namespace gpu { namespace cudev
}}} // namespace cv { namespace cuda { namespace cudev
#endif // __OPENCV_GPU_VEC_DISTANCE_HPP__
#endif // __OPENCV_CUDA_VEC_DISTANCE_HPP__

10
modules/core/include/opencv2/core/cuda/vec_math.hpp
View File

@@ -40,13 +40,13 @@
//
//M*/
#ifndef __OPENCV_GPU_VECMATH_HPP__
#define __OPENCV_GPU_VECMATH_HPP__
#ifndef __OPENCV_CUDA_VECMATH_HPP__
#define __OPENCV_CUDA_VECMATH_HPP__
#include "vec_traits.hpp"
#include "saturate_cast.hpp"
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
// saturate_cast
@@ -917,6 +917,6 @@ CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, double, double, double)
#undef CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC
}}} // namespace cv { namespace gpu { namespace device
}}} // namespace cv { namespace cuda { namespace device
#endif // __OPENCV_GPU_VECMATH_HPP__
#endif // __OPENCV_CUDA_VECMATH_HPP__

48
modules/core/include/opencv2/core/cuda/vec_traits.hpp
View File

@@ -40,12 +40,12 @@
//
//M*/
#ifndef __OPENCV_GPU_VEC_TRAITS_HPP__
#define __OPENCV_GPU_VEC_TRAITS_HPP__
#ifndef __OPENCV_CUDA_VEC_TRAITS_HPP__
#define __OPENCV_CUDA_VEC_TRAITS_HPP__
#include "common.hpp"
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
template<typename T, int N> struct TypeVec;
@@ -122,7 +122,7 @@ namespace cv { namespace gpu { namespace cudev
return val;
}
#define OPENCV_GPU_IMPLEMENT_TYPE_VEC(type) \
#define OPENCV_CUDA_IMPLEMENT_TYPE_VEC(type) \
template<> struct TypeVec<type, 1> { typedef type vec_type; }; \
template<> struct TypeVec<type ## 1, 1> { typedef type ## 1 vec_type; }; \
template<> struct TypeVec<type, 2> { typedef type ## 2 vec_type; }; \
@@ -134,16 +134,16 @@ namespace cv { namespace gpu { namespace cudev
template<> struct TypeVec<type, 8> { typedef type ## 8 vec_type; }; \
template<> struct TypeVec<type ## 8, 8> { typedef type ## 8 vec_type; };
OPENCV_GPU_IMPLEMENT_TYPE_VEC(uchar)
OPENCV_GPU_IMPLEMENT_TYPE_VEC(char)
OPENCV_GPU_IMPLEMENT_TYPE_VEC(ushort)
OPENCV_GPU_IMPLEMENT_TYPE_VEC(short)
OPENCV_GPU_IMPLEMENT_TYPE_VEC(int)
OPENCV_GPU_IMPLEMENT_TYPE_VEC(uint)
OPENCV_GPU_IMPLEMENT_TYPE_VEC(float)
OPENCV_GPU_IMPLEMENT_TYPE_VEC(double)
OPENCV_CUDA_IMPLEMENT_TYPE_VEC(uchar)
OPENCV_CUDA_IMPLEMENT_TYPE_VEC(char)
OPENCV_CUDA_IMPLEMENT_TYPE_VEC(ushort)
OPENCV_CUDA_IMPLEMENT_TYPE_VEC(short)
OPENCV_CUDA_IMPLEMENT_TYPE_VEC(int)
OPENCV_CUDA_IMPLEMENT_TYPE_VEC(uint)
OPENCV_CUDA_IMPLEMENT_TYPE_VEC(float)
OPENCV_CUDA_IMPLEMENT_TYPE_VEC(double)
#undef OPENCV_GPU_IMPLEMENT_TYPE_VEC
#undef OPENCV_CUDA_IMPLEMENT_TYPE_VEC
template<> struct TypeVec<schar, 1> { typedef schar vec_type; };
template<> struct TypeVec<schar, 2> { typedef char2 vec_type; };
@@ -159,7 +159,7 @@ namespace cv { namespace gpu { namespace cudev
template<typename T> struct VecTraits;
#define OPENCV_GPU_IMPLEMENT_VEC_TRAITS(type) \
#define OPENCV_CUDA_IMPLEMENT_VEC_TRAITS(type) \
template<> struct VecTraits<type> \
{ \
typedef type elem_type; \
@@ -209,15 +209,15 @@ namespace cv { namespace gpu { namespace cudev
static __device__ __host__ __forceinline__ type ## 8 make(const type* v) {return make_ ## type ## 8(v[0], v[1], v[2], v[3], v[4], v[5], v[6], v[7]);} \
};
OPENCV_GPU_IMPLEMENT_VEC_TRAITS(uchar)
OPENCV_GPU_IMPLEMENT_VEC_TRAITS(ushort)
OPENCV_GPU_IMPLEMENT_VEC_TRAITS(short)
OPENCV_GPU_IMPLEMENT_VEC_TRAITS(int)
OPENCV_GPU_IMPLEMENT_VEC_TRAITS(uint)
OPENCV_GPU_IMPLEMENT_VEC_TRAITS(float)
OPENCV_GPU_IMPLEMENT_VEC_TRAITS(double)
OPENCV_CUDA_IMPLEMENT_VEC_TRAITS(uchar)
OPENCV_CUDA_IMPLEMENT_VEC_TRAITS(ushort)
OPENCV_CUDA_IMPLEMENT_VEC_TRAITS(short)
OPENCV_CUDA_IMPLEMENT_VEC_TRAITS(int)
OPENCV_CUDA_IMPLEMENT_VEC_TRAITS(uint)
OPENCV_CUDA_IMPLEMENT_VEC_TRAITS(float)
OPENCV_CUDA_IMPLEMENT_VEC_TRAITS(double)
#undef OPENCV_GPU_IMPLEMENT_VEC_TRAITS
#undef OPENCV_CUDA_IMPLEMENT_VEC_TRAITS
template<> struct VecTraits<char>
{
@@ -275,6 +275,6 @@ namespace cv { namespace gpu { namespace cudev
static __device__ __host__ __forceinline__ char8 make(schar a0, schar a1, schar a2, schar a3, schar a4, schar a5, schar a6, schar a7) {return make_char8(a0, a1, a2, a3, a4, a5, a6, a7);}
static __device__ __host__ __forceinline__ char8 make(const schar* v) {return make_char8(v[0], v[1], v[2], v[3], v[4], v[5], v[6], v[7]);}
};
}}} // namespace cv { namespace gpu { namespace cudev
}}} // namespace cv { namespace cuda { namespace cudev
#endif // __OPENCV_GPU_VEC_TRAITS_HPP__
#endif // __OPENCV_CUDA_VEC_TRAITS_HPP__

10
modules/core/include/opencv2/core/cuda/warp.hpp
View File

@@ -40,10 +40,10 @@
//
//M*/
#ifndef __OPENCV_GPU_DEVICE_WARP_HPP__
#define __OPENCV_GPU_DEVICE_WARP_HPP__
#ifndef __OPENCV_CUDA_DEVICE_WARP_HPP__
#define __OPENCV_CUDA_DEVICE_WARP_HPP__
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
struct Warp
{
@@ -126,6 +126,6 @@ namespace cv { namespace gpu { namespace cudev
*t = value;
}
};
}}} // namespace cv { namespace gpu { namespace cudev
}}} // namespace cv { namespace cuda { namespace cudev
#endif /* __OPENCV_GPU_DEVICE_WARP_HPP__ */
#endif /* __OPENCV_CUDA_DEVICE_WARP_HPP__ */

10
modules/core/include/opencv2/core/cuda/warp_reduce.hpp
View File

@@ -40,10 +40,10 @@
//
//M*/
#ifndef OPENCV_GPU_WARP_REDUCE_HPP__
#define OPENCV_GPU_WARP_REDUCE_HPP__
#ifndef OPENCV_CUDA_WARP_REDUCE_HPP__
#define OPENCV_CUDA_WARP_REDUCE_HPP__
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
template <class T>
__device__ __forceinline__ T warp_reduce(volatile T *ptr , const unsigned int tid = threadIdx.x)
@@ -63,6 +63,6 @@ namespace cv { namespace gpu { namespace cudev
return ptr[tid - lane];
}
}}} // namespace cv { namespace gpu { namespace cudev {
}}} // namespace cv { namespace cuda { namespace cudev {
#endif /* OPENCV_GPU_WARP_REDUCE_HPP__ */
#endif /* OPENCV_CUDA_WARP_REDUCE_HPP__ */

8
modules/core/include/opencv2/core/cuda/warp_shuffle.hpp
View File

@@ -40,10 +40,10 @@
//
//M*/
#ifndef __OPENCV_GPU_WARP_SHUFFLE_HPP__
#define __OPENCV_GPU_WARP_SHUFFLE_HPP__
#ifndef __OPENCV_CUDA_WARP_SHUFFLE_HPP__
#define __OPENCV_CUDA_WARP_SHUFFLE_HPP__
namespace cv { namespace gpu { namespace cudev
namespace cv { namespace cuda { namespace device
{
template <typename T>
__device__ __forceinline__ T shfl(T val, int srcLane, int width = warpSize)
@@ -142,4 +142,4 @@ namespace cv { namespace gpu { namespace cudev
}
}}}
#endif // __OPENCV_GPU_WARP_SHUFFLE_HPP__
#endif // __OPENCV_CUDA_WARP_SHUFFLE_HPP__

12
modules/core/include/opencv2/core/gpu_stream_accessor.hpp → modules/core/include/opencv2/core/cuda_stream_accessor.hpp
View File

@@ -40,16 +40,16 @@
//
//M*/
#ifndef __OPENCV_CORE_GPU_STREAM_ACCESSOR_HPP__
#define __OPENCV_CORE_GPU_STREAM_ACCESSOR_HPP__
#ifndef __OPENCV_CORE_CUDA_STREAM_ACCESSOR_HPP__
#define __OPENCV_CORE_CUDA_STREAM_ACCESSOR_HPP__
#ifndef __cplusplus
# error gpu_stream_accessor.hpp header must be compiled as C++
# error cuda_stream_accessor.hpp header must be compiled as C++
#endif
// This is only header file that depends on Cuda. All other headers are independent.
// So if you use OpenCV binaries you do noot need to install Cuda Toolkit.
// But of you wanna use GPU by yourself, may get cuda stream instance using the class below.
// But of you wanna use CUDA by yourself, may get cuda stream instance using the class below.
// In this case you have to install Cuda Toolkit.
#include <cuda_runtime.h>
@@ -57,7 +57,7 @@
namespace cv
{
namespace gpu
namespace cuda
{
class Stream;
class Event;
@@ -74,4 +74,4 @@ namespace cv
}
}
#endif /* __OPENCV_CORE_GPU_STREAM_ACCESSOR_HPP__ */
#endif /* __OPENCV_CORE_CUDA_STREAM_ACCESSOR_HPP__ */

44
modules/core/include/opencv2/core/gpu_types.hpp → modules/core/include/opencv2/core/cuda_types.hpp
View File

@@ -40,22 +40,22 @@
//
//M*/
#ifndef __OPENCV_CORE_GPU_TYPES_HPP__
#define __OPENCV_CORE_GPU_TYPES_HPP__
#ifndef __OPENCV_CORE_CUDA_TYPES_HPP__
#define __OPENCV_CORE_CUDA_TYPES_HPP__
#ifndef __cplusplus
# error gpu_types.hpp header must be compiled as C++
# error cuda_types.hpp header must be compiled as C++
#endif
#ifdef __CUDACC__
#define __CV_GPU_HOST_DEVICE__ __host__ __device__ __forceinline__
#define __CV_CUDA_HOST_DEVICE__ __host__ __device__ __forceinline__
#else
#define __CV_GPU_HOST_DEVICE__
#define __CV_CUDA_HOST_DEVICE__
#endif
namespace cv
{
namespace gpu
namespace cuda
{
// Simple lightweight structures that encapsulates information about an image on device.
// It is intended to pass to nvcc-compiled code. GpuMat depends on headers that nvcc can't compile
@@ -69,41 +69,41 @@ namespace cv
T* data;
__CV_GPU_HOST_DEVICE__ DevPtr() : data(0) {}
__CV_GPU_HOST_DEVICE__ DevPtr(T* data_) : data(data_) {}
__CV_CUDA_HOST_DEVICE__ DevPtr() : data(0) {}
__CV_CUDA_HOST_DEVICE__ DevPtr(T* data_) : data(data_) {}
__CV_GPU_HOST_DEVICE__ size_t elemSize() const { return elem_size; }
__CV_GPU_HOST_DEVICE__ operator T*() { return data; }
__CV_GPU_HOST_DEVICE__ operator const T*() const { return data; }
__CV_CUDA_HOST_DEVICE__ size_t elemSize() const { return elem_size; }
__CV_CUDA_HOST_DEVICE__ operator T*() { return data; }
__CV_CUDA_HOST_DEVICE__ operator const T*() const { return data; }
};
template <typename T> struct PtrSz : public DevPtr<T>
{
__CV_GPU_HOST_DEVICE__ PtrSz() : size(0) {}
__CV_GPU_HOST_DEVICE__ PtrSz(T* data_, size_t size_) : DevPtr<T>(data_), size(size_) {}
__CV_CUDA_HOST_DEVICE__ PtrSz() : size(0) {}
__CV_CUDA_HOST_DEVICE__ PtrSz(T* data_, size_t size_) : DevPtr<T>(data_), size(size_) {}
size_t size;
};
template <typename T> struct PtrStep : public DevPtr<T>
{
__CV_GPU_HOST_DEVICE__ PtrStep() : step(0) {}
__CV_GPU_HOST_DEVICE__ PtrStep(T* data_, size_t step_) : DevPtr<T>(data_), step(step_) {}
__CV_CUDA_HOST_DEVICE__ PtrStep() : step(0) {}
__CV_CUDA_HOST_DEVICE__ PtrStep(T* data_, size_t step_) : DevPtr<T>(data_), step(step_) {}
//! stride between two consecutive rows in bytes. Step is stored always and everywhere in bytes!!!
size_t step;
__CV_GPU_HOST_DEVICE__ T* ptr(int y = 0) { return ( T*)( ( char*)DevPtr<T>::data + y * step); }
__CV_GPU_HOST_DEVICE__ const T* ptr(int y = 0) const { return (const T*)( (const char*)DevPtr<T>::data + y * step); }
__CV_CUDA_HOST_DEVICE__ T* ptr(int y = 0) { return ( T*)( ( char*)DevPtr<T>::data + y * step); }
__CV_CUDA_HOST_DEVICE__ const T* ptr(int y = 0) const { return (const T*)( (const char*)DevPtr<T>::data + y * step); }
__CV_GPU_HOST_DEVICE__ T& operator ()(int y, int x) { return ptr(y)[x]; }
__CV_GPU_HOST_DEVICE__ const T& operator ()(int y, int x) const { return ptr(y)[x]; }
__CV_CUDA_HOST_DEVICE__ T& operator ()(int y, int x) { return ptr(y)[x]; }
__CV_CUDA_HOST_DEVICE__ const T& operator ()(int y, int x) const { return ptr(y)[x]; }
};
template <typename T> struct PtrStepSz : public PtrStep<T>
{
__CV_GPU_HOST_DEVICE__ PtrStepSz() : cols(0), rows(0) {}
__CV_GPU_HOST_DEVICE__ PtrStepSz(int rows_, int cols_, T* data_, size_t step_)
__CV_CUDA_HOST_DEVICE__ PtrStepSz() : cols(0), rows(0) {}
__CV_CUDA_HOST_DEVICE__ PtrStepSz(int rows_, int cols_, T* data_, size_t step_)
: PtrStep<T>(data_, step_), cols(cols_), rows(rows_) {}
template <typename U>
@@ -123,4 +123,4 @@ namespace cv
}
}
#endif /* __OPENCV_CORE_GPU_TYPES_HPP__ */
#endif /* __OPENCV_CORE_CUDA_TYPES_HPP__ */

20
modules/core/include/opencv2/core/mat.hpp
View File

@@ -93,14 +93,14 @@ public:
template<typename _Tp> _InputArray(const _Tp* vec, int n);
template<typename _Tp, int m, int n> _InputArray(const Matx<_Tp, m, n>& matx);
_InputArray(const double& val);
_InputArray(const gpu::GpuMat& d_mat);
_InputArray(const cuda::GpuMat& d_mat);
_InputArray(const ogl::Buffer& buf);
_InputArray(const gpu::CudaMem& cuda_mem);
_InputArray(const cuda::CudaMem& cuda_mem);
template<typename _Tp> _InputArray(const cudev::GpuMat_<_Tp>& m);
virtual Mat getMat(int i=-1) const;
virtual void getMatVector(std::vector<Mat>& mv) const;
virtual gpu::GpuMat getGpuMat() const;
virtual cuda::GpuMat getGpuMat() const;
virtual ogl::Buffer getOGlBuffer() const;
virtual int kind() const;
@@ -142,9 +142,9 @@ public:
_OutputArray();
_OutputArray(Mat& m);
_OutputArray(std::vector<Mat>& vec);
_OutputArray(gpu::GpuMat& d_mat);
_OutputArray(cuda::GpuMat& d_mat);
_OutputArray(ogl::Buffer& buf);
_OutputArray(gpu::CudaMem& cuda_mem);
_OutputArray(cuda::CudaMem& cuda_mem);
template<typename _Tp> _OutputArray(cudev::GpuMat_<_Tp>& m);
template<typename _Tp> _OutputArray(std::vector<_Tp>& vec);
template<typename _Tp> _OutputArray(std::vector<std::vector<_Tp> >& vec);
@@ -155,9 +155,9 @@ public:
_OutputArray(const Mat& m);
_OutputArray(const std::vector<Mat>& vec);
_OutputArray(const gpu::GpuMat& d_mat);
_OutputArray(const cuda::GpuMat& d_mat);
_OutputArray(const ogl::Buffer& buf);
_OutputArray(const gpu::CudaMem& cuda_mem);
_OutputArray(const cuda::CudaMem& cuda_mem);
template<typename _Tp> _OutputArray(const cudev::GpuMat_<_Tp>& m);
template<typename _Tp> _OutputArray(const std::vector<_Tp>& vec);
template<typename _Tp> _OutputArray(const std::vector<std::vector<_Tp> >& vec);
@@ -170,9 +170,9 @@ public:
virtual bool fixedType() const;
virtual bool needed() const;
virtual Mat& getMatRef(int i=-1) const;
virtual gpu::GpuMat& getGpuMatRef() const;
virtual cuda::GpuMat& getGpuMatRef() const;
virtual ogl::Buffer& getOGlBufferRef() const;
virtual gpu::CudaMem& getCudaMemRef() const;
virtual cuda::CudaMem& getCudaMemRef() const;
virtual void create(Size sz, int type, int i=-1, bool allowTransposed=false, int fixedDepthMask=0) const;
virtual void create(int rows, int cols, int type, int i=-1, bool allowTransposed=false, int fixedDepthMask=0) const;
virtual void create(int dims, const int* size, int type, int i=-1, bool allowTransposed=false, int fixedDepthMask=0) const;
@@ -506,7 +506,7 @@ public:
//Mat(const void* img, bool copyData=false);
//! download data from GpuMat
explicit Mat(const gpu::GpuMat& m);
explicit Mat(const cuda::GpuMat& m);
//! destructor - calls release()
~Mat();

12
modules/core/include/opencv2/core/opengl.hpp
View File

@@ -99,12 +99,12 @@ public:
//! copy from host/device memory (blocking)
void copyFrom(InputArray arr, Target target = ARRAY_BUFFER, bool autoRelease = false);
//! copy from device memory (non blocking)
void copyFrom(InputArray arr, gpu::Stream& stream, Target target = ARRAY_BUFFER, bool autoRelease = false);
void copyFrom(InputArray arr, cuda::Stream& stream, Target target = ARRAY_BUFFER, bool autoRelease = false);
//! copy to host/device memory (blocking)
void copyTo(OutputArray arr) const;
//! copy to device memory (non blocking)
void copyTo(OutputArray arr, gpu::Stream& stream) const;
void copyTo(OutputArray arr, cuda::Stream& stream) const;
//! create copy of current buffer
Buffer clone(Target target = ARRAY_BUFFER, bool autoRelease = false) const;
@@ -120,12 +120,12 @@ public:
void unmapHost();
//! map to device memory (blocking)
gpu::GpuMat mapDevice();
cuda::GpuMat mapDevice();
void unmapDevice();
//! map to device memory (non blocking)
gpu::GpuMat mapDevice(gpu::Stream& stream);
void unmapDevice(gpu::Stream& stream);
cuda::GpuMat mapDevice(cuda::Stream& stream);
void unmapDevice(cuda::Stream& stream);
int rows() const;
int cols() const;
@@ -276,7 +276,7 @@ CV_EXPORTS void render(const Arrays& arr, InputArray indices, int mode = POINTS,
}} // namespace cv::ogl
namespace cv { namespace gpu {
namespace cv { namespace cuda {
//! set a CUDA device to use OpenGL interoperability
CV_EXPORTS void setGlDevice(int device = 0);

26
modules/core/include/opencv2/core/private.gpu.hpp → modules/core/include/opencv2/core/private.cuda.hpp
View File

@@ -41,8 +41,8 @@
//
//M*/
#ifndef __OPENCV_CORE_PRIVATE_GPU_HPP__
#define __OPENCV_CORE_PRIVATE_GPU_HPP__
#ifndef __OPENCV_CORE_PRIVATE_CUDA_HPP__
#define __OPENCV_CORE_PRIVATE_CUDA_HPP__
#ifndef __OPENCV_BUILD
# error this is a private header which should not be used from outside of the OpenCV library
@@ -53,13 +53,13 @@
#include "opencv2/core/cvdef.h"
#include "opencv2/core/base.hpp"
#include "opencv2/core/gpu.hpp"
#include "opencv2/core/cuda.hpp"
#ifdef HAVE_CUDA
# include <cuda.h>
# include <cuda_runtime.h>
# include <npp.h>
# include "opencv2/core/gpu_stream_accessor.hpp"
# include "opencv2/core/cuda_stream_accessor.hpp"
# include "opencv2/core/cuda/common.hpp"
# define NPP_VERSION (NPP_VERSION_MAJOR * 1000 + NPP_VERSION_MINOR * 100 + NPP_VERSION_BUILD)
@@ -71,24 +71,24 @@
# endif
# if defined(CUDA_ARCH_BIN_OR_PTX_10)
# error "OpenCV GPU module doesn't support NVIDIA compute capability 1.0"
# error "OpenCV CUDA module doesn't support NVIDIA compute capability 1.0"
# endif
#endif
namespace cv { namespace gpu {
namespace cv { namespace cuda {
CV_EXPORTS cv::String getNppErrorMessage(int code);
CV_EXPORTS cv::String getCudaDriverApiErrorMessage(int code);
}}
#ifndef HAVE_CUDA
static inline void throw_no_cuda() { CV_Error(cv::Error::GpuNotSupported, "The library is compiled without GPU support"); }
static inline void throw_no_cuda() { CV_Error(cv::Error::GpuNotSupported, "The library is compiled without CUDA support"); }
#else // HAVE_CUDA
static inline void throw_no_cuda() { CV_Error(cv::Error::StsNotImplemented, "The called functionality is disabled for current build or platform"); }
namespace cv { namespace gpu
namespace cv { namespace cuda
{
static inline void checkNppError(int code, const char* file, const int line, const char* func)
{
@@ -131,13 +131,13 @@ namespace cv { namespace gpu
}}
#if defined(__GNUC__)
#define nppSafeCall(expr) cv::gpu::checkNppError(expr, __FILE__, __LINE__, __func__)
#define cuSafeCall(expr) cv::gpu::checkCudaDriverApiError(expr, __FILE__, __LINE__, __func__)
#define nppSafeCall(expr) cv::cuda::checkNppError(expr, __FILE__, __LINE__, __func__)
#define cuSafeCall(expr) cv::cuda::checkCudaDriverApiError(expr, __FILE__, __LINE__, __func__)
#else /* defined(__CUDACC__) || defined(__MSVC__) */
#define nppSafeCall(expr) cv::gpu::checkNppError(expr, __FILE__, __LINE__, "")
#define cuSafeCall(expr) cv::gpu::checkCudaDriverApiError(expr, __FILE__, __LINE__, "")
#define nppSafeCall(expr) cv::cuda::checkNppError(expr, __FILE__, __LINE__, "")
#define cuSafeCall(expr) cv::cuda::checkCudaDriverApiError(expr, __FILE__, __LINE__, "")
#endif
#endif // HAVE_CUDA
#endif // __OPENCV_CORE_GPU_PRIVATE_HPP__
#endif // __OPENCV_CORE_CUDA_PRIVATE_HPP__