All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Pages
viennacl::linalg::cuda::amg Namespace Reference

## Functions

__global__ void | amg_influence_trivial_kernel (const unsigned int *row_indices, const unsigned int *column_indices, unsigned int size1, unsigned int nnz, unsigned int *influences_row, unsigned int *influences_id, unsigned int *influences_values)
template<typename NumericT >
void | amg_influence_trivial (compressed_matrix< NumericT > const &A, viennacl::linalg::detail::amg::amg_level_context &amg_context, viennacl::linalg::amg_tag &tag)
| Routine for taking all connections in the matrix as strong. More...
template<typename NumericT >
void | amg_influence_advanced (compressed_matrix< NumericT > const &A, viennacl::linalg::detail::amg::amg_level_context &amg_context, viennacl::linalg::amg_tag &tag)
| Routine for extracting strongly connected points considering a user-provided threshold value. More...
template<typename NumericT >
void | amg_influence (compressed_matrix< NumericT > const &A, viennacl::linalg::detail::amg::amg_level_context &amg_context, viennacl::linalg::amg_tag &tag)
| Dispatcher for influence processing. More...
void | enumerate_coarse_points (viennacl::linalg::detail::amg::amg_level_context &amg_context)
| Assign IDs to coarse points. More...
template<typename IndexT >
__global__ void | amg_pmis2_init_workdata (IndexT *work_state, IndexT *work_random, IndexT *work_index, IndexT const *point_types, IndexT const *random_weights, unsigned int size)
| CUDA kernel initializing the work vectors at each PMIS iteration. More...
template<typename IndexT >
__global__ void | amg_pmis2_max_neighborhood (IndexT const *work_state, IndexT const *work_random, IndexT const *work_index, IndexT *work_state2, IndexT *work_random2, IndexT *work_index2, IndexT const *influences_row, IndexT const *influences_id, unsigned int size)
| CUDA kernel propagating the state triple (status, weight, nodeID) to neighbors using a max()-operation. More...
template<typename IndexT >
__global__ void | amg_pmis2_mark_mis_nodes (IndexT const *work_state, IndexT const *work_index, IndexT *point_types, IndexT *undecided_buffer, unsigned int size)
| CUDA kernel for marking MIS and non-MIS nodes. More...
__global__ void | amg_pmis2_reset_state (unsigned int *point_types, unsigned int size)
| CUDA kernel for resetting non-MIS (i.e. coarse) points to undecided so that subsequent kernels work. More...
template<typename NumericT >
void | amg_coarse_ag_stage1_mis2 (compressed_matrix< NumericT > const &A, viennacl::linalg::detail::amg::amg_level_context &amg_context, viennacl::linalg::amg_tag &tag)
| AG (aggregation based) coarsening, single-threaded version of stage 1. More...
template<typename IndexT >
__global__ void | amg_agg_propagate_coarse_indices (IndexT *point_types, IndexT *coarse_ids, IndexT const *influences_row, IndexT const *influences_id, unsigned int size)
template<typename IndexT >
__global__ void | amg_agg_merge_undecided (IndexT *point_types, IndexT *coarse_ids, IndexT const *influences_row, IndexT const *influences_id, unsigned int size)
__global__ void | amg_agg_merge_undecided_2 (unsigned int *point_types, unsigned int size)
template<typename NumericT >
void | amg_coarse_ag (compressed_matrix< NumericT > const &A, viennacl::linalg::detail::amg::amg_level_context &amg_context, viennacl::linalg::amg_tag &tag)
| AG (aggregation based) coarsening. Partially single-threaded version (VIENNACL_AMG_COARSE_AG) More...
template<typename InternalT1 >
void | amg_coarse (InternalT1 &A, viennacl::linalg::detail::amg::amg_level_context &amg_context, viennacl::linalg::amg_tag &tag)
| Calls the right coarsening procedure. More...
template<typename NumericT >
__global__ void | amg_interpol_ag_kernel (unsigned int *P_row_buffer, unsigned int *P_col_buffer, NumericT *P_elements, unsigned int *coarse_ids, unsigned int size)
template<typename NumericT >
void | amg_interpol_ag (compressed_matrix< NumericT > const &A, compressed_matrix< NumericT > &P, viennacl::linalg::detail::amg::amg_level_context &amg_context, viennacl::linalg::amg_tag &tag)
| AG (aggregation based) interpolation. Multi-Threaded! (VIENNACL_INTERPOL_SA) More...
template<typename NumericT >
__global__ void | amg_interpol_sa_kernel (const unsigned int *A_row_indices, const unsigned int *A_col_indices, const NumericT *A_elements, unsigned int A_size1, unsigned int A_nnz, unsigned int *Jacobi_row_indices, unsigned int *Jacobi_col_indices, NumericT *Jacobi_elements, NumericT omega)
template<typename NumericT >
void | amg_interpol_sa (compressed_matrix< NumericT > const &A, compressed_matrix< NumericT > &P, viennacl::linalg::detail::amg::amg_level_context &amg_context, viennacl::linalg::amg_tag &tag)
| Smoothed aggregation interpolation. (VIENNACL_INTERPOL_SA) More...
template<typename MatrixT >
void | amg_interpol (MatrixT const &A, MatrixT &P, viennacl::linalg::detail::amg::amg_level_context &amg_context, viennacl::linalg::amg_tag &tag)
| Dispatcher for building the interpolation matrix. More...
template<typename NumericT >
__global__ void | compressed_matrix_assign_to_dense (const unsigned int *row_indices, const unsigned int *column_indices, const NumericT *elements, NumericT *B, unsigned int B_row_start, unsigned int B_col_start, unsigned int B_row_inc, unsigned int B_col_inc, unsigned int B_row_size, unsigned int B_col_size, unsigned int B_internal_rows, unsigned int B_internal_cols)
template<typename NumericT , unsigned int AlignmentV>
void | assign_to_dense (viennacl::compressed_matrix< NumericT, AlignmentV > const &A, viennacl::matrix_base< NumericT > &B)
template<typename NumericT >
__global__ void | compressed_matrix_smooth_jacobi_kernel (const unsigned int *row_indices, const unsigned int *column_indices, const NumericT *elements, NumericT weight, const NumericT *x_old, NumericT *x_new, const NumericT *rhs, unsigned int size)
template<typename NumericT >
void | smooth_jacobi (unsigned int iterations, compressed_matrix< NumericT > const &A, vector< NumericT > &x, vector< NumericT > &x_backup, vector< NumericT > const &rhs_smooth, NumericT weight)
| Damped Jacobi Smoother (CUDA version) More...

