1 /*M///////////////////////////////////////////////////////////////////////////////////////
2 //
3 //  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
4 //
5 //  By downloading, copying, installing or using the software you agree to this license.
6 //  If you do not agree to this license, do not download, install,
7 //  copy or use the software.
8 //
9 //
10 //                          License Agreement
11 //                For Open Source Computer Vision Library
12 //
13 // Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
14 // Copyright (C) 2009, Willow Garage Inc., all rights reserved.
15 // Copyright (C) 2013, OpenCV Foundation, all rights reserved.
16 // Third party copyrights are property of their respective owners.
17 //
18 // Redistribution and use in source and binary forms, with or without modification,
19 // are permitted provided that the following conditions are met:
20 //
21 //   * Redistribution's of source code must retain the above copyright notice,
22 //     this list of conditions and the following disclaimer.
23 //
24 //   * Redistribution's in binary form must reproduce the above copyright notice,
25 //     this list of conditions and the following disclaimer in the documentation
26 //     and/or other materials provided with the distribution.
27 //
28 //   * The name of the copyright holders may not be used to endorse or promote products
29 //     derived from this software without specific prior written permission.
30 //
31 // This software is provided by the copyright holders and contributors "as is" and
32 // any express or implied warranties, including, but not limited to, the implied
33 // warranties of merchantability and fitness for a particular purpose are disclaimed.
34 // In no event shall the Intel Corporation or contributors be liable for any direct,
35 // indirect, incidental, special, exemplary, or consequential damages
36 // (including, but not limited to, procurement of substitute goods or services;
37 // loss of use, data, or profits; or business interruption) however caused
38 // and on any theory of liability, whether in contract, strict liability,
39 // or tort (including negligence or otherwise) arising in any way out of
40 // the use of this software, even if advised of the possibility of such damage.
41 //
42 //M*/
43 
44 #pragma once
45 
46 #ifndef __OPENCV_CUDEV_PTR2D_TEXTURE_HPP__
47 #define __OPENCV_CUDEV_PTR2D_TEXTURE_HPP__
48 
49 #include <cstring>
50 #include "../common.hpp"
51 #include "glob.hpp"
52 #include "gpumat.hpp"
53 #include "traits.hpp"
54 
55 #if CUDART_VERSION >= 5050
56 
57 namespace
58 {
59     template <typename T> struct CvCudevTextureRef
60     {
61         typedef texture<T, cudaTextureType2D, cudaReadModeElementType> TexRef;
62 
63         static TexRef ref;
64 
bind__anon57003a530111::CvCudevTextureRef65         __host__ static void bind(const cv::cudev::GlobPtrSz<T>& mat,
66                                   bool normalizedCoords = false,
67                                   cudaTextureFilterMode filterMode = cudaFilterModePoint,
68                                   cudaTextureAddressMode addressMode = cudaAddressModeClamp)
69         {
70             ref.normalized = normalizedCoords;
71             ref.filterMode = filterMode;
72             ref.addressMode[0] = addressMode;
73             ref.addressMode[1] = addressMode;
74             ref.addressMode[2] = addressMode;
75 
76             cudaChannelFormatDesc desc = cudaCreateChannelDesc<T>();
77 
78             CV_CUDEV_SAFE_CALL( cudaBindTexture2D(0, &ref, mat.data, &desc, mat.cols, mat.rows, mat.step) );
79         }
80 
unbind__anon57003a530111::CvCudevTextureRef81         __host__ static void unbind()
82         {
83             cudaUnbindTexture(ref);
84         }
85     };
86 
87     template <typename T>
88     typename CvCudevTextureRef<T>::TexRef CvCudevTextureRef<T>::ref;
89 }
90 
91 #endif
92 
93 namespace cv { namespace cudev {
94 
95 //! @addtogroup cudev
96 //! @{
97 
98 #if CUDART_VERSION >= 5050
99 
100 template <typename T> struct TexturePtr
101 {
102     typedef T     value_type;
103     typedef float index_type;
104 
105     cudaTextureObject_t texObj;
106 
operator ()cv::cudev::TexturePtr107     __device__ __forceinline__ T operator ()(float y, float x) const
108     {
109     #if CV_CUDEV_ARCH < 300
110         // Use the texture reference
111         return tex2D(CvCudevTextureRef<T>::ref, x, y);
112     #else
113         // Use the texture object
114         return tex2D<T>(texObj, x, y);
115     #endif
116     }
117 };
118 
119 template <typename T> struct Texture : TexturePtr<T>
120 {
121     int rows, cols;
122     bool cc30;
123 
Texturecv::cudev::Texture124     __host__ explicit Texture(const GlobPtrSz<T>& mat,
125                               bool normalizedCoords = false,
126                               cudaTextureFilterMode filterMode = cudaFilterModePoint,
127                               cudaTextureAddressMode addressMode = cudaAddressModeClamp)
128     {
129         cc30 = deviceSupports(FEATURE_SET_COMPUTE_30);
130 
131         rows = mat.rows;
132         cols = mat.cols;
133 
134         if (cc30)
135         {
136             // Use the texture object
137             cudaResourceDesc texRes;
138             std::memset(&texRes, 0, sizeof(texRes));
139             texRes.resType = cudaResourceTypePitch2D;
140             texRes.res.pitch2D.devPtr = mat.data;
141             texRes.res.pitch2D.height = mat.rows;
142             texRes.res.pitch2D.width = mat.cols;
143             texRes.res.pitch2D.pitchInBytes = mat.step;
144             texRes.res.pitch2D.desc = cudaCreateChannelDesc<T>();
145 
146             cudaTextureDesc texDescr;
147             std::memset(&texDescr, 0, sizeof(texDescr));
148             texDescr.normalizedCoords = normalizedCoords;
149             texDescr.filterMode = filterMode;
150             texDescr.addressMode[0] = addressMode;
151             texDescr.addressMode[1] = addressMode;
152             texDescr.addressMode[2] = addressMode;
153             texDescr.readMode = cudaReadModeElementType;
154 
155             CV_CUDEV_SAFE_CALL( cudaCreateTextureObject(&this->texObj, &texRes, &texDescr, 0) );
156         }
157         else
158         {
159             // Use the texture reference
160             CvCudevTextureRef<T>::bind(mat, normalizedCoords, filterMode, addressMode);
161         }
162     }
163 
~Texturecv::cudev::Texture164     __host__ ~Texture()
165     {
166         if (cc30)
167         {
168             // Use the texture object
169             cudaDestroyTextureObject(this->texObj);
170         }
171         else
172         {
173             // Use the texture reference
174             CvCudevTextureRef<T>::unbind();
175         }
176     }
177 };
178 
179 template <typename T> struct PtrTraits< Texture<T> > : PtrTraitsBase<Texture<T>, TexturePtr<T> >
180 {
181 };
182 
183 #else
184 
185 template <typename T> struct TexturePtr
186 {
187     typedef T     value_type;
188     typedef float index_type;
189 
190     cudaTextureObject_t texObj;
191 
192     __device__ __forceinline__ T operator ()(float y, float x) const
193     {
194     #if CV_CUDEV_ARCH >= 300
195         // Use the texture object
196         return tex2D<T>(texObj, x, y);
197     #else
198         (void) y;
199         (void) x;
200         return T();
201     #endif
202     }
203 };
204 
205 template <typename T> struct Texture : TexturePtr<T>
206 {
207     int rows, cols;
208 
209     __host__ explicit Texture(const GlobPtrSz<T>& mat,
210                               bool normalizedCoords = false,
211                               cudaTextureFilterMode filterMode = cudaFilterModePoint,
212                               cudaTextureAddressMode addressMode = cudaAddressModeClamp)
213     {
214         CV_Assert( deviceSupports(FEATURE_SET_COMPUTE_30) );
215 
216         rows = mat.rows;
217         cols = mat.cols;
218 
219         // Use the texture object
220         cudaResourceDesc texRes;
221         std::memset(&texRes, 0, sizeof(texRes));
222         texRes.resType = cudaResourceTypePitch2D;
223         texRes.res.pitch2D.devPtr = mat.data;
224         texRes.res.pitch2D.height = mat.rows;
225         texRes.res.pitch2D.width = mat.cols;
226         texRes.res.pitch2D.pitchInBytes = mat.step;
227         texRes.res.pitch2D.desc = cudaCreateChannelDesc<T>();
228 
229         cudaTextureDesc texDescr;
230         std::memset(&texDescr, 0, sizeof(texDescr));
231         texDescr.normalizedCoords = normalizedCoords;
232         texDescr.filterMode = filterMode;
233         texDescr.addressMode[0] = addressMode;
234         texDescr.addressMode[1] = addressMode;
235         texDescr.addressMode[2] = addressMode;
236         texDescr.readMode = cudaReadModeElementType;
237 
238         CV_CUDEV_SAFE_CALL( cudaCreateTextureObject(&this->texObj, &texRes, &texDescr, 0) );
239     }
240 
241     __host__ ~Texture()
242     {
243         // Use the texture object
244         cudaDestroyTextureObject(this->texObj);
245     }
246 };
247 
248 template <typename T> struct PtrTraits< Texture<T> > : PtrTraitsBase<Texture<T>, TexturePtr<T> >
249 {
250 };
251 
252 #endif
253 
254 //! @}
255 
256 }}
257 
258 #endif
259