Merge moved code from opencv

modules/cudaarithm/src/core.cpp

@@ -57,8 +57,6 @@ void cv::cuda::transpose(InputArray, OutputArray, Stream&) { throw_no_cuda(); }
 
 void cv::cuda::flip(InputArray, OutputArray, int, Stream&) { throw_no_cuda(); }
 
-Ptr<LookUpTable> cv::cuda::createLookUpTable(InputArray) { throw_no_cuda(); return Ptr<LookUpTable>(); }
-
 void cv::cuda::copyMakeBorder(InputArray, OutputArray, int, int, int, int, int, Scalar, Stream&) { throw_no_cuda(); }
 
 #else /* !defined (HAVE_CUDA) */

modules/cudaarithm/src/cuda/lut.cu

@@ -48,6 +48,8 @@
 
 #else
 
+#include "../lut.hpp"
+
 #include "opencv2/cudaarithm.hpp"
 #include "opencv2/cudev.hpp"
 #include "opencv2/core/private.cuda.hpp"
@@ -56,23 +58,9 @@ using namespace cv;
 using namespace cv::cuda;
 using namespace cv::cudev;
 
-namespace
-{
-    texture<uchar, cudaTextureType1D, cudaReadModeElementType> texLutTable;
-
-    class LookUpTableImpl : public LookUpTable
-    {
-    public:
-        LookUpTableImpl(InputArray lut);
-        ~LookUpTableImpl();
-
-        void transform(InputArray src, OutputArray dst, Stream& stream = Stream::Null()) CV_OVERRIDE;
+namespace cv { namespace cuda {
 
-    private:
-        GpuMat d_lut;
-        cudaTextureObject_t texLutTableObj;
-        bool cc30;
-    };
+    texture<uchar, cudaTextureType1D, cudaReadModeElementType> texLutTable;
 
     LookUpTableImpl::LookUpTableImpl(InputArray _lut)
     {
@@ -200,11 +188,7 @@ namespace
 
         syncOutput(dst, _dst, stream);
     }
-}
 
-Ptr<LookUpTable> cv::cuda::createLookUpTable(InputArray lut)
-{
-    return makePtr<LookUpTableImpl>(lut);
-}
+} }
 
 #endif

modules/cudaarithm/src/lut.cpp

+// This file is part of OpenCV project.
+// It is subject to the license terms in the LICENSE file found in the top-level directory
+// of this distribution and at http://opencv.org/license.html.
+
+#include "precomp.hpp"
+
+#include "lut.hpp"
+
+using namespace cv;
+using namespace cv::cuda;
+
+#if !defined (HAVE_CUDA) || defined (CUDA_DISABLER)
+
+Ptr<LookUpTable> cv::cuda::createLookUpTable(InputArray) { throw_no_cuda(); return Ptr<LookUpTable>(); }
+
+#else /* !defined (HAVE_CUDA) || defined (CUDA_DISABLER) */
+
+Ptr<LookUpTable> cv::cuda::createLookUpTable(InputArray lut)
+{
+    return makePtr<LookUpTableImpl>(lut);
+}
+
+#endif

modules/cudaarithm/src/lut.hpp

+// This file is part of OpenCV project.
+// It is subject to the license terms in the LICENSE file found in the top-level directory
+// of this distribution and at http://opencv.org/license.html.
+
+#ifndef __CUDAARITHM_LUT_HPP__
+#define __CUDAARITHM_LUT_HPP__
+
+#include "opencv2/cudaarithm.hpp"
+
+#include <cuda_runtime.h>
+
+namespace cv { namespace cuda {
+
+class LookUpTableImpl : public LookUpTable
+{
+public:
+    LookUpTableImpl(InputArray lut);
+    ~LookUpTableImpl();
+
+    void transform(InputArray src, OutputArray dst, Stream& stream = Stream::Null()) CV_OVERRIDE;
+
+private:
+    GpuMat d_lut;
+    cudaTextureObject_t texLutTableObj;
+    bool cc30;
+};
+
+} }
+
+#endif // __CUDAARITHM_LUT_HPP__

modules/cudafilters/src/cuda/filter2d.cu

@@ -77,17 +77,17 @@ namespace cv { namespace cuda { namespace device
     template <typename T, typename D, template <typename> class Brd> struct Filter2DCaller;
 
     #define IMPLEMENT_FILTER2D_TEX_READER(type) \
-        texture< type , cudaTextureType2D, cudaReadModeElementType> tex_filter2D_ ## type (0, cudaFilterModePoint, cudaAddressModeClamp); \
         struct tex_filter2D_ ## type ## _reader \
         { \
+            PtrStepSz<type> dat; \
             typedef type elem_type; \
             typedef int index_type; \
             const int xoff; \
             const int yoff; \
-            tex_filter2D_ ## type ## _reader (int xoff_, int yoff_) : xoff(xoff_), yoff(yoff_) {} \
+            tex_filter2D_ ## type ## _reader (PtrStepSz<type> dat_, int xoff_, int yoff_) : dat(dat_), xoff(xoff_), yoff(yoff_) {} \
             __device__ __forceinline__ elem_type operator ()(index_type y, index_type x) const \
             { \
-                return tex2D(tex_filter2D_ ## type , x + xoff, y + yoff); \
+                return dat(y + yoff, x + xoff ); \
             } \
         }; \
         template <typename D, template <typename> class Brd> struct Filter2DCaller< type , D, Brd> \
@@ -98,8 +98,7 @@ namespace cv { namespace cuda { namespace device
                 typedef typename TypeVec<float, VecTraits< type >::cn>::vec_type work_type; \
                 dim3 block(16, 16); \
                 dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y)); \
-                bindTexture(&tex_filter2D_ ## type , srcWhole); \
-                tex_filter2D_ ## type ##_reader texSrc(xoff, yoff); \
+                tex_filter2D_ ## type ##_reader texSrc(srcWhole, xoff, yoff); \
                 Brd<work_type> brd(dst.rows, dst.cols, VecTraits<work_type>::make(borderValue)); \
                 BorderReader< tex_filter2D_ ## type ##_reader, Brd<work_type> > brdSrc(texSrc, brd); \
                 filter2D<<<grid, block, 0, stream>>>(brdSrc, dst, kernel, kWidth, kHeight, anchorX, anchorY); \