some device layer utility functions

diff --git a/modules/gpu/src/cuda/dynamic_smem.hpp b/modules/gpu/src/cuda/dynamic_smem.hpp
new file mode 100644 (file)
index 0000000..d420c1e
--- /dev/null
+++ b/modules/gpu/src/cuda/dynamic_smem.hpp
@@ -0,0 +1,83 @@
+/*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) 2000-2008, Intel Corporation, all rights reserved.\r
+// Copyright (C) 2009, Willow Garage Inc., 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
+namespace cv\r
+{\r
+    namespace gpu\r
+    {\r
+        namespace device\r
+        {   \r
+            template<class T> struct DynamicSharedMem\r
+            {\r
+                __device__ operator T*()\r
+                {\r
+                    extern __shared__ int __smem[];\r
+                    return (T*)__smem;\r
+                }\r
+\r
+                __device__ operator const T*() const\r
+                {\r
+                    extern __shared__ int __smem[];\r
+                    return (T*)__smem;\r
+                }\r
+            };\r
+\r
+            // specialize for double to avoid unaligned memory access compile errors\r
+            template<> struct DynamicSharedMem<double>\r
+            {\r
+                __device__ operator double*()\r
+                {\r
+                    extern __shared__ double __smem_d[];\r
+                    return (double*)__smem_d;\r
+                }\r
+\r
+                __device__ operator const double*() const\r
+                {\r
+                    extern __shared__ double __smem_d[];\r
+                    return (double*)__smem_d;\r
+                }\r
+            };\r
+        }\r
+\r
+    }\r
+}
\ No newline at end of file

diff --git a/modules/gpu/src/cuda/limits_gpu.hpp b/modules/gpu/src/cuda/limits_gpu.hpp
new file mode 100644 (file)
index 0000000..83a0743
--- /dev/null
+++ b/modules/gpu/src/cuda/limits_gpu.hpp
@@ -0,0 +1,208 @@
+/*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) 2000-2008, Intel Corporation, all rights reserved.\r
+// Copyright (C) 2009, Willow Garage Inc., 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
+namespace cv\r
+{\r
+    namespace gpu\r
+    {\r
+        namespace device\r
+        {     \r
+            template<class T> struct numeric_limits_gpu\r
+            {          \r
+                typedef T type;\r
+                __device__ static type min()  { return type(); };            \r
+                __device__ static type max() { return type(); };\r
+                __device__ static type epsilon() { return type(); }\r
+                __device__ static type round_error() { return type(); }\r
+                __device__ static type denorm_min()  { return type(); }\r
+                __device__ static type infinity() { return type(); }\r
+                __device__ static type quiet_NaN() { return type(); }\r
+                __device__ static type signaling_NaN() { return T(); }\r
+            };\r
+\r
+            template<> struct numeric_limits_gpu<bool>\r
+            {          \r
+                typedef bool type;\r
+                __device__ static type min() { return false; };            \r
+                __device__ static type max() { return true;  };\r
+                __device__ static type epsilon();\r
+                __device__ static type round_error();\r
+                __device__ static type denorm_min();\r
+                __device__ static type infinity();\r
+                __device__ static type quiet_NaN();\r
+                __device__ static type signaling_NaN();\r
+            };\r
+\r
+            template<> struct numeric_limits_gpu<char>\r
+            {          \r
+                typedef char type;\r
+                __device__ static type min() { return CHAR_MIN; };            \r
+                __device__ static type max() { return CHAR_MAX; };\r
+                __device__ static type epsilon();\r
+                __device__ static type round_error();\r
+                __device__ static type denorm_min();\r
+                __device__ static type infinity();\r
+                __device__ static type quiet_NaN();\r
+                __device__ static type signaling_NaN();\r
+            };\r
+\r
+            template<> struct numeric_limits_gpu<unsigned char>\r
+            {          \r
+                typedef unsigned char type;\r
+                __device__ static type min() { return 0; };            \r
+                __device__ static type max() { return UCHAR_MAX; };\r
+                __device__ static type epsilon();\r
+                __device__ static type round_error();\r
+                __device__ static type denorm_min();\r
+                __device__ static type infinity();\r
+                __device__ static type quiet_NaN();\r
+                __device__ static type signaling_NaN();\r
+            };\r
+\r
+            template<> struct numeric_limits_gpu<short>\r
+            {  \r
+                typedef short type;\r
+                __device__ static type min() { return SHRT_MIN; };            \r
+                __device__ static type max() { return SHRT_MAX; };\r
+                __device__ static type epsilon();\r
+                __device__ static type round_error();\r
+                __device__ static type denorm_min();\r
+                __device__ static type infinity();\r
+                __device__ static type quiet_NaN();\r
+                __device__ static type signaling_NaN();\r
+            };\r
+\r
+            template<> struct numeric_limits_gpu<unsigned short>\r
+            {          \r
+                typedef unsigned short type;\r
+                __device__ static type min() { return 0; };            \r
+                __device__ static type max() { return USHRT_MAX; };\r
+                __device__ static type epsilon();\r
+                __device__ static type round_error();\r
+                __device__ static type denorm_min();\r
+                __device__ static type infinity();\r
+                __device__ static type quiet_NaN();\r
+                __device__ static type signaling_NaN();\r
+            };\r
+\r
+            template<> struct numeric_limits_gpu<int>\r
+            {    \r
+                typedef int type;\r
+                __device__ static type min() { return INT_MIN; };            \r
+                __device__ static type max() { return INT_MAX; };\r
+                __device__ static type epsilon();\r
+                __device__ static type round_error();\r
+                __device__ static type denorm_min();\r
+                __device__ static type infinity();\r
+                __device__ static type quiet_NaN();\r
+                __device__ static type signaling_NaN();\r
+            };\r
+\r
+\r
+            template<> struct numeric_limits_gpu<unsigned int>\r
+            {          \r
+                typedef unsigned int type;\r
+                __device__ static type min() { return 0; };            \r
+                __device__ static type max() { return UINT_MAX; };\r
+                __device__ static type epsilon();\r
+                __device__ static type round_error();\r
+                __device__ static type denorm_min();\r
+                __device__ static type infinity();\r
+                __device__ static type quiet_NaN();\r
+                __device__ static type signaling_NaN();\r
+            };\r
+\r
+            template<> struct numeric_limits_gpu<long>\r
+            {          \r
+                typedef long type;\r
+                __device__ static type min() { return LONG_MIN; };            \r
+                __device__ static type max() { return LONG_MAX; };\r
+                __device__ static type epsilon();\r
+                __device__ static type round_error();\r
+                __device__ static type denorm_min();\r
+                __device__ static type infinity();\r
+                __device__ static type quiet_NaN();\r
+                __device__ static type signaling_NaN();\r
+            };\r
+\r
+            template<> struct numeric_limits_gpu<unsigned long>\r
+            {          \r
+                typedef unsigned long type;\r
+                __device__ static type min() { return 0; };            \r
+                __device__ static type max() { return ULONG_MAX; };\r
+                __device__ static type epsilon();\r
+                __device__ static type round_error();\r
+                __device__ static type denorm_min();\r
+                __device__ static type infinity();\r
+                __device__ static type quiet_NaN();\r
+                __device__ static type signaling_NaN();\r
+            };\r
+                        \r
+            template<> struct numeric_limits_gpu<float>\r
+            {          \r
+                typedef float type;\r
+                __device__ static type min() { return 1.175494351e-38f/*FLT_MIN*/; };            \r
+                __device__ static type max() { return 3.402823466e+38f/*FLT_MAX*/; };\r
+                __device__ static type epsilon();\r
+                __device__ static type round_error();\r
+                __device__ static type denorm_min();\r
+                __device__ static type infinity();\r
+                __device__ static type quiet_NaN();\r
+                __device__ static type signaling_NaN();\r
+            };\r
+\r
+            template<> struct numeric_limits_gpu<double>\r
+            {          \r
+                typedef double type;\r
+                __device__ static type min() { return 2.2250738585072014e-308/*DBL_MIN*/; };            \r
+                __device__ static type max() { return 1.7976931348623158e+308/*DBL_MAX*/; };\r
+                __device__ static type epsilon();\r
+                __device__ static type round_error();\r
+                __device__ static type denorm_min();\r
+                __device__ static type infinity();\r
+                __device__ static type quiet_NaN();\r
+                __device__ static type signaling_NaN();\r
+            };            \r
+        }\r
+    }\r
+}
\ No newline at end of file

