Chiebot-Mirror
/
opencv
8
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

fixes for the newly added gcc warning keys

Browse Source
pull/2/head
Marina Kolpakova 13 years ago
parent f6ef504ef0
commit b065c7a296
20 changed files with 113 additions and 67 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Unified View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 1
      modules/gpu/CMakeLists.txt
  2. 2
      modules/gpu/src/arithm.cpp
  3. 3
      modules/gpu/src/cuda/bf_knnmatch.cu
  4. 2
      modules/gpu/src/cuda/bf_match.cu
  5. 2
      modules/gpu/src/cuda/bf_radius_match.cu
  6. 2
      modules/gpu/src/cuda/hog.cu
  7. 3
      modules/gpu/src/cuda/remap.cu
  8. 7
      modules/gpu/src/cuda/resize.cu
  9. 4
      modules/gpu/src/cuda/warp.cu
  10. 10
      modules/gpu/src/nvidia/NCVBroxOpticalFlow.cu
  11. 3
      modules/gpu/src/nvidia/NCVHaarObjectDetection.cu
  12. 1
      modules/gpu/src/nvidia/core/NCV.cu
  13. 22
      modules/gpu/src/nvidia/core/NCV.hpp
  14. 1
      modules/gpu/src/nvidia/core/NCVRuntimeTemplates.hpp
  15. 22
      modules/gpu/src/opencv2/gpu/device/filters.hpp
  16. 6
      modules/gpu/src/opencv2/gpu/device/functional.hpp
  17. 34
      modules/gpu/test/nvidia/TestHypothesesFilter.cpp
  18. 24
      modules/gpu/test/nvidia/TestResize.cpp
  19. 5
      modules/gpu/test/nvidia/main_nvidia.cpp
  20. 26
      modules/gpu/test/test_nvidia.cpp

1
modules/gpu/CMakeLists.txt
Unescape View File

@ -30,6 +30,7 @@ if (HAVE_CUDA)
source_group("Src\\NVidia" FILES ${ncv_files}) source_group("Src\\NVidia" FILES ${ncv_files})
ocv_include_directories("src/nvidia" "src/nvidia/core" "src/nvidia/NPP_staging" ${CUDA_INCLUDE_DIRS}) 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) ocv_warnings_disable(CMAKE_CXX_FLAGS -Wundef -Wmissing-declarations /wd4211 /wd4201 /wd4100 /wd4505 /wd4408)
string(REPLACE "-Wsign-promo" "" CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}")
#set (CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} "-keep") #set (CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} "-keep")
#set (CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} "-Xcompiler;/EHsc-;") #set (CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} "-Xcompiler;/EHsc-;")

2
modules/gpu/src/arithm.cpp
Unescape View File

