From 8158e5b7a0000ff1a6f7d65763d2b1605844ba46 Mon Sep 17 00:00:00 2001
From: Richard Veale <richard.e.veale@gmail.com>
Date: Thu, 14 Mar 2019 02:53:59 +0900
Subject: [PATCH] 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 | 9 ++++-----
 1 file changed, 4 insertions(+), 5 deletions(-)

diff --git a/modules/cudafilters/src/cuda/filter2d.cu b/modules/cudafilters/src/cuda/filter2d.cu
index 67d75d4c78..e37e91dd2f 100644
--- a/modules/cudafilters/src/cuda/filter2d.cu
+++ b/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); \