## Function Documentation

template<typename IndexT >

__global__ void viennacl::linalg::cuda::amg::amg_agg_merge_undecided | ( | IndexT * | point_types,
|  | IndexT * | coarse_ids,
|  | IndexT const * | influences_row,
|  | IndexT const * | influences_id,
|  | unsigned int | size
| ) |  |

Definition at line 416 of file amg_operations.hpp.

__global__ void viennacl::linalg::cuda::amg::amg_agg_merge_undecided_2 | ( | unsigned int * | point_types,
|  | unsigned int | size
| ) |  |

Definition at line 445 of file amg_operations.hpp.

template<typename IndexT >

__global__ void viennacl::linalg::cuda::amg::amg_agg_propagate_coarse_indices | ( | IndexT * | point_types,
|  | IndexT * | coarse_ids,
|  | IndexT const * | influences_row,
|  | IndexT const * | influences_id,
|  | unsigned int | size
| ) |  |

Definition at line 386 of file amg_operations.hpp.

template<typename InternalT1 >

void viennacl::linalg::cuda::amg::amg_coarse | ( | InternalT1 & | A,
|  | viennacl::linalg::detail::amg::amg_level_context & | amg_context,
|  | viennacl::linalg::amg_tag & | tag
| ) |  |

Calls the right coarsening procedure.

Parameters

  A | Operator matrix on all levels
  amg_context | AMG hierarchy datastructures
  tag | AMG preconditioner tag

Definition at line 526 of file amg_operations.hpp.

template<typename NumericT >

void viennacl::linalg::cuda::amg::amg_coarse_ag | ( | compressed_matrix< NumericT > const & | A,
|  | viennacl::linalg::detail::amg::amg_level_context & | amg_context,
|  | viennacl::linalg::amg_tag & | tag
| ) |  |

AG (aggregation based) coarsening. Partially single-threaded version (VIENNACL_AMG_COARSE_AG)

Parameters

  A | Operator matrix
  amg_context | AMG hierarchy datastructures
  tag | AMG preconditioner tag

Definition at line 466 of file amg_operations.hpp.

template<typename NumericT >

void viennacl::linalg::cuda::amg::amg_coarse_ag_stage1_mis2 | ( | compressed_matrix< NumericT > const & | A,
|  | viennacl::linalg::detail::amg::amg_level_context & | amg_context,
|  | viennacl::linalg::amg_tag & | tag
| ) |  |

AG (aggregation based) coarsening, single-threaded version of stage 1.

