Commit 02645345 authored by Richard Veale's avatar Richard Veale Committed by Alexander Alekhin

[moved from opencv] Merge pull request opencv/opencv#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...

original commit: https://github.com/opencv/opencv/commit/8158e5b7a0000ff1a6f7d65763d2b1605844ba46
parent 90d16271
......@@ -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); \
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment