OpenFPM_pdata  4.1.0
Project that contain the implementation of distributed structures
Data Structures | Public Types | Public Member Functions | Data Fields
cub::AgentSpmv< AgentSpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA, PTX_ARCH > Struct Template Reference

AgentSpmv implements a stateful abstraction of CUDA thread blocks for participating in device-wide SpMV. More...

Detailed Description

template<typename AgentSpmvPolicyT, typename ValueT, typename OffsetT, bool HAS_ALPHA, bool HAS_BETA, int PTX_ARCH = CUB_PTX_ARCH>
struct cub::AgentSpmv< AgentSpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA, PTX_ARCH >

AgentSpmv implements a stateful abstraction of CUDA thread blocks for participating in device-wide SpMV.

< PTX compute capability

Definition at line 126 of file agent_spmv_orig.cuh.

Data Structures

struct  _TempStorage
 Shared memory type required by this thread block. More...
 
union  MergeItem
 Merge item type (either a non-zero value or a row-end offset) More...
 
struct  TempStorage
 Temporary storage type (unionable) More...
 

Public Types

enum  { BLOCK_THREADS = AgentSpmvPolicyT::BLOCK_THREADS, ITEMS_PER_THREAD = AgentSpmvPolicyT::ITEMS_PER_THREAD, TILE_ITEMS = BLOCK_THREADS * ITEMS_PER_THREAD }
 Constants.
 
typedef CubVector< OffsetT, 2 >::Type CoordinateT
 2D merge path coordinate type
 
typedef CacheModifiedInputIterator< AgentSpmvPolicyT::ROW_OFFSETS_SEARCH_LOAD_MODIFIER, OffsetT, OffsetT > RowOffsetsSearchIteratorT
 Input iterator wrapper types (for applying cache modifiers)
 
typedef CacheModifiedInputIterator< AgentSpmvPolicyT::ROW_OFFSETS_LOAD_MODIFIER, OffsetT, OffsetT > RowOffsetsIteratorT
 
typedef CacheModifiedInputIterator< AgentSpmvPolicyT::COLUMN_INDICES_LOAD_MODIFIER, OffsetT, OffsetT > ColumnIndicesIteratorT
 
typedef CacheModifiedInputIterator< AgentSpmvPolicyT::VALUES_LOAD_MODIFIER, ValueT, OffsetT > ValueIteratorT
 
typedef CacheModifiedInputIterator< AgentSpmvPolicyT::VECTOR_VALUES_LOAD_MODIFIER, ValueT, OffsetT > VectorValueIteratorT
 
typedef KeyValuePair< OffsetT, ValueT > KeyValuePairT
 
typedef ReduceByKeyOp< cub::Sum > ReduceBySegmentOpT
 
typedef BlockReduce< ValueT, BLOCK_THREADS, BLOCK_REDUCE_WARP_REDUCTIONS > BlockReduceT
 
typedef BlockScan< KeyValuePairT, BLOCK_THREADS, AgentSpmvPolicyT::SCAN_ALGORITHM > BlockScanT
 
typedef BlockScan< ValueT, BLOCK_THREADS, AgentSpmvPolicyT::SCAN_ALGORITHM > BlockPrefixSumT
 
typedef BlockExchange< ValueT, BLOCK_THREADS, ITEMS_PER_THREAD > BlockExchangeT
 

Public Member Functions

__device__ __forceinline__ AgentSpmv (TempStorage &temp_storage, SpmvParams< ValueT, OffsetT > &spmv_params)
 
__device__ __forceinline__ KeyValuePairT ConsumeTile (int tile_idx, CoordinateT tile_start_coord, CoordinateT tile_end_coord, Int2Type< true > is_direct_load)
 
__device__ __forceinline__ KeyValuePairT ConsumeTile (int tile_idx, CoordinateT tile_start_coord, CoordinateT tile_end_coord, Int2Type< false > is_direct_load)
 
__device__ __forceinline__ void ConsumeTile (CoordinateT *d_tile_coordinates, KeyValuePairT *d_tile_carry_pairs, int num_merge_tiles)
 

