diff --git a/modules/cudev/include/opencv2/cudev/ptr2d/texture.hpp b/modules/cudev/include/opencv2/cudev/ptr2d/texture.hpp
index 095864fcb0..6df4a783d8 100644
--- a/modules/cudev/include/opencv2/cudev/ptr2d/texture.hpp
+++ b/modules/cudev/include/opencv2/cudev/ptr2d/texture.hpp
@@ -52,6 +52,8 @@
 #include "gpumat.hpp"
 #include "traits.hpp"
 
+#if CUDART_VERSION >= 5050
+
 namespace
 {
     template <typename T> struct CvCudevTextureRef
@@ -78,7 +80,7 @@ namespace
 
         __host__ static void unbind()
         {
-            CV_CUDEV_SAFE_CALL( cudaUnbindTexture(ref) );
+            cudaUnbindTexture(ref);
         }
     };
 
@@ -86,8 +88,12 @@ namespace
     typename CvCudevTextureRef<T>::TexRef CvCudevTextureRef<T>::ref;
 }
 
+#endif
+
 namespace cv { namespace cudev {
 
+#if CUDART_VERSION >= 5050
+
 template <typename T> struct TexturePtr
 {
     typedef T     value_type;
@@ -171,6 +177,77 @@ template <typename T> struct PtrTraits< Texture<T> > : PtrTraitsBase<Texture<T>,
 {
 };
 
+#else
+
+template <typename T> struct TexturePtr
+{
+    typedef T     value_type;
+    typedef float index_type;
+
+    cudaTextureObject_t texObj;
+
+    __device__ __forceinline__ T operator ()(float y, float x) const
+    {
+    #if CV_CUDEV_ARCH >= 300
+        // Use the texture object
+        return tex2D<T>(texObj, x, y);
+    #else
+        (void) y;
+        (void) x;
+        return T();
+    #endif
+    }
+};
+
+template <typename T> struct Texture : TexturePtr<T>
+{
+    int rows, cols;
+
+    __host__ explicit Texture(const GlobPtrSz<T>& mat,
+                              bool normalizedCoords = false,
+                              cudaTextureFilterMode filterMode = cudaFilterModePoint,
+                              cudaTextureAddressMode addressMode = cudaAddressModeClamp)
+    {
+        CV_Assert( deviceSupports(FEATURE_SET_COMPUTE_30) );
+
+        rows = mat.rows;
+        cols = mat.cols;
+
+        // Use the texture object
+        cudaResourceDesc texRes;
+        std::memset(&texRes, 0, sizeof(texRes));
+        texRes.resType = cudaResourceTypePitch2D;
+        texRes.res.pitch2D.devPtr = mat.data;
+        texRes.res.pitch2D.height = mat.rows;
+        texRes.res.pitch2D.width = mat.cols;
+        texRes.res.pitch2D.pitchInBytes = mat.step;
+        texRes.res.pitch2D.desc = cudaCreateChannelDesc<T>();
+
+        cudaTextureDesc texDescr;
+        std::memset(&texDescr, 0, sizeof(texDescr));
+        texDescr.normalizedCoords = normalizedCoords;
+        texDescr.filterMode = filterMode;
+        texDescr.addressMode[0] = addressMode;
+        texDescr.addressMode[1] = addressMode;
+        texDescr.addressMode[2] = addressMode;
+        texDescr.readMode = cudaReadModeElementType;
+
+        CV_CUDEV_SAFE_CALL( cudaCreateTextureObject(&this->texObj, &texRes, &texDescr, 0) );
+    }
+
+    __host__ ~Texture()
+    {
+        // Use the texture object
+        cudaDestroyTextureObject(this->texObj);
+    }
+};
+
+template <typename T> struct PtrTraits< Texture<T> > : PtrTraitsBase<Texture<T>, TexturePtr<T> >
+{
+};
+
+#endif
+
 }}
 
 #endif