some utility for GPU module internal purposes

modules/gpu/include/opencv2/gpu/gpu.hpp

@@ -52,7 +52,7 @@ namespace cv
 {
     namespace gpu
     {
-        //////////////////////////////// Initialization ////////////////////////
+        //////////////////////////////// Initialization & Info ////////////////////////
 
         //! This is the only function that do not throw exceptions if the library is compiled without Cuda.
         CV_EXPORTS int getCudaEnabledDeviceCount();
@@ -67,6 +67,11 @@ namespace cv
 
         CV_EXPORTS void getGpuMemInfo(size_t& free, size_t& total);
 
+        //////////////////////////////// Error handling ////////////////////////
+        
+        CV_EXPORTS void error(const char *error_string, const char *file, const int line, const char *func);
+        CV_EXPORTS void nppError( int err, const char *file, const int line, const char *func);        
+
         //////////////////////////////// GpuMat ////////////////////////////////
         class Stream;
         class CudaMem;

modules/gpu/src/cuda/cuda_shared.hpp

@@ -58,8 +58,42 @@ namespace cv
 
         static inline int divUp(int total, int grain) { return (total + grain - 1) / grain; }
 
-        template<class T> 
+        template<class T> static inline void uploadConstant(const char* name, const T& value) 
-        static inline void uploadConstant(const char* name, const T& value) { cudaSafeCall( cudaMemcpyToSymbol(name, &value, sizeof(T)) ); }
+        { 
+            cudaSafeCall( cudaMemcpyToSymbol(name, &value, sizeof(T)) ); 
+        }
+
+        template<class T> static inline void uploadConstant(const char* name, const T& value, cudaStream_t stream) 
+        { 
+            cudaSafeCall( cudaMemcpyToSymbolAsyc(name, &value, sizeof(T), 0, cudaMemcpyHostToDevice, stream) ); 
+        }        
+
+        template<class T> static inline void bindTexture(const char* name, const DevMem2D_<T>& img/*, bool normalized = false,
+            enum cudaTextureFilterMode filterMode = cudaFilterModePoint, enum cudaTextureAddressMode addrMode = cudaAddressModeClamp*/)
+        {            
+            //!!!! const_cast is disabled!
+            //!!!! Please use constructor of 'class texture'  instead.
+
+            //textureReference* tex; 
+            //cudaSafeCall( cudaGetTextureReference((const textureReference**)&tex, name) ); 
+            //tex->normalized = normalized;
+            //tex->filterMode = filterMode;
+            //tex->addressMode[0] = addrMode;
+            //tex->addressMode[1] = addrMode;
+            
+            const textureReference* tex; 
+            cudaSafeCall( cudaGetTextureReference(&tex, name) ); 
+
+            cudaChannelFormatDesc desc = cudaCreateChannelDesc<T>();
+            cudaSafeCall( cudaBindTexture2D(0, tex, img.ptr(), &desc, img.cols, img.rows, img.step) );
+        }
+
+        static inline void unbindTexture(const char *name)
+        {
+            const textureReference* tex; 
+            cudaSafeCall( cudaGetTextureReference(&tex, name) ); 
+            cudaSafeCall( cudaUnbindTexture(tex) );
+        }        
     }
 }
 

modules/gpu/src/cuda/safe_call.hpp

@@ -44,7 +44,7 @@
 #define __OPENCV_CUDA_SAFE_CALL_HPP__
 
 #include "cuda_runtime_api.h"
-#include <nppdefs.h>
+//#include <nppdefs.h>
 
 #if defined(__GNUC__)
     #define cudaSafeCall(expr)  ___cudaSafeCall(expr, __FILE__, __LINE__, __func__)
@@ -58,8 +58,8 @@ namespace cv
 {
     namespace gpu
     {
-        extern "C" void error( const char *error_string, const char *file, const int line, const char *func = "");
+        void error( const char *error_string, const char *file, const int line, const char *func = "");
-        extern "C" void npp_error( int error, const char *file, const int line, const char *func = "");   
+        void nppError( int error, const char *file, const int line, const char *func = "");   
 
         static inline void ___cudaSafeCall(cudaError_t err, const char *file, const int line, const char *func = "")
         {
@@ -67,10 +67,10 @@ namespace cv
                 cv::gpu::error(cudaGetErrorString(err), file, line, func);
         }
 
-        static inline void ___nppSafeCall(NppStatus err, const char *file, const int line, const char *func = "")
+        static inline void ___nppSafeCall(int err, const char *file, const int line, const char *func = "")
         {
             if (err < 0)
-                cv::gpu::npp_error(err, file, line, func);
+                cv::gpu::nppError(err, file, line, func);
         }
     }
 }

modules/gpu/src/error.cpp

@@ -129,12 +129,12 @@ namespace cv
             return interpreter.str();
         }
 
-        extern "C" void npp_error( int err, const char *file, const int line, const char *func)
+        void nppError( int err, const char *file, const int line, const char *func)
         {                    
             cv::error( cv::Exception(CV_GpuNppCallError, getNppErrorString(err), func, file, line) );                
         }
 
-        extern "C" void error(const char *error_string, const char *file, const int line, const char *func)
+        void error(const char *error_string, const char *file, const int line, const char *func)
         {                       
             cv::error( cv::Exception(CV_GpuApiCallError, error_string, func, file, line) );
         }