private:\r
\r
\r
- static void NCVDebugOutputHandler(const char* msg) { CV_Error(CV_GpuApiCallError, msg); }\r
+ static void NCVDebugOutputHandler(const std::string &msg) { CV_Error(CV_GpuApiCallError, msg.c_str()); }\r
\r
\r
NCVStatus load(const string& classifierFile)\r
//M*/\r
\r
\r
-#include <ios>\r
-#include <stdarg.h>\r
+#include <iostream>\r
+#include <string>\r
#include <vector>\r
-#include <cstdio>\r
#include "NCV.hpp"\r
\r
using namespace std;\r
//==============================================================================\r
\r
\r
-static void stdioDebugOutput(const char *msg)\r
+static void stdDebugOutput(const string &msg)\r
{\r
- printf("%s", msg);\r
+ cout << msg;\r
}\r
\r
\r
-static NCVDebugOutputHandler *debugOutputHandler = stdioDebugOutput;\r
+static NCVDebugOutputHandler *debugOutputHandler = stdDebugOutput;\r
\r
\r
-void ncvDebugOutput(const char *msg, ...)\r
+void ncvDebugOutput(const string &msg)\r
{\r
- const int K_DEBUG_STRING_MAXLEN = 1024;\r
- char buffer[K_DEBUG_STRING_MAXLEN];\r
- va_list args;\r
- va_start(args, msg);\r
- vsnprintf(buffer, K_DEBUG_STRING_MAXLEN, msg, args);\r
- va_end (args);\r
- debugOutputHandler(buffer);\r
+ debugOutputHandler(msg);\r
}\r
\r
\r
\r
allocBegin = NULL;\r
\r
- if (reusePtr == NULL)\r
+ if (reusePtr == NULL && capacity != 0)\r
{\r
bReusesMemory = false;\r
switch (memT)\r
{\r
ncvAssertPrintCheck(currentSize == 0, "NCVMemStackAllocator dtor:: not all objects were deallocated properly, forcing destruction");\r
\r
- if (!bReusesMemory)\r
+ if (!bReusesMemory && (allocBegin != (Ncv8u *)(0x1)))\r
{\r
switch (_memType)\r
{\r
seg.clear();\r
ncvAssertReturn(isInitialized(), NCV_ALLOCATOR_BAD_ALLOC);\r
\r
- size = alignUp(static_cast<Ncv32u>(size), this->_alignment);\r
+ size = alignUp(size, this->_alignment);\r
this->currentSize += size;\r
this->_maxSize = std::max(this->_maxSize, this->currentSize);\r
\r
break;\r
}\r
\r
- this->currentSize += alignUp(static_cast<Ncv32u>(size), this->_alignment);\r
+ this->currentSize += alignUp(size, this->_alignment);\r
this->_maxSize = std::max(this->_maxSize, this->currentSize);\r
\r
seg.begin.memtype = this->_memType;\r
ncvAssertReturn(seg.begin.memtype == this->_memType, NCV_ALLOCATOR_BAD_DEALLOC);\r
ncvAssertReturn(seg.begin.ptr != NULL, NCV_ALLOCATOR_BAD_DEALLOC);\r
\r
- ncvAssertReturn(currentSize >= alignUp(static_cast<Ncv32u>(seg.size), this->_alignment), NCV_ALLOCATOR_BAD_DEALLOC);\r
- currentSize -= alignUp(static_cast<Ncv32u>(seg.size), this->_alignment);\r
+ ncvAssertReturn(currentSize >= alignUp(seg.size, this->_alignment), NCV_ALLOCATOR_BAD_DEALLOC);\r
+ currentSize -= alignUp(seg.size, this->_alignment);\r
\r
switch (this->_memType)\r
{\r
#ifndef _ncv_hpp_\r
#define _ncv_hpp_\r
\r
-#if (defined WIN32 || defined _WIN32 || defined WINCE) && defined CVAPI_EXPORTS //&& !defined(__CUDACC__) \r
+#if (defined WIN32 || defined _WIN32 || defined WINCE) && defined CVAPI_EXPORTS\r
#define NCV_EXPORTS __declspec(dllexport)\r
#else\r
#define NCV_EXPORTS\r
#endif\r
\r
#include <cuda_runtime.h>\r
+#include <sstream>\r
+#include <iostream>\r
\r
\r
//==============================================================================\r
}\r
\r
\r
-#define NCV_CT_PREP_PASTE_AUX(a,b) a##b ///< Concatenation indirection macro\r
+#define NCV_CT_PREP_PASTE_AUX(a,b) a##b ///< Concatenation indirection macro\r
#define NCV_CT_PREP_PASTE(a,b) NCV_CT_PREP_PASTE_AUX(a, b) ///< Concatenation macro\r
\r
\r
Ncv32u height; ///< Rectangle height.\r
__host__ __device__ NcvSize32u() : width(0), height(0) {};\r
__host__ __device__ NcvSize32u(Ncv32u width, Ncv32u height) : width(width), height(height) {}\r
+ __host__ __device__ bool operator == (const NcvSize32u &another) const {return this->width == another.width && this->height == another.height;}\r
+};\r
+\r
+\r
+struct NcvPoint2D32s\r
+{\r
+ Ncv32s x; ///< Point X.\r
+ Ncv32s y; ///< Point Y.\r
+ __host__ __device__ NcvPoint2D32s() : x(0), y(0) {};\r
+ __host__ __device__ NcvPoint2D32s(Ncv32s x, Ncv32s y) : x(x), y(y) {}\r
+};\r
+\r
+\r
+struct NcvPoint2D32u\r
+{\r
+ Ncv32u x; ///< Point X.\r
+ Ncv32u y; ///< Point Y.\r
+ __host__ __device__ NcvPoint2D32u() : x(0), y(0) {};\r
+ __host__ __device__ NcvPoint2D32u(Ncv32u x, Ncv32u y) : x(x), y(y) {}\r
};\r
\r
\r
NCV_CT_ASSERT(sizeof(NcvRect32s) == 4 * sizeof(Ncv32s));\r
NCV_CT_ASSERT(sizeof(NcvRect32u) == 4 * sizeof(Ncv32u));\r
NCV_CT_ASSERT(sizeof(NcvSize32u) == 2 * sizeof(Ncv32u));\r
+NCV_CT_ASSERT(sizeof(NcvPoint2D32u) == 2 * sizeof(Ncv32u));\r
\r
\r
//==============================================================================\r
//==============================================================================\r
\r
\r
-#define NCV_CT_PREP_STRINGIZE_AUX(x) #x\r
-#define NCV_CT_PREP_STRINGIZE(x) NCV_CT_PREP_STRINGIZE_AUX(x)\r
-\r
+NCV_EXPORTS void ncvDebugOutput(const std::string &msg);\r
\r
-NCV_EXPORTS void ncvDebugOutput(const char *msg, ...);\r
\r
-\r
-typedef void NCVDebugOutputHandler(const char* msg);\r
+typedef void NCVDebugOutputHandler(const std::string &msg);\r
\r
\r
NCV_EXPORTS void ncvSetDebugOutputHandler(NCVDebugOutputHandler* func);\r
\r
\r
#define ncvAssertPrintCheck(pred, msg) \\r
- ((pred) ? true : (ncvDebugOutput("\n%s\n", \\r
- "NCV Assertion Failed: " msg ", file=" __FILE__ ", line=" NCV_CT_PREP_STRINGIZE(__LINE__) \\r
- ), false))\r
-\r
-\r
-#define ncvAssertPrintReturn(pred, msg, err) \\r
- if (ncvAssertPrintCheck(pred, msg)) ; else return err\r
-\r
-\r
-#define ncvAssertReturn(pred, err) \\r
do \\r
{ \\r
if (!(pred)) \\r
{ \\r
- ncvDebugOutput("\n%s%d%s\n", "NCV Assertion Failed: retcode=", (int)err, ", file=" __FILE__ ", line=" NCV_CT_PREP_STRINGIZE(__LINE__)); \\r
- return err; \\r
+ std::ostringstream oss; \\r
+ oss << "NCV Assertion Failed: " << msg << ", file=" << __FILE__ << ", line=" << __LINE__ << std::endl; \\r
+ ncvDebugOutput(oss.str()); \\r
} \\r
} while (0)\r
\r
\r
+#define ncvAssertPrintReturn(pred, msg, err) \\r
+ do \\r
+ { \\r
+ ncvAssertPrintCheck(pred, msg); \\r
+ if (!(pred)) return err; \\r
+ } while (0)\r
+\r
+\r
+#define ncvAssertReturn(pred, err) \\r
+ ncvAssertPrintReturn(pred, "retcode=" << (int)err, err)\r
+\r
+\r
#define ncvAssertReturnNcvStat(ncvOp) \\r
do \\r
{ \\r
NCVStatus _ncvStat = ncvOp; \\r
- if (NCV_SUCCESS != _ncvStat) \\r
- { \\r
- ncvDebugOutput("\n%s%d%s\n", "NCV Assertion Failed: NcvStat=", (int)_ncvStat, ", file=" __FILE__ ", line=" NCV_CT_PREP_STRINGIZE(__LINE__)); \\r
- return _ncvStat; \\r
- } \\r
+ ncvAssertPrintReturn(NCV_SUCCESS==_ncvStat, "NcvStat=" << (int)_ncvStat, _ncvStat); \\r
} while (0)\r
\r
\r
{ \\r
cudaError_t resCall = cudacall; \\r
cudaError_t resGLE = cudaGetLastError(); \\r
- if (cudaSuccess != resCall || cudaSuccess != resGLE) \\r
- { \\r
- ncvDebugOutput("\n%s%d%s\n", "NCV CUDA Assertion Failed: cudaError_t=", (int)(resCall | resGLE), ", file=" __FILE__ ", line=" NCV_CT_PREP_STRINGIZE(__LINE__)); \\r
- return errCode; \\r
- } \\r
+ ncvAssertPrintReturn(cudaSuccess==resCall && cudaSuccess==resGLE, "cudaError_t=" << (int)(resCall | resGLE), errCode); \\r
} while (0)\r
\r
\r
/**\r
* Return-codes for status notification, errors and warnings\r
*/\r
-enum NCVStatus\r
+enum\r
{\r
//NCV statuses\r
NCV_SUCCESS,\r
NPPST_MEM_INSUFFICIENT_BUFFER, ///< Insufficient user-allocated buffer\r
NPPST_MEM_RESIDENCE_ERROR, ///< Memory residence error detected (check if pointers should be device or pinned)\r
NPPST_MEM_INTERNAL_ERROR, ///< Internal memory management error\r
+\r
+ NCV_LAST_STATUS ///< Marker to continue error numeration in other files\r
};\r
\r
\r
+typedef Ncv32u NCVStatus;\r
+\r
+\r
#define NCV_SET_SKIP_COND(x) \\r
bool __ncv_skip_cond = x\r
\r
return ncvStat;\r
}\r
\r
+ T &at(Ncv32u x, Ncv32u y) const\r
+ {\r
+ if (x >= this->_width || y >= this->_height)\r
+ {\r
+ printf("Error addressing matrix at [%d, %d]\n", x, y);\r
+ return *this->_ptr;\r
+ }\r
+ return ((T *)((Ncv8u *)this->_ptr + y * this->_pitch))[x];\r
+ }\r
+\r
T *ptr() const {return this->_ptr;}\r
Ncv32u width() const {return this->_width;}\r
Ncv32u height() const {return this->_height;}\r
+ NcvSize32u size() const {return NcvSize32u(this->_width, this->_height);}\r
Ncv32u pitch() const {return this->_pitch;}\r
NCVMemoryType memType() const {return this->_memtype;}\r
\r
this->_width = roi.width;\r
this->_height = roi.height;\r
this->_pitch = mat.pitch();\r
- this->_ptr = mat.ptr() + roi.y * mat.stride() + roi.x;\r
+ this->_ptr = &mat.at(roi.x, roi.y);\r
this->_memtype = mat.memType();\r
\r
this->bReused = true;\r
NCV_EXPORTS NCVStatus ncvDrawRects_32u_device(Ncv32u *d_dst, Ncv32u dstStride, Ncv32u dstWidth, Ncv32u dstHeight,\r
NcvRect32u *d_rects, Ncv32u numRects, Ncv32u color, cudaStream_t cuStream);\r
\r
+\r
+#define CLAMP(x,a,b) ( (x) > (b) ? (b) : ( (x) < (a) ? (a) : (x) ) )\r
+#define CLAMP_TOP(x, a) (((x) > (a)) ? (a) : (x))\r
+#define CLAMP_BOTTOM(x, a) (((x) < (a)) ? (a) : (x))\r
+#define CLAMP_0_255(x) CLAMP(x,0,255)\r
+\r
+\r
+#define SUB_BEGIN(type, name) struct { __inline type name\r
+#define SUB_END(name) } name;\r
+#define SUB_CALL(name) name.name\r
+\r
+#define SQR(x) ((x)*(x))\r
+\r
+\r
+#define ncvSafeMatAlloc(name, type, alloc, width, height, err) \\r
+ NCVMatrixAlloc<type> name(alloc, width, height); \\r
+ ncvAssertReturn(name.isMemAllocated(), err);\r
+\r
+\r
+\r
#endif // _ncv_hpp_\r
--- /dev/null
+/*M///////////////////////////////////////////////////////////////////////////////////////\r
+//\r
+// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. \r
+// \r
+// By downloading, copying, installing or using the software you agree to this license.\r
+// If you do not agree to this license, do not download, install,\r
+// copy or use the software.\r
+//\r
+//\r
+// License Agreement\r
+// For Open Source Computer Vision Library\r
+//\r
+// Copyright (C) 2009-2010, NVIDIA Corporation, all rights reserved.\r
+// Third party copyrights are property of their respective owners.\r
+//\r
+// Redistribution and use in source and binary forms, with or without modification,\r
+// are permitted provided that the following conditions are met:\r
+//\r
+// * Redistribution's of source code must retain the above copyright notice,\r
+// this list of conditions and the following disclaimer.\r
+//\r
+// * Redistribution's in binary form must reproduce the above copyright notice,\r
+// this list of conditions and the following disclaimer in the documentation\r
+// and/or other materials provided with the distribution.\r
+//\r
+// * The name of the copyright holders may not be used to endorse or promote products\r
+// derived from this software without specific prior written permission.\r
+//\r
+// This software is provided by the copyright holders and contributors "as is" and\r
+// any express or implied warranties, including, but not limited to, the implied\r
+// warranties of merchantability and fitness for a particular purpose are disclaimed.\r
+// In no event shall the Intel Corporation or contributors be liable for any direct,\r
+// indirect, incidental, special, exemplary, or consequential damages\r
+// (including, but not limited to, procurement of substitute goods or services;\r
+// loss of use, data, or profits; or business interruption) however caused\r
+// and on any theory of liability, whether in contract, strict liability,\r
+// or tort (including negligence or otherwise) arising in any way out of\r
+// the use of this software, even if advised of the possibility of such damage.\r
+//\r
+//M*/\r
+\r
+#ifndef _ncv_color_conversion_hpp_\r
+#define _ncv_color_conversion_hpp_\r
+\r
+#include "NCVPixelOperations.hpp"\r
+\r
+enum NCVColorSpace\r
+{\r
+ NCVColorSpaceGray,\r
+ NCVColorSpaceRGBA,\r
+};\r
+\r
+template<NCVColorSpace CSin, NCVColorSpace CSout, typename Tin, typename Tout> struct __pixColorConv {\r
+static void _pixColorConv(const Tin &pixIn, Tout &pixOut);\r
+};\r
+\r
+template<typename Tin, typename Tout> struct __pixColorConv<NCVColorSpaceRGBA, NCVColorSpaceGray, Tin, Tout> {\r
+static void _pixColorConv(const Tin &pixIn, Tout &pixOut)\r
+{\r
+ Ncv32f luma = 0.299f * pixIn.x + 0.587f * pixIn.y + 0.114f * pixIn.z;\r
+ _TDemoteClampNN(luma, pixOut.x);\r
+}};\r
+\r
+template<typename Tin, typename Tout> struct __pixColorConv<NCVColorSpaceGray, NCVColorSpaceRGBA, Tin, Tout> {\r
+static void _pixColorConv(const Tin &pixIn, Tout &pixOut)\r
+{\r
+ _TDemoteClampNN(pixIn.x, pixOut.x);\r
+ _TDemoteClampNN(pixIn.x, pixOut.y);\r
+ _TDemoteClampNN(pixIn.x, pixOut.z);\r
+ pixOut.w = 0;\r
+}};\r
+\r
+template<NCVColorSpace CSin, NCVColorSpace CSout, typename Tin, typename Tout>\r
+static\r
+NCVStatus _ncvColorConv_host(const NCVMatrix<Tin> &h_imgIn,\r
+ const NCVMatrix<Tout> &h_imgOut)\r
+{\r
+ ncvAssertReturn(h_imgIn.size() == h_imgOut.size(), NCV_DIMENSIONS_INVALID);\r
+ ncvAssertReturn(h_imgIn.memType() == h_imgOut.memType() &&\r
+ (h_imgIn.memType() == NCVMemoryTypeHostPinned || h_imgIn.memType() == NCVMemoryTypeNone), NCV_MEM_RESIDENCE_ERROR);\r
+ NCV_SET_SKIP_COND(h_imgIn.memType() == NCVMemoryTypeNone);\r
+ NCV_SKIP_COND_BEGIN\r
+\r
+ for (Ncv32u i=0; i<h_imgIn.height(); i++)\r
+ {\r
+ for (Ncv32u j=0; j<h_imgIn.width(); j++)\r
+ {\r
+ __pixColorConv<CSin, CSout, Tin, Tout>::_pixColorConv(h_imgIn.at(j,i), h_imgOut.at(j,i));\r
+ }\r
+ }\r
+\r
+ NCV_SKIP_COND_END\r
+ return NCV_SUCCESS;\r
+}\r
+\r
+#endif //_ncv_color_conversion_hpp_\r
--- /dev/null
+/*M///////////////////////////////////////////////////////////////////////////////////////\r
+//\r
+// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. \r
+// \r
+// By downloading, copying, installing or using the software you agree to this license.\r
+// If you do not agree to this license, do not download, install,\r
+// copy or use the software.\r
+//\r
+//\r
+// License Agreement\r
+// For Open Source Computer Vision Library\r
+//\r
+// Copyright (C) 2009-2010, NVIDIA Corporation, all rights reserved.\r
+// Third party copyrights are property of their respective owners.\r
+//\r
+// Redistribution and use in source and binary forms, with or without modification,\r
+// are permitted provided that the following conditions are met:\r
+//\r
+// * Redistribution's of source code must retain the above copyright notice,\r
+// this list of conditions and the following disclaimer.\r
+//\r
+// * Redistribution's in binary form must reproduce the above copyright notice,\r
+// this list of conditions and the following disclaimer in the documentation\r
+// and/or other materials provided with the distribution.\r
+//\r
+// * The name of the copyright holders may not be used to endorse or promote products\r
+// derived from this software without specific prior written permission.\r
+//\r
+// This software is provided by the copyright holders and contributors "as is" and\r
+// any express or implied warranties, including, but not limited to, the implied\r
+// warranties of merchantability and fitness for a particular purpose are disclaimed.\r
+// In no event shall the Intel Corporation or contributors be liable for any direct,\r
+// indirect, incidental, special, exemplary, or consequential damages\r
+// (including, but not limited to, procurement of substitute goods or services;\r
+// loss of use, data, or profits; or business interruption) however caused\r
+// and on any theory of liability, whether in contract, strict liability,\r
+// or tort (including negligence or otherwise) arising in any way out of\r
+// the use of this software, even if advised of the possibility of such damage.\r
+//\r
+//M*/\r
+\r
+#ifndef _ncv_pixel_operations_hpp_\r
+#define _ncv_pixel_operations_hpp_\r
+\r
+#include <limits.h>\r
+#include <float.h>\r
+#include "NCV.hpp"\r
+\r
+template<typename TBase> inline TBase _pixMaxVal();\r
+template<> static inline Ncv8u _pixMaxVal<Ncv8u>() {return UCHAR_MAX;}\r
+template<> static inline Ncv16u _pixMaxVal<Ncv16u>() {return USHRT_MAX;}\r
+template<> static inline Ncv32u _pixMaxVal<Ncv32u>() {return UINT_MAX;}\r
+template<> static inline Ncv8s _pixMaxVal<Ncv8s>() {return CHAR_MAX;}\r
+template<> static inline Ncv16s _pixMaxVal<Ncv16s>() {return SHRT_MAX;}\r
+template<> static inline Ncv32s _pixMaxVal<Ncv32s>() {return INT_MAX;}\r
+template<> static inline Ncv32f _pixMaxVal<Ncv32f>() {return FLT_MAX;}\r
+template<> static inline Ncv64f _pixMaxVal<Ncv64f>() {return DBL_MAX;}\r
+\r
+template<typename TBase> inline TBase _pixMinVal();\r
+template<> static inline Ncv8u _pixMinVal<Ncv8u>() {return 0;}\r
+template<> static inline Ncv16u _pixMinVal<Ncv16u>() {return 0;}\r
+template<> static inline Ncv32u _pixMinVal<Ncv32u>() {return 0;}\r
+template<> static inline Ncv8s _pixMinVal<Ncv8s>() {return CHAR_MIN;}\r
+template<> static inline Ncv16s _pixMinVal<Ncv16s>() {return SHRT_MIN;}\r
+template<> static inline Ncv32s _pixMinVal<Ncv32s>() {return INT_MIN;}\r
+template<> static inline Ncv32f _pixMinVal<Ncv32f>() {return FLT_MIN;}\r
+template<> static inline Ncv64f _pixMinVal<Ncv64f>() {return DBL_MIN;}\r
+\r
+template<typename Tvec> struct TConvVec2Base;\r
+template<> struct TConvVec2Base<uchar1> {typedef Ncv8u TBase;};\r
+template<> struct TConvVec2Base<uchar3> {typedef Ncv8u TBase;};\r
+template<> struct TConvVec2Base<uchar4> {typedef Ncv8u TBase;};\r
+template<> struct TConvVec2Base<ushort1> {typedef Ncv16u TBase;};\r
+template<> struct TConvVec2Base<ushort3> {typedef Ncv16u TBase;};\r
+template<> struct TConvVec2Base<ushort4> {typedef Ncv16u TBase;};\r
+template<> struct TConvVec2Base<uint1> {typedef Ncv32u TBase;};\r
+template<> struct TConvVec2Base<uint3> {typedef Ncv32u TBase;};\r
+template<> struct TConvVec2Base<uint4> {typedef Ncv32u TBase;};\r
+template<> struct TConvVec2Base<float1> {typedef Ncv32f TBase;};\r
+template<> struct TConvVec2Base<float3> {typedef Ncv32f TBase;};\r
+template<> struct TConvVec2Base<float4> {typedef Ncv32f TBase;};\r
+template<> struct TConvVec2Base<double1> {typedef Ncv64f TBase;};\r
+template<> struct TConvVec2Base<double3> {typedef Ncv64f TBase;};\r
+template<> struct TConvVec2Base<double4> {typedef Ncv64f TBase;};\r
+\r
+#define NC(T) (sizeof(T) / sizeof(TConvVec2Base<T>::TBase))\r
+\r
+template<typename TBase, Ncv32u NC> struct TConvBase2Vec;\r
+template<> struct TConvBase2Vec<Ncv8u, 1> {typedef uchar1 TVec;};\r
+template<> struct TConvBase2Vec<Ncv8u, 3> {typedef uchar3 TVec;};\r
+template<> struct TConvBase2Vec<Ncv8u, 4> {typedef uchar4 TVec;};\r
+template<> struct TConvBase2Vec<Ncv16u, 1> {typedef ushort1 TVec;};\r
+template<> struct TConvBase2Vec<Ncv16u, 3> {typedef ushort3 TVec;};\r
+template<> struct TConvBase2Vec<Ncv16u, 4> {typedef ushort4 TVec;};\r
+template<> struct TConvBase2Vec<Ncv32u, 1> {typedef uint1 TVec;};\r
+template<> struct TConvBase2Vec<Ncv32u, 3> {typedef uint3 TVec;};\r
+template<> struct TConvBase2Vec<Ncv32u, 4> {typedef uint4 TVec;};\r
+template<> struct TConvBase2Vec<Ncv32f, 1> {typedef float1 TVec;};\r
+template<> struct TConvBase2Vec<Ncv32f, 3> {typedef float3 TVec;};\r
+template<> struct TConvBase2Vec<Ncv32f, 4> {typedef float4 TVec;};\r
+template<> struct TConvBase2Vec<Ncv64f, 1> {typedef double1 TVec;};\r
+template<> struct TConvBase2Vec<Ncv64f, 3> {typedef double3 TVec;};\r
+template<> struct TConvBase2Vec<Ncv64f, 4> {typedef double4 TVec;};\r
+\r
+//TODO: consider using CUDA intrinsics to avoid branching\r
+template<typename Tin> static inline void _TDemoteClampZ(Tin &a, Ncv8u &out) {out = (Ncv8u)CLAMP_0_255(a);};\r
+template<typename Tin> static inline void _TDemoteClampZ(Tin &a, Ncv16u &out) {out = (Ncv16u)CLAMP(a, 0, USHRT_MAX);}\r
+template<typename Tin> static inline void _TDemoteClampZ(Tin &a, Ncv32u &out) {out = (Ncv32u)CLAMP(a, 0, UINT_MAX);}\r
+template<typename Tin> static inline void _TDemoteClampZ(Tin &a, Ncv32f &out) {out = (Ncv32f)a;}\r
+\r
+//TODO: consider using CUDA intrinsics to avoid branching\r
+template<typename Tin> static inline void _TDemoteClampNN(Tin &a, Ncv8u &out) {out = (Ncv8u)CLAMP_0_255(a+0.5f);}\r
+template<typename Tin> static inline void _TDemoteClampNN(Tin &a, Ncv16u &out) {out = (Ncv16u)CLAMP(a+0.5f, 0, USHRT_MAX);}\r
+template<typename Tin> static inline void _TDemoteClampNN(Tin &a, Ncv32u &out) {out = (Ncv32u)CLAMP(a+0.5f, 0, UINT_MAX);}\r
+template<typename Tin> static inline void _TDemoteClampNN(Tin &a, Ncv32f &out) {out = (Ncv32f)a;}\r
+\r
+template<typename Tout> inline Tout _pixMakeZero();\r
+template<> static inline uchar1 _pixMakeZero<uchar1>() {return make_uchar1(0);}\r
+template<> static inline uchar3 _pixMakeZero<uchar3>() {return make_uchar3(0,0,0);}\r
+template<> static inline uchar4 _pixMakeZero<uchar4>() {return make_uchar4(0,0,0,0);}\r
+template<> static inline ushort1 _pixMakeZero<ushort1>() {return make_ushort1(0);}\r
+template<> static inline ushort3 _pixMakeZero<ushort3>() {return make_ushort3(0,0,0);}\r
+template<> static inline ushort4 _pixMakeZero<ushort4>() {return make_ushort4(0,0,0,0);}\r
+template<> static inline uint1 _pixMakeZero<uint1>() {return make_uint1(0);}\r
+template<> static inline uint3 _pixMakeZero<uint3>() {return make_uint3(0,0,0);}\r
+template<> static inline uint4 _pixMakeZero<uint4>() {return make_uint4(0,0,0,0);}\r
+template<> static inline float1 _pixMakeZero<float1>() {return make_float1(0.f);}\r
+template<> static inline float3 _pixMakeZero<float3>() {return make_float3(0.f,0.f,0.f);}\r
+template<> static inline float4 _pixMakeZero<float4>() {return make_float4(0.f,0.f,0.f,0.f);}\r
+template<> static inline double1 _pixMakeZero<double1>() {return make_double1(0.);}\r
+template<> static inline double3 _pixMakeZero<double3>() {return make_double3(0.,0.,0.);}\r
+template<> static inline double4 _pixMakeZero<double4>() {return make_double4(0.,0.,0.,0.);}\r
+\r
+static inline uchar1 _pixMake(Ncv8u x) {return make_uchar1(x);}\r
+static inline uchar3 _pixMake(Ncv8u x, Ncv8u y, Ncv8u z) {return make_uchar3(x,y,z);}\r
+static inline uchar4 _pixMake(Ncv8u x, Ncv8u y, Ncv8u z, Ncv8u w) {return make_uchar4(x,y,z,w);}\r
+static inline ushort1 _pixMake(Ncv16u x) {return make_ushort1(x);}\r
+static inline ushort3 _pixMake(Ncv16u x, Ncv16u y, Ncv16u z) {return make_ushort3(x,y,z);}\r
+static inline ushort4 _pixMake(Ncv16u x, Ncv16u y, Ncv16u z, Ncv16u w) {return make_ushort4(x,y,z,w);}\r
+static inline uint1 _pixMake(Ncv32u x) {return make_uint1(x);}\r
+static inline uint3 _pixMake(Ncv32u x, Ncv32u y, Ncv32u z) {return make_uint3(x,y,z);}\r
+static inline uint4 _pixMake(Ncv32u x, Ncv32u y, Ncv32u z, Ncv32u w) {return make_uint4(x,y,z,w);}\r
+static inline float1 _pixMake(Ncv32f x) {return make_float1(x);}\r
+static inline float3 _pixMake(Ncv32f x, Ncv32f y, Ncv32f z) {return make_float3(x,y,z);}\r
+static inline float4 _pixMake(Ncv32f x, Ncv32f y, Ncv32f z, Ncv32f w) {return make_float4(x,y,z,w);}\r
+static inline double1 _pixMake(Ncv64f x) {return make_double1(x);}\r
+static inline double3 _pixMake(Ncv64f x, Ncv64f y, Ncv64f z) {return make_double3(x,y,z);}\r
+static inline double4 _pixMake(Ncv64f x, Ncv64f y, Ncv64f z, Ncv64f w) {return make_double4(x,y,z,w);}\r
+\r
+\r
+template<typename Tin, typename Tout, Ncv32u CN> struct __pixDemoteClampZ_CN {static Tout _pixDemoteClampZ_CN(Tin &pix);};\r
+\r
+template<typename Tin, typename Tout> struct __pixDemoteClampZ_CN<Tin, Tout, 1> {\r
+static Tout _pixDemoteClampZ_CN(Tin &pix)\r
+{\r
+ Tout out;\r
+ _TDemoteClampZ(pix.x, out.x);\r
+ return out;\r
+}};\r
+\r
+template<typename Tin, typename Tout> struct __pixDemoteClampZ_CN<Tin, Tout, 3> {\r
+static Tout _pixDemoteClampZ_CN(Tin &pix)\r
+{\r
+ Tout out;\r
+ _TDemoteClampZ(pix.x, out.x);\r
+ _TDemoteClampZ(pix.y, out.y);\r
+ _TDemoteClampZ(pix.z, out.z);\r
+ return out;\r
+}};\r
+\r
+template<typename Tin, typename Tout> struct __pixDemoteClampZ_CN<Tin, Tout, 4> {\r
+static Tout _pixDemoteClampZ_CN(Tin &pix)\r
+{\r
+ Tout out;\r
+ _TDemoteClampZ(pix.x, out.x);\r
+ _TDemoteClampZ(pix.y, out.y);\r
+ _TDemoteClampZ(pix.z, out.z);\r
+ _TDemoteClampZ(pix.w, out.w);\r
+ return out;\r
+}};\r
+\r
+template<typename Tin, typename Tout> static inline Tout _pixDemoteClampZ(Tin &pix)\r
+{\r
+ return __pixDemoteClampZ_CN<Tin, Tout, NC(Tin)>::_pixDemoteClampZ_CN(pix);\r
+}\r
+\r
+\r
+template<typename Tin, typename Tout, Ncv32u CN> struct __pixDemoteClampNN_CN {static Tout _pixDemoteClampNN_CN(Tin &pix);};\r
+\r
+template<typename Tin, typename Tout> struct __pixDemoteClampNN_CN<Tin, Tout, 1> {\r
+static Tout _pixDemoteClampNN_CN(Tin &pix)\r
+{\r
+ Tout out;\r
+ _TDemoteClampNN(pix.x, out.x);\r
+ return out;\r
+}};\r
+\r
+template<typename Tin, typename Tout> struct __pixDemoteClampNN_CN<Tin, Tout, 3> {\r
+static Tout _pixDemoteClampNN_CN(Tin &pix)\r
+{\r
+ Tout out;\r
+ _TDemoteClampNN(pix.x, out.x);\r
+ _TDemoteClampNN(pix.y, out.y);\r
+ _TDemoteClampNN(pix.z, out.z);\r
+ return out;\r
+}};\r
+\r
+template<typename Tin, typename Tout> struct __pixDemoteClampNN_CN<Tin, Tout, 4> {\r
+static Tout _pixDemoteClampNN_CN(Tin &pix)\r
+{\r
+ Tout out;\r
+ _TDemoteClampNN(pix.x, out.x);\r
+ _TDemoteClampNN(pix.y, out.y);\r
+ _TDemoteClampNN(pix.z, out.z);\r
+ _TDemoteClampNN(pix.w, out.w);\r
+ return out;\r
+}};\r
+\r
+template<typename Tin, typename Tout> static inline Tout _pixDemoteClampNN(Tin &pix)\r
+{\r
+ return __pixDemoteClampNN_CN<Tin, Tout, NC(Tin)>::_pixDemoteClampNN_CN(pix);\r
+}\r
+\r
+\r
+template<typename Tin, typename Tout, typename Tw, Ncv32u CN> struct __pixScale_CN {static Tout _pixScale_CN(Tin &pix, Tw w);};\r
+\r
+template<typename Tin, typename Tout, typename Tw> struct __pixScale_CN<Tin, Tout, Tw, 1> {\r
+static Tout _pixScale_CN(Tin &pix, Tw w)\r
+{\r
+ Tout out;\r
+ typedef typename TConvVec2Base<Tout>::TBase TBout;\r
+ out.x = (TBout)(pix.x * w);\r
+ return out;\r
+}};\r
+\r
+template<typename Tin, typename Tout, typename Tw> struct __pixScale_CN<Tin, Tout, Tw, 3> {\r
+static Tout _pixScale_CN(Tin &pix, Tw w)\r
+{\r
+ Tout out;\r
+ typedef typename TConvVec2Base<Tout>::TBase TBout;\r
+ out.x = (TBout)(pix.x * w);\r
+ out.y = (TBout)(pix.y * w);\r
+ out.z = (TBout)(pix.z * w);\r
+ return out;\r
+}};\r
+\r
+template<typename Tin, typename Tout, typename Tw> struct __pixScale_CN<Tin, Tout, Tw, 4> {\r
+static Tout _pixScale_CN(Tin &pix, Tw w)\r
+{\r
+ Tout out;\r
+ typedef typename TConvVec2Base<Tout>::TBase TBout;\r
+ out.x = (TBout)(pix.x * w);\r
+ out.y = (TBout)(pix.y * w);\r
+ out.z = (TBout)(pix.z * w);\r
+ out.w = (TBout)(pix.w * w);\r
+ return out;\r
+}};\r
+\r
+template<typename Tin, typename Tout, typename Tw> static Tout _pixScale(Tin &pix, Tw w)\r
+{\r
+ return __pixScale_CN<Tin, Tout, Tw, NC(Tin)>::_pixScale_CN(pix, w);\r
+}\r
+\r
+\r
+template<typename Tin, typename Tout, Ncv32u CN> struct __pixAdd_CN {static Tout _pixAdd_CN(Tout &pix1, Tin &pix2);};\r
+\r
+template<typename Tin, typename Tout> struct __pixAdd_CN<Tin, Tout, 1> {\r
+static Tout _pixAdd_CN(Tout &pix1, Tin &pix2)\r
+{\r
+ Tout out;\r
+ out.x = pix1.x + pix2.x;\r
+ return out;\r
+}};\r
+\r
+template<typename Tin, typename Tout> struct __pixAdd_CN<Tin, Tout, 3> {\r
+static Tout _pixAdd_CN(Tout &pix1, Tin &pix2)\r
+{\r
+ Tout out;\r
+ out.x = pix1.x + pix2.x;\r
+ out.y = pix1.y + pix2.y;\r
+ out.z = pix1.z + pix2.z;\r
+ return out;\r
+}};\r
+\r
+template<typename Tin, typename Tout> struct __pixAdd_CN<Tin, Tout, 4> {\r
+static Tout _pixAdd_CN(Tout &pix1, Tin &pix2)\r
+{\r
+ Tout out;\r
+ out.x = pix1.x + pix2.x;\r
+ out.y = pix1.y + pix2.y;\r
+ out.z = pix1.z + pix2.z;\r
+ out.w = pix1.w + pix2.w;\r
+ return out;\r
+}};\r
+\r
+template<typename Tin, typename Tout> static Tout _pixAdd(Tout &pix1, Tin &pix2)\r
+{\r
+ return __pixAdd_CN<Tin, Tout, NC(Tin)>::_pixAdd_CN(pix1, pix2);\r
+}\r
+\r
+\r
+template<typename Tin, typename Tout, Ncv32u CN> struct __pixDist_CN {static Tout _pixDist_CN(Tin &pix1, Tin &pix2);};\r
+\r
+template<typename Tin, typename Tout> struct __pixDist_CN<Tin, Tout, 1> {\r
+static Tout _pixDist_CN(Tin &pix1, Tin &pix2)\r
+{\r
+ return Tout(SQR(pix1.x - pix2.x));\r
+}};\r
+\r
+template<typename Tin, typename Tout> struct __pixDist_CN<Tin, Tout, 3> {\r
+static Tout _pixDist_CN(Tin &pix1, Tin &pix2)\r
+{\r
+ return Tout(SQR(pix1.x - pix2.x) + SQR(pix1.y - pix2.y) + SQR(pix1.z - pix2.z));\r
+}};\r
+\r
+template<typename Tin, typename Tout> struct __pixDist_CN<Tin, Tout, 4> {\r
+static Tout _pixDist_CN(Tin &pix1, Tin &pix2)\r
+{\r
+ return Tout(SQR(pix1.x - pix2.x) + SQR(pix1.y - pix2.y) + SQR(pix1.z - pix2.z) + SQR(pix1.w - pix2.w));\r
+}};\r
+\r
+template<typename Tin, typename Tout> static Tout _pixDist(Tin &pix1, Tin &pix2)\r
+{\r
+ return __pixDist_CN<Tin, Tout, NC(Tin)>::_pixDist_CN(pix1, pix2);\r
+}\r
+\r
+\r
+template <typename T> struct TAccPixWeighted;\r
+template<> struct TAccPixWeighted<uchar1> {typedef double1 type;};\r
+template<> struct TAccPixWeighted<uchar3> {typedef double3 type;};\r
+template<> struct TAccPixWeighted<uchar4> {typedef double4 type;};\r
+template<> struct TAccPixWeighted<ushort1> {typedef double1 type;};\r
+template<> struct TAccPixWeighted<ushort3> {typedef double3 type;};\r
+template<> struct TAccPixWeighted<ushort4> {typedef double4 type;};\r
+template<> struct TAccPixWeighted<float1> {typedef double1 type;};\r
+template<> struct TAccPixWeighted<float3> {typedef double3 type;};\r
+template<> struct TAccPixWeighted<float4> {typedef double4 type;};\r
+\r
+template<typename Tfrom> struct TAccPixDist {};\r
+template<> struct TAccPixDist<uchar1> {typedef Ncv32u type;};\r
+template<> struct TAccPixDist<uchar3> {typedef Ncv32u type;};\r
+template<> struct TAccPixDist<uchar4> {typedef Ncv32u type;};\r
+template<> struct TAccPixDist<ushort1> {typedef Ncv32u type;};\r
+template<> struct TAccPixDist<ushort3> {typedef Ncv32u type;};\r
+template<> struct TAccPixDist<ushort4> {typedef Ncv32u type;};\r
+template<> struct TAccPixDist<float1> {typedef Ncv32f type;};\r
+template<> struct TAccPixDist<float3> {typedef Ncv32f type;};\r
+template<> struct TAccPixDist<float4> {typedef Ncv32f type;};\r
+\r
+#endif //_ncv_pixel_operations_hpp_\r
--- /dev/null
+/*M///////////////////////////////////////////////////////////////////////////////////////\r
+//\r
+// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. \r
+// \r
+// By downloading, copying, installing or using the software you agree to this license.\r
+// If you do not agree to this license, do not download, install,\r
+// copy or use the software.\r
+//\r
+//\r
+// License Agreement\r
+// For Open Source Computer Vision Library\r
+//\r
+// Copyright (C) 2009-2010, NVIDIA Corporation, all rights reserved.\r
+// Third party copyrights are property of their respective owners.\r
+//\r
+// Redistribution and use in source and binary forms, with or without modification,\r
+// are permitted provided that the following conditions are met:\r
+//\r
+// * Redistribution's of source code must retain the above copyright notice,\r
+// this list of conditions and the following disclaimer.\r
+//\r
+// * Redistribution's in binary form must reproduce the above copyright notice,\r
+// this list of conditions and the following disclaimer in the documentation\r
+// and/or other materials provided with the distribution.\r
+//\r
+// * The name of the copyright holders may not be used to endorse or promote products\r
+// derived from this software without specific prior written permission.\r
+//\r
+// This software is provided by the copyright holders and contributors "as is" and\r
+// any express or implied warranties, including, but not limited to, the implied\r
+// warranties of merchantability and fitness for a particular purpose are disclaimed.\r
+// In no event shall the Intel Corporation or contributors be liable for any direct,\r
+// indirect, incidental, special, exemplary, or consequential damages\r
+// (including, but not limited to, procurement of substitute goods or services;\r
+// loss of use, data, or profits; or business interruption) however caused\r
+// and on any theory of liability, whether in contract, strict liability,\r
+// or tort (including negligence or otherwise) arising in any way out of\r
+// the use of this software, even if advised of the possibility of such damage.\r
+//\r
+//M*/\r
+\r
+\r
+#include <cuda_runtime.h>\r
+#include <stdio.h>\r
+#include "NCV.hpp"\r
+#include "NCVPyramid.hpp"\r
+#include "NCVPixelOperations.hpp"\r
+\r
+\r
+template<typename T, Ncv32u CN> struct __average4_CN {static T _average4_CN(const T &p00, const T &p01, const T &p10, const T &p11);};\r
+\r
+template<typename T> struct __average4_CN<T, 1> {\r
+static T _average4_CN(const T &p00, const T &p01, const T &p10, const T &p11)\r
+{\r
+ T out;\r
+ out.x = ((Ncv32s)p00.x + p01.x + p10.x + p11.x + 2) / 4;\r
+ return out;\r
+}};\r
+\r
+template<> struct __average4_CN<float1, 1> {\r
+static float1 _average4_CN(const float1 &p00, const float1 &p01, const float1 &p10, const float1 &p11)\r
+{\r
+ float1 out;\r
+ out.x = (p00.x + p01.x + p10.x + p11.x) / 4;\r
+ return out;\r
+}};\r
+\r
+template<> struct __average4_CN<double1, 1> {\r
+static double1 _average4_CN(const double1 &p00, const double1 &p01, const double1 &p10, const double1 &p11)\r
+{\r
+ double1 out;\r
+ out.x = (p00.x + p01.x + p10.x + p11.x) / 4;\r
+ return out;\r
+}};\r
+\r
+template<typename T> struct __average4_CN<T, 3> {\r
+static T _average4_CN(const T &p00, const T &p01, const T &p10, const T &p11)\r
+{\r
+ T out;\r
+ out.x = ((Ncv32s)p00.x + p01.x + p10.x + p11.x + 2) / 4;\r
+ out.y = ((Ncv32s)p00.y + p01.y + p10.y + p11.y + 2) / 4;\r
+ out.z = ((Ncv32s)p00.z + p01.z + p10.z + p11.z + 2) / 4;\r
+ return out;\r
+}};\r
+\r
+template<> struct __average4_CN<float3, 3> {\r
+static float3 _average4_CN(const float3 &p00, const float3 &p01, const float3 &p10, const float3 &p11)\r
+{\r
+ float3 out;\r
+ out.x = (p00.x + p01.x + p10.x + p11.x) / 4;\r
+ out.y = (p00.y + p01.y + p10.y + p11.y) / 4;\r
+ out.z = (p00.z + p01.z + p10.z + p11.z) / 4;\r
+ return out;\r
+}};\r
+\r
+template<> struct __average4_CN<double3, 3> {\r
+static double3 _average4_CN(const double3 &p00, const double3 &p01, const double3 &p10, const double3 &p11)\r
+{\r
+ double3 out;\r
+ out.x = (p00.x + p01.x + p10.x + p11.x) / 4;\r
+ out.y = (p00.y + p01.y + p10.y + p11.y) / 4;\r
+ out.z = (p00.z + p01.z + p10.z + p11.z) / 4;\r
+ return out;\r
+}};\r
+\r
+template<typename T> struct __average4_CN<T, 4> {\r
+static T _average4_CN(const T &p00, const T &p01, const T &p10, const T &p11)\r
+{\r
+ T out;\r
+ out.x = ((Ncv32s)p00.x + p01.x + p10.x + p11.x + 2) / 4;\r
+ out.y = ((Ncv32s)p00.y + p01.y + p10.y + p11.y + 2) / 4;\r
+ out.z = ((Ncv32s)p00.z + p01.z + p10.z + p11.z + 2) / 4;\r
+ out.w = ((Ncv32s)p00.w + p01.w + p10.w + p11.w + 2) / 4;\r
+ return out;\r
+}};\r
+\r
+template<> struct __average4_CN<float4, 4> {\r
+static float4 _average4_CN(const float4 &p00, const float4 &p01, const float4 &p10, const float4 &p11)\r
+{\r
+ float4 out;\r
+ out.x = (p00.x + p01.x + p10.x + p11.x) / 4;\r
+ out.y = (p00.y + p01.y + p10.y + p11.y) / 4;\r
+ out.z = (p00.z + p01.z + p10.z + p11.z) / 4;\r
+ out.w = (p00.w + p01.w + p10.w + p11.w) / 4;\r
+ return out;\r
+}};\r
+\r
+template<> struct __average4_CN<double4, 4> {\r
+static double4 _average4_CN(const double4 &p00, const double4 &p01, const double4 &p10, const double4 &p11)\r
+{\r
+ double4 out;\r
+ out.x = (p00.x + p01.x + p10.x + p11.x) / 4;\r
+ out.y = (p00.y + p01.y + p10.y + p11.y) / 4;\r
+ out.z = (p00.z + p01.z + p10.z + p11.z) / 4;\r
+ out.w = (p00.w + p01.w + p10.w + p11.w) / 4;\r
+ return out;\r
+}};\r
+\r
+template<typename T> static T _average4(const T &p00, const T &p01, const T &p10, const T &p11)\r
+{\r
+ return __average4_CN<T, NC(T)>::_average4_CN(p00, p01, p10, p11);\r
+}\r
+\r
+\r
+template<typename Tin, typename Tout, Ncv32u CN> struct __lerp_CN {static Tout _lerp_CN(const Tin &a, const Tin &b, Ncv32f d);};\r
+\r
+template<typename Tin, typename Tout> struct __lerp_CN<Tin, Tout, 1> {\r
+static Tout _lerp_CN(const Tin &a, const Tin &b, Ncv32f d)\r
+{\r
+ typedef typename TConvVec2Base<Tout>::TBase TB;\r
+ return _pixMake(TB(b.x * d + a.x * (1 - d)));\r
+}};\r
+\r
+template<typename Tin, typename Tout> struct __lerp_CN<Tin, Tout, 3> {\r
+static Tout _lerp_CN(const Tin &a, const Tin &b, Ncv32f d)\r
+{\r
+ typedef typename TConvVec2Base<Tout>::TBase TB;\r
+ return _pixMake(TB(b.x * d + a.x * (1 - d)),\r
+ TB(b.y * d + a.y * (1 - d)),\r
+ TB(b.z * d + a.z * (1 - d)));\r
+}};\r
+\r
+template<typename Tin, typename Tout> struct __lerp_CN<Tin, Tout, 4> {\r
+static Tout _lerp_CN(const Tin &a, const Tin &b, Ncv32f d)\r
+{\r
+ typedef typename TConvVec2Base<Tout>::TBase TB;\r
+ return _pixMake(TB(b.x * d + a.x * (1 - d)),\r
+ TB(b.y * d + a.y * (1 - d)),\r
+ TB(b.z * d + a.z * (1 - d)),\r
+ TB(b.w * d + a.w * (1 - d)));\r
+}};\r
+\r
+template<typename Tin, typename Tout> static Tout _lerp(const Tin &a, const Tin &b, Ncv32f d)\r
+{\r
+ return __lerp_CN<Tin, Tout, NC(Tin)>::_lerp_CN(a, b, d);\r
+}\r
+\r
+\r
+template<typename T>\r
+static T _interpLinear(const T &a, const T &b, Ncv32f d)\r
+{\r
+ typedef typename TConvBase2Vec<Ncv32f, NC(T)>::TVec TVFlt;\r
+ TVFlt tmp = _lerp<T, TVFlt>(a, b, d);\r
+ return _pixDemoteClampZ<TVFlt, T>(tmp);\r
+}\r
+\r
+\r
+template<typename T>\r
+static T _interpBilinear(const NCVMatrix<T> &refLayer, Ncv32f x, Ncv32f y)\r
+{\r
+ Ncv32u xl = (Ncv32u)x;\r
+ Ncv32u xh = xl+1;\r
+ Ncv32f dx = x - xl;\r
+ Ncv32u yl = (Ncv32u)y;\r
+ Ncv32u yh = yl+1;\r
+ Ncv32f dy = y - yl;\r
+ T p00, p01, p10, p11;\r
+ p00 = refLayer.at(xl, yl);\r
+ p01 = xh < refLayer.width() ? refLayer.at(xh, yl) : p00;\r
+ p10 = yh < refLayer.height() ? refLayer.at(xl, yh) : p00;\r
+ p11 = (xh < refLayer.width() && yh < refLayer.height()) ? refLayer.at(xh, yh) : p00;\r
+ typedef typename TConvBase2Vec<Ncv32f, NC(T)>::TVec TVFlt;\r
+ TVFlt m_00_01 = _lerp<T, TVFlt>(p00, p01, dx);\r
+ TVFlt m_10_11 = _lerp<T, TVFlt>(p10, p11, dx);\r
+ TVFlt mixture = _lerp<TVFlt, TVFlt>(m_00_01, m_10_11, dy);\r
+ return _pixDemoteClampZ<TVFlt, T>(mixture);\r
+}\r
+\r
+\r
+template <class T>\r
+NCVImagePyramid<T>::NCVImagePyramid(const NCVMatrix<T> &img,\r
+ Ncv8u numLayers,\r
+ INCVMemAllocator &alloc,\r
+ cudaStream_t cuStream)\r
+{\r
+ this->_isInitialized = false;\r
+ ncvAssertPrintReturn(img.memType() == alloc.memType(), "NCVImagePyramid_host::ctor error", );\r
+\r
+ this->layer0 = &img;\r
+ NcvSize32u szLastLayer(img.width(), img.height());\r
+ this->nLayers = 1;\r
+\r
+ NCV_SET_SKIP_COND(alloc.isCounting());\r
+ NcvBool bDeviceCode = alloc.memType() == NCVMemoryTypeDevice;\r
+\r
+ if (numLayers == 0)\r
+ {\r
+ numLayers = 255; //it will cut-off when any of the dimensions goes 1\r
+ }\r
+\r
+ for (Ncv32u i=0; i<(Ncv32u)numLayers-1; i++)\r
+ {\r
+ NcvSize32u szCurLayer(szLastLayer.width / 2, szLastLayer.height / 2);\r
+ if (szCurLayer.width == 0 || szCurLayer.height == 0)\r
+ {\r
+ break;\r
+ }\r
+\r
+ this->pyramid.push_back(new NCVMatrixAlloc<T>(alloc, szCurLayer.width, szCurLayer.height));\r
+ ncvAssertPrintReturn(((NCVMatrixAlloc<T> *)(this->pyramid[i]))->isMemAllocated(), "NCVImagePyramid_host::ctor error", );\r
+ this->nLayers++;\r
+\r
+ //fill in the layer\r
+ NCV_SKIP_COND_BEGIN\r
+\r
+ const NCVMatrix<T> *prevLayer = i == 0 ? this->layer0 : this->pyramid[i-1];\r
+ NCVMatrix<T> *curLayer = this->pyramid[i];\r
+\r
+ if (bDeviceCode)\r
+ {\r
+ //TODO: in cuStream\r
+ }\r
+ else\r
+ {\r
+ for (Ncv32u i=0; i<szCurLayer.height; i++)\r
+ {\r
+ for (Ncv32u j=0; j<szCurLayer.width; j++)\r
+ {\r
+ T p00 = prevLayer->at(2*j+0, 2*i+0);\r
+ T p01 = prevLayer->at(2*j+1, 2*i+0);\r
+ T p10 = prevLayer->at(2*j+0, 2*i+1);\r
+ T p11 = prevLayer->at(2*j+1, 2*i+1);\r
+ curLayer->at(j, i) = _average4(p00, p01, p10, p11);\r
+ }\r
+ }\r
+ }\r
+\r
+ NCV_SKIP_COND_END\r
+\r
+ szLastLayer = szCurLayer;\r
+ }\r
+\r
+ this->_isInitialized = true;\r
+}\r
+\r
+\r
+template <class T>\r
+NCVImagePyramid<T>::~NCVImagePyramid()\r
+{\r
+}\r
+\r
+\r
+template <class T>\r
+NcvBool NCVImagePyramid<T>::isInitialized() const\r
+{\r
+ return this->_isInitialized;\r
+}\r
+\r
+\r
+template <class T>\r
+NCVStatus NCVImagePyramid<T>::getLayer(NCVMatrix<T> &outImg,\r
+ NcvSize32u outRoi,\r
+ NcvBool bTrilinear,\r
+ cudaStream_t cuStream) const\r
+{\r
+ ncvAssertReturn(this->isInitialized(), NCV_UNKNOWN_ERROR);\r
+ ncvAssertReturn(outImg.memType() == this->layer0->memType(), NCV_MEM_RESIDENCE_ERROR);\r
+ ncvAssertReturn(outRoi.width <= this->layer0->width() && outRoi.height <= this->layer0->height() &&\r
+ outRoi.width > 0 && outRoi.height > 0, NCV_DIMENSIONS_INVALID);\r
+\r
+ if (outRoi.width == this->layer0->width() && outRoi.height == this->layer0->height())\r
+ {\r
+ ncvAssertReturnNcvStat(this->layer0->copy2D(outImg, NcvSize32u(this->layer0->width(), this->layer0->height()), cuStream));\r
+ return NCV_SUCCESS;\r
+ }\r
+\r
+ Ncv32f lastScale = 1.0f;\r
+ Ncv32f curScale;\r
+ const NCVMatrix<T> *lastLayer = this->layer0;\r
+ const NCVMatrix<T> *curLayer = NULL;\r
+ NcvBool bUse2Refs = false;\r
+\r
+ for (Ncv32u i=0; i<this->nLayers-1; i++)\r
+ {\r
+ curScale = lastScale * 0.5f;\r
+ curLayer = this->pyramid[i];\r
+\r
+ if (outRoi.width == curLayer->width() && outRoi.height == curLayer->height())\r
+ {\r
+ ncvAssertReturnNcvStat(this->pyramid[i]->copy2D(outImg, NcvSize32u(this->pyramid[i]->width(), this->pyramid[i]->height()), cuStream));\r
+ return NCV_SUCCESS;\r
+ }\r
+\r
+ if (outRoi.width >= curLayer->width() && outRoi.height >= curLayer->height())\r
+ {\r
+ if (outRoi.width < lastLayer->width() && outRoi.height < lastLayer->height())\r
+ {\r
+ bUse2Refs = true;\r
+ }\r
+ break;\r
+ }\r
+\r
+ lastScale = curScale;\r
+ lastLayer = curLayer;\r
+ }\r
+\r
+ bUse2Refs = bUse2Refs && bTrilinear;\r
+\r
+ NCV_SET_SKIP_COND(outImg.memType() == NCVMemoryTypeNone);\r
+ NcvBool bDeviceCode = this->layer0->memType() == NCVMemoryTypeDevice;\r
+\r
+ NCV_SKIP_COND_BEGIN\r
+\r
+ if (bDeviceCode)\r
+ {\r
+ //TODO: in cuStream\r
+ }\r
+ else\r
+ {\r
+ for (Ncv32u i=0; i<outRoi.height; i++)\r
+ {\r
+ for (Ncv32u j=0; j<outRoi.width; j++)\r
+ {\r
+ //top layer pixel (always exists)\r
+ NcvSize32u szTopLayer(lastLayer->width(), lastLayer->height());\r
+ Ncv32f ptTopX = 1.0f * (szTopLayer.width - 1) * j / (outRoi.width - 1);\r
+ Ncv32f ptTopY = 1.0f * (szTopLayer.height - 1) * i / (outRoi.height - 1);\r
+ T topPix = _interpBilinear(*lastLayer, ptTopX, ptTopY);\r
+ T trilinearPix = topPix;\r
+\r
+ if (bUse2Refs)\r
+ {\r
+ //bottom layer pixel (exists only if the requested scale is greater than the smallest layer scale)\r
+ NcvSize32u szBottomLayer(curLayer->width(), curLayer->height());\r
+ Ncv32f ptBottomX = 1.0f * (szBottomLayer.width - 1) * j / (outRoi.width - 1);\r
+ Ncv32f ptBottomY = 1.0f * (szBottomLayer.height - 1) * i / (outRoi.height - 1);\r
+ T bottomPix = _interpBilinear(*curLayer, ptBottomX, ptBottomY);\r
+\r
+ Ncv32f scale = (1.0f * outRoi.width / layer0->width() + 1.0f * outRoi.height / layer0->height()) / 2;\r
+ Ncv32f dl = (scale - curScale) / (lastScale - curScale);\r
+ dl = CLAMP(dl, 0.0f, 1.0f);\r
+ trilinearPix = _interpLinear(bottomPix, topPix, dl);\r
+ }\r
+\r
+ outImg.at(j, i) = trilinearPix;\r
+ }\r
+ }\r
+ }\r
+\r
+ NCV_SKIP_COND_END\r
+\r
+ return NCV_SUCCESS;\r
+}\r
+\r
+\r
+template class NCVImagePyramid<uchar1>;\r
+template class NCVImagePyramid<uchar3>;\r
+template class NCVImagePyramid<uchar4>;\r
+template class NCVImagePyramid<ushort1>;\r
+template class NCVImagePyramid<ushort3>;\r
+template class NCVImagePyramid<ushort4>;\r
+template class NCVImagePyramid<uint1>;\r
+template class NCVImagePyramid<uint3>;\r
+template class NCVImagePyramid<uint4>;\r
+template class NCVImagePyramid<float1>;\r
+template class NCVImagePyramid<float3>;\r
+template class NCVImagePyramid<float4>;\r
--- /dev/null
+/*M///////////////////////////////////////////////////////////////////////////////////////\r
+//\r
+// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. \r
+// \r
+// By downloading, copying, installing or using the software you agree to this license.\r
+// If you do not agree to this license, do not download, install,\r
+// copy or use the software.\r
+//\r
+//\r
+// License Agreement\r
+// For Open Source Computer Vision Library\r
+//\r
+// Copyright (C) 2009-2010, NVIDIA Corporation, all rights reserved.\r
+// Third party copyrights are property of their respective owners.\r
+//\r
+// Redistribution and use in source and binary forms, with or without modification,\r
+// are permitted provided that the following conditions are met:\r
+//\r
+// * Redistribution's of source code must retain the above copyright notice,\r
+// this list of conditions and the following disclaimer.\r
+//\r
+// * Redistribution's in binary form must reproduce the above copyright notice,\r
+// this list of conditions and the following disclaimer in the documentation\r
+// and/or other materials provided with the distribution.\r
+//\r
+// * The name of the copyright holders may not be used to endorse or promote products\r
+// derived from this software without specific prior written permission.\r
+//\r
+// This software is provided by the copyright holders and contributors "as is" and\r
+// any express or implied warranties, including, but not limited to, the implied\r
+// warranties of merchantability and fitness for a particular purpose are disclaimed.\r
+// In no event shall the Intel Corporation or contributors be liable for any direct,\r
+// indirect, incidental, special, exemplary, or consequential damages\r
+// (including, but not limited to, procurement of substitute goods or services;\r
+// loss of use, data, or profits; or business interruption) however caused\r
+// and on any theory of liability, whether in contract, strict liability,\r
+// or tort (including negligence or otherwise) arising in any way out of\r
+// the use of this software, even if advised of the possibility of such damage.\r
+//\r
+//M*/\r
+\r
+\r
+#ifndef _ncvpyramid_hpp_\r
+#define _ncvpyramid_hpp_\r
+\r
+#include <memory>\r
+#include <vector>\r
+#include "NCV.hpp"\r
+\r
+\r
+template <class T>\r
+class NCV_EXPORTS NCVMatrixStack\r
+{\r
+public:\r
+ NCVMatrixStack() {this->_arr.clear();}\r
+ ~NCVMatrixStack()\r
+ {\r
+ const Ncv32u nElem = this->_arr.size();\r
+ for (Ncv32u i=0; i<nElem; i++)\r
+ {\r
+ pop_back();\r
+ }\r
+ }\r
+ void push_back(NCVMatrix<T> *elem) {this->_arr.push_back(std::tr1::shared_ptr< NCVMatrix<T> >(elem));}\r
+ void pop_back() {this->_arr.pop_back();}\r
+ NCVMatrix<T> * operator [] (int i) const {return this->_arr[i].get();}\r
+private:\r
+ std::vector< std::tr1::shared_ptr< NCVMatrix<T> > > _arr;\r
+};\r
+\r
+\r
+template <class T>\r
+class NCV_EXPORTS NCVImagePyramid\r
+{\r
+public:\r
+\r
+ NCVImagePyramid(const NCVMatrix<T> &img,\r
+ Ncv8u nLayers,\r
+ INCVMemAllocator &alloc,\r
+ cudaStream_t cuStream);\r
+ ~NCVImagePyramid();\r
+ NcvBool isInitialized() const;\r
+ NCVStatus getLayer(NCVMatrix<T> &outImg,\r
+ NcvSize32u outRoi,\r
+ NcvBool bTrilinear,\r
+ cudaStream_t cuStream) const;\r
+\r
+private:\r
+\r
+ NcvBool _isInitialized;\r
+ const NCVMatrix<T> *layer0;\r
+ NCVMatrixStack<T> pyramid;\r
+ Ncv32u nLayers;\r
+};\r
+\r
+\r
+#endif //_ncvpyramid_hpp_\r
\r
namespace\r
{\r
- void outputHandler(const char* msg) \r
- { \r
- CV_Error(CV_GpuApiCallError, msg); \r
- }\r
+ static void outputHandler(const std::string &msg) { CV_Error(CV_GpuApiCallError, msg.c_str()); }\r
}\r
\r
void cv::gpu::BroxOpticalFlow::operator ()(const GpuMat& frame0, const GpuMat& frame1, GpuMat& u, GpuMat& v, Stream& s) \r
projects / profile / ivi / opencv.git / commitdiff