template<typename T, int LOGICAL_WARP_THREADS, int PTX_ARCH>
struct cub::WarpScanSmem< T, LOGICAL_WARP_THREADS, PTX_ARCH >
WarpScanSmem provides smem-based variants of parallel prefix scan of items partitioned across a CUDA thread warp.
< The PTX compute capability for which to to specialize this collective
Definition at line 55 of file warp_scan_smem.cuh.
|
__device__ __forceinline__ | WarpScanSmem (TempStorage &temp_storage) |
| Constructor.
|
|
template<bool HAS_IDENTITY, int STEP, typename ScanOp > |
__device__ __forceinline__ void | ScanStep (T &partial, ScanOp scan_op, Int2Type< STEP >) |
| Basic inclusive scan iteration (template unrolled, inductive-case specialization)
|
|
template<bool HAS_IDENTITY, typename ScanOp > |
__device__ __forceinline__ void | ScanStep (T &, ScanOp, Int2Type< STEPS >) |
| Basic inclusive scan iteration(template unrolled, base-case specialization)
|
|
__device__ __forceinline__ void | InclusiveScan (T input, T &output, Sum scan_op, Int2Type< true >) |
| Inclusive prefix scan (specialized for summation across primitive types) More...
|
|
template<typename ScanOp , int IS_PRIMITIVE> |
__device__ __forceinline__ void | InclusiveScan (T input, T &output, ScanOp scan_op, Int2Type< IS_PRIMITIVE >) |
| Inclusive prefix scan. More...
|
|
__device__ __forceinline__ T | Broadcast (T input, unsigned int src_lane) |
| Broadcast. More...
|
|
template<typename ScanOp > |
__device__ __forceinline__ void | InclusiveScan (T input, T &inclusive_output, ScanOp scan_op) |
| Inclusive scan. More...
|
|
template<typename ScanOp > |
__device__ __forceinline__ void | InclusiveScan (T input, T &inclusive_output, ScanOp scan_op, T &warp_aggregate) |
| Inclusive scan with aggregate. More...
|
|
template<typename ScanOpT , typename IsIntegerT > |
__device__ __forceinline__ void | Update (T, T &inclusive, T &exclusive, ScanOpT, IsIntegerT) |
| Update inclusive and exclusive using input and inclusive. More...
|
|
__device__ __forceinline__ void | Update (T input, T &inclusive, T &exclusive, cub::Sum, Int2Type< true >) |
| Update inclusive and exclusive using input and inclusive (specialized for summation of integer types)
|
|
template<typename ScanOpT , typename IsIntegerT > |
__device__ __forceinline__ void | Update (T, T &inclusive, T &exclusive, ScanOpT scan_op, T initial_value, IsIntegerT) |
| Update inclusive and exclusive using initial value using input, inclusive, and initial value.
|
|
__device__ __forceinline__ void | Update (T input, T &inclusive, T &exclusive, cub::Sum scan_op, T initial_value, Int2Type< true >) |
| Update inclusive and exclusive using initial value using input and inclusive (specialized for summation of integer types)
|
|
template<typename ScanOpT , typename IsIntegerT > |
__device__ __forceinline__ void | Update (T, T &inclusive, T &exclusive, T &warp_aggregate, ScanOpT, IsIntegerT) |
| Update inclusive, exclusive, and warp aggregate using input and inclusive.
|
|
__device__ __forceinline__ void | Update (T input, T &inclusive, T &exclusive, T &warp_aggregate, cub::Sum, Int2Type< true >) |
| Update inclusive, exclusive, and warp aggregate using input and inclusive (specialized for summation of integer types)
|
|
template<typename ScanOpT , typename IsIntegerT > |
__device__ __forceinline__ void | Update (T, T &inclusive, T &exclusive, T &warp_aggregate, ScanOpT scan_op, T initial_value, IsIntegerT) |
| Update inclusive, exclusive, and warp aggregate using input, inclusive, and initial value.
|
|