All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Pages
iterative_operations.hpp
Go to the documentation of this file.
#ifndef VIENNACL_LINALG_ITERATIVE_OPERATIONS_HPP_
#define VIENNACL_LINALG_ITERATIVE_OPERATIONS_HPP_
/* =========================================================================
Copyright (c) 2010-2016, Institute for Microelectronics,
Institute for Analysis and Scientific Computing,
TU Wien.
Portions of this software are copyright by UChicago Argonne, LLC.
-----------------
ViennaCL - The Vienna Computing Library
-----------------
Project Head: Karl Rupp rupp@iue.tuwien.ac.at
(A list of authors and contributors can be found in the manual)
License: MIT (X11), see file LICENSE in the base directory
============================================================================= */
#include "viennacl/forwards.h"
#include "viennacl/range.hpp"
#include "viennacl/scalar.hpp"
#include "viennacl/tools/tools.hpp"
#include "viennacl/meta/predicate.hpp"
#include "viennacl/meta/enable_if.hpp"
#include "viennacl/traits/size.hpp"
#include "viennacl/traits/start.hpp"
#include "viennacl/traits/handle.hpp"
#include "viennacl/traits/stride.hpp"
#include "viennacl/linalg/host_based/iterative_operations.hpp"
#ifdef VIENNACL_WITH_OPENCL
#include "viennacl/linalg/opencl/iterative_operations.hpp"
#endif
#ifdef VIENNACL_WITH_CUDA
#include "viennacl/linalg/cuda/iterative_operations.hpp"
#endif
namespace viennacl
{
namespace linalg
{
template<typename NumericT>
void pipelined_cg_vector_update(vector_base<NumericT> & result,
NumericT alpha,
vector_base<NumericT> & p,
vector_base<NumericT> & r,
vector_base<NumericT> const & Ap,
NumericT beta,
vector_base<NumericT> & inner_prod_buffer)
{
switch (viennacl::traits::handle(result).get_active_handle_id())
{
case viennacl::MAIN_MEMORY:
viennacl::linalg::host_based::pipelined_cg_vector_update(result, alpha, p, r, Ap, beta, inner_prod_buffer);
break;
#ifdef VIENNACL_WITH_OPENCL
case viennacl::OPENCL_MEMORY:
viennacl::linalg::opencl::pipelined_cg_vector_update(result, alpha, p, r, Ap, beta, inner_prod_buffer);
break;
#endif
#ifdef VIENNACL_WITH_CUDA
case viennacl::CUDA_MEMORY:
viennacl::linalg::cuda::pipelined_cg_vector_update(result, alpha, p, r, Ap, beta, inner_prod_buffer);
break;
#endif
case viennacl::MEMORY_NOT_INITIALIZED:
throw memory_exception("not initialised!");
default:
throw memory_exception("not implemented");
}
}
template<typename MatrixT, typename NumericT>
void pipelined_cg_prod(MatrixT const & A,
vector_base<NumericT> const & p,
vector_base<NumericT> & Ap,
vector_base<NumericT> & inner_prod_buffer)
{
switch (viennacl::traits::handle(p).get_active_handle_id())
{
case viennacl::MAIN_MEMORY:
viennacl::linalg::host_based::pipelined_cg_prod(A, p, Ap, inner_prod_buffer);
break;
#ifdef VIENNACL_WITH_OPENCL
case viennacl::OPENCL_MEMORY:
viennacl::linalg::opencl::pipelined_cg_prod(A, p, Ap, inner_prod_buffer);
break;
#endif
#ifdef VIENNACL_WITH_CUDA
case viennacl::CUDA_MEMORY:
viennacl::linalg::cuda::pipelined_cg_prod(A, p, Ap, inner_prod_buffer);
break;
#endif
case viennacl::MEMORY_NOT_INITIALIZED:
throw memory_exception("not initialised!");
default:
throw memory_exception("not implemented");
}
}
template<typename NumericT>
void pipelined_bicgstab_update_s(vector_base<NumericT> & s,
vector_base<NumericT> & r,
vector_base<NumericT> const & Ap,
vector_base<NumericT> & inner_prod_buffer,
vcl_size_t buffer_chunk_size,
vcl_size_t buffer_chunk_offset)
{
switch (viennacl::traits::handle(s).get_active_handle_id())
{
case viennacl::MAIN_MEMORY:
viennacl::linalg::host_based::pipelined_bicgstab_update_s(s, r, Ap, inner_prod_buffer, buffer_chunk_size, buffer_chunk_offset);
break;
#ifdef VIENNACL_WITH_OPENCL
case viennacl::OPENCL_MEMORY:
viennacl::linalg::opencl::pipelined_bicgstab_update_s(s, r, Ap, inner_prod_buffer, buffer_chunk_size, buffer_chunk_offset);
break;
#endif
#ifdef VIENNACL_WITH_CUDA
case viennacl::CUDA_MEMORY:
viennacl::linalg::cuda::pipelined_bicgstab_update_s(s, r, Ap, inner_prod_buffer, buffer_chunk_size, buffer_chunk_offset);
break;
#endif
case viennacl::MEMORY_NOT_INITIALIZED:
throw memory_exception("not initialised!");
default:
throw memory_exception("not implemented");
}
}
template<typename NumericT>
void pipelined_bicgstab_vector_update(vector_base<NumericT> & result, NumericT alpha, vector_base<NumericT> & p, NumericT omega, vector_base<NumericT> const & s,
vector_base<NumericT> & residual, vector_base<NumericT> const & As,
NumericT beta, vector_base<NumericT> const & Ap,
vector_base<NumericT> const & r0star,
vector_base<NumericT> & inner_prod_buffer,
vcl_size_t buffer_chunk_size)
{
switch (viennacl::traits::handle(s).get_active_handle_id())
{
case viennacl::MAIN_MEMORY:
viennacl::linalg::host_based::pipelined_bicgstab_vector_update(result, alpha, p, omega, s, residual, As, beta, Ap, r0star, inner_prod_buffer, buffer_chunk_size);
break;
#ifdef VIENNACL_WITH_OPENCL
case viennacl::OPENCL_MEMORY:
viennacl::linalg::opencl::pipelined_bicgstab_vector_update(result, alpha, p, omega, s, residual, As, beta, Ap, r0star, inner_prod_buffer, buffer_chunk_size);
break;
#endif
#ifdef VIENNACL_WITH_CUDA
case viennacl::CUDA_MEMORY:
viennacl::linalg::cuda::pipelined_bicgstab_vector_update(result, alpha, p, omega, s, residual, As, beta, Ap, r0star, inner_prod_buffer, buffer_chunk_size);
break;
#endif
case viennacl::MEMORY_NOT_INITIALIZED:
throw memory_exception("not initialised!");
default:
throw memory_exception("not implemented");
}
}
template<typename MatrixT, typename NumericT>
void pipelined_bicgstab_prod(MatrixT const & A,
vector_base<NumericT> const & p,
vector_base<NumericT> & Ap,
vector_base<NumericT> const & r0star,
vector_base<NumericT> & inner_prod_buffer,
vcl_size_t buffer_chunk_size,
vcl_size_t buffer_chunk_offset)
{
switch (viennacl::traits::handle(p).get_active_handle_id())
{
case viennacl::MAIN_MEMORY:
viennacl::linalg::host_based::pipelined_bicgstab_prod(A, p, Ap, r0star, inner_prod_buffer, buffer_chunk_size, buffer_chunk_offset);
break;
#ifdef VIENNACL_WITH_OPENCL
case viennacl::OPENCL_MEMORY:
viennacl::linalg::opencl::pipelined_bicgstab_prod(A, p, Ap, r0star, inner_prod_buffer, buffer_chunk_size, buffer_chunk_offset);
break;
#endif
#ifdef VIENNACL_WITH_CUDA
case viennacl::CUDA_MEMORY:
viennacl::linalg::cuda::pipelined_bicgstab_prod(A, p, Ap, r0star, inner_prod_buffer, buffer_chunk_size, buffer_chunk_offset);
break;
#endif
case viennacl::MEMORY_NOT_INITIALIZED:
throw memory_exception("not initialised!");
default:
throw memory_exception("not implemented");
}
}
template <typename T>
void pipelined_gmres_normalize_vk(vector_base<T> & v_k,
vector_base<T> const & residual,
vector_base<T> & R_buffer,
vcl_size_t offset_in_R,
vector_base<T> const & inner_prod_buffer,
vector_base<T> & r_dot_vk_buffer,
vcl_size_t buffer_chunk_size,
vcl_size_t buffer_chunk_offset)
{
switch (viennacl::traits::handle(v_k).get_active_handle_id())
{
case viennacl::MAIN_MEMORY:
viennacl::linalg::host_based::pipelined_gmres_normalize_vk(v_k, residual, R_buffer, offset_in_R, inner_prod_buffer, r_dot_vk_buffer, buffer_chunk_size, buffer_chunk_offset);
break;
#ifdef VIENNACL_WITH_OPENCL
case viennacl::OPENCL_MEMORY:
viennacl::linalg::opencl::pipelined_gmres_normalize_vk(v_k, residual, R_buffer, offset_in_R, inner_prod_buffer, r_dot_vk_buffer, buffer_chunk_size, buffer_chunk_offset);
break;
#endif
#ifdef VIENNACL_WITH_CUDA
case viennacl::CUDA_MEMORY:
viennacl::linalg::cuda::pipelined_gmres_normalize_vk(v_k, residual, R_buffer, offset_in_R, inner_prod_buffer, r_dot_vk_buffer, buffer_chunk_size, buffer_chunk_offset);
break;
#endif
case viennacl::MEMORY_NOT_INITIALIZED:
throw memory_exception("not initialised!");
default:
throw memory_exception("not implemented");
}
}
template <typename T>
void pipelined_gmres_gram_schmidt_stage1(vector_base<T> const & device_krylov_basis,
vcl_size_t v_k_size,
vcl_size_t v_k_internal_size,
vcl_size_t k,
vector_base<T> & vi_in_vk_buffer,
vcl_size_t buffer_chunk_size)
{
switch (viennacl::traits::handle(device_krylov_basis).get_active_handle_id())
{
case viennacl::MAIN_MEMORY:
viennacl::linalg::host_based::pipelined_gmres_gram_schmidt_stage1(device_krylov_basis, v_k_size, v_k_internal_size, k, vi_in_vk_buffer, buffer_chunk_size);
break;
#ifdef VIENNACL_WITH_OPENCL
case viennacl::OPENCL_MEMORY:
viennacl::linalg::opencl::pipelined_gmres_gram_schmidt_stage1(device_krylov_basis, v_k_size, v_k_internal_size, k, vi_in_vk_buffer, buffer_chunk_size);
break;
#endif
#ifdef VIENNACL_WITH_CUDA
case viennacl::CUDA_MEMORY:
viennacl::linalg::cuda::pipelined_gmres_gram_schmidt_stage1(device_krylov_basis, v_k_size, v_k_internal_size, k, vi_in_vk_buffer, buffer_chunk_size);
break;
#endif
case viennacl::MEMORY_NOT_INITIALIZED:
throw memory_exception("not initialised!");
default:
throw memory_exception("not implemented");
}
}
template <typename T>
void pipelined_gmres_gram_schmidt_stage2(vector_base<T> & device_krylov_basis,
vcl_size_t v_k_size,
vcl_size_t v_k_internal_size,
vcl_size_t k,
vector_base<T> const & vi_in_vk_buffer,
vector_base<T> & R_buffer,
vcl_size_t krylov_dim,
vector_base<T> & inner_prod_buffer,
vcl_size_t buffer_chunk_size)
{
switch (viennacl::traits::handle(device_krylov_basis).get_active_handle_id())
{
case viennacl::MAIN_MEMORY:
viennacl::linalg::host_based::pipelined_gmres_gram_schmidt_stage2(device_krylov_basis, v_k_size, v_k_internal_size, k, vi_in_vk_buffer, R_buffer, krylov_dim, inner_prod_buffer, buffer_chunk_size);
break;
#ifdef VIENNACL_WITH_OPENCL
case viennacl::OPENCL_MEMORY:
viennacl::linalg::opencl::pipelined_gmres_gram_schmidt_stage2(device_krylov_basis, v_k_size, v_k_internal_size, k, vi_in_vk_buffer, R_buffer, krylov_dim, inner_prod_buffer, buffer_chunk_size);
break;
#endif
#ifdef VIENNACL_WITH_CUDA
case viennacl::CUDA_MEMORY:
viennacl::linalg::cuda::pipelined_gmres_gram_schmidt_stage2(device_krylov_basis, v_k_size, v_k_internal_size, k, vi_in_vk_buffer, R_buffer, krylov_dim, inner_prod_buffer, buffer_chunk_size);
break;
#endif
case viennacl::MEMORY_NOT_INITIALIZED:
throw memory_exception("not initialised!");
default:
throw memory_exception("not implemented");
}
}
template <typename T>
void pipelined_gmres_update_result(vector_base<T> & result,
vector_base<T> const & residual,
vector_base<T> const & krylov_basis,
vcl_size_t v_k_size,
vcl_size_t v_k_internal_size,
vector_base<T> const & coefficients,
vcl_size_t k)
{
switch (viennacl::traits::handle(result).get_active_handle_id())
{
case viennacl::MAIN_MEMORY:
viennacl::linalg::host_based::pipelined_gmres_update_result(result, residual, krylov_basis, v_k_size, v_k_internal_size, coefficients, k);
break;
#ifdef VIENNACL_WITH_OPENCL
case viennacl::OPENCL_MEMORY:
viennacl::linalg::opencl::pipelined_gmres_update_result(result, residual, krylov_basis, v_k_size, v_k_internal_size, coefficients, k);
break;
#endif
#ifdef VIENNACL_WITH_CUDA
case viennacl::CUDA_MEMORY:
viennacl::linalg::cuda::pipelined_gmres_update_result(result, residual, krylov_basis, v_k_size, v_k_internal_size, coefficients, k);
break;
#endif
case viennacl::MEMORY_NOT_INITIALIZED:
throw memory_exception("not initialised!");
default:
throw memory_exception("not implemented");
}
}
template <typename MatrixType, typename T>
void pipelined_gmres_prod(MatrixType const & A,
vector_base<T> const & p,
vector_base<T> & Ap,
vector_base<T> & inner_prod_buffer)
{
switch (viennacl::traits::handle(p).get_active_handle_id())
{
case viennacl::MAIN_MEMORY:
viennacl::linalg::host_based::pipelined_gmres_prod(A, p, Ap, inner_prod_buffer);
break;
#ifdef VIENNACL_WITH_OPENCL
case viennacl::OPENCL_MEMORY:
viennacl::linalg::opencl::pipelined_gmres_prod(A, p, Ap, inner_prod_buffer);
break;
#endif
#ifdef VIENNACL_WITH_CUDA
case viennacl::CUDA_MEMORY:
viennacl::linalg::cuda::pipelined_gmres_prod(A, p, Ap, inner_prod_buffer);
break;
#endif
case viennacl::MEMORY_NOT_INITIALIZED:
throw memory_exception("not initialised!");
default:
throw memory_exception("not implemented");
}
}
} //namespace linalg
} //namespace viennacl
#endif
viennacl::linalg::host_based::pipelined_gmres_normalize_vk
void pipelined_gmres_normalize_vk(vector_base< T > &v_k, vector_base< T > const &residual, vector_base< T > &R_buffer, vcl_size_t offset_in_R, vector_base< T > const &inner_prod_buffer, vector_base< T > &r_dot_vk_buffer, vcl_size_t buffer_chunk_size, vcl_size_t buffer_chunk_offset)
Performs a vector normalization needed for an efficient pipelined GMRES algorithm.
Definition: iterative_operations.hpp:711
viennacl::linalg::opencl::pipelined_bicgstab_prod
void pipelined_bicgstab_prod(compressed_matrix< NumericT > const &A, vector_base< NumericT > const &p, vector_base< NumericT > &Ap, vector_base< NumericT > const &r0star, vector_base< NumericT > &inner_prod_buffer, vcl_size_t buffer_chunk_size, vcl_size_t buffer_chunk_offset)
Definition: iterative_operations.hpp:360
viennacl::linalg::cuda::pipelined_gmres_gram_schmidt_stage1
void pipelined_gmres_gram_schmidt_stage1(vector_base< T > const &device_krylov_basis, vcl_size_t v_k_size, vcl_size_t v_k_internal_size, vcl_size_t param_k, vector_base< T > &vi_in_vk_buffer, vcl_size_t buffer_chunk_size)
Definition: iterative_operations.hpp:1738
viennacl::memory_exception
Exception class in case of memory errors.
Definition: forwards.h:572
size.hpp
Generic size and resize functionality for different vector and matrix types.
start.hpp
Extracts the underlying OpenCL start index handle from a vector, a matrix, an expression etc...
tools.hpp
Various little tools used here and there in ViennaCL.
viennacl::linalg::pipelined_gmres_gram_schmidt_stage1
void pipelined_gmres_gram_schmidt_stage1(vector_base< T > const &device_krylov_basis, vcl_size_t v_k_size, vcl_size_t v_k_internal_size, vcl_size_t k, vector_base< T > &vi_in_vk_buffer, vcl_size_t buffer_chunk_size)
Computes the first reduction stage for multiple inner products <v_i, v_k>, i=0..k-1.
Definition: iterative_operations.hpp:286
viennacl::linalg::pipelined_gmres_update_result
void pipelined_gmres_update_result(vector_base< T > &result, vector_base< T > const &residual, vector_base< T > const &krylov_basis, vcl_size_t v_k_size, vcl_size_t v_k_internal_size, vector_base< T > const &coefficients, vcl_size_t k)
Computes x += eta_0 r + sum_{i=1}^{k-1} eta_i v_{i-1}.
Definition: iterative_operations.hpp:356
viennacl::linalg::pipelined_gmres_gram_schmidt_stage2
void pipelined_gmres_gram_schmidt_stage2(vector_base< T > &device_krylov_basis, vcl_size_t v_k_size, vcl_size_t v_k_internal_size, vcl_size_t k, vector_base< T > const &vi_in_vk_buffer, vector_base< T > &R_buffer, vcl_size_t krylov_dim, vector_base< T > &inner_prod_buffer, vcl_size_t buffer_chunk_size)
Computes the second reduction stage for multiple inner products <v_i, v_k>, i=0..k-1, then updates v_k -= <v_i, v_k> v_i and computes the first reduction stage for ||v_k||.
Definition: iterative_operations.hpp:321
viennacl::OPENCL_MEMORY
Definition: forwards.h:349
viennacl::linalg::host_based::pipelined_cg_vector_update
void pipelined_cg_vector_update(vector_base< NumericT > &result, NumericT alpha, vector_base< NumericT > &p, vector_base< NumericT > &r, vector_base< NumericT > const &Ap, NumericT beta, vector_base< NumericT > &inner_prod_buffer)
Performs a joint vector update operation needed for an efficient pipelined CG algorithm.
Definition: iterative_operations.hpp:367
viennacl::linalg::cuda::pipelined_cg_prod
void pipelined_cg_prod(compressed_matrix< NumericT > const &A, vector_base< NumericT > const &p, vector_base< NumericT > &Ap, vector_base< NumericT > &inner_prod_buffer)
Definition: iterative_operations.hpp:280
viennacl::linalg::host_based::pipelined_gmres_gram_schmidt_stage2
void pipelined_gmres_gram_schmidt_stage2(vector_base< T > &device_krylov_basis, vcl_size_t v_k_size, vcl_size_t v_k_internal_size, vcl_size_t k, vector_base< T > const &vi_in_vk_buffer, vector_base< T > &R_buffer, vcl_size_t krylov_dim, vector_base< T > &inner_prod_buffer, vcl_size_t buffer_chunk_size)
Computes the second reduction stage for multiple inner products <v_i, v_k>, i=0..k-1, then updates v_k -= <v_i, v_k> v_i and computes the first reduction stage for ||v_k||.
Definition: iterative_operations.hpp:792
iterative_operations.hpp
Implementations of specialized kernels for fast iterative solvers using OpenMP on the CPU...
viennacl::linalg::opencl::pipelined_gmres_gram_schmidt_stage2
void pipelined_gmres_gram_schmidt_stage2(vector_base< T > &device_krylov_basis, vcl_size_t v_k_size, vcl_size_t v_k_internal_size, vcl_size_t param_k, vector_base< T > const &vi_in_vk_buffer, vector_base< T > &R_buffer, vcl_size_t krylov_dim, vector_base< T > &inner_prod_buffer, vcl_size_t buffer_chunk_size)
Definition: iterative_operations.hpp:667
forwards.h
This file provides the forward declarations for the main types used within ViennaCL.
viennacl::linalg::cuda::pipelined_gmres_normalize_vk
void pipelined_gmres_normalize_vk(vector_base< T > &v_k, vector_base< T > const &residual, vector_base< T > &R_buffer, vcl_size_t offset_in_R, vector_base< T > const &inner_prod_buffer, vector_base< T > &r_dot_vk_buffer, vcl_size_t buffer_chunk_size, vcl_size_t buffer_chunk_offset)
Performs a vector normalization needed for an efficient pipelined GMRES algorithm.
Definition: iterative_operations.hpp:1660
stride.hpp
Determines row and column increments for matrices and matrix proxies.
viennacl::linalg::pipelined_cg_vector_update
void pipelined_cg_vector_update(vector_base< NumericT > &result, NumericT alpha, vector_base< NumericT > &p, vector_base< NumericT > &r, vector_base< NumericT > const &Ap, NumericT beta, vector_base< NumericT > &inner_prod_buffer)
Performs a joint vector update operation needed for an efficient pipelined CG algorithm.
Definition: iterative_operations.hpp:59
viennacl::linalg::cuda::pipelined_bicgstab_prod
void pipelined_bicgstab_prod(compressed_matrix< NumericT > const &A, vector_base< NumericT > const &p, vector_base< NumericT > &Ap, vector_base< NumericT > const &r0star, vector_base< NumericT > &inner_prod_buffer, vcl_size_t buffer_chunk_size, vcl_size_t buffer_chunk_offset)
Definition: iterative_operations.hpp:1080
viennacl::linalg::opencl::pipelined_gmres_normalize_vk
void pipelined_gmres_normalize_vk(vector_base< T > &v_k, vector_base< T > const &residual, vector_base< T > &R_buffer, vcl_size_t offset_in_R, vector_base< T > const &inner_prod_buffer, vector_base< T > &r_dot_vk_buffer, vcl_size_t buffer_chunk_size, vcl_size_t buffer_chunk_offset)
Performs a vector normalization needed for an efficient pipelined GMRES algorithm.
Definition: iterative_operations.hpp:609
viennacl::linalg::host_based::pipelined_bicgstab_prod
void pipelined_bicgstab_prod(compressed_matrix< NumericT > const &A, vector_base< NumericT > const &p, vector_base< NumericT > &Ap, vector_base< NumericT > const &r0star, vector_base< NumericT > &inner_prod_buffer, vcl_size_t buffer_chunk_size, vcl_size_t buffer_chunk_offset)
Performs a fused matrix-vector product with a compressed_matrix for an efficient pipelined BiCGStab a...
Definition: iterative_operations.hpp:607
viennacl::linalg::pipelined_cg_prod
void pipelined_cg_prod(MatrixT const &A, vector_base< NumericT > const &p, vector_base< NumericT > &Ap, vector_base< NumericT > &inner_prod_buffer)
Performs a joint vector update operation needed for an efficient pipelined CG algorithm.
Definition: iterative_operations.hpp:97
NumericT
float NumericT
Definition: bisect.cpp:40
viennacl::linalg::host_based::pipelined_bicgstab_vector_update
void pipelined_bicgstab_vector_update(vector_base< NumericT > &result, NumericT alpha, vector_base< NumericT > &p, NumericT omega, vector_base< NumericT > const &s, vector_base< NumericT > &residual, vector_base< NumericT > const &As, NumericT beta, vector_base< NumericT > const &Ap, vector_base< NumericT > const &r0star, vector_base< NumericT > &inner_prod_buffer, vcl_size_t buffer_chunk_size)
Performs a joint vector update operation needed for an efficient pipelined BiCGStab algorithm...
Definition: iterative_operations.hpp:554
viennacl::linalg::cuda::pipelined_gmres_update_result
void pipelined_gmres_update_result(vector_base< T > &result, vector_base< T > const &residual, vector_base< T > const &krylov_basis, vcl_size_t v_k_size, vcl_size_t v_k_internal_size, vector_base< T > const &coefficients, vcl_size_t param_k)
Definition: iterative_operations.hpp:1882
viennacl::linalg::pipelined_bicgstab_vector_update
void pipelined_bicgstab_vector_update(vector_base< NumericT > &result, NumericT alpha, vector_base< NumericT > &p, NumericT omega, vector_base< NumericT > const &s, vector_base< NumericT > &residual, vector_base< NumericT > const &As, NumericT beta, vector_base< NumericT > const &Ap, vector_base< NumericT > const &r0star, vector_base< NumericT > &inner_prod_buffer, vcl_size_t buffer_chunk_size)
Performs a joint vector update operation needed for an efficient pipelined BiCGStab algorithm...
Definition: iterative_operations.hpp:171
iterative_operations.hpp
Implementations of operations using sparse matrices using CUDA.
viennacl::MEMORY_NOT_INITIALIZED
Definition: forwards.h:347
viennacl::CUDA_MEMORY
Definition: forwards.h:350
viennacl::linalg::pipelined_gmres_prod
void pipelined_gmres_prod(MatrixType const &A, vector_base< T > const &p, vector_base< T > &Ap, vector_base< T > &inner_prod_buffer)
Performs a joint vector update operation needed for an efficient pipelined GMRES algorithm.
Definition: iterative_operations.hpp:393
iterative_operations.hpp
Implementations of specialized kernels for fast iterative solvers using OpenCL.
viennacl::linalg::opencl::pipelined_cg_vector_update
void pipelined_cg_vector_update(vector_base< NumericT > &result, NumericT alpha, vector_base< NumericT > &p, vector_base< NumericT > &r, vector_base< NumericT > const &Ap, NumericT beta, vector_base< NumericT > &inner_prod_buffer)
Definition: iterative_operations.hpp:51
viennacl::linalg::opencl::pipelined_bicgstab_vector_update
void pipelined_bicgstab_vector_update(vector_base< NumericT > &result, NumericT alpha, vector_base< NumericT > &p, NumericT omega, vector_base< NumericT > const &s, vector_base< NumericT > &residual, vector_base< NumericT > const &As, NumericT beta, vector_base< NumericT > const &Ap, vector_base< NumericT > const &r0star, vector_base< NumericT > &inner_prod_buffer, vcl_size_t buffer_chunk_size)
Definition: iterative_operations.hpp:326
viennacl::linalg::pipelined_gmres_normalize_vk
void pipelined_gmres_normalize_vk(vector_base< T > &v_k, vector_base< T > const &residual, vector_base< T > &R_buffer, vcl_size_t offset_in_R, vector_base< T > const &inner_prod_buffer, vector_base< T > &r_dot_vk_buffer, vcl_size_t buffer_chunk_size, vcl_size_t buffer_chunk_offset)
Performs a vector normalization needed for an efficient pipelined GMRES algorithm.
Definition: iterative_operations.hpp:248
viennacl::vector_base< NumericT >
viennacl::vcl_size_t
std::size_t vcl_size_t
Definition: forwards.h:75
viennacl::linalg::cuda::pipelined_cg_vector_update
void pipelined_cg_vector_update(vector_base< NumericT > &result, NumericT alpha, vector_base< NumericT > &p, vector_base< NumericT > &r, vector_base< NumericT > const &Ap, NumericT beta, vector_base< NumericT > &inner_prod_buffer)
Definition: iterative_operations.hpp:85
viennacl::linalg::cuda::pipelined_bicgstab_vector_update
void pipelined_bicgstab_vector_update(vector_base< NumericT > &result, NumericT alpha, vector_base< NumericT > &p, NumericT omega, vector_base< NumericT > const &s, vector_base< NumericT > &residual, vector_base< NumericT > const &As, NumericT beta, vector_base< NumericT > const &Ap, vector_base< NumericT > const &r0star, vector_base< NumericT > &inner_prod_buffer, vcl_size_t buffer_chunk_size)
Definition: iterative_operations.hpp:864
viennacl::linalg::opencl::pipelined_gmres_gram_schmidt_stage1
void pipelined_gmres_gram_schmidt_stage1(vector_base< T > const &device_krylov_basis, vcl_size_t v_k_size, vcl_size_t v_k_internal_size, vcl_size_t param_k, vector_base< T > &vi_in_vk_buffer, vcl_size_t buffer_chunk_size)
Definition: iterative_operations.hpp:642
predicate.hpp
All the predicates used within ViennaCL. Checks for expressions to be vectors, etc.
viennacl::linalg::opencl::pipelined_bicgstab_update_s
void pipelined_bicgstab_update_s(vector_base< NumericT > &s, vector_base< NumericT > &r, vector_base< NumericT > const &Ap, vector_base< NumericT > &inner_prod_buffer, vcl_size_t buffer_chunk_size, vcl_size_t buffer_chunk_offset)
Definition: iterative_operations.hpp:295
viennacl::MAIN_MEMORY
Definition: forwards.h:348
viennacl::linalg::pipelined_bicgstab_prod
void pipelined_bicgstab_prod(MatrixT const &A, vector_base< NumericT > const &p, vector_base< NumericT > &Ap, vector_base< NumericT > const &r0star, vector_base< NumericT > &inner_prod_buffer, vcl_size_t buffer_chunk_size, vcl_size_t buffer_chunk_offset)
Performs a joint vector update operation needed for an efficient pipelined CG algorithm.
Definition: iterative_operations.hpp:208
viennacl::linalg::host_based::pipelined_cg_prod
void pipelined_cg_prod(compressed_matrix< NumericT > const &A, vector_base< NumericT > const &p, vector_base< NumericT > &Ap, vector_base< NumericT > &inner_prod_buffer)
Performs a fused matrix-vector product with a compressed_matrix for an efficient pipelined CG algorit...
Definition: iterative_operations.hpp:413
viennacl::linalg::cuda::pipelined_gmres_gram_schmidt_stage2
void pipelined_gmres_gram_schmidt_stage2(vector_base< T > &device_krylov_basis, vcl_size_t v_k_size, vcl_size_t v_k_internal_size, vcl_size_t param_k, vector_base< T > const &vi_in_vk_buffer, vector_base< T > &R_buffer, vcl_size_t krylov_dim, vector_base< T > &inner_prod_buffer, vcl_size_t buffer_chunk_size)
Definition: iterative_operations.hpp:1830
viennacl::linalg::host_based::pipelined_bicgstab_update_s
void pipelined_bicgstab_update_s(vector_base< NumericT > &s, vector_base< NumericT > &r, vector_base< NumericT > const &Ap, vector_base< NumericT > &inner_prod_buffer, vcl_size_t buffer_chunk_size, vcl_size_t buffer_chunk_offset)
Performs a joint vector update operation needed for an efficient pipelined BiCGStab algorithm...
Definition: iterative_operations.hpp:504
viennacl::linalg::cuda::pipelined_bicgstab_update_s
void pipelined_bicgstab_update_s(vector_base< NumericT > &s, vector_base< NumericT > &r, vector_base< NumericT > const &Ap, vector_base< NumericT > &inner_prod_buffer, vcl_size_t buffer_chunk_size, vcl_size_t buffer_chunk_offset)
Definition: iterative_operations.hpp:791
viennacl::linalg::host_based::pipelined_gmres_update_result
void pipelined_gmres_update_result(vector_base< T > &result, vector_base< T > const &residual, vector_base< T > const &krylov_basis, vcl_size_t v_k_size, vcl_size_t v_k_internal_size, vector_base< T > const &coefficients, vcl_size_t k)
Computes x += eta_0 r + sum_{i=1}^{k-1} eta_i v_{i-1}.
Definition: iterative_operations.hpp:836
range.hpp
Implementation of a range object for use with proxy objects.
viennacl::linalg::opencl::pipelined_cg_prod
void pipelined_cg_prod(compressed_matrix< NumericT > const &A, vector_base< NumericT > const &p, vector_base< NumericT > &Ap, vector_base< NumericT > &inner_prod_buffer)
Definition: iterative_operations.hpp:78
viennacl::linalg::pipelined_bicgstab_update_s
void pipelined_bicgstab_update_s(vector_base< NumericT > &s, vector_base< NumericT > &r, vector_base< NumericT > const &Ap, vector_base< NumericT > &inner_prod_buffer, vcl_size_t buffer_chunk_size, vcl_size_t buffer_chunk_offset)
Performs a joint vector update operation needed for an efficient pipelined CG algorithm.
Definition: iterative_operations.hpp:134
handle.hpp
Extracts the underlying OpenCL handle from a vector, a matrix, an expression etc. ...
viennacl::traits::handle
viennacl::backend::mem_handle & handle(T &obj)
Returns the generic memory handle of an object. Non-const version.
Definition: handle.hpp:41
viennacl::linalg::opencl::pipelined_gmres_update_result
void pipelined_gmres_update_result(vector_base< T > &result, vector_base< T > const &residual, vector_base< T > const &krylov_basis, vcl_size_t v_k_size, vcl_size_t v_k_internal_size, vector_base< T > const &coefficients, vcl_size_t param_k)
Definition: iterative_operations.hpp:699
viennacl::linalg::host_based::pipelined_gmres_gram_schmidt_stage1
void pipelined_gmres_gram_schmidt_stage1(vector_base< T > const &device_krylov_basis, vcl_size_t v_k_size, vcl_size_t v_k_internal_size, vcl_size_t k, vector_base< T > &vi_in_vk_buffer, vcl_size_t buffer_chunk_size)
Computes first reduction stage for multiple inner products <v_i, v_k>, i=0..k-1.
Definition: iterative_operations.hpp:760
scalar.hpp
Implementation of the ViennaCL scalar class.
viennacl::linalg::opencl::pipelined_gmres_prod
void pipelined_gmres_prod(compressed_matrix< T > const &A, vector_base< T > const &p, vector_base< T > &Ap, vector_base< T > &inner_prod_buffer)
Definition: iterative_operations.hpp:727
viennacl::linalg::cuda::pipelined_gmres_prod
void pipelined_gmres_prod(compressed_matrix< NumericT > const &A, vector_base< NumericT > const &p, vector_base< NumericT > &Ap, vector_base< NumericT > &inner_prod_buffer)
Definition: iterative_operations.hpp:1907
viennacl::linalg::host_based::pipelined_gmres_prod
void pipelined_gmres_prod(MatrixType const &A, vector_base< T > const &p, vector_base< T > &Ap, vector_base< T > &inner_prod_buffer)
Definition: iterative_operations.hpp:866
enable_if.hpp
Simple enable-if variant that uses the SFINAE pattern.