Parameters

  A | Operator matrix on all levels
  amg_context | AMG hierarchy datastructures
  tag | AMG preconditioner tag

Definition at line 290 of file amg_operations.hpp.

template<typename NumericT >

void viennacl::linalg::cuda::amg::amg_influence | ( | compressed_matrix< NumericT > const & | A,
|  | viennacl::linalg::detail::amg::amg_level_context & | amg_context,
|  | viennacl::linalg::amg_tag & | tag
| ) |  |

Dispatcher for influence processing.

Definition at line 103 of file amg_operations.hpp.

template<typename NumericT >

void viennacl::linalg::cuda::amg::amg_influence_advanced | ( | compressed_matrix< NumericT > const & | A,
|  | viennacl::linalg::detail::amg::amg_level_context & | amg_context,
|  | viennacl::linalg::amg_tag & | tag
| ) |  |

Routine for extracting strongly connected points considering a user-provided threshold value.

Definition at line 94 of file amg_operations.hpp.

template<typename NumericT >

void viennacl::linalg::cuda::amg::amg_influence_trivial | ( | compressed_matrix< NumericT > const & | A,
|  | viennacl::linalg::detail::amg::amg_level_context & | amg_context,
|  | viennacl::linalg::amg_tag & | tag
| ) |  |

Routine for taking all connections in the matrix as strong.

Definition at line 74 of file amg_operations.hpp.

__global__ void viennacl::linalg::cuda::amg::amg_influence_trivial_kernel | ( | const unsigned int * | row_indices,
|  | const unsigned int * | column_indices,
|  | unsigned int | size1,
|  | unsigned int | nnz,
|  | unsigned int * | influences_row,
|  | unsigned int * | influences_id,
|  | unsigned int * | influences_values
| ) |  |

Definition at line 44 of file amg_operations.hpp.

template<typename MatrixT >

void viennacl::linalg::cuda::amg::amg_interpol | ( | MatrixT const & | A,
|  | MatrixT & | P,
|  | viennacl::linalg::detail::amg::amg_level_context & | amg_context,
|  | viennacl::linalg::amg_tag & | tag
| ) |  |

Dispatcher for building the interpolation matrix.

Parameters

  A | Operator matrix
  P | Prolongation matrix
  amg_context | AMG hierarchy datastructures
  tag | AMG configuration tag

Definition at line 693 of file amg_operations.hpp.

template<typename NumericT >

void viennacl::linalg::cuda::amg::amg_interpol_ag | ( | compressed_matrix< NumericT > const & | A,
|  | compressed_matrix< NumericT > & | P,
|  | viennacl::linalg::detail::amg::amg_level_context & | amg_context,
|  | viennacl::linalg::amg_tag & | tag
| ) |  |

AG (aggregation based) interpolation. Multi-Threaded! (VIENNACL_INTERPOL_SA)

Parameters

  A | Operator matrix
  P | Prolongation matrix
  amg_context | AMG hierarchy datastructures
  tag | AMG configuration tag

Definition at line 572 of file amg_operations.hpp.

template<typename NumericT >

__global__ void viennacl::linalg::cuda::amg::amg_interpol_ag_kernel | ( | unsigned int * | P_row_buffer,
|  | unsigned int * | P_col_buffer,
|  | NumericT * | P_elements,
|  | unsigned int * | coarse_ids,
|  | unsigned int | size
| ) |  |

Definition at line 543 of file amg_operations.hpp.

template<typename NumericT >

void viennacl::linalg::cuda::amg::amg_interpol_sa | ( | compressed_matrix< NumericT > const & | A,
|  | compressed_matrix< NumericT > & | P,
|  | viennacl::linalg::detail::amg::amg_level_context & | amg_context,
|  | viennacl::linalg::amg_tag & | tag
| ) |  |

Smoothed aggregation interpolation. (VIENNACL_INTERPOL_SA)

Parameters

  A | Operator matrix
  P | Prolongation matrix
  amg_context | AMG hierarchy datastructures
  tag | AMG configuration tag

Definition at line 654 of file amg_operations.hpp.

template<typename NumericT >

__global__ void viennacl::linalg::cuda::amg::amg_interpol_sa_kernel | ( | const unsigned int * | A_row_indices,
|  | const unsigned int * | A_col_indices,
|  | const NumericT * | A_elements,
|  | unsigned int | A_size1,
|  | unsigned int | A_nnz,
|  | unsigned int * | Jacobi_row_indices,
|  | unsigned int * | Jacobi_col_indices,
|  | NumericT * | Jacobi_elements,
|  | NumericT | omega
| ) |  |