@ -68,7 +68,7 @@ void cv::gpu::polarToCart(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, bool,
void cv::gpu::gemm(const GpuMat& src1, const GpuMat& src2, double alpha, const GpuMat& src3, double beta, GpuMat& dst, int flags, Stream& stream) void cv::gpu::gemm(const GpuMat& src1, const GpuMat& src2, double alpha, const GpuMat& src3, double beta, GpuMat& dst, int flags, Stream& stream)
{ {
#ifndef HAVE_CUBLAS #ifndef HAVE_CUBLAS
(void)src1; (void)src2; (void)alpha; (void)src3; (void)beta; (void)dst; (void)flags; (void)stream;
CV_Error(CV_StsNotImplemented, "The library was build without CUBLAS"); CV_Error(CV_StsNotImplemented, "The library was build without CUBLAS");
#else #else

3
modules/gpu/src/cuda/bf_knnmatch.cu
Unescape View File

@ -748,6 +748,7 @@ namespace cv { namespace gpu { namespace device
const DevMem2Db& trainIdx, const DevMem2Db& distance, const DevMem2Db& trainIdx, const DevMem2Db& distance,
int cc, cudaStream_t stream) int cc, cudaStream_t stream)
{ {
(void)cc;
if (query.cols <= 64) if (query.cols <= 64)
{ {
matchUnrolledCached<16, 64, Dist>(query, train, mask, static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<float2> > (distance), stream); matchUnrolledCached<16, 64, Dist>(query, train, mask, static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<float2> > (distance), stream);
@ -779,6 +780,7 @@ namespace cv { namespace gpu { namespace device
const DevMem2Db& trainIdx, const DevMem2Db& imgIdx, const DevMem2Db& distance, const DevMem2Db& trainIdx, const DevMem2Db& imgIdx, const DevMem2Db& distance,
int cc, cudaStream_t stream) int cc, cudaStream_t stream)
{ {
(void)cc;
if (query.cols <= 64) if (query.cols <= 64)
{ {
matchUnrolledCached<16, 64, Dist>(query, trains, n, mask, static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<int2> >(imgIdx), static_cast< DevMem2D_<float2> > (distance), stream); matchUnrolledCached<16, 64, Dist>(query, trains, n, mask, static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<int2> >(imgIdx), static_cast< DevMem2D_<float2> > (distance), stream);
@ -943,6 +945,7 @@ namespace cv { namespace gpu { namespace device
const DevMem2Df& allDist, const DevMem2Df& allDist,
int cc, cudaStream_t stream) int cc, cudaStream_t stream)
{ {
(void)cc;
if (query.cols <= 64) if (query.cols <= 64)
{ {
calcDistanceUnrolled<16, 64, Dist>(query, train, mask, allDist, stream); calcDistanceUnrolled<16, 64, Dist>(query, train, mask, allDist, stream);

2
modules/gpu/src/cuda/bf_match.cu
Unescape View File

@ -567,6 +567,7 @@ namespace cv { namespace gpu { namespace device
const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Di& trainIdx, const DevMem2Df& distance,
int cc, cudaStream_t stream) int cc, cudaStream_t stream)
{ {
(void)cc;
if (query.cols <= 64) if (query.cols <= 64)
{ {
matchUnrolledCached<16, 64, Dist>(query, train, mask, trainIdx, distance, stream); matchUnrolledCached<16, 64, Dist>(query, train, mask, trainIdx, distance, stream);
@ -598,6 +599,7 @@ namespace cv { namespace gpu { namespace device
const DevMem2Di& trainIdx, const DevMem2Di& imgIdx, const DevMem2Df& distance, const DevMem2Di& trainIdx, const DevMem2Di& imgIdx, const DevMem2Df& distance,
int cc, cudaStream_t stream) int cc, cudaStream_t stream)
{ {
(void)cc;
if (query.cols <= 64) if (query.cols <= 64)
{ {
matchUnrolledCached<16, 64, Dist>(query, trains, n, mask, trainIdx, imgIdx, distance, stream); matchUnrolledCached<16, 64, Dist>(query, trains, n, mask, trainIdx, imgIdx, distance, stream);

2
modules/gpu/src/cuda/bf_radius_match.cu
Unescape View File

@ -281,6 +281,7 @@ namespace cv { namespace gpu { namespace device
const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2D_<unsigned int>& nMatches, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2D_<unsigned int>& nMatches,
int cc, cudaStream_t stream) int cc, cudaStream_t stream)
{ {
(void)cc;
if (query.cols <= 64) if (query.cols <= 64)
{ {
matchUnrolled<16, 64, Dist>(query, train, maxDistance, mask, trainIdx, distance, nMatches, stream); matchUnrolled<16, 64, Dist>(query, train, maxDistance, mask, trainIdx, distance, nMatches, stream);
@ -312,6 +313,7 @@ namespace cv { namespace gpu { namespace device
const DevMem2Di& trainIdx, const DevMem2Di& imgIdx, const DevMem2Df& distance, const DevMem2D_<unsigned int>& nMatches, const DevMem2Di& trainIdx, const DevMem2Di& imgIdx, const DevMem2Df& distance, const DevMem2D_<unsigned int>& nMatches,
int cc, cudaStream_t stream) int cc, cudaStream_t stream)
{ {
(void)cc;
if (query.cols <= 64) if (query.cols <= 64)
{ {
matchUnrolled<16, 64, Dist>(query, trains, n, maxDistance, masks, trainIdx, imgIdx, distance, nMatches, stream); matchUnrolled<16, 64, Dist>(query, trains, n, maxDistance, masks, trainIdx, imgIdx, distance, nMatches, stream);

2
modules/gpu/src/cuda/hog.cu
Unescape View File

@ -619,6 +619,7 @@ namespace cv { namespace gpu { namespace device
void compute_gradients_8UC4(int nbins, int height, int width, const DevMem2Db& img, void compute_gradients_8UC4(int nbins, int height, int width, const DevMem2Db& img,
float angle_scale, DevMem2Df grad, DevMem2Db qangle, bool correct_gamma) float angle_scale, DevMem2Df grad, DevMem2Db qangle, bool correct_gamma)
{ {
(void)nbins;
const int nthreads = 256; const int nthreads = 256;
dim3 bdim(nthreads, 1); dim3 bdim(nthreads, 1);
@ -691,6 +692,7 @@ namespace cv { namespace gpu { namespace device
void compute_gradients_8UC1(int nbins, int height, int width, const DevMem2Db& img, void compute_gradients_8UC1(int nbins, int height, int width, const DevMem2Db& img,
float angle_scale, DevMem2Df grad, DevMem2Db qangle, bool correct_gamma) float angle_scale, DevMem2Df grad, DevMem2Db qangle, bool correct_gamma)
{ {
(void)nbins;
const int nthreads = 256; const int nthreads = 256;
dim3 bdim(nthreads, 1); dim3 bdim(nthreads, 1);

3
modules/gpu/src/cuda/remap.cu
Unescape View File

@ -87,6 +87,9 @@ namespace cv { namespace gpu { namespace device
{ {
static void call(DevMem2D_<T> src, DevMem2D_<T> srcWhole, int xoff, int yoff, DevMem2Df mapx, DevMem2Df mapy, DevMem2D_<T> dst, const float* borderValue, int) static void call(DevMem2D_<T> src, DevMem2D_<T> srcWhole, int xoff, int yoff, DevMem2Df mapx, DevMem2Df mapy, DevMem2D_<T> dst, const float* borderValue, int)
{ {
(void)srcWhole;
(void)xoff;
(void)yoff;
typedef typename TypeVec<float, VecTraits<T>::cn>::vec_type work_type; typedef typename TypeVec<float, VecTraits<T>::cn>::vec_type work_type;
dim3 block(32, 8); dim3 block(32, 8);

7
modules/gpu/src/cuda/resize.cu
Unescape View File

@ -131,6 +131,10 @@ namespace cv { namespace gpu { namespace device
{ {
static void call(DevMem2D_<T> src, DevMem2D_<T> srcWhole, int xoff, int yoff, float fx, float fy, DevMem2D_<T> dst) static void call(DevMem2D_<T> src, DevMem2D_<T> srcWhole, int xoff, int yoff, float fx, float fy, DevMem2D_<T> dst)
{ {
(void)srcWhole;
(void)xoff;
(void)yoff;
dim3 block(32, 8); dim3 block(32, 8);
dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y)); dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y));
@ -219,6 +223,9 @@ namespace cv { namespace gpu { namespace device
{ {
static void call(DevMem2D_<T> src, DevMem2D_<T> srcWhole, int xoff, int yoff, float fx, float fy, DevMem2D_<T> dst, cudaStream_t stream) static void call(DevMem2D_<T> src, DevMem2D_<T> srcWhole, int xoff, int yoff, float fx, float fy, DevMem2D_<T> dst, cudaStream_t stream)
{ {
(void)srcWhole;
(void)xoff;
(void)yoff;
int iscale_x = round(fx); int iscale_x = round(fx);
int iscale_y = round(fy); int iscale_y = round(fy);

4
modules/gpu/src/cuda/warp.cu
Unescape View File

@ -158,6 +158,10 @@ namespace cv { namespace gpu { namespace device
{ {
static void call(DevMem2D_<T> src, DevMem2D_<T> srcWhole, int xoff, int yoff, DevMem2D_<T> dst, const float* borderValue, int) static void call(DevMem2D_<T> src, DevMem2D_<T> srcWhole, int xoff, int yoff, DevMem2D_<T> dst, const float* borderValue, int)
{ {
(void)xoff;
(void)yoff;
(void)srcWhole;
typedef typename TypeVec<float, VecTraits<T>::cn>::vec_type work_type; typedef typename TypeVec<float, VecTraits<T>::cn>::vec_type work_type;
dim3 block(32, 8); dim3 block(32, 8);

10
modules/gpu/src/nvidia/NCVBroxOpticalFlow.cu
Unescape View File

@ -1136,7 +1136,7 @@ NCVStatus NCVBroxOpticalFlow(const NCVBroxOpticalFlowDescriptor desc,
ptrVNew->ptr(), dstSize, ns * sizeof (float), dstROI, 1.0f/scale_factor, 1.0f/scale_factor, nppStBicubic) ); ptrVNew->ptr(), dstSize, ns * sizeof (float), dstROI, 1.0f/scale_factor, 1.0f/scale_factor, nppStBicubic) );
ScaleVector(ptrVNew->ptr(), ptrVNew->ptr(), 1.0f/scale_factor, ns * nh, stream); ScaleVector(ptrVNew->ptr(), ptrVNew->ptr(), 1.0f/scale_factor, ns * nh, stream);
ncvAssertCUDALastErrorReturn(NCV_CUDA_ERROR); ncvAssertCUDALastErrorReturn((int)NCV_CUDA_ERROR);
cv::gpu::device::swap<FloatVector*>(ptrU, ptrUNew); cv::gpu::device::swap<FloatVector*>(ptrU, ptrUNew);
cv::gpu::device::swap<FloatVector*>(ptrV, ptrVNew); cv::gpu::device::swap<FloatVector*>(ptrV, ptrVNew);
@ -1145,17 +1145,17 @@ NCVStatus NCVBroxOpticalFlow(const NCVBroxOpticalFlowDescriptor desc,
} }
// end of warping iterations // end of warping iterations
ncvAssertCUDAReturn(cudaStreamSynchronize(stream), NCV_CUDA_ERROR); ncvAssertCUDAReturn(cudaStreamSynchronize(stream), (int)NCV_CUDA_ERROR);
ncvAssertCUDAReturn( cudaMemcpy2DAsync ncvAssertCUDAReturn( cudaMemcpy2DAsync
(uOut.ptr(), uOut.pitch(), ptrU->ptr(), (uOut.ptr(), uOut.pitch(), ptrU->ptr(),
kSourcePitch, kSourceWidth*sizeof(float), kSourceHeight, cudaMemcpyDeviceToDevice, stream), NCV_CUDA_ERROR ); kSourcePitch, kSourceWidth*sizeof(float), kSourceHeight, cudaMemcpyDeviceToDevice, stream), (int)NCV_CUDA_ERROR );
ncvAssertCUDAReturn( cudaMemcpy2DAsync ncvAssertCUDAReturn( cudaMemcpy2DAsync
(vOut.ptr(), vOut.pitch(), ptrV->ptr(), (vOut.ptr(), vOut.pitch(), ptrV->ptr(),
kSourcePitch, kSourceWidth*sizeof(float), kSourceHeight, cudaMemcpyDeviceToDevice, stream), NCV_CUDA_ERROR ); kSourcePitch, kSourceWidth*sizeof(float), kSourceHeight, cudaMemcpyDeviceToDevice, stream), (int)NCV_CUDA_ERROR );
ncvAssertCUDAReturn(cudaStreamSynchronize(stream), NCV_CUDA_ERROR); ncvAssertCUDAReturn(cudaStreamSynchronize(stream), (int)NCV_CUDA_ERROR);
} }
return NCV_SUCCESS; return NCV_SUCCESS;

3
modules/gpu/src/nvidia/NCVHaarObjectDetection.cu
Unescape View File

@ -687,6 +687,7 @@ struct applyHaarClassifierAnchorParallelFunctor
template<class TList> template<class TList>
void call(TList tl) void call(TList tl)
{ {
(void)tl;
applyHaarClassifierAnchorParallel < applyHaarClassifierAnchorParallel <
Loki::TL::TypeAt<TList, 0>::Result::value, Loki::TL::TypeAt<TList, 0>::Result::value,
Loki::TL::TypeAt<TList, 1>::Result::value, Loki::TL::TypeAt<TList, 1>::Result::value,
@ -796,6 +797,7 @@ struct applyHaarClassifierClassifierParallelFunctor
template<class TList> template<class TList>
void call(TList tl) void call(TList tl)
{ {
(void)tl;
applyHaarClassifierClassifierParallel < applyHaarClassifierClassifierParallel <
Loki::TL::TypeAt<TList, 0>::Result::value, Loki::TL::TypeAt<TList, 0>::Result::value,
Loki::TL::TypeAt<TList, 1>::Result::value, Loki::TL::TypeAt<TList, 1>::Result::value,
@ -876,6 +878,7 @@ struct initializeMaskVectorFunctor
template<class TList> template<class TList>
void call(TList tl) void call(TList tl)
{ {
(void)tl;
initializeMaskVector < initializeMaskVector <
Loki::TL::TypeAt<TList, 0>::Result::value, Loki::TL::TypeAt<TList, 0>::Result::value,
Loki::TL::TypeAt<TList, 1>::Result::value > Loki::TL::TypeAt<TList, 1>::Result::value >

1
modules/gpu/src/nvidia/core/NCV.cu
Unescape View File

@ -854,6 +854,7 @@ static NCVStatus drawRectsWrapperDevice(T *d_dst,
T color, T color,
cudaStream_t cuStream) cudaStream_t cuStream)
{ {
(void)cuStream;
ncvAssertReturn(d_dst != NULL && d_rects != NULL, NCV_NULL_PTR); ncvAssertReturn(d_dst != NULL && d_rects != NULL, NCV_NULL_PTR);
ncvAssertReturn(dstWidth > 0 && dstHeight > 0, NCV_DIMENSIONS_INVALID); ncvAssertReturn(dstWidth > 0 && dstHeight > 0, NCV_DIMENSIONS_INVALID);
ncvAssertReturn(dstStride >= dstWidth, NCV_INVALID_STEP); ncvAssertReturn(dstStride >= dstWidth, NCV_INVALID_STEP);

22
modules/gpu/src/nvidia/core/NCV.hpp
Unescape View File

@ -1,7 +1,7 @@
/*M/////////////////////////////////////////////////////////////////////////////////////// /*M///////////////////////////////////////////////////////////////////////////////////////
// //
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
// //
// By downloading, copying, installing or using the software you agree to this license. // 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, // If you do not agree to this license, do not download, install,
// copy or use the software. // copy or use the software.
@ -461,7 +461,7 @@ public:
virtual NcvBool isInitialized(void) const = 0; virtual NcvBool isInitialized(void) const = 0;
virtual NcvBool isCounting(void) const = 0; virtual NcvBool isCounting(void) const = 0;
virtual NCVMemoryType memType(void) const = 0; virtual NCVMemoryType memType(void) const = 0;
virtual Ncv32u alignment(void) const = 0; virtual Ncv32u alignment(void) const = 0;
virtual size_t maxSize(void) const = 0; virtual size_t maxSize(void) const = 0;
@ -585,11 +585,11 @@ public:
} }
else else
{ {
ncvAssertReturn(dst._length * sizeof(T) >= howMuch && ncvAssertReturn(dst._length * sizeof(T) >= howMuch &&
this->_length * sizeof(T) >= howMuch && this->_length * sizeof(T) >= howMuch &&
howMuch > 0, NCV_MEM_COPY_ERROR); howMuch > 0, NCV_MEM_COPY_ERROR);
} }
ncvAssertReturn((this->_ptr != NULL || this->_memtype == NCVMemoryTypeNone) && ncvAssertReturn((this->_ptr != NULL || this->_memtype == NCVMemoryTypeNone) &&
(dst._ptr != NULL || dst._memtype == NCVMemoryTypeNone), NCV_NULL_PTR); (dst._ptr != NULL || dst._memtype == NCVMemoryTypeNone), NCV_NULL_PTR);
NCVStatus ncvStat = NCV_SUCCESS; NCVStatus ncvStat = NCV_SUCCESS;
@ -766,18 +766,18 @@ public:
} }
else else
{ {
ncvAssertReturn(dst._pitch * dst._height >= howMuch && ncvAssertReturn(dst._pitch * dst._height >= howMuch &&
this->_pitch * this->_height >= howMuch && this->_pitch * this->_height >= howMuch &&
howMuch > 0, NCV_MEM_COPY_ERROR); howMuch > 0, NCV_MEM_COPY_ERROR);
} }
ncvAssertReturn((this->_ptr != NULL || this->_memtype == NCVMemoryTypeNone) && ncvAssertReturn((this->_ptr != NULL || this->_memtype == NCVMemoryTypeNone) &&
(dst._ptr != NULL || dst._memtype == NCVMemoryTypeNone), NCV_NULL_PTR); (dst._ptr != NULL || dst._memtype == NCVMemoryTypeNone), NCV_NULL_PTR);
NCVStatus ncvStat = NCV_SUCCESS; NCVStatus ncvStat = NCV_SUCCESS;
if (this->_memtype != NCVMemoryTypeNone) if (this->_memtype != NCVMemoryTypeNone)
{ {
ncvStat = memSegCopyHelper(dst._ptr, dst._memtype, ncvStat = memSegCopyHelper(dst._ptr, dst._memtype,
this->_ptr, this->_memtype, this->_ptr, this->_memtype,
howMuch, cuStream); howMuch, cuStream);
} }
@ -788,7 +788,7 @@ public:
{ {
ncvAssertReturn(this->width() >= roi.width && this->height() >= roi.height && ncvAssertReturn(this->width() >= roi.width && this->height() >= roi.height &&
dst.width() >= roi.width && dst.height() >= roi.height, NCV_MEM_COPY_ERROR); dst.width() >= roi.width && dst.height() >= roi.height, NCV_MEM_COPY_ERROR);
ncvAssertReturn((this->_ptr != NULL || this->_memtype == NCVMemoryTypeNone) && ncvAssertReturn((this->_ptr != NULL || this->_memtype == NCVMemoryTypeNone) &&
(dst._ptr != NULL || dst._memtype == NCVMemoryTypeNone), NCV_NULL_PTR); (dst._ptr != NULL || dst._memtype == NCVMemoryTypeNone), NCV_NULL_PTR);
NCVStatus ncvStat = NCV_SUCCESS; NCVStatus ncvStat = NCV_SUCCESS;
@ -802,7 +802,7 @@ public:
return ncvStat; return ncvStat;
} }
T &at(Ncv32u x, Ncv32u y) const T& at(Ncv32u x, Ncv32u y) const
{ {
NcvBool bOutRange = (x >= this->_width || y >= this->_height); NcvBool bOutRange = (x >= this->_width || y >= this->_height);
ncvAssertPrintCheck(!bOutRange, "Error addressing matrix at [" << x << ", " << y << "]"); ncvAssertPrintCheck(!bOutRange, "Error addressing matrix at [" << x << ", " << y << "]");

1
modules/gpu/src/nvidia/core/NCVRuntimeTemplates.hpp
Unescape View File

@ -211,6 +211,7 @@ namespace NCVRuntimeTemplateBool
static void call(Func &functor, std::vector<int> &templateParams) static void call(Func &functor, std::vector<int> &templateParams)
{ {
(void)templateParams;
functor.call(TList()); functor.call(TList());
} }
}; };

22
modules/gpu/src/opencv2/gpu/device/filters.hpp
Unescape View File

@ -55,7 +55,12 @@ namespace cv { namespace gpu { namespace device
typedef typename Ptr2D::elem_type elem_type; typedef typename Ptr2D::elem_type elem_type;
typedef float index_type; typedef float index_type;
explicit __host__ __device__ __forceinline__ PointFilter(const Ptr2D& src_, float fx = 0.f, float fy = 0.f) : src(src_) {} explicit __host__ __device__ __forceinline__ PointFilter(const Ptr2D& src_, float fx = 0.f, float fy = 0.f)
: src(src_)
{
(void)fx;
(void)fy;
}
__device__ __forceinline__ elem_type operator ()(float y, float x) const __device__ __forceinline__ elem_type operator ()(float y, float x) const
{ {
@ -70,8 +75,12 @@ namespace cv { namespace gpu { namespace device
typedef typename Ptr2D::elem_type elem_type; typedef typename Ptr2D::elem_type elem_type;
typedef float index_type; typedef float index_type;
explicit __host__ __device__ __forceinline__ LinearFilter(const Ptr2D& src_, float fx = 0.f, float fy = 0.f) : src(src_) {} explicit __host__ __device__ __forceinline__ LinearFilter(const Ptr2D& src_, float fx = 0.f, float fy = 0.f)
: src(src_)
{
(void)fx;
(void)fy;
}
__device__ __forceinline__ elem_type operator ()(float y, float x) const __device__ __forceinline__ elem_type operator ()(float y, float x) const
{ {
typedef typename TypeVec<float, VecTraits<elem_type>::cn>::vec_type work_type; typedef typename TypeVec<float, VecTraits<elem_type>::cn>::vec_type work_type;
@ -107,7 +116,12 @@ namespace cv { namespace gpu { namespace device
typedef float index_type; typedef float index_type;
typedef typename TypeVec<float, VecTraits<elem_type>::cn>::vec_type work_type; typedef typename TypeVec<float, VecTraits<elem_type>::cn>::vec_type work_type;
explicit __host__ __device__ __forceinline__ CubicFilter(const Ptr2D& src_, float fx = 0.f, float fy = 0.f) : src(src_) {} explicit __host__ __device__ __forceinline__ CubicFilter(const Ptr2D& src_, float fx = 0.f, float fy = 0.f)
: src(src_)
{
(void)fx;
(void)fy;
}
static __device__ __forceinline__ float bicubicCoeff(float x_) static __device__ __forceinline__ float bicubicCoeff(float x_)
{ {

6
modules/gpu/src/opencv2/gpu/device/functional.hpp
Unescape View File

@ -470,7 +470,7 @@ namespace cv { namespace gpu { namespace device
template <typename T> struct thresh_trunc_func : unary_function<T, T> 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_) {} explicit __host__ __device__ __forceinline__ thresh_trunc_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {(void)maxVal_;}
__device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const
{ {
@ -487,7 +487,7 @@ namespace cv { namespace gpu { namespace device
template <typename T> struct thresh_to_zero_func : unary_function<T, T> 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_) {} explicit __host__ __device__ __forceinline__ thresh_to_zero_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {(void)maxVal_;}
__device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const
{ {
@ -503,7 +503,7 @@ namespace cv { namespace gpu { namespace device
template <typename T> struct thresh_to_zero_inv_func : unary_function<T, T> 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_) {} explicit __host__ __device__ __forceinline__ thresh_to_zero_inv_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {(void)maxVal_;}
__device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const
{ {

34
modules/gpu/test/nvidia/TestHypothesesFilter.cpp
Unescape View File

@ -1,11 +1,11 @@
/* /*
* Copyright 1993-2010 NVIDIA Corporation. All rights reserved. * Copyright 1993-2010 NVIDIA Corporation. All rights reserved.
* *
* NVIDIA Corporation and its licensors retain all intellectual * NVIDIA Corporation and its licensors retain all intellectual
* property and proprietary rights in and to this software and * property and proprietary rights in and to this software and
* related documentation and any modifications thereto. * related documentation and any modifications thereto.
* Any use, reproduction, disclosure, or distribution of this * Any use, reproduction, disclosure, or distribution of this
* software and related documentation without an express license * software and related documentation without an express license
* agreement from NVIDIA Corporation is strictly prohibited. * agreement from NVIDIA Corporation is strictly prohibited.
*/ */
@ -13,14 +13,14 @@
#include "NCVHaarObjectDetection.hpp" #include "NCVHaarObjectDetection.hpp"
TestHypothesesFilter::TestHypothesesFilter(std::string testName, NCVTestSourceProvider<Ncv32u> &src, TestHypothesesFilter::TestHypothesesFilter(std::string testName, NCVTestSourceProvider<Ncv32u> &src_,
Ncv32u numDstRects, Ncv32u minNeighbors, Ncv32f eps) Ncv32u numDstRects_, Ncv32u minNeighbors_, Ncv32f eps_)
: :
NCVTestProvider(testName), NCVTestProvider(testName),
src(src), src(src_),
numDstRects(numDstRects), numDstRects(numDstRects_),
minNeighbors(minNeighbors), minNeighbors(minNeighbors_),
eps(eps) eps(eps_)
{ {
} }
@ -94,11 +94,11 @@ bool TestHypothesesFilter::process()
for (Ncv32u j=0; j<numNeighbors; j++) for (Ncv32u j=0; j<numNeighbors; j++)
{ {
randVal = (1.0 * h_random32u.ptr()[randCnt++]) / 0xFFFFFFFF; randCnt = randCnt % h_random32u.length(); randVal = (1.0 * h_random32u.ptr()[randCnt++]) / 0xFFFFFFFF; randCnt = randCnt % h_random32u.length();
h_vecSrc.ptr()[srcSlotSize * i + j].x = h_vecSrc.ptr()[srcSlotSize * i + j].x =
h_vecDst_groundTruth.ptr()[i].x + h_vecDst_groundTruth.ptr()[i].x +
(Ncv32s)(h_vecDst_groundTruth.ptr()[i].width * this->eps * (randVal - 0.5)); (Ncv32s)(h_vecDst_groundTruth.ptr()[i].width * this->eps * (randVal - 0.5));
randVal = (1.0 * h_random32u.ptr()[randCnt++]) / 0xFFFFFFFF; randCnt = randCnt % h_random32u.length(); randVal = (1.0 * h_random32u.ptr()[randCnt++]) / 0xFFFFFFFF; randCnt = randCnt % h_random32u.length();
h_vecSrc.ptr()[srcSlotSize * i + j].y = h_vecSrc.ptr()[srcSlotSize * i + j].y =
h_vecDst_groundTruth.ptr()[i].y + h_vecDst_groundTruth.ptr()[i].y +
(Ncv32s)(h_vecDst_groundTruth.ptr()[i].height * this->eps * (randVal - 0.5)); (Ncv32s)(h_vecDst_groundTruth.ptr()[i].height * this->eps * (randVal - 0.5));
h_vecSrc.ptr()[srcSlotSize * i + j].width = h_vecDst_groundTruth.ptr()[i].width; h_vecSrc.ptr()[srcSlotSize * i + j].width = h_vecDst_groundTruth.ptr()[i].width;
@ -109,11 +109,11 @@ bool TestHypothesesFilter::process()
for (Ncv32u j=numNeighbors; j<srcSlotSize; j++) for (Ncv32u j=numNeighbors; j<srcSlotSize; j++)
{ {
randVal = (1.0 * h_random32u.ptr()[randCnt++]) / 0xFFFFFFFF; randCnt = randCnt % h_random32u.length(); randVal = (1.0 * h_random32u.ptr()[randCnt++]) / 0xFFFFFFFF; randCnt = randCnt % h_random32u.length();
h_vecSrc.ptr()[srcSlotSize * i + j].x = h_vecSrc.ptr()[srcSlotSize * i + j].x =
this->canvasWidth + h_vecDst_groundTruth.ptr()[i].x + this->canvasWidth + h_vecDst_groundTruth.ptr()[i].x +
(Ncv32s)(h_vecDst_groundTruth.ptr()[i].width * this->eps * (randVal - 0.5)); (Ncv32s)(h_vecDst_groundTruth.ptr()[i].width * this->eps * (randVal - 0.5));
randVal = (1.0 * h_random32u.ptr()[randCnt++]) / 0xFFFFFFFF; randCnt = randCnt % h_random32u.length(); randVal = (1.0 * h_random32u.ptr()[randCnt++]) / 0xFFFFFFFF; randCnt = randCnt % h_random32u.length();
h_vecSrc.ptr()[srcSlotSize * i + j].y = h_vecSrc.ptr()[srcSlotSize * i + j].y =
this->canvasHeight + h_vecDst_groundTruth.ptr()[i].y + this->canvasHeight + h_vecDst_groundTruth.ptr()[i].y +
(Ncv32s)(h_vecDst_groundTruth.ptr()[i].height * this->eps * (randVal - 0.5)); (Ncv32s)(h_vecDst_groundTruth.ptr()[i].height * this->eps * (randVal - 0.5));
h_vecSrc.ptr()[srcSlotSize * i + j].width = h_vecDst_groundTruth.ptr()[i].width; h_vecSrc.ptr()[srcSlotSize * i + j].width = h_vecDst_groundTruth.ptr()[i].width;
@ -124,8 +124,8 @@ bool TestHypothesesFilter::process()
//shuffle //shuffle
for (Ncv32u i=0; i<this->numDstRects*srcSlotSize-1; i++) for (Ncv32u i=0; i<this->numDstRects*srcSlotSize-1; i++)
{ {
Ncv32u randVal = h_random32u.ptr()[randCnt++]; randCnt = randCnt % h_random32u.length(); Ncv32u randValLocal = h_random32u.ptr()[randCnt++]; randCnt = randCnt % h_random32u.length();
Ncv32u secondSwap = randVal % (this->numDstRects*srcSlotSize-1 - i); Ncv32u secondSwap = randValLocal % (this->numDstRects*srcSlotSize-1 - i);
NcvRect32u tmp = h_vecSrc.ptr()[i + secondSwap]; NcvRect32u tmp = h_vecSrc.ptr()[i + secondSwap];
h_vecSrc.ptr()[i + secondSwap] = h_vecSrc.ptr()[i]; h_vecSrc.ptr()[i + secondSwap] = h_vecSrc.ptr()[i];
h_vecSrc.ptr()[i] = tmp; h_vecSrc.ptr()[i] = tmp;

24
modules/gpu/test/nvidia/TestResize.cpp
Unescape View File

@ -1,11 +1,11 @@
/* /*
* Copyright 1993-2010 NVIDIA Corporation. All rights reserved. * Copyright 1993-2010 NVIDIA Corporation. All rights reserved.
* *
* NVIDIA Corporation and its licensors retain all intellectual * NVIDIA Corporation and its licensors retain all intellectual
* property and proprietary rights in and to this software and * property and proprietary rights in and to this software and
* related documentation and any modifications thereto. * related documentation and any modifications thereto.
* Any use, reproduction, disclosure, or distribution of this * Any use, reproduction, disclosure, or distribution of this
* software and related documentation without an express license * software and related documentation without an express license
* agreement from NVIDIA Corporation is strictly prohibited. * agreement from NVIDIA Corporation is strictly prohibited.
*/ */
@ -15,15 +15,15 @@
template <class T> template <class T>
TestResize<T>::TestResize(std::string testName, NCVTestSourceProvider<T> &src, TestResize<T>::TestResize(std::string testName, NCVTestSourceProvider<T> &src_,
Ncv32u width, Ncv32u height, Ncv32u scaleFactor, NcvBool bTextureCache) Ncv32u width_, Ncv32u height_, Ncv32u scaleFactor_, NcvBool bTextureCache_)
: :
NCVTestProvider(testName), NCVTestProvider(testName),
src(src), src(src_),
width(width), width(width_),
height(height), height(height_),
scaleFactor(scaleFactor), scaleFactor(scaleFactor_),
bTextureCache(bTextureCache) bTextureCache(bTextureCache_)
{ {
} }

5
modules/gpu/test/nvidia/main_nvidia.cpp
Unescape View File

@ -248,6 +248,7 @@ void generateHaarLoaderTests(NCVAutoTestLister &testLister)
void generateHaarApplicationTests(NCVAutoTestLister &testLister, NCVTestSourceProvider<Ncv8u> &src, void generateHaarApplicationTests(NCVAutoTestLister &testLister, NCVTestSourceProvider<Ncv8u> &src,
Ncv32u maxWidth, Ncv32u maxHeight) Ncv32u maxWidth, Ncv32u maxHeight)
{ {
(void)maxHeight;
for (Ncv32u i=20; i<512; i+=11) for (Ncv32u i=20; i<512; i+=11)
{ {
for (Ncv32u j=20; j<128; j+=5) for (Ncv32u j=20; j<128; j+=5)
@ -268,11 +269,12 @@ void generateHaarApplicationTests(NCVAutoTestLister &testLister, NCVTestSourcePr
static void devNullOutput(const std::string& msg) static void devNullOutput(const std::string& msg)
{ {
(void)msg;
} }
bool nvidia_NPPST_Integral_Image(const std::string& test_data_path, OutputLevel outputLevel) bool nvidia_NPPST_Integral_Image(const std::string& test_data_path, OutputLevel outputLevel)
{ {
path = test_data_path; path = test_data_path.c_str();
ncvSetDebugOutputHandler(devNullOutput); ncvSetDebugOutputHandler(devNullOutput);
NCVAutoTestLister testListerII("NPPST Integral Image", outputLevel); NCVAutoTestLister testListerII("NPPST Integral Image", outputLevel);
@ -374,6 +376,7 @@ bool nvidia_NCV_Vector_Operations(const std::string& test_data_path, OutputLevel
generateVectorTests(testListerVectorOperations, testSrcRandom_32u, 4096*4096); generateVectorTests(testListerVectorOperations, testSrcRandom_32u, 4096*4096);
return testListerVectorOperations.invoke(); return testListerVectorOperations.invoke();
} }
bool nvidia_NCV_Haar_Cascade_Loader(const std::string& test_data_path, OutputLevel outputLevel) bool nvidia_NCV_Haar_Cascade_Loader(const std::string& test_data_path, OutputLevel outputLevel)

26
modules/gpu/test/test_nvidia.cpp
Unescape View File

@ -58,15 +58,15 @@ struct NVidiaTest : TestWithParam<cv::gpu::DeviceInfo>
{ {
cv::gpu::DeviceInfo devInfo; cv::gpu::DeviceInfo devInfo;
std::string path; std::string _path;
virtual void SetUp() virtual void SetUp()
{ {
devInfo = GetParam(); devInfo = GetParam();
cv::gpu::setDevice(devInfo.deviceID()); cv::gpu::setDevice(devInfo.deviceID());
_path = TS::ptr()->get_data_path().c_str();
path = std::string(TS::ptr()->get_data_path()) + "haarcascade/"; _path = _path + "haarcascade/";
} }
}; };
@ -84,63 +84,63 @@ OutputLevel nvidiaTestOutputLevel = OutputLevelCompact;
TEST_P(NPPST, SquaredIntegral) TEST_P(NPPST, SquaredIntegral)
{ {
bool res = nvidia_NPPST_Squared_Integral_Image(path, nvidiaTestOutputLevel); bool res = nvidia_NPPST_Squared_Integral_Image(_path, nvidiaTestOutputLevel);
ASSERT_TRUE(res); ASSERT_TRUE(res);
} }
TEST_P(NPPST, RectStdDev) TEST_P(NPPST, RectStdDev)
{ {
bool res = nvidia_NPPST_RectStdDev(path, nvidiaTestOutputLevel); bool res = nvidia_NPPST_RectStdDev(_path, nvidiaTestOutputLevel);
ASSERT_TRUE(res); ASSERT_TRUE(res);
} }
TEST_P(NPPST, Resize) TEST_P(NPPST, Resize)
{ {
bool res = nvidia_NPPST_Resize(path, nvidiaTestOutputLevel); bool res = nvidia_NPPST_Resize(_path, nvidiaTestOutputLevel);
ASSERT_TRUE(res); ASSERT_TRUE(res);
} }
TEST_P(NPPST, VectorOperations) TEST_P(NPPST, VectorOperations)
{ {
bool res = nvidia_NPPST_Vector_Operations(path, nvidiaTestOutputLevel); bool res = nvidia_NPPST_Vector_Operations(_path, nvidiaTestOutputLevel);
ASSERT_TRUE(res); ASSERT_TRUE(res);
} }
TEST_P(NPPST, Transpose) TEST_P(NPPST, Transpose)
{ {
bool res = nvidia_NPPST_Transpose(path, nvidiaTestOutputLevel); bool res = nvidia_NPPST_Transpose(_path, nvidiaTestOutputLevel);
ASSERT_TRUE(res); ASSERT_TRUE(res);
} }
TEST_P(NCV, VectorOperations) TEST_P(NCV, VectorOperations)
{ {
bool res = nvidia_NCV_Vector_Operations(path, nvidiaTestOutputLevel); bool res = nvidia_NCV_Vector_Operations(_path, nvidiaTestOutputLevel);
ASSERT_TRUE(res); ASSERT_TRUE(res);
} }
TEST_P(NCV, HaarCascadeLoader) TEST_P(NCV, HaarCascadeLoader)
{ {
bool res = nvidia_NCV_Haar_Cascade_Loader(path, nvidiaTestOutputLevel); bool res = nvidia_NCV_Haar_Cascade_Loader(_path, nvidiaTestOutputLevel);
ASSERT_TRUE(res); ASSERT_TRUE(res);
} }
TEST_P(NCV, HaarCascadeApplication) TEST_P(NCV, HaarCascadeApplication)
{ {
bool res = nvidia_NCV_Haar_Cascade_Application(path, nvidiaTestOutputLevel); bool res = nvidia_NCV_Haar_Cascade_Application(_path, nvidiaTestOutputLevel);
ASSERT_TRUE(res); ASSERT_TRUE(res);
} }
TEST_P(NCV, HypothesesFiltration) TEST_P(NCV, HypothesesFiltration)
{ {
bool res = nvidia_NCV_Hypotheses_Filtration(path, nvidiaTestOutputLevel); bool res = nvidia_NCV_Hypotheses_Filtration(_path, nvidiaTestOutputLevel);
ASSERT_TRUE(res); ASSERT_TRUE(res);
} }
@ -148,7 +148,7 @@ TEST_P(NCV, HypothesesFiltration)
TEST_P(NCV, Visualization) TEST_P(NCV, Visualization)
{ {
// this functionality doesn't used in gpu module // this functionality doesn't used in gpu module
bool res = nvidia_NCV_Visualization(path, nvidiaTestOutputLevel); bool res = nvidia_NCV_Visualization(_path, nvidiaTestOutputLevel);
ASSERT_TRUE(res); ASSERT_TRUE(res);
} }

Write Preview
Loading…
Cancel
Save