39 #include "../thread/thread_load.cuh"
40 #include "../thread/thread_store.cuh"
41 #include "../util_device.cuh"
42 #include "../util_namespace.cuh"
44 #if (THRUST_VERSION >= 100700)
46 #include <thrust/iterator/iterator_facade.h>
47 #include <thrust/iterator/iterator_traits.h>
48 #endif // THRUST_VERSION
108 typename OffsetT = ptrdiff_t>
119 __host__ __device__ __forceinline__ Reference(ValueType* ptr) : ptr(ptr) {}
122 __device__ __forceinline__ ValueType operator =(ValueType val)
124 ThreadStore<MODIFIER>(ptr, val);
138 #if (THRUST_VERSION >= 100700)
140 typedef typename thrust::detail::iterator_facade_category<
141 thrust::device_system_tag,
142 thrust::random_access_traversal_tag,
148 #endif // THRUST_VERSION
157 template <
typename QualifiedValueType>
159 QualifiedValueType* ptr)
183 return Reference(ptr);
187 template <
typename Distance>
195 template <
typename Distance>
203 template <
typename Distance>
211 template <
typename Distance>
221 return ptr - other.ptr;
225 template <
typename Distance>
228 return Reference(ptr + n);
234 return (ptr == rhs.ptr);
240 return (ptr != rhs.ptr);