diff --git a/modules/cudafilters/src/cuda/filter2d.cu b/modules/cudafilters/src/cuda/filter2d.cu index 67d75d4c7..e37e91dd2 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 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 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 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 class Brd> struct Filter2DCaller< type , D, Brd> \ @@ -98,8 +98,7 @@ namespace cv { namespace cuda { namespace device typedef typename TypeVec::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 brd(dst.rows, dst.cols, VecTraits::make(borderValue)); \ BorderReader< tex_filter2D_ ## type ##_reader, Brd > brdSrc(texSrc, brd); \ filter2D<<>>(brdSrc, dst, kernel, kWidth, kHeight, anchorX, anchorY); \