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. More...

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

Member Typedef Documentation

◆ BlockExchangeT

template<typename AgentSpmvPolicyT , typename ValueT , typename OffsetT , bool HAS_ALPHA, bool HAS_BETA, int PTX_ARCH = CUB_PTX_ARCH>

typedef BlockExchange< ValueT, BLOCK_THREADS, ITEMS_PER_THREAD> cub::AgentSpmv< AgentSpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA, PTX_ARCH >::BlockExchangeT

Definition at line 207 of file agent_spmv_orig.cuh.

◆ BlockPrefixSumT

template<typename AgentSpmvPolicyT , typename ValueT , typename OffsetT , bool HAS_ALPHA, bool HAS_BETA, int PTX_ARCH = CUB_PTX_ARCH>

typedef BlockScan< ValueT, BLOCK_THREADS, AgentSpmvPolicyT::SCAN_ALGORITHM> cub::AgentSpmv< AgentSpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA, PTX_ARCH >::BlockPrefixSumT

Definition at line 200 of file agent_spmv_orig.cuh.

◆ BlockReduceT

template<typename AgentSpmvPolicyT , typename ValueT , typename OffsetT , bool HAS_ALPHA, bool HAS_BETA, int PTX_ARCH = CUB_PTX_ARCH>

typedef BlockReduce< ValueT, BLOCK_THREADS, BLOCK_REDUCE_WARP_REDUCTIONS> cub::AgentSpmv< AgentSpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA, PTX_ARCH >::BlockReduceT

Definition at line 186 of file agent_spmv_orig.cuh.

◆ BlockScanT

template<typename AgentSpmvPolicyT , typename ValueT , typename OffsetT , bool HAS_ALPHA, bool HAS_BETA, int PTX_ARCH = CUB_PTX_ARCH>

typedef BlockScan< KeyValuePairT, BLOCK_THREADS, AgentSpmvPolicyT::SCAN_ALGORITHM> cub::AgentSpmv< AgentSpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA, PTX_ARCH >::BlockScanT

Definition at line 193 of file agent_spmv_orig.cuh.

◆ ColumnIndicesIteratorT

template<typename AgentSpmvPolicyT , typename ValueT , typename OffsetT , bool HAS_ALPHA, bool HAS_BETA, int PTX_ARCH = CUB_PTX_ARCH>

typedef CacheModifiedInputIterator< AgentSpmvPolicyT::COLUMN_INDICES_LOAD_MODIFIER, OffsetT, OffsetT> cub::AgentSpmv< AgentSpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA, PTX_ARCH >::ColumnIndicesIteratorT

Definition at line 161 of file agent_spmv_orig.cuh.

◆ CoordinateT

template<typename AgentSpmvPolicyT , typename ValueT , typename OffsetT , bool HAS_ALPHA, bool HAS_BETA, int PTX_ARCH = CUB_PTX_ARCH>

typedef CubVector<OffsetT,2>::Type cub::AgentSpmv< AgentSpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA, PTX_ARCH >::CoordinateT

2D merge path coordinate type

Definition at line 141 of file agent_spmv_orig.cuh.

◆ KeyValuePairT

template<typename AgentSpmvPolicyT , typename ValueT , typename OffsetT , bool HAS_ALPHA, bool HAS_BETA, int PTX_ARCH = CUB_PTX_ARCH>

typedef KeyValuePair<OffsetT, ValueT> cub::AgentSpmv< AgentSpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA, PTX_ARCH >::KeyValuePairT

Definition at line 176 of file agent_spmv_orig.cuh.

◆ ReduceBySegmentOpT

template<typename AgentSpmvPolicyT , typename ValueT , typename OffsetT , bool HAS_ALPHA, bool HAS_BETA, int PTX_ARCH = CUB_PTX_ARCH>

typedef ReduceByKeyOp<cub::Sum> cub::AgentSpmv< AgentSpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA, PTX_ARCH >::ReduceBySegmentOpT

Definition at line 179 of file agent_spmv_orig.cuh.

◆ RowOffsetsIteratorT

template<typename AgentSpmvPolicyT , typename ValueT , typename OffsetT , bool HAS_ALPHA, bool HAS_BETA, int PTX_ARCH = CUB_PTX_ARCH>

typedef CacheModifiedInputIterator< AgentSpmvPolicyT::ROW_OFFSETS_LOAD_MODIFIER, OffsetT, OffsetT> cub::AgentSpmv< AgentSpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA, PTX_ARCH >::RowOffsetsIteratorT

Definition at line 155 of file agent_spmv_orig.cuh.

◆ RowOffsetsSearchIteratorT

template<typename AgentSpmvPolicyT , typename ValueT , typename OffsetT , bool HAS_ALPHA, bool HAS_BETA, int PTX_ARCH = CUB_PTX_ARCH>

typedef CacheModifiedInputIterator< AgentSpmvPolicyT::ROW_OFFSETS_SEARCH_LOAD_MODIFIER, OffsetT, OffsetT> cub::AgentSpmv< AgentSpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA, PTX_ARCH >::RowOffsetsSearchIteratorT

Input iterator wrapper types (for applying cache modifiers)

Definition at line 149 of file agent_spmv_orig.cuh.

◆ ValueIteratorT

template<typename AgentSpmvPolicyT , typename ValueT , typename OffsetT , bool HAS_ALPHA, bool HAS_BETA, int PTX_ARCH = CUB_PTX_ARCH>

typedef CacheModifiedInputIterator< AgentSpmvPolicyT::VALUES_LOAD_MODIFIER, ValueT, OffsetT> cub::AgentSpmv< AgentSpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA, PTX_ARCH >::ValueIteratorT

