Merge pull request #13695 from flyingfalling:3.4

* Fixed bug that made cuda::filter give corrupted output when different filters called in different threads. This was due to use of global texture variable shared by all filters.

* REV: fixed tab issue for opencv coding style...

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); \