40 #include "../util_ptx.cuh"
41 #include "../util_type.cuh"
42 #include "../util_namespace.cuh"
109 typename InputIteratorT>
110 __device__ __forceinline__
typename std::iterator_traits<InputIteratorT>::value_type
ThreadLoad(InputIteratorT itr);
116 #ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document
120 template <
int COUNT,
int MAX>
121 struct IterateThreadLoad
123 template <CacheLoadModifier MODIFIER,
typename T>
124 static __device__ __forceinline__
void Load(T
const *ptr, T *vals)
126 vals[COUNT] = ThreadLoad<MODIFIER>(ptr + COUNT);
127 IterateThreadLoad<COUNT + 1, MAX>::template Load<MODIFIER>(ptr, vals);
130 template <
typename InputIteratorT,
typename T>
131 static __device__ __forceinline__
void Dereference(InputIteratorT itr, T *vals)
133 vals[COUNT] = itr[COUNT];
134 IterateThreadLoad<COUNT + 1, MAX>::Dereference(itr, vals);
141 struct IterateThreadLoad<MAX, MAX>
143 template <CacheLoadModifier MODIFIER,
typename T>
144 static __device__ __forceinline__
void Load(T
const *ptr, T *vals) {}
146 template <
typename InputIteratorT,
typename T>
147 static __device__ __forceinline__
void Dereference(InputIteratorT itr, T *vals) {}
154 #define _CUB_LOAD_16(cub_modifier, ptx_modifier) \
156 __device__ __forceinline__ uint4 ThreadLoad<cub_modifier, uint4 const *>(uint4 const *ptr) \
159 asm volatile ("ld."#ptx_modifier".v4.u32 {%0, %1, %2, %3}, [%4];" : \
164 _CUB_ASM_PTR_(ptr)); \
168 __device__ __forceinline__ ulonglong2 ThreadLoad<cub_modifier, ulonglong2 const *>(ulonglong2 const *ptr) \
171 asm volatile ("ld."#ptx_modifier".v2.u64 {%0, %1}, [%2];" : \
174 _CUB_ASM_PTR_(ptr)); \
181 #define _CUB_LOAD_8(cub_modifier, ptx_modifier) \
183 __device__ __forceinline__ ushort4 ThreadLoad<cub_modifier, ushort4 const *>(ushort4 const *ptr) \
186 asm volatile ("ld."#ptx_modifier".v4.u16 {%0, %1, %2, %3}, [%4];" : \
191 _CUB_ASM_PTR_(ptr)); \
195 __device__ __forceinline__ uint2 ThreadLoad<cub_modifier, uint2 const *>(uint2 const *ptr) \
198 asm volatile ("ld."#ptx_modifier".v2.u32 {%0, %1}, [%2];" : \
201 _CUB_ASM_PTR_(ptr)); \
205 __device__ __forceinline__ unsigned long long ThreadLoad<cub_modifier, unsigned long long const *>(unsigned long long const *ptr) \
207 unsigned long long retval; \
208 asm volatile ("ld."#ptx_modifier".u64 %0, [%1];" : \
210 _CUB_ASM_PTR_(ptr)); \
217 #define _CUB_LOAD_4(cub_modifier, ptx_modifier) \
219 __device__ __forceinline__ unsigned int ThreadLoad<cub_modifier, unsigned int const *>(unsigned int const *ptr) \
221 unsigned int retval; \
222 asm volatile ("ld."#ptx_modifier".u32 %0, [%1];" : \
224 _CUB_ASM_PTR_(ptr)); \
232 #define _CUB_LOAD_2(cub_modifier, ptx_modifier) \
234 __device__ __forceinline__ unsigned short ThreadLoad<cub_modifier, unsigned short const *>(unsigned short const *ptr) \
236 unsigned short retval; \
237 asm volatile ("ld."#ptx_modifier".u16 %0, [%1];" : \
239 _CUB_ASM_PTR_(ptr)); \
247 #define _CUB_LOAD_1(cub_modifier, ptx_modifier) \
249 __device__ __forceinline__ unsigned char ThreadLoad<cub_modifier, unsigned char const *>(unsigned char const *ptr) \
251 unsigned short retval; \
255 " ld."#ptx_modifier".u8 datum, [%1];" \
256 " cvt.u16.u8 %0, datum;" \
259 _CUB_ASM_PTR_(ptr)); \
260 return (unsigned char) retval; \
267 #define _CUB_LOAD_ALL(cub_modifier, ptx_modifier) \
268 _CUB_LOAD_16(cub_modifier, ptx_modifier) \
269 _CUB_LOAD_8(cub_modifier, ptx_modifier) \
270 _CUB_LOAD_4(cub_modifier, ptx_modifier) \
271 _CUB_LOAD_2(cub_modifier, ptx_modifier) \
272 _CUB_LOAD_1(cub_modifier, ptx_modifier) \
278 #if CUB_PTX_ARCH >= 200
286 _CUB_LOAD_ALL(LOAD_CG, volatile.global)
287 _CUB_LOAD_ALL(LOAD_CS, global)
288 _CUB_LOAD_ALL(LOAD_CV, volatile.global)
291 #if CUB_PTX_ARCH >= 350
311 template <
typename InputIteratorT>
312 __device__ __forceinline__
typename std::iterator_traits<InputIteratorT>::value_type
ThreadLoad(
314 Int2Type<LOAD_DEFAULT> modifier,
315 Int2Type<false> is_pointer)
324 template <
typename T>
327 Int2Type<LOAD_DEFAULT> modifier,
328 Int2Type<true> is_pointer)
337 template <
typename T>
338 __device__ __forceinline__ T ThreadLoadVolatilePointer(
340 Int2Type<true> is_primitive)
342 T retval = *
reinterpret_cast<volatile T*
>(ptr);
344 #if (CUB_PTX_ARCH <= 130)
345 if (
sizeof(T) == 1) __threadfence_block();
355 template <
typename T>
356 __device__ __forceinline__ T ThreadLoadVolatilePointer(
358 Int2Type<false> is_primitive)
361 #if CUB_PTX_ARCH <= 130
364 __threadfence_block();
369 typedef typename UnitWord<T>::VolatileWord VolatileWord;
371 const int VOLATILE_MULTIPLE =
sizeof(T) /
sizeof(VolatileWord);
383 VolatileWord *words =
reinterpret_cast<VolatileWord*
>(&retval);
384 IterateThreadLoad<0, VOLATILE_MULTIPLE>::Dereference(
385 reinterpret_cast<volatile VolatileWord*>(ptr),
389 #endif // CUB_PTX_ARCH <= 130
396 template <
typename T>
399 Int2Type<LOAD_VOLATILE> modifier,
400 Int2Type<true> is_pointer)
403 return ThreadLoadVolatilePointer(ptr, Int2Type<Traits<T>::PRIMITIVE>());
410 template <
typename T,
int MODIFIER>
413 Int2Type<MODIFIER> modifier,
414 Int2Type<true> is_pointer)
416 typedef typename UnitWord<T>::DeviceWord DeviceWord;
418 const int DEVICE_MULTIPLE =
sizeof(T) /
sizeof(DeviceWord);
420 DeviceWord words[DEVICE_MULTIPLE];
422 IterateThreadLoad<0, DEVICE_MULTIPLE>::template Load<CacheLoadModifier(MODIFIER)>(
423 reinterpret_cast<DeviceWord*
>(
const_cast<T*
>(ptr)),
426 return *
reinterpret_cast<T*
>(words);
435 typename InputIteratorT>
436 __device__ __forceinline__
typename std::iterator_traits<InputIteratorT>::value_type
ThreadLoad(InputIteratorT itr)
441 Int2Type<MODIFIER>(),
442 Int2Type<IsPointer<InputIteratorT>::VALUE>());
447 #endif // DOXYGEN_SHOULD_SKIP_THIS