← Back to branch summary

~paparazzi-uav/paparazzi/v5.0-manual

« back to all changes in this revision

Viewing changes to sw/ext/opencv_bebop/opencv/modules/cudaarithm/src/cuda/threshold.cu

  • Committer: Paparazzi buildbot
  • Date: 2016-05-18 15:00:29 UTC
  • Revision ID: felix.ruess+docbot@gmail.com-20160518150029-e8lgzi5kvb4p7un9
Manual import commit 4b8bbb730080dac23cf816b98908dacfabe2a8ec from v5.0 branch.

Show diffs side-by-side

added added

removed removed

Lines of Context:
sw/ext/opencv_bebop/opencv/modules/cudaarithm/src/cuda/threshold.cu
 
1
/*M///////////////////////////////////////////////////////////////////////////////////////
 
2
//
 
3
//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
 
4
//
 
5
//  By downloading, copying, installing or using the software you agree to this license.
 
6
//  If you do not agree to this license, do not download, install,
 
7
//  copy or use the software.
 
8
//
 
9
//
 
10
//                           License Agreement
 
11
//                For Open Source Computer Vision Library
 
12
//
 
13
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
 
14
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
 
15
// Third party copyrights are property of their respective owners.
 
16
//
 
17
// Redistribution and use in source and binary forms, with or without modification,
 
18
// are permitted provided that the following conditions are met:
 
19
//
 
20
//   * Redistribution's of source code must retain the above copyright notice,
 
21
//     this list of conditions and the following disclaimer.
 
22
//
 
23
//   * Redistribution's in binary form must reproduce the above copyright notice,
 
24
//     this list of conditions and the following disclaimer in the documentation
 
25
//     and/or other materials provided with the distribution.
 
26
//
 
27
//   * The name of the copyright holders may not be used to endorse or promote products
 
28
//     derived from this software without specific prior written permission.
 
29
//
 
30
// This software is provided by the copyright holders and contributors "as is" and
 
31
// any express or implied warranties, including, but not limited to, the implied
 
32
// warranties of merchantability and fitness for a particular purpose are disclaimed.
 
33
// In no event shall the Intel Corporation or contributors be liable for any direct,
 
34
// indirect, incidental, special, exemplary, or consequential damages
 
35
// (including, but not limited to, procurement of substitute goods or services;
 
36
// loss of use, data, or profits; or business interruption) however caused
 
37
// and on any theory of liability, whether in contract, strict liability,
 
38
// or tort (including negligence or otherwise) arising in any way out of
 
39
// the use of this software, even if advised of the possibility of such damage.
 
40
//
 
41
//M*/
 
42
 
 
43
#include "opencv2/opencv_modules.hpp"
 
44
 
 
45
#ifndef HAVE_OPENCV_CUDEV
 
46
 
 
47
#error "opencv_cudev is required"
 
48
 
 
49
#else
 
50
 
 
51
#include "opencv2/cudaarithm.hpp"
 
52
#include "opencv2/cudev.hpp"
 
53
#include "opencv2/core/private.cuda.hpp"
 
54
 
 
55
using namespace cv;
 
56
using namespace cv::cuda;
 
57
using namespace cv::cudev;
 
58
 
 
59
namespace
 
60
{
 
61
    template <typename ScalarDepth> struct TransformPolicy : DefaultTransformPolicy
 
62
    {
 
63
    };
 
64
    template <> struct TransformPolicy<double> : DefaultTransformPolicy
 
65
    {
 
66
        enum {
 
67
            shift = 1
 
68
        };
 
69
    };
 
70
 
 
71
    template <typename T>
 
72
    void thresholdImpl(const GpuMat& src, GpuMat& dst, double thresh, double maxVal, int type, Stream& stream)
 
73
    {
 
74
        const T thresh_ = static_cast<T>(thresh);
 
75
        const T maxVal_ = static_cast<T>(maxVal);
 
76
 
 
77
        switch (type)
 
78
        {
 
79
        case 0:
 
80
            gridTransformUnary_< TransformPolicy<T> >(globPtr<T>(src), globPtr<T>(dst), thresh_binary_func(thresh_, maxVal_), stream);
 
81
            break;
 
82
        case 1:
 
83
            gridTransformUnary_< TransformPolicy<T> >(globPtr<T>(src), globPtr<T>(dst), thresh_binary_inv_func(thresh_, maxVal_), stream);
 
84
            break;
 
85
        case 2:
 
86
            gridTransformUnary_< TransformPolicy<T> >(globPtr<T>(src), globPtr<T>(dst), thresh_trunc_func(thresh_), stream);
 
87
            break;
 
88
        case 3:
 
89
            gridTransformUnary_< TransformPolicy<T> >(globPtr<T>(src), globPtr<T>(dst), thresh_to_zero_func(thresh_), stream);
 
90
            break;
 
91
        case 4:
 
92
            gridTransformUnary_< TransformPolicy<T> >(globPtr<T>(src), globPtr<T>(dst), thresh_to_zero_inv_func(thresh_), stream);
 
93
            break;
 
94
        };
 
95
    }
 
96
}
 
97
 
 
98
double cv::cuda::threshold(InputArray _src, OutputArray _dst, double thresh, double maxVal, int type, Stream& stream)
 
99
{
 
100
    GpuMat src = getInputMat(_src, stream);
 
101
 
 
102
    const int depth = src.depth();
 
103
 
 
104
    CV_Assert( src.channels() == 1 && depth <= CV_64F );
 
105
    CV_Assert( type <= 4 /*THRESH_TOZERO_INV*/ );
 
106
 
 
107
    GpuMat dst = getOutputMat(_dst, src.size(), src.type(), stream);
 
108
 
 
109
    if (depth == CV_32F && type == 2 /*THRESH_TRUNC*/)
 
110
    {
 
111
        NppStreamHandler h(StreamAccessor::getStream(stream));
 
112
 
 
113
        NppiSize sz;
 
114
        sz.width  = src.cols;
 
115
        sz.height = src.rows;
 
116
 
 
117
        nppSafeCall( nppiThreshold_32f_C1R(src.ptr<Npp32f>(), static_cast<int>(src.step),
 
118
            dst.ptr<Npp32f>(), static_cast<int>(dst.step), sz, static_cast<Npp32f>(thresh), NPP_CMP_GREATER) );
 
119
 
 
120
        if (!stream)
 
121
            CV_CUDEV_SAFE_CALL( cudaDeviceSynchronize() );
 
122
    }
 
123
    else
 
124
    {
 
125
        typedef void (*func_t)(const GpuMat& src, GpuMat& dst, double thresh, double maxVal, int type, Stream& stream);
 
126
        static const func_t funcs[] =
 
127
        {
 
128
            thresholdImpl<uchar>,
 
129
            thresholdImpl<schar>,
 
130
            thresholdImpl<ushort>,
 
131
            thresholdImpl<short>,
 
132
            thresholdImpl<int>,
 
133
            thresholdImpl<float>,
 
134
            thresholdImpl<double>
 
135
        };
 
136
 
 
137
        if (depth != CV_32F && depth != CV_64F)
 
138
        {
 
139
            thresh = cvFloor(thresh);
 
140
            maxVal = cvRound(maxVal);
 
141
        }
 
142
 
 
143
        funcs[depth](src, dst, thresh, maxVal, type, stream);
 
144
    }
 
145
 
 
146
    syncOutput(dst, _dst, stream);
 
147
 
 
148
    return thresh;
 
149
}
 
150
 
 
151
#endif