Definition at line 594 of file amg_operations.hpp.

template<typename IndexT >

__global__ void viennacl::linalg::cuda::amg::amg_pmis2_init_workdata | ( | IndexT * | work_state,
|  | IndexT * | work_random,
|  | IndexT * | work_index,
|  | IndexT const * | point_types,
|  | IndexT const * | random_weights,
|  | unsigned int | size
| ) |  |

CUDA kernel initializing the work vectors at each PMIS iteration.

Definition at line 139 of file amg_operations.hpp.

template<typename IndexT >

__global__ void viennacl::linalg::cuda::amg::amg_pmis2_mark_mis_nodes | ( | IndexT const * | work_state,
|  | IndexT const * | work_index,
|  | IndexT * | point_types,
|  | IndexT * | undecided_buffer,
|  | unsigned int | size
| ) |  |

CUDA kernel for marking MIS and non-MIS nodes.

Definition at line 229 of file amg_operations.hpp.

template<typename IndexT >

__global__ void viennacl::linalg::cuda::amg::amg_pmis2_max_neighborhood | ( | IndexT const * | work_state,
|  | IndexT const * | work_random,
|  | IndexT const * | work_index,
|  | IndexT * | work_state2,
|  | IndexT * | work_random2,
|  | IndexT * | work_index2,
|  | IndexT const * | influences_row,
|  | IndexT const * | influences_id,
|  | unsigned int | size
| ) |  |

CUDA kernel propagating the state triple (status, weight, nodeID) to neighbors using a max()-operation.

Definition at line 167 of file amg_operations.hpp.

__global__ void viennacl::linalg::cuda::amg::amg_pmis2_reset_state | ( | unsigned int * | point_types,
|  | unsigned int | size
| ) |  |

CUDA kernel for resetting non-MIS (i.e. coarse) points to undecided so that subsequent kernels work.

Definition at line 271 of file amg_operations.hpp.

template<typename NumericT , unsigned int AlignmentV>

void viennacl::linalg::cuda::amg::assign_to_dense | ( | viennacl::compressed_matrix< NumericT, AlignmentV > const & | A,
|  | viennacl::matrix_base< NumericT > & | B
| ) |  |

Definition at line 730 of file amg_operations.hpp.

template<typename NumericT >

__global__ void viennacl::linalg::cuda::amg::compressed_matrix_assign_to_dense | ( | const unsigned int * | row_indices,
|  | const unsigned int * | column_indices,
|  | const NumericT * | elements,
|  | NumericT * | B,
|  | unsigned int | B_row_start,
|  | unsigned int | B_col_start,
|  | unsigned int | B_row_inc,
|  | unsigned int | B_col_inc,
|  | unsigned int | B_row_size,
|  | unsigned int | B_col_size,
|  | unsigned int | B_internal_rows,
|  | unsigned int | B_internal_cols
| ) |  |

Definition at line 708 of file amg_operations.hpp.

template<typename NumericT >

__global__ void viennacl::linalg::cuda::amg::compressed_matrix_smooth_jacobi_kernel | ( | const unsigned int * | row_indices,
|  | const unsigned int * | column_indices,
|  | const NumericT * | elements,
|  | NumericT | weight,
|  | const NumericT * | x_old,
|  | NumericT * | x_new,
|  | const NumericT * | rhs,
|  | unsigned int | size
| ) |  |

Definition at line 749 of file amg_operations.hpp.

void viennacl::linalg::cuda::amg::enumerate_coarse_points | ( | viennacl::linalg::detail::amg::amg_level_context & | amg_context | ) |

inline

Assign IDs to coarse points.

TODO: Use exclusive_scan on GPU for this.

Definition at line 115 of file amg_operations.hpp.

template<typename NumericT >

void viennacl::linalg::cuda::amg::smooth_jacobi | ( | unsigned int | iterations,
|  | compressed_matrix< NumericT > const & | A,
|  | vector< NumericT > & | x,
|  | vector< NumericT > & | x_backup,
|  | vector< NumericT > const & | rhs_smooth,
|  | NumericT | weight
| ) |  |

Damped Jacobi Smoother (CUDA version)

Parameters

  iterations | Number of smoother iterations
  A | Operator matrix for the smoothing
  x | The vector smoothing is applied to
  x_backup | (Different) Vector holding the same values as x
  rhs_smooth | The right hand side of the equation for the smoother
  weight | Damping factor. 0: No effect of smoother. 1: Undamped Jacobi iteration

Definition at line 791 of file amg_operations.hpp.