<!DOCTYPE html
PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml" lang="en-us" xml:lang="en-us">
<head>
<meta http-equiv="Content-Type" content="text/html; charset=utf-8"></meta>
<meta http-equiv="X-UA-Compatible" content="IE=edge"></meta>
<meta name="copyright" content="(C) Copyright 2005"></meta>
<meta name="DC.rights.owner" content="(C) Copyright 2005"></meta>
<meta name="DC.Type" content="concept"></meta>
<meta name="DC.Title" content="Incomplete-LU and Cholesky Preconditioned Iterative Methods Using CUSPARSE and CUBLAS"></meta>
<meta name="abstract" content="White paper describing how to use the CUSPARSE and CUBLAS libraries to achieve a 2x speedup over CPU in the incomplete-LU and Cholesky preconditioned iterative methods."></meta>
<meta name="description" content="White paper describing how to use the CUSPARSE and CUBLAS libraries to achieve a 2x speedup over CPU in the incomplete-LU and Cholesky preconditioned iterative methods."></meta>
<meta name="DC.Coverage" content="White Papers"></meta>
<meta name="DC.subject" content="CUSPARSE and CUBLAS, CUSPARSE and CUBLAS speedup, CUSPARSE and CUBLAS CG method, CUSPARSE and CUBLAS libraries, CUSPARSE and CUBLAS sparse triangle solve, CUSPARSE and CUBLAS BiCGStab method, CUSPARSE and CUBLAS implementation comparison, CUSPARSE and CUBLAS performance, CUSPARSE and CUBLAS right-hand sides, CUSPARSE and CUBLAS sparsity pattern"></meta>
<meta name="keywords" content="CUSPARSE and CUBLAS, CUSPARSE and CUBLAS speedup, CUSPARSE and CUBLAS CG method, CUSPARSE and CUBLAS libraries, CUSPARSE and CUBLAS sparse triangle solve, CUSPARSE and CUBLAS BiCGStab method, CUSPARSE and CUBLAS implementation comparison, CUSPARSE and CUBLAS performance, CUSPARSE and CUBLAS right-hand sides, CUSPARSE and CUBLAS sparsity pattern"></meta>
<meta name="DC.Format" content="XHTML"></meta>
<meta name="DC.Identifier" content="abstract"></meta>
<link rel="stylesheet" type="text/css" href="../common/formatting/commonltr.css"></link>
<link rel="stylesheet" type="text/css" href="../common/formatting/site.css"></link>
<title>Incomplete-LU and Cholesky Preconditioned Iterative Methods :: CUDA Toolkit Documentation</title>
<!--[if lt IE 9]>
<script src="../common/formatting/html5shiv-printshiv.min.js"></script>
<![endif]-->
<script type="text/javascript" charset="utf-8" src="../common/formatting/jquery.min.js"></script>
<script type="text/javascript" charset="utf-8" src="../common/formatting/jquery.ba-hashchange.min.js"></script>
<link rel="canonical" href="http://docs.nvidia.com/cuda/incomplete-lu-cholesky/index.html"></link>
<link rel="stylesheet" type="text/css" href="../common/formatting/qwcode.highlight.css"></link>
</head>
<body>
<article id="contents">
<div id="eqn-warning">This document includes math equations
(highlighted in red) which are best viewed with <a target="_blank" href="https://www.mozilla.org/firefox">Firefox</a> version 4.0
or higher, or another <a target="_blank" href="http://www.w3.org/Math/Software/mathml_software_cat_browsers.html">MathML-aware
browser</a>. There is also a <a href="../../pdf/Incomplete_LU_Cholesky.pdf">PDF version of this
document</a>.
</div>
<div id="eqn-warning-buf"></div>
<div id="release-info">Incomplete-LU and Cholesky Preconditioned Iterative Methods
(<a href="../../pdf/Incomplete_LU_Cholesky.pdf">PDF</a>)
-
CUDA Toolkit v5.5
(<a href="https://developer.nvidia.com/cuda-toolkit-archive">older</a>)
-
Last updated
July 19, 2013
-
<a href="mailto:cudatools@nvidia.com?subject=CUDA Tools Documentation Feedback: incomplete-lu-cholesky">Send Feedback</a></div>
<div class="topic nested0" id="abstract"><a name="abstract" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#abstract" name="abstract" shape="rect">Incomplete-LU and Cholesky Preconditioned Iterative Methods Using
CUSPARSE and CUBLAS</a></h2>
<div class="body conbody"></div>
</div>
<div class="topic concept nested0" id="introduction"><a name="introduction" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#introduction" name="introduction" shape="rect">1. Introduction</a></h2>
<div class="body conbody">
<p class="p">The solution of large sparse linear systems is an important problem in
computational mechanics, atmospheric modeling, geophysics, biology,
circuit simulation and many other applications in the field of
computational science and engineering. In general, these linear systems
can be solved using direct or preconditioned iterative methods. Although
the direct methods are often more reliable, they usually have large
memory requirements and do not scale well on massively parallel computer
platforms.
</p>
<p class="p">The iterative methods are more amenable to parallelism and therefore can
be used to solve larger problems. Currently, the most popular iterative
schemes belong to the Krylov subspace family of methods. They include
<dfn class="term">Bi-Conjugate Gradient Stabilized</dfn> (BiCGStab) and
<dfn class="term">Conjugate Gradient</dfn> (CG) iterative methods for nonsymmetric
and <dfn class="term">symmetric positive definite</dfn> (s.p.d.) linear systems,
respectively <a class="xref" href="index.html#references__2" shape="rect">[2]</a>, <a class="xref" href="index.html#references__11" shape="rect">[11]</a>. We describe these methods in more
detail in the next section.
</p>
<p class="p">In practice, we often use a variety of preconditioning techniques to
improve the convergence of the iterative methods. In this white paper we
focus on the incomplete-LU and Cholesky preconditioning <a class="xref" href="index.html#references__11" shape="rect">[11]</a>, which is one of the most popular of
these preconditioning techniques. It computes an incomplete factorization
of the coefficient matrix and requires a solution of lower and upper
triangular linear systems in every iteration of the iterative method.
</p>
<p class="p">In order to implement the preconditioned BiCGStab and CG we use the
sparse matrix-vector multiplication <a class="xref" href="index.html#references__3" shape="rect">[3]</a>, <a class="xref" href="index.html#references__15" shape="rect">[15]</a> and the sparse triangular solve
<a class="xref" href="index.html#references__8" shape="rect">[8]</a>, <a class="xref" href="index.html#references__16" shape="rect">[16]</a> implemented in the CUSPARSE library.
We point out that the underlying implementation of these algorithms takes
advantage of the CUDA parallel programming paradigm <a class="xref" href="index.html#references__5" shape="rect">[5]</a>, <a class="xref" href="index.html#references__9" shape="rect">[9]</a>, <a class="xref" href="index.html#references__13" shape="rect">[13]</a>, which allows us to explore the
computational resources of the graphical processing unit (GPU). In our
numerical experiments the incomplete factorization is performed on the
CPU (host) and the resulting lower and upper triangular factors are then
transferred to the GPU (device) memory before starting the iterative
method. However, the computation of the incomplete factorization could
also be accelerated on the GPU.
</p>
<p class="p">We point out that the parallelism available in these iterative methods
depends highly on the sparsity pattern of the coefficient matrix at hand.
In our numerical experiments the incomplete-LU and Cholesky
preconditioned iterative methods achieve on average more than 2x speedup
using the CUSPARSE and CUBLAS libraries on the GPU over the MKL <a class="xref" href="index.html#references__17" shape="rect">[17]</a> implementation on the CPU. For
example, the speedup for the preconditioned iterative methods with the
incomplete-LU and Cholesky factorization with 0 fill-in (ilu0) is shown
in <a class="xref" href="index.html#introduction__speedup-of-incomplete-lu-cholesky-with-0-fill-in" shape="rect">Figure 1</a> for matrices resulting from a variety of applications.
It will be described in more detail in the last section.
</p>
<div class="fig fignone" id="introduction__speedup-of-incomplete-lu-cholesky-with-0-fill-in"><a name="introduction__speedup-of-incomplete-lu-cholesky-with-0-fill-in" shape="rect">
<!-- --></a><span class="figcap">Figure 1. Speedup of the Incomplete-LU Cholesky (with 0 fill-in) Prec.
Iterative Methods</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/speedup-of-incomplete-lu-cholesky.png" alt="Figure of the speedup of the incomplete-LU Cholesky (with 0 fill-in) Prec. Iterative Methods."></img></div><br clear="none"></br></div>
<p class="p">In the next sections we briefly describe the methods of interest and
comment on the role played in them by the parallel sparse matrix-vector
multiplication and triangular solve algorithms.
</p>
</div>
</div>
<div class="topic concept nested0" id="preconditioned-iterative-methods"><a name="preconditioned-iterative-methods" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#preconditioned-iterative-methods" name="preconditioned-iterative-methods" shape="rect">2. Preconditioned Iterative Methods</a></h2>
<div class="body conbody">
<div class="section">
<p class="p">Let us consider the linear system</p>
<div class="tablenoborder"><a name="preconditioned-iterative-methods__eq-1" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="preconditioned-iterative-methods__eq-1" class="table" frame="void" border="0" rules="none">
<tbody class="tbody">
<tr class="row">
<td class="entry" align="center" valign="top" width="95.23809523809523%" rowspan="1" colspan="1">
<p class="p d4p_eqn_block">
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mi>A</mi>
<mo></mo>
<mi mathvariant="bold">x</mi>
<mo>=</mo>
<mi mathvariant="bold">f</mi>
</mrow>
</math>
</p>
</td>
<td class="entry" align="right" valign="top" width="4.761904761904762%" rowspan="1" colspan="1">
<p class="p">(1)</p>
</td>
</tr>
</tbody>
</table>
</div>
<p class="p">where
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mi>A</mi>
<mo>∈</mo>
<msup>
<mo>ℝ</mo>
<mrow>
<mi>n</mi>
<mo>×</mo>
<mi>n</mi>
</mrow>
</msup>
</mrow>
</math> is a nonsingular
coefficient matrix and
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mi mathvariant="bold">x</mi>
<mo>,</mo>
<mi mathvariant="bold">f</mi>
<mo>∈</mo>
<msup>
<mo>ℝ</mo>
<mi>n</mi>
</msup>
</mrow>
</math> are
the solution and right-hand-side vectors.
</p>
<p class="p">In general, the iterative methods start with an initial guess and
perform a series of steps that find more accurate approximations to the
solution. There are two types of iterative methods: (i) the stationary
iterative methods, such as the splitting-based <dfn class="term">Jacobi</dfn> and
<dfn class="term">Gauss-Seidel</dfn> (GS), and (ii) the nonstationary iterative
methods, such as the <dfn class="term">Krylov</dfn> subspace family of methods,
which includes <dfn class="term">CG</dfn> and <dfn class="term">BiCGStab</dfn>. As we
mentioned earlier we focus on the latter in this white paper.
</p>
<p class="p">The convergence of the iterative methods depends highly on the
spectrum of the coefficient matrix and can be significantly improved
using preconditioning. The preconditioning modifies the spectrum of the
coefficient matrix of the linear system in order to reduce the number
of iterative steps required for convergence. It often involves finding
a preconditioning matrix
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mi>M</mi>
</mrow>
</math>,
such that
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<msup>
<mi>M</mi>
<mrow>
<mo>−</mo>
<mn>1</mn>
</mrow>
</msup>
</mrow>
</math> is a good
approximation of
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<msup>
<mi>A</mi>
<mrow>
<mo>−</mo>
<mn>1</mn>
</mrow>
</msup>
</mrow>
</math> and the
systems with
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mi>M</mi>
</mrow>
</math> are relatively
easy to solve.
</p>
<p class="p">For the s.p.d. matrix
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mi>A</mi>
</mrow>
</math> we can
let
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mi>M</mi>
</mrow>
</math> be its
incomplete-Cholesky factorization, so that
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mi>A</mi>
<mo>≈</mo>
<mi>M</mi>
<mo>=</mo>
<msup>
<mover accent="true">
<mi>R</mi>
<mo>˜</mo>
</mover>
<mi>T</mi>
</msup>
<mover accent="true">
<mi>R</mi>
<mo>˜</mo>
</mover>
</mrow>
</math>, where
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mover>
<mi>R</mi>
<mo>˜</mo>
</mover>
</mrow>
</math> is an upper triangular matrix. Let
us assume that
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mi>M</mi>
</mrow>
</math> is
nonsingular, then
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<msup>
<mover accent="true">
<mi>R</mi>
<mo>˜</mo>
</mover>
<mrow>
<mo>−</mo>
<mi>T</mi>
</mrow>
</msup>
<mi>A</mi>
<msup>
<mover accent="true">
<mi>R</mi>
<mo>˜</mo>
</mover>
<mrow>
<mo>−</mo>
<mn>1</mn>
</mrow>
</msup>
</mrow>
</math> is s.p.d.
and instead of solving the linear system <a class="xref" href="index.html#preconditioned-iterative-methods__eq-1" shape="rect">(1)</a>, we can
solve the preconditioned linear system
</p>
<div class="tablenoborder"><a name="preconditioned-iterative-methods__eq-2" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="preconditioned-iterative-methods__eq-2" class="table" frame="void" border="0" rules="none">
<tbody class="tbody">
<tr class="row">
<td class="entry" align="center" valign="top" width="95.23809523809523%" rowspan="1" colspan="1">
<p class="p d4p_eqn_block">
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mrow>
<mo>(</mo>
<mrow>
<msup>
<mover accent="true">
<mi>R</mi>
<mo>˜</mo>
</mover>
<mrow>
<mo>−</mo>
<mi>T</mi>
</mrow>
</msup>
<mo></mo>
<mi>A</mi>
<mo></mo>
<msup>
<mover accent="true">
<mi>R</mi>
<mo>˜</mo>
</mover>
<mrow>
<mo>−</mo>
<mn>1</mn>
</mrow>
</msup>
</mrow>
<mo>)</mo>
</mrow>
<mo></mo>
<mrow>
<mo>(</mo>
<mrow>
<mover accent="true">
<mi>R</mi>
<mo>˜</mo>
</mover>
<mo></mo>
<mi mathvariant="bold">x</mi>
</mrow>
<mo>)</mo>
</mrow>
<mo>=</mo>
<msup>
<mover accent="true">
<mi>R</mi>
<mo>˜</mo>
</mover>
<mrow>
<mo>−</mo>
<mi>T</mi>
</mrow>
</msup>
<mo></mo>
<mi mathvariant="bold">f</mi>
</mrow>
</math>
</p>
</td>
<td class="entry" align="right" valign="top" width="4.761904761904762%" rowspan="1" colspan="1">
<p class="p">(2)</p>
</td>
</tr>
</tbody>
</table>
</div>
<p class="p">The pseudocode for the preconditioned CG iterative method is shown in
<a class="xref" href="index.html#preconditioned-iterative-methods__algorithm-1-conjugate-gradient-cg" shape="rect">Algorithm 1</a>.
</p>
</div>
<div class="example" id="preconditioned-iterative-methods__algorithm-1-conjugate-gradient-cg"><a name="preconditioned-iterative-methods__algorithm-1-conjugate-gradient-cg" shape="rect">
<!-- --></a><h2 class="title sectiontitle">Algorithm 1 Conjugate Gradient (CG)</h2>
<p class="p d4p_eqn_block">
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mtable displaystyle="true" equalrows="true" columnalign="right left">
<mtr>
<mtd>
<mi mathvariant="monospace"> </mi>
<mn mathvariant="monospace">1</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext>Letting initial guess be </mtext>
<msub>
<mi mathvariant="bold">x</mi>
<mn>0</mn>
</msub>
<mtext>, compute </mtext>
<mi mathvariant="bold">r</mi>
<mo>←</mo>
<mi mathvariant="bold">f</mi>
<mo>−</mo>
<mi>A</mi>
<mo></mo>
<msub>
<mi mathvariant="bold">x</mi>
<mn>0</mn>
</msub>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mtext> </mtext>
</mtd>
</mtr>
<mtr>
<mtd>
<mtext> </mtext>
</mtd>
</mtr>
</mtable>
</math>
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mtable displaystyle="true" equalrows="true" columnalign="right left left">
<mtr>
<mtd>
<mn mathvariant="monospace">2</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext mathvariant="bold">for </mtext>
<mi>i</mi>
<mo>←</mo>
<mn>1</mn>
<mo>,</mo>
<mn>2</mn>
<mo>,</mo>
<mn>...</mn>
<mtext> until convergence </mtext>
<mtext mathvariant="bold">do</mtext>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">3</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mtext>Solve </mtext>
<mi>M</mi>
<mo></mo>
<mi mathvariant="bold">z</mi>
<mo>←</mo>
<mi mathvariant="bold">r</mi>
</mrow>
</mtd>
<mtd>
<mo>⊳</mo>
<mtext>Sparse lower and upper triangular solves</mtext>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">4</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<msub>
<mi>ρ</mi>
<mi>i</mi>
</msub>
<mo>←</mo>
<msup>
<mi mathvariant="bold">r</mi>
<mi>T</mi>
</msup>
<mo></mo>
<mi mathvariant="bold">z</mi>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">5</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mtext mathvariant="bold">if </mtext>
<mi>i</mi>
<mo>==</mo>
<mn>1</mn>
<mtext mathvariant="bold"> then</mtext>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">6</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mi mathvariant="bold">p</mi>
<mo>←</mo>
<mi mathvariant="bold">z</mi>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">7</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mtext mathvariant="bold">else</mtext>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">8</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mi>β</mi>
<mo>←</mo>
<mfrac>
<mrow>
<msub>
<mi>ρ</mi>
<mi>i</mi>
</msub>
</mrow>
<mrow>
<msub>
<mi>ρ</mi>
<mrow>
<mi>i</mi>
<mo>−</mo>
<mn>1</mn>
</mrow>
</msub>
</mrow>
</mfrac>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">9</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mi mathvariant="bold">p</mi>
<mo>←</mo>
<mi mathvariant="bold">z</mi>
<mo>+</mo>
<mi>β</mi>
<mo></mo>
<mi mathvariant="bold">p</mi>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">10</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mtext mathvariant="bold">end if</mtext>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">11</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mtext>Compute </mtext>
<mi mathvariant="bold">q</mi>
<mo>←</mo>
<mi>A</mi>
<mo></mo>
<mi mathvariant="bold">p</mi>
</mrow>
</mtd>
<mtd>
<mo>⊳</mo>
<mtext>Sparse matrix-vector multiplication</mtext>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">12</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mi>α</mi>
<mo>←</mo>
<mfrac>
<mrow>
<msub>
<mi>ρ</mi>
<mi>i</mi>
</msub>
</mrow>
<mrow>
<msup>
<mi mathvariant="bold">p</mi>
<mi>T</mi>
</msup>
<mo></mo>
<mi mathvariant="bold">q</mi>
</mrow>
</mfrac>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">13</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mi mathvariant="bold">x</mi>
<mo>←</mo>
<mi mathvariant="bold">x</mi>
<mo>+</mo>
<mi>α</mi>
<mo></mo>
<mi mathvariant="bold">p</mi>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">14</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mi mathvariant="bold">r</mi>
<mo>←</mo>
<mi mathvariant="bold">r</mi>
<mo>−</mo>
<mi>α</mi>
<mo></mo>
<mi mathvariant="bold">q</mi>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">15</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext mathvariant="bold">end for</mtext>
</mrow>
</mtd>
</mtr>
</mtable>
</math>
</p>
</div>
<div class="section">
<p class="p">Notice that in every iteration of the incomplete-Cholesky
preconditioned CG iterative method we need to perform one sparse
matrix-vector multiplication and two triangular solves. The
corresponding CG code using the CUSPARSE and CUBLAS libraries in C
programming language is shown below.
</p><pre xml:space="preserve"><span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-doccomment">/***** CG Code *****/</span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">/* ASSUMPTIONS:
1. The CUSPARSE and CUBLAS libraries have been initialized.
2. The appropriate memory has been allocated and set to zero.
3. The matrix A (valA, csrRowPtrA, csrColIndA) and the incomplete-
Cholesky upper triangular factor R (valR, csrRowPtrR, csrColIndR)
have been computed and are present in the device (GPU) memory. */</span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//create the info and analyse the lower and upper triangular factors</span>
cusparseCreateSolveAnalysisInfo(&inforRt);
cusparseCreateSolveAnalysisInfo(&inforR);
cusparseDcsrsv_analysis(handle,CUSPARSE_OPERATION_TRANSPOSE,
n, descrR, valR, csrRowPtrR, csrColIndR, inforRt);
cusparseDcsrsv_analysis(handle,CUSPARSE_OPERATION_NON_TRANSPOSE,
n, descrR, valR, csrRowPtrR, csrColIndR, inforR);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//1: compute initial residual r = f - A x0 (using initial guess in x)</span>
cusparseDcsrmv(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, n, n, 1.0,
descrA, valA, csrRowPtrA, csrColIndA, x, 0.0, r);
cublasDscal(n,-1.0, r, 1);
cublasDaxpy(n, 1.0, f, 1, r, 1);
nrmr0 = cublasDnrm2(n, r, 1);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//2: repeat until convergence (based on max. it. and relative residual)</span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">for</span> (i=0; i<maxit; i++){
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//3: Solve M z = r (sparse lower and upper triangular solves)</span>
cusparseDcsrsv_solve(handle, CUSPARSE_OPERATION_TRANSPOSE,
n, 1.0, descrpR, valR, csrRowPtrR, csrColIndR,
inforRt, r, t);
cusparseDcsrsv_solve(handle, CUSPARSE_OPERATION_NON_TRANSPOSE,
n, 1.0, descrpR, valR, csrRowPtrR, csrColIndR,
inforR, t, z);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//4: \rho = r^{T} z </span>
rhop= rho;
rho = cublasDdot(n, r, 1, z, 1);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">if</span> (i == 0){
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//6: p = z</span>
cublasDcopy(n, z, 1, p, 1);
}
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">else</span>{
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//8: \beta = rho_{i} / \rho_{i-1}</span>
beta= rho/rhop;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//9: p = z + \beta p</span>
cublasDaxpy(n, beta, p, 1, z, 1);
cublasDcopy(n, z, 1, p, 1);
}
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//11: Compute q = A p (sparse matrix-vector multiplication)</span>
cusparseDcsrmv(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, n, n, 1.0,
descrA, valA, csrRowPtrA, csrColIndA, p, 0.0, q);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//12: \alpha = \rho_{i} / (p^{T} q) </span>
temp = cublasDdot(n, p, 1, q, 1);
alpha= rho/temp;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//13: x = x + \alpha p</span>
cublasDaxpy(n, alpha, p, 1, x, 1);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//14: r = r - \alpha q</span>
cublasDaxpy(n,-alpha, q, 1, r, 1);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//check for convergence </span>
nrmr = cublasDnrm2(n, r, 1);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">if</span> (nrmr/nrmr0 < tol){
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">break</span>;
}
}
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//destroy the analysis info (for lower and upper triangular factors)</span>
cusparseDestroySolveAnalysisInfo(inforRt);
cusparseDestroySolveAnalysisInfo(inforR);
</pre><p class="p">For the nonsymmetric matrix
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mi>A</mi>
</mrow>
</math>
we can let
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mi>M</mi>
</mrow>
</math> be its
incomplete-LU factorization, so that
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mi>A</mi>
<mo>⊬</mo>
<mi>M</mi>
<mo>=</mo>
<mover>
<mi>L</mi>
<mo>˜</mo>
</mover>
<mover>
<mi>U</mi>
<mo>˜</mo>
</mover>
</mrow>
</math>, where
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mover>
<mi>L</mi>
<mo>˜</mo>
</mover>
</mrow>
</math> and
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mover>
<mi>U</mi>
<mo>˜</mo>
</mover>
</mrow>
</math> are lower and upper triangular
matrices, respectively. Let us assume that
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mi>M</mi>
</mrow>
</math> is nonsingular, then
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<msup>
<mi>M</mi>
<mrow>
<mo>−</mo>
<mn>1</mn>
</mrow>
</msup>
<mi>A</mi>
</mrow>
</math> is nonsingular and instead of
solving the linear system <a class="xref" href="index.html#preconditioned-iterative-methods__eq-1" shape="rect">(1)</a>, we can
solve the preconditioned linear system
</p>
<div class="tablenoborder"><a name="preconditioned-iterative-methods__eq-3" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="preconditioned-iterative-methods__eq-3" class="table" frame="void" border="0" rules="none">
<tbody class="tbody">
<tr class="row">
<td class="entry" align="center" valign="top" width="95.23809523809523%" rowspan="1" colspan="1">
<p class="p d4p_eqn_block">
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mrow>
<mo>(</mo>
<mrow>
<msup>
<mi>M</mi>
<mrow>
<mo>−</mo>
<mn>1</mn>
</mrow>
</msup>
<mo></mo>
<mi>A</mi>
</mrow>
<mo>)</mo>
</mrow>
<mo></mo>
<mi mathvariant="bold">x</mi>
<mo>=</mo>
<msup>
<mi>M</mi>
<mrow>
<mo>−</mo>
<mn>1</mn>
</mrow>
</msup>
<mo></mo>
<mi mathvariant="bold">f</mi>
</mrow>
</math>
</p>
</td>
<td class="entry" align="right" valign="top" width="4.761904761904762%" rowspan="1" colspan="1">
<p class="p">(3)</p>
</td>
</tr>
</tbody>
</table>
</div>
<p class="p">The pseudocode for the preconditioned BiCGStab iterative method is
shown in <a class="xref" href="index.html#preconditioned-iterative-methods__algorithm-2-bi-conjugate-gradient-stabilized-bicgstab" shape="rect">Algorithm 2</a>.
</p>
</div>
<div class="example" id="preconditioned-iterative-methods__algorithm-2-bi-conjugate-gradient-stabilized-bicgstab"><a name="preconditioned-iterative-methods__algorithm-2-bi-conjugate-gradient-stabilized-bicgstab" shape="rect">
<!-- --></a><h2 class="title sectiontitle">Algorithm 2 Bi-Conjugate Gradient Stabilized (BiCGStab)</h2>
<p class="p d4p_eqn_block">
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mtable displaystyle="true" equalrows="true" columnalign="right left">
<mtr>
<mtd>
<mi mathvariant="monospace"> </mi>
<mn mathvariant="monospace">1</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext>Letting initial guess be </mtext>
<msub>
<mi mathvariant="bold">x</mi>
<mn>0</mn>
</msub>
<mtext>, compute </mtext>
<mi mathvariant="bold">r</mi>
<mo>←</mo>
<mi mathvariant="bold">f</mi>
<mo>−</mo>
<mi>A</mi>
<msub>
<mi mathvariant="bold">x</mi>
<mn>0</mn>
</msub>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mi mathvariant="monospace"> </mi>
<mn mathvariant="monospace">2</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext>Set </mtext>
<mi mathvariant="bold">p</mi>
<mo>←</mo>
<mi mathvariant="bold">r</mi>
<mtext> and choose </mtext>
<mover accent="true">
<mi mathvariant="bold">r</mi>
<mo>˜</mo>
</mover>
<mtext>, for example you can set </mtext>
<mover accent="true">
<mi mathvariant="bold">r</mi>
<mo>˜</mo>
</mover>
<mo>←</mo>
<mi mathvariant="bold">r</mi>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mtext> </mtext>
</mtd>
</mtr>
<mtr>
<mtd>
<mtext> </mtext>
</mtd>
</mtr>
</mtable>
</math>
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mtable displaystyle="true" equalrows="true" columnalign="right left left">
<mtr>
<mtd>
<mn mathvariant="monospace">3</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext mathvariant="bold">for </mtext>
<mi>i</mi>
<mo>←</mo>
<mn>1</mn>
<mo>,</mo>
<mn>2</mn>
<mo>,</mo>
<mn>...</mn>
<mtext> until convergence </mtext>
<mtext mathvariant="bold">do</mtext>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">4</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<msub>
<mi>ρ</mi>
<mi>i</mi>
</msub>
<mo>←</mo>
<msup>
<mover accent="true">
<mi mathvariant="bold">r</mi>
<mo>˜</mo>
</mover>
<mi>T</mi>
</msup>
<mi mathvariant="bold">r</mi>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">5</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mtext mathvariant="bold">if</mtext>
<mtext> </mtext>
<msub>
<mi>ρ</mi>
<mi>i</mi>
</msub>
<mo>==</mo>
<mn>0.0</mn>
<mtext mathvariant="bold"> then</mtext>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">6</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mtext>method failed</mtext>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">7</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mtext mathvariant="bold">end if</mtext>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">8</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mtext mathvariant="bold">if </mtext>
<mi>i</mi>
<mo>></mo>
<mn>1</mn>
<mtext mathvariant="bold"> then</mtext>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">9</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mtext mathvariant="bold">if </mtext>
<mi>ω</mi>
<mo>==</mo>
<mn>0.0</mn>
<mtext mathvariant="bold"> then</mtext>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">10</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mtext>method failed</mtext>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">11</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mtext mathvariant="bold">end if</mtext>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">12</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mi>β</mi>
<mo>←</mo>
<mfrac>
<mrow>
<msub>
<mi>ρ</mi>
<mi>i</mi>
</msub>
</mrow>
<mrow>
<msub>
<mi>ρ</mi>
<mrow>
<mi>i</mi>
<mo>−</mo>
<mn>1</mn>
</mrow>
</msub>
</mrow>
</mfrac>
<mo>×</mo>
<mrow>
<mo>(</mo>
<mrow>
<mfrac>
<mi>α</mi>
<mi>ω</mi>
</mfrac>
</mrow>
<mo>)</mo>
</mrow>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">13</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mi mathvariant="bold">p</mi>
<mo>←</mo>
<mi mathvariant="bold">r</mi>
<mo>+</mo>
<mtext> </mtext>
<mi>β</mi>
<mo></mo>
<mrow>
<mo>(</mo>
<mrow>
<mi mathvariant="bold">p</mi>
<mo>−</mo>
<mi>ω</mi>
<mo></mo>
<mi mathvariant="bold">v</mi>
</mrow>
<mo>)</mo>
</mrow>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">14</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mtext mathvariant="bold">end if</mtext>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">15</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mtext>Solve </mtext>
<mi>M</mi>
<mo></mo>
<mover accent="true">
<mi mathvariant="bold">p</mi>
<mo>^</mo>
</mover>
<mo>←</mo>
<mi mathvariant="bold">p</mi>
</mrow>
</mtd>
<mtd>
<mrow>
<mo>⊳</mo>
<mtext>Sparse lower and upper triangular solves</mtext>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">16</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mtext>Compute </mtext>
<mi mathvariant="bold">q</mi>
<mo>←</mo>
<mi>A</mi>
<mo></mo>
<mover accent="true">
<mi mathvariant="bold">p</mi>
<mo>^</mo>
</mover>
</mrow>
</mtd>
<mtd>
<mrow>
<mo>⊳</mo>
<mtext>Sparse matrix-vector multiplication</mtext>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">17</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mi>α</mi>
<mo>←</mo>
<mfrac>
<mrow>
<msub>
<mi>ρ</mi>
<mi>i</mi>
</msub>
</mrow>
<mrow>
<msup>
<mover accent="true">
<mi mathvariant="bold">r</mi>
<mo>˜</mo>
</mover>
<mi>T</mi>
</msup>
<mo></mo>
<mi mathvariant="bold">q</mi>
</mrow>
</mfrac>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">18</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mi mathvariant="bold">s</mi>
<mo>←</mo>
<mi mathvariant="bold">r</mi>
<mo>−</mo>
<mi>α</mi>
<mo></mo>
<mi mathvariant="bold">q</mi>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">19</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mi mathvariant="bold">x</mi>
<mo>←</mo>
<mi mathvariant="bold">x</mi>
<mo>+</mo>
<mi>α</mi>
<mo></mo>
<mover accent="true">
<mi mathvariant="bold">p</mi>
<mo>^</mo>
</mover>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">20</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mtext mathvariant="bold">if </mtext>
<msub>
<mrow>
<mrow>
<mo>‖</mo>
<mi>s</mi>
<mo>‖</mo>
</mrow>
</mrow>
<mn>2</mn>
</msub>
<mo>≤</mo>
<mi mathvariant="italic">tol</mi>
<mtext mathvariant="bold"> then</mtext>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">21</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mtext>method converged</mtext>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">22</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mtext mathvariant="bold">end if</mtext>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">23</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mtext>Solve </mtext>
<mi>M</mi>
<mover accent="true">
<mi mathvariant="bold">s</mi>
<mo>^</mo>
</mover>
<mo>←</mo>
<mi mathvariant="bold">s</mi>
</mrow>
</mtd>
<mtd>
<mrow>
<mo>⊳</mo>
<mtext>Sparse lower and upper triangular solves</mtext>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">24</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mtext>Compute </mtext>
<mi mathvariant="bold">t</mi>
<mo>←</mo>
<mi>A</mi>
<mo></mo>
<mover accent="true">
<mi mathvariant="bold">s</mi>
<mo>^</mo>
</mover>
</mrow>
</mtd>
<mtd>
<mrow>
<mo>⊳</mo>
<mtext>Sparse matrix-vector multiplication</mtext>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">25</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mi>ω</mi>
<mo>←</mo>
<mfrac>
<mrow>
<msup>
<mi mathvariant="bold">t</mi>
<mi>T</mi>
</msup>
<mo></mo>
<mi mathvariant="bold">s</mi>
</mrow>
<mrow>
<msup>
<mi mathvariant="bold">t</mi>
<mi>T</mi>
</msup>
<mo></mo>
<mi mathvariant="bold">t</mi>
</mrow>
</mfrac>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">26</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mi mathvariant="bold">x</mi>
<mo>←</mo>
<mi mathvariant="bold">x</mi>
<mo>+</mo>
<mi>ω</mi>
<mo></mo>
<mover accent="true">
<mi mathvariant="bold">s</mi>
<mo>^</mo>
</mover>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">27</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext> </mtext>
<mi mathvariant="bold">r</mi>
<mo>←</mo>
<mi mathvariant="bold">s</mi>
<mo>−</mo>
<mi>ω</mi>
<mo></mo>
<mi mathvariant="bold">t</mi>
</mrow>
</mtd>
</mtr>
<mtr>
<mtd>
<mn mathvariant="monospace">28</mn>
<mo mathvariant="monospace">:</mo>
</mtd>
<mtd>
<mrow>
<mtext mathvariant="bold">end for</mtext>
</mrow>
</mtd>
</mtr>
</mtable>
</math>
</p>
</div>
<div class="section">
<p class="p">Notice that in every iteration of the incomplete-LU preconditioned
BiCGStab iterative method we need to perform two sparse matrix-vector
multiplications and four triangular solves. The corresponding BiCGStab
code using the CUSPARSE and CUBLAS libraries in C programming language
is shown below.
</p><pre xml:space="preserve"><span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-doccomment">/***** BiCGStab Code *****/</span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">/* ASSUMPTIONS:
1. The CUSPARSE and CUBLAS libraries have been initialized.
2. The appropriate memory has been allocated and set to zero.
3. The matrix A (valA, csrRowPtrA, csrColIndA) and the incomplete-
LU lower L (valL, csrRowPtrL, csrColIndL) and upper U (valU,
csrRowPtrU, csrColIndU) triangular factors have been
computed and are present in the device (GPU) memory. */</span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//create the info and analyse the lower and upper triangular factors</span>
cusparseCreateSolveAnalysisInfo(&infoL);
cusparseCreateSolveAnalysisInfo(&infoU);
cusparseDcsrsv_analysis(handle, CUSPARSE_OPERATION_NON_TRANSPOSE,
n, descrL, valL, csrRowPtrL, csrColIndL, infoL);
cusparseDcsrsv_analysis(handle, CUSPARSE_OPERATION_NON_TRANSPOSE,
n, descrU, valU, csrRowPtrU, csrColIndU, infoU);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//1: compute initial residual r = b - A x0 (using initial guess in x)</span>
cusparseDcsrmv(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, n, n, 1.0,
descrA, valA, csrRowPtrA, csrColIndA, x, 0.0, r);
cublasDscal(n,-1.0, r, 1);
cublasDaxpy(n, 1.0, f, 1, r, 1);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//2: Set p=r and \tilde{r}=r</span>
cublasDcopy(n, r, 1, p, 1);
cublasDcopy(n, r, 1, rw,1);
nrmr0 = cublasDnrm2(n, r, 1);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//3: repeat until convergence (based on max. it. and relative residual)</span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">for</span> (i=0; i<maxit; i++){
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//4: \rho = \tilde{r}^{T} r</span>
rhop= rho;
rho = cublasDdot(n, rw, 1, r, 1);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">if</span> (i > 0){
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//12: \beta = (\rho_{i} / \rho_{i-1}) ( \alpha / \omega )</span>
beta= (rho/rhop)*(alpha/omega);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//13: p = r + \beta (p - \omega v)</span>
cublasDaxpy(n,-omega,q, 1, p, 1);
cublasDscal(n, beta, p, 1);
cublasDaxpy(n, 1.0, r, 1, p, 1);
}
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//15: M \hat{p} = p (sparse lower and upper triangular solves)</span>
cusparseDcsrsv_solve(handle, CUSPARSE_OPERATION_NON_TRANSPOSE,
n, 1.0, descrL, valL, csrRowPtrL, csrColIndL,
infoL, p, t);
cusparseDcsrsv_solve(handle, CUSPARSE_OPERATION_NON_TRANSPOSE,
n, 1.0, descrU, valU, csrRowPtrU, csrColIndU,
infoU, t, ph);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//16: q = A \hat{p} (sparse matrix-vector multiplication)</span>
cusparseDcsrmv(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, n, n, 1.0,
descrA, valA, csrRowPtrA, csrColIndA, ph, 0.0, q);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//17: \alpha = \rho_{i} / (\tilde{r}^{T} q)</span>
temp = cublasDdot(n, rw, 1, q, 1);
alpha= rho/temp;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//18: s = r - \alpha q</span>
cublasDaxpy(n,-alpha, q, 1, r, 1);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//19: x = x + \alpha \hat{p}</span>
cublasDaxpy(n, alpha, ph,1, x, 1);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//20: check for convergence</span>
nrmr = cublasDnrm2(n, r, 1);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">if</span> (nrmr/nrmr0 < tol){
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">break</span>;
}
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//23: M \hat{s} = r (sparse lower and upper triangular solves)</span>
cusparseDcsrsv_solve(handle, CUSPARSE_OPERATION_NON_TRANSPOSE,
n, 1.0, descrL, valL, csrRowPtrL, csrColIndL,
infoL, r, t);
cusparseDcsrsv_solve(handle, CUSPARSE_OPERATION_NON_TRANSPOSE,
n, 1.0, descrU, valU, csrRowPtrU, csrColIndU,
infoU, t, s);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//24: t = A \hat{s} (sparse matrix-vector multiplication)</span>
cusparseDcsrmv(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, n, n, 1.0,
descrA, valA, csrRowPtrA, csrColIndA, s, 0.0, t);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//25: \omega = (t^{T} s) / (t^{T} t)</span>
temp = cublasDdot(n, t, 1, r, 1);
temp2= cublasDdot(n, t, 1, t, 1);
omega= temp/temp2;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//26: x = x + \omega \hat{s}</span>
cublasDaxpy(n, omega, s, 1, x, 1);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//27: r = s - \omega t</span>
cublasDaxpy(n,-omega, t, 1, r, 1);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//check for convergence </span>
nrmr = cublasDnrm2(n, r, 1);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">if</span> (nrmr/nrmr0 < tol){
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">break</span>;
}
}
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">//destroy the analysis info (for lower and upper triangular factors)</span>
cusparseDestroySolveAnalysisInfo(infoL);
cusparseDestroySolveAnalysisInfo(infoU);</pre><p class="p">As shown in <a class="xref" href="index.html#preconditioned-iterative-methods__splitting-of-total-time-taken-on-gpu-by-preconditioned-iterative-method" shape="rect">Figure 2</a> the majority of time in each iteration of the
incomplete-LU and Cholesky preconditioned iterative methods is spent in
the sparse matrix-vector multiplication and triangular solve. The
sparse matrix-vector multiplication has already been extensively
studied in the following references <a class="xref" href="index.html#references__3" shape="rect">[3]</a>, <a class="xref" href="index.html#references__15" shape="rect">[15]</a>. The sparse triangular solve is
not as well known, so we briefly point out the strategy used to explore
parallelism in it and refer the reader to the NVIDIA technical report
<a class="xref" href="index.html#references__8" shape="rect">[8]</a> for further details.
</p>
<div class="fig fignone" id="preconditioned-iterative-methods__splitting-of-total-time-taken-on-gpu-by-preconditioned-iterative-method"><a name="preconditioned-iterative-methods__splitting-of-total-time-taken-on-gpu-by-preconditioned-iterative-method" shape="rect">
<!-- --></a><span class="figcap">Figure 2. The Splitting of Total Time Taken on the GPU by the
Preconditioned Iterative Method</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/splitting-of-total-time-taken.png" alt="Figure of the splitting of total time taken on the GPU by the Preconditioned Iterative Method."></img></div><br clear="none"></br></div>
<p class="p">To understand the main ideas behind the sparse triangular solve,
notice that although the forward and back substitution is an inherently
sequential algorithm for dense triangular systems, the dependencies on
the previously obtained elements of the solution do not necessarily
exist for the sparse triangular systems. We pursue the strategy that
takes advantage of the lack of these dependencies and split the
solution process into two phases as mentioned in <a class="xref" href="index.html#references__1" shape="rect">[1]</a>, <a class="xref" href="index.html#references__4" shape="rect">[4]</a>, <a class="xref" href="index.html#references__6" shape="rect">[6]</a>, <a class="xref" href="index.html#references__7" shape="rect">[7]</a>, <a class="xref" href="index.html#references__8" shape="rect">[8]</a>, <a class="xref" href="index.html#references__10" shape="rect">[10]</a>, <a class="xref" href="index.html#references__12" shape="rect">[12]</a>, <a class="xref" href="index.html#references__14" shape="rect">[14]</a>.
</p>
<p class="p">The <dfn class="term">analysis</dfn> phase builds the data dependency graph that
groups independent rows into levels based on the matrix sparsity
pattern. The <dfn class="term">solve</dfn> phase iterates across the constructed
levels one-by-one and computes all elements of the solution
corresponding to the rows at a single level in parallel. Notice that by
construction the rows within each level are independent of each other,
but are dependent on at least one row from the previous level.
</p>
<p class="p">The <dfn class="term">analysis</dfn> phase needs to be performed only once and is
usually significantly slower than the <dfn class="term">solve</dfn> phase, which
can be performed multiple times. This arrangement is ideally suited for
the incomplete-LU and Cholesky preconditioned iterative methods.
</p>
</div>
</div>
</div>
<div class="topic concept nested0" id="numerical-experiments"><a name="numerical-experiments" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#numerical-experiments" name="numerical-experiments" shape="rect">3. Numerical Experiments</a></h2>
<div class="body conbody">
<p class="p">In this section we study the performance of the incomplete-LU and
Cholesky preconditioned <dfn class="term">BiCGStab</dfn> and CG iterative methods.
We use twelve matrices selected from The University of Florida Sparse
Matrix Collection <a class="xref" href="index.html#references__18" shape="rect">[18]</a>
in our numerical experiments. The seven s.p.d. and five nonsymmetric
matrices with the respective number of rows (m), columns (n=m) and
non-zero elements (nnz) are grouped and shown according to their
increasing order in <a class="xref" href="index.html#numerical-experiments__symmetric-positive-definite-spd-and-nonsymmetric-test-matrices" shape="rect">Table 1</a>.
</p>
<div class="tablenoborder"><a name="numerical-experiments__symmetric-positive-definite-spd-and-nonsymmetric-test-matrices" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="numerical-experiments__symmetric-positive-definite-spd-and-nonsymmetric-test-matrices" class="table" frame="border" border="1" rules="all">
<caption><span class="tablecap">Table 1. Symmetric Positive Definite (s.p.d.) and Nonsymmetric Test
Matrices</span></caption>
<thead class="thead" align="left">
<tr class="row">
<th class="entry" valign="top" width="8.333333333333332%" id="d54e3109" rowspan="1" colspan="1">#</th>
<th class="entry" valign="top" width="25%" id="d54e3112" rowspan="1" colspan="1">Matrix</th>
<th class="entry" valign="top" width="16.666666666666664%" id="d54e3115" rowspan="1" colspan="1">m,n</th>
<th class="entry" valign="top" width="16.666666666666664%" id="d54e3118" rowspan="1" colspan="1">nnz</th>
<th class="entry" valign="top" width="8.333333333333332%" id="d54e3121" rowspan="1" colspan="1">s.p.d.</th>
<th class="entry" valign="top" width="25%" id="d54e3125" rowspan="1" colspan="1">Application</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" width="8.333333333333332%" headers="d54e3109" rowspan="1" colspan="1">1.</td>
<td class="entry" valign="top" width="25%" headers="d54e3112" rowspan="1" colspan="1">offshore</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e3115" rowspan="1" colspan="1">259,789</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e3118" rowspan="1" colspan="1">4,242,673</td>
<td class="entry" valign="top" width="8.333333333333332%" headers="d54e3121" rowspan="1" colspan="1">yes</td>
<td class="entry" valign="top" width="25%" headers="d54e3125" rowspan="1" colspan="1">Geophysics</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="8.333333333333332%" headers="d54e3109" rowspan="1" colspan="1">2.</td>
<td class="entry" valign="top" width="25%" headers="d54e3112" rowspan="1" colspan="1">af_shell3</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e3115" rowspan="1" colspan="1">504,855</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e3118" rowspan="1" colspan="1">17,562,051</td>
<td class="entry" valign="top" width="8.333333333333332%" headers="d54e3121" rowspan="1" colspan="1">yes</td>
<td class="entry" valign="top" width="25%" headers="d54e3125" rowspan="1" colspan="1">Mechanics</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="8.333333333333332%" headers="d54e3109" rowspan="1" colspan="1">3.</td>
<td class="entry" valign="top" width="25%" headers="d54e3112" rowspan="1" colspan="1">parabolic_fem</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e3115" rowspan="1" colspan="1">525,825</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e3118" rowspan="1" colspan="1">3,674,625</td>
<td class="entry" valign="top" width="8.333333333333332%" headers="d54e3121" rowspan="1" colspan="1">yes</td>
<td class="entry" valign="top" width="25%" headers="d54e3125" rowspan="1" colspan="1">General</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="8.333333333333332%" headers="d54e3109" rowspan="1" colspan="1">4.</td>
<td class="entry" valign="top" width="25%" headers="d54e3112" rowspan="1" colspan="1">apache2</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e3115" rowspan="1" colspan="1">715,176</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e3118" rowspan="1" colspan="1">4,817,870</td>
<td class="entry" valign="top" width="8.333333333333332%" headers="d54e3121" rowspan="1" colspan="1">yes</td>
<td class="entry" valign="top" width="25%" headers="d54e3125" rowspan="1" colspan="1">Mechanics</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="8.333333333333332%" headers="d54e3109" rowspan="1" colspan="1">5.</td>
<td class="entry" valign="top" width="25%" headers="d54e3112" rowspan="1" colspan="1">ecology2</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e3115" rowspan="1" colspan="1">999,999</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e3118" rowspan="1" colspan="1">4,995,991</td>
<td class="entry" valign="top" width="8.333333333333332%" headers="d54e3121" rowspan="1" colspan="1">yes</td>
<td class="entry" valign="top" width="25%" headers="d54e3125" rowspan="1" colspan="1">Biology</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="8.333333333333332%" headers="d54e3109" rowspan="1" colspan="1">6.</td>
<td class="entry" valign="top" width="25%" headers="d54e3112" rowspan="1" colspan="1">thermal2</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e3115" rowspan="1" colspan="1">1,228,045</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e3118" rowspan="1" colspan="1">8,580,313</td>
<td class="entry" valign="top" width="8.333333333333332%" headers="d54e3121" rowspan="1" colspan="1">yes</td>
<td class="entry" valign="top" width="25%" headers="d54e3125" rowspan="1" colspan="1">Thermal Simulation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="8.333333333333332%" headers="d54e3109" rowspan="1" colspan="1">7.</td>
<td class="entry" valign="top" width="25%" headers="d54e3112" rowspan="1" colspan="1">G3_circuit</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e3115" rowspan="1" colspan="1">1,585,478</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e3118" rowspan="1" colspan="1">7,660,826</td>
<td class="entry" valign="top" width="8.333333333333332%" headers="d54e3121" rowspan="1" colspan="1">yes</td>
<td class="entry" valign="top" width="25%" headers="d54e3125" rowspan="1" colspan="1">Circuit Simulation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="8.333333333333332%" headers="d54e3109" rowspan="1" colspan="1">8.</td>
<td class="entry" valign="top" width="25%" headers="d54e3112" rowspan="1" colspan="1">FEM_3D_thermal2</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e3115" rowspan="1" colspan="1">147,900</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e3118" rowspan="1" colspan="1">3,489,300</td>
<td class="entry" valign="top" width="8.333333333333332%" headers="d54e3121" rowspan="1" colspan="1">no</td>
<td class="entry" valign="top" width="25%" headers="d54e3125" rowspan="1" colspan="1">Mechanics</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="8.333333333333332%" headers="d54e3109" rowspan="1" colspan="1">9.</td>
<td class="entry" valign="top" width="25%" headers="d54e3112" rowspan="1" colspan="1">thermomech_dK</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e3115" rowspan="1" colspan="1">204,316</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e3118" rowspan="1" colspan="1">2,846,228</td>
<td class="entry" valign="top" width="8.333333333333332%" headers="d54e3121" rowspan="1" colspan="1">no</td>
<td class="entry" valign="top" width="25%" headers="d54e3125" rowspan="1" colspan="1">Mechanics</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="8.333333333333332%" headers="d54e3109" rowspan="1" colspan="1">10.</td>
<td class="entry" valign="top" width="25%" headers="d54e3112" rowspan="1" colspan="1">ASIC_320ks</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e3115" rowspan="1" colspan="1">321,671</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e3118" rowspan="1" colspan="1">1,316,08511</td>
<td class="entry" valign="top" width="8.333333333333332%" headers="d54e3121" rowspan="1" colspan="1">no</td>
<td class="entry" valign="top" width="25%" headers="d54e3125" rowspan="1" colspan="1">Circuit Simulation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="8.333333333333332%" headers="d54e3109" rowspan="1" colspan="1">11.</td>
<td class="entry" valign="top" width="25%" headers="d54e3112" rowspan="1" colspan="1">cage13</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e3115" rowspan="1" colspan="1">445,315</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e3118" rowspan="1" colspan="1">7,479,343</td>
<td class="entry" valign="top" width="8.333333333333332%" headers="d54e3121" rowspan="1" colspan="1">no</td>
<td class="entry" valign="top" width="25%" headers="d54e3125" rowspan="1" colspan="1">Biology</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="8.333333333333332%" headers="d54e3109" rowspan="1" colspan="1">12.</td>
<td class="entry" valign="top" width="25%" headers="d54e3112" rowspan="1" colspan="1">atmosmodd</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e3115" rowspan="1" colspan="1">1,270,432</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e3118" rowspan="1" colspan="1">8,814,880</td>
<td class="entry" valign="top" width="8.333333333333332%" headers="d54e3121" rowspan="1" colspan="1">no</td>
<td class="entry" valign="top" width="25%" headers="d54e3125" rowspan="1" colspan="1">Atmospheric Model</td>
</tr>
</tbody>
</table>
</div>
<p class="p">In the following experiments we use the hardware system with NVIDIA
C2050 (ECC on) GPU and Intel Core i7 CPU 950 @ 3.07GHz, using the 64-bit
Linux operating system Ubuntu 10.04 LTS, CUSPARSE library 4.0 and MKL
10.2.3.029. The MKL_NUM_THREADS and MKL_DYNAMIC environment variables are
left unset to allow MKL to use the optimal number of threads.
</p>
<p class="p">We
compute the incomplete-LU and Cholesky factorizations using the MKL
routines <samp class="ph codeph">csrilu0</samp> and <samp class="ph codeph">csrilut</samp> with 0 and
threshold fill-in, respectively. In the <samp class="ph codeph">csrilut</samp> routine
we allow three different levels of fill-in denoted by
(5,10<sup class="ph sup">-3</sup>), (10,10<sup class="ph sup">-5</sup>) and (20,10<sup class="ph sup">-7</sup>). In
general, the
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mo>(</mo>
<mi>k</mi>
<mo>,</mo>
<mi mathvariant="italic">tol</mi>
<mo>)</mo>
</mrow>
</math> fill-in is based on
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mi>n</mi>
<mi>n</mi>
<mi>z</mi>
<mo>/</mo>
<mi>n</mi>
<mo>+</mo>
<mi>k</mi>
</mrow>
</math> maximum allowed number of
elements per row and the dropping of elements with magnitude
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mo>|</mo>
<msub>
<mi>l</mi>
<mrow>
<mi>i</mi>
<mi>j</mi>
</mrow>
</msub>
<mo>|</mo>
<mo>,</mo>
<mo>|</mo>
<msub>
<mi>u</mi>
<mrow>
<mi>i</mi>
<mi>j</mi>
</mrow>
</msub>
<mo>|</mo>
<mo><</mo>
<mi mathvariant="italic">tol</mi>
<mo>×</mo>
<msub>
<mrow>
<mo>‖</mo>
<msubsup>
<mi mathvariant="bold">a</mi>
<mi>i</mi>
<mi>T</mi>
</msubsup>
<mo>‖</mo>
</mrow>
<mn>2</mn>
</msub>
</mrow>
</math>, where
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<msub>
<mi>l</mi>
<mrow>
<mi>i</mi>
<mi>j</mi>
</mrow>
</msub>
</mrow>
</math>,
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<msub>
<mi>u</mi>
<mrow>
<mi>i</mi>
<mi>j</mi>
</mrow>
</msub>
</mrow>
</math> and
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<msubsup>
<mi mathvariant="bold">a</mi>
<mi>i</mi>
<mi>T</mi>
</msubsup>
</mrow>
</math> are the elements of the lower
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mi>L</mi>
</mrow>
</math>, upper
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mi>U</mi>
</mrow>
</math> triangular factors and the <em class="ph i">i</em>-th
row of the coefficient matrix
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mi>A</mi>
</mrow>
</math>,
respectively.
</p>
<p class="p">We compare the implementation of the BiCGStab and CG iterative methods
using the CUSPARSE and CUBLAS libraries on the GPU and MKL on the CPU. In
our experiments we let the initial guess be zero, the right-hand-side
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mi mathvariant="bold">f</mi>
<mo>=</mo>
<mi>A</mi>
<mi mathvariant="bold">e</mi>
</mrow>
</math> where
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<msup>
<mi mathvariant="bold">e</mi>
<mi>T</mi>
</msup>
<msup>
<mrow>
<mo>=</mo>
<mo>(</mo>
<mn>1</mn>
<mo>,</mo>
<mo>…</mo>
<mo>,</mo>
<mn>1</mn>
<mo>)</mo>
</mrow>
<mi>T</mi>
</msup>
</mrow>
</math>, and the stopping criteria be the
maximum number of iterations 2000 or relative residual
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<msub>
<mrow>
<mo>‖</mo>
<msub>
<mi mathvariant="bold">r</mi>
<mi>i</mi>
</msub>
<mo>‖</mo>
</mrow>
<mn>2</mn>
</msub>
<mo>/</mo>
<msub>
<mrow>
<mo>‖</mo>
<msub>
<mi mathvariant="bold">r</mi>
<mn>0</mn>
</msub>
<mo>‖</mo>
</mrow>
<mn>2</mn>
</msub>
<mo><</mo>
<msup>
<mn>10</mn>
<mrow>
<mo>−</mo>
<mn>7</mn>
</mrow>
</msup>
</mrow>
</math>, where
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<msub>
<mi mathvariant="bold">r</mi>
<mi>i</mi>
</msub>
<mo>=</mo>
<mi mathvariant="bold">f</mi>
<mo>−</mo>
<mi>A</mi>
<msub>
<mi mathvariant="bold">x</mi>
<mi>i</mi>
</msub>
</mrow>
</math> is the residual at <em class="ph i">i</em>-th
iteration.
</p>
<div class="tablenoborder"><a name="numerical-experiments__csrilu0-preconditioned-cg-and-bicgstab-methods" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="numerical-experiments__csrilu0-preconditioned-cg-and-bicgstab-methods" class="table" frame="border" border="1" rules="all">
<caption><span class="tablecap">Table 2. <samp class="ph codeph">csrilu0</samp> Preconditioned CG and BiCGStab
Methods</span></caption>
<thead class="thead" align="left">
<tr class="row">
<th class="entry" valign="top" width="5.263157894736842%" id="d54e3900" rowspan="1" colspan="1"> </th>
<th class="entry" colspan="2" align="center" valign="top" id="d54e3902" rowspan="1">ilu0</th>
<th class="entry" colspan="3" align="center" valign="top" id="d54e3905" rowspan="1">CPU</th>
<th class="entry" colspan="3" align="center" valign="top" id="d54e3908" rowspan="1">GPU</th>
<th class="entry" align="center" valign="top" width="10.526315789473683%" id="d54e3911" rowspan="1" colspan="1">Speedup</th>
</tr>
<tr class="row">
<th class="entry" valign="top" width="5.263157894736842%" id="d54e3917" rowspan="1" colspan="1">#</th>
<th class="entry" valign="top" width="10.526315789473683%" id="d54e3920" rowspan="1" colspan="1">fact. time(s)</th>
<th class="entry" valign="top" width="10.526315789473683%" id="d54e3923" rowspan="1" colspan="1">copy time(s)</th>
<th class="entry" valign="top" width="10.526315789473683%" id="d54e3926" rowspan="1" colspan="1">solve time(s)</th>
<th class="entry" valign="top" width="10.526315789473683%" id="d54e3929" rowspan="1" colspan="1">
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mfrac>
<mrow>
<msub>
<mrow>
<mo>‖</mo>
<msub>
<mi mathvariant="bold">r</mi>
<mi>i</mi>
</msub>
<mo>‖</mo>
</mrow>
<mn>2</mn>
</msub>
</mrow>
<mrow>
<msub>
<mrow>
<mo>‖</mo>
<msub>
<mi mathvariant="bold">r</mi>
<mn>0</mn>
</msub>
<mo>‖</mo>
</mrow>
<mn>2</mn>
</msub>
</mrow>
</mfrac>
</mrow>
</math>
</th>
<th class="entry" valign="top" width="10.526315789473683%" id="d54e3996" rowspan="1" colspan="1"># it.</th>
<th class="entry" valign="top" width="10.526315789473683%" id="d54e3999" rowspan="1" colspan="1">solve time(s)</th>
<th class="entry" valign="top" width="10.526315789473683%" id="d54e4002" rowspan="1" colspan="1">
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mfrac>
<mrow>
<msub>
<mrow>
<mo>‖</mo>
<msub>
<mi mathvariant="bold">r</mi>
<mi>i</mi>
</msub>
<mo>‖</mo>
</mrow>
<mn>2</mn>
</msub>
</mrow>
<mrow>
<msub>
<mrow>
<mo>‖</mo>
<msub>
<mi mathvariant="bold">r</mi>
<mn>0</mn>
</msub>
<mo>‖</mo>
</mrow>
<mn>2</mn>
</msub>
</mrow>
</mfrac>
</mrow>
</math>
</th>
<th class="entry" valign="top" width="10.526315789473683%" id="d54e4068" rowspan="1" colspan="1"># it.</th>
<th class="entry" valign="top" width="10.526315789473683%" id="d54e4071" rowspan="1" colspan="1">vs. ilu0</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" width="5.263157894736842%" headers="d54e3900 d54e3917" rowspan="1" colspan="1">1</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3902 d54e3920" rowspan="1" colspan="1">0.38</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3902 d54e3923" rowspan="1" colspan="1">0.02</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3926" rowspan="1" colspan="1">0.72</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3929" rowspan="1" colspan="1">8.83E-08</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3996" rowspan="1" colspan="1">25</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e3999" rowspan="1" colspan="1">1.52</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e4002" rowspan="1" colspan="1">8.83E-08</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e4068" rowspan="1" colspan="1">25</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3911 d54e4071" rowspan="1" colspan="1">0.57</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="5.263157894736842%" headers="d54e3900 d54e3917" rowspan="1" colspan="1">2</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3902 d54e3920" rowspan="1" colspan="1">1.62</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3902 d54e3923" rowspan="1" colspan="1">0.04</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3926" rowspan="1" colspan="1">38.5</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3929" rowspan="1" colspan="1">1.00E-07</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3996" rowspan="1" colspan="1">569</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e3999" rowspan="1" colspan="1">33.9</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e4002" rowspan="1" colspan="1">9.69E-08</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e4068" rowspan="1" colspan="1">571</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3911 d54e4071" rowspan="1" colspan="1">1.13</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="5.263157894736842%" headers="d54e3900 d54e3917" rowspan="1" colspan="1">3</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3902 d54e3920" rowspan="1" colspan="1">0.13</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3902 d54e3923" rowspan="1" colspan="1">0.01</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3926" rowspan="1" colspan="1">39.2</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3929" rowspan="1" colspan="1">9.84E-08</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3996" rowspan="1" colspan="1">1044</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e3999" rowspan="1" colspan="1">6.91</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e4002" rowspan="1" colspan="1">9.84E-08</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e4068" rowspan="1" colspan="1">1044</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3911 d54e4071" rowspan="1" colspan="1">5.59</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="5.263157894736842%" headers="d54e3900 d54e3917" rowspan="1" colspan="1">4</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3902 d54e3920" rowspan="1" colspan="1">0.12</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3902 d54e3923" rowspan="1" colspan="1">0.01</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3926" rowspan="1" colspan="1">35.0</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3929" rowspan="1" colspan="1">9.97E-08</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3996" rowspan="1" colspan="1">713</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e3999" rowspan="1" colspan="1">12.8</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e4002" rowspan="1" colspan="1">9.97E-08</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e4068" rowspan="1" colspan="1">713</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3911 d54e4071" rowspan="1" colspan="1">2.72</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="5.263157894736842%" headers="d54e3900 d54e3917" rowspan="1" colspan="1">5</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3902 d54e3920" rowspan="1" colspan="1">0.09</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3902 d54e3923" rowspan="1" colspan="1">0.01</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3926" rowspan="1" colspan="1">107</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3929" rowspan="1" colspan="1">9.98E-08</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3996" rowspan="1" colspan="1">1746</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e3999" rowspan="1" colspan="1">55.3</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e4002" rowspan="1" colspan="1">9.98E-08</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e4068" rowspan="1" colspan="1">1746</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3911 d54e4071" rowspan="1" colspan="1">1.92</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="5.263157894736842%" headers="d54e3900 d54e3917" rowspan="1" colspan="1">6</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3902 d54e3920" rowspan="1" colspan="1">0.40</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3902 d54e3923" rowspan="1" colspan="1">0.02</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3926" rowspan="1" colspan="1">155.</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3929" rowspan="1" colspan="1">9.96E-08</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3996" rowspan="1" colspan="1">1656</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e3999" rowspan="1" colspan="1">54.4</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e4002" rowspan="1" colspan="1">9.79E-08</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e4068" rowspan="1" colspan="1">1656</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3911 d54e4071" rowspan="1" colspan="1">2.83</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="5.263157894736842%" headers="d54e3900 d54e3917" rowspan="1" colspan="1">7</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3902 d54e3920" rowspan="1" colspan="1">0.16</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3902 d54e3923" rowspan="1" colspan="1">0.02</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3926" rowspan="1" colspan="1">20.2</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3929" rowspan="1" colspan="1">8.70E-08</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3996" rowspan="1" colspan="1">183</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e3999" rowspan="1" colspan="1">8.61</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e4002" rowspan="1" colspan="1">8.22E-08</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e4068" rowspan="1" colspan="1">183</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3911 d54e4071" rowspan="1" colspan="1">2.32</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="5.263157894736842%" headers="d54e3900 d54e3917" rowspan="1" colspan="1">8</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3902 d54e3920" rowspan="1" colspan="1">0.32</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3902 d54e3923" rowspan="1" colspan="1">0.02</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3926" rowspan="1" colspan="1">0.13</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3929" rowspan="1" colspan="1">5.25E-08</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3996" rowspan="1" colspan="1">4</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e3999" rowspan="1" colspan="1">0.52</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e4002" rowspan="1" colspan="1">5.25E-08</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e4068" rowspan="1" colspan="1">4</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3911 d54e4071" rowspan="1" colspan="1">0.53</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="5.263157894736842%" headers="d54e3900 d54e3917" rowspan="1" colspan="1">9</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3902 d54e3920" rowspan="1" colspan="1">0.20</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3902 d54e3923" rowspan="1" colspan="1">0.01</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3926" rowspan="1" colspan="1">72.7</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3929" rowspan="1" colspan="1">1.96E-04</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3996" rowspan="1" colspan="1">2000</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e3999" rowspan="1" colspan="1">40.0</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e4002" rowspan="1" colspan="1">2.08E-04</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e4068" rowspan="1" colspan="1">2000</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3911 d54e4071" rowspan="1" colspan="1">1.80</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="5.263157894736842%" headers="d54e3900 d54e3917" rowspan="1" colspan="1">10</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3902 d54e3920" rowspan="1" colspan="1">0.11</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3902 d54e3923" rowspan="1" colspan="1">0.01</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3926" rowspan="1" colspan="1">0.27</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3929" rowspan="1" colspan="1">6.33E-08</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3996" rowspan="1" colspan="1">6</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e3999" rowspan="1" colspan="1">0.12</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e4002" rowspan="1" colspan="1">6.33E-08</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e4068" rowspan="1" colspan="1">6</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3911 d54e4071" rowspan="1" colspan="1">1.59</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="5.263157894736842%" headers="d54e3900 d54e3917" rowspan="1" colspan="1">11</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3902 d54e3920" rowspan="1" colspan="1">0.70</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3902 d54e3923" rowspan="1" colspan="1">0.03</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3926" rowspan="1" colspan="1">0.28</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3929" rowspan="1" colspan="1">2.52E-08</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3996" rowspan="1" colspan="1">2.5</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e3999" rowspan="1" colspan="1">0.15</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e4002" rowspan="1" colspan="1">2.52E-08</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e4068" rowspan="1" colspan="1">2.5</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3911 d54e4071" rowspan="1" colspan="1">1.10</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="5.263157894736842%" headers="d54e3900 d54e3917" rowspan="1" colspan="1">12</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3902 d54e3920" rowspan="1" colspan="1">0.25</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3902 d54e3923" rowspan="1" colspan="1">0.04</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3926" rowspan="1" colspan="1">12.5</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3929" rowspan="1" colspan="1">7.33E-08</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3905 d54e3996" rowspan="1" colspan="1">76.5</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e3999" rowspan="1" colspan="1">4.30</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e4002" rowspan="1" colspan="1">9.69E-08</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3908 d54e4068" rowspan="1" colspan="1">74.5</td>
<td class="entry" valign="top" width="10.526315789473683%" headers="d54e3911 d54e4071" rowspan="1" colspan="1">2.79</td>
</tr>
</tbody>
</table>
</div>
<div class="tablenoborder">
<table cellpadding="4" cellspacing="0" summary="" class="table" frame="border" border="1" rules="all">
<caption><span class="tablecap">Table 3. <samp class="ph codeph">csrilut</samp>(5,10<sup class="ph sup">-3</sup>) Preconditioned CG and
BiCGStab Methods</span></caption>
<thead class="thead" align="left">
<tr class="row">
<th class="entry" valign="top" width="4.854368932038835%" id="d54e4532" rowspan="1" colspan="1"> </th>
<th class="entry" colspan="2" align="center" valign="top" id="d54e4534" rowspan="1">ilut(5,10<sup class="ph sup">-3</sup>)
</th>
<th class="entry" colspan="3" align="center" valign="top" id="d54e4540" rowspan="1">CPU</th>
<th class="entry" colspan="3" align="center" valign="top" id="d54e4543" rowspan="1">GPU</th>
<th class="entry" colspan="2" align="center" valign="top" id="d54e4546" rowspan="1">Speedup</th>
</tr>
<tr class="row">
<th class="entry" valign="top" width="4.854368932038835%" id="d54e4552" rowspan="1" colspan="1">#</th>
<th class="entry" valign="top" width="9.70873786407767%" id="d54e4555" rowspan="1" colspan="1">fact. time(s)</th>
<th class="entry" valign="top" width="9.70873786407767%" id="d54e4558" rowspan="1" colspan="1">copy time(s)</th>
<th class="entry" valign="top" width="9.70873786407767%" id="d54e4561" rowspan="1" colspan="1">solve time(s)</th>
<th class="entry" valign="top" width="9.70873786407767%" id="d54e4564" rowspan="1" colspan="1">
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mfrac>
<mrow>
<msub>
<mrow>
<mo>‖</mo>
<msub>
<mi mathvariant="bold">r</mi>
<mi>i</mi>
</msub>
<mo>‖</mo>
</mrow>
<mn>2</mn>
</msub>
</mrow>
<mrow>
<msub>
<mrow>
<mo>‖</mo>
<msub>
<mi mathvariant="bold">r</mi>
<mn>0</mn>
</msub>
<mo>‖</mo>
</mrow>
<mn>2</mn>
</msub>
</mrow>
</mfrac>
</mrow>
</math>
</th>
<th class="entry" valign="top" width="7.281553398058252%" id="d54e4631" rowspan="1" colspan="1"># it.</th>
<th class="entry" valign="top" width="9.70873786407767%" id="d54e4634" rowspan="1" colspan="1">solve time(s)</th>
<th class="entry" valign="top" width="9.70873786407767%" id="d54e4637" rowspan="1" colspan="1">
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mfrac>
<mrow>
<msub>
<mrow>
<mo>‖</mo>
<msub>
<mi mathvariant="bold">r</mi>
<mi>i</mi>
</msub>
<mo>‖</mo>
</mrow>
<mn>2</mn>
</msub>
</mrow>
<mrow>
<msub>
<mrow>
<mo>‖</mo>
<msub>
<mi mathvariant="bold">r</mi>
<mn>0</mn>
</msub>
<mo>‖</mo>
</mrow>
<mn>2</mn>
</msub>
</mrow>
</mfrac>
</mrow>
</math>
</th>
<th class="entry" valign="top" width="7.281553398058252%" id="d54e4703" rowspan="1" colspan="1"># it.</th>
<th class="entry" valign="top" width="11.165048543689318%" id="d54e4706" rowspan="1" colspan="1">vs. ilut (5,10<sup class="ph sup">-3</sup>)
</th>
<th class="entry" valign="top" width="11.165048543689318%" id="d54e4712" rowspan="1" colspan="1">vs. ilu0</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" width="4.854368932038835%" headers="d54e4532 d54e4552" rowspan="1" colspan="1">1</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4534 d54e4555" rowspan="1" colspan="1">0.14</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4534 d54e4558" rowspan="1" colspan="1">0.01</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4540 d54e4561" rowspan="1" colspan="1">1.17</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4540 d54e4564" rowspan="1" colspan="1">9.70E-08</td>
<td class="entry" valign="top" width="7.281553398058252%" headers="d54e4540 d54e4631" rowspan="1" colspan="1">32</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4543 d54e4634" rowspan="1" colspan="1">1.82</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4543 d54e4637" rowspan="1" colspan="1">9.70E-08</td>
<td class="entry" valign="top" width="7.281553398058252%" headers="d54e4543 d54e4703" rowspan="1" colspan="1">32</td>
<td class="entry" valign="top" width="11.165048543689318%" headers="d54e4546 d54e4706" rowspan="1" colspan="1">0.67</td>
<td class="entry" valign="top" width="11.165048543689318%" headers="d54e4546 d54e4712" rowspan="1" colspan="1">0.69</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.854368932038835%" headers="d54e4532 d54e4552" rowspan="1" colspan="1">2</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4534 d54e4555" rowspan="1" colspan="1">0.51</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4534 d54e4558" rowspan="1" colspan="1">0.03</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4540 d54e4561" rowspan="1" colspan="1">49.1</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4540 d54e4564" rowspan="1" colspan="1">9.89E-08</td>
<td class="entry" valign="top" width="7.281553398058252%" headers="d54e4540 d54e4631" rowspan="1" colspan="1">748</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4543 d54e4634" rowspan="1" colspan="1">33.6</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4543 d54e4637" rowspan="1" colspan="1">9.89E-08</td>
<td class="entry" valign="top" width="7.281553398058252%" headers="d54e4543 d54e4703" rowspan="1" colspan="1">748</td>
<td class="entry" valign="top" width="11.165048543689318%" headers="d54e4546 d54e4706" rowspan="1" colspan="1">1.45</td>
<td class="entry" valign="top" width="11.165048543689318%" headers="d54e4546 d54e4712" rowspan="1" colspan="1">1.39</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.854368932038835%" headers="d54e4532 d54e4552" rowspan="1" colspan="1">3</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4534 d54e4555" rowspan="1" colspan="1">1.47</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4534 d54e4558" rowspan="1" colspan="1">0.02</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4540 d54e4561" rowspan="1" colspan="1">11.7</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4540 d54e4564" rowspan="1" colspan="1">9.72E-08</td>
<td class="entry" valign="top" width="7.281553398058252%" headers="d54e4540 d54e4631" rowspan="1" colspan="1">216</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4543 d54e4634" rowspan="1" colspan="1">6.93</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4543 d54e4637" rowspan="1" colspan="1">9.72E-08</td>
<td class="entry" valign="top" width="7.281553398058252%" headers="d54e4543 d54e4703" rowspan="1" colspan="1">216</td>
<td class="entry" valign="top" width="11.165048543689318%" headers="d54e4546 d54e4706" rowspan="1" colspan="1">1.56</td>
<td class="entry" valign="top" width="11.165048543689318%" headers="d54e4546 d54e4712" rowspan="1" colspan="1">1.86</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.854368932038835%" headers="d54e4532 d54e4552" rowspan="1" colspan="1">4</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4534 d54e4555" rowspan="1" colspan="1">0.17</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4534 d54e4558" rowspan="1" colspan="1">0.01</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4540 d54e4561" rowspan="1" colspan="1">67.9</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4540 d54e4564" rowspan="1" colspan="1">9.96E-08</td>
<td class="entry" valign="top" width="7.281553398058252%" headers="d54e4540 d54e4631" rowspan="1" colspan="1">1495</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4543 d54e4634" rowspan="1" colspan="1">26.5</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4543 d54e4637" rowspan="1" colspan="1">9.96E-08</td>
<td class="entry" valign="top" width="7.281553398058252%" headers="d54e4543 d54e4703" rowspan="1" colspan="1">1495</td>
<td class="entry" valign="top" width="11.165048543689318%" headers="d54e4546 d54e4706" rowspan="1" colspan="1">2.56</td>
<td class="entry" valign="top" width="11.165048543689318%" headers="d54e4546 d54e4712" rowspan="1" colspan="1">5.27</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.854368932038835%" headers="d54e4532 d54e4552" rowspan="1" colspan="1">5</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4534 d54e4555" rowspan="1" colspan="1">0.55</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4534 d54e4558" rowspan="1" colspan="1">0.04</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4540 d54e4561" rowspan="1" colspan="1">59.5</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4540 d54e4564" rowspan="1" colspan="1">9.22E-08</td>
<td class="entry" valign="top" width="7.281553398058252%" headers="d54e4540 d54e4631" rowspan="1" colspan="1">653</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4543 d54e4634" rowspan="1" colspan="1">71.6</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4543 d54e4637" rowspan="1" colspan="1">9.22E-08</td>
<td class="entry" valign="top" width="7.281553398058252%" headers="d54e4543 d54e4703" rowspan="1" colspan="1">653</td>
<td class="entry" valign="top" width="11.165048543689318%" headers="d54e4546 d54e4706" rowspan="1" colspan="1">0.83</td>
<td class="entry" valign="top" width="11.165048543689318%" headers="d54e4546 d54e4712" rowspan="1" colspan="1">1.08</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.854368932038835%" headers="d54e4532 d54e4552" rowspan="1" colspan="1">6</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4534 d54e4555" rowspan="1" colspan="1">3.59</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4534 d54e4558" rowspan="1" colspan="1">0.05</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4540 d54e4561" rowspan="1" colspan="1">47.0</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4540 d54e4564" rowspan="1" colspan="1">9.50E-08</td>
<td class="entry" valign="top" width="7.281553398058252%" headers="d54e4540 d54e4631" rowspan="1" colspan="1">401</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4543 d54e4634" rowspan="1" colspan="1">90.1</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4543 d54e4637" rowspan="1" colspan="1">9.64E-08</td>
<td class="entry" valign="top" width="7.281553398058252%" headers="d54e4543 d54e4703" rowspan="1" colspan="1">401</td>
<td class="entry" valign="top" width="11.165048543689318%" headers="d54e4546 d54e4706" rowspan="1" colspan="1">0.54</td>
<td class="entry" valign="top" width="11.165048543689318%" headers="d54e4546 d54e4712" rowspan="1" colspan="1">0.92</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.854368932038835%" headers="d54e4532 d54e4552" rowspan="1" colspan="1">7</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4534 d54e4555" rowspan="1" colspan="1">1.24</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4534 d54e4558" rowspan="1" colspan="1">0.05</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4540 d54e4561" rowspan="1" colspan="1">23.1</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4540 d54e4564" rowspan="1" colspan="1">8.08E-08</td>
<td class="entry" valign="top" width="7.281553398058252%" headers="d54e4540 d54e4631" rowspan="1" colspan="1">153</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4543 d54e4634" rowspan="1" colspan="1">24.8</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4543 d54e4637" rowspan="1" colspan="1">8.08E-08</td>
<td class="entry" valign="top" width="7.281553398058252%" headers="d54e4543 d54e4703" rowspan="1" colspan="1">153</td>
<td class="entry" valign="top" width="11.165048543689318%" headers="d54e4546 d54e4706" rowspan="1" colspan="1">0.93</td>
<td class="entry" valign="top" width="11.165048543689318%" headers="d54e4546 d54e4712" rowspan="1" colspan="1">2.77</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.854368932038835%" headers="d54e4532 d54e4552" rowspan="1" colspan="1">8</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4534 d54e4555" rowspan="1" colspan="1">0.82</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4534 d54e4558" rowspan="1" colspan="1">0.03</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4540 d54e4561" rowspan="1" colspan="1">0.12</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4540 d54e4564" rowspan="1" colspan="1">3.97E-08</td>
<td class="entry" valign="top" width="7.281553398058252%" headers="d54e4540 d54e4631" rowspan="1" colspan="1">2</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4543 d54e4634" rowspan="1" colspan="1">1.12</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4543 d54e4637" rowspan="1" colspan="1">3.97E-08</td>
<td class="entry" valign="top" width="7.281553398058252%" headers="d54e4543 d54e4703" rowspan="1" colspan="1">2</td>
<td class="entry" valign="top" width="11.165048543689318%" headers="d54e4546 d54e4706" rowspan="1" colspan="1">0.48</td>
<td class="entry" valign="top" width="11.165048543689318%" headers="d54e4546 d54e4712" rowspan="1" colspan="1">1.10</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.854368932038835%" headers="d54e4532 d54e4552" rowspan="1" colspan="1">9</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4534 d54e4555" rowspan="1" colspan="1">0.10</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4534 d54e4558" rowspan="1" colspan="1">0.01</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4540 d54e4561" rowspan="1" colspan="1">54.3</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4540 d54e4564" rowspan="1" colspan="1">5.68E-04</td>
<td class="entry" valign="top" width="7.281553398058252%" headers="d54e4540 d54e4631" rowspan="1" colspan="1">2000</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4543 d54e4634" rowspan="1" colspan="1">24.5</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4543 d54e4637" rowspan="1" colspan="1">1.58E-04</td>
<td class="entry" valign="top" width="7.281553398058252%" headers="d54e4543 d54e4703" rowspan="1" colspan="1">2000</td>
<td class="entry" valign="top" width="11.165048543689318%" headers="d54e4546 d54e4706" rowspan="1" colspan="1">2.21</td>
<td class="entry" valign="top" width="11.165048543689318%" headers="d54e4546 d54e4712" rowspan="1" colspan="1">1.34</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.854368932038835%" headers="d54e4532 d54e4552" rowspan="1" colspan="1">10</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4534 d54e4555" rowspan="1" colspan="1">0.12</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4534 d54e4558" rowspan="1" colspan="1">0.01</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4540 d54e4561" rowspan="1" colspan="1">0.16</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4540 d54e4564" rowspan="1" colspan="1">4.89E-08</td>
<td class="entry" valign="top" width="7.281553398058252%" headers="d54e4540 d54e4631" rowspan="1" colspan="1">4</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4543 d54e4634" rowspan="1" colspan="1">0.08</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4543 d54e4637" rowspan="1" colspan="1">6.45E-08</td>
<td class="entry" valign="top" width="7.281553398058252%" headers="d54e4543 d54e4703" rowspan="1" colspan="1">4</td>
<td class="entry" valign="top" width="11.165048543689318%" headers="d54e4546 d54e4706" rowspan="1" colspan="1">1.37</td>
<td class="entry" valign="top" width="11.165048543689318%" headers="d54e4546 d54e4712" rowspan="1" colspan="1">1.15</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.854368932038835%" headers="d54e4532 d54e4552" rowspan="1" colspan="1">11</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4534 d54e4555" rowspan="1" colspan="1">4.99</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4534 d54e4558" rowspan="1" colspan="1">0.07</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4540 d54e4561" rowspan="1" colspan="1">0.36</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4540 d54e4564" rowspan="1" colspan="1">1.40E-08</td>
<td class="entry" valign="top" width="7.281553398058252%" headers="d54e4540 d54e4631" rowspan="1" colspan="1">2.5</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4543 d54e4634" rowspan="1" colspan="1">0.37</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4543 d54e4637" rowspan="1" colspan="1">1.40E-08</td>
<td class="entry" valign="top" width="7.281553398058252%" headers="d54e4543 d54e4703" rowspan="1" colspan="1">2.5</td>
<td class="entry" valign="top" width="11.165048543689318%" headers="d54e4546 d54e4706" rowspan="1" colspan="1">0.99</td>
<td class="entry" valign="top" width="11.165048543689318%" headers="d54e4546 d54e4712" rowspan="1" colspan="1">6.05</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.854368932038835%" headers="d54e4532 d54e4552" rowspan="1" colspan="1">12</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4534 d54e4555" rowspan="1" colspan="1">0.32</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4534 d54e4558" rowspan="1" colspan="1">0.03</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4540 d54e4561" rowspan="1" colspan="1">39.2</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4540 d54e4564" rowspan="1" colspan="1">7.05E-08</td>
<td class="entry" valign="top" width="7.281553398058252%" headers="d54e4540 d54e4631" rowspan="1" colspan="1">278.5</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4543 d54e4634" rowspan="1" colspan="1">10.6</td>
<td class="entry" valign="top" width="9.70873786407767%" headers="d54e4543 d54e4637" rowspan="1" colspan="1">8.82E-08</td>
<td class="entry" valign="top" width="7.281553398058252%" headers="d54e4543 d54e4703" rowspan="1" colspan="1">270.5</td>
<td class="entry" valign="top" width="11.165048543689318%" headers="d54e4546 d54e4706" rowspan="1" colspan="1">3.60</td>
<td class="entry" valign="top" width="11.165048543689318%" headers="d54e4546 d54e4712" rowspan="1" colspan="1">8.60</td>
</tr>
</tbody>
</table>
</div>
<p class="p">The results of the numerical experiments are shown in <a class="xref" href="index.html#numerical-experiments__csrilu0-preconditioned-cg-and-bicgstab-methods" shape="rect">Table 2</a> through <a class="xref" href="index.html#numerical-experiments__csrilut-20-10-preconditioned-cg-and-bicgstab-methods" shape="rect">Table 5</a>, where we state the speedup obtained by the iterative
method on the GPU over CPU (speedup), number of iterations required for
convergence (# it.), achieved relative residual (
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mfrac>
<mrow>
<msub>
<mrow>
<mo>‖</mo>
<msub>
<mi mathvariant="bold">r</mi>
<mi>i</mi>
</msub>
<mo>‖</mo>
</mrow>
<mn>2</mn>
</msub>
</mrow>
<mrow>
<msub>
<mrow>
<mo>‖</mo>
<msub>
<mi mathvariant="bold">r</mi>
<mn>0</mn>
</msub>
<mo>‖</mo>
</mrow>
<mn>2</mn>
</msub>
</mrow>
</mfrac>
</mrow>
</math>) and time in seconds taken by the
factorization (fact.), iterative solution of the linear system (solve),
and <samp class="ph codeph">cudaMemcpy</samp> of the lower and upper triangular factors
to the GPU (copy). We include the time taken to compute the incomplete-LU
and Cholesky factorization as well as to transfer the triangular factors
from the CPU to the GPU memory in the computed speedup.
</p>
<div class="tablenoborder">
<table cellpadding="4" cellspacing="0" summary="" class="table" frame="border" border="1" rules="all">
<caption><span class="tablecap">Table 4.
<samp class="ph codeph">csrilut</samp>(10,10<sup class="ph sup">-5</sup>) Preconditioned CG and
BiCGStab Methods</span></caption>
<thead class="thead" align="left">
<tr class="row">
<th class="entry" valign="top" width="4.8543689320388355%" id="d54e5286" rowspan="1" colspan="1"> </th>
<th class="entry" colspan="2" align="center" valign="top" id="d54e5288" rowspan="1">ilut(10,10<sup class="ph sup">-5</sup>)
</th>
<th class="entry" colspan="3" align="center" valign="top" id="d54e5294" rowspan="1">CPU</th>
<th class="entry" colspan="3" align="center" valign="top" id="d54e5297" rowspan="1">GPU</th>
<th class="entry" colspan="2" align="center" valign="top" id="d54e5300" rowspan="1">Speedup</th>
</tr>
<tr class="row">
<th class="entry" valign="top" width="4.8543689320388355%" id="d54e5306" rowspan="1" colspan="1">#</th>
<th class="entry" valign="top" width="9.708737864077671%" id="d54e5309" rowspan="1" colspan="1">fact. time(s)</th>
<th class="entry" valign="top" width="9.708737864077671%" id="d54e5312" rowspan="1" colspan="1">copy time(s)</th>
<th class="entry" valign="top" width="9.708737864077671%" id="d54e5315" rowspan="1" colspan="1">solve time(s)</th>
<th class="entry" valign="top" width="9.708737864077671%" id="d54e5318" rowspan="1" colspan="1">
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mfrac>
<mrow>
<msub>
<mrow>
<mo>‖</mo>
<msub>
<mi mathvariant="bold">r</mi>
<mi>i</mi>
</msub>
<mo>‖</mo>
</mrow>
<mn>2</mn>
</msub>
</mrow>
<mrow>
<msub>
<mrow>
<mo>‖</mo>
<msub>
<mi mathvariant="bold">r</mi>
<mn>0</mn>
</msub>
<mo>‖</mo>
</mrow>
<mn>2</mn>
</msub>
</mrow>
</mfrac>
</mrow>
</math>
</th>
<th class="entry" valign="top" width="6.796116504854369%" id="d54e5385" rowspan="1" colspan="1"># it.</th>
<th class="entry" valign="top" width="9.708737864077671%" id="d54e5388" rowspan="1" colspan="1">solve time(s)</th>
<th class="entry" valign="top" width="9.708737864077671%" id="d54e5391" rowspan="1" colspan="1">
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mfrac>
<mrow>
<msub>
<mrow>
<mo>‖</mo>
<msub>
<mi mathvariant="bold">r</mi>
<mi>i</mi>
</msub>
<mo>‖</mo>
</mrow>
<mn>2</mn>
</msub>
</mrow>
<mrow>
<msub>
<mrow>
<mo>‖</mo>
<msub>
<mi mathvariant="bold">r</mi>
<mn>0</mn>
</msub>
<mo>‖</mo>
</mrow>
<mn>2</mn>
</msub>
</mrow>
</mfrac>
</mrow>
</math>
</th>
<th class="entry" valign="top" width="6.796116504854369%" id="d54e5457" rowspan="1" colspan="1"># it.</th>
<th class="entry" valign="top" width="11.650485436893204%" id="d54e5460" rowspan="1" colspan="1">vs. ilut (10,10<sup class="ph sup">-5</sup>)
</th>
<th class="entry" valign="top" width="11.650485436893204%" id="d54e5466" rowspan="1" colspan="1">vs. ilu0</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" width="4.8543689320388355%" headers="d54e5286 d54e5306" rowspan="1" colspan="1">1</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5288 d54e5309" rowspan="1" colspan="1">0.15</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5288 d54e5312" rowspan="1" colspan="1">0.01</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5294 d54e5315" rowspan="1" colspan="1">1.06</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5294 d54e5318" rowspan="1" colspan="1">8.79E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5294 d54e5385" rowspan="1" colspan="1">34</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5297 d54e5388" rowspan="1" colspan="1">1.96</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5297 d54e5391" rowspan="1" colspan="1">8.79E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5297 d54e5457" rowspan="1" colspan="1">34</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5300 d54e5460" rowspan="1" colspan="1">0.57</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5300 d54e5466" rowspan="1" colspan="1">0.63</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.8543689320388355%" headers="d54e5286 d54e5306" rowspan="1" colspan="1">2</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5288 d54e5309" rowspan="1" colspan="1">0.52</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5288 d54e5312" rowspan="1" colspan="1">0.03</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5294 d54e5315" rowspan="1" colspan="1">60.0</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5294 d54e5318" rowspan="1" colspan="1">9.86E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5294 d54e5385" rowspan="1" colspan="1">748</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5297 d54e5388" rowspan="1" colspan="1">38.7</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5297 d54e5391" rowspan="1" colspan="1">9.86E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5297 d54e5457" rowspan="1" colspan="1">748</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5300 d54e5460" rowspan="1" colspan="1">1.54</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5300 d54e5466" rowspan="1" colspan="1">1.70</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.8543689320388355%" headers="d54e5286 d54e5306" rowspan="1" colspan="1">3</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5288 d54e5309" rowspan="1" colspan="1">3.89</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5288 d54e5312" rowspan="1" colspan="1">0.03</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5294 d54e5315" rowspan="1" colspan="1">9.02</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5294 d54e5318" rowspan="1" colspan="1">9.79E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5294 d54e5385" rowspan="1" colspan="1">147</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5297 d54e5388" rowspan="1" colspan="1">5.42</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5297 d54e5391" rowspan="1" colspan="1">9.78E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5297 d54e5457" rowspan="1" colspan="1">147</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5300 d54e5460" rowspan="1" colspan="1">1.38</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5300 d54e5466" rowspan="1" colspan="1">1.83</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.8543689320388355%" headers="d54e5286 d54e5306" rowspan="1" colspan="1">4</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5288 d54e5309" rowspan="1" colspan="1">1.09</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5288 d54e5312" rowspan="1" colspan="1">0.03</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5294 d54e5315" rowspan="1" colspan="1">34.5</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5294 d54e5318" rowspan="1" colspan="1">9.83E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5294 d54e5385" rowspan="1" colspan="1">454</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5297 d54e5388" rowspan="1" colspan="1">38.2</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5297 d54e5391" rowspan="1" colspan="1">9.83E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5297 d54e5457" rowspan="1" colspan="1">454</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5300 d54e5460" rowspan="1" colspan="1">0.91</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5300 d54e5466" rowspan="1" colspan="1">2.76</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.8543689320388355%" headers="d54e5286 d54e5306" rowspan="1" colspan="1">5</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5288 d54e5309" rowspan="1" colspan="1">3.25</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5288 d54e5312" rowspan="1" colspan="1">0.06</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5294 d54e5315" rowspan="1" colspan="1">26.3</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5294 d54e5318" rowspan="1" colspan="1">9.71E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5294 d54e5385" rowspan="1" colspan="1">272</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5297 d54e5388" rowspan="1" colspan="1">55.2</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5297 d54e5391" rowspan="1" colspan="1">9.71E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5297 d54e5457" rowspan="1" colspan="1">272</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5300 d54e5460" rowspan="1" colspan="1">0.51</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5300 d54e5466" rowspan="1" colspan="1">0.53</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.8543689320388355%" headers="d54e5286 d54e5306" rowspan="1" colspan="1">6</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5288 d54e5309" rowspan="1" colspan="1">11.0</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5288 d54e5312" rowspan="1" colspan="1">0.07</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5294 d54e5315" rowspan="1" colspan="1">44.7</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5294 d54e5318" rowspan="1" colspan="1">9.42E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5294 d54e5385" rowspan="1" colspan="1">263</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5297 d54e5388" rowspan="1" colspan="1">84.0</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5297 d54e5391" rowspan="1" colspan="1">9.44E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5297 d54e5457" rowspan="1" colspan="1">263</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5300 d54e5460" rowspan="1" colspan="1">0.59</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5300 d54e5466" rowspan="1" colspan="1">1.02</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.8543689320388355%" headers="d54e5286 d54e5306" rowspan="1" colspan="1">7</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5288 d54e5309" rowspan="1" colspan="1">5.95</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5288 d54e5312" rowspan="1" colspan="1">0.09</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5294 d54e5315" rowspan="1" colspan="1">8.84</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5294 d54e5318" rowspan="1" colspan="1">8.53E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5294 d54e5385" rowspan="1" colspan="1">43</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5297 d54e5388" rowspan="1" colspan="1">17.0</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5297 d54e5391" rowspan="1" colspan="1">8.53E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5297 d54e5457" rowspan="1" colspan="1">43</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5300 d54e5460" rowspan="1" colspan="1">0.64</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5300 d54e5466" rowspan="1" colspan="1">1.68</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.8543689320388355%" headers="d54e5286 d54e5306" rowspan="1" colspan="1">8</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5288 d54e5309" rowspan="1" colspan="1">2.94</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5288 d54e5312" rowspan="1" colspan="1">0.04</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5294 d54e5315" rowspan="1" colspan="1">0.09</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5294 d54e5318" rowspan="1" colspan="1">2.10E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5294 d54e5385" rowspan="1" colspan="1">1.5</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5297 d54e5388" rowspan="1" colspan="1">1.75</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5297 d54e5391" rowspan="1" colspan="1">2.10E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5297 d54e5457" rowspan="1" colspan="1">1.5</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5300 d54e5460" rowspan="1" colspan="1">0.64</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5300 d54e5466" rowspan="1" colspan="1">3.54</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.8543689320388355%" headers="d54e5286 d54e5306" rowspan="1" colspan="1">9</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5288 d54e5309" rowspan="1" colspan="1">0.11</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5288 d54e5312" rowspan="1" colspan="1">0.01</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5294 d54e5315" rowspan="1" colspan="1">53.2</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5294 d54e5318" rowspan="1" colspan="1">4.24E-03</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5294 d54e5385" rowspan="1" colspan="1">2000</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5297 d54e5388" rowspan="1" colspan="1">24.4</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5297 d54e5391" rowspan="1" colspan="1">4.92E-03</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5297 d54e5457" rowspan="1" colspan="1">2000</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5300 d54e5460" rowspan="1" colspan="1">2.18</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5300 d54e5466" rowspan="1" colspan="1">1.31</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.8543689320388355%" headers="d54e5286 d54e5306" rowspan="1" colspan="1">10</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5288 d54e5309" rowspan="1" colspan="1">0.12</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5288 d54e5312" rowspan="1" colspan="1">0.01</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5294 d54e5315" rowspan="1" colspan="1">0.16</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5294 d54e5318" rowspan="1" colspan="1">4.89E-11</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5294 d54e5385" rowspan="1" colspan="1">4</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5297 d54e5388" rowspan="1" colspan="1">0.08</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5297 d54e5391" rowspan="1" colspan="1">6.45E-11</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5297 d54e5457" rowspan="1" colspan="1">4</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5300 d54e5460" rowspan="1" colspan="1">1.36</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5300 d54e5466" rowspan="1" colspan="1">1.18</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.8543689320388355%" headers="d54e5286 d54e5306" rowspan="1" colspan="1">11</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5288 d54e5309" rowspan="1" colspan="1">2.89</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5288 d54e5312" rowspan="1" colspan="1">0.09</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5294 d54e5315" rowspan="1" colspan="1">0.44</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5294 d54e5318" rowspan="1" colspan="1">6.10E-09</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5294 d54e5385" rowspan="1" colspan="1">2.5</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5297 d54e5388" rowspan="1" colspan="1">0.48</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5297 d54e5391" rowspan="1" colspan="1">6.10E-09</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5297 d54e5457" rowspan="1" colspan="1">2.5</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5300 d54e5460" rowspan="1" colspan="1">1.00</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5300 d54e5466" rowspan="1" colspan="1">33.2</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.8543689320388355%" headers="d54e5286 d54e5306" rowspan="1" colspan="1">12</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5288 d54e5309" rowspan="1" colspan="1">0.36</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5288 d54e5312" rowspan="1" colspan="1">0.03</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5294 d54e5315" rowspan="1" colspan="1">36.6</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5294 d54e5318" rowspan="1" colspan="1">7.05E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5294 d54e5385" rowspan="1" colspan="1">278.5</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5297 d54e5388" rowspan="1" colspan="1">10.6</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5297 d54e5391" rowspan="1" colspan="1">8.82E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5297 d54e5457" rowspan="1" colspan="1">270.5</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5300 d54e5460" rowspan="1" colspan="1">3.35</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5300 d54e5466" rowspan="1" colspan="1">8.04</td>
</tr>
</tbody>
</table>
</div>
<div class="tablenoborder"><a name="numerical-experiments__csrilut-20-10-preconditioned-cg-and-bicgstab-methods" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="numerical-experiments__csrilut-20-10-preconditioned-cg-and-bicgstab-methods" class="table" frame="border" border="1" rules="all">
<caption><span class="tablecap">Table 5. <samp class="ph codeph">csrilut</samp>(20,10<sup class="ph sup">-7</sup>) Preconditioned CG
and BiCGStab Methods</span></caption>
<thead class="thead" align="left">
<tr class="row">
<th class="entry" valign="top" width="4.8543689320388355%" id="d54e5962" rowspan="1" colspan="1"> </th>
<th class="entry" colspan="2" align="center" valign="top" id="d54e5964" rowspan="1">ilut(20,10<sup class="ph sup">-7</sup>)
</th>
<th class="entry" colspan="3" align="center" valign="top" id="d54e5970" rowspan="1">CPU</th>
<th class="entry" colspan="3" align="center" valign="top" id="d54e5973" rowspan="1">GPU</th>
<th class="entry" colspan="2" align="center" valign="top" id="d54e5976" rowspan="1">Speedup</th>
</tr>
<tr class="row">
<th class="entry" valign="top" width="4.8543689320388355%" id="d54e5982" rowspan="1" colspan="1">#</th>
<th class="entry" valign="top" width="9.708737864077671%" id="d54e5985" rowspan="1" colspan="1">fact. time(s)</th>
<th class="entry" valign="top" width="9.708737864077671%" id="d54e5988" rowspan="1" colspan="1">copy time(s)</th>
<th class="entry" valign="top" width="9.708737864077671%" id="d54e5991" rowspan="1" colspan="1">solve time(s)</th>
<th class="entry" valign="top" width="9.708737864077671%" id="d54e5994" rowspan="1" colspan="1">
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mfrac>
<mrow>
<msub>
<mrow>
<mo>‖</mo>
<msub>
<mi mathvariant="bold">r</mi>
<mi>i</mi>
</msub>
<mo>‖</mo>
</mrow>
<mn>2</mn>
</msub>
</mrow>
<mrow>
<msub>
<mrow>
<mo>‖</mo>
<msub>
<mi mathvariant="bold">r</mi>
<mn>0</mn>
</msub>
<mo>‖</mo>
</mrow>
<mn>2</mn>
</msub>
</mrow>
</mfrac>
</mrow>
</math>
</th>
<th class="entry" valign="top" width="6.796116504854369%" id="d54e6061" rowspan="1" colspan="1"># it.</th>
<th class="entry" valign="top" width="9.708737864077671%" id="d54e6064" rowspan="1" colspan="1">solve time(s)</th>
<th class="entry" valign="top" width="9.708737864077671%" id="d54e6067" rowspan="1" colspan="1">
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mrow>
<mfrac>
<mrow>
<msub>
<mrow>
<mo>‖</mo>
<msub>
<mi mathvariant="bold">r</mi>
<mi>i</mi>
</msub>
<mo>‖</mo>
</mrow>
<mn>2</mn>
</msub>
</mrow>
<mrow>
<msub>
<mrow>
<mo>‖</mo>
<msub>
<mi mathvariant="bold">r</mi>
<mn>0</mn>
</msub>
<mo>‖</mo>
</mrow>
<mn>2</mn>
</msub>
</mrow>
</mfrac>
</mrow>
</math>
</th>
<th class="entry" valign="top" width="6.796116504854369%" id="d54e6133" rowspan="1" colspan="1"># it.</th>
<th class="entry" valign="top" width="11.650485436893204%" id="d54e6136" rowspan="1" colspan="1">vs. ilut (20,10<sup class="ph sup">-7</sup>)
</th>
<th class="entry" valign="top" width="11.650485436893204%" id="d54e6142" rowspan="1" colspan="1">vs. ilu0</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" width="4.8543689320388355%" headers="d54e5962 d54e5982" rowspan="1" colspan="1">1</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5964 d54e5985" rowspan="1" colspan="1">0.82</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5964 d54e5988" rowspan="1" colspan="1">0.02</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5970 d54e5991" rowspan="1" colspan="1">47.6</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5970 d54e5994" rowspan="1" colspan="1">9.90E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5970 d54e6061" rowspan="1" colspan="1">1297</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5973 d54e6064" rowspan="1" colspan="1">159.</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5973 d54e6067" rowspan="1" colspan="1">9.86E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5973 d54e6133" rowspan="1" colspan="1">1292</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5976 d54e6136" rowspan="1" colspan="1">0.30</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5976 d54e6142" rowspan="1" colspan="1">25.2</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.8543689320388355%" headers="d54e5962 d54e5982" rowspan="1" colspan="1">2</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5964 d54e5985" rowspan="1" colspan="1">9.21</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5964 d54e5988" rowspan="1" colspan="1">0.11</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5970 d54e5991" rowspan="1" colspan="1">32.1</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5970 d54e5994" rowspan="1" colspan="1">8.69E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5970 d54e6061" rowspan="1" colspan="1">193</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5973 d54e6064" rowspan="1" colspan="1">84.6</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5973 d54e6067" rowspan="1" colspan="1">8.67E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5973 d54e6133" rowspan="1" colspan="1">193</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5976 d54e6136" rowspan="1" colspan="1">0.44</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5976 d54e6142" rowspan="1" colspan="1">1.16</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.8543689320388355%" headers="d54e5962 d54e5982" rowspan="1" colspan="1">3</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5964 d54e5985" rowspan="1" colspan="1">10.04</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5964 d54e5988" rowspan="1" colspan="1">0.04</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5970 d54e5991" rowspan="1" colspan="1">6.26</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5970 d54e5994" rowspan="1" colspan="1">9.64E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5970 d54e6061" rowspan="1" colspan="1">90</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5973 d54e6064" rowspan="1" colspan="1">4.75</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5973 d54e6067" rowspan="1" colspan="1">9.64E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5973 d54e6133" rowspan="1" colspan="1">90</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5976 d54e6136" rowspan="1" colspan="1">1.10</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5976 d54e6142" rowspan="1" colspan="1">2.36</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.8543689320388355%" headers="d54e5962 d54e5982" rowspan="1" colspan="1">4</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5964 d54e5985" rowspan="1" colspan="1">8.12</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5964 d54e5988" rowspan="1" colspan="1">0.10</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5970 d54e5991" rowspan="1" colspan="1">15.7</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5970 d54e5994" rowspan="1" colspan="1">9.02E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5970 d54e6061" rowspan="1" colspan="1">148</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5973 d54e6064" rowspan="1" colspan="1">22.5</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5973 d54e6067" rowspan="1" colspan="1">9.02E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5973 d54e6133" rowspan="1" colspan="1">148</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5976 d54e6136" rowspan="1" colspan="1">0.78</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5976 d54e6142" rowspan="1" colspan="1">1.84</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.8543689320388355%" headers="d54e5962 d54e5982" rowspan="1" colspan="1">5</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5964 d54e5985" rowspan="1" colspan="1">8.60</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5964 d54e5988" rowspan="1" colspan="1">0.10</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5970 d54e5991" rowspan="1" colspan="1">21.2</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5970 d54e5994" rowspan="1" colspan="1">9.52E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5970 d54e6061" rowspan="1" colspan="1">158</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5973 d54e6064" rowspan="1" colspan="1">53.6</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5973 d54e6067" rowspan="1" colspan="1">9.52E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5973 d54e6133" rowspan="1" colspan="1">158</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5976 d54e6136" rowspan="1" colspan="1">0.48</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5976 d54e6142" rowspan="1" colspan="1">0.54</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.8543689320388355%" headers="d54e5962 d54e5982" rowspan="1" colspan="1">6</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5964 d54e5985" rowspan="1" colspan="1">35.2</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5964 d54e5988" rowspan="1" colspan="1">0.11</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5970 d54e5991" rowspan="1" colspan="1">29.2</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5970 d54e5994" rowspan="1" colspan="1">9.88E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5970 d54e6061" rowspan="1" colspan="1">162</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5973 d54e6064" rowspan="1" colspan="1">80.5</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5973 d54e6067" rowspan="1" colspan="1">9.88E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5973 d54e6133" rowspan="1" colspan="1">162</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5976 d54e6136" rowspan="1" colspan="1">0.56</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5976 d54e6142" rowspan="1" colspan="1">1.18</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.8543689320388355%" headers="d54e5962 d54e5982" rowspan="1" colspan="1">7</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5964 d54e5985" rowspan="1" colspan="1">23.1</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5964 d54e5988" rowspan="1" colspan="1">0.14</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5970 d54e5991" rowspan="1" colspan="1">3.79</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5970 d54e5994" rowspan="1" colspan="1">7.50E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5970 d54e6061" rowspan="1" colspan="1">14</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5973 d54e6064" rowspan="1" colspan="1">12.1</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5973 d54e6067" rowspan="1" colspan="1">7.50E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5973 d54e6133" rowspan="1" colspan="1">14</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5976 d54e6136" rowspan="1" colspan="1">0.76</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5976 d54e6142" rowspan="1" colspan="1">3.06</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.8543689320388355%" headers="d54e5962 d54e5982" rowspan="1" colspan="1">8</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5964 d54e5985" rowspan="1" colspan="1">5.23</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5964 d54e5988" rowspan="1" colspan="1">0.05</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5970 d54e5991" rowspan="1" colspan="1">0.14</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5970 d54e5994" rowspan="1" colspan="1">1.19E-09</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5970 d54e6061" rowspan="1" colspan="1">1.5</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5973 d54e6064" rowspan="1" colspan="1">2.37</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5973 d54e6067" rowspan="1" colspan="1">1.19E-09</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5973 d54e6133" rowspan="1" colspan="1">1.5</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5976 d54e6136" rowspan="1" colspan="1">0.70</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5976 d54e6142" rowspan="1" colspan="1">6.28</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.8543689320388355%" headers="d54e5962 d54e5982" rowspan="1" colspan="1">9</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5964 d54e5985" rowspan="1" colspan="1">0.12</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5964 d54e5988" rowspan="1" colspan="1">0.01</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5970 d54e5991" rowspan="1" colspan="1">55.1</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5970 d54e5994" rowspan="1" colspan="1">3.91E-03</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5970 d54e6061" rowspan="1" colspan="1">2000</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5973 d54e6064" rowspan="1" colspan="1">24.4</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5973 d54e6067" rowspan="1" colspan="1">2.27E-03</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5973 d54e6133" rowspan="1" colspan="1">2000</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5976 d54e6136" rowspan="1" colspan="1">2.25</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5976 d54e6142" rowspan="1" colspan="1">1.36</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.8543689320388355%" headers="d54e5962 d54e5982" rowspan="1" colspan="1">10</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5964 d54e5985" rowspan="1" colspan="1">0.14</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5964 d54e5988" rowspan="1" colspan="1">0.01</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5970 d54e5991" rowspan="1" colspan="1">0.14</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5970 d54e5994" rowspan="1" colspan="1">9.35E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5970 d54e6061" rowspan="1" colspan="1">3.5</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5973 d54e6064" rowspan="1" colspan="1">0.07</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5973 d54e6067" rowspan="1" colspan="1">7.19E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5973 d54e6133" rowspan="1" colspan="1">3.5</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5976 d54e6136" rowspan="1" colspan="1">1.28</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5976 d54e6142" rowspan="1" colspan="1">1.18</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.8543689320388355%" headers="d54e5962 d54e5982" rowspan="1" colspan="1">11</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5964 d54e5985" rowspan="1" colspan="1">218.</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5964 d54e5988" rowspan="1" colspan="1">0.12</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5970 d54e5991" rowspan="1" colspan="1">0.43</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5970 d54e5994" rowspan="1" colspan="1">9.80E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5970 d54e6061" rowspan="1" colspan="1">2</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5973 d54e6064" rowspan="1" colspan="1">0.66</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5973 d54e6067" rowspan="1" colspan="1">9.80E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5973 d54e6133" rowspan="1" colspan="1">2</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5976 d54e6136" rowspan="1" colspan="1">1.00</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5976 d54e6142" rowspan="1" colspan="1">247.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="4.8543689320388355%" headers="d54e5962 d54e5982" rowspan="1" colspan="1">12</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5964 d54e5985" rowspan="1" colspan="1">15.0</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5964 d54e5988" rowspan="1" colspan="1">0.21</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5970 d54e5991" rowspan="1" colspan="1">12.2</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5970 d54e5994" rowspan="1" colspan="1">3.45E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5970 d54e6061" rowspan="1" colspan="1">31</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5973 d54e6064" rowspan="1" colspan="1">4.95</td>
<td class="entry" valign="top" width="9.708737864077671%" headers="d54e5973 d54e6067" rowspan="1" colspan="1">3.45E-08</td>
<td class="entry" valign="top" width="6.796116504854369%" headers="d54e5973 d54e6133" rowspan="1" colspan="1">31</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5976 d54e6136" rowspan="1" colspan="1">1.35</td>
<td class="entry" valign="top" width="11.650485436893204%" headers="d54e5976 d54e6142" rowspan="1" colspan="1">5.93</td>
</tr>
</tbody>
</table>
</div>
<p class="p">The summary of performance of BiCGStab and CG iterative methods
preconditioned with different incomplete factorizations on the GPU is
shown in <a class="xref" href="index.html#numerical-experiments__performance-of-bicgstab-and-cg-with-incomplete-lu-cholesky" shape="rect">Figure 3</a>, where "*" indicates that the method did not converge
to the required tolerance. Notice that in general in our numerical
experiments the performance for the incomplete factorizations decreases
as the threshold parameters are relaxed and the factorization becomes
more dense, thus inhibiting parallelism due to data dependencies between
rows in the sparse triangular solve. For this reason, the best
performance on the GPU is obtained for the incomplete-LU and Cholesky
factorization with 0 fill-in, which will be our point of reference.
</p>
<div class="fig fignone" id="numerical-experiments__performance-of-bicgstab-and-cg-with-incomplete-lu-cholesky"><a name="numerical-experiments__performance-of-bicgstab-and-cg-with-incomplete-lu-cholesky" shape="rect">
<!-- --></a><span class="figcap">Figure 3. Performance of BiCGStab and CG with Incomplete-LU Cholesky
Preconditioning</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/performance-of-bicgstab-and-cg.png" alt="Figure of the performance of BiCGStab and CG with incomplete-LU Cholesky preconditioning."></img></div><br clear="none"></br></div>
<p class="p">Although the incomplete factorizations with a more relaxed threshold are
often closer to the exact factorization and thus result in fewer
iterative steps, they are also much more expensive to compute. Moreover,
notice that even though the number of iterative steps decreases, each
step is more computationally expensive. As a result of these tradeoffs
the total time, the sum of the time taken by the factorization and the
iterative solve, for the iterative method does not necessarily decrease
with a more relaxed threshold in our numerical experiments.
</p>
<p class="p">The speedup based on the total time taken by the preconditioned
iterative method on the GPU with <samp class="ph codeph">csrilu0</samp> preconditioner
and CPU with all four preconditioners is shown in <a class="xref" href="index.html#numerical-experiments__speedup-of-prec-bicgstab-and-cg-on-gpu-with-csrilu0-vs-cpu-with-all" shape="rect">Figure 4</a>. Notice that for majority of matrices in our numerical
experiments the implementation of the iterative method using the CUSPARSE
and CUBLAS libraries does indeed outperform the MKL.
</p>
<div class="fig fignone" id="numerical-experiments__speedup-of-prec-bicgstab-and-cg-on-gpu-with-csrilu0-vs-cpu-with-all"><a name="numerical-experiments__speedup-of-prec-bicgstab-and-cg-on-gpu-with-csrilu0-vs-cpu-with-all" shape="rect">
<!-- --></a><span class="figcap">Figure 4. Speedup of prec. BiCGStab and CG on GPU (with
<samp class="ph codeph">csrilu0</samp>) vs. CPU (with all)</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/speedup-of-prec-bicgstab-and-cg.png" alt="Figure showing speedup of prec. BiCGStab and CG on GPU with (csrilu0) vs CPU (with all)."></img></div><br clear="none"></br></div>
<p class="p">Finally, the average of the obtained speedups is shown in <a class="xref" href="index.html#numerical-experiments__average-speedup-of-bicgstab-and-cg-on-gpu-with-csrilu0-and-cpu-with-all" shape="rect">Figure 5</a>, where we have excluded the runs with cage13 matrix for
<samp class="ph codeph">ilut</samp>(10,10<sup class="ph sup">-5</sup>) and runs with offshore and
cage13 matrices for <samp class="ph codeph">ilut</samp>(20,10<sup class="ph sup">-7</sup>) incomplete
factorizations because of their disproportional speedup. However, the
speedup including these runs is shown in parenthesis on the same plot.
Consequently, we can conclude that the incomplete-LU and Cholesky
preconditioned BiCGStab and CG iterative methods obtain on average more
than 2x speedup on the GPU over their CPU implementation.
</p>
<div class="fig fignone" id="numerical-experiments__average-speedup-of-bicgstab-and-cg-on-gpu-with-csrilu0-and-cpu-with-all"><a name="numerical-experiments__average-speedup-of-bicgstab-and-cg-on-gpu-with-csrilu0-and-cpu-with-all" shape="rect">
<!-- --></a><span class="figcap">Figure 5. Average Speedup of BiCGStab and CG on GPU (with
<samp class="ph codeph">csrilu0</samp>) and CPU (with all)</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/average-speedup-of-bicgstab-and-cg.png" alt="Figure showing the average speedup of BiCGStab and CG on GPU (with csrilu0) and CPU (with all)."></img></div><br clear="none"></br></div>
</div>
</div>
<div class="topic concept nested0" id="conclusion"><a name="conclusion" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#conclusion" name="conclusion" shape="rect">4. Conclusion</a></h2>
<div class="body conbody">
<p class="p">The performance of the iterative methods depends highly on the sparsity
pattern of the coefficient matrix at hand. In our numerical experiments
the incomplete-LU and Cholesky preconditioned BiCGStab and CG iterative
methods implemented on the GPU using the CUSPARSE and CUBLAS libraries
achieved an average of 2x speedup over their MKL implementation.
</p>
<p class="p">The sparse matrix-vector multiplication and triangular solve, which is
split into a slower <dfn class="term">analysis</dfn> phase that needs to be
performed only once and a faster <dfn class="term">solve</dfn> phase that can be
performed multiple times, were the essential building blocks of these
iterative methods. In fact the obtained speedup was usually mostly
influenced by the time taken by the <dfn class="term">solve</dfn> phase of the
algorithm.
</p>
<p class="p">Finally, we point out that the use of multiple-right-hand-sides would
increase the available parallelism and can result in a significant
relative performance improvement in the preconditioned iterative methods.
Also, the development of incomplete-LU and Cholesky factorizations using
CUDA parallel programming paradigm can further improve the obtained
speedup.
</p>
</div>
</div>
<div class="topic reference nested0" id="acknowledgements"><a name="acknowledgements" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#acknowledgements" name="acknowledgements" shape="rect">A. Acknowledgements</a></h2>
<div class="body refbody">
<div class="section">
<p class="p">This white paper was authored by Maxim Naumov for NVIDIA Corporation.</p>
<p class="p">Permission to make digital or hard copies of all or part of this work for any use is
granted without fee provided that copies bear this notice and the full citation on the first
page.
</p>
</div>
</div>
</div>
<div class="topic reference nested0" id="references"><a name="references" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#references" name="references" shape="rect">B. References</a></h2>
<div class="body refbody">
<div class="section" id="references__1"><a name="references__1" shape="rect">
<!-- --></a><p class="p">[1] <cite class="cite">E. Anderson and Y. Saad Solving Sparse Triangular Linear
Systems on Parallel Computers, Int. J. High Speed Comput., pp. 73-95,
1989.</cite></p>
</div>
<div class="section" id="references__2"><a name="references__2" shape="rect">
<!-- --></a><p class="p">[2] <cite class="cite">R. Barrett, M. Berry, T. F. Chan, J. Demmel, J. Donato, J.
Dongarra, V. Eijkhout, R. Pozo, C. Romine, H. van der Vorst,
Templates for the Solution of Linear Systems: Building Blocks for
Iterative Methods, SIAM, Philadelphia, PA, 1994.</cite></p>
</div>
<div class="section" id="references__3"><a name="references__3" shape="rect">
<!-- --></a><p class="p">[3] <cite class="cite">N. Bell and M. Garland, Implementing Sparse Matrix-Vector
Multiplication on Throughput-Oriented Processors, Proc. Conf. HPC
Networking, Storage and Analysis (SC09), ACM, pp. 1-11,
2009.</cite></p>
</div>
<div class="section" id="references__4"><a name="references__4" shape="rect">
<!-- --></a><p class="p">[4] <cite class="cite">A. Greenbaum, Solving Sparse Triangular Linear Systems using
Fortran with Parallel Extensions on the NYU Ultracomputer Prototype,
Report 99, NYU Ultracomputer Note, New York University, NY, April,
1986.</cite></p>
</div>
<div class="section" id="references__5"><a name="references__5" shape="rect">
<!-- --></a><p class="p">[5] <cite class="cite">D. B. Kirk and W. W. Hwu, Programming Massively Parallel
Processors: A Hands-on Approach, Elsevier, 2010.</cite></p>
</div>
<div class="section" id="references__6"><a name="references__6" shape="rect">
<!-- --></a><p class="p">[6] <cite class="cite">J. Mayer, Parallel Algorithms for Solving Linear Systems
with Sparse Triangular Matrices, Computing, pp. 291-312 (86),
2009.</cite></p>
</div>
<div class="section" id="references__7"><a name="references__7" shape="rect">
<!-- --></a><p class="p">[7] <cite class="cite">R. Mirchandaney, J. H. Saltz and D. Baxter, Run-Time
Parallelization and Scheduling of Loops, IEEE Transactions on
Computers, pp. (40), 1991.</cite></p>
</div>
<div class="section" id="references__8"><a name="references__8" shape="rect">
<!-- --></a><p class="p">[8] <cite class="cite">M. Naumov, Parallel Solution of Sparse Triangular Linear
Systems in the Preconditioned Iterative Methods on the GPU, NVIDIA
Technical Report, NVR-2011-001, 2011.</cite></p>
</div>
<div class="section" id="references__9"><a name="references__9" shape="rect">
<!-- --></a><p class="p">[9] <cite class="cite">J. Nickolls, I. Buck, M. Garland and K. Skadron, Scalable
Parallel Programming with CUDA, Queue, pp. 40-53 (6-2),
2008.</cite></p>
</div>
<div class="section" id="references__10"><a name="references__10" shape="rect">
<!-- --></a><p class="p">[10] <cite class="cite">E. Rothberg and A. Gupta, Parallel ICCG on a Hierarchical
Memory Multiprocessor - Addressing the Triangular Solve Bottleneck,
Parallel Comput., pp. 719-741 (18), 1992.</cite></p>
</div>
<div class="section" id="references__11"><a name="references__11" shape="rect">
<!-- --></a><p class="p">[11] <cite class="cite">Y. Saad, Iterative Methods for Sparse Linear Systems, SIAM,
Philadelphia, PA, 2nd Ed., 2003.</cite></p>
</div>
<div class="section" id="references__12"><a name="references__12" shape="rect">
<!-- --></a><p class="p">[12] <cite class="cite">J. H. Saltz, Aggregation Methods for Solving Sparse
Triangular Systems on Multiprocessors, SIAM J. Sci. Statist. Comput.,
pp. 123-144 (11), 1990.</cite></p>
</div>
<div class="section" id="references__13"><a name="references__13" shape="rect">
<!-- --></a><p class="p">[13] <cite class="cite">J. Sanders and E. Kandrot, CUDA by Example: An Introduction
to General-Purpose GPU Programming, Addison-Wesley, 2010.</cite></p>
</div>
<div class="section" id="references__14"><a name="references__14" shape="rect">
<!-- --></a><p class="p">[14] <cite class="cite">M. Wolf, M. Heroux and E. Boman, Factors Impacting
Performance of Multithreaded Sparse Triangular Solve, 9th Int. Meet.
HPC Comput. Sci. (VECPAR), 2010.</cite></p>
</div>
<div class="section" id="references__15"><a name="references__15" shape="rect">
<!-- --></a><p class="p">[15] <cite class="cite">S. Williams, L. Oliker, R. Vuduc, J. Shalf, K. Yelick and
J. Demmel, Optimization of Sparse Matrix-Vector Multiplication on
Emerging Multicore Platforms, Parallel Comput., pp. 178-194 (35-3),
2009.</cite></p>
</div>
<div class="section" id="references__16"><a name="references__16" shape="rect">
<!-- --></a><p class="p">[16] <cite class="cite">NVIDIA CUSPARSE and CUBLAS Libraries,</cite><a class="xref" href="http://www.nvidia.com/object/cuda_develop.html" target="_blank" shape="rect">http://www.nvidia.com/object/cuda_develop.html</a></p>
</div>
<div class="section" id="references__17"><a name="references__17" shape="rect">
<!-- --></a><p class="p">[17] <cite class="cite">Intel Math Kernel Library,</cite><a class="xref" href="http://software.intel.com/en-us/articles/intel-mkl" target="_blank" shape="rect">http://software.intel.com/en-us/articles/intel-mkl</a></p>
</div>
<div class="section" id="references__18"><a name="references__18" shape="rect">
<!-- --></a><p class="p">[18] <cite class="cite">The University of Florida Sparse Matrix Collection, </cite><a class="xref" href="http://www.cise.ufl.edu/research/sparse/matrices/" target="_blank" shape="rect">http://www.cise.ufl.edu/research/sparse/matrices/</a>.
</p>
</div>
</div>
</div>
<div class="topic concept nested0" id="notices-header"><a name="notices-header" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#notices-header" name="notices-header" shape="rect">Notices</a></h2>
<div class="topic reference nested1" id="notice"><a name="notice" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#notice" name="notice" shape="rect"></a></h3>
<div class="body refbody">
<div class="section">
<h3 class="title sectiontitle">Notice</h3>
<p class="p">ALL NVIDIA DESIGN SPECIFICATIONS, REFERENCE BOARDS, FILES, DRAWINGS, DIAGNOSTICS, LISTS, AND OTHER DOCUMENTS (TOGETHER AND
SEPARATELY, "MATERIALS") ARE BEING PROVIDED "AS IS." NVIDIA MAKES NO WARRANTIES, EXPRESSED, IMPLIED, STATUTORY, OR OTHERWISE
WITH RESPECT TO THE MATERIALS, AND EXPRESSLY DISCLAIMS ALL IMPLIED WARRANTIES OF NONINFRINGEMENT, MERCHANTABILITY, AND FITNESS
FOR A PARTICULAR PURPOSE.
</p>
<p class="p">Information furnished is believed to be accurate and reliable. However, NVIDIA Corporation assumes no responsibility for the
consequences of use of such information or for any infringement of patents or other rights of third parties that may result
from its use. No license is granted by implication of otherwise under any patent rights of NVIDIA Corporation. Specifications
mentioned in this publication are subject to change without notice. This publication supersedes and replaces all other information
previously supplied. NVIDIA Corporation products are not authorized as critical components in life support devices or systems
without express written approval of NVIDIA Corporation.
</p>
</div>
</div>
</div>
<div class="topic reference nested1" id="trademarks"><a name="trademarks" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#trademarks" name="trademarks" shape="rect"></a></h3>
<div class="body refbody">
<div class="section">
<h3 class="title sectiontitle">Trademarks</h3>
<p class="p">NVIDIA and the NVIDIA logo are trademarks or registered trademarks of NVIDIA Corporation
in the U.S. and other countries. Other company and product names may be trademarks of
the respective companies with which they are associated.
</p>
</div>
</div>
</div>
<div class="topic reference nested1" id="copyright-past-to-present"><a name="copyright-past-to-present" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#copyright-past-to-present" name="copyright-past-to-present" shape="rect"></a></h3>
<div class="body refbody">
<div class="section">
<h3 class="title sectiontitle">Copyright</h3>
<p class="p">© <span class="ph">2011</span>-<span class="ph">2013</span> NVIDIA
Corporation. All rights reserved.
</p>
</div>
</div>
</div>
</div>
<hr id="contents-end"></hr>
<div id="release-info">Incomplete-LU and Cholesky Preconditioned Iterative Methods
(<a href="../../pdf/Incomplete_LU_Cholesky.pdf">PDF</a>)
-
CUDA Toolkit v5.5
(<a href="https://developer.nvidia.com/cuda-toolkit-archive">older</a>)
-
Last updated
July 19, 2013
-
<a href="mailto:cudatools@nvidia.com?subject=CUDA Tools Documentation Feedback: incomplete-lu-cholesky">Send Feedback</a></div>
</article>
<header id="header"><span id="company">NVIDIA</span><span id="site-title">CUDA Toolkit Documentation</span><form id="search" method="get" action="search">
<input type="text" name="search-text"></input><fieldset id="search-location">
<legend>Search In:</legend>
<label><input type="radio" name="search-type" value="site"></input>Entire Site</label>
<label><input type="radio" name="search-type" value="document"></input>Just This Document</label></fieldset>
<button type="reset">clear search</button>
<button id="submit" type="submit">search</button></form>
</header>
<nav id="site-nav">
<div class="category closed"><span class="twiddle">▷</span><a href="../index.html" title="The root of the site.">CUDA Toolkit</a></div>
<ul class="closed">
<li><a href="../cuda-toolkit-release-notes/index.html" title="The Release Notes for the CUDA Toolkit from v4.0 to today.">Release Notes</a></li>
<li><a href="../eula/index.html" title="The End User License Agreements for the NVIDIA CUDA Toolkit, the NVIDIA CUDA Samples, the NVIDIA Display Driver, and NVIDIA NSight (Visual Studio Edition).">EULA</a></li>
<li><a href="../cuda-getting-started-guide-for-linux/index.html" title="This guide discusses how to install and check for correct operation of the CUDA Development Tools on GNU/Linux systems.">Getting Started Linux</a></li>
<li><a href="../cuda-getting-started-guide-for-mac-os-x/index.html" title="This guide discusses how to install and check for correct operation of the CUDA Development Tools on Mac OS X systems.">Getting Started Mac OS X</a></li>
<li><a href="../cuda-getting-started-guide-for-microsoft-windows/index.html" title="This guide discusses how to install and check for correct operation of the CUDA Development Tools on Microsoft Windows systems.">Getting Started Windows</a></li>
<li><a href="../cuda-c-programming-guide/index.html" title="This guide provides a detailed discussion of the CUDA programming model and programming interface. It then describes the hardware implementation, and provides guidance on how to achieve maximum performance. The Appendixes include a list of all CUDA-enabled devices, detailed description of all extensions to the C language, listings of supported mathematical functions, C++ features supported in host and device code, details on texture fetching, technical specifications of various devices, and concludes by introducing the low-level driver API.">Programming Guide</a></li>
<li><a href="../cuda-c-best-practices-guide/index.html" title="This guide presents established parallelization and optimization techniques and explains coding metaphors and idioms that can greatly simplify programming for CUDA-capable GPU architectures. The intent is to provide guidelines for obtaining the best performance from NVIDIA GPUs using the CUDA Toolkit.">Best Practices Guide</a></li>
<li><a href="../kepler-compatibility-guide/index.html" title="This application note is intended to help developers ensure that their NVIDIA CUDA applications will run effectively on GPUs based on the NVIDIA Kepler Architecture. This document provides guidance to ensure that your software applications are compatible with Kepler.">Kepler Compatibility Guide</a></li>
<li><a href="../kepler-tuning-guide/index.html" title="Kepler is NVIDIA's next-generation architecture for CUDA compute applications. Applications that follow the best practices for the Fermi architecture should typically see speedups on the Kepler architecture without any code changes. This guide summarizes the ways that an application can be fine-tuned to gain additional speedups by leveraging Kepler architectural features.">Kepler Tuning Guide</a></li>
<li><a href="../parallel-thread-execution/index.html" title="This guide provides detailed instructions on the use of PTX, a low-level parallel thread execution virtual machine and instruction set architecture (ISA). PTX exposes the GPU as a data-parallel computing device.">PTX ISA</a></li>
<li><a href="../optimus-developer-guide/index.html" title="This document explains how CUDA APIs can be used to query for GPU capabilities in NVIDIA Optimus systems.">Developer Guide for Optimus</a></li>
<li><a href="../video-decoder/index.html" title="This document provides the video decoder API specification and the format conversion and display using DirectX or OpenGL following decode.">Video Decoder</a></li>
<li><a href="../video-encoder/index.html" title="This document provides the CUDA video encoder specifications, including the C-library API functions and encoder query parameters.">Video Encoder</a></li>
<li><a href="../inline-ptx-assembly/index.html" title="This document shows how to inline PTX (parallel thread execution) assembly language statements into CUDA code. It describes available assembler statement parameters and constraints, and the document also provides a list of some pitfalls that you may encounter.">Inline PTX Assembly</a></li>
<li><a href="../cuda-runtime-api/index.html" title="The CUDA runtime API.">CUDA Runtime API</a></li>
<li><a href="../cuda-driver-api/index.html" title="The CUDA driver API.">CUDA Driver API</a></li>
<li><a href="../cuda-math-api/index.html" title="The CUDA math API.">CUDA Math API</a></li>
<li><a href="../cublas/index.html" title="The CUBLAS library is an implementation of BLAS (Basic Linear Algebra Subprograms) on top of the NVIDIA CUDA runtime. It allows the user to access the computational resources of NVIDIA Graphical Processing Unit (GPU), but does not auto-parallelize across multiple GPUs.">CUBLAS</a></li>
<li><a href="../cufft/index.html" title="The CUFFT library user guide.">CUFFT</a></li>
<li><a href="../curand/index.html" title="The CURAND library user guide.">CURAND</a></li>
<li><a href="../cusparse/index.html" title="The CUSPARSE library user guide.">CUSPARSE</a></li>
<li><a href="../npp/index.html" title="NVIDIA NPP is a library of functions for performing CUDA accelerated processing. The initial set of functionality in the library focuses on imaging and video processing and is widely applicable for developers in these areas. NPP will evolve over time to encompass more of the compute heavy tasks in a variety of problem domains. The NPP library is written to maximize flexibility, while maintaining high performance.">NPP</a></li>
<li><a href="../thrust/index.html" title="The Thrust getting started guide.">Thrust</a></li>
<li><a href="../cuda-samples/index.html" title="This document contains a complete listing of the code samples that are included with the NVIDIA CUDA Toolkit. It describes each code sample, lists the minimum GPU specification, and provides links to the source code and white papers if available.">CUDA Samples</a></li>
<li><a href="../cuda-compiler-driver-nvcc/index.html" title="This document is a reference guide on the use of the CUDA compiler driver nvcc. Instead of being a specific CUDA compilation driver, nvcc mimics the behavior of the GNU compiler gcc, accepting a range of conventional compiler options, such as for defining macros and include/library paths, and for steering the compilation process.">NVCC</a></li>
<li><a href="../cuda-gdb/index.html" title="The NVIDIA tool for debugging CUDA applications running on Linux and Mac, providing developers with a mechanism for debugging CUDA applications running on actual hardware. CUDA-GDB is an extension to the x86-64 port of GDB, the GNU Project debugger.">CUDA-GDB</a></li>
<li><a href="../cuda-memcheck/index.html" title="CUDA-MEMCHECK is a suite of run time tools capable of precisely detecting out of bounds and misaligned memory access errors, checking device allocation leaks, reporting hardware errors and identifying shared memory data access hazards.">CUDA-MEMCHECK</a></li>
<li><a href="../nsight-eclipse-edition-getting-started-guide/index.html" title="Nsight Eclipse Edition getting started guide">Nsight Eclipse Edition</a></li>
<li><a href="../profiler-users-guide/index.html" title="This is the guide to the Profiler.">Profiler</a></li>
<li><a href="../cuda-binary-utilities/index.html" title="The application notes for cuobjdump and nvdisasm.">CUDA Binary Utilities</a></li>
<li><a href="../floating-point/index.html" title="A number of issues related to floating point accuracy and compliance are a frequent source of confusion on both CPUs and GPUs. The purpose of this white paper is to discuss the most common issues related to NVIDIA GPUs and to supplement the documentation in the CUDA C Programming Guide.">Floating Point and IEEE 754</a></li>
<li><a href="../incomplete-lu-cholesky/index.html" title="In this white paper we show how to use the CUSPARSE and CUBLAS libraries to achieve a 2x speedup over CPU in the incomplete-LU and Cholesky preconditioned iterative methods. We focus on the Bi-Conjugate Gradient Stabilized and Conjugate Gradient iterative methods, that can be used to solve large sparse nonsymmetric and symmetric positive definite linear systems, respectively. Also, we comment on the parallel sparse triangular solve, which is an essential building block in these algorithms.">Incomplete-LU and Cholesky Preconditioned Iterative Methods</a></li>
<li><a href="../libnvvm-api/index.html" title="The libNVVM API.">libNVVM API</a></li>
<li><a href="../libdevice-users-guide/index.html" title="The libdevice library is an LLVM bitcode library that implements common functions for GPU kernels.">libdevice User's Guide</a></li>
<li><a href="../nvvm-ir-spec/index.html" title="NVVM IR is a compiler IR (internal representation) based on the LLVM IR. The NVVM IR is designed to represent GPU compute kernels (for example, CUDA kernels). High-level language front-ends, like the CUDA C compiler front-end, can generate NVVM IR.">NVVM IR</a></li>
<li><a href="../cupti/index.html" title="The CUPTI API.">CUPTI</a></li>
<li><a href="../debugger-api/index.html" title="The CUDA debugger API.">Debugger API</a></li>
<li><a href="../gpudirect-rdma/index.html" title="A tool for Kepler-class GPUs and CUDA 5.0 enabling a direct path for communication between the GPU and a peer device on the PCI Express bus when the devices share the same upstream root complex using standard features of PCI Express. This document introduces the technology and describes the steps necessary to enable a RDMA for GPUDirect connection to NVIDIA GPUs within the Linux device driver model.">RDMA for GPUDirect</a></li>
</ul>
<div class="category"><span class="twiddle">▼</span><a href="index.html" title="Incomplete-LU and Cholesky Preconditioned Iterative Methods">Incomplete-LU and Cholesky Preconditioned Iterative Methods</a></div>
<ul>
<li><a href="#introduction">1. Introduction</a></li>
<li><a href="#preconditioned-iterative-methods">2. Preconditioned Iterative Methods</a></li>
<li><a href="#numerical-experiments">3. Numerical Experiments</a></li>
<li><a href="#conclusion">4. Conclusion</a></li>
<li><a href="#acknowledgements">A. Acknowledgements</a></li>
<li><a href="#references">B. References</a></li>
</ul>
</nav>
<nav id="search-results">
<h2>Search Results</h2>
<ol></ol>
</nav>
<script language="JavaScript" type="text/javascript" charset="utf-8" src="../common/formatting/common.min.js"></script>
<script language="JavaScript" type="text/javascript" charset="utf-8" src="../common/scripts/omniture/s_code_us_dev_aut1-nolinktrackin.js"></script>
<script language="JavaScript" type="text/javascript" charset="utf-8" src="../common/scripts/omniture/omniture.js"></script>
<noscript><a href="http://www.omniture.com" title="Web Analytics"><img src="http://omniture.nvidia.com/b/ss/nvidiacudadocs/1/H.17--NS/0" height="1" width="1" border="0" alt=""></img></a></noscript>
<script language="JavaScript" type="text/javascript" charset="utf-8" src="../common/scripts/google-analytics/google-analytics-write.js"></script>
<script language="JavaScript" type="text/javascript" charset="utf-8" src="../common/scripts/google-analytics/google-analytics-tracker.js"></script>
</body>
</html>
