提交 6aa7a86a 编写于 作者: M Maksim Shabunin

Doxygen documentation for core module

上级 bea86dee
......@@ -147,14 +147,18 @@ endif()
# ========= Doxygen docs =========
if(BUILD_DOCS AND HAVE_DOXYGEN)
# documented modules list
set(candidates)
set(all_headers)
set(all_images)
list(APPEND candidates ${BASE_MODULES} ${EXTRA_MODULES})
# blacklisted modules
ocv_list_filterout(candidates "^ts$")
# gathering headers
set(all_headers) # files and dirs to process
set(all_images) # image search paths
set(reflist) # modules reference
foreach(m ${candidates})
set(reflist "${reflist} \n- @subpage ${m}")
set(all_headers ${all_headers} "${OPENCV_MODULE_opencv_${m}_HEADERS}")
set(docs_dir "${OPENCV_MODULE_opencv_${m}_LOCATION}/doc")
if(EXISTS ${docs_dir})
......@@ -164,15 +168,20 @@ if(BUILD_DOCS AND HAVE_DOXYGEN)
endforeach()
# additional config
set(doxyfile "${CMAKE_CURRENT_BINARY_DIR}/Doxyfile")
set(rootfile "${CMAKE_CURRENT_BINARY_DIR}/root.markdown")
set(all_headers ${all_headers} ${rootfile})
string(REGEX REPLACE ";" " \\\\\\n" CMAKE_DOXYGEN_INPUT_LIST "${all_headers}")
string(REGEX REPLACE ";" " \\\\\\n" CMAKE_DOXYGEN_IMAGE_PATH "${all_images}")
set(CMAKE_DOXYGEN_INDEX_MD "${CMAKE_SOURCE_DIR}/README.md")
set(CMAKE_DOXYGEN_LAYOUT "${CMAKE_CURRENT_SOURCE_DIR}/DoxygenLayout.xml")
set(CMAKE_DOXYGEN_OUTPUT_PATH "doxygen")
set(CMAKE_DOXYGEN_MODULES_REFERENCE "${reflist}")
set(CMAKE_DOXYGEN_EXAMPLE_PATH "${CMAKE_SOURCE_DIR}/samples/cpp")
# writing file
set(doxyfile "${CMAKE_CURRENT_BINARY_DIR}/Doxyfile")
configure_file(Doxyfile.in ${doxyfile} @ONLY)
configure_file(root.markdown.in ${rootfile} @ONLY)
configure_file(mymath.sty "${CMAKE_DOXYGEN_OUTPUT_PATH}/html/mymath.sty" @ONLY)
add_custom_target(doxygen
COMMAND ${DOXYGEN_BUILD} ${doxyfile}
......
......@@ -21,8 +21,8 @@ ABBREVIATE_BRIEF = "The $name class" \
the
ALWAYS_DETAILED_SEC = NO
INLINE_INHERITED_MEMB = NO
FULL_PATH_NAMES = NO
STRIP_FROM_PATH =
FULL_PATH_NAMES = YES
STRIP_FROM_PATH = @CMAKE_SOURCE_DIR@/modules
STRIP_FROM_INC_PATH =
SHORT_NAMES = NO
JAVADOC_AUTOBRIEF = NO
......@@ -39,7 +39,7 @@ OPTIMIZE_FOR_FORTRAN = NO
OPTIMIZE_OUTPUT_VHDL = NO
EXTENSION_MAPPING =
MARKDOWN_SUPPORT = YES
AUTOLINK_SUPPORT = NO
AUTOLINK_SUPPORT = YES
BUILTIN_STL_SUPPORT = YES
CPP_CLI_SUPPORT = NO
SIP_SUPPORT = NO
......@@ -53,7 +53,7 @@ LOOKUP_CACHE_SIZE = 0
EXTRACT_ALL = YES
EXTRACT_PRIVATE = NO
EXTRACT_PACKAGE = NO
EXTRACT_STATIC = NO
EXTRACT_STATIC = YES
EXTRACT_LOCAL_CLASSES = NO
EXTRACT_LOCAL_METHODS = NO
EXTRACT_ANON_NSPACES = NO
......@@ -66,11 +66,11 @@ CASE_SENSE_NAMES = YES
HIDE_SCOPE_NAMES = NO
SHOW_INCLUDE_FILES = YES
SHOW_GROUPED_MEMB_INC = NO
FORCE_LOCAL_INCLUDES = NO
FORCE_LOCAL_INCLUDES = YES
INLINE_INFO = YES
SORT_MEMBER_DOCS = YES
SORT_BRIEF_DOCS = NO
SORT_MEMBERS_CTORS_1ST = NO
SORT_BRIEF_DOCS = YES
SORT_MEMBERS_CTORS_1ST = YES
SORT_GROUP_NAMES = NO
SORT_BY_SCOPE_NAME = NO
STRICT_PROTO_MATCHING = NO
......@@ -86,7 +86,7 @@ SHOW_NAMESPACES = YES
FILE_VERSION_FILTER =
LAYOUT_FILE = @CMAKE_DOXYGEN_LAYOUT@
CITE_BIB_FILES =
QUIET = NO
QUIET = YES
WARNINGS = YES
WARN_IF_UNDOCUMENTED = YES
WARN_IF_DOC_ERROR = YES
......@@ -100,19 +100,16 @@ RECURSIVE = YES
EXCLUDE =
EXCLUDE_SYMLINKS = NO
EXCLUDE_PATTERNS =
EXCLUDE_SYMBOLS = CV_WRAP \
CV_EXPORTS \
CV_EXPORTS_W \
CV_WRAP_AS
EXAMPLE_PATH =
EXCLUDE_SYMBOLS = cv::DataType<*>
EXAMPLE_PATH = @CMAKE_DOXYGEN_EXAMPLE_PATH@
EXAMPLE_PATTERNS = *
EXAMPLE_RECURSIVE = NO
EXAMPLE_RECURSIVE = YES
IMAGE_PATH = @CMAKE_DOXYGEN_IMAGE_PATH@
INPUT_FILTER =
FILTER_PATTERNS =
FILTER_SOURCE_FILES = NO
FILTER_SOURCE_PATTERNS =
USE_MDFILE_AS_MAINPAGE = @CMAKE_DOXYGEN_INDEX_MD@
USE_MDFILE_AS_MAINPAGE =
SOURCE_BROWSER = NO
INLINE_SOURCES = NO
STRIP_CODE_COMMENTS = YES
......@@ -161,18 +158,18 @@ QHP_SECT_FILTER_ATTRS =
QHG_LOCATION =
GENERATE_ECLIPSEHELP = NO
ECLIPSE_DOC_ID = org.doxygen.Project
DISABLE_INDEX = YES
DISABLE_INDEX = NO
GENERATE_TREEVIEW = YES
ENUM_VALUES_PER_LINE = 4
ENUM_VALUES_PER_LINE = 0
TREEVIEW_WIDTH = 250
EXT_LINKS_IN_WINDOW = YES
FORMULA_FONTSIZE = 10
FORMULA_FONTSIZE = 14
FORMULA_TRANSPARENT = YES
USE_MATHJAX = NO
USE_MATHJAX = YES
MATHJAX_FORMAT = HTML-CSS
MATHJAX_RELPATH = http://cdn.mathjax.org/mathjax/latest
MATHJAX_EXTENSIONS =
MATHJAX_CODEFILE =
MATHJAX_EXTENSIONS = TeX/AMSmath TeX/AMSsymbols
MATHJAX_CODEFILE = @CMAKE_CURRENT_SOURCE_DIR@/mymath.js
SEARCHENGINE = YES
SERVER_BASED_SEARCH = NO
EXTERNAL_SEARCH = NO
......@@ -180,13 +177,13 @@ SEARCHENGINE_URL =
SEARCHDATA_FILE = searchdata.xml
EXTERNAL_SEARCH_ID =
EXTRA_SEARCH_MAPPINGS =
GENERATE_LATEX = YES
GENERATE_LATEX = NO
LATEX_OUTPUT = latex
LATEX_CMD_NAME = latex
MAKEINDEX_CMD_NAME = makeindex
COMPACT_LATEX = NO
PAPER_TYPE = a4
EXTRA_PACKAGES =
EXTRA_PACKAGES = mymath
LATEX_HEADER =
LATEX_FOOTER =
LATEX_EXTRA_FILES =
......@@ -222,12 +219,29 @@ EXPAND_ONLY_PREDEF = NO
SEARCH_INCLUDES = YES
INCLUDE_PATH =
INCLUDE_FILE_PATTERNS =
PREDEFINED = CV_WRAP= \
__cplusplus=1 \
PREDEFINED = __cplusplus=1 \
HAVE_IPP_A=1 \
CVAPI(x)=x \
CV_PROP_RW= \
CV_EXPORTS= \
CV_EXPORTS_W=
CV_EXPORTS_W= \
CV_EXPORTS_W_SIMPLE= \
CV_EXPORTS_AS(x)= \
CV_EXPORTS_W_MAP= \
CV_IN_OUT= \
CV_OUT= \
CV_PROP= \
CV_PROP_RW= \
CV_WRAP= \
CV_WRAP_AS(x)= \
CV_CDECL= \
CV_Func = \
CV_DO_PRAGMA(x)= \
CV_SUPPRESS_DEPRECATED_START= \
CV_SUPPRESS_DEPRECATED_END= \
CV_INLINE= \
CV_NORETURN= \
CV_DEFAULT(x)=" = x" \
CV_NEON=1
EXPAND_AS_DEFINED =
SKIP_FUNCTION_MACROS = YES
TAGFILES =
......
MathJax.Hub.Config({
TeX: {
Macros: {
matTT: [ "\\[ \\left|\\begin{array}{ccc} #1 & #2 & #3\\\\ #4 & #5 & #6\\\\ #7 & #8 & #9 \\end{array}\\right| \\]", 9],
fork: ["\\left\\{ \\begin{array}{l l} #1 & \\mbox{#2}\\\\ #3 & \\mbox{#4}\\\\ \\end{array} \\right.", 4],
forkthree: ["\\left\\{ \\begin{array}{l l} #1 & \\mbox{#2}\\\\ #3 & \\mbox{#4}\\\\ #5 & \\mbox{#6}\\\\ \\end{array} \\right.", 6],
vecthree: ["\\begin{bmatrix} #1\\\\ #2\\\\ #3 \\end{bmatrix}", 3],
vecthreethree: ["\\begin{bmatrix} #1 & #2 & #3\\\\ #4 & #5 & #6\\\\ #7 & #8 & #9 \\end{bmatrix}", 9]
}
}
});
\ProvidesPackage{mymath}
\usepackage{euler}
\usepackage{amssymb}
\usepackage{amsmath}
\newcommand{\matTT}[9]{
\[
\left|\begin{array}{ccc}
......
OpenCV modules {#mainpage}
==============
- @subpage intro
- @subpage core
<!-- @CMAKE_DOXYGEN_MODULES_REFERENCE@ -->
Introduction {#intro}
============
OpenCV (Open Source Computer Vision Library: <http://opencv.org>) is an open-source BSD-licensed
library that includes several hundreds of computer vision algorithms. The document describes the
so-called OpenCV 2.x API, which is essentially a C++ API, as opposite to the C-based OpenCV 1.x API.
The latter is described in opencv1x.pdf.
OpenCV has a modular structure, which means that the package includes several shared or static
libraries. The following modules are available:
- @ref core - a compact module defining basic data structures, including the dense
multi-dimensional array Mat and basic functions used by all other modules.
- **imgproc** - an image processing module that includes linear and non-linear image filtering,
geometrical image transformations (resize, affine and perspective warping, generic table-based
remapping), color space conversion, histograms, and so on.
- **video** - a video analysis module that includes motion estimation, background subtraction,
and object tracking algorithms.
- **calib3d** - basic multiple-view geometry algorithms, single and stereo camera calibration,
object pose estimation, stereo correspondence algorithms, and elements of 3D reconstruction.
- **features2d** - salient feature detectors, descriptors, and descriptor matchers.
- **objdetect** - detection of objects and instances of the predefined classes (for example,
faces, eyes, mugs, people, cars, and so on).
- **highgui** - an easy-to-use interface to simple UI capabilities.
- **videoio** - an easy-to-use interface to video capturing and video codecs.
- **gpu** - GPU-accelerated algorithms from different OpenCV modules.
- ... some other helper modules, such as FLANN and Google test wrappers, Python bindings, and
others.
The further chapters of the document describe functionality of each module. But first, make sure to
get familiar with the common API concepts used thoroughly in the library.
API Concepts
------------
### cv Namespace
All the OpenCV classes and functions are placed into the cv namespace. Therefore, to access this
functionality from your code, use the cv:: specifier or using namespace cv; directive:
@code
#include "opencv2/core.hpp"
...
cv::Mat H = cv::findHomography(points1, points2, CV_RANSAC, 5);
...
@endcode
or :
~~~
#include "opencv2/core.hpp"
using namespace cv;
...
Mat H = findHomography(points1, points2, CV_RANSAC, 5 );
...
~~~
Some of the current or future OpenCV external names may conflict with STL or other libraries. In
this case, use explicit namespace specifiers to resolve the name conflicts:
@code
Mat a(100, 100, CV_32F);
randu(a, Scalar::all(1), Scalar::all(std::rand()));
cv::log(a, a);
a /= std::log(2.);
@endcode
### Automatic Memory Management
OpenCV handles all the memory automatically.
First of all, std::vector, Mat, and other data structures used by the functions and methods have
destructors that deallocate the underlying memory buffers when needed. This means that the
destructors do not always deallocate the buffers as in case of Mat. They take into account possible
data sharing. A destructor decrements the reference counter associated with the matrix data buffer.
The buffer is deallocated if and only if the reference counter reaches zero, that is, when no other
structures refer to the same buffer. Similarly, when a Mat instance is copied, no actual data is
really copied. Instead, the reference counter is incremented to memorize that there is another owner
of the same data. There is also the Mat::clone method that creates a full copy of the matrix data.
See the example below:
@code
// create a big 8Mb matrix
Mat A(1000, 1000, CV_64F);
// create another header for the same matrix;
// this is an instant operation, regardless of the matrix size.
Mat B = A;
// create another header for the 3-rd row of A; no data is copied either
Mat C = B.row(3);
// now create a separate copy of the matrix
Mat D = B.clone();
// copy the 5-th row of B to C, that is, copy the 5-th row of A
// to the 3-rd row of A.
B.row(5).copyTo(C);
// now let A and D share the data; after that the modified version
// of A is still referenced by B and C.
A = D;
// now make B an empty matrix (which references no memory buffers),
// but the modified version of A will still be referenced by C,
// despite that C is just a single row of the original A
B.release();
// finally, make a full copy of C. As a result, the big modified
// matrix will be deallocated, since it is not referenced by anyone
C = C.clone();
@endcode
You see that the use of Mat and other basic structures is simple. But what about high-level classes
or even user data types created without taking automatic memory management into account? For them,
OpenCV offers the Ptr template class that is similar to std::shared\_ptr from C++11. So, instead of
using plain pointers:
@code
T* ptr = new T(...);
@endcode
you can use:
@code
Ptr<T> ptr(new T(...));
@endcode
or:
@code
Ptr<T> ptr = makePtr<T>(...);
@endcode
Ptr\<T\> encapsulates a pointer to a T instance and a reference counter associated with the pointer.
See the Ptr description for details.
### Automatic Allocation of the Output Data
OpenCV deallocates the memory automatically, as well as automatically allocates the memory for
output function parameters most of the time. So, if a function has one or more input arrays (cv::Mat
instances) and some output arrays, the output arrays are automatically allocated or reallocated. The
size and type of the output arrays are determined from the size and type of input arrays. If needed,
the functions take extra parameters that help to figure out the output array properties.
Example:
@code
#include "opencv2/imgproc.hpp"
#include "opencv2/highgui.hpp"
using namespace cv;
int main(int, char**)
{
VideoCapture cap(0);
if(!cap.isOpened()) return -1;
Mat frame, edges;
namedWindow("edges",1);
for(;;)
{
cap >> frame;
cvtColor(frame, edges, COLOR_BGR2GRAY);
GaussianBlur(edges, edges, Size(7,7), 1.5, 1.5);
Canny(edges, edges, 0, 30, 3);
imshow("edges", edges);
if(waitKey(30) >= 0) break;
}
return 0;
}
@endcode
The array frame is automatically allocated by the \>\> operator since the video frame resolution and
the bit-depth is known to the video capturing module. The array edges is automatically allocated by
the cvtColor function. It has the same size and the bit-depth as the input array. The number of
channels is 1 because the color conversion code COLOR\_BGR2GRAY is passed, which means a color to
grayscale conversion. Note that frame and edges are allocated only once during the first execution
of the loop body since all the next video frames have the same resolution. If you somehow change the
video resolution, the arrays are automatically reallocated.
The key component of this technology is the Mat::create method. It takes the desired array size and
type. If the array already has the specified size and type, the method does nothing. Otherwise, it
releases the previously allocated data, if any (this part involves decrementing the reference
counter and comparing it with zero), and then allocates a new buffer of the required size. Most
functions call the Mat::create method for each output array, and so the automatic output data
allocation is implemented.
Some notable exceptions from this scheme are cv::mixChannels, cv::RNG::fill, and a few other
functions and methods. They are not able to allocate the output array, so you have to do this in
advance.
### Saturation Arithmetics
As a computer vision library, OpenCV deals a lot with image pixels that are often encoded in a
compact, 8- or 16-bit per channel, form and thus have a limited value range. Furthermore, certain
operations on images, like color space conversions, brightness/contrast adjustments, sharpening,
complex interpolation (bi-cubic, Lanczos) can produce values out of the available range. If you just
store the lowest 8 (16) bits of the result, this results in visual artifacts and may affect a
further image analysis. To solve this problem, the so-called *saturation* arithmetics is used. For
example, to store r, the result of an operation, to an 8-bit image, you find the nearest value
within the 0..255 range:
\f[I(x,y)= \min ( \max (\textrm{round}(r), 0), 255)\f]
Similar rules are applied to 8-bit signed, 16-bit signed and unsigned types. This semantics is used
everywhere in the library. In C++ code, it is done using the saturate\_cast\<\> functions that
resemble standard C++ cast operations. See below the implementation of the formula provided above:
@code
I.at<uchar>(y, x) = saturate_cast<uchar>(r);
@endcode
where cv::uchar is an OpenCV 8-bit unsigned integer type. In the optimized SIMD code, such SSE2
instructions as paddusb, packuswb, and so on are used. They help achieve exactly the same behavior
as in C++ code.
@note Saturation is not applied when the result is 32-bit integer.
### Fixed Pixel Types. Limited Use of Templates
Templates is a great feature of C++ that enables implementation of very powerful, efficient and yet
safe data structures and algorithms. However, the extensive use of templates may dramatically
increase compilation time and code size. Besides, it is difficult to separate an interface and
implementation when templates are used exclusively. This could be fine for basic algorithms but not
good for computer vision libraries where a single algorithm may span thousands lines of code.
Because of this and also to simplify development of bindings for other languages, like Python, Java,
Matlab that do not have templates at all or have limited template capabilities, the current OpenCV
implementation is based on polymorphism and runtime dispatching over templates. In those places
where runtime dispatching would be too slow (like pixel access operators), impossible (generic
Ptr\<\> implementation), or just very inconvenient (saturate\_cast\<\>()) the current implementation
introduces small template classes, methods, and functions. Anywhere else in the current OpenCV
version the use of templates is limited.
Consequently, there is a limited fixed set of primitive data types the library can operate on. That
is, array elements should have one of the following types:
- 8-bit unsigned integer (uchar)
- 8-bit signed integer (schar)
- 16-bit unsigned integer (ushort)
- 16-bit signed integer (short)
- 32-bit signed integer (int)
- 32-bit floating-point number (float)
- 64-bit floating-point number (double)
- a tuple of several elements where all elements have the same type (one of the above). An array
whose elements are such tuples, are called multi-channel arrays, as opposite to the
single-channel arrays, whose elements are scalar values. The maximum possible number of
channels is defined by the CV\_CN\_MAX constant, which is currently set to 512.
For these basic types, the following enumeration is applied:
@code
enum { CV_8U=0, CV_8S=1, CV_16U=2, CV_16S=3, CV_32S=4, CV_32F=5, CV_64F=6 };
@endcode
Multi-channel (n-channel) types can be specified using the following options:
- CV_8UC1 ... CV_64FC4 constants (for a number of channels from 1 to 4)
- CV_8UC(n) ... CV_64FC(n) or CV_MAKETYPE(CV_8U, n) ... CV_MAKETYPE(CV_64F, n) macros when
the number of channels is more than 4 or unknown at the compilation time.
@note `CV_32FC1 == CV_32F, CV_32FC2 == CV_32FC(2) == CV_MAKETYPE(CV_32F, 2)`, and
`CV_MAKETYPE(depth, n) == ((x&7)<<3) + (n-1)``. This means that the constant type is formed from the
depth, taking the lowest 3 bits, and the number of channels minus 1, taking the next
`log2(CV_CN_MAX)`` bits.
Examples:
@code
Mat mtx(3, 3, CV_32F); // make a 3x3 floating-point matrix
Mat cmtx(10, 1, CV_64FC2); // make a 10x1 2-channel floating-point
// matrix (10-element complex vector)
Mat img(Size(1920, 1080), CV_8UC3); // make a 3-channel (color) image
// of 1920 columns and 1080 rows.
Mat grayscale(image.size(), CV_MAKETYPE(image.depth(), 1)); // make a 1-channel image of
// the same size and same
// channel type as img
@endcode
Arrays with more complex elements cannot be constructed or processed using OpenCV. Furthermore, each
function or method can handle only a subset of all possible array types. Usually, the more complex
the algorithm is, the smaller the supported subset of formats is. See below typical examples of such
limitations:
- The face detection algorithm only works with 8-bit grayscale or color images.
- Linear algebra functions and most of the machine learning algorithms work with floating-point
arrays only.
- Basic functions, such as cv::add, support all types.
- Color space conversion functions support 8-bit unsigned, 16-bit unsigned, and 32-bit
floating-point types.
The subset of supported types for each function has been defined from practical needs and could be
extended in future based on user requests.
### InputArray and OutputArray
Many OpenCV functions process dense 2-dimensional or multi-dimensional numerical arrays. Usually,
such functions take cppMat as parameters, but in some cases it's more convenient to use
std::vector\<\> (for a point set, for example) or Matx\<\> (for 3x3 homography matrix and such). To
avoid many duplicates in the API, special "proxy" classes have been introduced. The base "proxy"
class is InputArray. It is used for passing read-only arrays on a function input. The derived from
InputArray class OutputArray is used to specify an output array for a function. Normally, you should
not care of those intermediate types (and you should not declare variables of those types
explicitly) - it will all just work automatically. You can assume that instead of
InputArray/OutputArray you can always use Mat, std::vector\<\>, Matx\<\>, Vec\<\> or Scalar. When a
function has an optional input or output array, and you do not have or do not want one, pass
cv::noArray().
### Error Handling
OpenCV uses exceptions to signal critical errors. When the input data has a correct format and
belongs to the specified value range, but the algorithm cannot succeed for some reason (for example,
the optimization algorithm did not converge), it returns a special error code (typically, just a
boolean variable).
The exceptions can be instances of the cv::Exception class or its derivatives. In its turn,
cv::Exception is a derivative of std::exception. So it can be gracefully handled in the code using
other standard C++ library components.
The exception is typically thrown either using the CV\_Error(errcode, description) macro, or its
printf-like CV\_Error\_(errcode, printf-spec, (printf-args)) variant, or using the
CV\_Assert(condition) macro that checks the condition and throws an exception when it is not
satisfied. For performance-critical code, there is CV\_DbgAssert(condition) that is only retained in
the Debug configuration. Due to the automatic memory management, all the intermediate buffers are
automatically deallocated in case of a sudden error. You only need to add a try statement to catch
exceptions, if needed: :
@code
try
{
... // call OpenCV
}
catch( cv::Exception& e )
{
const char* err_msg = e.what();
std::cout << "exception caught: " << err_msg << std::endl;
}
@endcode
### Multi-threading and Re-enterability
The current OpenCV implementation is fully re-enterable. That is, the same function, the same
*constant* method of a class instance, or the same *non-constant* method of different class
instances can be called from different threads. Also, the same cv::Mat can be used in different
threads because the reference-counting operations use the architecture-specific atomic instructions.
......@@ -48,10 +48,15 @@
#include <opencv2/core.hpp>
/*! @file */
namespace cv
{
//! @addtogroup core
//! @{
/** @brief Affine transform
@todo document
*/
template<typename T>
class Affine3
{
......@@ -63,30 +68,31 @@ namespace cv
Affine3();
//Augmented affine matrix
//! Augmented affine matrix
Affine3(const Mat4& affine);
//Rotation matrix
//! Rotation matrix
Affine3(const Mat3& R, const Vec3& t = Vec3::all(0));
//Rodrigues vector
//! Rodrigues vector
Affine3(const Vec3& rvec, const Vec3& t = Vec3::all(0));
//Combines all contructors above. Supports 4x4, 4x3, 3x3, 1x3, 3x1 sizes of data matrix
//! Combines all contructors above. Supports 4x4, 4x3, 3x3, 1x3, 3x1 sizes of data matrix
explicit Affine3(const Mat& data, const Vec3& t = Vec3::all(0));
//From 16th element array
//! From 16th element array
explicit Affine3(const float_type* vals);
//! Create identity transform
static Affine3 Identity();
//Rotation matrix
//! Rotation matrix
void rotation(const Mat3& R);
//Rodrigues vector
//! Rodrigues vector
void rotation(const Vec3& rvec);
//Combines rotation methods above. Suports 3x3, 1x3, 3x1 sizes of data matrix;
//! Combines rotation methods above. Suports 3x3, 1x3, 3x1 sizes of data matrix;
void rotation(const Mat& data);
void linear(const Mat3& L);
......@@ -96,21 +102,21 @@ namespace cv
Mat3 linear() const;
Vec3 translation() const;
//Rodrigues vector
//! Rodrigues vector
Vec3 rvec() const;
Affine3 inv(int method = cv::DECOMP_SVD) const;
// a.rotate(R) is equivalent to Affine(R, 0) * a;
//! a.rotate(R) is equivalent to Affine(R, 0) * a;
Affine3 rotate(const Mat3& R) const;
// a.rotate(R) is equivalent to Affine(rvec, 0) * a;
//! a.rotate(R) is equivalent to Affine(rvec, 0) * a;
Affine3 rotate(const Vec3& rvec) const;
// a.translate(t) is equivalent to Affine(E, t) * a;
//! a.translate(t) is equivalent to Affine(E, t) * a;
Affine3 translate(const Vec3& t) const;
// a.concatenate(affine) is equivalent to affine * a;
//! a.concatenate(affine) is equivalent to affine * a;
Affine3 concatenate(const Affine3& affine) const;
template <typename Y> operator Affine3<Y>() const;
......@@ -155,11 +161,15 @@ namespace cv
typedef Vec<channel_type, channels> vec_type;
};
//! @} core
}
//! @cond IGNORED
///////////////////////////////////////////////////////////////////////////////////
/// Implementaiton
// Implementaiton
template<typename T> inline
cv::Affine3<T>::Affine3()
......@@ -431,7 +441,6 @@ cv::Affine3<Y> cv::Affine3<T>::cast() const
return Affine3<Y>(matrix);
}
/** @cond IGNORED */
template<typename T> inline
cv::Affine3<T> cv::operator*(const cv::Affine3<T>& affine1, const cv::Affine3<T>& affine2)
{
......@@ -449,7 +458,6 @@ V cv::operator*(const cv::Affine3<T>& affine, const V& v)
r.z = m.val[8] * v.x + m.val[9] * v.y + m.val[10] * v.z + m.val[11];
return r;
}
/** @endcond */
static inline
cv::Vec3f cv::operator*(const cv::Affine3f& affine, const cv::Vec3f& v)
......@@ -507,6 +515,7 @@ cv::Affine3<T>::operator Eigen::Transform<T, 3, Eigen::Affine>() const
#endif /* defined EIGEN_WORLD_VERSION && defined EIGEN_GEOMETRY_MODULE_H */
//! @endcond
#endif /* __cplusplus */
......
......@@ -10,6 +10,9 @@
namespace cv
{
//! @addtogroup core
//! @{
class BufferPoolController
{
protected:
......@@ -21,6 +24,8 @@ public:
virtual void freeAllReservedBuffers() = 0;
};
//! @}
}
#endif // __OPENCV_CORE_BUFFER_POOL_HPP__
......@@ -51,13 +51,22 @@
#include "opencv2/core.hpp"
#include "opencv2/core/cuda_types.hpp"
/**
@defgroup cuda CUDA-accelerated Computer Vision
@{
@defgroup cuda_struct Data structures
@}
*/
namespace cv { namespace cuda {
//////////////////////////////// GpuMat ///////////////////////////////
//! @addtogroup cuda_struct
//! @{
// Smart pointer for GPU memory with reference counting.
// Its interface is mostly similar with cv::Mat.
//////////////////////////////// GpuMat ///////////////////////////////
//! Smart pointer for GPU memory with reference counting.
//! Its interface is mostly similar with cv::Mat.
class CV_EXPORTS GpuMat
{
public:
......@@ -283,11 +292,10 @@ CV_EXPORTS void setBufferPoolConfig(int deviceId, size_t stackSize, int stackCou
//////////////////////////////// CudaMem ////////////////////////////////
// CudaMem is limited cv::Mat with page locked memory allocation.
// Page locked memory is only needed for async and faster coping to GPU.
// It is convertable to cv::Mat header without reference counting
// so you can use it with other opencv functions.
//! CudaMem is limited cv::Mat with page locked memory allocation.
//! Page locked memory is only needed for async and faster coping to GPU.
//! It is convertable to cv::Mat header without reference counting
//! so you can use it with other opencv functions.
class CV_EXPORTS CudaMem
{
public:
......@@ -363,10 +371,9 @@ CV_EXPORTS void unregisterPageLocked(Mat& m);
///////////////////////////////// Stream //////////////////////////////////
// Encapculates Cuda Stream. Provides interface for async coping.
// Passed to each function that supports async kernel execution.
// Reference counting is enabled.
//! Encapculates Cuda Stream. Provides interface for async coping.
//! Passed to each function that supports async kernel execution.
//! Reference counting is enabled.
class CV_EXPORTS Stream
{
typedef void (Stream::*bool_type)() const;
......@@ -563,10 +570,10 @@ public:
enum ComputeMode
{
ComputeModeDefault, /**< default compute mode (Multiple threads can use ::cudaSetDevice() with this device) */
ComputeModeExclusive, /**< compute-exclusive-thread mode (Only one thread in one process will be able to use ::cudaSetDevice() with this device) */
ComputeModeProhibited, /**< compute-prohibited mode (No threads can use ::cudaSetDevice() with this device) */
ComputeModeExclusiveProcess /**< compute-exclusive-process mode (Many threads in one process will be able to use ::cudaSetDevice() with this device) */
ComputeModeDefault, /**< default compute mode (Multiple threads can use cudaSetDevice with this device) */
ComputeModeExclusive, /**< compute-exclusive-thread mode (Only one thread in one process will be able to use cudaSetDevice with this device) */
ComputeModeProhibited, /**< compute-prohibited mode (No threads can use cudaSetDevice with this device) */
ComputeModeExclusiveProcess /**< compute-exclusive-process mode (Many threads in one process will be able to use cudaSetDevice with this device) */
};
//! compute mode
......@@ -686,6 +693,8 @@ private:
CV_EXPORTS void printCudaDeviceInfo(int device);
CV_EXPORTS void printShortCudaDeviceInfo(int device);
//! @}
}} // namespace cv { namespace cuda {
......
......@@ -46,6 +46,8 @@
#include "opencv2/core/cuda.hpp"
//! @cond IGNORED
namespace cv { namespace cuda {
//////////////////////////////// GpuMat ///////////////////////////////
......@@ -224,7 +226,6 @@ const _Tp* GpuMat::ptr(int y) const
return (const _Tp*)ptr(y);
}
/** @cond IGNORED */
template <class T> inline
GpuMat::operator PtrStepSz<T>() const
{
......@@ -236,7 +237,6 @@ GpuMat::operator PtrStep<T>() const
{
return PtrStep<T>((T*)data, step);
}
/** @endcond */
inline
GpuMat GpuMat::row(int y) const
......@@ -589,6 +589,7 @@ bool DeviceInfo::supports(FeatureSet feature_set) const
return version >= feature_set;
}
}} // namespace cv { namespace cuda {
//////////////////////////////// Mat ////////////////////////////////
......@@ -604,4 +605,6 @@ Mat::Mat(const cuda::GpuMat& m)
}
//! @endcond
#endif // __OPENCV_CORE_CUDAINL_HPP__
......@@ -43,8 +43,11 @@
#ifndef __OPENCV_CUDA_DEVICE_BLOCK_HPP__
#define __OPENCV_CUDA_DEVICE_BLOCK_HPP__
namespace cv { namespace cuda { namespace device
{
//! @addtogroup cuda
//! @{
struct Block
{
static __device__ __forceinline__ unsigned int id()
......@@ -198,6 +201,7 @@ namespace cv { namespace cuda { namespace device
}
}
};
//!@}
}}}
#endif /* __OPENCV_CUDA_DEVICE_BLOCK_HPP__ */
......@@ -49,6 +49,9 @@
namespace cv { namespace cuda { namespace device
{
//! @addtogroup cuda
//! @{
//////////////////////////////////////////////////////////////
// BrdConstant
......@@ -709,6 +712,7 @@ namespace cv { namespace cuda { namespace device
int width;
D val;
};
//! @}
}}} // namespace cv { namespace cuda { namespace cudev
#endif // __OPENCV_CUDA_BORDER_INTERPOLATE_HPP__
......@@ -47,6 +47,8 @@
namespace cv { namespace cuda { namespace device
{
//! @addtogroup cuda
//! @{
// All OPENCV_CUDA_IMPLEMENT_*_TRAITS(ColorSpace1_to_ColorSpace2, ...) macros implements
// template <typename T> class ColorSpace1_to_ColorSpace2_traits
// {
......@@ -296,6 +298,7 @@ namespace cv { namespace cuda { namespace device
OPENCV_CUDA_IMPLEMENT_Luv2RGB_TRAITS(luv4_to_lbgra, 4, 4, false, 0)
#undef OPENCV_CUDA_IMPLEMENT_Luv2RGB_TRAITS
//! @}
}}} // namespace cv { namespace cuda { namespace cudev
#endif // __OPENCV_CUDA_BORDER_INTERPOLATE_HPP__
......@@ -48,6 +48,7 @@
#include "opencv2/core/cvdef.h"
#include "opencv2/core/base.hpp"
#ifndef CV_PI_F
#ifndef CV_PI
#define CV_PI_F 3.14159265f
......@@ -57,11 +58,14 @@
#endif
namespace cv { namespace cuda {
//! @addtogroup cuda
//! @{
static inline void checkCudaError(cudaError_t err, const char* file, const int line, const char* func)
{
if (cudaSuccess != err)
cv::error(cv::Error::GpuApiCallError, cudaGetErrorString(err), func, file, line);
}
//! @}
}}
#ifndef cudaSafeCall
......@@ -70,6 +74,8 @@ namespace cv { namespace cuda {
namespace cv { namespace cuda
{
//! @addtogroup cuda
//! @{
template <typename T> static inline bool isAligned(const T* ptr, size_t size)
{
return reinterpret_cast<size_t>(ptr) % size == 0;
......@@ -79,12 +85,15 @@ namespace cv { namespace cuda
{
return step % size == 0;
}
//! @}
}}
namespace cv { namespace cuda
{
namespace device
{
//! @addtogroup cuda
//! @{
__host__ __device__ __forceinline__ int divUp(int total, int grain)
{
return (total + grain - 1) / grain;
......@@ -95,9 +104,8 @@ namespace cv { namespace cuda
cudaChannelFormatDesc desc = cudaCreateChannelDesc<T>();
cudaSafeCall( cudaBindTexture2D(0, tex, img.ptr(), &desc, img.cols, img.rows, img.step) );
}
//! @}
}
}}
#endif // __OPENCV_CUDA_COMMON_HPP__
......@@ -47,6 +47,9 @@
namespace cv { namespace cuda { namespace device
{
//! @addtogroup cuda
//! @{
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 200
// for Fermi memory space is detected automatically
......@@ -100,6 +103,7 @@ namespace cv { namespace cuda { namespace device
#undef OPENCV_CUDA_ASM_PTR
#endif // __CUDA_ARCH__ >= 200
//! @}
}}} // namespace cv { namespace cuda { namespace cudev
#endif // __OPENCV_CUDA_DATAMOV_UTILS_HPP__
......@@ -49,6 +49,8 @@
#include "../limits.hpp"
#include "../functional.hpp"
//! @cond IGNORED
namespace cv { namespace cuda { namespace device
{
#ifndef CV_DESCALE
......@@ -1973,4 +1975,6 @@ namespace cv { namespace cuda { namespace device
}}} // namespace cv { namespace cuda { namespace cudev
//! @endcond
#endif // __OPENCV_CUDA_COLOR_DETAIL_HPP__
......@@ -47,6 +47,8 @@
#include "../warp.hpp"
#include "../warp_shuffle.hpp"
//! @cond IGNORED
namespace cv { namespace cuda { namespace device
{
namespace reduce_detail
......@@ -358,4 +360,6 @@ namespace cv { namespace cuda { namespace device
}
}}}
//! @endcond
#endif // __OPENCV_CUDA_REDUCE_DETAIL_HPP__
......@@ -47,6 +47,8 @@
#include "../warp.hpp"
#include "../warp_shuffle.hpp"
//! @cond IGNORED
namespace cv { namespace cuda { namespace device
{
namespace reduce_key_val_detail
......@@ -495,4 +497,6 @@ namespace cv { namespace cuda { namespace device
}
}}}
//! @endcond
#endif // __OPENCV_CUDA_PRED_VAL_REDUCE_DETAIL_HPP__
......@@ -47,6 +47,8 @@
#include "../vec_traits.hpp"
#include "../functional.hpp"
//! @cond IGNORED
namespace cv { namespace cuda { namespace device
{
namespace transform_detail
......@@ -392,4 +394,6 @@ namespace cv { namespace cuda { namespace device
} // namespace transform_detail
}}} // namespace cv { namespace cuda { namespace cudev
//! @endcond
#endif // __OPENCV_CUDA_TRANSFORM_DETAIL_HPP__
......@@ -46,6 +46,8 @@
#include "../common.hpp"
#include "../vec_traits.hpp"
//! @cond IGNORED
namespace cv { namespace cuda { namespace device
{
namespace type_traits_detail
......@@ -184,4 +186,6 @@ namespace cv { namespace cuda { namespace device
} // namespace type_traits_detail
}}} // namespace cv { namespace cuda { namespace cudev
//! @endcond
#endif // __OPENCV_CUDA_TYPE_TRAITS_DETAIL_HPP__
......@@ -45,6 +45,8 @@
#include "../datamov_utils.hpp"
//! @cond IGNORED
namespace cv { namespace cuda { namespace device
{
namespace vec_distance_detail
......@@ -114,4 +116,6 @@ namespace cv { namespace cuda { namespace device
} // namespace vec_distance_detail
}}} // namespace cv { namespace cuda { namespace cudev
//! @endcond
#endif // __OPENCV_CUDA_VEC_DISTANCE_DETAIL_HPP__
......@@ -45,6 +45,8 @@
namespace cv { namespace cuda { namespace device
{
//! @addtogroup cuda
//! @{
template<class T> struct DynamicSharedMem
{
__device__ __forceinline__ operator T*()
......@@ -75,6 +77,7 @@ namespace cv { namespace cuda { namespace device
return (double*)__smem_d;
}
};
//! @}
}}}
#endif // __OPENCV_CUDA_DYNAMIC_SMEM_HPP__
......@@ -48,6 +48,8 @@
namespace cv { namespace cuda { namespace device
{
//! @addtogroup cuda
//! @{
struct Emulation
{
......@@ -256,6 +258,7 @@ namespace cv { namespace cuda { namespace device
}
};
}; //struct Emulation
//!@}
}}} // namespace cv { namespace cuda { namespace cudev
#endif /* OPENCV_CUDA_EMULATION_HPP_ */
......@@ -50,6 +50,8 @@
namespace cv { namespace cuda { namespace device
{
//! @addtogroup cuda
//! @{
template <typename Ptr2D> struct PointFilter
{
typedef typename Ptr2D::elem_type elem_type;
......@@ -273,6 +275,7 @@ namespace cv { namespace cuda { namespace device
float scale_x, scale_y;
int width, haight;
};
//! @}
}}} // namespace cv { namespace cuda { namespace cudev
#endif // __OPENCV_CUDA_FILTERS_HPP__
......@@ -47,6 +47,8 @@
namespace cv { namespace cuda { namespace device
{
//! @addtogroup cuda
//! @{
template<class Func>
void printFuncAttrib(Func& func)
{
......@@ -66,6 +68,7 @@ namespace cv { namespace cuda { namespace device
printf("\n");
fflush(stdout);
}
//! @}
}}} // namespace cv { namespace cuda { namespace cudev
#endif /* __OPENCV_CUDA_DEVICE_FUNCATTRIB_HPP_ */
......@@ -51,6 +51,8 @@
namespace cv { namespace cuda { namespace device
{
//! @addtogroup cuda
//! @{
// Function Objects
template<typename Argument, typename Result> struct unary_function : public std::unary_function<Argument, Result> {};
template<typename Argument1, typename Argument2, typename Result> struct binary_function : public std::binary_function<Argument1, Argument2, Result> {};
......@@ -784,6 +786,7 @@ namespace cv { namespace cuda { namespace device
#define OPENCV_CUDA_TRANSFORM_FUNCTOR_TRAITS(type) \
template <> struct TransformFunctorTraits< type > : DefaultTransformFunctorTraits< type >
//! @}
}}} // namespace cv { namespace cuda { namespace cudev
#endif // __OPENCV_CUDA_FUNCTIONAL_HPP__
......@@ -49,7 +49,8 @@
namespace cv { namespace cuda { namespace device
{
//! @addtogroup cuda
//! @{
template <class T> struct numeric_limits;
template <> struct numeric_limits<bool>
......@@ -116,7 +117,7 @@ template <> struct numeric_limits<double>
__device__ __forceinline__ static double epsilon() { return DBL_EPSILON; }
static const bool is_signed = true;
};
//! @}
}}} // namespace cv { namespace cuda { namespace cudev {
#endif // __OPENCV_CUDA_LIMITS_HPP__
......@@ -49,6 +49,8 @@
namespace cv { namespace cuda { namespace device
{
//! @addtogroup cuda
//! @{
template <int N, typename T, class Op>
__device__ __forceinline__ void reduce(volatile T* smem, T& val, unsigned int tid, const Op& op)
{
......@@ -192,6 +194,7 @@ namespace cv { namespace cuda { namespace device
{
return thrust::make_tuple((volatile T0*) t0, (volatile T1*) t1, (volatile T2*) t2, (volatile T3*) t3, (volatile T4*) t4, (volatile T5*) t5, (volatile T6*) t6, (volatile T7*) t7, (volatile T8*) t8, (volatile T9*) t9);
}
//! @}
}}}
#endif // __OPENCV_CUDA_UTILITY_HPP__
......@@ -47,6 +47,8 @@
namespace cv { namespace cuda { namespace device
{
//! @addtogroup cuda
//! @{
template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(uchar v) { return _Tp(v); }
template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(schar v) { return _Tp(v); }
template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(ushort v) { return _Tp(v); }
......@@ -279,6 +281,7 @@ namespace cv { namespace cuda { namespace device
return saturate_cast<uint>((float)v);
#endif
}
//! @}
}}}
#endif /* __OPENCV_CUDA_SATURATE_CAST_HPP__ */
......@@ -50,6 +50,8 @@
namespace cv { namespace cuda { namespace device
{
//! @addtogroup cuda
//! @{
enum ScanKind { EXCLUSIVE = 0, INCLUSIVE = 1 };
template <ScanKind Kind, typename T, typename F> struct WarpScan
......@@ -245,6 +247,7 @@ namespace cv { namespace cuda { namespace device
return warpScanInclusive(idata, s_Data, tid);
}
}
//! @}
}}}
#endif // __OPENCV_CUDA_SCAN_HPP__
......@@ -75,7 +75,7 @@
#include "common.hpp"
/*
/** @file
This header file contains inline functions that implement intra-word SIMD
operations, that are hardware accelerated on sm_3x (Kepler) GPUs. Efficient
emulation code paths are provided for earlier architectures (sm_1x, sm_2x)
......@@ -125,6 +125,8 @@
namespace cv { namespace cuda { namespace device
{
//! @addtogroup cuda
//! @{
// 2
static __device__ __forceinline__ unsigned int vadd2(unsigned int a, unsigned int b)
......@@ -904,6 +906,7 @@ namespace cv { namespace cuda { namespace device
return r;
}
//! @}
}}}
#endif // __OPENCV_CUDA_SIMD_FUNCTIONS_HPP__
......@@ -49,6 +49,8 @@
namespace cv { namespace cuda { namespace device
{
//! @addtogroup cuda
//! @{
template <typename T, typename D, typename UnOp, typename Mask>
static inline void transform(PtrStepSz<T> src, PtrStepSz<D> dst, UnOp op, const Mask& mask, cudaStream_t stream)
{
......@@ -62,6 +64,7 @@ namespace cv { namespace cuda { namespace device
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);
}
//! @}
}}}
#endif // __OPENCV_CUDA_TRANSFORM_HPP__
......@@ -47,6 +47,8 @@
namespace cv { namespace cuda { namespace device
{
//! @addtogroup cuda
//! @{
template <typename T> struct IsSimpleParameter
{
enum {value = type_traits_detail::IsIntegral<T>::value || type_traits_detail::IsFloat<T>::value ||
......@@ -77,6 +79,7 @@ namespace cv { namespace cuda { namespace device
typedef typename type_traits_detail::Select<IsSimpleParameter<UnqualifiedType>::value,
T, typename type_traits_detail::AddParameterType<T>::type>::type ParameterType;
};
//! @}
}}}
#endif // __OPENCV_CUDA_TYPE_TRAITS_HPP__
......@@ -48,6 +48,8 @@
namespace cv { namespace cuda { namespace device
{
//! @addtogroup cuda
//! @{
#define OPENCV_CUDA_LOG_WARP_SIZE (5)
#define OPENCV_CUDA_WARP_SIZE (1 << OPENCV_CUDA_LOG_WARP_SIZE)
#define OPENCV_CUDA_LOG_MEM_BANKS ((__CUDA_ARCH__ >= 200) ? 5 : 4) // 32 banks on fermi, 16 on tesla
......@@ -208,6 +210,7 @@ namespace cv { namespace cuda { namespace device
return false;
}
//! @}
}}} // namespace cv { namespace cuda { namespace cudev
#endif // __OPENCV_CUDA_UTILITY_HPP__
......@@ -49,6 +49,8 @@
namespace cv { namespace cuda { namespace device
{
//! @addtogroup cuda
//! @{
template <typename T> struct L1Dist
{
typedef int value_type;
......@@ -219,6 +221,7 @@ namespace cv { namespace cuda { namespace device
U vec1Vals[MAX_LEN / THREAD_DIM];
};
//! @}
}}} // namespace cv { namespace cuda { namespace cudev
#endif // __OPENCV_CUDA_VEC_DISTANCE_HPP__
......@@ -49,6 +49,9 @@
namespace cv { namespace cuda { namespace device
{
//! @addtogroup cuda
//! @{
// saturate_cast
namespace vec_math_detail
......@@ -917,6 +920,8 @@ CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, double, double, double)
#undef CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC
//! @}
}}} // namespace cv { namespace cuda { namespace device
#endif // __OPENCV_CUDA_VECMATH_HPP__
......@@ -47,6 +47,8 @@
namespace cv { namespace cuda { namespace device
{
//! @addtogroup cuda
//! @{
template<typename T, int N> struct TypeVec;
struct __align__(8) uchar8
......@@ -275,6 +277,7 @@ namespace cv { namespace cuda { namespace device
static __device__ __host__ __forceinline__ char8 make(schar a0, schar a1, schar a2, schar a3, schar a4, schar a5, schar a6, schar a7) {return make_char8(a0, a1, a2, a3, a4, a5, a6, a7);}
static __device__ __host__ __forceinline__ char8 make(const schar* v) {return make_char8(v[0], v[1], v[2], v[3], v[4], v[5], v[6], v[7]);}
};
//! @}
}}} // namespace cv { namespace cuda { namespace cudev
#endif // __OPENCV_CUDA_VEC_TRAITS_HPP__
......@@ -45,6 +45,8 @@
namespace cv { namespace cuda { namespace device
{
//! @addtogroup cuda
//! @{
struct Warp
{
enum
......@@ -126,6 +128,7 @@ namespace cv { namespace cuda { namespace device
*t = value;
}
};
//! @}
}}} // namespace cv { namespace cuda { namespace cudev
#endif /* __OPENCV_CUDA_DEVICE_WARP_HPP__ */
Markdown is supported
0% .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑!
想要评论请 注册