ViennaCL - The Vienna Computing Library  1.7.1
Free open-source GPU-accelerated linear algebra and solver library.
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Pages
iterative_operations.hpp
Go to the documentation of this file.
1 #ifndef VIENNACL_LINALG_ITERATIVE_OPERATIONS_HPP_
2 #define VIENNACL_LINALG_ITERATIVE_OPERATIONS_HPP_
3 
4 /* =========================================================================
5  Copyright (c) 2010-2016, Institute for Microelectronics,
6  Institute for Analysis and Scientific Computing,
7  TU Wien.
8  Portions of this software are copyright by UChicago Argonne, LLC.
9 
10  -----------------
11  ViennaCL - The Vienna Computing Library
12  -----------------
13 
14  Project Head: Karl Rupp rupp@iue.tuwien.ac.at
15 
16  (A list of authors and contributors can be found in the manual)
17 
18  License: MIT (X11), see file LICENSE in the base directory
19 ============================================================================= */
20 
25 #include "viennacl/forwards.h"
26 #include "viennacl/range.hpp"
27 #include "viennacl/scalar.hpp"
28 #include "viennacl/tools/tools.hpp"
29 #include "viennacl/meta/predicate.hpp"
30 #include "viennacl/meta/enable_if.hpp"
31 #include "viennacl/traits/size.hpp"
32 #include "viennacl/traits/start.hpp"
33 #include "viennacl/traits/handle.hpp"
34 #include "viennacl/traits/stride.hpp"
35 #include "viennacl/linalg/host_based/iterative_operations.hpp"
36 
37 #ifdef VIENNACL_WITH_OPENCL
38  #include "viennacl/linalg/opencl/iterative_operations.hpp"
39 #endif
40 
41 #ifdef VIENNACL_WITH_CUDA
42  #include "viennacl/linalg/cuda/iterative_operations.hpp"
43 #endif
44 
45 namespace viennacl
46 {
47 namespace linalg
48 {
49 
58 template<typename NumericT>
59 void pipelined_cg_vector_update(vector_base<NumericT> & result,
60  NumericT alpha,
61  vector_base<NumericT> & p,
62  vector_base<NumericT> & r,
63  vector_base<NumericT> const & Ap,
64  NumericT beta,
65  vector_base<NumericT> & inner_prod_buffer)
66 {
67  switch (viennacl::traits::handle(result).get_active_handle_id())
68  {
69  case viennacl::MAIN_MEMORY:
70  viennacl::linalg::host_based::pipelined_cg_vector_update(result, alpha, p, r, Ap, beta, inner_prod_buffer);
71  break;
72 #ifdef VIENNACL_WITH_OPENCL
73  case viennacl::OPENCL_MEMORY:
74  viennacl::linalg::opencl::pipelined_cg_vector_update(result, alpha, p, r, Ap, beta, inner_prod_buffer);
75  break;
76 #endif
77 #ifdef VIENNACL_WITH_CUDA
78  case viennacl::CUDA_MEMORY:
79  viennacl::linalg::cuda::pipelined_cg_vector_update(result, alpha, p, r, Ap, beta, inner_prod_buffer);
80  break;
81 #endif
82  case viennacl::MEMORY_NOT_INITIALIZED:
83  throw memory_exception("not initialised!");
84  default:
85  throw memory_exception("not implemented");
86  }
87 }
88 
89 
96 template<typename MatrixT, typename NumericT>
97 void pipelined_cg_prod(MatrixT const & A,
98  vector_base<NumericT> const & p,
99  vector_base<NumericT> & Ap,
100  vector_base<NumericT> & inner_prod_buffer)
101 {
102  switch (viennacl::traits::handle(p).get_active_handle_id())
103  {
104  case viennacl::MAIN_MEMORY:
105  viennacl::linalg::host_based::pipelined_cg_prod(A, p, Ap, inner_prod_buffer);
106  break;
107 #ifdef VIENNACL_WITH_OPENCL
108  case viennacl::OPENCL_MEMORY:
109  viennacl::linalg::opencl::pipelined_cg_prod(A, p, Ap, inner_prod_buffer);
110  break;
111 #endif
112 #ifdef VIENNACL_WITH_CUDA
113  case viennacl::CUDA_MEMORY:
114  viennacl::linalg::cuda::pipelined_cg_prod(A, p, Ap, inner_prod_buffer);
115  break;
116 #endif
117  case viennacl::MEMORY_NOT_INITIALIZED:
118  throw memory_exception("not initialised!");
119  default:
120  throw memory_exception("not implemented");
121  }
122 }
123 
125 
133 template<typename NumericT>
134 void pipelined_bicgstab_update_s(vector_base<NumericT> & s,
135  vector_base<NumericT> & r,
136  vector_base<NumericT> const & Ap,
137  vector_base<NumericT> & inner_prod_buffer,
138  vcl_size_t buffer_chunk_size,
139  vcl_size_t buffer_chunk_offset)
140 {
141  switch (viennacl::traits::handle(s).get_active_handle_id())
142  {
143  case viennacl::MAIN_MEMORY:
144  viennacl::linalg::host_based::pipelined_bicgstab_update_s(s, r, Ap, inner_prod_buffer, buffer_chunk_size, buffer_chunk_offset);
145  break;
146 #ifdef VIENNACL_WITH_OPENCL
147  case viennacl::OPENCL_MEMORY:
148  viennacl::linalg::opencl::pipelined_bicgstab_update_s(s, r, Ap, inner_prod_buffer, buffer_chunk_size, buffer_chunk_offset);
149  break;
150 #endif
151 #ifdef VIENNACL_WITH_CUDA
152  case viennacl::CUDA_MEMORY:
153  viennacl::linalg::cuda::pipelined_bicgstab_update_s(s, r, Ap, inner_prod_buffer, buffer_chunk_size, buffer_chunk_offset);
154  break;
155 #endif
156  case viennacl::MEMORY_NOT_INITIALIZED:
157  throw memory_exception("not initialised!");
158  default:
159  throw memory_exception("not implemented");
160  }
161 }
162 
170 template<typename NumericT>
171 void pipelined_bicgstab_vector_update(vector_base<NumericT> & result, NumericT alpha, vector_base<NumericT> & p, NumericT omega, vector_base<NumericT> const & s,
172  vector_base<NumericT> & residual, vector_base<NumericT> const & As,
173  NumericT beta, vector_base<NumericT> const & Ap,
174  vector_base<NumericT> const & r0star,
175  vector_base<NumericT> & inner_prod_buffer,
176  vcl_size_t buffer_chunk_size)
177 {
178  switch (viennacl::traits::handle(s).get_active_handle_id())
179  {
180  case viennacl::MAIN_MEMORY:
181  viennacl::linalg::host_based::pipelined_bicgstab_vector_update(result, alpha, p, omega, s, residual, As, beta, Ap, r0star, inner_prod_buffer, buffer_chunk_size);
182  break;
183  #ifdef VIENNACL_WITH_OPENCL
184  case viennacl::OPENCL_MEMORY:
185  viennacl::linalg::opencl::pipelined_bicgstab_vector_update(result, alpha, p, omega, s, residual, As, beta, Ap, r0star, inner_prod_buffer, buffer_chunk_size);
186  break;
187  #endif
188  #ifdef VIENNACL_WITH_CUDA
189  case viennacl::CUDA_MEMORY:
190  viennacl::linalg::cuda::pipelined_bicgstab_vector_update(result, alpha, p, omega, s, residual, As, beta, Ap, r0star, inner_prod_buffer, buffer_chunk_size);
191  break;
192  #endif
193  case viennacl::MEMORY_NOT_INITIALIZED:
194  throw memory_exception("not initialised!");
195  default:
196  throw memory_exception("not implemented");
197  }
198 }
199 
200 
207 template<typename MatrixT, typename NumericT>
208 void pipelined_bicgstab_prod(MatrixT const & A,
209  vector_base<NumericT> const & p,
210  vector_base<NumericT> & Ap,
211  vector_base<NumericT> const & r0star,
212  vector_base<NumericT> & inner_prod_buffer,
213  vcl_size_t buffer_chunk_size,
214  vcl_size_t buffer_chunk_offset)
215 {
216  switch (viennacl::traits::handle(p).get_active_handle_id())
217  {
218  case viennacl::MAIN_MEMORY:
219  viennacl::linalg::host_based::pipelined_bicgstab_prod(A, p, Ap, r0star, inner_prod_buffer, buffer_chunk_size, buffer_chunk_offset);
220  break;
221 #ifdef VIENNACL_WITH_OPENCL
222  case viennacl::OPENCL_MEMORY:
223  viennacl::linalg::opencl::pipelined_bicgstab_prod(A, p, Ap, r0star, inner_prod_buffer, buffer_chunk_size, buffer_chunk_offset);
224  break;
225 #endif
226 #ifdef VIENNACL_WITH_CUDA
227  case viennacl::CUDA_MEMORY:
228  viennacl::linalg::cuda::pipelined_bicgstab_prod(A, p, Ap, r0star, inner_prod_buffer, buffer_chunk_size, buffer_chunk_offset);
229  break;
230 #endif
231  case viennacl::MEMORY_NOT_INITIALIZED:
232  throw memory_exception("not initialised!");
233  default:
234  throw memory_exception("not implemented");
235  }
236 }
237 
239 
247 template <typename T>
248 void pipelined_gmres_normalize_vk(vector_base<T> & v_k,
249  vector_base<T> const & residual,
250  vector_base<T> & R_buffer,
251  vcl_size_t offset_in_R,
252  vector_base<T> const & inner_prod_buffer,
253  vector_base<T> & r_dot_vk_buffer,
254  vcl_size_t buffer_chunk_size,
255  vcl_size_t buffer_chunk_offset)
256 {
257  switch (viennacl::traits::handle(v_k).get_active_handle_id())
258  {
259  case viennacl::MAIN_MEMORY:
260  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);
261  break;
262 #ifdef VIENNACL_WITH_OPENCL
263  case viennacl::OPENCL_MEMORY:
264  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);
265  break;
266 #endif
267 #ifdef VIENNACL_WITH_CUDA
268  case viennacl::CUDA_MEMORY:
269  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);
270  break;
271 #endif
272  case viennacl::MEMORY_NOT_INITIALIZED:
273  throw memory_exception("not initialised!");
274  default:
275  throw memory_exception("not implemented");
276  }
277 }
278 
279 
280 
285 template <typename T>
286 void pipelined_gmres_gram_schmidt_stage1(vector_base<T> const & device_krylov_basis,
287  vcl_size_t v_k_size,
288  vcl_size_t v_k_internal_size,
289  vcl_size_t k,
290  vector_base<T> & vi_in_vk_buffer,
291  vcl_size_t buffer_chunk_size)
292 {
293  switch (viennacl::traits::handle(device_krylov_basis).get_active_handle_id())
294  {
295  case viennacl::MAIN_MEMORY:
296  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);
297  break;
298 #ifdef VIENNACL_WITH_OPENCL
299  case viennacl::OPENCL_MEMORY:
300  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);
301  break;
302 #endif
303 #ifdef VIENNACL_WITH_CUDA
304  case viennacl::CUDA_MEMORY:
305  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);
306  break;
307 #endif
308  case viennacl::MEMORY_NOT_INITIALIZED:
309  throw memory_exception("not initialised!");
310  default:
311  throw memory_exception("not implemented");
312  }
313 }
314 
315 
320 template <typename T>
321 void pipelined_gmres_gram_schmidt_stage2(vector_base<T> & device_krylov_basis,
322  vcl_size_t v_k_size,
323  vcl_size_t v_k_internal_size,
324  vcl_size_t k,
325  vector_base<T> const & vi_in_vk_buffer,
326  vector_base<T> & R_buffer,
327  vcl_size_t krylov_dim,
328  vector_base<T> & inner_prod_buffer,
329  vcl_size_t buffer_chunk_size)
330 {
331  switch (viennacl::traits::handle(device_krylov_basis).get_active_handle_id())
332  {
333  case viennacl::MAIN_MEMORY:
334  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);
335  break;
336 #ifdef VIENNACL_WITH_OPENCL
337  case viennacl::OPENCL_MEMORY:
338  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);
339  break;
340 #endif
341 #ifdef VIENNACL_WITH_CUDA
342  case viennacl::CUDA_MEMORY:
343  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);
344  break;
345 #endif
346  case viennacl::MEMORY_NOT_INITIALIZED:
347  throw memory_exception("not initialised!");
348  default:
349  throw memory_exception("not implemented");
350  }
351 }
352 
353 
355 template <typename T>
356 void pipelined_gmres_update_result(vector_base<T> & result,
357  vector_base<T> const & residual,
358  vector_base<T> const & krylov_basis,
359  vcl_size_t v_k_size,
360  vcl_size_t v_k_internal_size,
361  vector_base<T> const & coefficients,
362  vcl_size_t k)
363 {
364  switch (viennacl::traits::handle(result).get_active_handle_id())
365  {
366  case viennacl::MAIN_MEMORY:
367  viennacl::linalg::host_based::pipelined_gmres_update_result(result, residual, krylov_basis, v_k_size, v_k_internal_size, coefficients, k);
368  break;
369 #ifdef VIENNACL_WITH_OPENCL
370  case viennacl::OPENCL_MEMORY:
371  viennacl::linalg::opencl::pipelined_gmres_update_result(result, residual, krylov_basis, v_k_size, v_k_internal_size, coefficients, k);
372  break;
373 #endif
374 #ifdef VIENNACL_WITH_CUDA
375  case viennacl::CUDA_MEMORY:
376  viennacl::linalg::cuda::pipelined_gmres_update_result(result, residual, krylov_basis, v_k_size, v_k_internal_size, coefficients, k);
377  break;
378 #endif
379  case viennacl::MEMORY_NOT_INITIALIZED:
380  throw memory_exception("not initialised!");
381  default:
382  throw memory_exception("not implemented");
383  }
384 }
385 
392 template <typename MatrixType, typename T>
393 void pipelined_gmres_prod(MatrixType const & A,
394  vector_base<T> const & p,
395  vector_base<T> & Ap,
396  vector_base<T> & inner_prod_buffer)
397 {
398  switch (viennacl::traits::handle(p).get_active_handle_id())
399  {
400  case viennacl::MAIN_MEMORY:
401  viennacl::linalg::host_based::pipelined_gmres_prod(A, p, Ap, inner_prod_buffer);
402  break;
403 #ifdef VIENNACL_WITH_OPENCL
404  case viennacl::OPENCL_MEMORY:
405  viennacl::linalg::opencl::pipelined_gmres_prod(A, p, Ap, inner_prod_buffer);
406  break;
407 #endif
408 #ifdef VIENNACL_WITH_CUDA
409  case viennacl::CUDA_MEMORY:
410  viennacl::linalg::cuda::pipelined_gmres_prod(A, p, Ap, inner_prod_buffer);
411  break;
412 #endif
413  case viennacl::MEMORY_NOT_INITIALIZED:
414  throw memory_exception("not initialised!");
415  default:
416  throw memory_exception("not implemented");
417  }
418 }
419 
420 
421 } //namespace linalg
422 } //namespace viennacl
423 
424 
425 #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 , 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 , i=0..k-1, then updates v_k -= v_i and computes the first reduction stage for ||v_k||.
Definition: iterative_operations.hpp:321
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 , i=0..k-1, then updates 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::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::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 , 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.