extern "C" magma_int_t
magma_dsgesv_gpu(char trans, magma_int_t n, magma_int_t nrhs,
double *dA, magma_int_t ldda,
magma_int_t *ipiv, magma_int_t *dipiv,
double *dB, magma_int_t lddb,
double *dX, magma_int_t lddx,
double *dworkd, float *dworks,
magma_int_t *iter, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
DSGESV computes the solution to a real system of linear equations
A * X = B or A' * X = B
where A is an N-by-N matrix and X and B are N-by-NRHS matrices.
DSGESV first attempts to factorize the matrix in real SINGLE PRECISION
and use this factorization within an iterative refinement procedure
to produce a solution with real DOUBLE PRECISION norm-wise backward error
quality (see below). If the approach fails the method switches to a
real DOUBLE PRECISION factorization and solve.
The iterative refinement is not going to be a winning strategy if
the ratio real SINGLE PRECISION performance over real DOUBLE PRECISION
performance is too small. A reasonable strategy should take the
number of right-hand sides and the size of the matrix into account.
This might be done with a call to ILAENV in the future. Up to now, we
always try iterative refinement.
The iterative refinement process is stopped if
ITER > ITERMAX
or for all the RHS we have:
RNRM < SQRT(N)*XNRM*ANRM*EPS*BWDMAX
where
o ITER is the number of the current iteration in the iterative
refinement process
o RNRM is the infinity-norm of the residual
o XNRM is the infinity-norm of the solution
o ANRM is the infinity-operator-norm of the matrix A
o EPS is the machine epsilon returned by DLAMCH('Epsilon')
The value ITERMAX and BWDMAX are fixed to 30 and 1.0D+00 respectively.
Arguments
=========
TRANS (input) CHARACTER*1
Specifies the form of the system of equations:
= 'N': A * X = B (No transpose)
= 'T': A'* X = B (Transpose)
= 'C': A'* X = B (Conjugate transpose = Transpose)
N (input) INTEGER
The number of linear equations, i.e., the order of the
matrix A. N >= 0.
NRHS (input) INTEGER
The number of right hand sides, i.e., the number of columns
of the matrix B. NRHS >= 0.
dA (input or input/output) DOUBLE PRECISION array on the GPU, dimension (ldda,N)
On entry, the N-by-N coefficient matrix A.
On exit, if iterative refinement has been successfully used
(info.EQ.0 and ITER.GE.0, see description below), A is
unchanged. If double precision factorization has been used
(info.EQ.0 and ITER.LT.0, see description below), then the
array dA contains the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
ldda (input) INTEGER
The leading dimension of the array dA. ldda >= max(1,N).
IPIV (output) INTEGER array, dimension (N)
The pivot indices that define the permutation matrix P;
row i of the matrix was interchanged with row IPIV(i).
Corresponds either to the single precision factorization
(if info.EQ.0 and ITER.GE.0) or the double precision
factorization (if info.EQ.0 and ITER.LT.0).
dIPIV (output) INTEGER array on the GPU, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was moved to row IPIV(i).
dB (input) DOUBLE PRECISION array on the GPU, dimension (lddb,NRHS)
The N-by-NRHS right hand side matrix B.
lddb (input) INTEGER
The leading dimension of the array dB. lddb >= max(1,N).
dX (output) DOUBLE PRECISION array on the GPU, dimension (lddx,NRHS)
If info = 0, the N-by-NRHS solution matrix X.
lddx (input) INTEGER
The leading dimension of the array dX. lddx >= max(1,N).
dworkd (workspace) DOUBLE PRECISION array on the GPU, dimension (N*NRHS)
This array is used to hold the residual vectors.
dworks (workspace) SINGLE PRECISION array on the GPU, dimension (N*(N+NRHS))
This array is used to store the real single precision matrix
and the right-hand sides or solutions in single precision.
iter (output) INTEGER
< 0: iterative refinement has failed, double precision
factorization has been performed
-1 : the routine fell back to full precision for
implementation- or machine-specific reasons
-2 : narrowing the precision induced an overflow,
the routine fell back to full precision
-3 : failure of SGETRF
-31: stop the iterative refinement after the 30th iteration
> 0: iterative refinement has been successfully used.
Returns the number of iterations
info (output) INTEGER
= 0: successful exit
< 0: if info = -i, the i-th argument had an illegal value
> 0: if info = i, U(i,i) computed in DOUBLE PRECISION is
exactly zero. The factorization has been completed,
but the factor U is exactly singular, so the solution
could not be computed.
===================================================================== */
#define dB(i,j) (dB + (i) + (j)*lddb)
#define dX(i,j) (dX + (i) + (j)*lddx)
#define dR(i,j) (dR + (i) + (j)*lddr)
double c_neg_one = MAGMA_D_NEG_ONE;
double c_one = MAGMA_D_ONE;
magma_int_t ione = 1;
double *dR;
float *dSA, *dSX;
double Xnrmv, Rnrmv;
double Anrm, Xnrm, Rnrm, cte, eps;
magma_int_t i, j, iiter, lddsa, lddr;
/* Check arguments */
*iter = 0;
*info = 0;
if ( n < 0 )
*info = -1;
else if ( nrhs < 0 )
*info = -2;
else if ( ldda < max(1,n))
*info = -4;
else if ( lddb < max(1,n))
*info = -8;
else if ( lddx < max(1,n))
*info = -10;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if ( n == 0 || nrhs == 0 )
return *info;
lddsa = n;
lddr = n;
dSA = dworks;
dSX = dSA + lddsa*n;
dR = dworkd;
eps = lapackf77_dlamch("Epsilon");
Anrm = magmablas_dlange('I', n, n, dA, ldda, (double*)dworkd );
cte = Anrm * eps * pow((double)n, 0.5) * BWDMAX;
/*
* Convert to single precision
*/
//magmablas_dlag2s( n, nrhs, dB, lddb, dSX, lddsx, info ); // done inside dsgetrs with pivots
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
magmablas_dlag2s( n, n, dA, ldda, dSA, lddsa, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
// factor dSA in single precision
magma_sgetrf_gpu( n, n, dSA, lddsa, ipiv, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// Generate parallel pivots
{
magma_int_t *newipiv;
magma_imalloc_cpu( &newipiv, n );
if ( newipiv == NULL ) {
*iter = -3;
goto FALLBACK;
}
swp2pswp( trans, n, ipiv, newipiv );
magma_setvector( n, sizeof(magma_int_t), newipiv, 1, dipiv, 1 );
magma_free_cpu( newipiv );
}
// solve dSA*dSX = dB in single precision
// converts dB to dSX and applies pivots, solves, then converts result back to dX
magma_dsgetrs_gpu( trans, n, nrhs, dSA, lddsa, dipiv, dB, lddb, dX, lddx, dSX, info );
// residual dR = dB - dA*dX in double precision
magmablas_dlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dR, lddr );
if ( nrhs == 1 ) {
magma_dgemv( trans, n, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_dgemm( trans, MagmaNoTrans, n, nrhs, n,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
// TODO: use MAGMA_D_ABS( dX(i,j) ) instead of dlange?
for( j=0; j < nrhs; j++ ) {
i = magma_idamax( n, dX(0,j), 1) - 1;
magma_dgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_dlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_idamax ( n, dR(0,j), 1 ) - 1;
magma_dgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_dlange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
if ( Rnrm > Xnrm*cte ) {
goto REFINEMENT;
}
}
*iter = 0;
return *info;
REFINEMENT:
for( iiter=1; iiter < ITERMAX; ) {
*info = 0;
// convert residual dR to single precision dSX
// solve dSA*dSX = R in single precision
// convert result back to double precision dR
// it's okay that dR is used for both dB input and dX output.
magma_dsgetrs_gpu( trans, n, nrhs, dSA, lddsa, dipiv, dR, lddr, dR, lddr, dSX, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// Add correction and setup residual
// dX += dR --and--
// dR = dB
// This saves going through dR a second time (if done with one more kernel).
// -- not really: first time is read, second time is write.
for( j=0; j < nrhs; j++ ) {
magmablas_daxpycp( n, dR(0,j), dX(0,j), dB(0,j) );
}
// residual dR = dB - dA*dX in double precision
if ( nrhs == 1 ) {
magma_dgemv( trans, n, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_dgemm( trans, MagmaNoTrans, n, nrhs, n,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
/* Check whether the nrhs normwise backward errors satisfy the
* stopping criterion. If yes, set ITER=IITER>0 and return. */
for( j=0; j < nrhs; j++ ) {
i = magma_idamax( n, dX(0,j), 1) - 1;
magma_dgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_dlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_idamax ( n, dR(0,j), 1 ) - 1;
magma_dgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_dlange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
if ( Rnrm > Xnrm*cte ) {
goto L20;
}
}
/* If we are here, the nrhs normwise backward errors satisfy
* the stopping criterion, we are good to exit. */
*iter = iiter;
return *info;
L20:
iiter++;
}
/* If we are at this place of the code, this is because we have
* performed ITER=ITERMAX iterations and never satisified the
* stopping criterion. Set up the ITER flag accordingly and follow
* up on double precision routine. */
*iter = -ITERMAX - 1;
FALLBACK:
/* Single-precision iterative refinement failed to converge to a
* satisfactory solution, so we resort to double precision. */
magma_dgetrf_gpu( n, n, dA, ldda, ipiv, info );
if (*info == 0) {
magmablas_dlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dX, lddx );
magma_dgetrs_gpu( trans, n, nrhs, dA, ldda, ipiv, dX, lddx, info );
}
return *info;
}
/**
Purpose
-------
DLAHR2 reduces the first NB columns of a real general n-BY-(n-k+1)
matrix A so that elements below the k-th subdiagonal are zero. The
reduction is performed by an orthogonal similarity transformation
Q' * A * Q. The routine returns the matrices V and T which determine
Q as a block reflector I - V*T*V', and also the matrix Y = A * V.
(Note this is different than LAPACK, which computes Y = A * V * T.)
This is an auxiliary routine called by DGEHRD.
Arguments
---------
@param[in]
n INTEGER
The order of the matrix A.
@param[in]
k INTEGER
The offset for the reduction. Elements below the k-th
subdiagonal in the first NB columns are reduced to zero.
K < N.
@param[in]
nb INTEGER
The number of columns to be reduced.
@param[in,out]
A DOUBLE_PRECISION array, dimension (LDA,N-K+1)
On entry, the n-by-(n-k+1) general matrix A.
On exit, the elements on and above the k-th subdiagonal in
the first NB columns are overwritten with the corresponding
elements of the reduced matrix; the elements below the k-th
subdiagonal, with the array TAU, represent the matrix Q as a
product of elementary reflectors. The other columns of A are
unchanged. See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
tau DOUBLE_PRECISION array, dimension (NB)
The scalar factors of the elementary reflectors. See Further
Details.
@param[out]
T DOUBLE_PRECISION array, dimension (LDT,NB)
The upper triangular matrix T.
@param[in]
ldt INTEGER
The leading dimension of the array T. LDT >= NB.
@param[out]
Y DOUBLE_PRECISION array, dimension (LDY,NB)
The n-by-nb matrix Y.
@param[in]
ldy INTEGER
The leading dimension of the array Y. LDY >= N.
@param[in,out]
data Structure with pointers to dA, dT, dV, dW, dY
which are distributed across multiple GPUs.
Further Details
---------------
The matrix Q is represented as a product of nb elementary reflectors
Q = H(1) H(2) . . . H(nb).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i+k-1) = 0, v(i+k) = 1; v(i+k+1:n) is stored on exit in
A(i+k+1:n,i), and tau in TAU(i).
The elements of the vectors v together form the (n-k+1)-by-nb matrix
V which is needed, with T and Y, to apply the transformation to the
unreduced part of the matrix, using an update of the form:
A := (I - V*T*V') * (A - Y*T*V').
The contents of A on exit are illustrated by the following example
with n = 7, k = 3 and nb = 2:
@verbatim
( a a a a a )
( a a a a a )
( a a a a a )
( h h a a a )
( v1 h a a a )
( v1 v2 a a a )
( v1 v2 a a a )
@endverbatim
where "a" denotes an element of the original matrix A, h denotes a
modified element of the upper Hessenberg matrix H, and vi denotes an
element of the vector defining H(i).
This implementation follows the hybrid algorithm and notations described in
S. Tomov and J. Dongarra, "Accelerating the reduction to upper Hessenberg
form through hybrid GPU-based computing," University of Tennessee Computer
Science Technical Report, UT-CS-09-642 (also LAPACK Working Note 219),
May 24, 2009.
@ingroup magma_dgeev_aux
********************************************************************/
extern "C" magma_int_t
magma_dlahr2_m(
magma_int_t n, magma_int_t k, magma_int_t nb,
double *A, magma_int_t lda,
double *tau,
double *T, magma_int_t ldt,
double *Y, magma_int_t ldy,
struct dgehrd_data* data )
{
#define A( i, j ) ( A + (i) + (j)*lda)
#define Y( i, j ) ( Y + (i) + (j)*ldy)
#define T( i, j ) ( T + (i) + (j)*ldt)
#define dA( d, i, j ) (data->A [d] + (i) + (j)*ldda)
#define dTi( d ) (data->Ti[d])
#define dV( d, i, j ) (data->V [d] + (i) + (j)*ldv )
#define dVd( d, i, j ) (data->Vd[d] + (i) + (j)*ldvd)
#define dY( d, i, j ) (data->Y [d] + (i) + (j)*ldda)
double c_zero = MAGMA_D_ZERO;
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
double tmp;
magma_int_t ngpu = data->ngpu;
magma_int_t ldda = data->ldda;
magma_int_t ldv = data->ldv;
magma_int_t ldvd = data->ldvd;
magma_int_t ione = 1;
magma_int_t d, dki1, dn, nblocks, gblock, lblock, lgid;
magma_int_t n_k_i_1, n_k;
double scale;
magma_int_t i;
double ei = MAGMA_D_ZERO;
magma_int_t info_data = 0;
magma_int_t *info = &info_data;
if (n < 0) {
*info = -1;
} else if (k < 0 || k >= n) {
*info = -2;
} else if (nb < 1 || nb > n) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
} else if (ldt < nb) {
*info = -8;
} else if (ldy < max(1,n)) {
*info = -10;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
// adjust from 1-based indexing
k -= 1;
// Function Body
if (n <= 1)
return 0;
// zero out current top block of V on all GPUs
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magmablasSetKernelStream( data->streams[d] );
magmablas_dlaset( MagmaFull, nb, nb, c_zero, c_zero, dV(d,k,0), ldv );
}
// set all Y=0
lapackf77_dlaset( "Full", &n, &nb, &c_zero, &c_zero, Y, &ldy );
for (i = 0; i < nb; ++i) {
n_k_i_1 = n - k - i - 1;
n_k = n - k;
if (i > 0) {
// Finish applying I - V * T * V' on right
tmp = MAGMA_D_NEGATE( tau[i-1] );
blasf77_daxpy( &n_k, &tmp, Y(k,i-1), &ione, A(k,i), &ione );
// Apply I - V * T' * V' to this column (call it b) from the
// left, using the last column of T as workspace, w.
//
// Let V = ( V1 ) and b = ( b1 ) (first i-1 rows)
// ( V2 ) ( b2 )
// where V1 is unit lower triangular
// w := b1 = A(k+1:k+i, i)
blasf77_dcopy( &i,
A(k+1,i), &ione,
T(0,nb-1), &ione );
// w := V1' * b1 = VA(k+1:k+i, 0:i-1)' * w
blasf77_dtrmv( "Lower", "Conj", "Unit", &i,
A(k+1,0), &lda,
T(0,nb-1), &ione );
// w := w + V2'*b2 = w + VA(k+i+1:n-1, 0:i-1)' * A(k+i+1:n-1, i)
blasf77_dgemv( "Conj", &n_k_i_1, &i,
&c_one, A(k+i+1,0), &lda,
A(k+i+1,i), &ione,
&c_one, T(0,nb-1), &ione );
// w := T'*w = T(0:i-1, 0:i-1)' * w
blasf77_dtrmv( "Upper", "Conj", "Non-unit", &i,
T(0,0), &ldt,
T(0,nb-1), &ione );
// b2 := b2 - V2*w = A(k+i+1:n-1, i) - VA(k+i+1:n-1, 0:i-1) * w
blasf77_dgemv( "No trans", &n_k_i_1, &i,
&c_neg_one, A(k+i+1,0), &lda,
T(0,nb-1), &ione,
&c_one, A(k+i+1,i), &ione );
// w := V1*w = VA(k+1:k+i, 0:i-1) * w
blasf77_dtrmv( "Lower", "No trans", "Unit", &i,
A(k+1,0), &lda,
T(0,nb-1), &ione );
// b1 := b1 - w = A(k+1:k+i-1, i) - w
blasf77_daxpy( &i,
&c_neg_one, T(0,nb-1), &ione,
A(k+1,i), &ione );
// Restore diagonal element, saved below during previous iteration
*A(k+i,i-1) = ei;
}
// Generate the elementary reflector H(i) to annihilate A(k+i+1:n-1,i)
lapackf77_dlarfg( &n_k_i_1,
A(k+i+1,i),
A(k+i+2,i), &ione, &tau[i] );
// Save diagonal element and set to one, to simplify multiplying by V
ei = *A(k+i+1,i);
*A(k+i+1,i) = c_one;
// compute yi = A vi = sum_g A{d} vi{d}
nblocks = (n-1) / nb / ngpu + 1;
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magmablasSetKernelStream( data->streams[d] );
// dV(k+i+1:n-1, i) = VA(k+i:n, i)
magma_dsetvector_async( n_k_i_1,
A(k+i+1,i), 1,
dV(d, k+i+1, i), 1, data->streams[d] );
// copy column of dV -> dVd, using block cyclic distribution.
// This assumes V and Vd have been padded so that
// a 2D matrix copy doesn't access them out-of-bounds
gblock = k / nb;
lblock = gblock / ngpu;
lgid = gblock % ngpu;
if ( d < lgid ) {
lblock += 1;
}
// treat V as (nb*ngpu) x nblock matrix, and Vd as nb x nblock matrix
magmablas_dlacpy( MagmaFull, nb, nblocks-lblock,
dV (d, d*nb + lblock*nb*ngpu, i), nb*ngpu,
dVd(d, 0 + lblock*nb, i), nb );
// convert global indices (k) to local indices (dk)
magma_indices_1D_bcyclic( nb, ngpu, d, k+i+1, n, &dki1, &dn );
// dY(k:n, i) = dA(k:n, k+i+1:n) * dV(k+i+1:n, i)
// skip if matrix is empty
// each GPU copies to different temporary vector in Y,
// which are summed in separate loop below
if ( dn-dki1 > 0 ) {
magma_dgemv( MagmaNoTrans, n-k, dn-dki1,
c_one, dA (d, k, dki1), ldda,
dVd(d, dki1, i), 1,
c_zero, dY (d, k, i), 1 );
// copy vector to host, storing in column nb+d of Y
// as temporary space (Y has >= nb+ngpu columns)
magma_dgetvector_async( n-k,
dY(d, k, i), 1,
Y(k, nb+d), 1, data->streams[d] );
}
}
// while GPU is doing above Ag*v...
// Compute T(0:i,i) = [ -tau T V' vi ]
// [ tau ]
// T(0:i-1, i) = -tau VA(k+i+1:n-1, 0:i-1)' VA(k+i+1:n-1, i)
scale = MAGMA_D_NEGATE( tau[i] );
blasf77_dgemv( "Conj", &n_k_i_1, &i,
&scale, A(k+i+1,0), &lda,
A(k+i+1,i), &ione,
&c_zero, T(0,i), &ione );
// T(0:i-1, i) = T(0:i-1, 0:i-1) * T(0:i-1, i)
blasf77_dtrmv( "Upper", "No trans", "Non-unit", &i,
T(0,0), &ldt,
T(0,i), &ione );
*T(i,i) = tau[i];
// apply reflectors to next column, A(i+1), on right only.
// one axpy will be required to finish this, in the next iteration above
if ( i > 0 && i+1 < nb ) {
// Update next column, A(k:n,i+1), applying Q on right.
// One axpy will be required to finish this, in the next iteration
// above, after yi is computed.
// This updates one more row than LAPACK does (row k),
// making block above panel an even multiple of nb.
// Use last column of T as workspace, w.
magma_int_t i1 = i+1;
// If real, conjugate row of V, and undo afterwards
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i1, A(k+i1,0), &lda );
#endif
// w = T(0:i, 0:i+1) * VA(k+i+1, 0:i+1)'
// T is now rectangular, so we use gemv instead of trmv as in lapack.
blasf77_dgemv( "No trans", &i, &i1,
&c_one, T(0,0), &ldt,
A(k+i1,0), &lda,
&c_zero, T(0,nb-1), &ione );
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i1, A(k+i1,0), &lda );
#endif
// A(k:n, i+1) -= Y(k:n, 0:i) * w
blasf77_dgemv( "No trans", &n_k, &i,
&c_neg_one, Y(k,0), &ldy,
T(0,nb-1), &ione,
&c_one, A(k,i1), &ione );
}
// yi = sum_g yi{d}
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magma_queue_sync( data->streams[d] );
magma_indices_1D_bcyclic( nb, ngpu, d, k+i+1, n, &dki1, &dn );
if ( dn-dki1 > 0 ) {
// yi = yi + yi{d}
blasf77_daxpy( &n_k, &c_one, Y(k,nb+d), &ione, Y(k,i), &ione );
}
}
}
// Restore diagonal element
*A(k+nb,nb-1) = ei;
// compute Y = Am V = sum_g Am{d} V{d} --- top part, Y(0:k-1,:)
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magmablasSetKernelStream( data->streams[d] );
// convert global indices (k) to local indices (dk)
magma_indices_1D_bcyclic( nb, ngpu, d, k+1, n, &dki1, &dn );
// dY(0:k, :) = dA(0:k, k+i+1:n-1) * dV(k+i+1:n-1, :)
// skip if matrix is empty
// each GPU copies to different temporary block in Y,
// which are summed in separate loop below
if ( dn-dki1 > 0 ) {
magma_dgemm( MagmaNoTrans, MagmaNoTrans, k, nb, dn-dki1,
c_one, dA (d, 0, dki1), ldda,
dVd(d, dki1, 0), ldvd,
c_zero, dY (d, 0, 0), ldda );
// copy result to host, storing in columns [nb + nb*d : nb + nb*(d+1)] of Y
// as temporary space (Y has nb + nb*ngpu columns)
magma_dgetmatrix_async( k, nb,
dY(d, 0, 0), ldda,
Y(0,nb+nb*d), ldy, data->streams[d] );
}
}
// Y = sum_g Y{d}
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magma_queue_sync( 0 );
magma_indices_1D_bcyclic( nb, ngpu, d, k+1, n, &dki1, &dn );
if ( dn-dki1 > 0 ) {
// Y = Y + Am V
for( i = 0; i < nb; ++i ) {
blasf77_daxpy( &k, &c_one, Y(0,nb+nb*d+i), &ione, Y(0,i), &ione );
}
}
}
// copy Y and T matrices to GPUs
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magma_dsetmatrix_async( n, nb, Y, ldy, dY(d, 0, 0), ldda, data->streams[d] );
magma_dsetmatrix_async( nb, nb, T, nb, dTi(d), nb, data->streams[d] );
}
return 0;
} /* magma_dlahr2 */
extern "C" magma_int_t
magma_dlobpcg( magma_d_sparse_matrix A, magma_d_solver_par *solver_par ){
#define residualNorms(i,iter) ( residualNorms + (i) + (iter)*n )
#define magmablas_swap(x, y) { pointer = x; x = y; y = pointer; }
#define hresidualNorms(i,iter) (hresidualNorms + (i) + (iter)*n )
#define gramA( m, n) (gramA + (m) + (n)*ldgram)
#define gramB( m, n) (gramB + (m) + (n)*ldgram)
#define gevectors(m, n) (gevectors + (m) + (n)*ldgram)
#define h_gramB( m, n) (h_gramB + (m) + (n)*ldgram)
#define magma_d_bspmv_tuned(m, n, alpha, A, X, beta, AX) { \
magmablas_dtranspose( m, n, X, m, blockW, n ); \
magma_d_vector x, ax; \
x.memory_location = Magma_DEV; x.num_rows = m*n; x.nnz = m*n; x.val = blockW; \
ax.memory_location= Magma_DEV; ax.num_rows = m*n; ax.nnz = m*n; ax.val = AX; \
magma_d_spmv(alpha, A, x, beta, ax ); \
magmablas_dtranspose( n, m, blockW, n, X, m ); \
}
//**************************************************************
// Memory allocation for the eigenvectors, eigenvalues, and workspace
solver_par->solver = Magma_LOBPCG;
magma_int_t m = A.num_rows;
magma_int_t n =(solver_par->num_eigenvalues);
double *blockX = solver_par->eigenvectors;
double *evalues = solver_par->eigenvalues;
double *dwork, *hwork;
double *blockP, *blockAP, *blockR, *blockAR, *blockAX, *blockW;
double *gramA, *gramB, *gramM;
double *gevectors, *h_gramB;
double *pointer, *origX = blockX;
double *eval_gpu;
magma_int_t lwork = max( 2*n+n*magma_get_dsytrd_nb(n),
1 + 6*3*n + 2* 3*n* 3*n);
magma_dmalloc_pinned( &hwork , lwork );
magma_dmalloc( &blockAX , m*n );
magma_dmalloc( &blockAR , m*n );
magma_dmalloc( &blockAP , m*n );
magma_dmalloc( &blockR , m*n );
magma_dmalloc( &blockP , m*n );
magma_dmalloc( &blockW , m*n );
magma_dmalloc( &dwork , m*n );
magma_dmalloc( &eval_gpu , 3*n );
//**********************************************************+
magma_int_t verbosity = 1;
magma_int_t *iwork, liwork = 15*n+9;
// === Set solver parameters ===
double residualTolerance = solver_par->epsilon;
magma_int_t maxIterations = solver_par->maxiter;
// === Set some constants & defaults ===
double c_one = MAGMA_D_ONE, c_zero = MAGMA_D_ZERO;
double *residualNorms, *condestGhistory, condestG;
double *gevalues;
magma_int_t *activeMask;
// === Check some parameters for possible quick exit ===
solver_par->info = 0;
if (m < 2)
solver_par->info = -1;
else if (n > m)
solver_par->info = -2;
if (solver_par->info != 0) {
magma_xerbla( __func__, -(solver_par->info) );
return solver_par->info;
}
magma_int_t *info = &(solver_par->info); // local info variable;
// === Allocate GPU memory for the residual norms' history ===
magma_dmalloc(&residualNorms, (maxIterations+1) * n);
magma_malloc( (void **)&activeMask, (n+1) * sizeof(magma_int_t) );
// === Allocate CPU work space ===
magma_dmalloc_cpu(&condestGhistory, maxIterations+1);
magma_dmalloc_cpu(&gevalues, 3 * n);
magma_malloc_cpu((void **)&iwork, liwork * sizeof(magma_int_t));
double *hW;
magma_dmalloc_pinned(&hW, n*n);
magma_dmalloc_pinned(&gevectors, 9*n*n);
magma_dmalloc_pinned(&h_gramB , 9*n*n);
// === Allocate GPU workspace ===
magma_dmalloc(&gramM, n * n);
magma_dmalloc(&gramA, 9 * n * n);
magma_dmalloc(&gramB, 9 * n * n);
#if defined(PRECISION_z) || defined(PRECISION_c)
double *rwork;
magma_int_t lrwork = 1 + 5*(3*n) + 2*(3*n)*(3*n);
magma_dmalloc_cpu(&rwork, lrwork);
#endif
// === Set activemask to one ===
for(int k =0; k<n; k++)
iwork[k]=1;
magma_setmatrix(n, 1, sizeof(magma_int_t), iwork, n ,activeMask, n);
magma_int_t gramDim, ldgram = 3*n, ikind = 4;
// === Make the initial vectors orthonormal ===
magma_dgegqr_gpu(ikind, m, n, blockX, m, dwork, hwork, info );
//magma_dorthomgs( m, n, blockX );
magma_d_bspmv_tuned(m, n, c_one, A, blockX, c_zero, blockAX );
// === Compute the Gram matrix = (X, AX) & its eigenstates ===
magma_dgemm(MagmaConjTrans, MagmaNoTrans, n, n, m,
c_one, blockX, m, blockAX, m, c_zero, gramM, n);
magma_dsyevd_gpu( MagmaVec, MagmaUpper,
n, gramM, n, evalues, hW, n, hwork, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork, info );
// === Update X = X * evectors ===
magma_dgemm(MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, blockX, m, gramM, n, c_zero, blockW, m);
magmablas_swap(blockW, blockX);
// === Update AX = AX * evectors ===
magma_dgemm(MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, blockAX, m, gramM, n, c_zero, blockW, m);
magmablas_swap(blockW, blockAX);
condestGhistory[1] = 7.82;
magma_int_t iterationNumber, cBlockSize, restart = 1, iter;
//Chronometry
real_Double_t tempo1, tempo2;
magma_device_sync(); tempo1=magma_wtime();
// === Main LOBPCG loop ============================================================
for(iterationNumber = 1; iterationNumber < maxIterations; iterationNumber++)
{
// === compute the residuals (R = Ax - x evalues )
magmablas_dlacpy( MagmaUpperLower, m, n, blockAX, m, blockR, m);
/*
for(int i=0; i<n; i++){
magma_daxpy(m, MAGMA_D_MAKE(-evalues[i],0), blockX+i*m, 1, blockR+i*m, 1);
}
*/
#if defined(PRECISION_z) || defined(PRECISION_d)
magma_dsetmatrix( 3*n, 1, evalues, 3*n, eval_gpu, 3*n );
#else
magma_ssetmatrix( 3*n, 1, evalues, 3*n, eval_gpu, 3*n );
#endif
magma_dlobpcg_res( m, n, eval_gpu, blockX, blockR, eval_gpu);
magmablas_dnrm2_cols(m, n, blockR, m, residualNorms(0, iterationNumber));
// === remove the residuals corresponding to already converged evectors
magma_dcompact(m, n, blockR, m,
residualNorms(0, iterationNumber), residualTolerance,
activeMask, &cBlockSize);
if (cBlockSize == 0)
break;
// === apply a preconditioner P to the active residulas: R_new = P R_old
// === for now set P to be identity (no preconditioner => nothing to be done )
// magmablas_dlacpy( MagmaUpperLower, m, cBlockSize, blockR, m, blockW, m);
/*
// === make the preconditioned residuals orthogonal to X
magma_dgemm(MagmaConjTrans, MagmaNoTrans, n, cBlockSize, m,
c_one, blockX, m, blockR, m, c_zero, gramB(0,0), ldgram);
magma_dgemm(MagmaNoTrans, MagmaNoTrans, m, cBlockSize, n,
c_mone, blockX, m, gramB(0,0), ldgram, c_one, blockR, m);
*/
// === make the active preconditioned residuals orthonormal
magma_dgegqr_gpu(ikind, m, cBlockSize, blockR, m, dwork, hwork, info );
//magma_dorthomgs( m, cBlockSize, blockR );
// === compute AR
magma_d_bspmv_tuned(m, cBlockSize, c_one, A, blockR, c_zero, blockAR );
if (!restart) {
// === compact P & AP as well
magma_dcompactActive(m, n, blockP, m, activeMask);
magma_dcompactActive(m, n, blockAP, m, activeMask);
/*
// === make P orthogonal to X ?
magma_dgemm(MagmaConjTrans, MagmaNoTrans, n, cBlockSize, m,
c_one, blockX, m, blockP, m, c_zero, gramB(0,0), ldgram);
magma_dgemm(MagmaNoTrans, MagmaNoTrans, m, cBlockSize, n,
c_mone, blockX, m, gramB(0,0), ldgram, c_one, blockP, m);
// === make P orthogonal to R ?
magma_dgemm(MagmaConjTrans, MagmaNoTrans, cBlockSize, cBlockSize, m,
c_one, blockR, m, blockP, m, c_zero, gramB(0,0), ldgram);
magma_dgemm(MagmaNoTrans, MagmaNoTrans, m, cBlockSize, cBlockSize,
c_mone, blockR, m, gramB(0,0), ldgram, c_one, blockP, m);
*/
// === Make P orthonormal & properly change AP (without multiplication by A)
magma_dgegqr_gpu(ikind, m, cBlockSize, blockP, m, dwork, hwork, info );
//magma_dorthomgs( m, cBlockSize, blockP );
//magma_d_bspmv_tuned(m, cBlockSize, c_one, A, blockP, c_zero, blockAP );
magma_dsetmatrix( cBlockSize, cBlockSize, hwork, cBlockSize, dwork, cBlockSize);
// magma_dtrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
// m, cBlockSize, c_one, dwork, cBlockSize, blockAP, m);
// replacement according to Stan
#if defined(PRECISION_s) || defined(PRECISION_d)
magmablas_dtrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
m, cBlockSize, c_one, dwork, cBlockSize, blockAP, m);
#else
magma_dtrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit, m,
cBlockSize, c_one, dwork, cBlockSize, blockAP, m);
#endif
}
iter = max(1,iterationNumber-10- (int)(log(1.*cBlockSize)));
double condestGmean = 0.;
for(int i = 0; i<iterationNumber-iter+1; i++)
condestGmean += condestGhistory[i];
condestGmean = condestGmean / (iterationNumber-iter+1);
if (restart)
gramDim = n+cBlockSize;
else
gramDim = n+2*cBlockSize;
/* --- The Raileight-Ritz method for [X R P] -----------------------
[ X R P ]' [AX AR AP] y = evalues [ X R P ]' [ X R P ], i.e.,
GramA GramB
/ X'AX X'AR X'AP \ / X'X X'R X'P \
| R'AX R'AR R'AP | y = evalues | R'X R'R R'P |
\ P'AX P'AR P'AP / \ P'X P'R P'P /
----------------------------------------------------------------- */
// === assemble GramB; first, set it to I
magmablas_dlaset(MagmaFull, ldgram, ldgram, c_zero, c_one, gramB, ldgram); // identity
if (!restart) {
magma_dgemm(MagmaConjTrans, MagmaNoTrans, cBlockSize, n, m,
c_one, blockP, m, blockX, m, c_zero, gramB(n+cBlockSize,0), ldgram);
magma_dgemm(MagmaConjTrans, MagmaNoTrans, cBlockSize, cBlockSize, m,
c_one, blockP, m, blockR, m, c_zero, gramB(n+cBlockSize,n), ldgram);
}
magma_dgemm(MagmaConjTrans, MagmaNoTrans, cBlockSize, n, m,
c_one, blockR, m, blockX, m, c_zero, gramB(n,0), ldgram);
// === get GramB from the GPU to the CPU and compute its eigenvalues only
magma_dgetmatrix(gramDim, gramDim, gramB, ldgram, h_gramB, ldgram);
lapackf77_dsyev("N", "L", &gramDim, h_gramB, &ldgram, gevalues,
hwork, &lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork,
#endif
info);
// === check stability criteria if we need to restart
condestG = log10( gevalues[gramDim-1]/gevalues[0] ) + 1.;
if ((condestG/condestGmean>2 && condestG>2) || condestG>8) {
// Steepest descent restart for stability
restart=1;
printf("restart at step #%d\n", (int) iterationNumber);
}
// === assemble GramA; first, set it to I
magmablas_dlaset(MagmaFull, ldgram, ldgram, c_zero, c_one, gramA, ldgram); // identity
magma_dgemm(MagmaConjTrans, MagmaNoTrans, cBlockSize, n, m,
c_one, blockR, m, blockAX, m, c_zero, gramA(n,0), ldgram);
magma_dgemm(MagmaConjTrans, MagmaNoTrans, cBlockSize, cBlockSize, m,
c_one, blockR, m, blockAR, m, c_zero, gramA(n,n), ldgram);
if (!restart) {
magma_dgemm(MagmaConjTrans, MagmaNoTrans, cBlockSize, n, m,
c_one, blockP, m, blockAX, m, c_zero,
gramA(n+cBlockSize,0), ldgram);
magma_dgemm(MagmaConjTrans, MagmaNoTrans, cBlockSize, cBlockSize, m,
c_one, blockP, m, blockAR, m, c_zero,
gramA(n+cBlockSize,n), ldgram);
magma_dgemm(MagmaConjTrans, MagmaNoTrans, cBlockSize, cBlockSize, m,
c_one, blockP, m, blockAP, m, c_zero,
gramA(n+cBlockSize,n+cBlockSize), ldgram);
}
/*
// === Compute X' AX or just use the eigenvalues below ?
magma_dgemm(MagmaConjTrans, MagmaNoTrans, n, n, m,
c_one, blockX, m, blockAX, m, c_zero,
gramA(0,0), ldgram);
*/
if (restart==0) {
magma_dgetmatrix(gramDim, gramDim, gramA, ldgram, gevectors, ldgram);
}
else {
gramDim = n+cBlockSize;
magma_dgetmatrix(gramDim, gramDim, gramA, ldgram, gevectors, ldgram);
}
for(int k=0; k<n; k++)
*gevectors(k,k) = MAGMA_D_MAKE(evalues[k], 0);
// === the previous eigensolver destroyed what is in h_gramB => must copy it again
magma_dgetmatrix(gramDim, gramDim, gramB, ldgram, h_gramB, ldgram);
magma_int_t itype = 1;
lapackf77_dsygvd(&itype, "V", "L", &gramDim,
gevectors, &ldgram, h_gramB, &ldgram,
gevalues, hwork, &lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, &lrwork,
#endif
iwork, &liwork, info);
for(int k =0; k<n; k++)
evalues[k] = gevalues[k];
// === copy back the result to gramA on the GPU and use it for the updates
magma_dsetmatrix(gramDim, gramDim, gevectors, ldgram, gramA, ldgram);
if (restart == 0) {
// === contribution from P to the new X (in new search direction P)
magma_dgemm(MagmaNoTrans, MagmaNoTrans, m, n, cBlockSize,
c_one, blockP, m, gramA(n+cBlockSize,0), ldgram, c_zero, dwork, m);
magmablas_swap(dwork, blockP);
// === contribution from R to the new X (in new search direction P)
magma_dgemm(MagmaNoTrans, MagmaNoTrans, m, n, cBlockSize,
c_one, blockR, m, gramA(n,0), ldgram, c_one, blockP, m);
// === corresponding contribution from AP to the new AX (in AP)
magma_dgemm(MagmaNoTrans, MagmaNoTrans, m, n, cBlockSize,
c_one, blockAP, m, gramA(n+cBlockSize,0), ldgram, c_zero, dwork, m);
magmablas_swap(dwork, blockAP);
// === corresponding contribution from AR to the new AX (in AP)
magma_dgemm(MagmaNoTrans, MagmaNoTrans, m, n, cBlockSize,
c_one, blockAR, m, gramA(n,0), ldgram, c_one, blockAP, m);
}
else {
// === contribution from R (only) to the new X
magma_dgemm(MagmaNoTrans, MagmaNoTrans, m, n, cBlockSize,
c_one, blockR, m, gramA(n,0), ldgram, c_zero, blockP, m);
// === corresponding contribution from AR (only) to the new AX
magma_dgemm(MagmaNoTrans, MagmaNoTrans,m, n, cBlockSize,
c_one, blockAR, m, gramA(n,0), ldgram, c_zero, blockAP, m);
}
// === contribution from old X to the new X + the new search direction P
magma_dgemm(MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, blockX, m, gramA, ldgram, c_zero, dwork, m);
magmablas_swap(dwork, blockX);
//magma_daxpy(m*n, c_one, blockP, 1, blockX, 1);
magma_dlobpcg_maxpy( m, n, blockP, blockX );
// === corresponding contribution from old AX to new AX + AP
magma_dgemm(MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, blockAX, m, gramA, ldgram, c_zero, dwork, m);
magmablas_swap(dwork, blockAX);
//magma_daxpy(m*n, c_one, blockAP, 1, blockAX, 1);
magma_dlobpcg_maxpy( m, n, blockAP, blockAX );
condestGhistory[iterationNumber+1]=condestG;
if (verbosity==1) {
// double res;
// magma_dgetmatrix(1, 1,
// (double*)residualNorms(0, iterationNumber), 1,
// (double*)&res, 1);
//
// printf("Iteration %4d, CBS %4d, Residual: %10.7f\n",
// iterationNumber, cBlockSize, res);
printf("%4d-%2d ", (int) iterationNumber, (int) cBlockSize);
magma_dprint_gpu(1, n, residualNorms(0, iterationNumber), 1);
}
restart = 0;
} // === end for iterationNumber = 1,maxIterations =======================
// fill solver info
magma_device_sync(); tempo2=magma_wtime();
solver_par->runtime = (real_Double_t) tempo2-tempo1;
solver_par->numiter = iterationNumber;
if( solver_par->numiter < solver_par->maxiter){
solver_par->info = 0;
}else if( solver_par->init_res > solver_par->final_res )
solver_par->info = -2;
else
solver_par->info = -1;
// =============================================================================
// === postprocessing;
// =============================================================================
// === compute the real AX and corresponding eigenvalues
magma_d_bspmv_tuned(m, n, c_one, A, blockX, c_zero, blockAX );
magma_dgemm(MagmaConjTrans, MagmaNoTrans, n, n, m,
c_one, blockX, m, blockAX, m, c_zero, gramM, n);
magma_dsyevd_gpu( MagmaVec, MagmaUpper,
n, gramM, n, gevalues, dwork, n, hwork, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork, info );
for(int k =0; k<n; k++)
evalues[k] = gevalues[k];
// === update X = X * evectors
magmablas_swap(blockX, dwork);
magma_dgemm(MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, dwork, m, gramM, n, c_zero, blockX, m);
// === update AX = AX * evectors to compute the final residual
magmablas_swap(blockAX, dwork);
magma_dgemm(MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, dwork, m, gramM, n, c_zero, blockAX, m);
// === compute R = AX - evalues X
magmablas_dlacpy( MagmaUpperLower, m, n, blockAX, m, blockR, m);
for(int i=0; i<n; i++)
magma_daxpy(m, MAGMA_D_MAKE(-evalues[i], 0), blockX+i*m, 1, blockR+i*m, 1);
// === residualNorms[iterationNumber] = || R ||
magmablas_dnrm2_cols(m, n, blockR, m, residualNorms(0, iterationNumber));
// === restore blockX if needed
if (blockX != origX)
magmablas_dlacpy( MagmaUpperLower, m, n, blockX, m, origX, m);
printf("Eigenvalues:\n");
for(int i =0; i<n; i++)
printf("%e ", evalues[i]);
printf("\n\n");
printf("Final residuals:\n");
magma_dprint_gpu(1, n, residualNorms(0, iterationNumber), 1);
printf("\n\n");
//=== Print residual history in a file for plotting ====
double *hresidualNorms;
magma_dmalloc_cpu(&hresidualNorms, (iterationNumber+1) * n);
magma_dgetmatrix(n, iterationNumber,
(double*)residualNorms, n,
(double*)hresidualNorms, n);
printf("Residuals are stored in file residualNorms\n");
printf("Plot the residuals using: myplot \n");
FILE *residuals_file;
residuals_file = fopen("residualNorms", "w");
for(int i =1; i<iterationNumber; i++) {
for(int j = 0; j<n; j++)
fprintf(residuals_file, "%f ", *hresidualNorms(j,i));
fprintf(residuals_file, "\n");
}
fclose(residuals_file);
magma_free_cpu(hresidualNorms);
// === free work space
magma_free( residualNorms );
magma_free_cpu( condestGhistory );
magma_free_cpu( gevalues );
magma_free_cpu( iwork );
magma_free_pinned( hW );
magma_free_pinned( gevectors );
magma_free_pinned( h_gramB );
magma_free( gramM );
magma_free( gramA );
magma_free( gramB );
magma_free( activeMask );
magma_free( blockAX );
magma_free( blockAR );
magma_free( blockAP );
magma_free( blockR );
magma_free( blockP );
magma_free( blockW );
magma_free( dwork );
magma_free( eval_gpu );
magma_free_pinned( hwork );
#if defined(PRECISION_z) || defined(PRECISION_c)
magma_free_cpu( rwork );
#endif
return MAGMA_SUCCESS;
}
/**
Purpose
-------
DSPOSV computes the solution to a real system of linear equations
A * X = B,
where A is an N-by-N symmetric positive definite matrix and X and B
are N-by-NRHS matrices.
DSPOSV first attempts to factorize the matrix in real SINGLE PRECISION
and use this factorization within an iterative refinement procedure
to produce a solution with real DOUBLE PRECISION norm-wise backward error
quality (see below). If the approach fails the method switches to a
real DOUBLE PRECISION factorization and solve.
The iterative refinement is not going to be a winning strategy if
the ratio real SINGLE PRECISION performance over real DOUBLE PRECISION
performance is too small. A reasonable strategy should take the
number of right-hand sides and the size of the matrix into account.
This might be done with a call to ILAENV in the future. Up to now, we
always try iterative refinement.
The iterative refinement process is stopped if
ITER > ITERMAX
or for all the RHS we have:
RNRM < SQRT(N)*XNRM*ANRM*EPS*BWDMAX
where
o ITER is the number of the current iteration in the iterative
refinement process
o RNRM is the infinity-norm of the residual
o XNRM is the infinity-norm of the solution
o ANRM is the infinity-operator-norm of the matrix A
o EPS is the machine epsilon returned by DLAMCH('Epsilon')
The value ITERMAX and BWDMAX are fixed to 30 and 1.0D+00 respectively.
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored;
- = MagmaLower: Lower triangle of A is stored.
@param[in]
n INTEGER
The number of linear equations, i.e., the order of the
matrix A. N >= 0.
@param[in]
nrhs INTEGER
The number of right hand sides, i.e., the number of columns
of the matrix B. NRHS >= 0.
@param[in,out]
dA DOUBLE PRECISION array on the GPU, dimension (LDDA,N)
On entry, the symmetric matrix A. If UPLO = MagmaUpper, the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit, if iterative refinement has been successfully used
(INFO.EQ.0 and ITER.GE.0, see description below), then A is
unchanged, if double factorization has been used
(INFO.EQ.0 and ITER.LT.0, see description below), then the
array dA contains the factor U or L from the Cholesky
factorization A = U**T*U or A = L*L**T.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
@param[in]
dB DOUBLE PRECISION array on the GPU, dimension (LDDB,NRHS)
The N-by-NRHS right hand side matrix B.
@param[in]
lddb INTEGER
The leading dimension of the array dB. LDDB >= max(1,N).
@param[out]
dX DOUBLE PRECISION array on the GPU, dimension (LDDX,NRHS)
If INFO = 0, the N-by-NRHS solution matrix X.
@param[in]
lddx INTEGER
The leading dimension of the array dX. LDDX >= max(1,N).
@param
dworkd (workspace) DOUBLE PRECISION array on the GPU, dimension (N*NRHS)
This array is used to hold the residual vectors.
@param
dworks (workspace) SINGLE PRECISION array on the GPU, dimension (N*(N+NRHS))
This array is used to store the real single precision matrix
and the right-hand sides or solutions in single precision.
@param[out]
iter INTEGER
- < 0: iterative refinement has failed, double precision
factorization has been performed
+ -1 : the routine fell back to full precision for
implementation- or machine-specific reasons
+ -2 : narrowing the precision induced an overflow,
the routine fell back to full precision
+ -3 : failure of SPOTRF
+ -31: stop the iterative refinement after the 30th iteration
- > 0: iterative refinement has been successfully used.
Returns the number of iterations
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: if INFO = i, the leading minor of order i of (DOUBLE
PRECISION) A is not positive definite, so the
factorization could not be completed, and the solution
has not been computed.
@ingroup magma_dposv_driver
********************************************************************/
extern "C" magma_int_t
magma_dsposv_gpu(
magma_uplo_t uplo, magma_int_t n, magma_int_t nrhs,
magmaDouble_ptr dA, magma_int_t ldda,
magmaDouble_ptr dB, magma_int_t lddb,
magmaDouble_ptr dX, magma_int_t lddx,
magmaDouble_ptr dworkd, magmaFloat_ptr dworks,
magma_int_t *iter,
magma_int_t *info)
{
#define dB(i,j) (dB + (i) + (j)*lddb)
#define dX(i,j) (dX + (i) + (j)*lddx)
#define dR(i,j) (dR + (i) + (j)*lddr)
#define dSX(i,j) (dSX + (i) + (j)*lddsx)
// Constants
const double BWDMAX = 1.0;
const magma_int_t ITERMAX = 30;
const double c_neg_one = MAGMA_D_NEG_ONE;
const double c_one = MAGMA_D_ONE;
const magma_int_t ione = 1;
// Local variables
magmaDouble_ptr dR;
magmaFloat_ptr dSA, dSX;
double Xnrmv, Rnrmv;
double Anrm, Xnrm, Rnrm, cte, eps;
magma_int_t i, j, iiter, lddsa, lddsx, lddr;
/* Check arguments */
*iter = 0;
*info = 0;
if ( n < 0 )
*info = -1;
else if ( nrhs < 0 )
*info = -2;
else if ( ldda < max(1,n))
*info = -4;
else if ( lddb < max(1,n))
*info = -7;
else if ( lddx < max(1,n))
*info = -9;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if ( n == 0 || nrhs == 0 )
return *info;
lddsa = n;
lddsx = n;
lddr = n;
dSA = dworks;
dSX = dSA + lddsa*n;
dR = dworkd;
magma_queue_t queue;
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queue );
eps = lapackf77_dlamch("Epsilon");
Anrm = magmablas_dlansy( MagmaInfNorm, uplo, n, dA, ldda, (double*)dworkd, n*nrhs, queue );
cte = Anrm * eps * magma_dsqrt( n ) * BWDMAX;
/*
* Convert to single precision
*/
magmablas_dlag2s( n, nrhs, dB, lddb, dSX, lddsx, queue, info );
if (*info != 0) {
*iter = -2;
goto fallback;
}
magmablas_dlat2s( uplo, n, dA, ldda, dSA, lddsa, queue, info );
if (*info != 0) {
*iter = -2;
goto fallback;
}
// factor dSA in single precision
magma_spotrf_gpu( uplo, n, dSA, lddsa, info );
if (*info != 0) {
*iter = -3;
goto fallback;
}
// solve dSA*dSX = dB in single precision
magma_spotrs_gpu( uplo, n, nrhs, dSA, lddsa, dSX, lddsx, info );
// residual dR = dB - dA*dX in double precision
magmablas_slag2d( n, nrhs, dSX, lddsx, dX, lddx, queue, info );
magmablas_dlacpy( MagmaFull, n, nrhs, dB, lddb, dR, lddr, queue );
if ( nrhs == 1 ) {
magma_dsymv( uplo, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1, queue );
}
else {
magma_dsymm( MagmaLeft, uplo, n, nrhs,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr, queue );
}
// TODO: use MAGMA_D_ABS( dX(i,j) ) instead of dlange?
for( j=0; j < nrhs; j++ ) {
i = magma_idamax( n, dX(0,j), 1, queue ) - 1;
magma_dgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1, queue );
Xnrm = lapackf77_dlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_idamax( n, dR(0,j), 1, queue ) - 1;
magma_dgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1, queue );
Rnrm = lapackf77_dlange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
if ( Rnrm > Xnrm*cte ) {
goto refinement;
}
}
*iter = 0;
goto cleanup;
//return *info;
refinement:
for( iiter=1; iiter < ITERMAX; ) {
*info = 0;
// convert residual dR to single precision dSX
magmablas_dlag2s( n, nrhs, dR, lddr, dSX, lddsx, queue, info );
if (*info != 0) {
*iter = -2;
goto fallback;
}
// solve dSA*dSX = R in single precision
magma_spotrs_gpu( uplo, n, nrhs, dSA, lddsa, dSX, lddsx, info );
// Add correction and setup residual
// dX += dSX [including conversion] --and--
// dR = dB
for( j=0; j < nrhs; j++ ) {
magmablas_dsaxpycp( n, dSX(0,j), dX(0,j), dB(0,j), dR(0,j), queue );
}
// residual dR = dB - dA*dX in double precision
if ( nrhs == 1 ) {
magma_dsymv( uplo, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1, queue );
}
else {
magma_dsymm( MagmaLeft, uplo, n, nrhs,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr, queue );
}
// TODO: use MAGMA_D_ABS( dX(i,j) ) instead of dlange?
/* Check whether the nrhs normwise backward errors satisfy the
* stopping criterion. If yes, set ITER=IITER > 0 and return. */
for( j=0; j < nrhs; j++ ) {
i = magma_idamax( n, dX(0,j), 1, queue ) - 1;
magma_dgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1, queue );
Xnrm = lapackf77_dlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_idamax( n, dR(0,j), 1, queue ) - 1;
magma_dgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1, queue );
Rnrm = lapackf77_dlange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
if ( Rnrm > Xnrm*cte ) {
goto L20;
}
}
/* If we are here, the nrhs normwise backward errors satisfy
* the stopping criterion, we are good to exit. */
*iter = iiter;
goto cleanup;
//return *info;
L20:
iiter++;
}
/* If we are at this place of the code, this is because we have
* performed ITER=ITERMAX iterations and never satisified the
* stopping criterion. Set up the ITER flag accordingly and follow
* up on double precision routine. */
*iter = -ITERMAX - 1;
fallback:
/* Single-precision iterative refinement failed to converge to a
* satisfactory solution, so we resort to double precision. */
magma_dpotrf_gpu( uplo, n, dA, ldda, info );
if (*info == 0) {
magmablas_dlacpy( MagmaFull, n, nrhs, dB, lddb, dX, lddx, queue );
magma_dpotrs_gpu( uplo, n, nrhs, dA, ldda, dX, lddx, info );
}
cleanup:
magma_queue_destroy( queue );
return *info;
}
extern "C" magma_int_t
magma_dgetri_gpu( magma_int_t n, double *dA, magma_int_t lda,
magma_int_t *ipiv, double *dwork, magma_int_t lwork,
magma_int_t *info )
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
Purpose
=======
DGETRI computes the inverse of a matrix using the LU factorization
computed by DGETRF. This method inverts U and then computes inv(A) by
solving the system inv(A)*L = inv(U) for inv(A).
Note that it is generally both faster and more accurate to use DGESV,
or DGETRF and DGETRS, to solve the system AX = B, rather than inverting
the matrix and multiplying to form X = inv(A)*B. Only in special
instances should an explicit inverse be computed with this routine.
Arguments
=========
N (input) INTEGER
The order of the matrix A. N >= 0.
dA (input/output) DOUBLE_PRECISION array on the GPU, dimension (LDA,N)
On entry, the factors L and U from the factorization
A = P*L*U as computed by DGETRF_GPU.
On exit, if INFO = 0, the inverse of the original matrix A.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
IPIV (input) INTEGER array, dimension (N)
The pivot indices from DGETRF; for 1<=i<=N, row i of the
matrix was interchanged with row IPIV(i).
DWORK (workspace/output) COMPLEX*16 array on the GPU, dimension (MAX(1,LWORK))
LWORK (input) INTEGER
The dimension of the array DWORK. LWORK >= N*NB, where NB is
the optimal blocksize returned by magma_get_dgetri_nb(n).
Unlike LAPACK, this version does not currently support a
workspace query, because the workspace is on the GPU.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
> 0: if INFO = i, U(i,i) is exactly zero; the matrix is
singular and its cannot be computed.
===================================================================== */
/* Local variables */
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
double *dL = dwork;
magma_int_t ldl = n;
magma_int_t nb = magma_get_dgetri_nb(n);
magma_int_t j, jmax, jb, jp;
*info = 0;
if (n < 0)
*info = -1;
else if (lda < max(1,n))
*info = -3;
else if ( lwork < n*nb )
*info = -6;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if ( n == 0 )
return *info;
/* Invert the triangular factor U */
magma_dtrtri_gpu( MagmaUpper, MagmaNonUnit, n, dA, lda, info );
if ( *info != 0 )
return *info;
jmax = ((n-1) / nb)*nb;
for( j = jmax; j >= 0; j -= nb ) {
jb = min( nb, n-j );
// copy current block column of L to work space,
// then replace with zeros in A.
magmablas_dlacpy( MagmaUpperLower, n-j, jb,
&dA[j + j*lda], lda,
&dL[j ], ldl );
magmablas_dlaset( MagmaLower, n-j, jb, &dA[j + j*lda], lda );
// compute current block column of Ainv
// Ainv(:, j:j+jb-1)
// = ( U(:, j:j+jb-1) - Ainv(:, j+jb:n) L(j+jb:n, j:j+jb-1) )
// * L(j:j+jb-1, j:j+jb-1)^{-1}
// where L(:, j:j+jb-1) is stored in dL.
if ( j+jb < n ) {
magma_dgemm( MagmaNoTrans, MagmaNoTrans, n, jb, n-j-jb,
c_neg_one, &dA[(j+jb)*lda], lda,
&dL[ j+jb ], ldl,
c_one, &dA[ j*lda], lda );
}
magma_dtrsm( MagmaRight, MagmaLower, MagmaNoTrans, MagmaUnit,
n, jb, c_one,
&dL[j ], ldl,
&dA[j*lda], lda );
}
// Apply column interchanges
for( j = n-2; j >= 0; --j ) {
jp = ipiv[j] - 1;
if ( jp != j ) {
magmablas_dswap( n, &dA[ j*lda ], 1, &dA[ jp*lda ], 1 );
}
}
return *info;
}
/**
Purpose
-------
DSGEQRSV solves the least squares problem
min || A*X - B ||,
where A is an M-by-N matrix and X and B are M-by-NRHS matrices.
DSGEQRSV first attempts to factorize the matrix in real SINGLE PRECISION
and use this factorization within an iterative refinement procedure
to produce a solution with real DOUBLE PRECISION norm-wise backward error
quality (see below). If the approach fails the method switches to a
real DOUBLE PRECISION factorization and solve.
The iterative refinement is not going to be a winning strategy if
the ratio real SINGLE PRECISION performance over real DOUBLE PRECISION
performance is too small. A reasonable strategy should take the
number of right-hand sides and the size of the matrix into account.
This might be done with a call to ILAENV in the future. Up to now, we
always try iterative refinement.
The iterative refinement process is stopped if
ITER > ITERMAX
or for all the RHS we have:
RNRM < SQRT(N)*XNRM*ANRM*EPS*BWDMAX
where
o ITER is the number of the current iteration in the iterative
refinement process
o RNRM is the infinity-norm of the residual
o XNRM is the infinity-norm of the solution
o ANRM is the infinity-operator-norm of the matrix A
o EPS is the machine epsilon returned by DLAMCH('Epsilon')
The value ITERMAX and BWDMAX are fixed to 30 and 1.0D+00 respectively.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. M >= N >= 0.
@param[in]
nrhs INTEGER
The number of right hand sides, i.e., the number of columns
of the matrix B. NRHS >= 0.
@param[in,out]
dA DOUBLE PRECISION array on the GPU, dimension (LDDA,N)
On entry, the M-by-N coefficient matrix A.
On exit, if iterative refinement has been successfully used
(info.EQ.0 and ITER.GE.0, see description below), A is
unchanged. If double precision factorization has been used
(info.EQ.0 and ITER.LT.0, see description below), then the
array dA contains the QR factorization of A as returned by
function DGEQRF_GPU.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
@param[in,out]
dB DOUBLE PRECISION array on the GPU, dimension (LDDB,NRHS)
The M-by-NRHS right hand side matrix B.
May be overwritten (e.g., if refinement fails).
@param[in]
lddb INTEGER
The leading dimension of the array dB. LDDB >= max(1,M).
@param[out]
dX DOUBLE PRECISION array on the GPU, dimension (LDDX,NRHS)
If info = 0, the N-by-NRHS solution matrix X.
@param[in]
lddx INTEGER
The leading dimension of the array dX. LDDX >= max(1,N).
@param[out]
iter INTEGER
- < 0: iterative refinement has failed, double precision
factorization has been performed
+ -1 : the routine fell back to full precision for
implementation- or machine-specific reasons
+ -2 : narrowing the precision induced an overflow,
the routine fell back to full precision
+ -3 : failure of SGEQRF
+ -31: stop the iterative refinement after the 30th iteration
- > 0: iterative refinement has been successfully used.
Returns the number of iterations
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if info = -i, the i-th argument had an illegal value
@ingroup magma_dgels_driver
********************************************************************/
extern "C" magma_int_t
magma_dsgeqrsv_gpu(
magma_int_t m, magma_int_t n, magma_int_t nrhs,
magmaDouble_ptr dA, magma_int_t ldda,
magmaDouble_ptr dB, magma_int_t lddb,
magmaDouble_ptr dX, magma_int_t lddx,
magma_int_t *iter,
magma_int_t *info)
{
#define dB(i,j) (dB + (i) + (j)*lddb)
#define dX(i,j) (dX + (i) + (j)*lddx)
#define dR(i,j) (dR + (i) + (j)*lddr)
#define dSX(i,j) (dSX + (i) + (j)*lddsx)
double c_neg_one = MAGMA_D_NEG_ONE;
double c_one = MAGMA_D_ONE;
magma_int_t ione = 1;
double *hworkd;
float *hworks;
double *tau;
float *stau;
magmaDouble_ptr dworkd;
magmaFloat_ptr dworks;
magmaDouble_ptr dR, dT;
magmaFloat_ptr dSA, dSX, dST;
double Xnrmv, Rnrmv;
double Anrm, Xnrm, Rnrm, cte, eps;
magma_int_t i, j, iiter, lddsa, lddsx, lddr, nb, lhwork, minmn, size, ldworkd;
/* Check arguments */
*iter = 0;
*info = 0;
if ( m < 0 )
*info = -1;
else if ( n < 0 || n > m )
*info = -2;
else if ( nrhs < 0 )
*info = -3;
else if ( ldda < max(1,m))
*info = -5;
else if ( lddb < max(1,m))
*info = -7;
else if ( lddx < max(1,n))
*info = -9;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if ( m == 0 || n == 0 || nrhs == 0 )
return *info;
nb = magma_get_sgeqrf_nb(m);
minmn= min(m, n);
/* dSX contains both B and X, so must be max(m or lddb,n). */
lddsa = ldda;
lddsx = max(lddb,n);
lddr = lddb;
/*
* Allocate temporary buffers
*/
/* dworks(dSA + dSX + dST) */
size = lddsa*n + lddsx*nrhs + ( 2*minmn + ((n+31)/32)*32 )*nb;
if (MAGMA_SUCCESS != magma_smalloc( &dworks, size )) {
fprintf(stderr, "Allocation of dworks failed (%d)\n", (int) size);
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dSA = dworks;
dSX = dSA + lddsa*n;
dST = dSX + lddsx*nrhs;
/* dworkd(dR) = lddr*nrhs */
ldworkd = lddr*nrhs;
if (MAGMA_SUCCESS != magma_dmalloc( &dworkd, ldworkd )) {
magma_free( dworks );
fprintf(stderr, "Allocation of dworkd failed\n");
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dR = dworkd;
/* hworks(workspace for cgeqrs + stau) = min(m,n) + lhworks */
lhwork = (m - n + nb)*(nrhs + nb) + nrhs*nb;
size = lhwork + minmn;
magma_smalloc_cpu( &hworks, size );
if ( hworks == NULL ) {
magma_free( dworks );
magma_free( dworkd );
fprintf(stderr, "Allocation of hworks failed\n");
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
stau = hworks + lhwork;
eps = lapackf77_dlamch("Epsilon");
Anrm = magmablas_dlange(MagmaInfNorm, m, n, dA, ldda, (double*)dworkd );
cte = Anrm * eps * pow((double)n, 0.5) * BWDMAX;
/*
* Convert to single precision
*/
magmablas_dlag2s( m, nrhs, dB, lddb, dSX, lddsx, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
magmablas_dlag2s( m, n, dA, ldda, dSA, lddsa, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
// factor dSA in single precision
magma_sgeqrf_gpu( m, n, dSA, lddsa, stau, dST, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// solve dSA*dSX = dB in single precision
magma_sgeqrs_gpu( m, n, nrhs, dSA, lddsa, stau, dST, dSX, lddsx, hworks, lhwork, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// residual dR = dB - dA*dX in double precision
magmablas_slag2d( n, nrhs, dSX, lddsx, dX, lddx, info );
magmablas_dlacpy( MagmaUpperLower, m, nrhs, dB, lddb, dR, lddr );
if ( nrhs == 1 ) {
magma_dgemv( MagmaNoTrans, m, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_dgemm( MagmaNoTrans, MagmaNoTrans, m, nrhs, n,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
// TODO: use MAGMA_D_ABS( dX(i,j) ) instead of dlange?
for( j=0; j < nrhs; j++ ) {
i = magma_idamax( n, dX(0,j), 1) - 1;
magma_dgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_dlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_idamax ( m, dR(0,j), 1 ) - 1;
magma_dgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_dlange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
if ( Rnrm > Xnrm*cte ) {
goto REFINEMENT;
}
}
*iter = 0;
/* Free workspaces */
magma_free( dworks );
magma_free( dworkd );
magma_free_cpu( hworks );
return *info;
REFINEMENT:
/* TODO: this iterative refinement algorithm works only for compatibile
* systems (B in colspan of A).
* See Matrix Computations (3rd ed) p. 267 for correct algorithm. */
for( iiter=1; iiter < ITERMAX; ) {
*info = 0;
// convert residual dR to single precision dSX
magmablas_dlag2s( m, nrhs, dR, lddr, dSX, lddsx, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
// solve dSA*dSX = R in single precision
magma_sgeqrs_gpu( m, n, nrhs, dSA, lddsa, stau, dST, dSX, lddsx, hworks, lhwork, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// Add correction and setup residual
// dX += dSX [including conversion] --and--
// dR[1:n] = dB[1:n] (only n rows, not whole m rows! -- useless if m > n)
for( j=0; j < nrhs; j++ ) {
magmablas_dsaxpycp( n, dSX(0,j), dX(0,j), dB(0,j), dR(0,j) );
}
// dR = dB (whole m rows)
magmablas_dlacpy( MagmaUpperLower, m, nrhs, dB, lddb, dR, lddr );
// residual dR = dB - dA*dX in double precision
if ( nrhs == 1 ) {
magma_dgemv( MagmaNoTrans, m, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_dgemm( MagmaNoTrans, MagmaNoTrans, m, nrhs, n,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
/* Check whether the nrhs normwise backward errors satisfy the
* stopping criterion. If yes, set ITER=IITER > 0 and return. */
for( j=0; j < nrhs; j++ ) {
i = magma_idamax( n, dX(0,j), 1) - 1;
magma_dgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_dlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_idamax ( m, dR(0,j), 1 ) - 1;
magma_dgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_dlange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
if ( Rnrm > Xnrm*cte ) {
goto L20;
}
}
/* If we are here, the nrhs normwise backward errors satisfy
* the stopping criterion, we are good to exit. */
*iter = iiter;
/* Free workspaces */
magma_free( dworks );
magma_free( dworkd );
magma_free_cpu( hworks );
return *info;
L20:
iiter++;
}
/* If we are at this place of the code, this is because we have
* performed ITER=ITERMAX iterations and never satisified the
* stopping criterion. Set up the ITER flag accordingly and follow
* up on double precision routine. */
*iter = -ITERMAX - 1;
FALLBACK:
/* Single-precision iterative refinement failed to converge to a
* satisfactory solution, so we resort to double precision. */
magma_free( dworks );
magma_free_cpu( hworks );
/*
* Allocate temporary buffers
*/
/* dworkd = dT for dgeqrf */
nb = magma_get_dgeqrf_nb( m );
size = (2*min(m, n) + (n+31)/32*32 )*nb;
if ( size > ldworkd ) {
magma_free( dworkd );
if (MAGMA_SUCCESS != magma_dmalloc( &dworkd, size )) {
fprintf(stderr, "Allocation of dworkd2 failed\n");
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
}
dT = dworkd;
/* hworkd(dtau + workspace for dgeqrs) = min(m,n) + lhwork */
size = lhwork + minmn;
magma_dmalloc_cpu( &hworkd, size );
if ( hworkd == NULL ) {
magma_free( dworkd );
fprintf(stderr, "Allocation of hworkd2 failed\n");
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
tau = hworkd + lhwork;
magma_dgeqrf_gpu( m, n, dA, ldda, tau, dT, info );
if (*info == 0) {
// if m > n, then dB won't fit in dX, so solve with dB and copy n rows to dX
magma_dgeqrs_gpu( m, n, nrhs, dA, ldda, tau, dT, dB, lddb, hworkd, lhwork, info );
magmablas_dlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dX, lddx );
}
magma_free( dworkd );
magma_free_cpu( hworkd );
return *info;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dlacpy
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gbytes, gpu_perf, gpu_time, cpu_perf, cpu_time;
double error, work[1];
double c_neg_one = MAGMA_D_NEG_ONE;
double *h_A, *h_B, *h_R;
magmaDouble_ptr d_A, d_B;
magma_int_t M, N, size, lda, ldb, ldda, lddb;
magma_int_t ione = 1;
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
magma_uplo_t uplo[] = { MagmaLower, MagmaUpper, MagmaFull };
printf("%% uplo M N CPU GByte/s (ms) GPU GByte/s (ms) check\n");
printf("%%================================================================\n");
for( int iuplo = 0; iuplo < 3; ++iuplo ) {
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
lda = M;
ldb = lda;
ldda = magma_roundup( M, opts.align ); // multiple of 32 by default
lddb = ldda;
size = lda*N;
if ( uplo[iuplo] == MagmaLower ) {
// load & save lower trapezoid (with diagonal)
if ( M > N ) {
gbytes = 2. * sizeof(double) * (1.*M*N - 0.5*N*(N-1)) / 1e9;
} else {
gbytes = 2. * sizeof(double) * 0.5*M*(M+1) / 1e9;
}
}
else if ( uplo[iuplo] == MagmaUpper ) {
// load & save upper trapezoid (with diagonal)
if ( N > M ) {
gbytes = 2. * sizeof(double) * (1.*M*N - 0.5*M*(M-1)) / 1e9;
} else {
gbytes = 2. * sizeof(double) * 0.5*N*(N+1) / 1e9;
}
}
else {
// load & save entire matrix
gbytes = 2. * sizeof(double) * 1.*M*N / 1e9;
}
TESTING_MALLOC_CPU( h_A, double, size );
TESTING_MALLOC_CPU( h_B, double, size );
TESTING_MALLOC_CPU( h_R, double, size );
TESTING_MALLOC_DEV( d_A, double, ldda*N );
TESTING_MALLOC_DEV( d_B, double, lddb*N );
/* Initialize the matrix */
for( int j = 0; j < N; ++j ) {
for( int i = 0; i < M; ++i ) {
h_A[i + j*lda] = MAGMA_D_MAKE( i + j/10000., j );
h_B[i + j*ldb] = MAGMA_D_MAKE( i - j/10000. + 10000., j );
}
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_dsetmatrix( M, N, h_A, lda, d_A, ldda, opts.queue );
magma_dsetmatrix( M, N, h_B, ldb, d_B, lddb, opts.queue );
gpu_time = magma_sync_wtime( opts.queue );
//magmablas_dlacpy( uplo[iuplo], M-2, N-2, d_A+1+ldda, ldda, d_B+1+lddb, lddb, opts.queue ); // inset by 1 row & col
magmablas_dlacpy( uplo[iuplo], M, N, d_A, ldda, d_B, lddb, opts.queue );
gpu_time = magma_sync_wtime( opts.queue ) - gpu_time;
gpu_perf = gbytes / gpu_time;
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
//magma_int_t M2 = M-2; // inset by 1 row & col
//magma_int_t N2 = N-2;
//lapackf77_dlacpy( lapack_uplo_const(uplo[iuplo]), &M2, &N2, h_A+1+lda, &lda, h_B+1+ldb, &ldb );
lapackf77_dlacpy( lapack_uplo_const(uplo[iuplo]), &M, &N, h_A, &lda, h_B, &ldb );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gbytes / cpu_time;
if ( opts.verbose ) {
printf( "A= " ); magma_dprint( M, N, h_A, lda );
printf( "B= " ); magma_dprint( M, N, h_B, ldb );
printf( "dA=" ); magma_dprint_gpu( M, N, d_A, ldda );
printf( "dB=" ); magma_dprint_gpu( M, N, d_B, lddb );
}
/* =====================================================================
Check the result
=================================================================== */
magma_dgetmatrix( M, N, d_B, lddb, h_R, lda, opts.queue );
blasf77_daxpy(&size, &c_neg_one, h_B, &ione, h_R, &ione);
error = lapackf77_dlange("f", &M, &N, h_R, &lda, work);
printf("%5s %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %s\n",
lapack_uplo_const(uplo[iuplo]), (int) M, (int) N,
cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000.,
(error == 0. ? "ok" : "failed") );
status += ! (error == 0.);
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
printf( "\n" );
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
void magmablas_dsymm_mgpu_spec(
char side, char uplo, magma_int_t m, magma_int_t n,
double alpha,
double *dA[], magma_int_t ldda, magma_int_t offset,
double *dB[], magma_int_t lddb,
double beta,
double *dC[], magma_int_t lddc,
double *dwork[], magma_int_t dworksiz,
double *C, magma_int_t ldc,
double *work[], magma_int_t ldwork,
magma_int_t ngpu, magma_int_t nb,
magma_queue_t streams[][20], magma_int_t nstream,
magma_event_t redevents[][MagmaMaxGPUs*MagmaMaxGPUs+10],magma_int_t nbevents,
magma_int_t gnode[MagmaMaxGPUs][MagmaMaxGPUs+2], magma_int_t nbcmplx )
{
#define dA(dev, i, j) (dA[dev] + (i) + (j)*ldda)
#define dB(dev, i, j) (dB[dev] + (i) + (j)*lddb)
#define dC(dev, i, j) (dC[dev] + (i) + (j)*lddc)
#define dwork(dev, i, j) (dwork[dev] + (i) + (j)*lddwork)
#define C(i, j) (C + (i) + (j)*ldc)
assert( ldda >= m );
assert( lddb >= m );
assert( lddc >= m );
assert( nstream >= ngpu );
assert( nbevents >= ngpu*ngpu );
double *dwork1[MagmaMaxGPUs];
double *dwork2[MagmaMaxGPUs];
magma_int_t lddwork = lddc;
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
dwork1[dev] = dwork[dev];
dwork2[dev] = dwork[dev]+n*lddwork;
}
assert( dworksiz >= (2*n*lddwork) );
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_t cstream;
magmablasGetKernelStream(&cstream);
magma_int_t dev,devperm,myblk,mycolsize,myblkoffst;
magma_int_t gdev,gcolsize,gmaster,gngpu;
magma_int_t masterdev,lcdev,lccolsize,myngpu;
magma_int_t stdev = (offset/nb)%ngpu;
magma_int_t blockoffset = offset % nb;
magma_int_t fstblksiz = 0;
if(blockoffset>0){
fstblksiz = min(m, (nb - blockoffset));
}
//magma_int_t nbblk = magma_ceildiv(m,nb);
magma_int_t nbblk = magma_ceildiv((m+blockoffset),nb);
magma_int_t maxgsize = n*nb*magma_ceildiv(nbblk,ngpu);
magma_int_t remm = m- fstblksiz;
magma_int_t nbblkoffst = offset/nb;
magma_int_t nblstblks = -1;
magma_int_t devlstblk = -1;
magma_int_t lstblksiz = remm%nb;
if(lstblksiz>0){
nblstblks = nbblk%ngpu;
devlstblk = (nblstblks-1+ngpu)%ngpu;
}
magma_int_t nbcmplxactive = 0;
magma_int_t cmplxisactive[MagmaMaxGPUs];
magma_int_t gpuisactive[MagmaMaxGPUs];
memset(gpuisactive, 0, MagmaMaxGPUs*sizeof(magma_int_t));
memset(cmplxisactive, 0, MagmaMaxGPUs*sizeof(magma_int_t));
//*******************************
// each GPU make a GEMM with the
// transpose of its blocks to compute
// a final portion of X=A*VT
//*******************************
/* dB = V*T already ==> dB' = T'*V'
* compute T'*V'*X is equal to compute locally (VT)'_i*X_i
* then each GPU broadcast its X_i to assemble the full X which is used
* to compute W = X - 0.5 * V * T'*V'*X = X - 0.5 * V *dwork3
*/
if(ngpu ==1){
magma_setdevice( 0 );
magmablasSetKernelStream( streams[ 0 ][ 0 ] );
// compute X[me] = A*VT = A[me]^tr *VT;
magma_dgemm( MagmaTrans, MagmaNoTrans, m, n, m,
alpha, dA(0,offset,offset), ldda,
dB[0], lddb,
beta, dC[0], lddc );
return;
}
//ngpu>1
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
masterdev = -1;
gnode[cmplxid][MagmaMaxGPUs+1] = -1;
myngpu = gnode[cmplxid][MagmaMaxGPUs];
for( magma_int_t idev = 0; idev < myngpu; ++idev ) {
dev = gnode[cmplxid][idev];
devperm = (dev-stdev+ngpu)%ngpu;
myblk = (nbblk/ngpu) + (nbblk%ngpu > devperm ? 1:0 );
mycolsize = myblk*nb;
myblkoffst = nb*((nbblkoffst/ngpu)+(nbblkoffst%ngpu > dev?1:0));
if(dev==stdev){
mycolsize -= blockoffset;
myblkoffst += blockoffset; // local index in parent matrix
}
if((devperm==devlstblk)&&(lstblksiz>0)){
mycolsize -= (nb-(remm%nb));
}
mycolsize = min(mycolsize,m);
if(mycolsize>0){
if(masterdev==-1) masterdev = dev;
//printf("dev %d devperm %d on cmplx %d master %d nbblk %d myblk %d m %d n %d mycolsize %d stdev %d fstblksize %d lastdev %d lastsize %d dA(%d,%d,%d) ==> dwork(%d,%d)\n",dev,devperm,cmplxid,masterdev,nbblk,myblk,m,n,mycolsize,stdev,fstblksiz,devlstblk,remm%nb,dev,offset,myblkoffst,dev,maxgsize*dev);
gpuisactive[dev] = mycolsize;
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ dev ] );
magma_dgemm( MagmaTrans, MagmaNoTrans, mycolsize, n, m,
alpha, dA(dev,offset,myblkoffst), ldda,
dB(dev,0,0), lddb,
beta, &dwork[dev][maxgsize*dev], mycolsize );
magma_event_record(redevents[dev][dev*ngpu+dev], streams[dev][dev]);
}
if(dev == masterdev){
nbcmplxactive = nbcmplxactive +1;
cmplxisactive[cmplxid] = 1;
gnode[cmplxid][MagmaMaxGPUs+1] = masterdev;
}
}
}
/*
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_queue_sync( streams[ dev ][ dev ] );
}
*/
//*******************************
// each Master GPU has the final
// result either by receiving
// from CPU of by making the add
// by himself, so now it is time
// to broadcast over the GPUs of
// its board.
//*******************************
//printf("=======================================================================\n");
//printf(" sending \n");
//printf("=======================================================================\n");
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
myngpu = gnode[cmplxid][MagmaMaxGPUs];
masterdev = gnode[cmplxid][MagmaMaxGPUs+1];
for( magma_int_t idev = 0; idev < myngpu; ++idev ) {
dev = gnode[cmplxid][idev];
mycolsize = gpuisactive[dev];
if(mycolsize>0){
// I am an active GPU send my portion local
// to all active gpu of my cmplex and global to the
// active master of the other real and they should
// send it out to their actives slaves.
magma_setdevice( dev );
//==============================================
// sending to the master of the active real
//==============================================
//printf ("\n\n**************GPU %d\n ",dev);
//printf (" GPU %d sending to cmplx masters\n",dev);
for( magma_int_t k = 0; k < nbcmplx; ++k ) {
if(k!=cmplxid){
gmaster = gnode[k][MagmaMaxGPUs+1];
if(gmaster!=-1){ //real is active
//printf (" device %d from cmplx %d is sending to master %d on cmplx %d block of size %d event %d\n",dev,cmplxid,gmaster,k,mycolsize,redevents[dev][gmaster*ngpu+dev]);
magma_queue_wait_event(streams[ dev ][ gmaster ], redevents[dev][dev*ngpu+dev]);
cudaMemcpy2DAsync(&dwork[gmaster][maxgsize*dev], mycolsize*sizeof(double),
&dwork[dev][maxgsize*dev], mycolsize*sizeof(double),
mycolsize*sizeof(double), n,
cudaMemcpyDeviceToDevice, streams[dev][gmaster]);
magma_event_record(redevents[dev][gmaster*ngpu+dev], streams[dev][gmaster]);
}
}
}
//==============================================
//
//==============================================
// sending to the active GPUs of my real
//==============================================
//printf (" GPU %d sending internal\n",dev);
for( magma_int_t l = 0; l < myngpu; ++l ) {
lcdev = gnode[cmplxid][l];
lccolsize = gpuisactive[lcdev];
if((lcdev!=dev)&&(lccolsize>0)){
//printf (" device %d from cmplx %d is sending internal to dev %d block of size %d event %d\n",dev,cmplxid,lcdev,mycolsize,redevents[dev][lcdev*ngpu+dev]);
magma_queue_wait_event(streams[ dev ][ lcdev ], redevents[dev][dev*ngpu+dev]);
cudaMemcpy2DAsync(&dwork[lcdev][maxgsize*dev], mycolsize*sizeof(double),
&dwork[dev][maxgsize*dev], mycolsize*sizeof(double),
mycolsize*sizeof(double), n,
cudaMemcpyDeviceToDevice, streams[dev][lcdev]);
magma_event_record(redevents[dev][lcdev*ngpu+dev], streams[dev][lcdev]);
}
}
//==============================================
}// end if mycolsize>0
}// for idev
}// for cmplxid
//printf("=======================================================================\n");
//printf(" master wait and resend internally \n");
//printf("=======================================================================\n");
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
myngpu = gnode[cmplxid][MagmaMaxGPUs];
masterdev = gnode[cmplxid][MagmaMaxGPUs+1];
//==============================================
// if I am active master so wait receiving contribution
// of the GPUs of other real and send it locally
//==============================================
if(masterdev != -1){
mycolsize = gpuisactive[masterdev];
magma_setdevice( masterdev );
//printf(" GPU %d distributing internal\n",masterdev);
for( magma_int_t k = 0; k < nbcmplx; ++k ) {
if(k!=cmplxid){
gngpu = gnode[k][MagmaMaxGPUs];
for( magma_int_t g = 0; g < gngpu; ++g ) {
gdev = gnode[k][g];
gcolsize = gpuisactive[gdev];
// check if I received from this GPU,
// if yes send it to my group
if(gcolsize>0){
magma_queue_wait_event(streams[ masterdev ][ gdev ], redevents[gdev][masterdev*ngpu+gdev]);
for( magma_int_t l = 0; l < myngpu; ++l ) {
lcdev = gnode[cmplxid][l];
lccolsize = gpuisactive[lcdev];
if((lcdev!=masterdev)&&(lccolsize>0)){
//printf(" Master %d on cmplx %d waiting on event %d is distributing internal results of %d to lcdev %d block of size %d event %d\n", masterdev,cmplxid,redevents[gdev][masterdev*ngpu+gdev],gdev,lcdev,gcolsize,redevents[masterdev][lcdev*ngpu+gdev]);
cudaMemcpy2DAsync(&dwork[lcdev][maxgsize*gdev], gcolsize*sizeof(double),
&dwork[masterdev][maxgsize*gdev], gcolsize*sizeof(double),
gcolsize*sizeof(double), n,
cudaMemcpyDeviceToDevice, streams[masterdev][gdev]);
magma_event_record(redevents[masterdev][lcdev*ngpu+gdev], streams[masterdev][gdev]);
}
}
}
}
}
}
}// if active master
//==============================================
}// for cmplxid
/*
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_queue_sync( streams[ dev ][ 0 ] );
for( magma_int_t s = 0; s < ngpu; ++s ) {
magma_queue_sync( streams[ dev ][ s ] );
}
}
*/
//printf("=======================================================================\n");
//printf(" distributing \n");
//printf("=======================================================================\n");
magma_int_t lcblki,gbblki,gblk,ib;
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
myngpu = gnode[cmplxid][MagmaMaxGPUs];
masterdev = gnode[cmplxid][MagmaMaxGPUs+1];
for( magma_int_t idev = 0; idev < myngpu; ++idev ) {
dev = gnode[cmplxid][idev];
mycolsize = gpuisactive[dev];
if(mycolsize>0){ // I am an active GPU
//printf("\n\n==============GPU %d collecting\n",dev);
magma_setdevice( dev );
// collect my results first as tyhere is no need to wait to
// receive nothing, just wait that my gemm are done.
// in theory this should be inside the loop but cuda was not
// able to run it first for all gpu and on gpu>0 it was waiting
// however it was on different stream so it should run. but maybe
// this is because there are too many function call and this make
// cuda not handleit so nice. anyway it coul dbe removed when cuda
// is able to lunch it first without wait.
gdev = dev;
gcolsize = gpuisactive[gdev];
if(gcolsize>0){
devperm = (gdev-stdev+ngpu)%ngpu;
gblk = (nbblk/ngpu) + (nbblk%ngpu > devperm ? 1:0 );
magmablasSetKernelStream( streams[ dev ][ gdev ] );
magma_queue_wait_event(streams[ dev ][ gdev ], redevents[gdev][dev*ngpu+gdev]);
//printf (" GPU %d stream %d doing dlacpy\n",dev,streams[ dev ][ gdev ]);
for( magma_int_t blki = 0; blki < gblk; ++blki){
gbblki = (blki*ngpu + devperm)*nb - blockoffset;
lcblki = blki*nb;
ib = nb;//min(nb,m-gbblki);
if(gdev==stdev){
lcblki = blki*nb-blockoffset;
if(blki==0){
gbblki = 0;
lcblki = 0;
ib = nb-blockoffset;
}
}
ib = min(ib,m-gbblki);
//printf(" blockoffset %d nbblk %d stdev %d receiving from gdev %d gblk %d gcolsize %d copying blki %d of size ibxn %dx%d from work[%d] to C[%d]\n", blockoffset,nbblk,stdev,gdev,gblk,gcolsize,blki,ib,n,lcblki,gbblki);
magmablas_dlacpy( 'A', ib, n, &dwork[dev][maxgsize*gdev+lcblki], gcolsize, &dC[dev][gbblki], lddc);
}// end blki
}
for( magma_int_t k = 0; k < nbcmplx; ++k ) {
gngpu = gnode[k][MagmaMaxGPUs];
for( magma_int_t g = 0; g < gngpu; ++g ) {
gdev = gnode[k][g];
gcolsize = gpuisactive[gdev];
// if gcolsize>0, ==> gpu gdev was active and so
// I received from him/computed a portion of dwork,
// so go over its gblk and distribute it on dC.
if(gdev!=dev){
if(gcolsize>0){
devperm = (gdev-stdev+ngpu)%ngpu;
gblk = (nbblk/ngpu) + (nbblk%ngpu > devperm ? 1:0 );
magmablasSetKernelStream( streams[ dev ][ gdev ] );
if(k==cmplxid){
//we are on the same group so wait on event issued by gdev for me citing his id
magma_queue_wait_event(streams[ dev ][ gdev ], redevents[gdev][dev*ngpu+gdev]);
//printf (" GPU %d stream %d waiting on event %d to collecte from %d the size of gcolsize %d\n",dev,streams[ dev ][ gdev ],redevents[gdev][dev*ngpu+gdev],gdev,gcolsize);
}else{
//we are on different group so:
//if I am the master wait on the event issued by gdev for me citing his id
//else wait event issued by my master for me on the behalf of gdev
//printf (" GPU %d stream %d waiting on event %d to collecte from %d the size of gcolsize %d\n",dev,streams[ dev ][ gdev ],redevents[masterdev][dev*ngpu+gdev],gdev,gcolsize);
if(dev==masterdev)
magma_queue_wait_event(streams[ dev ][ gdev ], redevents[gdev][dev*ngpu+gdev]);
else
magma_queue_wait_event(streams[ dev ][ gdev ], redevents[masterdev][dev*ngpu+gdev]);
}
//printf (" GPU %d stream %d doing dlacpy\n",dev,streams[ dev ][ gdev ]);
for( magma_int_t blki = 0; blki < gblk; ++blki){
gbblki = (blki*ngpu + devperm)*nb - blockoffset;
lcblki = blki*nb;
ib = nb;//min(nb,m-gbblki);
if(gdev==stdev){
lcblki = blki*nb-blockoffset;
if(blki==0){
gbblki = 0;
lcblki = 0;
ib = nb-blockoffset;
}
}
ib = min(ib,m-gbblki);
//printf(" blockoffset %d nbblk %d stdev %d receiving from gdev %d gblk %d gcolsize %d copying blki %d of size ibxn %dx%d from work[%d] to C[%d]\n", blockoffset,nbblk,stdev,gdev,gblk,gcolsize,blki,ib,n,lcblki,gbblki);
magmablas_dlacpy( 'A', ib, n, &dwork[dev][maxgsize*gdev+lcblki], gcolsize, &dC[dev][gbblki], lddc);
}// end blki
}// en gcolsize>0 meaning gdev is active
} // end if gdev != dev
}// end loop over the g gpus of the cmplx k
}//end loop over the real k
}// end mycolsize>0 meaning that I am active
}// end loop over idev of cmplxid
}// end loop of the cmplx
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
cudaDeviceSynchronize();
}
// put back the input gpu and its input stream
magma_setdevice( cdev );
magmablasSetKernelStream( cstream );
}
/**
Purpose
-------
DSGESV computes the solution to a real system of linear equations
A * X = B, A**T * X = B, or A**H * X = B,
where A is an N-by-N matrix and X and B are N-by-NRHS matrices.
DSGESV first attempts to factorize the matrix in real SINGLE PRECISION
and use this factorization within an iterative refinement procedure
to produce a solution with real DOUBLE PRECISION norm-wise backward error
quality (see below). If the approach fails the method switches to a
real DOUBLE PRECISION factorization and solve.
The iterative refinement is not going to be a winning strategy if
the ratio real SINGLE PRECISION performance over real DOUBLE PRECISION
performance is too small. A reasonable strategy should take the
number of right-hand sides and the size of the matrix into account.
This might be done with a call to ILAENV in the future. Up to now, we
always try iterative refinement.
The iterative refinement process is stopped if
ITER > ITERMAX
or for all the RHS we have:
RNRM < SQRT(N)*XNRM*ANRM*EPS*BWDMAX
where
o ITER is the number of the current iteration in the iterative
refinement process
o RNRM is the infinity-norm of the residual
o XNRM is the infinity-norm of the solution
o ANRM is the infinity-operator-norm of the matrix A
o EPS is the machine epsilon returned by DLAMCH('Epsilon')
The value ITERMAX and BWDMAX are fixed to 30 and 1.0D+00 respectively.
Arguments
---------
@param[in]
trans magma_trans_t
Specifies the form of the system of equations:
- = MagmaNoTrans: A * X = B (No transpose)
- = MagmaTrans: A**T * X = B (Transpose)
- = MagmaConjTrans: A**H * X = B (Conjugate transpose)
@param[in]
n INTEGER
The number of linear equations, i.e., the order of the
matrix A. N >= 0.
@param[in]
nrhs INTEGER
The number of right hand sides, i.e., the number of columns
of the matrix B. NRHS >= 0.
@param[in,out]
dA DOUBLE PRECISION array on the GPU, dimension (ldda,N)
On entry, the N-by-N coefficient matrix A.
On exit, if iterative refinement has been successfully used
(info.EQ.0 and ITER.GE.0, see description below), A is
unchanged. If double precision factorization has been used
(info.EQ.0 and ITER.LT.0, see description below), then the
array dA contains the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
@param[in]
ldda INTEGER
The leading dimension of the array dA. ldda >= max(1,N).
@param[out]
ipiv INTEGER array, dimension (N)
The pivot indices that define the permutation matrix P;
row i of the matrix was interchanged with row IPIV(i).
Corresponds either to the single precision factorization
(if info.EQ.0 and ITER.GE.0) or the double precision
factorization (if info.EQ.0 and ITER.LT.0).
@param[out]
dipiv INTEGER array on the GPU, dimension (N)
The pivot indices; for 1 <= i <= N, after permuting, row i of the
matrix was moved to row dIPIV(i).
Note this is different than IPIV, where interchanges
are applied one-after-another.
@param[in]
dB DOUBLE PRECISION array on the GPU, dimension (lddb,NRHS)
The N-by-NRHS right hand side matrix B.
@param[in]
lddb INTEGER
The leading dimension of the array dB. lddb >= max(1,N).
@param[out]
dX DOUBLE PRECISION array on the GPU, dimension (lddx,NRHS)
If info = 0, the N-by-NRHS solution matrix X.
@param[in]
lddx INTEGER
The leading dimension of the array dX. lddx >= max(1,N).
@param
dworkd (workspace) DOUBLE PRECISION array on the GPU, dimension (N*NRHS)
This array is used to hold the residual vectors.
@param
dworks (workspace) SINGLE PRECISION array on the GPU, dimension (N*(N+NRHS))
This array is used to store the real single precision matrix
and the right-hand sides or solutions in single precision.
@param[out]
iter INTEGER
- < 0: iterative refinement has failed, double precision
factorization has been performed
+ -1 : the routine fell back to full precision for
implementation- or machine-specific reasons
+ -2 : narrowing the precision induced an overflow,
the routine fell back to full precision
+ -3 : failure of SGETRF
+ -31: stop the iterative refinement after the 30th iteration
- > 0: iterative refinement has been successfully used.
Returns the number of iterations
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if info = -i, the i-th argument had an illegal value
- > 0: if info = i, U(i,i) computed in DOUBLE PRECISION is
exactly zero. The factorization has been completed,
but the factor U is exactly singular, so the solution
could not be computed.
@ingroup magma_dgesv_driver
********************************************************************/
extern "C" magma_int_t
magma_dsgesv_gpu(magma_trans_t trans, magma_int_t n, magma_int_t nrhs,
double *dA, magma_int_t ldda,
magma_int_t *ipiv, magma_int_t *dipiv,
double *dB, magma_int_t lddb,
double *dX, magma_int_t lddx,
double *dworkd, float *dworks,
magma_int_t *iter, magma_int_t *info)
{
#define dB(i,j) (dB + (i) + (j)*lddb)
#define dX(i,j) (dX + (i) + (j)*lddx)
#define dR(i,j) (dR + (i) + (j)*lddr)
double c_neg_one = MAGMA_D_NEG_ONE;
double c_one = MAGMA_D_ONE;
magma_int_t ione = 1;
double *dR;
float *dSA, *dSX;
double Xnrmv, Rnrmv;
double Anrm, Xnrm, Rnrm, cte, eps;
magma_int_t i, j, iiter, lddsa, lddr;
/* Check arguments */
*iter = 0;
*info = 0;
if ( n < 0 )
*info = -1;
else if ( nrhs < 0 )
*info = -2;
else if ( ldda < max(1,n))
*info = -4;
else if ( lddb < max(1,n))
*info = -8;
else if ( lddx < max(1,n))
*info = -10;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if ( n == 0 || nrhs == 0 )
return *info;
lddsa = n;
lddr = n;
dSA = dworks;
dSX = dSA + lddsa*n;
dR = dworkd;
eps = lapackf77_dlamch("Epsilon");
Anrm = magmablas_dlange(MagmaInfNorm, n, n, dA, ldda, (double*)dworkd );
cte = Anrm * eps * pow((double)n, 0.5) * BWDMAX;
/*
* Convert to single precision
*/
//magmablas_dlag2s( n, nrhs, dB, lddb, dSX, lddsx, info ); // done inside dsgetrs with pivots
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
magmablas_dlag2s( n, n, dA, ldda, dSA, lddsa, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
// factor dSA in single precision
magma_sgetrf_gpu( n, n, dSA, lddsa, ipiv, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// Generate parallel pivots
{
magma_int_t *newipiv;
magma_imalloc_cpu( &newipiv, n );
if ( newipiv == NULL ) {
*iter = -3;
goto FALLBACK;
}
swp2pswp( trans, n, ipiv, newipiv );
magma_setvector( n, sizeof(magma_int_t), newipiv, 1, dipiv, 1 );
magma_free_cpu( newipiv );
}
// solve dSA*dSX = dB in single precision
// converts dB to dSX and applies pivots, solves, then converts result back to dX
magma_dsgetrs_gpu( trans, n, nrhs, dSA, lddsa, dipiv, dB, lddb, dX, lddx, dSX, info );
// residual dR = dB - dA*dX in double precision
magmablas_dlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dR, lddr );
if ( nrhs == 1 ) {
magma_dgemv( trans, n, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_dgemm( trans, MagmaNoTrans, n, nrhs, n,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
// TODO: use MAGMA_D_ABS( dX(i,j) ) instead of dlange?
for( j=0; j < nrhs; j++ ) {
i = magma_idamax( n, dX(0,j), 1) - 1;
magma_dgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_dlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_idamax ( n, dR(0,j), 1 ) - 1;
magma_dgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_dlange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
if ( Rnrm > Xnrm*cte ) {
goto REFINEMENT;
}
}
*iter = 0;
return *info;
REFINEMENT:
for( iiter=1; iiter < ITERMAX; ) {
*info = 0;
// convert residual dR to single precision dSX
// solve dSA*dSX = R in single precision
// convert result back to double precision dR
// it's okay that dR is used for both dB input and dX output.
magma_dsgetrs_gpu( trans, n, nrhs, dSA, lddsa, dipiv, dR, lddr, dR, lddr, dSX, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// Add correction and setup residual
// dX += dR --and--
// dR = dB
// This saves going through dR a second time (if done with one more kernel).
// -- not really: first time is read, second time is write.
for( j=0; j < nrhs; j++ ) {
magmablas_daxpycp( n, dR(0,j), dX(0,j), dB(0,j) );
}
// residual dR = dB - dA*dX in double precision
if ( nrhs == 1 ) {
magma_dgemv( trans, n, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_dgemm( trans, MagmaNoTrans, n, nrhs, n,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
/* Check whether the nrhs normwise backward errors satisfy the
* stopping criterion. If yes, set ITER=IITER > 0 and return. */
for( j=0; j < nrhs; j++ ) {
i = magma_idamax( n, dX(0,j), 1) - 1;
magma_dgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_dlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_idamax ( n, dR(0,j), 1 ) - 1;
magma_dgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_dlange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
if ( Rnrm > Xnrm*cte ) {
goto L20;
}
}
/* If we are here, the nrhs normwise backward errors satisfy
* the stopping criterion, we are good to exit. */
*iter = iiter;
return *info;
L20:
iiter++;
}
/* If we are at this place of the code, this is because we have
* performed ITER=ITERMAX iterations and never satisified the
* stopping criterion. Set up the ITER flag accordingly and follow
* up on double precision routine. */
*iter = -ITERMAX - 1;
FALLBACK:
/* Single-precision iterative refinement failed to converge to a
* satisfactory solution, so we resort to double precision. */
magma_dgetrf_gpu( n, n, dA, ldda, ipiv, info );
if (*info == 0) {
magmablas_dlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dX, lddx );
magma_dgetrs_gpu( trans, n, nrhs, dA, ldda, ipiv, dX, lddx, info );
}
return *info;
}
/**
Purpose
-------
DGETRI computes the inverse of a matrix using the LU factorization
computed by DGETRF. This method inverts U and then computes inv(A) by
solving the system inv(A)*L = inv(U) for inv(A).
Note that it is generally both faster and more accurate to use DGESV,
or DGETRF and DGETRS, to solve the system AX = B, rather than inverting
the matrix and multiplying to form X = inv(A)*B. Only in special
instances should an explicit inverse be computed with this routine.
Arguments
---------
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
dA DOUBLE_PRECISION array on the GPU, dimension (LDDA,N)
On entry, the factors L and U from the factorization
A = P*L*U as computed by DGETRF_GPU.
On exit, if INFO = 0, the inverse of the original matrix A.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,N).
@param[in]
ipiv INTEGER array, dimension (N)
The pivot indices from DGETRF; for 1 <= i <= N, row i of the
matrix was interchanged with row IPIV(i).
@param[out]
dwork (workspace) DOUBLE_PRECISION array on the GPU, dimension (MAX(1,LWORK))
@param[in]
lwork INTEGER
The dimension of the array DWORK. LWORK >= N*NB, where NB is
the optimal blocksize returned by magma_get_dgetri_nb(n).
\n
Unlike LAPACK, this version does not currently support a
workspace query, because the workspace is on the GPU.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: if INFO = i, U(i,i) is exactly zero; the matrix is
singular and its cannot be computed.
@ingroup magma_dgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_dgetri_gpu( magma_int_t n, double *dA, magma_int_t ldda,
magma_int_t *ipiv, double *dwork, magma_int_t lwork,
magma_int_t *info )
{
#define dA(i, j) (dA + (i) + (j)*ldda)
#define dL(i, j) (dL + (i) + (j)*lddl)
/* Local variables */
double c_zero = MAGMA_D_ZERO;
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
double *dL = dwork;
magma_int_t lddl = n;
magma_int_t nb = magma_get_dgetri_nb(n);
magma_int_t j, jmax, jb, jp;
*info = 0;
if (n < 0)
*info = -1;
else if (ldda < max(1,n))
*info = -3;
else if ( lwork < n*nb )
*info = -6;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if ( n == 0 )
return *info;
/* Invert the triangular factor U */
magma_dtrtri_gpu( MagmaUpper, MagmaNonUnit, n, dA, ldda, info );
if ( *info != 0 )
return *info;
jmax = ((n-1) / nb)*nb;
for( j = jmax; j >= 0; j -= nb ) {
jb = min( nb, n-j );
// copy current block column of A to work space dL
// (only needs lower trapezoid, but we also copy upper triangle),
// then zero the strictly lower trapezoid block column of A.
magmablas_dlacpy( MagmaFull, n-j, jb,
dA(j,j), ldda,
dL(j,0), lddl );
magmablas_dlaset( MagmaLower, n-j-1, jb, c_zero, c_zero, dA(j+1,j), ldda );
// compute current block column of Ainv
// Ainv(:, j:j+jb-1)
// = ( U(:, j:j+jb-1) - Ainv(:, j+jb:n) L(j+jb:n, j:j+jb-1) )
// * L(j:j+jb-1, j:j+jb-1)^{-1}
// where L(:, j:j+jb-1) is stored in dL.
if ( j+jb < n ) {
magma_dgemm( MagmaNoTrans, MagmaNoTrans, n, jb, n-j-jb,
c_neg_one, dA(0,j+jb), ldda,
dL(j+jb,0), lddl,
c_one, dA(0,j), ldda );
}
// TODO use magmablas work interface
magma_dtrsm( MagmaRight, MagmaLower, MagmaNoTrans, MagmaUnit,
n, jb, c_one,
dL(j,0), lddl,
dA(0,j), ldda );
}
// Apply column interchanges
for( j = n-2; j >= 0; --j ) {
jp = ipiv[j] - 1;
if ( jp != j ) {
magmablas_dswap( n, dA(0,j), 1, dA(0,jp), 1 );
}
}
return *info;
}
