|
CUB
|
BlockDiscontinuity provides operations for flagging discontinuities within a list of data items partitioned across a CUDA threadblock.
| T | The data type to be exchanged. |
| BLOCK_THREADS | The threadblock size in threads. |
tile_predecessor and last_tile_item) are only considered valid in thread0__syncthreads() barrier is required if the supplied BlockDiscontinuity::SmemStorage is to be reused or repurposed by the threadblockPublic Types | |
| typedef _SmemStorage | SmemStorage |
The operations exposed by BlockDiscontinuity require shared memory of this type. This opaque storage can be allocated directly using the __shared__ keyword. Alternatively, it can be aliased to externally allocated shared memory or union'd with other types to facilitate shared memory reuse. | |
Static Public Methods | |
| template<int ITEMS_PER_THREAD, typename FlagT , typename FlagOp > | |
| static __device__ __forceinline__ void | Flag (SmemStorage &smem_storage, T(&input)[ITEMS_PER_THREAD], FlagOp flag_op, FlagT(&flags)[ITEMS_PER_THREAD], T &last_tile_item) |
| Sets discontinuity flags for a tile of threadblock items, for which the first item has no reference (and is always flagged). The last tile item of the last thread is also returned to thread0. More... | |
| template<int ITEMS_PER_THREAD, typename FlagT , typename FlagOp > | |
| static __device__ __forceinline__ void | Flag (SmemStorage &smem_storage, T(&input)[ITEMS_PER_THREAD], FlagOp flag_op, FlagT(&flags)[ITEMS_PER_THREAD]) |
| Sets discontinuity flags for a tile of threadblock items, for which the first item has no reference (and is always flagged). More... | |
| template<int ITEMS_PER_THREAD, typename FlagT , typename FlagOp > | |
| static __device__ __forceinline__ void | Flag (SmemStorage &smem_storage, T(&input)[ITEMS_PER_THREAD], T tile_predecessor, FlagOp flag_op, FlagT(&flags)[ITEMS_PER_THREAD], T &last_tile_item) |
| Sets discontinuity flags for a tile of threadblock items. The last tile item of the last thread is also returned to thread0. More... | |
| template<int ITEMS_PER_THREAD, typename FlagT , typename FlagOp > | |
| static __device__ __forceinline__ void | Flag (SmemStorage &smem_storage, T(&input)[ITEMS_PER_THREAD], T tile_predecessor, FlagOp flag_op, FlagT(&flags)[ITEMS_PER_THREAD]) |
| Sets discontinuity flags for a tile of threadblock items. More... | |
|
inlinestatic |
Sets discontinuity flags for a tile of threadblock items, for which the first item has no reference (and is always flagged). The last tile item of the last thread is also returned to thread0.
Assuming a blocked arrangement of elements across threads, flagsi is set non-zero for item inputi when scan_op(previous-item, inputi) is true (where previous-item is either inputi-1, or inputITEMS_PER_THREAD-1 in the previous thread). Furthermore, flagsi is always non-zero for input0 in thread0.
The last_tile_item is undefined in threads other than thread0.
A subsequent __syncthreads() threadblock barrier should be invoked after calling this method if the supplied smem_storage is to be reused or repurposed by the threadblock.
| ITEMS_PER_THREAD | [inferred] The number of consecutive items partitioned onto each thread. |
| FlagT | [inferred] The flag type (must be an integer type) |
| FlagOp | [inferred] Binary boolean functor type, having input parameters (const T &a, const T &b) and returning true if a discontinuity exists between a and b, otherwise false. |
| [in] | smem_storage | Shared reference to opaque SmemStorage layout |
| [in] | input | Input items |
| [in] | flag_op | Binary boolean flag predicate |
| [out] | flags | Discontinuity flags |
| [out] | last_tile_item | [thread0 only] The last tile item (inputITEMS_PER_THREAD-1 from threadBLOCK_THREADS-1) |
|
inlinestatic |
Sets discontinuity flags for a tile of threadblock items, for which the first item has no reference (and is always flagged).
Assuming a blocked arrangement of elements across threads, flagsi is set non-zero for item inputi when scan_op(previous-item, inputi) is true (where previous-item is either inputi-1, or inputITEMS_PER_THREAD-1 in the previous thread). Furthermore, flagsi is always non-zero for input0 in thread0.
The last_tile_item is undefined in threads other than thread0.
A subsequent __syncthreads() threadblock barrier should be invoked after calling this method if the supplied smem_storage is to be reused or repurposed by the threadblock.
| ITEMS_PER_THREAD | [inferred] The number of consecutive items partitioned onto each thread. |
| FlagT | [inferred] The flag type (must be an integer type) |
| FlagOp | [inferred] Binary boolean functor type, having input parameters (const T &a, const T &b) and returning true if a discontinuity exists between a and b, otherwise false. |
| [in] | smem_storage | Shared reference to opaque SmemStorage layout |
| [in] | input | Input items |
| [in] | flag_op | Binary boolean flag predicate |
| [out] | flags | Discontinuity flags |
|
inlinestatic |
Sets discontinuity flags for a tile of threadblock items. The last tile item of the last thread is also returned to thread0.
Assuming a blocked arrangement of elements across threads, flagsi is set non-zero for item inputi when scan_op(previous-item, inputi) is true (where previous-item is either inputi-1, or inputITEMS_PER_THREAD-1 in the previous thread). For thread0, item input0 is compared against /p tile_predecessor.
The tile_predecessor and last_tile_item are undefined in threads other than thread0.
A subsequent __syncthreads() threadblock barrier should be invoked after calling this method if the supplied smem_storage is to be reused or repurposed by the threadblock.
| ITEMS_PER_THREAD | [inferred] The number of consecutive items partitioned onto each thread. |
| FlagT | [inferred] The flag type (must be an integer type) |
| FlagOp | [inferred] Binary boolean functor type, having input parameters (const T &a, const T &b) and returning true if a discontinuity exists between a and b, otherwise false. |
| [in] | smem_storage | Shared reference to opaque SmemStorage layout |
| [in] | input | Input items |
| [in] | tile_predecessor | [thread0 only] Item with which to compare the first tile item (input0from thread0). |
| [in] | flag_op | Binary boolean flag predicate |
| [out] | flags | Discontinuity flags |
| [out] | last_tile_item | [thread0 only] The last tile item (inputITEMS_PER_THREAD-1 from threadBLOCK_THREADS-1) |
|
inlinestatic |
Sets discontinuity flags for a tile of threadblock items.
Assuming a blocked arrangement of elements across threads, flagsi is set non-zero for item inputi when scan_op(previous-item, inputi) is true (where previous-item is either inputi-1, or inputITEMS_PER_THREAD-1 in the previous thread). For thread0, item input0 is compared against /p tile_predecessor.
The tile_predecessor and last_tile_item are undefined in threads other than thread0.
A subsequent __syncthreads() threadblock barrier should be invoked after calling this method if the supplied smem_storage is to be reused or repurposed by the threadblock.
| ITEMS_PER_THREAD | [inferred] The number of consecutive items partitioned onto each thread. |
| FlagT | [inferred] The flag type (must be an integer type) |
| FlagOp | [inferred] Binary boolean functor type, having input parameters (const T &a, const T &b) and returning true if a discontinuity exists between a and b, otherwise false. |
| [in] | smem_storage | Shared reference to opaque SmemStorage layout |
| [in] | input | Input items |
| [in] | tile_predecessor | [thread0 only] Item with which to compare the first tile item (input0from thread0). |
| [in] | flag_op | Binary boolean flag predicate |
| [out] | flags | Discontinuity flags |
1.8.3.1