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 (CUDA_VERSION >= 5050) || defined(DOXYGEN_ACTIVE) // This iterator is compatible with CUDA 5.5 and newer
47 #if (THRUST_VERSION >= 100700) // This iterator is compatible with Thrust API 1.7 and newer
48 #include <thrust/iterator/iterator_facade.h>
49 #include <thrust/iterator/iterator_traits.h>
50 #endif // THRUST_VERSION
64 #ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document
74 template <
int UNIQUE_ID>
78 typedef typename UnitWord<T>::DeviceWord DeviceWord;
79 typedef typename UnitWord<T>::TextureWord TextureWord;
83 DEVICE_MULTIPLE =
sizeof(T) /
sizeof(DeviceWord),
84 TEXTURE_MULTIPLE =
sizeof(T) /
sizeof(TextureWord)
88 typedef texture<TextureWord> TexRef;
94 static cudaError_t BindTexture(
void *d_in,
size_t &offset)
98 cudaChannelFormatDesc tex_desc = cudaCreateChannelDesc<TextureWord>();
99 ref.channelDesc = tex_desc;
100 return (
CubDebug(cudaBindTexture(&offset, ref, d_in)));
107 static cudaError_t UnbindTexture()
109 return CubDebug(cudaUnbindTexture(ref));
113 template <
typename Distance>
114 static __device__ __forceinline__ T Fetch(Distance tex_offset)
116 DeviceWord temp[DEVICE_MULTIPLE];
117 TextureWord *words =
reinterpret_cast<TextureWord*
>(temp);
120 for (
int i = 0; i < TEXTURE_MULTIPLE; ++i)
122 words[i] = tex1Dfetch(ref, (tex_offset * TEXTURE_MULTIPLE) + i);
125 return reinterpret_cast<T&
>(temp);
131 template <
typename T>
132 template <
int UNIQUE_ID>
133 typename IteratorTexRef<T>::template TexId<UNIQUE_ID>::TexRef IteratorTexRef<T>::template TexId<UNIQUE_ID>::ref = 0;
139 #endif // DOXYGEN_SHOULD_SKIP_THIS
203 typename OffsetT = ptrdiff_t>
215 #if (THRUST_VERSION >= 100700)
217 typedef typename thrust::detail::iterator_facade_category<
218 thrust::device_system_tag,
219 thrust::random_access_traversal_tag,
225 #endif // THRUST_VERSION
233 typedef typename IteratorTexRef<T>::template TexId<UNIQUE_ID> TexId;
245 template <
typename QualifiedT>
248 size_t bytes =
size_t(-1),
249 size_t tex_offset = 0)
253 cudaError_t retval = TexId::BindTexture(this->ptr + tex_offset, offset);
261 return TexId::UnbindTexture();
282 #if (CUB_PTX_ARCH == 0)
284 return ptr[tex_offset];
287 return TexId::Fetch(tex_offset);
292 template <
typename Distance>
297 retval.tex_offset = tex_offset + n;
302 template <
typename Distance>
310 template <
typename Distance>
315 retval.tex_offset = tex_offset - n;
320 template <
typename Distance>
330 return tex_offset - other.tex_offset;
334 template <
typename Distance>
350 return ((ptr == rhs.ptr) && (tex_offset == rhs.tex_offset));
356 return ((ptr != rhs.ptr) || (tex_offset != rhs.tex_offset));
374 #endif // CUDA_VERSION