Definition at line 167 of file agent_spmv_orig.cuh.

◆ VectorValueIteratorT

template<typename AgentSpmvPolicyT , typename ValueT , typename OffsetT , bool HAS_ALPHA, bool HAS_BETA, int PTX_ARCH = CUB_PTX_ARCH>

typedef CacheModifiedInputIterator< AgentSpmvPolicyT::VECTOR_VALUES_LOAD_MODIFIER, ValueT, OffsetT> cub::AgentSpmv< AgentSpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA, PTX_ARCH >::VectorValueIteratorT

Definition at line 173 of file agent_spmv_orig.cuh.

Member Enumeration Documentation

◆ anonymous enum

template<typename AgentSpmvPolicyT , typename ValueT , typename OffsetT , bool HAS_ALPHA, bool HAS_BETA, int PTX_ARCH = CUB_PTX_ARCH>

anonymous enum

Constants.

Definition at line 133 of file agent_spmv_orig.cuh.

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_storage	Reference to temp_storage
spmv_params	SpMV 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__ 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_coordinates	Pointer to the temporary array of tile starting coordinates
[out]	d_tile_carry_pairs	Pointer to the temporary array carry-out dot product row-ids, one per block
[in]	num_merge_tiles	Number of merge tiles

Definition at line 602 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_load Marker 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__ 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_load Marker type indicating whether to load nonzeros directly during path-discovery or beforehand in batch

Definition at line 290 of file agent_spmv_orig.cuh.

Field Documentation

◆ spmv_params

template<typename AgentSpmvPolicyT , typename ValueT , typename OffsetT , bool HAS_ALPHA, bool HAS_BETA, int PTX_ARCH = CUB_PTX_ARCH>

SpmvParams<ValueT, OffsetT>& cub::AgentSpmv< AgentSpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA, PTX_ARCH >::spmv_params

Reference to temp_storage.

Definition at line 255 of file agent_spmv_orig.cuh.

◆ temp_storage

template<typename AgentSpmvPolicyT , typename ValueT , typename OffsetT , bool HAS_ALPHA, bool HAS_BETA, int PTX_ARCH = CUB_PTX_ARCH>

_TempStorage& cub::AgentSpmv< AgentSpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA, PTX_ARCH >::temp_storage

Definition at line 253 of file agent_spmv_orig.cuh.

◆ wd_column_indices

template<typename AgentSpmvPolicyT , typename ValueT , typename OffsetT , bool HAS_ALPHA, bool HAS_BETA, int PTX_ARCH = CUB_PTX_ARCH>

ColumnIndicesIteratorT cub::AgentSpmv< AgentSpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA, PTX_ARCH >::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.)

Definition at line 259 of file agent_spmv_orig.cuh.

◆ wd_row_end_offsets

template<typename AgentSpmvPolicyT , typename ValueT , typename OffsetT , bool HAS_ALPHA, bool HAS_BETA, int PTX_ARCH = CUB_PTX_ARCH>

RowOffsetsIteratorT cub::AgentSpmv< AgentSpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA, PTX_ARCH >::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.

Definition at line 258 of file agent_spmv_orig.cuh.

◆ wd_values

template<typename AgentSpmvPolicyT , typename ValueT , typename OffsetT , bool HAS_ALPHA, bool HAS_BETA, int PTX_ARCH = CUB_PTX_ARCH>

ValueIteratorT cub::AgentSpmv< AgentSpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA, PTX_ARCH >::wd_values

Wrapped pointer to the array of num_nonzeros values of the corresponding nonzero elements of matrix A.

Definition at line 257 of file agent_spmv_orig.cuh.

◆ wd_vector_x

template<typename AgentSpmvPolicyT , typename ValueT , typename OffsetT , bool HAS_ALPHA, bool HAS_BETA, int PTX_ARCH = CUB_PTX_ARCH>

VectorValueIteratorT cub::AgentSpmv< AgentSpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA, PTX_ARCH >::wd_vector_x

Wrapped Pointer to the array of num_cols values corresponding to the dense input vector x

Definition at line 260 of file agent_spmv_orig.cuh.

◆ wd_vector_y

template<typename AgentSpmvPolicyT , typename ValueT , typename OffsetT , bool HAS_ALPHA, bool HAS_BETA, int PTX_ARCH = CUB_PTX_ARCH>

VectorValueIteratorT cub::AgentSpmv< AgentSpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA, PTX_ARCH >::wd_vector_y

Wrapped Pointer to the array of num_cols values corresponding to the dense input vector x

Definition at line 261 of file agent_spmv_orig.cuh.

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

openfpm_data/src/util/cuda/cub_old/agent/agent_spmv_orig.cuh

Detailed Description

Data Structures

Public Types

Public Member Functions

Data Fields

Member Typedef Documentation

◆ BlockExchangeT

◆ BlockPrefixSumT

◆ BlockReduceT

◆ BlockScanT

◆ ColumnIndicesIteratorT

◆ CoordinateT

◆ KeyValuePairT

◆ ReduceBySegmentOpT

◆ RowOffsetsIteratorT

◆ RowOffsetsSearchIteratorT

◆ ValueIteratorT

◆ VectorValueIteratorT

Member Enumeration Documentation

◆ anonymous enum

Constructor & Destructor Documentation

◆ AgentSpmv()

Member Function Documentation

◆ ConsumeTile() [1/3]

◆ ConsumeTile() [2/3]

◆ ConsumeTile() [3/3]

Field Documentation

◆ spmv_params

◆ temp_storage

◆ wd_column_indices

◆ wd_row_end_offsets

◆ wd_values

◆ wd_vector_x

◆ wd_vector_y