diff --git a/modules/gpu/src/error.cpp b/modules/gpu/src/error.cpp
index 432e9302b936bad110395afccc229ad4077c6a18..a30679e93384e5f3c679302ddcdccc218f05182a 100644 (file)
--- a/modules/gpu/src/error.cpp
+++ b/modules/gpu/src/error.cpp
@@ -135,7 +135,8 @@ namespace cv
         }\r
 \r
         void error(const char *error_string, const char *file, const int line, const char *func)\r
-        {                       \r
+        {          \r
+            //if (uncaught_exception())\r
             cv::error( cv::Exception(CV_GpuApiCallError, error_string, func, file, line) );\r
         }\r
     }\r

diff --git a/modules/gpu/src/matrix_operations.cpp b/modules/gpu/src/matrix_operations.cpp
index 99db93af5df849b3fce0f731c075c65ed02ef748..153c524528fab2211a2b24848a5803ae99e6d8e4 100644 (file)
--- a/modules/gpu/src/matrix_operations.cpp
+++ b/modules/gpu/src/matrix_operations.cpp
@@ -572,7 +572,7 @@ void cv::gpu::GpuMat::release()
 //////////////////////////////// CudaMem //////////////////////////////\r
 ///////////////////////////////////////////////////////////////////////\r
 \r
-bool cv::gpu::CudaMem::can_device_map_to_host()\r
+bool cv::gpu::CudaMem::canMapHostMemory()\r
 {\r
     cudaDeviceProp prop;\r
     cudaGetDeviceProperties(&prop, 0);\r
@@ -581,7 +581,7 @@ bool cv::gpu::CudaMem::can_device_map_to_host()
 \r
 void cv::gpu::CudaMem::create(int _rows, int _cols, int _type, int _alloc_type)\r
 {\r
-    if (_alloc_type == ALLOC_ZEROCOPY && !can_device_map_to_host())\r
+    if (_alloc_type == ALLOC_ZEROCOPY && !canMapHostMemory())\r
             cv::gpu::error("ZeroCopy is not supported by current device", __FILE__, __LINE__);\r
 \r
     _type &= TYPE_MASK;\r