39 #include "../thread/thread_load.cuh"
40 #include "../thread/thread_store.cuh"
41 #include "../util_device.cuh"
42 #include "../util_debug.cuh"
43 #include "../util_namespace.cuh"
45 #if (THRUST_VERSION >= 100700)
47 #include <thrust/iterator/iterator_facade.h>
48 #include <thrust/iterator/iterator_traits.h>
49 #endif // THRUST_VERSION
110 typename OffsetT = ptrdiff_t>
122 #if (THRUST_VERSION >= 100700)
124 typedef typename thrust::detail::iterator_facade_category<
125 thrust::device_system_tag,
126 thrust::random_access_traversal_tag,
132 #endif // THRUST_VERSION
137 typedef typename UnitWord<T>::TextureWord TextureWord;
141 TEXTURE_MULTIPLE =
sizeof(T) /
sizeof(TextureWord)
148 cudaTextureObject_t tex_obj;
161 template <
typename QualifiedT>
164 size_t bytes =
size_t(-1),
165 size_t tex_offset = 0)
168 this->tex_offset = tex_offset;
170 cudaChannelFormatDesc channel_desc = cudaCreateChannelDesc<TextureWord>();
171 cudaResourceDesc res_desc;
172 cudaTextureDesc tex_desc;
173 memset(&res_desc, 0,
sizeof(cudaResourceDesc));
174 memset(&tex_desc, 0,
sizeof(cudaTextureDesc));
175 res_desc.resType = cudaResourceTypeLinear;
176 res_desc.res.linear.devPtr = this->ptr;
177 res_desc.res.linear.desc = channel_desc;
178 res_desc.res.linear.sizeInBytes = bytes;
179 tex_desc.readMode = cudaReadModeElementType;
180 return cudaCreateTextureObject(&tex_obj, &res_desc, &tex_desc, NULL);
186 return cudaDestroyTextureObject(tex_obj);
207 #if (CUB_PTX_ARCH == 0)
209 return ptr[tex_offset];
212 TextureWord words[TEXTURE_MULTIPLE];
215 for (
int i = 0; i < TEXTURE_MULTIPLE; ++i)
217 words[i] = tex1Dfetch<TextureWord>(
219 (tex_offset * TEXTURE_MULTIPLE) + i);
223 return *
reinterpret_cast<T*
>(words);
228 template <
typename Distance>
233 retval.tex_obj = tex_obj;
234 retval.tex_offset = tex_offset + n;
239 template <
typename Distance>
247 template <
typename Distance>
252 retval.tex_obj = tex_obj;
253 retval.tex_offset = tex_offset - n;
258 template <
typename Distance>
268 return tex_offset - other.tex_offset;
272 template <
typename Distance>
288 return ((ptr == rhs.ptr) && (tex_offset == rhs.tex_offset) && (tex_obj == rhs.tex_obj));
294 return ((ptr != rhs.ptr) || (tex_offset != rhs.tex_offset) || (tex_obj != rhs.tex_obj));