Merged the trunk r8589:8653 - all changes related to build warnings
This commit is contained in:
Andrey Kamaev
@@ -22,32 +22,31 @@ source_group("Device" FILES ${lib_device_hdrs})
|
||||
source_group("Device\\Detail" FILES ${lib_device_hdrs_detail})
|
||||
|
||||
if (HAVE_CUDA)
|
||||
file(GLOB_RECURSE ncv_srcs "src/nvidia/*.cpp")
|
||||
file(GLOB_RECURSE ncv_srcs "src/nvidia/*.cpp")
|
||||
file(GLOB_RECURSE ncv_cuda "src/nvidia/*.cu")
|
||||
file(GLOB_RECURSE ncv_hdrs "src/nvidia/*.hpp" "src/nvidia/*.h")
|
||||
set(ncv_files ${ncv_srcs} ${ncv_hdrs} ${ncv_cuda})
|
||||
|
||||
source_group("Src\\NVidia" FILES ${ncv_files})
|
||||
ocv_include_directories("src/nvidia" "src/nvidia/core" "src/nvidia/NPP_staging" ${CUDA_INCLUDE_DIRS})
|
||||
ocv_warnings_disable(CMAKE_CXX_FLAGS -Wundef -Wmissing-declarations /wd4211 /wd4201 /wd4100 /wd4505 /wd4408)
|
||||
|
||||
#set (CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} "-keep")
|
||||
#set (CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} "-Xcompiler;/EHsc-;")
|
||||
|
||||
|
||||
if(MSVC)
|
||||
if(NOT ENABLE_NOISY_WARNINGS)
|
||||
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /wd4211 /wd4201 /wd4100 /wd4505 /wd4408")
|
||||
|
||||
foreach(var CMAKE_CXX_FLAGS CMAKE_CXX_FLAGS_RELEASE CMAKE_CXX_FLAGS_DEBUG)
|
||||
string(REPLACE "/W4" "/W3" ${var} "${${var}}")
|
||||
endforeach()
|
||||
|
||||
|
||||
set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} -Xcompiler /wd4251)
|
||||
endif()
|
||||
endif()
|
||||
|
||||
OCV_CUDA_COMPILE(cuda_objs ${lib_cuda} ${ncv_cuda})
|
||||
ocv_cuda_compile(cuda_objs ${lib_cuda} ${ncv_cuda})
|
||||
#CUDA_BUILD_CLEAN_TARGET()
|
||||
|
||||
|
||||
set(cuda_link_libs ${CUDA_LIBRARIES} ${CUDA_npp_LIBRARY})
|
||||
else()
|
||||
set(lib_cuda "")
|
||||
@@ -60,9 +59,9 @@ ocv_set_module_sources(
|
||||
HEADERS ${lib_hdrs}
|
||||
SOURCES ${lib_int_hdrs} ${lib_cuda_hdrs} ${lib_device_hdrs} ${lib_device_hdrs_detail} ${lib_srcs} ${lib_cuda} ${ncv_files} ${cuda_objs}
|
||||
)
|
||||
|
||||
|
||||
ocv_create_module(${cuda_link_libs})
|
||||
|
||||
|
||||
if(HAVE_CUDA)
|
||||
if(HAVE_CUFFT)
|
||||
CUDA_ADD_CUFFT_TO_TARGET(${the_module})
|
||||
@@ -71,10 +70,10 @@ if(HAVE_CUDA)
|
||||
if(HAVE_CUBLAS)
|
||||
CUDA_ADD_CUBLAS_TO_TARGET(${the_module})
|
||||
endif()
|
||||
|
||||
|
||||
install(FILES src/nvidia/NPP_staging/NPP_staging.hpp src/nvidia/core/NCV.hpp
|
||||
DESTINATION ${OPENCV_INCLUDE_PREFIX}/opencv2/${name}
|
||||
COMPONENT main)
|
||||
DESTINATION ${OPENCV_INCLUDE_PREFIX}/opencv2/${name}
|
||||
COMPONENT main)
|
||||
endif()
|
||||
|
||||
ocv_add_precompiled_headers(${the_module})
|
||||
@@ -84,11 +83,11 @@ ocv_add_precompiled_headers(${the_module})
|
||||
################################################################################################################
|
||||
file(GLOB test_srcs "test/*.cpp")
|
||||
file(GLOB test_hdrs "test/*.hpp" "test/*.h")
|
||||
|
||||
set(nvidia "")
|
||||
if(HAVE_CUDA)
|
||||
file(GLOB nvidia "test/nvidia/*.cpp" "test/nvidia/*.hpp" "test/nvidia/*.h")
|
||||
file(GLOB nvidia "test/nvidia/*.cpp" "test/nvidia/*.hpp" "test/nvidia/*.h")
|
||||
set(nvidia FILES "Src\\\\\\\\NVidia" ${nvidia}) # 8 ugly backslashes :'(
|
||||
else()
|
||||
set(nvidia "")
|
||||
endif()
|
||||
|
||||
ocv_add_accuracy_tests(FILES "Include" ${test_hdrs}
|
||||
|
||||
@@ -1698,15 +1698,7 @@ class CV_EXPORTS GoodFeaturesToTrackDetector_GPU
|
||||
{
|
||||
public:
|
||||
explicit GoodFeaturesToTrackDetector_GPU(int maxCorners = 1000, double qualityLevel = 0.01, double minDistance = 0.0,
|
||||
int blockSize = 3, bool useHarrisDetector = false, double harrisK = 0.04)
|
||||
{
|
||||
this->maxCorners = maxCorners;
|
||||
this->qualityLevel = qualityLevel;
|
||||
this->minDistance = minDistance;
|
||||
this->blockSize = blockSize;
|
||||
this->useHarrisDetector = useHarrisDetector;
|
||||
this->harrisK = harrisK;
|
||||
}
|
||||
int blockSize = 3, bool useHarrisDetector = false, double harrisK = 0.04);
|
||||
|
||||
//! return 1 rows matrix with CV_32FC2 type
|
||||
void operator ()(const GpuMat& image, GpuMat& corners, const GpuMat& mask = GpuMat());
|
||||
@@ -1738,6 +1730,18 @@ private:
|
||||
GpuMat tmpCorners_;
|
||||
};
|
||||
|
||||
inline GoodFeaturesToTrackDetector_GPU::GoodFeaturesToTrackDetector_GPU(int maxCorners_, double qualityLevel_, double minDistance_,
|
||||
int blockSize_, bool useHarrisDetector_, double harrisK_)
|
||||
{
|
||||
maxCorners = maxCorners_;
|
||||
qualityLevel = qualityLevel_;
|
||||
minDistance = minDistance_;
|
||||
blockSize = blockSize_;
|
||||
useHarrisDetector = useHarrisDetector_;
|
||||
harrisK = harrisK_;
|
||||
}
|
||||
|
||||
|
||||
class CV_EXPORTS PyrLKOpticalFlow
|
||||
{
|
||||
public:
|
||||
|
||||
@@ -1,3 +1,7 @@
|
||||
#ifdef __GNUC__
|
||||
# pragma GCC diagnostic ignored "-Wmissing-declarations"
|
||||
#endif
|
||||
|
||||
#ifndef __OPENCV_PERF_PRECOMP_HPP__
|
||||
#define __OPENCV_PERF_PRECOMP_HPP__
|
||||
|
||||
@@ -11,7 +15,7 @@
|
||||
#include "opencv2/gpu/gpu.hpp"
|
||||
#include "perf_utility.hpp"
|
||||
|
||||
#if GTEST_CREATE_SHARED_LIBRARY
|
||||
#ifdef GTEST_CREATE_SHARED_LIBRARY
|
||||
#error no modules except ts should have GTEST_CREATE_SHARED_LIBRARY defined
|
||||
#endif
|
||||
|
||||
|
||||
@@ -52,157 +52,7 @@ void cv::gpu::swapChannels(GpuMat&, const int[], Stream&) { throw_nogpu(); }
|
||||
|
||||
#else /* !defined (HAVE_CUDA) */
|
||||
|
||||
namespace cv { namespace gpu { namespace device
|
||||
{
|
||||
#define OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name) \
|
||||
void name(const DevMem2Db& src, const DevMem2Db& dst, cudaStream_t stream);
|
||||
|
||||
#define OPENCV_GPU_DECLARE_CVTCOLOR_ALL(name) \
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name ## _8u) \
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name ## _16u) \
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name ## _32f)
|
||||
|
||||
#define OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(name) \
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name ## _8u) \
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name ## _32f) \
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name ## _full_8u) \
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name ## _full_32f)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_bgra)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_rgba)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr_to_bgr555)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr_to_bgr565)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(rgb_to_bgr555)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(rgb_to_bgr565)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgra_to_bgr555)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgra_to_bgr565)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(rgba_to_bgr555)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(rgba_to_bgr565)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr555_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr565_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr555_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr565_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr555_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr565_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr555_to_bgra)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr565_to_bgra)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(gray_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(gray_to_bgra)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(gray_to_bgr555)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(gray_to_bgr565)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr555_to_gray)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr565_to_gray)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgb_to_gray)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_gray)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgba_to_gray)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_gray)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgb_to_yuv)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgba_to_yuv)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgb_to_yuv4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgba_to_yuv4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_yuv)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_yuv)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_yuv4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_yuv4)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(yuv_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(yuv_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(yuv4_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(yuv4_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(yuv_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(yuv_to_bgra)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(yuv4_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(yuv4_to_bgra)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgb_to_YCrCb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgba_to_YCrCb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgb_to_YCrCb4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgba_to_YCrCb4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_YCrCb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_YCrCb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_YCrCb4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_YCrCb4)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(YCrCb_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(YCrCb_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(YCrCb4_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(YCrCb4_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(YCrCb_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(YCrCb_to_bgra)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(YCrCb4_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(YCrCb4_to_bgra)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgb_to_xyz)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgba_to_xyz)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgb_to_xyz4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgba_to_xyz4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_xyz)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_xyz)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_xyz4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_xyz4)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(xyz_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(xyz4_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(xyz_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(xyz4_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(xyz_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(xyz4_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(xyz_to_bgra)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(xyz4_to_bgra)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(rgb_to_hsv)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(rgba_to_hsv)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(rgb_to_hsv4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(rgba_to_hsv4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(bgr_to_hsv)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(bgra_to_hsv)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(bgr_to_hsv4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(bgra_to_hsv4)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hsv_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hsv_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hsv4_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hsv4_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hsv_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hsv_to_bgra)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hsv4_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hsv4_to_bgra)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(rgb_to_hls)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(rgba_to_hls)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(rgb_to_hls4)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(rgba_to_hls4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(bgr_to_hls)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(bgra_to_hls)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(bgr_to_hls4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(bgra_to_hls4)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hls_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hls_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hls4_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hls4_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hls_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hls_to_bgra)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hls4_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hls4_to_bgra)
|
||||
|
||||
#undef OPENCV_GPU_DECLARE_CVTCOLOR_ONE
|
||||
#undef OPENCV_GPU_DECLARE_CVTCOLOR_ALL
|
||||
#undef OPENCV_GPU_DECLARE_CVTCOLOR_8U32F
|
||||
}}}
|
||||
|
||||
#include <cvt_colot_internal.h>
|
||||
using namespace ::cv::gpu::device;
|
||||
|
||||
namespace
|
||||
|
||||
@@ -40,9 +40,10 @@
|
||||
//
|
||||
//M*/
|
||||
|
||||
#include "internal_shared.hpp"
|
||||
#include "opencv2/gpu/device/transform.hpp"
|
||||
#include "opencv2/gpu/device/color.hpp"
|
||||
#include <internal_shared.hpp>
|
||||
#include <opencv2/gpu/device/transform.hpp>
|
||||
#include <opencv2/gpu/device/color.hpp>
|
||||
#include <cvt_colot_internal.h>
|
||||
|
||||
namespace cv { namespace gpu { namespace device
|
||||
{
|
||||
|
||||
@@ -87,7 +87,9 @@ namespace cv { namespace gpu { namespace device
|
||||
__device__ __forceinline__ bool operator()(int y, int x) const
|
||||
{
|
||||
return true;
|
||||
}
|
||||
}
|
||||
__device__ __forceinline__ MaskTrue(){}
|
||||
__device__ __forceinline__ MaskTrue(const MaskTrue& mask_){}
|
||||
};
|
||||
|
||||
//////////////////////////////////////////////////////////////////////////////
|
||||
@@ -1795,6 +1797,9 @@ namespace cv { namespace gpu { namespace device
|
||||
return 0;
|
||||
}
|
||||
|
||||
__device__ __forceinline__ SumReductor(const SumReductor& other){}
|
||||
__device__ __forceinline__ SumReductor(){}
|
||||
|
||||
__device__ __forceinline__ S operator ()(volatile S a, volatile S b) const
|
||||
{
|
||||
return a + b;
|
||||
@@ -1813,6 +1818,9 @@ namespace cv { namespace gpu { namespace device
|
||||
return 0;
|
||||
}
|
||||
|
||||
__device__ __forceinline__ AvgReductor(const AvgReductor& other){}
|
||||
__device__ __forceinline__ AvgReductor(){}
|
||||
|
||||
__device__ __forceinline__ S operator ()(volatile S a, volatile S b) const
|
||||
{
|
||||
return a + b;
|
||||
@@ -1831,6 +1839,9 @@ namespace cv { namespace gpu { namespace device
|
||||
return numeric_limits<S>::max();
|
||||
}
|
||||
|
||||
__device__ __forceinline__ MinReductor(const MinReductor& other){}
|
||||
__device__ __forceinline__ MinReductor(){}
|
||||
|
||||
template <typename T> __device__ __forceinline__ T operator ()(volatile T a, volatile T b) const
|
||||
{
|
||||
return saturate_cast<T>(::min(a, b));
|
||||
@@ -1853,6 +1864,9 @@ namespace cv { namespace gpu { namespace device
|
||||
return numeric_limits<S>::min();
|
||||
}
|
||||
|
||||
__device__ __forceinline__ MaxReductor(const MaxReductor& other){}
|
||||
__device__ __forceinline__ MaxReductor(){}
|
||||
|
||||
template <typename T> __device__ __forceinline__ int operator ()(volatile T a, volatile T b) const
|
||||
{
|
||||
return ::max(a, b);
|
||||
|
||||
@@ -116,7 +116,7 @@ namespace cv { namespace gpu { namespace device
|
||||
|
||||
template <int N> __device__ float icvCalcHaarPatternSum(const float src[][5], int oldSize, int newSize, int y, int x)
|
||||
{
|
||||
#if __CUDA_ARCH__ >= 200
|
||||
#if __CUDA_ARCH__ && __CUDA_ARCH__ >= 200
|
||||
typedef double real_t;
|
||||
#else
|
||||
typedef float real_t;
|
||||
@@ -248,7 +248,7 @@ namespace cv { namespace gpu { namespace device
|
||||
template <typename Mask>
|
||||
__global__ void icvFindMaximaInLayer(const PtrStepf det, const PtrStepf trace, int4* maxPosBuffer, unsigned int* maxCounter)
|
||||
{
|
||||
#if __CUDA_ARCH__ >= 110
|
||||
#if __CUDA_ARCH__ && __CUDA_ARCH__ >= 110
|
||||
|
||||
extern __shared__ float N9[];
|
||||
|
||||
@@ -371,7 +371,7 @@ namespace cv { namespace gpu { namespace device
|
||||
float* featureX, float* featureY, int* featureLaplacian, int* featureOctave, float* featureSize, float* featureHessian,
|
||||
unsigned int* featureCounter)
|
||||
{
|
||||
#if __CUDA_ARCH__ >= 110
|
||||
#if __CUDA_ARCH__ && __CUDA_ARCH__ >= 110
|
||||
|
||||
const int4 maxPos = maxPosBuffer[blockIdx.x];
|
||||
|
||||
|
||||
197
modules/gpu/src/cvt_colot_internal.h
Normal file
197
modules/gpu/src/cvt_colot_internal.h
Normal file
@@ -0,0 +1,197 @@
|
||||
/*M///////////////////////////////////////////////////////////////////////////////////////
|
||||
//
|
||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
|
||||
//
|
||||
// By downloading, copying, installing or using the software you agree to this license.
|
||||
// If you do not agree to this license, do not download, install,
|
||||
// copy or use the software.
|
||||
//
|
||||
//
|
||||
// License Agreement
|
||||
// For Open Source Computer Vision Library
|
||||
//
|
||||
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
|
||||
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// Redistribution and use in source and binary forms, with or without modification,
|
||||
// are permitted provided that the following conditions are met:
|
||||
//
|
||||
// * Redistribution's of source code must retain the above copyright notice,
|
||||
// this list of conditions and the following disclaimer.
|
||||
//
|
||||
// * Redistribution's in binary form must reproduce the above copyright notice,
|
||||
// this list of conditions and the following disclaimer in the documentation
|
||||
// and/or other materials provided with the distribution.
|
||||
//
|
||||
// * The name of the copyright holders may not be used to endorse or promote products
|
||||
// derived from this software without specific prior written permission.
|
||||
//
|
||||
// This software is provided by the copyright holders and contributors "as is" and
|
||||
// any express or implied warranties, including, but not limited to, the implied
|
||||
// warranties of merchantability and fitness for a particular purpose are disclaimed.
|
||||
// In no event shall the Intel Corporation or contributors be liable for any direct,
|
||||
// indirect, incidental, special, exemplary, or consequential damages
|
||||
// (including, but not limited to, procurement of substitute goods or services;
|
||||
// loss of use, data, or profits; or business interruption) however caused
|
||||
// and on any theory of liability, whether in contract, strict liability,
|
||||
// or tort (including negligence or otherwise) arising in any way out of
|
||||
// the use of this software, even if advised of the possibility of such damage.
|
||||
//
|
||||
//M*/
|
||||
|
||||
#ifndef __cvt_color_internal_h__
|
||||
#define __cvt_color_internal_h__
|
||||
|
||||
namespace cv { namespace gpu { namespace device
|
||||
{
|
||||
#define OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name) \
|
||||
void name(const DevMem2Db& src, const DevMem2Db& dst, cudaStream_t stream);
|
||||
|
||||
#define OPENCV_GPU_DECLARE_CVTCOLOR_ALL(name) \
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name ## _8u) \
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name ## _16u) \
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name ## _32f)
|
||||
|
||||
#define OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(name) \
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name ## _8u) \
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name ## _32f) \
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name ## _full_8u) \
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name ## _full_32f)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_bgra)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_rgba)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr_to_bgr555)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr_to_bgr565)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(rgb_to_bgr555)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(rgb_to_bgr565)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgra_to_bgr555)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgra_to_bgr565)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(rgba_to_bgr555)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(rgba_to_bgr565)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr555_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr565_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr555_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr565_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr555_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr565_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr555_to_bgra)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr565_to_bgra)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(gray_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(gray_to_bgra)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(gray_to_bgr555)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(gray_to_bgr565)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr555_to_gray)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr565_to_gray)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgb_to_gray)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_gray)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgba_to_gray)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_gray)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgb_to_yuv)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgba_to_yuv)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgb_to_yuv4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgba_to_yuv4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_yuv)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_yuv)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_yuv4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_yuv4)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(yuv_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(yuv_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(yuv4_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(yuv4_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(yuv_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(yuv_to_bgra)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(yuv4_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(yuv4_to_bgra)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgb_to_YCrCb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgba_to_YCrCb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgb_to_YCrCb4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgba_to_YCrCb4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_YCrCb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_YCrCb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_YCrCb4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_YCrCb4)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(YCrCb_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(YCrCb_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(YCrCb4_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(YCrCb4_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(YCrCb_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(YCrCb_to_bgra)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(YCrCb4_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(YCrCb4_to_bgra)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgb_to_xyz)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgba_to_xyz)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgb_to_xyz4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgba_to_xyz4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_xyz)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_xyz)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_xyz4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_xyz4)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(xyz_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(xyz4_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(xyz_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(xyz4_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(xyz_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(xyz4_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(xyz_to_bgra)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_ALL(xyz4_to_bgra)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(rgb_to_hsv)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(rgba_to_hsv)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(rgb_to_hsv4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(rgba_to_hsv4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(bgr_to_hsv)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(bgra_to_hsv)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(bgr_to_hsv4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(bgra_to_hsv4)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hsv_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hsv_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hsv4_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hsv4_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hsv_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hsv_to_bgra)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hsv4_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hsv4_to_bgra)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(rgb_to_hls)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(rgba_to_hls)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(rgb_to_hls4)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(rgba_to_hls4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(bgr_to_hls)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(bgra_to_hls)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(bgr_to_hls4)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(bgra_to_hls4)
|
||||
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hls_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hls_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hls4_to_rgb)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hls4_to_rgba)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hls_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hls_to_bgra)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hls4_to_bgr)
|
||||
OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hls4_to_bgra)
|
||||
|
||||
#undef OPENCV_GPU_DECLARE_CVTCOLOR_ONE
|
||||
#undef OPENCV_GPU_DECLARE_CVTCOLOR_ALL
|
||||
#undef OPENCV_GPU_DECLARE_CVTCOLOR_8U32F
|
||||
}}}
|
||||
|
||||
#endif
|
||||
@@ -231,7 +231,7 @@ __device__ Ncv32u d_outMaskPosition;
|
||||
|
||||
__device__ void compactBlockWriteOutAnchorParallel(Ncv32u threadPassFlag, Ncv32u threadElem, Ncv32u *vectorOut)
|
||||
{
|
||||
#if __CUDA_ARCH__ >= 110
|
||||
#if __CUDA_ARCH__ && __CUDA_ARCH__ >= 110
|
||||
|
||||
__shared__ Ncv32u shmem[NUM_THREADS_ANCHORSPARALLEL * 2];
|
||||
__shared__ Ncv32u numPassed;
|
||||
@@ -587,7 +587,7 @@ __global__ void applyHaarClassifierClassifierParallel(Ncv32u *d_IImg, Ncv32u IIm
|
||||
}
|
||||
else
|
||||
{
|
||||
#if __CUDA_ARCH__ >= 110
|
||||
#if __CUDA_ARCH__ && __CUDA_ARCH__ >= 110
|
||||
if (bPass && !threadIdx.x)
|
||||
{
|
||||
Ncv32u outMaskOffset = atomicAdd(&d_outMaskPosition, 1);
|
||||
|
||||
@@ -142,7 +142,7 @@ struct NcvRect8u
|
||||
Ncv8u width;
|
||||
Ncv8u height;
|
||||
__host__ __device__ NcvRect8u() : x(0), y(0), width(0), height(0) {};
|
||||
__host__ __device__ NcvRect8u(Ncv8u x, Ncv8u y, Ncv8u width, Ncv8u height) : x(x), y(y), width(width), height(height) {}
|
||||
__host__ __device__ NcvRect8u(Ncv8u x_, Ncv8u y_, Ncv8u width_, Ncv8u height_) : x(x_), y(y_), width(width_), height(height_) {}
|
||||
};
|
||||
|
||||
|
||||
@@ -153,7 +153,8 @@ struct NcvRect32s
|
||||
Ncv32s width; ///< Rectangle width.
|
||||
Ncv32s height; ///< Rectangle height.
|
||||
__host__ __device__ NcvRect32s() : x(0), y(0), width(0), height(0) {};
|
||||
__host__ __device__ NcvRect32s(Ncv32s x, Ncv32s y, Ncv32s width, Ncv32s height) : x(x), y(y), width(width), height(height) {}
|
||||
__host__ __device__ NcvRect32s(Ncv32s x_, Ncv32s y_, Ncv32s width_, Ncv32s height_)
|
||||
: x(x_), y(y_), width(width_), height(height_) {}
|
||||
};
|
||||
|
||||
|
||||
@@ -164,7 +165,8 @@ struct NcvRect32u
|
||||
Ncv32u width; ///< Rectangle width.
|
||||
Ncv32u height; ///< Rectangle height.
|
||||
__host__ __device__ NcvRect32u() : x(0), y(0), width(0), height(0) {};
|
||||
__host__ __device__ NcvRect32u(Ncv32u x, Ncv32u y, Ncv32u width, Ncv32u height) : x(x), y(y), width(width), height(height) {}
|
||||
__host__ __device__ NcvRect32u(Ncv32u x_, Ncv32u y_, Ncv32u width_, Ncv32u height_)
|
||||
: x(x_), y(y_), width(width_), height(height_) {}
|
||||
};
|
||||
|
||||
|
||||
@@ -173,7 +175,7 @@ struct NcvSize32s
|
||||
Ncv32s width; ///< Rectangle width.
|
||||
Ncv32s height; ///< Rectangle height.
|
||||
__host__ __device__ NcvSize32s() : width(0), height(0) {};
|
||||
__host__ __device__ NcvSize32s(Ncv32s width, Ncv32s height) : width(width), height(height) {}
|
||||
__host__ __device__ NcvSize32s(Ncv32s width_, Ncv32s height_) : width(width_), height(height_) {}
|
||||
};
|
||||
|
||||
|
||||
@@ -182,7 +184,7 @@ struct NcvSize32u
|
||||
Ncv32u width; ///< Rectangle width.
|
||||
Ncv32u height; ///< Rectangle height.
|
||||
__host__ __device__ NcvSize32u() : width(0), height(0) {};
|
||||
__host__ __device__ NcvSize32u(Ncv32u width, Ncv32u height) : width(width), height(height) {}
|
||||
__host__ __device__ NcvSize32u(Ncv32u width_, Ncv32u height_) : width(width_), height(height_) {}
|
||||
__host__ __device__ bool operator == (const NcvSize32u &another) const {return this->width == another.width && this->height == another.height;}
|
||||
};
|
||||
|
||||
@@ -192,7 +194,7 @@ struct NcvPoint2D32s
|
||||
Ncv32s x; ///< Point X.
|
||||
Ncv32s y; ///< Point Y.
|
||||
__host__ __device__ NcvPoint2D32s() : x(0), y(0) {};
|
||||
__host__ __device__ NcvPoint2D32s(Ncv32s x, Ncv32s y) : x(x), y(y) {}
|
||||
__host__ __device__ NcvPoint2D32s(Ncv32s x_, Ncv32s y_) : x(x_), y(y_) {}
|
||||
};
|
||||
|
||||
|
||||
@@ -201,7 +203,7 @@ struct NcvPoint2D32u
|
||||
Ncv32u x; ///< Point X.
|
||||
Ncv32u y; ///< Point Y.
|
||||
__host__ __device__ NcvPoint2D32u() : x(0), y(0) {};
|
||||
__host__ __device__ NcvPoint2D32u(Ncv32u x, Ncv32u y) : x(x), y(y) {}
|
||||
__host__ __device__ NcvPoint2D32u(Ncv32u x_, Ncv32u y_) : x(x_), y(y_) {}
|
||||
};
|
||||
|
||||
|
||||
@@ -625,9 +627,9 @@ class NCVVectorAlloc : public NCVVector<T>
|
||||
|
||||
public:
|
||||
|
||||
NCVVectorAlloc(INCVMemAllocator &allocator, Ncv32u length)
|
||||
NCVVectorAlloc(INCVMemAllocator &allocator_, Ncv32u length)
|
||||
:
|
||||
allocator(allocator)
|
||||
allocator(allocator_)
|
||||
{
|
||||
NCVStatus ncvStat;
|
||||
|
||||
@@ -839,7 +841,7 @@ class NCVMatrixAlloc : public NCVMatrix<T>
|
||||
NCVMatrixAlloc& operator=(const NCVMatrixAlloc &);
|
||||
public:
|
||||
|
||||
NCVMatrixAlloc(INCVMemAllocator &allocator, Ncv32u width, Ncv32u height, Ncv32u pitch=0)
|
||||
NCVMatrixAlloc(INCVMemAllocator &allocator, Ncv32u width, Ncv32u height, Ncv32u _pitch=0)
|
||||
:
|
||||
allocator(allocator)
|
||||
{
|
||||
@@ -851,12 +853,12 @@ public:
|
||||
Ncv32u widthBytes = width * sizeof(T);
|
||||
Ncv32u pitchBytes = alignUp(widthBytes, allocator.alignment());
|
||||
|
||||
if (pitch != 0)
|
||||
if (_pitch != 0)
|
||||
{
|
||||
ncvAssertPrintReturn(pitch >= pitchBytes &&
|
||||
(pitch & (allocator.alignment() - 1)) == 0,
|
||||
ncvAssertPrintReturn(_pitch >= pitchBytes &&
|
||||
(_pitch & (allocator.alignment() - 1)) == 0,
|
||||
"NCVMatrixAlloc ctor:: incorrect pitch passed", );
|
||||
pitchBytes = pitch;
|
||||
pitchBytes = _pitch;
|
||||
}
|
||||
|
||||
Ncv32u requiredAllocSize = pitchBytes * height;
|
||||
@@ -1020,4 +1022,4 @@ NCV_EXPORTS NCVStatus ncvDrawRects_32u_device(Ncv32u *d_dst, Ncv32u dstStride, N
|
||||
|
||||
|
||||
|
||||
#endif // _ncv_hpp_
|
||||
#endif // _ncv_hpp_
|
||||
|
||||
@@ -41,7 +41,7 @@
|
||||
|
||||
#ifndef _ncvruntimetemplates_hpp_
|
||||
#define _ncvruntimetemplates_hpp_
|
||||
#if _MSC_VER >= 1200
|
||||
#if defined _MSC_VER &&_MSC_VER >= 1200
|
||||
#pragma warning( disable: 4800 )
|
||||
#endif
|
||||
|
||||
|
||||
@@ -47,7 +47,7 @@
|
||||
|
||||
namespace cv { namespace gpu { namespace device
|
||||
{
|
||||
#if __CUDA_ARCH__ >= 200
|
||||
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 200
|
||||
|
||||
// for Fermi memory space is detected automatically
|
||||
template <typename T> struct ForceGlob
|
||||
|
||||
@@ -63,6 +63,7 @@ namespace cv { namespace gpu { namespace device
|
||||
static __device__ __forceinline__ T max() { return numeric_limits<T>::max(); }
|
||||
static __device__ __forceinline__ T half() { return (T)(max()/2 + 1); }
|
||||
};
|
||||
|
||||
template<> struct ColorChannel<float>
|
||||
{
|
||||
typedef float worktype_f;
|
||||
@@ -73,14 +74,17 @@ namespace cv { namespace gpu { namespace device
|
||||
template <typename T> static __device__ __forceinline__ void setAlpha(typename TypeVec<T, 3>::vec_type& vec, T val)
|
||||
{
|
||||
}
|
||||
|
||||
template <typename T> static __device__ __forceinline__ void setAlpha(typename TypeVec<T, 4>::vec_type& vec, T val)
|
||||
{
|
||||
vec.w = val;
|
||||
}
|
||||
|
||||
template <typename T> static __device__ __forceinline__ T getAlpha(const typename TypeVec<T, 3>::vec_type& vec)
|
||||
{
|
||||
return ColorChannel<T>::max();
|
||||
}
|
||||
|
||||
template <typename T> static __device__ __forceinline__ T getAlpha(const typename TypeVec<T, 4>::vec_type& vec)
|
||||
{
|
||||
return vec.w;
|
||||
@@ -101,7 +105,8 @@ namespace cv { namespace gpu { namespace device
|
||||
|
||||
namespace color_detail
|
||||
{
|
||||
template <typename T, int scn, int dcn, int bidx> struct RGB2RGB : unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>
|
||||
template <typename T, int scn, int dcn, int bidx> struct RGB2RGB
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>
|
||||
{
|
||||
__device__ typename TypeVec<T, dcn>::vec_type operator()(const typename TypeVec<T, scn>::vec_type& src) const
|
||||
{
|
||||
@@ -114,6 +119,12 @@ namespace cv { namespace gpu { namespace device
|
||||
|
||||
return dst;
|
||||
}
|
||||
|
||||
__device__ __forceinline__ RGB2RGB()
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
|
||||
|
||||
__device__ __forceinline__ RGB2RGB(const RGB2RGB& other_)
|
||||
:unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
|
||||
};
|
||||
|
||||
template <> struct RGB2RGB<uchar, 4, 4, 2> : unary_function<uint, uint>
|
||||
@@ -129,6 +140,9 @@ namespace cv { namespace gpu { namespace device
|
||||
|
||||
return dst;
|
||||
}
|
||||
|
||||
__device__ __forceinline__ RGB2RGB():unary_function<uint, uint>(){}
|
||||
__device__ __forceinline__ RGB2RGB(const RGB2RGB& other_):unary_function<uint, uint>(){}
|
||||
};
|
||||
}
|
||||
|
||||
@@ -153,6 +167,7 @@ namespace cv { namespace gpu { namespace device
|
||||
{
|
||||
return (ushort)(((&src.x)[bidx] >> 3) | ((src.y & ~3) << 3) | (((&src.x)[bidx^2] & ~7) << 8));
|
||||
}
|
||||
|
||||
static __device__ __forceinline__ ushort cvt(uint src)
|
||||
{
|
||||
uint b = 0xffu & (src >> (bidx * 8));
|
||||
@@ -161,12 +176,14 @@ namespace cv { namespace gpu { namespace device
|
||||
return (ushort)((b >> 3) | ((g & ~3) << 3) | ((r & ~7) << 8));
|
||||
}
|
||||
};
|
||||
|
||||
template<int bidx> struct RGB2RGB5x5Converter<5, bidx>
|
||||
{
|
||||
static __device__ __forceinline__ ushort cvt(const uchar3& src)
|
||||
{
|
||||
return (ushort)(((&src.x)[bidx] >> 3) | ((src.y & ~7) << 2) | (((&src.x)[bidx^2] & ~7) << 7));
|
||||
}
|
||||
|
||||
static __device__ __forceinline__ ushort cvt(uint src)
|
||||
{
|
||||
uint b = 0xffu & (src >> (bidx * 8));
|
||||
@@ -178,19 +195,27 @@ namespace cv { namespace gpu { namespace device
|
||||
};
|
||||
|
||||
template<int scn, int bidx, int green_bits> struct RGB2RGB5x5;
|
||||
|
||||
template<int bidx, int green_bits> struct RGB2RGB5x5<3, bidx,green_bits> : unary_function<uchar3, ushort>
|
||||
{
|
||||
__device__ __forceinline__ ushort operator()(const uchar3& src) const
|
||||
{
|
||||
return RGB2RGB5x5Converter<green_bits, bidx>::cvt(src);
|
||||
}
|
||||
|
||||
__device__ __forceinline__ RGB2RGB5x5():unary_function<uchar3, ushort>(){}
|
||||
__device__ __forceinline__ RGB2RGB5x5(const RGB2RGB5x5& other_):unary_function<uchar3, ushort>(){}
|
||||
};
|
||||
|
||||
template<int bidx, int green_bits> struct RGB2RGB5x5<4, bidx,green_bits> : unary_function<uint, ushort>
|
||||
{
|
||||
__device__ __forceinline__ ushort operator()(uint src) const
|
||||
{
|
||||
return RGB2RGB5x5Converter<green_bits, bidx>::cvt(src);
|
||||
}
|
||||
|
||||
__device__ __forceinline__ RGB2RGB5x5():unary_function<uint, ushort>(){}
|
||||
__device__ __forceinline__ RGB2RGB5x5(const RGB2RGB5x5& other_):unary_function<uint, ushort>(){}
|
||||
};
|
||||
}
|
||||
|
||||
@@ -207,6 +232,7 @@ namespace cv { namespace gpu { namespace device
|
||||
namespace color_detail
|
||||
{
|
||||
template <int green_bits, int bidx> struct RGB5x52RGBConverter;
|
||||
|
||||
template <int bidx> struct RGB5x52RGBConverter<5, bidx>
|
||||
{
|
||||
static __device__ __forceinline__ void cvt(uint src, uchar3& dst)
|
||||
@@ -215,6 +241,7 @@ namespace cv { namespace gpu { namespace device
|
||||
dst.y = (src >> 2) & ~7;
|
||||
(&dst.x)[bidx ^ 2] = (src >> 7) & ~7;
|
||||
}
|
||||
|
||||
static __device__ __forceinline__ void cvt(uint src, uint& dst)
|
||||
{
|
||||
dst = 0;
|
||||
@@ -225,6 +252,7 @@ namespace cv { namespace gpu { namespace device
|
||||
dst |= ((src & 0x8000) * 0xffu) << 24;
|
||||
}
|
||||
};
|
||||
|
||||
template <int bidx> struct RGB5x52RGBConverter<6, bidx>
|
||||
{
|
||||
static __device__ __forceinline__ void cvt(uint src, uchar3& dst)
|
||||
@@ -233,6 +261,7 @@ namespace cv { namespace gpu { namespace device
|
||||
dst.y = (src >> 3) & ~3;
|
||||
(&dst.x)[bidx ^ 2] = (src >> 8) & ~7;
|
||||
}
|
||||
|
||||
static __device__ __forceinline__ void cvt(uint src, uint& dst)
|
||||
{
|
||||
dst = 0xffu << 24;
|
||||
@@ -244,6 +273,7 @@ namespace cv { namespace gpu { namespace device
|
||||
};
|
||||
|
||||
template <int dcn, int bidx, int green_bits> struct RGB5x52RGB;
|
||||
|
||||
template <int bidx, int green_bits> struct RGB5x52RGB<3, bidx, green_bits> : unary_function<ushort, uchar3>
|
||||
{
|
||||
__device__ __forceinline__ uchar3 operator()(ushort src) const
|
||||
@@ -252,7 +282,11 @@ namespace cv { namespace gpu { namespace device
|
||||
RGB5x52RGBConverter<green_bits, bidx>::cvt(src, dst);
|
||||
return dst;
|
||||
}
|
||||
__device__ __forceinline__ RGB5x52RGB():unary_function<ushort, uchar3>(){}
|
||||
__device__ __forceinline__ RGB5x52RGB(const RGB5x52RGB& other_):unary_function<ushort, uchar3>(){}
|
||||
|
||||
};
|
||||
|
||||
template <int bidx, int green_bits> struct RGB5x52RGB<4, bidx, green_bits> : unary_function<ushort, uint>
|
||||
{
|
||||
__device__ __forceinline__ uint operator()(ushort src) const
|
||||
@@ -261,6 +295,8 @@ namespace cv { namespace gpu { namespace device
|
||||
RGB5x52RGBConverter<green_bits, bidx>::cvt(src, dst);
|
||||
return dst;
|
||||
}
|
||||
__device__ __forceinline__ RGB5x52RGB():unary_function<ushort, uint>(){}
|
||||
__device__ __forceinline__ RGB5x52RGB(const RGB5x52RGB& other_):unary_function<ushort, uint>(){}
|
||||
};
|
||||
}
|
||||
|
||||
@@ -289,7 +325,11 @@ namespace cv { namespace gpu { namespace device
|
||||
|
||||
return dst;
|
||||
}
|
||||
__device__ __forceinline__ Gray2RGB():unary_function<T, typename TypeVec<T, dcn>::vec_type>(){}
|
||||
__device__ __forceinline__ Gray2RGB(const Gray2RGB& other_)
|
||||
: unary_function<T, typename TypeVec<T, dcn>::vec_type>(){}
|
||||
};
|
||||
|
||||
template <> struct Gray2RGB<uchar, 4> : unary_function<uchar, uint>
|
||||
{
|
||||
__device__ __forceinline__ uint operator()(uint src) const
|
||||
@@ -302,6 +342,8 @@ namespace cv { namespace gpu { namespace device
|
||||
|
||||
return dst;
|
||||
}
|
||||
__device__ __forceinline__ Gray2RGB():unary_function<uchar, uint>(){}
|
||||
__device__ __forceinline__ Gray2RGB(const Gray2RGB& other_):unary_function<uchar, uint>(){}
|
||||
};
|
||||
}
|
||||
|
||||
@@ -325,6 +367,7 @@ namespace cv { namespace gpu { namespace device
|
||||
return (ushort)((t >> 3) | ((t & ~3) << 3) | ((t & ~7) << 8));
|
||||
}
|
||||
};
|
||||
|
||||
template<> struct Gray2RGB5x5Converter<5>
|
||||
{
|
||||
static __device__ __forceinline__ ushort cvt(uint t)
|
||||
@@ -340,6 +383,9 @@ namespace cv { namespace gpu { namespace device
|
||||
{
|
||||
return Gray2RGB5x5Converter<green_bits>::cvt(src);
|
||||
}
|
||||
|
||||
__device__ __forceinline__ Gray2RGB5x5():unary_function<uchar, ushort>(){}
|
||||
__device__ __forceinline__ Gray2RGB5x5(const Gray2RGB5x5& other_):unary_function<uchar, ushort>(){}
|
||||
};
|
||||
}
|
||||
|
||||
@@ -365,6 +411,7 @@ namespace cv { namespace gpu { namespace device
|
||||
return (uchar)CV_DESCALE(((t << 3) & 0xf8) * B2Y + ((t >> 3) & 0xfc) * G2Y + ((t >> 8) & 0xf8) * R2Y, yuv_shift);
|
||||
}
|
||||
};
|
||||
|
||||
template <> struct RGB5x52GrayConverter<5>
|
||||
{
|
||||
static __device__ __forceinline__ uchar cvt(uint t)
|
||||
@@ -379,6 +426,8 @@ namespace cv { namespace gpu { namespace device
|
||||
{
|
||||
return RGB5x52GrayConverter<green_bits>::cvt(src);
|
||||
}
|
||||
__device__ __forceinline__ RGB5x52Gray() : unary_function<ushort, uchar>(){}
|
||||
__device__ __forceinline__ RGB5x52Gray(const RGB5x52Gray& other_) : unary_function<ushort, uchar>(){}
|
||||
};
|
||||
}
|
||||
|
||||
@@ -398,6 +447,7 @@ namespace cv { namespace gpu { namespace device
|
||||
{
|
||||
return (T)CV_DESCALE((unsigned)(src[bidx] * B2Y + src[1] * G2Y + src[bidx^2] * R2Y), yuv_shift);
|
||||
}
|
||||
|
||||
template <int bidx> static __device__ __forceinline__ uchar RGB2GrayConvert(uint src)
|
||||
{
|
||||
uint b = 0xffu & (src >> (bidx * 8));
|
||||
@@ -405,6 +455,7 @@ namespace cv { namespace gpu { namespace device
|
||||
uint r = 0xffu & (src >> ((bidx ^ 2) * 8));
|
||||
return CV_DESCALE((uint)(b * B2Y + g * G2Y + r * R2Y), yuv_shift);
|
||||
}
|
||||
|
||||
template <int bidx> static __device__ __forceinline__ float RGB2GrayConvert(const float* src)
|
||||
{
|
||||
return src[bidx] * 0.114f + src[1] * 0.587f + src[bidx^2] * 0.299f;
|
||||
@@ -416,13 +467,19 @@ namespace cv { namespace gpu { namespace device
|
||||
{
|
||||
return RGB2GrayConvert<bidx>(&src.x);
|
||||
}
|
||||
__device__ __forceinline__ RGB2Gray() : unary_function<typename TypeVec<T, scn>::vec_type, T>(){}
|
||||
__device__ __forceinline__ RGB2Gray(const RGB2Gray& other_)
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, T>(){}
|
||||
};
|
||||
|
||||
template <int bidx> struct RGB2Gray<uchar, 4, bidx> : unary_function<uint, uchar>
|
||||
{
|
||||
__device__ __forceinline__ uchar operator()(uint src) const
|
||||
{
|
||||
return RGB2GrayConvert<bidx>(src);
|
||||
}
|
||||
__device__ __forceinline__ RGB2Gray() : unary_function<uint, uchar>(){}
|
||||
__device__ __forceinline__ RGB2Gray(const RGB2Gray& other_) : unary_function<uint, uchar>(){}
|
||||
};
|
||||
}
|
||||
|
||||
@@ -463,7 +520,8 @@ namespace cv { namespace gpu { namespace device
|
||||
dst.z = (src[bidx] - dst.x) * c_RGB2YUVCoeffs_f[4] + ColorChannel<float>::half();
|
||||
}
|
||||
|
||||
template <typename T, int scn, int dcn, int bidx> struct RGB2YUV : unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>
|
||||
template <typename T, int scn, int dcn, int bidx> struct RGB2YUV
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>
|
||||
{
|
||||
__device__ __forceinline__ typename TypeVec<T, dcn>::vec_type operator ()(const typename TypeVec<T, scn>::vec_type& src) const
|
||||
{
|
||||
@@ -471,6 +529,10 @@ namespace cv { namespace gpu { namespace device
|
||||
RGB2YUVConvert<bidx>(&src.x, dst);
|
||||
return dst;
|
||||
}
|
||||
__device__ __forceinline__ RGB2YUV()
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
|
||||
__device__ __forceinline__ RGB2YUV(const RGB2YUV& other_)
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
|
||||
};
|
||||
}
|
||||
|
||||
@@ -492,13 +554,17 @@ namespace cv { namespace gpu { namespace device
|
||||
template <int bidx, typename T, typename D> static __device__ void YUV2RGBConvert(const T& src, D* dst)
|
||||
{
|
||||
const int b = src.x + CV_DESCALE((src.z - ColorChannel<D>::half()) * c_YUV2RGBCoeffs_i[3], yuv_shift);
|
||||
const int g = src.x + CV_DESCALE((src.z - ColorChannel<D>::half()) * c_YUV2RGBCoeffs_i[2] + (src.y - ColorChannel<D>::half()) * c_YUV2RGBCoeffs_i[1], yuv_shift);
|
||||
|
||||
const int g = src.x + CV_DESCALE((src.z - ColorChannel<D>::half()) * c_YUV2RGBCoeffs_i[2]
|
||||
+ (src.y - ColorChannel<D>::half()) * c_YUV2RGBCoeffs_i[1], yuv_shift);
|
||||
|
||||
const int r = src.x + CV_DESCALE((src.y - ColorChannel<D>::half()) * c_YUV2RGBCoeffs_i[0], yuv_shift);
|
||||
|
||||
dst[bidx] = saturate_cast<D>(b);
|
||||
dst[1] = saturate_cast<D>(g);
|
||||
dst[bidx^2] = saturate_cast<D>(r);
|
||||
}
|
||||
|
||||
template <int bidx> static __device__ uint YUV2RGBConvert(uint src)
|
||||
{
|
||||
const int x = 0xff & (src);
|
||||
@@ -506,7 +572,10 @@ namespace cv { namespace gpu { namespace device
|
||||
const int z = 0xff & (src >> 16);
|
||||
|
||||
const int b = x + CV_DESCALE((z - ColorChannel<uchar>::half()) * c_YUV2RGBCoeffs_i[3], yuv_shift);
|
||||
const int g = x + CV_DESCALE((z - ColorChannel<uchar>::half()) * c_YUV2RGBCoeffs_i[2] + (y - ColorChannel<uchar>::half()) * c_YUV2RGBCoeffs_i[1], yuv_shift);
|
||||
|
||||
const int g = x + CV_DESCALE((z - ColorChannel<uchar>::half()) * c_YUV2RGBCoeffs_i[2]
|
||||
+ (y - ColorChannel<uchar>::half()) * c_YUV2RGBCoeffs_i[1], yuv_shift);
|
||||
|
||||
const int r = x + CV_DESCALE((y - ColorChannel<uchar>::half()) * c_YUV2RGBCoeffs_i[0], yuv_shift);
|
||||
|
||||
uint dst = 0xffu << 24;
|
||||
@@ -517,14 +586,19 @@ namespace cv { namespace gpu { namespace device
|
||||
|
||||
return dst;
|
||||
}
|
||||
|
||||
template <int bidx, typename T> static __device__ __forceinline__ void YUV2RGBConvert(const T& src, float* dst)
|
||||
{
|
||||
dst[bidx] = src.x + (src.z - ColorChannel<float>::half()) * c_YUV2RGBCoeffs_f[3];
|
||||
dst[1] = src.x + (src.z - ColorChannel<float>::half()) * c_YUV2RGBCoeffs_f[2] + (src.y - ColorChannel<float>::half()) * c_YUV2RGBCoeffs_f[1];
|
||||
|
||||
dst[1] = src.x + (src.z - ColorChannel<float>::half()) * c_YUV2RGBCoeffs_f[2]
|
||||
+ (src.y - ColorChannel<float>::half()) * c_YUV2RGBCoeffs_f[1];
|
||||
|
||||
dst[bidx^2] = src.x + (src.y - ColorChannel<float>::half()) * c_YUV2RGBCoeffs_f[0];
|
||||
}
|
||||
|
||||
template <typename T, int scn, int dcn, int bidx> struct YUV2RGB : unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>
|
||||
template <typename T, int scn, int dcn, int bidx> struct YUV2RGB
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>
|
||||
{
|
||||
__device__ __forceinline__ typename TypeVec<T, dcn>::vec_type operator ()(const typename TypeVec<T, scn>::vec_type& src) const
|
||||
{
|
||||
@@ -535,13 +609,20 @@ namespace cv { namespace gpu { namespace device
|
||||
|
||||
return dst;
|
||||
}
|
||||
__device__ __forceinline__ YUV2RGB()
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
|
||||
__device__ __forceinline__ YUV2RGB(const YUV2RGB& other_)
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
|
||||
};
|
||||
|
||||
template <int bidx> struct YUV2RGB<uchar, 4, 4, bidx> : unary_function<uint, uint>
|
||||
{
|
||||
__device__ __forceinline__ uint operator ()(uint src) const
|
||||
{
|
||||
return YUV2RGBConvert<bidx>(src);
|
||||
}
|
||||
__device__ __forceinline__ YUV2RGB() : unary_function<uint, uint>(){}
|
||||
__device__ __forceinline__ YUV2RGB(const YUV2RGB& other_) : unary_function<uint, uint>(){}
|
||||
};
|
||||
}
|
||||
|
||||
@@ -574,6 +655,7 @@ namespace cv { namespace gpu { namespace device
|
||||
dst.y = saturate_cast<T>(Cr);
|
||||
dst.z = saturate_cast<T>(Cb);
|
||||
}
|
||||
|
||||
template <int bidx> static __device__ uint RGB2YCrCbConvert(uint src)
|
||||
{
|
||||
const int delta = ColorChannel<uchar>::half() * (1 << yuv_shift);
|
||||
@@ -590,6 +672,7 @@ namespace cv { namespace gpu { namespace device
|
||||
|
||||
return dst;
|
||||
}
|
||||
|
||||
template <int bidx, typename D> static __device__ __forceinline__ void RGB2YCrCbConvert(const float* src, D& dst)
|
||||
{
|
||||
dst.x = src[0] * c_RGB2YCrCbCoeffs_f[bidx^2] + src[1] * c_RGB2YCrCbCoeffs_f[1] + src[2] * c_RGB2YCrCbCoeffs_f[bidx];
|
||||
@@ -597,7 +680,8 @@ namespace cv { namespace gpu { namespace device
|
||||
dst.z = (src[bidx] - dst.x) * c_RGB2YCrCbCoeffs_f[4] + ColorChannel<float>::half();
|
||||
}
|
||||
|
||||
template <typename T, int scn, int dcn, int bidx> struct RGB2YCrCb : unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>
|
||||
template <typename T, int scn, int dcn, int bidx> struct RGB2YCrCb
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>
|
||||
{
|
||||
__device__ __forceinline__ typename TypeVec<T, dcn>::vec_type operator ()(const typename TypeVec<T, scn>::vec_type& src) const
|
||||
{
|
||||
@@ -605,13 +689,21 @@ namespace cv { namespace gpu { namespace device
|
||||
RGB2YCrCbConvert<bidx>(&src.x, dst);
|
||||
return dst;
|
||||
}
|
||||
__device__ __forceinline__ RGB2YCrCb()
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
|
||||
__device__ __forceinline__ RGB2YCrCb(const RGB2YCrCb& other_)
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
|
||||
};
|
||||
|
||||
template <int bidx> struct RGB2YCrCb<uchar, 4, 4, bidx> : unary_function<uint, uint>
|
||||
{
|
||||
__device__ __forceinline__ uint operator ()(uint src) const
|
||||
{
|
||||
return RGB2YCrCbConvert<bidx>(src);
|
||||
}
|
||||
|
||||
__device__ __forceinline__ RGB2YCrCb() : unary_function<uint, uint>(){}
|
||||
__device__ __forceinline__ RGB2YCrCb(const RGB2YCrCb& other_) : unary_function<uint, uint>(){}
|
||||
};
|
||||
}
|
||||
|
||||
@@ -640,6 +732,7 @@ namespace cv { namespace gpu { namespace device
|
||||
dst[1] = saturate_cast<D>(g);
|
||||
dst[bidx^2] = saturate_cast<D>(r);
|
||||
}
|
||||
|
||||
template <int bidx> static __device__ uint YCrCb2RGBConvert(uint src)
|
||||
{
|
||||
const int x = 0xff & (src);
|
||||
@@ -658,6 +751,7 @@ namespace cv { namespace gpu { namespace device
|
||||
|
||||
return dst;
|
||||
}
|
||||
|
||||
template <int bidx, typename T> __device__ __forceinline__ void YCrCb2RGBConvert(const T& src, float* dst)
|
||||
{
|
||||
dst[bidx] = src.x + (src.z - ColorChannel<float>::half()) * c_YCrCb2RGBCoeffs_f[3];
|
||||
@@ -665,7 +759,8 @@ namespace cv { namespace gpu { namespace device
|
||||
dst[bidx^2] = src.x + (src.y - ColorChannel<float>::half()) * c_YCrCb2RGBCoeffs_f[0];
|
||||
}
|
||||
|
||||
template <typename T, int scn, int dcn, int bidx> struct YCrCb2RGB : unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>
|
||||
template <typename T, int scn, int dcn, int bidx> struct YCrCb2RGB
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>
|
||||
{
|
||||
__device__ __forceinline__ typename TypeVec<T, dcn>::vec_type operator ()(const typename TypeVec<T, scn>::vec_type& src) const
|
||||
{
|
||||
@@ -676,13 +771,20 @@ namespace cv { namespace gpu { namespace device
|
||||
|
||||
return dst;
|
||||
}
|
||||
__device__ __forceinline__ YCrCb2RGB()
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
|
||||
__device__ __forceinline__ YCrCb2RGB(const YCrCb2RGB& other_)
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
|
||||
};
|
||||
|
||||
template <int bidx> struct YCrCb2RGB<uchar, 4, 4, bidx> : unary_function<uint, uint>
|
||||
{
|
||||
__device__ __forceinline__ uint operator ()(uint src) const
|
||||
{
|
||||
return YCrCb2RGBConvert<bidx>(src);
|
||||
}
|
||||
__device__ __forceinline__ YCrCb2RGB() : unary_function<uint, uint>(){}
|
||||
__device__ __forceinline__ YCrCb2RGB(const YCrCb2RGB& other_) : unary_function<uint, uint>(){}
|
||||
};
|
||||
}
|
||||
|
||||
@@ -709,6 +811,7 @@ namespace cv { namespace gpu { namespace device
|
||||
dst.y = saturate_cast<T>(CV_DESCALE(src[bidx^2] * c_RGB2XYZ_D65i[3] + src[1] * c_RGB2XYZ_D65i[4] + src[bidx] * c_RGB2XYZ_D65i[5], xyz_shift));
|
||||
dst.z = saturate_cast<T>(CV_DESCALE(src[bidx^2] * c_RGB2XYZ_D65i[6] + src[1] * c_RGB2XYZ_D65i[7] + src[bidx] * c_RGB2XYZ_D65i[8], xyz_shift));
|
||||
}
|
||||
|
||||
template <int bidx> static __device__ __forceinline__ uint RGB2XYZConvert(uint src)
|
||||
{
|
||||
const uint b = 0xffu & (src >> (bidx * 8));
|
||||
@@ -727,6 +830,7 @@ namespace cv { namespace gpu { namespace device
|
||||
|
||||
return dst;
|
||||
}
|
||||
|
||||
template <int bidx, typename D> static __device__ __forceinline__ void RGB2XYZConvert(const float* src, D& dst)
|
||||
{
|
||||
dst.x = src[bidx^2] * c_RGB2XYZ_D65f[0] + src[1] * c_RGB2XYZ_D65f[1] + src[bidx] * c_RGB2XYZ_D65f[2];
|
||||
@@ -734,7 +838,8 @@ namespace cv { namespace gpu { namespace device
|
||||
dst.z = src[bidx^2] * c_RGB2XYZ_D65f[6] + src[1] * c_RGB2XYZ_D65f[7] + src[bidx] * c_RGB2XYZ_D65f[8];
|
||||
}
|
||||
|
||||
template <typename T, int scn, int dcn, int bidx> struct RGB2XYZ : unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>
|
||||
template <typename T, int scn, int dcn, int bidx> struct RGB2XYZ
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>
|
||||
{
|
||||
__device__ __forceinline__ typename TypeVec<T, dcn>::vec_type operator()(const typename TypeVec<T, scn>::vec_type& src) const
|
||||
{
|
||||
@@ -744,13 +849,20 @@ namespace cv { namespace gpu { namespace device
|
||||
|
||||
return dst;
|
||||
}
|
||||
__device__ __forceinline__ RGB2XYZ()
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
|
||||
__device__ __forceinline__ RGB2XYZ(const RGB2XYZ& other_)
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
|
||||
};
|
||||
|
||||
template <int bidx> struct RGB2XYZ<uchar, 4, 4, bidx> : unary_function<uint, uint>
|
||||
{
|
||||
__device__ __forceinline__ uint operator()(uint src) const
|
||||
{
|
||||
return RGB2XYZConvert<bidx>(src);
|
||||
}
|
||||
__device__ __forceinline__ RGB2XYZ() : unary_function<uint, uint>(){}
|
||||
__device__ __forceinline__ RGB2XYZ(const RGB2XYZ& other_) : unary_function<uint, uint>(){}
|
||||
};
|
||||
}
|
||||
|
||||
@@ -775,6 +887,7 @@ namespace cv { namespace gpu { namespace device
|
||||
dst[1] = saturate_cast<D>(CV_DESCALE(src.x * c_XYZ2sRGB_D65i[3] + src.y * c_XYZ2sRGB_D65i[4] + src.z * c_XYZ2sRGB_D65i[5], xyz_shift));
|
||||
dst[bidx] = saturate_cast<D>(CV_DESCALE(src.x * c_XYZ2sRGB_D65i[6] + src.y * c_XYZ2sRGB_D65i[7] + src.z * c_XYZ2sRGB_D65i[8], xyz_shift));
|
||||
}
|
||||
|
||||
template <int bidx> static __device__ __forceinline__ uint XYZ2RGBConvert(uint src)
|
||||
{
|
||||
const int x = 0xff & src;
|
||||
@@ -793,6 +906,7 @@ namespace cv { namespace gpu { namespace device
|
||||
|
||||
return dst;
|
||||
}
|
||||
|
||||
template <int bidx, typename T> static __device__ __forceinline__ void XYZ2RGBConvert(const T& src, float* dst)
|
||||
{
|
||||
dst[bidx^2] = src.x * c_XYZ2sRGB_D65f[0] + src.y * c_XYZ2sRGB_D65f[1] + src.z * c_XYZ2sRGB_D65f[2];
|
||||
@@ -800,7 +914,8 @@ namespace cv { namespace gpu { namespace device
|
||||
dst[bidx] = src.x * c_XYZ2sRGB_D65f[6] + src.y * c_XYZ2sRGB_D65f[7] + src.z * c_XYZ2sRGB_D65f[8];
|
||||
}
|
||||
|
||||
template <typename T, int scn, int dcn, int bidx> struct XYZ2RGB : unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>
|
||||
template <typename T, int scn, int dcn, int bidx> struct XYZ2RGB
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>
|
||||
{
|
||||
__device__ __forceinline__ typename TypeVec<T, dcn>::vec_type operator()(const typename TypeVec<T, scn>::vec_type& src) const
|
||||
{
|
||||
@@ -811,13 +926,20 @@ namespace cv { namespace gpu { namespace device
|
||||
|
||||
return dst;
|
||||
}
|
||||
__device__ __forceinline__ XYZ2RGB()
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
|
||||
__device__ __forceinline__ XYZ2RGB(const XYZ2RGB& other_)
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
|
||||
};
|
||||
|
||||
template <int bidx> struct XYZ2RGB<uchar, 4, 4, bidx> : unary_function<uint, uint>
|
||||
{
|
||||
__device__ __forceinline__ uint operator()(uint src) const
|
||||
{
|
||||
return XYZ2RGBConvert<bidx>(src);
|
||||
}
|
||||
__device__ __forceinline__ XYZ2RGB() : unary_function<uint, uint>(){}
|
||||
__device__ __forceinline__ XYZ2RGB(const XYZ2RGB& other_) : unary_function<uint, uint>(){}
|
||||
};
|
||||
}
|
||||
|
||||
@@ -867,6 +989,7 @@ namespace cv { namespace gpu { namespace device
|
||||
dst.y = (uchar)s;
|
||||
dst.z = (uchar)v;
|
||||
}
|
||||
|
||||
template <int bidx, int hr> static __device__ uint RGB2HSVConvert(uint src)
|
||||
{
|
||||
const int hsv_shift = 12;
|
||||
@@ -902,6 +1025,7 @@ namespace cv { namespace gpu { namespace device
|
||||
|
||||
return dst;
|
||||
}
|
||||
|
||||
template <int bidx, int hr, typename D> static __device__ void RGB2HSVConvert(const float* src, D& dst)
|
||||
{
|
||||
const float hscale = hr * (1.f / 360.f);
|
||||
@@ -931,7 +1055,8 @@ namespace cv { namespace gpu { namespace device
|
||||
dst.z = v;
|
||||
}
|
||||
|
||||
template <typename T, int scn, int dcn, int bidx, int hr> struct RGB2HSV : unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>
|
||||
template <typename T, int scn, int dcn, int bidx, int hr> struct RGB2HSV
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>
|
||||
{
|
||||
__device__ __forceinline__ typename TypeVec<T, dcn>::vec_type operator()(const typename TypeVec<T, scn>::vec_type& src) const
|
||||
{
|
||||
@@ -941,13 +1066,20 @@ namespace cv { namespace gpu { namespace device
|
||||
|
||||
return dst;
|
||||
}
|
||||
__device__ __forceinline__ RGB2HSV()
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
|
||||
__device__ __forceinline__ RGB2HSV(const RGB2HSV& other_)
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
|
||||
};
|
||||
|
||||
template <int bidx, int hr> struct RGB2HSV<uchar, 4, 4, bidx, hr> : unary_function<uint, uint>
|
||||
{
|
||||
__device__ __forceinline__ uint operator()(uint src) const
|
||||
{
|
||||
return RGB2HSVConvert<bidx, hr>(src);
|
||||
}
|
||||
__device__ __forceinline__ RGB2HSV():unary_function<uint, uint>(){}
|
||||
__device__ __forceinline__ RGB2HSV(const RGB2HSV& other_):unary_function<uint, uint>(){}
|
||||
};
|
||||
}
|
||||
|
||||
@@ -1023,6 +1155,7 @@ namespace cv { namespace gpu { namespace device
|
||||
dst[1] = g;
|
||||
dst[bidx^2] = r;
|
||||
}
|
||||
|
||||
template <int bidx, int HR, typename T> static __device__ void HSV2RGBConvert(const T& src, uchar* dst)
|
||||
{
|
||||
float3 buf;
|
||||
@@ -1037,6 +1170,7 @@ namespace cv { namespace gpu { namespace device
|
||||
dst[1] = saturate_cast<uchar>(buf.y * 255.f);
|
||||
dst[2] = saturate_cast<uchar>(buf.z * 255.f);
|
||||
}
|
||||
|
||||
template <int bidx, int hr> static __device__ uint HSV2RGBConvert(uint src)
|
||||
{
|
||||
float3 buf;
|
||||
@@ -1056,7 +1190,8 @@ namespace cv { namespace gpu { namespace device
|
||||
return dst;
|
||||
}
|
||||
|
||||
template <typename T, int scn, int dcn, int bidx, int hr> struct HSV2RGB : unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>
|
||||
template <typename T, int scn, int dcn, int bidx, int hr> struct HSV2RGB
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>
|
||||
{
|
||||
__device__ __forceinline__ typename TypeVec<T, dcn>::vec_type operator()(const typename TypeVec<T, scn>::vec_type& src) const
|
||||
{
|
||||
@@ -1067,13 +1202,20 @@ namespace cv { namespace gpu { namespace device
|
||||
|
||||
return dst;
|
||||
}
|
||||
__device__ __forceinline__ HSV2RGB()
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
|
||||
__device__ __forceinline__ HSV2RGB(const HSV2RGB& other_)
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
|
||||
};
|
||||
|
||||
template <int bidx, int hr> struct HSV2RGB<uchar, 4, 4, bidx, hr> : unary_function<uint, uint>
|
||||
{
|
||||
__device__ __forceinline__ uint operator()(uint src) const
|
||||
{
|
||||
return HSV2RGBConvert<bidx, hr>(src);
|
||||
}
|
||||
__device__ __forceinline__ HSV2RGB():unary_function<uint, uint>(){}
|
||||
__device__ __forceinline__ HSV2RGB(const HSV2RGB& other_):unary_function<uint, uint>(){}
|
||||
};
|
||||
}
|
||||
|
||||
@@ -1149,6 +1291,7 @@ namespace cv { namespace gpu { namespace device
|
||||
dst.y = l;
|
||||
dst.z = s;
|
||||
}
|
||||
|
||||
template <int bidx, int hr, typename D> static __device__ void RGB2HLSConvert(const uchar* src, D& dst)
|
||||
{
|
||||
float3 buf;
|
||||
@@ -1163,6 +1306,7 @@ namespace cv { namespace gpu { namespace device
|
||||
dst.y = saturate_cast<uchar>(buf.y*255.f);
|
||||
dst.z = saturate_cast<uchar>(buf.z*255.f);
|
||||
}
|
||||
|
||||
template <int bidx, int hr> static __device__ uint RGB2HLSConvert(uint src)
|
||||
{
|
||||
float3 buf;
|
||||
@@ -1182,7 +1326,8 @@ namespace cv { namespace gpu { namespace device
|
||||
return dst;
|
||||
}
|
||||
|
||||
template <typename T, int scn, int dcn, int bidx, int hr> struct RGB2HLS : unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>
|
||||
template <typename T, int scn, int dcn, int bidx, int hr> struct RGB2HLS
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>
|
||||
{
|
||||
__device__ __forceinline__ typename TypeVec<T, dcn>::vec_type operator()(const typename TypeVec<T, scn>::vec_type& src) const
|
||||
{
|
||||
@@ -1192,13 +1337,20 @@ namespace cv { namespace gpu { namespace device
|
||||
|
||||
return dst;
|
||||
}
|
||||
__device__ __forceinline__ RGB2HLS()
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
|
||||
__device__ __forceinline__ RGB2HLS(const RGB2HLS& other_)
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
|
||||
};
|
||||
|
||||
template <int bidx, int hr> struct RGB2HLS<uchar, 4, 4, bidx, hr> : unary_function<uint, uint>
|
||||
{
|
||||
__device__ __forceinline__ uint operator()(uint src) const
|
||||
{
|
||||
return RGB2HLSConvert<bidx, hr>(src);
|
||||
}
|
||||
__device__ __forceinline__ RGB2HLS() : unary_function<uint, uint>(){}
|
||||
__device__ __forceinline__ RGB2HLS(const RGB2HLS& other_) : unary_function<uint, uint>(){}
|
||||
};
|
||||
}
|
||||
|
||||
@@ -1280,6 +1432,7 @@ namespace cv { namespace gpu { namespace device
|
||||
dst[1] = g;
|
||||
dst[bidx^2] = r;
|
||||
}
|
||||
|
||||
template <int bidx, int hr, typename T> static __device__ void HLS2RGBConvert(const T& src, uchar* dst)
|
||||
{
|
||||
float3 buf;
|
||||
@@ -1294,6 +1447,7 @@ namespace cv { namespace gpu { namespace device
|
||||
dst[1] = saturate_cast<uchar>(buf.y * 255.f);
|
||||
dst[2] = saturate_cast<uchar>(buf.z * 255.f);
|
||||
}
|
||||
|
||||
template <int bidx, int hr> static __device__ uint HLS2RGBConvert(uint src)
|
||||
{
|
||||
float3 buf;
|
||||
@@ -1313,7 +1467,8 @@ namespace cv { namespace gpu { namespace device
|
||||
return dst;
|
||||
}
|
||||
|
||||
template <typename T, int scn, int dcn, int bidx, int hr> struct HLS2RGB : unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>
|
||||
template <typename T, int scn, int dcn, int bidx, int hr> struct HLS2RGB
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>
|
||||
{
|
||||
__device__ __forceinline__ typename TypeVec<T, dcn>::vec_type operator()(const typename TypeVec<T, scn>::vec_type& src) const
|
||||
{
|
||||
@@ -1324,13 +1479,20 @@ namespace cv { namespace gpu { namespace device
|
||||
|
||||
return dst;
|
||||
}
|
||||
__device__ __forceinline__ HLS2RGB()
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
|
||||
__device__ __forceinline__ HLS2RGB(const HLS2RGB& other_)
|
||||
: unary_function<typename TypeVec<T, scn>::vec_type, typename TypeVec<T, dcn>::vec_type>(){}
|
||||
};
|
||||
|
||||
template <int bidx, int hr> struct HLS2RGB<uchar, 4, 4, bidx, hr> : unary_function<uint, uint>
|
||||
{
|
||||
__device__ __forceinline__ uint operator()(uint src) const
|
||||
{
|
||||
return HLS2RGBConvert<bidx, hr>(src);
|
||||
}
|
||||
__device__ __forceinline__ HLS2RGB() : unary_function<uint, uint>(){}
|
||||
__device__ __forceinline__ HLS2RGB(const HLS2RGB& other_) : unary_function<uint, uint>(){}
|
||||
};
|
||||
}
|
||||
|
||||
|
||||
@@ -56,158 +56,224 @@ namespace cv { namespace gpu { namespace device
|
||||
using thrust::binary_function;
|
||||
|
||||
// Arithmetic Operations
|
||||
|
||||
template <typename T> struct plus : binary_function<T, T, T>
|
||||
{
|
||||
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
|
||||
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,
|
||||
typename TypeTraits<T>::ParameterType b) const
|
||||
{
|
||||
return a + b;
|
||||
}
|
||||
__device__ __forceinline__ plus(const plus& other):binary_function<T,T,T>(){}
|
||||
__device__ __forceinline__ plus():binary_function<T,T,T>(){}
|
||||
};
|
||||
|
||||
template <typename T> struct minus : binary_function<T, T, T>
|
||||
{
|
||||
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
|
||||
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,
|
||||
typename TypeTraits<T>::ParameterType b) const
|
||||
{
|
||||
return a - b;
|
||||
}
|
||||
__device__ __forceinline__ minus(const minus& other):binary_function<T,T,T>(){}
|
||||
__device__ __forceinline__ minus():binary_function<T,T,T>(){}
|
||||
};
|
||||
|
||||
template <typename T> struct multiplies : binary_function<T, T, T>
|
||||
{
|
||||
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
|
||||
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,
|
||||
typename TypeTraits<T>::ParameterType b) const
|
||||
{
|
||||
return a * b;
|
||||
}
|
||||
__device__ __forceinline__ multiplies(const multiplies& other):binary_function<T,T,T>(){}
|
||||
__device__ __forceinline__ multiplies():binary_function<T,T,T>(){}
|
||||
};
|
||||
|
||||
template <typename T> struct divides : binary_function<T, T, T>
|
||||
{
|
||||
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
|
||||
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,
|
||||
typename TypeTraits<T>::ParameterType b) const
|
||||
{
|
||||
return a / b;
|
||||
}
|
||||
__device__ __forceinline__ divides(const divides& other):binary_function<T,T,T>(){}
|
||||
__device__ __forceinline__ divides():binary_function<T,T,T>(){}
|
||||
};
|
||||
|
||||
template <typename T> struct modulus : binary_function<T, T, T>
|
||||
{
|
||||
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
|
||||
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,
|
||||
typename TypeTraits<T>::ParameterType b) const
|
||||
{
|
||||
return a % b;
|
||||
}
|
||||
__device__ __forceinline__ modulus(const modulus& other):binary_function<T,T,T>(){}
|
||||
__device__ __forceinline__ modulus():binary_function<T,T,T>(){}
|
||||
};
|
||||
|
||||
template <typename T> struct negate : unary_function<T, T>
|
||||
{
|
||||
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a) const
|
||||
{
|
||||
return -a;
|
||||
}
|
||||
__device__ __forceinline__ negate(const negate& other):unary_function<T,T>(){}
|
||||
__device__ __forceinline__ negate():unary_function<T,T>(){}
|
||||
};
|
||||
|
||||
// Comparison Operations
|
||||
|
||||
template <typename T> struct equal_to : binary_function<T, T, bool>
|
||||
{
|
||||
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
|
||||
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,
|
||||
typename TypeTraits<T>::ParameterType b) const
|
||||
{
|
||||
return a == b;
|
||||
}
|
||||
__device__ __forceinline__ equal_to(const equal_to& other):binary_function<T,T,bool>(){}
|
||||
__device__ __forceinline__ equal_to():binary_function<T,T,bool>(){}
|
||||
};
|
||||
|
||||
template <typename T> struct not_equal_to : binary_function<T, T, bool>
|
||||
{
|
||||
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
|
||||
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,
|
||||
typename TypeTraits<T>::ParameterType b) const
|
||||
{
|
||||
return a != b;
|
||||
}
|
||||
__device__ __forceinline__ not_equal_to(const not_equal_to& other):binary_function<T,T,bool>(){}
|
||||
__device__ __forceinline__ not_equal_to():binary_function<T,T,bool>(){}
|
||||
};
|
||||
|
||||
template <typename T> struct greater : binary_function<T, T, bool>
|
||||
{
|
||||
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
|
||||
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,
|
||||
typename TypeTraits<T>::ParameterType b) const
|
||||
{
|
||||
return a > b;
|
||||
}
|
||||
__device__ __forceinline__ greater(const greater& other):binary_function<T,T,bool>(){}
|
||||
__device__ __forceinline__ greater():binary_function<T,T,bool>(){}
|
||||
};
|
||||
|
||||
template <typename T> struct less : binary_function<T, T, bool>
|
||||
{
|
||||
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
|
||||
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,
|
||||
typename TypeTraits<T>::ParameterType b) const
|
||||
{
|
||||
return a < b;
|
||||
}
|
||||
__device__ __forceinline__ less(const less& other):binary_function<T,T,bool>(){}
|
||||
__device__ __forceinline__ less():binary_function<T,T,bool>(){}
|
||||
};
|
||||
|
||||
template <typename T> struct greater_equal : binary_function<T, T, bool>
|
||||
{
|
||||
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
|
||||
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,
|
||||
typename TypeTraits<T>::ParameterType b) const
|
||||
{
|
||||
return a >= b;
|
||||
}
|
||||
__device__ __forceinline__ greater_equal(const greater_equal& other):binary_function<T,T,bool>(){}
|
||||
__device__ __forceinline__ greater_equal():binary_function<T,T,bool>(){}
|
||||
};
|
||||
|
||||
template <typename T> struct less_equal : binary_function<T, T, bool>
|
||||
{
|
||||
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
|
||||
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,
|
||||
typename TypeTraits<T>::ParameterType b) const
|
||||
{
|
||||
return a <= b;
|
||||
}
|
||||
__device__ __forceinline__ less_equal(const less_equal& other):binary_function<T,T,bool>(){}
|
||||
__device__ __forceinline__ less_equal():binary_function<T,T,bool>(){}
|
||||
};
|
||||
|
||||
// Logical Operations
|
||||
|
||||
template <typename T> struct logical_and : binary_function<T, T, bool>
|
||||
{
|
||||
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
|
||||
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,
|
||||
typename TypeTraits<T>::ParameterType b) const
|
||||
{
|
||||
return a && b;
|
||||
}
|
||||
__device__ __forceinline__ logical_and(const logical_and& other):binary_function<T,T,bool>(){}
|
||||
__device__ __forceinline__ logical_and():binary_function<T,T,bool>(){}
|
||||
};
|
||||
|
||||
template <typename T> struct logical_or : binary_function<T, T, bool>
|
||||
{
|
||||
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
|
||||
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,
|
||||
typename TypeTraits<T>::ParameterType b) const
|
||||
{
|
||||
return a || b;
|
||||
}
|
||||
__device__ __forceinline__ logical_or(const logical_or& other):binary_function<T,T,bool>(){}
|
||||
__device__ __forceinline__ logical_or():binary_function<T,T,bool>(){}
|
||||
};
|
||||
|
||||
template <typename T> struct logical_not : unary_function<T, bool>
|
||||
{
|
||||
__device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a) const
|
||||
{
|
||||
return !a;
|
||||
}
|
||||
__device__ __forceinline__ logical_not(const logical_not& other):unary_function<T,bool>(){}
|
||||
__device__ __forceinline__ logical_not():unary_function<T,bool>(){}
|
||||
};
|
||||
|
||||
// Bitwise Operations
|
||||
|
||||
template <typename T> struct bit_and : binary_function<T, T, T>
|
||||
{
|
||||
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
|
||||
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,
|
||||
typename TypeTraits<T>::ParameterType b) const
|
||||
{
|
||||
return a & b;
|
||||
}
|
||||
__device__ __forceinline__ bit_and(const bit_and& other):binary_function<T,T,T>(){}
|
||||
__device__ __forceinline__ bit_and():binary_function<T,T,T>(){}
|
||||
};
|
||||
|
||||
template <typename T> struct bit_or : binary_function<T, T, T>
|
||||
{
|
||||
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
|
||||
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,
|
||||
typename TypeTraits<T>::ParameterType b) const
|
||||
{
|
||||
return a | b;
|
||||
}
|
||||
__device__ __forceinline__ bit_or(const bit_or& other):binary_function<T,T,T>(){}
|
||||
__device__ __forceinline__ bit_or():binary_function<T,T,T>(){}
|
||||
};
|
||||
|
||||
template <typename T> struct bit_xor : binary_function<T, T, T>
|
||||
{
|
||||
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, typename TypeTraits<T>::ParameterType b) const
|
||||
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,
|
||||
typename TypeTraits<T>::ParameterType b) const
|
||||
{
|
||||
return a ^ b;
|
||||
}
|
||||
__device__ __forceinline__ bit_xor(const bit_xor& other):binary_function<T,T,T>(){}
|
||||
__device__ __forceinline__ bit_xor():binary_function<T,T,T>(){}
|
||||
};
|
||||
|
||||
template <typename T> struct bit_not : unary_function<T, T>
|
||||
{
|
||||
__device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType v) const
|
||||
{
|
||||
return ~v;
|
||||
}
|
||||
__device__ __forceinline__ bit_not(const bit_not& other):unary_function<T,T>(){}
|
||||
__device__ __forceinline__ bit_not():unary_function<T,T>(){}
|
||||
};
|
||||
|
||||
// Generalized Identity Operations
|
||||
|
||||
template <typename T> struct identity : unary_function<T, T>
|
||||
{
|
||||
__device__ __forceinline__ typename TypeTraits<T>::ParameterType operator()(typename TypeTraits<T>::ParameterType x) const
|
||||
{
|
||||
return x;
|
||||
}
|
||||
__device__ __forceinline__ identity(const identity& other):unary_function<T,T>(){}
|
||||
__device__ __forceinline__ identity():unary_function<T,T>(){}
|
||||
};
|
||||
|
||||
template <typename T1, typename T2> struct project1st : binary_function<T1, T2, T1>
|
||||
@@ -216,13 +282,18 @@ namespace cv { namespace gpu { namespace device
|
||||
{
|
||||
return lhs;
|
||||
}
|
||||
__device__ __forceinline__ project1st(const project1st& other):binary_function<T1,T2,T1>(){}
|
||||
__device__ __forceinline__ project1st():binary_function<T1,T2,T1>(){}
|
||||
};
|
||||
|
||||
template <typename T1, typename T2> struct project2nd : binary_function<T1, T2, T2>
|
||||
{
|
||||
__device__ __forceinline__ typename TypeTraits<T2>::ParameterType operator()(typename TypeTraits<T1>::ParameterType lhs, typename TypeTraits<T2>::ParameterType rhs) const
|
||||
{
|
||||
return rhs;
|
||||
}
|
||||
__device__ __forceinline__ project2nd(const project2nd& other):binary_function<T1,T2,T2>(){}
|
||||
__device__ __forceinline__ project2nd():binary_function<T1,T2,T2>(){}
|
||||
};
|
||||
|
||||
// Min/Max Operations
|
||||
@@ -231,6 +302,8 @@ namespace cv { namespace gpu { namespace device
|
||||
template <> struct name<type> : binary_function<type, type, type> \
|
||||
{ \
|
||||
__device__ __forceinline__ type operator()(type lhs, type rhs) const {return op(lhs, rhs);} \
|
||||
__device__ __forceinline__ name(const name& other):binary_function<type, type, type>(){}\
|
||||
__device__ __forceinline__ name():binary_function<type, type, type>(){}\
|
||||
};
|
||||
|
||||
template <typename T> struct maximum : binary_function<T, T, T>
|
||||
@@ -239,6 +312,8 @@ namespace cv { namespace gpu { namespace device
|
||||
{
|
||||
return lhs < rhs ? rhs : lhs;
|
||||
}
|
||||
__device__ __forceinline__ maximum(const maximum& other):binary_function<T, T, T>(){}
|
||||
__device__ __forceinline__ maximum():binary_function<T, T, T>(){}
|
||||
};
|
||||
|
||||
OPENCV_GPU_IMPLEMENT_MINMAX(maximum, uchar, ::max)
|
||||
@@ -257,6 +332,8 @@ namespace cv { namespace gpu { namespace device
|
||||
{
|
||||
return lhs < rhs ? lhs : rhs;
|
||||
}
|
||||
__device__ __forceinline__ minimum(const minimum& other):binary_function<T, T, T>(){}
|
||||
__device__ __forceinline__ minimum():binary_function<T, T, T>(){}
|
||||
};
|
||||
|
||||
OPENCV_GPU_IMPLEMENT_MINMAX(minimum, uchar, ::min)
|
||||
@@ -272,7 +349,7 @@ namespace cv { namespace gpu { namespace device
|
||||
#undef OPENCV_GPU_IMPLEMENT_MINMAX
|
||||
|
||||
// Math functions
|
||||
|
||||
///bound=========================================
|
||||
#define OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(name, func) \
|
||||
template <typename T> struct name ## _func : unary_function<T, float> \
|
||||
{ \
|
||||
@@ -342,17 +419,17 @@ namespace cv { namespace gpu { namespace device
|
||||
};
|
||||
|
||||
// Saturate Cast Functor
|
||||
|
||||
template <typename T, typename D> struct saturate_cast_func : unary_function<T, D>
|
||||
{
|
||||
__device__ __forceinline__ D operator ()(typename TypeTraits<T>::ParameterType v) const
|
||||
{
|
||||
return saturate_cast<D>(v);
|
||||
}
|
||||
__device__ __forceinline__ saturate_cast_func(const saturate_cast_func& other):unary_function<T, D>(){}
|
||||
__device__ __forceinline__ saturate_cast_func():unary_function<T, D>(){}
|
||||
};
|
||||
|
||||
// Threshold Functors
|
||||
|
||||
template <typename T> struct thresh_binary_func : unary_function<T, T>
|
||||
{
|
||||
__host__ __device__ __forceinline__ thresh_binary_func(T thresh_, T maxVal_) : thresh(thresh_), maxVal(maxVal_) {}
|
||||
@@ -361,10 +438,15 @@ namespace cv { namespace gpu { namespace device
|
||||
{
|
||||
return (src > thresh) * maxVal;
|
||||
}
|
||||
__device__ __forceinline__ thresh_binary_func(const thresh_binary_func& other)
|
||||
: unary_function<T, T>(), thresh(other.thresh), maxVal(other.maxVal){}
|
||||
|
||||
__device__ __forceinline__ thresh_binary_func():unary_function<T, T>(){}
|
||||
|
||||
const T thresh;
|
||||
const T maxVal;
|
||||
};
|
||||
|
||||
template <typename T> struct thresh_binary_inv_func : unary_function<T, T>
|
||||
{
|
||||
__host__ __device__ __forceinline__ thresh_binary_inv_func(T thresh_, T maxVal_) : thresh(thresh_), maxVal(maxVal_) {}
|
||||
@@ -373,10 +455,15 @@ namespace cv { namespace gpu { namespace device
|
||||
{
|
||||
return (src <= thresh) * maxVal;
|
||||
}
|
||||
__device__ __forceinline__ thresh_binary_inv_func(const thresh_binary_inv_func& other)
|
||||
: unary_function<T, T>(), thresh(other.thresh), maxVal(other.maxVal){}
|
||||
|
||||
__device__ __forceinline__ thresh_binary_inv_func():unary_function<T, T>(){}
|
||||
|
||||
const T thresh;
|
||||
const T maxVal;
|
||||
};
|
||||
|
||||
template <typename T> struct thresh_trunc_func : unary_function<T, T>
|
||||
{
|
||||
explicit __host__ __device__ __forceinline__ thresh_trunc_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {}
|
||||
@@ -386,8 +473,14 @@ namespace cv { namespace gpu { namespace device
|
||||
return minimum<T>()(src, thresh);
|
||||
}
|
||||
|
||||
__device__ __forceinline__ thresh_trunc_func(const thresh_trunc_func& other)
|
||||
: unary_function<T, T>(), thresh(other.thresh){}
|
||||
|
||||
__device__ __forceinline__ thresh_trunc_func():unary_function<T, T>(){}
|
||||
|
||||
const T thresh;
|
||||
};
|
||||
|
||||
template <typename T> struct thresh_to_zero_func : unary_function<T, T>
|
||||
{
|
||||
explicit __host__ __device__ __forceinline__ thresh_to_zero_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {}
|
||||
@@ -396,9 +489,14 @@ namespace cv { namespace gpu { namespace device
|
||||
{
|
||||
return (src > thresh) * src;
|
||||
}
|
||||
__device__ __forceinline__ thresh_to_zero_func(const thresh_to_zero_func& other)
|
||||
: unary_function<T, T>(), thresh(other.thresh){}
|
||||
|
||||
__device__ __forceinline__ thresh_to_zero_func():unary_function<T, T>(){}
|
||||
|
||||
const T thresh;
|
||||
};
|
||||
|
||||
template <typename T> struct thresh_to_zero_inv_func : unary_function<T, T>
|
||||
{
|
||||
explicit __host__ __device__ __forceinline__ thresh_to_zero_inv_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {}
|
||||
@@ -407,12 +505,15 @@ namespace cv { namespace gpu { namespace device
|
||||
{
|
||||
return (src <= thresh) * src;
|
||||
}
|
||||
__device__ __forceinline__ thresh_to_zero_inv_func(const thresh_to_zero_inv_func& other)
|
||||
: unary_function<T, T>(), thresh(other.thresh){}
|
||||
|
||||
__device__ __forceinline__ thresh_to_zero_inv_func():unary_function<T, T>(){}
|
||||
|
||||
const T thresh;
|
||||
};
|
||||
|
||||
//bound!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! ============>
|
||||
// Function Object Adaptors
|
||||
|
||||
template <typename Predicate> struct unary_negate : unary_function<typename Predicate::argument_type, bool>
|
||||
{
|
||||
explicit __host__ __device__ __forceinline__ unary_negate(const Predicate& p) : pred(p) {}
|
||||
|
||||
@@ -84,7 +84,7 @@ namespace cv { namespace gpu { namespace device
|
||||
}
|
||||
template<> __device__ __forceinline__ uchar saturate_cast<uchar>(double v)
|
||||
{
|
||||
#if __CUDA_ARCH__ >= 130
|
||||
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
|
||||
int iv = __double2int_rn(v);
|
||||
return saturate_cast<uchar>(iv);
|
||||
#else
|
||||
@@ -120,7 +120,7 @@ namespace cv { namespace gpu { namespace device
|
||||
}
|
||||
template<> __device__ __forceinline__ schar saturate_cast<schar>(double v)
|
||||
{
|
||||
#if __CUDA_ARCH__ >= 130
|
||||
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
|
||||
int iv = __double2int_rn(v);
|
||||
return saturate_cast<schar>(iv);
|
||||
#else
|
||||
@@ -151,7 +151,7 @@ namespace cv { namespace gpu { namespace device
|
||||
}
|
||||
template<> __device__ __forceinline__ ushort saturate_cast<ushort>(double v)
|
||||
{
|
||||
#if __CUDA_ARCH__ >= 130
|
||||
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
|
||||
int iv = __double2int_rn(v);
|
||||
return saturate_cast<ushort>(iv);
|
||||
#else
|
||||
@@ -178,7 +178,7 @@ namespace cv { namespace gpu { namespace device
|
||||
}
|
||||
template<> __device__ __forceinline__ short saturate_cast<short>(double v)
|
||||
{
|
||||
#if __CUDA_ARCH__ >= 130
|
||||
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
|
||||
int iv = __double2int_rn(v);
|
||||
return saturate_cast<short>(iv);
|
||||
#else
|
||||
@@ -192,7 +192,7 @@ namespace cv { namespace gpu { namespace device
|
||||
}
|
||||
template<> __device__ __forceinline__ int saturate_cast<int>(double v)
|
||||
{
|
||||
#if __CUDA_ARCH__ >= 130
|
||||
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
|
||||
return __double2int_rn(v);
|
||||
#else
|
||||
return saturate_cast<int>((float)v);
|
||||
@@ -205,7 +205,7 @@ namespace cv { namespace gpu { namespace device
|
||||
}
|
||||
template<> __device__ __forceinline__ uint saturate_cast<uint>(double v)
|
||||
{
|
||||
#if __CUDA_ARCH__ >= 130
|
||||
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
|
||||
return __double2uint_rn(v);
|
||||
#else
|
||||
return saturate_cast<uint>((float)v);
|
||||
@@ -213,4 +213,4 @@ namespace cv { namespace gpu { namespace device
|
||||
}
|
||||
}}}
|
||||
|
||||
#endif /* __OPENCV_GPU_SATURATE_CAST_HPP__ */
|
||||
#endif /* __OPENCV_GPU_SATURATE_CAST_HPP__ */
|
||||
|
||||
@@ -50,14 +50,14 @@
|
||||
namespace cv { namespace gpu { namespace device
|
||||
{
|
||||
template <typename T, typename D, typename UnOp, typename Mask>
|
||||
static inline void transform(DevMem2D_<T> src, DevMem2D_<D> dst, UnOp op, Mask mask, cudaStream_t stream)
|
||||
static inline void transform(DevMem2D_<T> src, DevMem2D_<D> dst, UnOp op, const Mask& mask, cudaStream_t stream)
|
||||
{
|
||||
typedef TransformFunctorTraits<UnOp> ft;
|
||||
transform_detail::TransformDispatcher<VecTraits<T>::cn == 1 && VecTraits<D>::cn == 1 && ft::smart_shift != 1>::call(src, dst, op, mask, stream);
|
||||
}
|
||||
|
||||
template <typename T1, typename T2, typename D, typename BinOp, typename Mask>
|
||||
static inline void transform(DevMem2D_<T1> src1, DevMem2D_<T2> src2, DevMem2D_<D> dst, BinOp op, Mask mask, cudaStream_t stream)
|
||||
static inline void transform(DevMem2D_<T1> src1, DevMem2D_<T2> src2, DevMem2D_<D> dst, BinOp op, const Mask& mask, cudaStream_t stream)
|
||||
{
|
||||
typedef TransformFunctorTraits<BinOp> ft;
|
||||
transform_detail::TransformDispatcher<VecTraits<T1>::cn == 1 && VecTraits<T2>::cn == 1 && VecTraits<D>::cn == 1 && ft::smart_shift != 1>::call(src1, src2, dst, op, mask, stream);
|
||||
|
||||
@@ -70,6 +70,7 @@ namespace cv { namespace gpu { namespace device
|
||||
struct SingleMask
|
||||
{
|
||||
explicit __host__ __device__ __forceinline__ SingleMask(PtrStepb mask_) : mask(mask_) {}
|
||||
__host__ __device__ __forceinline__ SingleMask(const SingleMask& mask_): mask(mask_.mask){}
|
||||
|
||||
__device__ __forceinline__ bool operator()(int y, int x) const
|
||||
{
|
||||
@@ -81,7 +82,10 @@ namespace cv { namespace gpu { namespace device
|
||||
|
||||
struct SingleMaskChannels
|
||||
{
|
||||
__host__ __device__ __forceinline__ SingleMaskChannels(PtrStepb mask_, int channels_) : mask(mask_), channels(channels_) {}
|
||||
__host__ __device__ __forceinline__ SingleMaskChannels(PtrStepb mask_, int channels_)
|
||||
: mask(mask_), channels(channels_) {}
|
||||
__host__ __device__ __forceinline__ SingleMaskChannels(const SingleMaskChannels& mask_)
|
||||
:mask(mask_.mask), channels(mask_.channels){}
|
||||
|
||||
__device__ __forceinline__ bool operator()(int y, int x) const
|
||||
{
|
||||
@@ -94,7 +98,11 @@ namespace cv { namespace gpu { namespace device
|
||||
|
||||
struct MaskCollection
|
||||
{
|
||||
explicit __host__ __device__ __forceinline__ MaskCollection(PtrStepb* maskCollection_) : maskCollection(maskCollection_) {}
|
||||
explicit __host__ __device__ __forceinline__ MaskCollection(PtrStepb* maskCollection_)
|
||||
: maskCollection(maskCollection_) {}
|
||||
|
||||
__device__ __forceinline__ MaskCollection(const MaskCollection& masks_)
|
||||
: maskCollection(masks_.maskCollection), curMask(masks_.curMask){}
|
||||
|
||||
__device__ __forceinline__ void next()
|
||||
{
|
||||
@@ -117,6 +125,9 @@ namespace cv { namespace gpu { namespace device
|
||||
|
||||
struct WithOutMask
|
||||
{
|
||||
__device__ __forceinline__ WithOutMask(){}
|
||||
__device__ __forceinline__ WithOutMask(const WithOutMask& mask){}
|
||||
|
||||
__device__ __forceinline__ void next() const
|
||||
{
|
||||
}
|
||||
|
||||
@@ -43,7 +43,7 @@
|
||||
#ifndef __OPENCV_PRECOMP_H__
|
||||
#define __OPENCV_PRECOMP_H__
|
||||
|
||||
#if _MSC_VER >= 1200
|
||||
#if defined _MSC_VER && _MSC_VER >= 1200
|
||||
#pragma warning( disable: 4251 4710 4711 4514 4996 )
|
||||
#endif
|
||||
|
||||
|
||||
25
modules/gpu/test/main_test_nvidia.h
Normal file
25
modules/gpu/test/main_test_nvidia.h
Normal file
@@ -0,0 +1,25 @@
|
||||
#ifndef __main_test_nvidia_h__
|
||||
#define __main_test_nvidia_h__
|
||||
|
||||
#include<string>
|
||||
|
||||
enum OutputLevel
|
||||
{
|
||||
OutputLevelNone,
|
||||
OutputLevelCompact,
|
||||
OutputLevelFull
|
||||
};
|
||||
|
||||
bool nvidia_NPPST_Integral_Image(const std::string& test_data_path, OutputLevel outputLevel);
|
||||
bool nvidia_NPPST_Squared_Integral_Image(const std::string& test_data_path, OutputLevel outputLevel);
|
||||
bool nvidia_NPPST_RectStdDev(const std::string& test_data_path, OutputLevel outputLevel);
|
||||
bool nvidia_NPPST_Resize(const std::string& test_data_path, OutputLevel outputLevel);
|
||||
bool nvidia_NPPST_Vector_Operations(const std::string& test_data_path, OutputLevel outputLevel);
|
||||
bool nvidia_NPPST_Transpose(const std::string& test_data_path, OutputLevel outputLevel);
|
||||
bool nvidia_NCV_Vector_Operations(const std::string& test_data_path, OutputLevel outputLevel);
|
||||
bool nvidia_NCV_Haar_Cascade_Loader(const std::string& test_data_path, OutputLevel outputLevel);
|
||||
bool nvidia_NCV_Haar_Cascade_Application(const std::string& test_data_path, OutputLevel outputLevel);
|
||||
bool nvidia_NCV_Hypotheses_Filtration(const std::string& test_data_path, OutputLevel outputLevel);
|
||||
bool nvidia_NCV_Visualization(const std::string& test_data_path, OutputLevel outputLevel);
|
||||
|
||||
#endif
|
||||
@@ -14,13 +14,13 @@
|
||||
#include <vector>
|
||||
|
||||
#include "NCVTest.hpp"
|
||||
|
||||
enum OutputLevel
|
||||
{
|
||||
OutputLevelNone,
|
||||
OutputLevelCompact,
|
||||
OutputLevelFull
|
||||
};
|
||||
#include <main_test_nvidia.h>
|
||||
//enum OutputLevel
|
||||
//{
|
||||
// OutputLevelNone,
|
||||
// OutputLevelCompact,
|
||||
// OutputLevelFull
|
||||
//};
|
||||
|
||||
class NCVAutoTestLister
|
||||
{
|
||||
|
||||
@@ -11,7 +11,9 @@
|
||||
#ifndef _ncvtest_hpp_
|
||||
#define _ncvtest_hpp_
|
||||
|
||||
#pragma warning( disable : 4201 4408 4127 4100)
|
||||
#if defined _MSC_VER
|
||||
# pragma warning( disable : 4201 4408 4127 4100)
|
||||
#endif
|
||||
|
||||
#include <string>
|
||||
#include <vector>
|
||||
@@ -36,6 +38,7 @@ class INCVTest
|
||||
public:
|
||||
virtual bool executeTest(NCVTestReport &report) = 0;
|
||||
virtual std::string getName() const = 0;
|
||||
virtual ~INCVTest(){}
|
||||
};
|
||||
|
||||
|
||||
|
||||
@@ -21,11 +21,15 @@
|
||||
#include "NCVAutoTestLister.hpp"
|
||||
#include "NCVTestSourceProvider.hpp"
|
||||
|
||||
#include <main_test_nvidia.h>
|
||||
|
||||
static std::string path;
|
||||
|
||||
namespace {
|
||||
|
||||
template <class T_in, class T_out>
|
||||
void generateIntegralTests(NCVAutoTestLister &testLister, NCVTestSourceProvider<T_in> &src,
|
||||
void generateIntegralTests(NCVAutoTestLister &testLister,
|
||||
NCVTestSourceProvider<T_in> &src,
|
||||
Ncv32u maxWidth, Ncv32u maxHeight)
|
||||
{
|
||||
for (Ncv32f _i=1.0; _i<maxWidth; _i*=1.2f)
|
||||
@@ -43,13 +47,9 @@ void generateIntegralTests(NCVAutoTestLister &testLister, NCVTestSourceProvider<
|
||||
testLister.add(new TestIntegralImage<T_in, T_out>(testName, src, 2, i));
|
||||
}
|
||||
|
||||
//test VGA
|
||||
testLister.add(new TestIntegralImage<T_in, T_out>("LinIntImg_VGA", src, 640, 480));
|
||||
|
||||
//TODO: add tests of various resolutions up to 4096x4096
|
||||
}
|
||||
|
||||
|
||||
void generateSquaredIntegralTests(NCVAutoTestLister &testLister, NCVTestSourceProvider<Ncv8u> &src,
|
||||
Ncv32u maxWidth, Ncv32u maxHeight)
|
||||
{
|
||||
@@ -68,13 +68,9 @@ void generateSquaredIntegralTests(NCVAutoTestLister &testLister, NCVTestSourcePr
|
||||
testLister.add(new TestIntegralImageSquared(testName, src, 32, i));
|
||||
}
|
||||
|
||||
//test VGA
|
||||
testLister.add(new TestIntegralImageSquared("SqLinIntImg_VGA", src, 640, 480));
|
||||
|
||||
//TODO: add tests of various resolutions up to 4096x4096
|
||||
}
|
||||
|
||||
|
||||
void generateRectStdDevTests(NCVAutoTestLister &testLister, NCVTestSourceProvider<Ncv8u> &src,
|
||||
Ncv32u maxWidth, Ncv32u maxHeight)
|
||||
{
|
||||
@@ -91,17 +87,12 @@ void generateRectStdDevTests(NCVAutoTestLister &testLister, NCVTestSourceProvide
|
||||
testLister.add(new TestRectStdDev(testName, src, i-1, i*2-1, rect, 2.5, true));
|
||||
}
|
||||
|
||||
//test VGA
|
||||
testLister.add(new TestRectStdDev("RectStdDev_VGA", src, 640, 480, rect, 1, true));
|
||||
|
||||
//TODO: add tests of various resolutions up to 4096x4096
|
||||
}
|
||||
|
||||
|
||||
template <class T>
|
||||
void generateResizeTests(NCVAutoTestLister &testLister, NCVTestSourceProvider<T> &src)
|
||||
{
|
||||
//test VGA
|
||||
for (Ncv32u i=1; i<480; i+=3)
|
||||
{
|
||||
char testName[80];
|
||||
@@ -110,7 +101,6 @@ void generateResizeTests(NCVAutoTestLister &testLister, NCVTestSourceProvider<T>
|
||||
testLister.add(new TestResize<T>(testName, src, 640, 480, i, false));
|
||||
}
|
||||
|
||||
//test HD
|
||||
for (Ncv32u i=1; i<1080; i+=5)
|
||||
{
|
||||
char testName[80];
|
||||
@@ -118,11 +108,8 @@ void generateResizeTests(NCVAutoTestLister &testLister, NCVTestSourceProvider<T>
|
||||
testLister.add(new TestResize<T>(testName, src, 1920, 1080, i, true));
|
||||
testLister.add(new TestResize<T>(testName, src, 1920, 1080, i, false));
|
||||
}
|
||||
|
||||
//TODO: add tests of various resolutions up to 4096x4096
|
||||
}
|
||||
|
||||
|
||||
void generateNPPSTVectorTests(NCVAutoTestLister &testLister, NCVTestSourceProvider<Ncv32u> &src, Ncv32u maxLength)
|
||||
{
|
||||
//compaction
|
||||
@@ -186,9 +173,10 @@ void generateTransposeTests(NCVAutoTestLister &testLister, NCVTestSourceProvider
|
||||
testLister.add(new TestTranspose<T>("TestTranspose_reg_0", src, 1072, 375));
|
||||
}
|
||||
|
||||
|
||||
template <class T>
|
||||
void generateDrawRectsTests(NCVAutoTestLister &testLister, NCVTestSourceProvider<T> &src, NCVTestSourceProvider<Ncv32u> &src32u,
|
||||
void generateDrawRectsTests(NCVAutoTestLister &testLister,
|
||||
NCVTestSourceProvider<T> &src,
|
||||
NCVTestSourceProvider<Ncv32u> &src32u,
|
||||
Ncv32u maxWidth, Ncv32u maxHeight)
|
||||
{
|
||||
for (Ncv32f _i=16.0; _i<maxWidth; _i*=1.1f)
|
||||
@@ -215,11 +203,8 @@ void generateDrawRectsTests(NCVAutoTestLister &testLister, NCVTestSourceProvider
|
||||
|
||||
//test VGA
|
||||
testLister.add(new TestDrawRects<T>("DrawRects_VGA", src, src32u, 640, 480, 640*480/1000, (T)0xFF));
|
||||
|
||||
//TODO: add tests of various resolutions up to 4096x4096
|
||||
}
|
||||
|
||||
|
||||
void generateVectorTests(NCVAutoTestLister &testLister, NCVTestSourceProvider<Ncv32u> &src, Ncv32u maxLength)
|
||||
{
|
||||
//growth
|
||||
@@ -237,7 +222,6 @@ void generateVectorTests(NCVAutoTestLister &testLister, NCVTestSourceProvider<Nc
|
||||
testLister.add(new TestHypothesesGrow("VectorGrow00b", src, 10, 42, 1.2f, 10, 0, 10, 0));
|
||||
}
|
||||
|
||||
|
||||
void generateHypothesesFiltrationTests(NCVAutoTestLister &testLister, NCVTestSourceProvider<Ncv32u> &src, Ncv32u maxLength)
|
||||
{
|
||||
for (Ncv32f _i=1.0; _i<maxLength; _i*=1.1f)
|
||||
@@ -261,7 +245,6 @@ void generateHaarLoaderTests(NCVAutoTestLister &testLister)
|
||||
testLister.add(new TestHaarCascadeLoader("haarcascade_eye_tree_eyeglasses.xml", path + "haarcascade_eye_tree_eyeglasses.xml"));
|
||||
}
|
||||
|
||||
|
||||
void generateHaarApplicationTests(NCVAutoTestLister &testLister, NCVTestSourceProvider<Ncv8u> &src,
|
||||
Ncv32u maxWidth, Ncv32u maxHeight)
|
||||
{
|
||||
@@ -285,7 +268,6 @@ void generateHaarApplicationTests(NCVAutoTestLister &testLister, NCVTestSourcePr
|
||||
|
||||
static void devNullOutput(const std::string& msg)
|
||||
{
|
||||
|
||||
}
|
||||
|
||||
bool nvidia_NPPST_Integral_Image(const std::string& test_data_path, OutputLevel outputLevel)
|
||||
@@ -304,6 +286,8 @@ bool nvidia_NPPST_Integral_Image(const std::string& test_data_path, OutputLevel
|
||||
return testListerII.invoke();
|
||||
}
|
||||
|
||||
}
|
||||
|
||||
bool nvidia_NPPST_Squared_Integral_Image(const std::string& test_data_path, OutputLevel outputLevel)
|
||||
{
|
||||
path = test_data_path;
|
||||
|
||||
@@ -39,6 +39,10 @@
|
||||
//
|
||||
//M*/
|
||||
|
||||
#ifdef __GNUC__
|
||||
# pragma GCC diagnostic ignored "-Wmissing-declarations"
|
||||
#endif
|
||||
|
||||
#ifndef __OPENCV_TEST_PRECOMP_HPP__
|
||||
#define __OPENCV_TEST_PRECOMP_HPP__
|
||||
|
||||
|
||||
@@ -513,7 +513,6 @@ PARAM_TEST_CASE(Filter2D, cv::gpu::DeviceInfo, cv::Size, MatType, KSize, Anchor,
|
||||
bool useRoi;
|
||||
|
||||
cv::Mat img;
|
||||
cv::Mat kernel;
|
||||
|
||||
virtual void SetUp()
|
||||
{
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
@@ -39,6 +39,7 @@
|
||||
//
|
||||
//M*/
|
||||
|
||||
#include <main_test_nvidia.h>
|
||||
#include "precomp.hpp"
|
||||
|
||||
#ifdef HAVE_CUDA
|
||||
@@ -46,24 +47,12 @@
|
||||
using namespace cvtest;
|
||||
using namespace testing;
|
||||
|
||||
enum OutputLevel
|
||||
{
|
||||
OutputLevelNone,
|
||||
OutputLevelCompact,
|
||||
OutputLevelFull
|
||||
};
|
||||
|
||||
bool nvidia_NPPST_Integral_Image(const std::string& test_data_path, OutputLevel outputLevel);
|
||||
bool nvidia_NPPST_Squared_Integral_Image(const std::string& test_data_path, OutputLevel outputLevel);
|
||||
bool nvidia_NPPST_RectStdDev(const std::string& test_data_path, OutputLevel outputLevel);
|
||||
bool nvidia_NPPST_Resize(const std::string& test_data_path, OutputLevel outputLevel);
|
||||
bool nvidia_NPPST_Vector_Operations(const std::string& test_data_path, OutputLevel outputLevel);
|
||||
bool nvidia_NPPST_Transpose(const std::string& test_data_path, OutputLevel outputLevel);
|
||||
bool nvidia_NCV_Vector_Operations(const std::string& test_data_path, OutputLevel outputLevel);
|
||||
bool nvidia_NCV_Haar_Cascade_Loader(const std::string& test_data_path, OutputLevel outputLevel);
|
||||
bool nvidia_NCV_Haar_Cascade_Application(const std::string& test_data_path, OutputLevel outputLevel);
|
||||
bool nvidia_NCV_Hypotheses_Filtration(const std::string& test_data_path, OutputLevel outputLevel);
|
||||
bool nvidia_NCV_Visualization(const std::string& test_data_path, OutputLevel outputLevel);
|
||||
//enum OutputLevel
|
||||
//{
|
||||
// OutputLevelNone,
|
||||
// OutputLevelCompact,
|
||||
// OutputLevelFull
|
||||
//};
|
||||
|
||||
struct NVidiaTest : TestWithParam<cv::gpu::DeviceInfo>
|
||||
{
|
||||
@@ -86,12 +75,12 @@ struct NCV : NVidiaTest {};
|
||||
|
||||
OutputLevel nvidiaTestOutputLevel = OutputLevelCompact;
|
||||
|
||||
TEST_P(NPPST, Integral)
|
||||
{
|
||||
bool res = nvidia_NPPST_Integral_Image(path, nvidiaTestOutputLevel);
|
||||
//TEST_P(NPPST, Integral)
|
||||
//{
|
||||
// bool res = nvidia_NPPST_Integral_Image(path, nvidiaTestOutputLevel);
|
||||
|
||||
ASSERT_TRUE(res);
|
||||
}
|
||||
// ASSERT_TRUE(res);
|
||||
//}
|
||||
|
||||
TEST_P(NPPST, SquaredIntegral)
|
||||
{
|
||||
|
||||
@@ -69,16 +69,16 @@ struct HOG : testing::TestWithParam<cv::gpu::DeviceInfo>, cv::gpu::HOGDescriptor
|
||||
}
|
||||
|
||||
#ifdef DUMP
|
||||
void dump(const cv::Mat& block_hists, const std::vector<cv::Point>& locations)
|
||||
void dump(const cv::Mat& blockHists, const std::vector<cv::Point>& locations)
|
||||
{
|
||||
f.write((char*)&block_hists.rows, sizeof(block_hists.rows));
|
||||
f.write((char*)&block_hists.cols, sizeof(block_hists.cols));
|
||||
f.write((char*)&blockHists.rows, sizeof(blockHists.rows));
|
||||
f.write((char*)&blockHists.cols, sizeof(blockHists.cols));
|
||||
|
||||
for (int i = 0; i < block_hists.rows; ++i)
|
||||
for (int i = 0; i < blockHists.rows; ++i)
|
||||
{
|
||||
for (int j = 0; j < block_hists.cols; ++j)
|
||||
for (int j = 0; j < blockHists.cols; ++j)
|
||||
{
|
||||
float val = block_hists.at<float>(i, j);
|
||||
float val = blockHists.at<float>(i, j);
|
||||
f.write((char*)&val, sizeof(val));
|
||||
}
|
||||
}
|
||||
@@ -90,21 +90,21 @@ struct HOG : testing::TestWithParam<cv::gpu::DeviceInfo>, cv::gpu::HOGDescriptor
|
||||
f.write((char*)&locations[i], sizeof(locations[i]));
|
||||
}
|
||||
#else
|
||||
void compare(const cv::Mat& block_hists, const std::vector<cv::Point>& locations)
|
||||
void compare(const cv::Mat& blockHists, const std::vector<cv::Point>& locations)
|
||||
{
|
||||
int rows, cols;
|
||||
f.read((char*)&rows, sizeof(rows));
|
||||
f.read((char*)&cols, sizeof(cols));
|
||||
ASSERT_EQ(rows, block_hists.rows);
|
||||
ASSERT_EQ(cols, block_hists.cols);
|
||||
ASSERT_EQ(rows, blockHists.rows);
|
||||
ASSERT_EQ(cols, blockHists.cols);
|
||||
|
||||
for (int i = 0; i < block_hists.rows; ++i)
|
||||
for (int i = 0; i < blockHists.rows; ++i)
|
||||
{
|
||||
for (int j = 0; j < block_hists.cols; ++j)
|
||||
for (int j = 0; j < blockHists.cols; ++j)
|
||||
{
|
||||
float val;
|
||||
f.read((char*)&val, sizeof(val));
|
||||
ASSERT_NEAR(val, block_hists.at<float>(i, j), 1e-3);
|
||||
ASSERT_NEAR(val, blockHists.at<float>(i, j), 1e-3);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
Reference in New Issue
Block a user