Data Fields

_TempStorage & temp_storage
 
SpmvParams< ValueT, OffsetT > & spmv_params
 Reference to temp_storage.
 
ValueIteratorT wd_values
 Wrapped pointer to the array of num_nonzeros values of the corresponding nonzero elements of matrix A.
 
RowOffsetsIteratorT wd_row_end_offsets
 Wrapped Pointer to the array of m offsets demarcating the end of every row in d_column_indices and d_values.
 
ColumnIndicesIteratorT wd_column_indices
 Wrapped Pointer to the array of num_nonzeros column-indices of the corresponding nonzero elements of matrix A. (Indices are zero-valued.)
 
VectorValueIteratorT wd_vector_x
 Wrapped Pointer to the array of num_cols values corresponding to the dense input vector x
 
VectorValueIteratorT wd_vector_y
 Wrapped Pointer to the array of num_cols values corresponding to the dense input vector x
 

Constructor & Destructor Documentation

◆ AgentSpmv()

template<typename AgentSpmvPolicyT , typename ValueT , typename OffsetT , bool HAS_ALPHA, bool HAS_BETA, int PTX_ARCH = CUB_PTX_ARCH>
__device__ __forceinline__ cub::AgentSpmv< AgentSpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA, PTX_ARCH >::AgentSpmv ( TempStorage &  temp_storage,
SpmvParams< ValueT, OffsetT > &  spmv_params 
)
inline

Constructor

Parameters
temp_storageReference to temp_storage
spmv_paramsSpMV input parameter bundle

Definition at line 271 of file agent_spmv_orig.cuh.

Member Function Documentation

◆ ConsumeTile() [1/3]

template<typename AgentSpmvPolicyT , typename ValueT , typename OffsetT , bool HAS_ALPHA, bool HAS_BETA, int PTX_ARCH = CUB_PTX_ARCH>
__device__ __forceinline__ KeyValuePairT cub::AgentSpmv< AgentSpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA, PTX_ARCH >::ConsumeTile ( int  tile_idx,
CoordinateT  tile_start_coord,
CoordinateT  tile_end_coord,
Int2Type< true >  is_direct_load 
)
inline

Consume a merge tile, specialized for direct-load of nonzeros

Parameters
is_direct_loadMarker type indicating whether to load nonzeros directly during path-discovery or beforehand in batch

Definition at line 290 of file agent_spmv_orig.cuh.

◆ ConsumeTile() [2/3]

template<typename AgentSpmvPolicyT , typename ValueT , typename OffsetT , bool HAS_ALPHA, bool HAS_BETA, int PTX_ARCH = CUB_PTX_ARCH>
__device__ __forceinline__ KeyValuePairT cub::AgentSpmv< AgentSpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA, PTX_ARCH >::ConsumeTile ( int  tile_idx,
CoordinateT  tile_start_coord,
CoordinateT  tile_end_coord,
Int2Type< false >  is_direct_load 
)
inline

Consume a merge tile, specialized for indirect load of nonzeros

Parameters
is_direct_loadMarker type indicating whether to load nonzeros directly during path-discovery or beforehand in batch

Definition at line 414 of file agent_spmv_orig.cuh.

◆ ConsumeTile() [3/3]

template<typename AgentSpmvPolicyT , typename ValueT , typename OffsetT , bool HAS_ALPHA, bool HAS_BETA, int PTX_ARCH = CUB_PTX_ARCH>
__device__ __forceinline__ void cub::AgentSpmv< AgentSpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA, PTX_ARCH >::ConsumeTile ( CoordinateT *  d_tile_coordinates,
KeyValuePairT *  d_tile_carry_pairs,
int  num_merge_tiles 
)
inline

Consume input tile

Parameters
[in]d_tile_coordinatesPointer to the temporary array of tile starting coordinates
[out]d_tile_carry_pairsPointer to the temporary array carry-out dot product row-ids, one per block
[in]num_merge_tilesNumber of merge tiles

Definition at line 602 of file agent_spmv_orig.cuh.

The documentation for this struct was generated from the following file: