/**
Purpose
-------
ZCGESV computes the solution to a complex 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.
ZCGESV first attempts to factorize the matrix in complex SINGLE PRECISION
and use this factorization within an iterative refinement procedure
to produce a solution with complex DOUBLE PRECISION norm-wise backward error
quality (see below). If the approach fails the method switches to a
complex DOUBLE PRECISION factorization and solve.
The iterative refinement is not going to be a winning strategy if
the ratio complex SINGLE PRECISION performance over complex 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 COMPLEX_16 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 COMPLEX_16 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 COMPLEX_16 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) COMPLEX_16 array on the GPU, dimension (N*NRHS)
This array is used to hold the residual vectors.
@param
dworks (workspace) COMPLEX array on the GPU, dimension (N*(N+NRHS))
This array is used to store the complex 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_zgesv_driver
********************************************************************/
extern "C" magma_int_t
magma_zcgesv_gpu(magma_trans_t trans, magma_int_t n, magma_int_t nrhs,
magmaDoubleComplex *dA, magma_int_t ldda,
magma_int_t *ipiv, magma_int_t *dipiv,
magmaDoubleComplex *dB, magma_int_t lddb,
magmaDoubleComplex *dX, magma_int_t lddx,
magmaDoubleComplex *dworkd, magmaFloatComplex *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)
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magma_int_t ione = 1;
magmaDoubleComplex *dR;
magmaFloatComplex *dSA, *dSX;
magmaDoubleComplex 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_zlange(MagmaInfNorm, n, n, dA, ldda, (double*)dworkd );
cte = Anrm * eps * pow((double)n, 0.5) * BWDMAX;
/*
* Convert to single precision
*/
//magmablas_zlag2c( n, nrhs, dB, lddb, dSX, lddsx, info ); // done inside zcgetrs with pivots
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
magmablas_zlag2c( n, n, dA, ldda, dSA, lddsa, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
// factor dSA in single precision
magma_cgetrf_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_zcgetrs_gpu( trans, n, nrhs, dSA, lddsa, dipiv, dB, lddb, dX, lddx, dSX, info );
// residual dR = dB - dA*dX in double precision
magmablas_zlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dR, lddr );
if ( nrhs == 1 ) {
magma_zgemv( trans, n, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_zgemm( trans, MagmaNoTrans, n, nrhs, n,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
// TODO: use MAGMA_Z_ABS( dX(i,j) ) instead of zlange?
for( j=0; j < nrhs; j++ ) {
i = magma_izamax( n, dX(0,j), 1) - 1;
magma_zgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_zlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_izamax ( n, dR(0,j), 1 ) - 1;
magma_zgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_zlange( "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_zcgetrs_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_zaxpycp( n, dR(0,j), dX(0,j), dB(0,j) );
}
// residual dR = dB - dA*dX in double precision
if ( nrhs == 1 ) {
magma_zgemv( trans, n, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_zgemm( 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_izamax( n, dX(0,j), 1) - 1;
magma_zgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_zlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_izamax ( n, dR(0,j), 1 ) - 1;
magma_zgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_zlange( "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_zgetrf_gpu( n, n, dA, ldda, ipiv, info );
if (*info == 0) {
magmablas_zlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dX, lddx );
magma_zgetrs_gpu( trans, n, nrhs, dA, ldda, ipiv, dX, lddx, info );
}
return *info;
}
extern "C" magma_int_t
magma_zcposv_gpu(char uplo, magma_int_t n, magma_int_t nrhs,
magmaDoubleComplex *dA, magma_int_t ldda,
magmaDoubleComplex *dB, magma_int_t lddb,
magmaDoubleComplex *dX, magma_int_t lddx,
magmaDoubleComplex *dworkd, magmaFloatComplex *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
=======
ZCPOSV computes the solution to a complex system of linear equations
A * X = B,
where A is an N-by-N Hermitian positive definite matrix and X and B
are N-by-NRHS matrices.
ZCPOSV first attempts to factorize the matrix in complex SINGLE PRECISION
and use this factorization within an iterative refinement procedure
to produce a solution with complex DOUBLE PRECISION norm-wise backward error
quality (see below). If the approach fails the method switches to a
complex DOUBLE PRECISION factorization and solve.
The iterative refinement is not going to be a winning strategy if
the ratio complex SINGLE PRECISION performance over complex 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
=========
UPLO (input) CHARACTER
= 'U': Upper triangle of A is stored;
= 'L': Lower triangle of A is stored.
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) COMPLEX_16 array on the GPU, dimension (LDDA,N)
On entry, the Hermitian matrix A. If UPLO = 'U', 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 = 'L', 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.
LDDA (input) INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
dB (input) COMPLEX_16 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) COMPLEX_16 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) COMPLEX_16 array on the GPU, dimension (N*NRHS)
This array is used to hold the residual vectors.
dworks (workspace) COMPLEX array on the GPU, dimension (N*(N+NRHS))
This array is used to store the complex 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 SPOTRF
-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, 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.
===================================================================== */
#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)
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magma_int_t ione = 1;
magmaDoubleComplex *dR;
magmaFloatComplex *dSA, *dSX;
magmaDoubleComplex 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;
eps = lapackf77_dlamch("Epsilon");
Anrm = magmablas_zlanhe('I', uplo, n, dA, ldda, (double*)dworkd );
cte = Anrm * eps * pow((double)n, 0.5) * BWDMAX;
/*
* Convert to single precision
*/
magmablas_zlag2c( n, nrhs, dB, lddb, dSX, lddsx, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
magmablas_zlat2c( uplo, n, dA, ldda, dSA, lddsa, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
// factor dSA in single precision
magma_cpotrf_gpu( uplo, n, dSA, lddsa, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// solve dSA*dSX = dB in single precision
magma_cpotrs_gpu( uplo, n, nrhs, dSA, lddsa, dSX, lddsx, info );
// residual dR = dB - dA*dX in double precision
magmablas_clag2z( n, nrhs, dSX, lddsx, dX, lddx, info );
magmablas_zlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dR, lddr );
if ( nrhs == 1 ) {
magma_zhemv( uplo, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_zhemm( MagmaLeft, uplo, n, nrhs,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
// TODO: use MAGMA_Z_ABS( dX(i,j) ) instead of zlange?
for( j=0; j < nrhs; j++ ) {
i = magma_izamax( n, dX(0,j), 1) - 1;
magma_zgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_zlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_izamax ( n, dR(0,j), 1 ) - 1;
magma_zgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_zlange( "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
magmablas_zlag2c( n, nrhs, dR, lddr, dSX, lddsx, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
// solve dSA*dSX = R in single precision
magma_cpotrs_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_zcaxpycp( n, dSX(0,j), dX(0,j), dB(0,j), dR(0,j) );
}
// residual dR = dB - dA*dX in double precision
if ( nrhs == 1 ) {
magma_zhemv( uplo, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_zhemm( MagmaLeft, uplo, n, nrhs,
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_izamax( n, dX(0,j), 1) - 1;
magma_zgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_zlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_izamax ( n, dR(0,j), 1 ) - 1;
magma_zgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_zlange( "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_zpotrf_gpu( uplo, n, dA, ldda, info );
if (*info == 0) {
magmablas_zlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dX, lddx );
magma_zpotrs_gpu( uplo, n, nrhs, dA, ldda, dX, lddx, info );
}
return *info;
}
/**
Purpose
-------
ZCGEQRSV 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.
ZCGEQRSV first attempts to factorize the matrix in complex SINGLE PRECISION
and use this factorization within an iterative refinement procedure
to produce a solution with complex DOUBLE PRECISION norm-wise backward error
quality (see below). If the approach fails the method switches to a
complex DOUBLE PRECISION factorization and solve.
The iterative refinement is not going to be a winning strategy if
the ratio complex SINGLE PRECISION performance over complex 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 COMPLEX_16 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 COMPLEX_16 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 COMPLEX_16 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_zgels_driver
********************************************************************/
extern "C" magma_int_t
magma_zcgeqrsv_gpu(
magma_int_t m, magma_int_t n, magma_int_t nrhs,
magmaDoubleComplex_ptr dA, magma_int_t ldda,
magmaDoubleComplex_ptr dB, magma_int_t lddb,
magmaDoubleComplex_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)
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magma_int_t ione = 1;
magmaDoubleComplex *hworkd;
magmaFloatComplex *hworks;
magmaDoubleComplex *tau;
magmaFloatComplex *stau;
magmaDoubleComplex_ptr dworkd;
magmaFloatComplex_ptr dworks;
magmaDoubleComplex_ptr dR, dT;
magmaFloatComplex_ptr dSA, dSX, dST;
magmaDoubleComplex 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_cgeqrf_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_cmalloc( &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_zmalloc( &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_cmalloc_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_zlange(MagmaInfNorm, m, n, dA, ldda, (double*)dworkd );
cte = Anrm * eps * pow((double)n, 0.5) * BWDMAX;
/*
* Convert to single precision
*/
magmablas_zlag2c( m, nrhs, dB, lddb, dSX, lddsx, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
magmablas_zlag2c( m, n, dA, ldda, dSA, lddsa, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
// factor dSA in single precision
magma_cgeqrf_gpu( m, n, dSA, lddsa, stau, dST, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// solve dSA*dSX = dB in single precision
magma_cgeqrs_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_clag2z( n, nrhs, dSX, lddsx, dX, lddx, info );
magmablas_zlacpy( MagmaUpperLower, m, nrhs, dB, lddb, dR, lddr );
if ( nrhs == 1 ) {
magma_zgemv( MagmaNoTrans, m, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, nrhs, n,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
// TODO: use MAGMA_Z_ABS( dX(i,j) ) instead of zlange?
for( j=0; j < nrhs; j++ ) {
i = magma_izamax( n, dX(0,j), 1) - 1;
magma_zgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_zlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_izamax ( m, dR(0,j), 1 ) - 1;
magma_zgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_zlange( "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_zlag2c( m, nrhs, dR, lddr, dSX, lddsx, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
// solve dSA*dSX = R in single precision
magma_cgeqrs_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_zcaxpycp( n, dSX(0,j), dX(0,j), dB(0,j), dR(0,j) );
}
// dR = dB (whole m rows)
magmablas_zlacpy( MagmaUpperLower, m, nrhs, dB, lddb, dR, lddr );
// residual dR = dB - dA*dX in double precision
if ( nrhs == 1 ) {
magma_zgemv( MagmaNoTrans, m, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_zgemm( 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_izamax( n, dX(0,j), 1) - 1;
magma_zgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_zlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_izamax ( m, dR(0,j), 1 ) - 1;
magma_zgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_zlange( "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 zgeqrf */
nb = magma_get_zgeqrf_nb( m );
size = (2*min(m, n) + (n+31)/32*32 )*nb;
if ( size > ldworkd ) {
magma_free( dworkd );
if (MAGMA_SUCCESS != magma_zmalloc( &dworkd, size )) {
fprintf(stderr, "Allocation of dworkd2 failed\n");
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
}
dT = dworkd;
/* hworkd(dtau + workspace for zgeqrs) = min(m,n) + lhwork */
size = lhwork + minmn;
magma_zmalloc_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_zgeqrf_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_zgeqrs_gpu( m, n, nrhs, dA, ldda, tau, dT, dB, lddb, hworkd, lhwork, info );
magmablas_zlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dX, lddx );
}
magma_free( dworkd );
magma_free_cpu( hworkd );
return *info;
}
extern "C" magma_int_t
magma_zqr(
magma_int_t m, magma_int_t n,
magma_z_matrix A,
magma_int_t lda,
magma_z_matrix *Q,
magma_z_matrix *R,
magma_queue_t queue )
{
magma_int_t info = 0;
// local constants
const magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
// local variables
magma_int_t inc = 1;
magma_int_t k = min(m,n);
magma_int_t ldt;
magma_int_t nb;
magmaDoubleComplex *tau = NULL;
magmaDoubleComplex *dT = NULL;
magmaDoubleComplex *dA = NULL;
magma_z_matrix dR1 = {Magma_CSR};
// allocate CPU resources
CHECK( magma_zmalloc_pinned( &tau, k ) );
// query number of blocks required for QR factorization
nb = magma_get_zgeqrf_nb( m, n );
ldt = (2 * k + magma_roundup(n, 32)) * nb;
CHECK( magma_zmalloc( &dT, ldt ) );
// get copy of matrix array
if ( A.memory_location == Magma_DEV ) {
dA = A.dval;
} else {
CHECK( magma_zmalloc( &dA, lda * n ) );
magma_zsetvector( lda * n, A.val, inc, dA, inc, queue );
}
// QR factorization
magma_zgeqrf_gpu( m, n, dA, lda, tau, dT, &info );
// construct R matrix
if ( R != NULL ) {
if ( A.memory_location == Magma_DEV ) {
CHECK( magma_zvinit( R, Magma_DEV, lda, n, c_zero, queue ) );
magmablas_zlacpy( MagmaUpper, k, n, dA, lda, R->dval, lda, queue );
} else {
CHECK( magma_zvinit( &dR1, Magma_DEV, lda, n, c_zero, queue ) );
magmablas_zlacpy( MagmaUpper, k, n, dA, lda, dR1.dval, lda, queue );
CHECK( magma_zvinit( R, Magma_CPU, lda, n, c_zero, queue ) );
magma_zgetvector( lda * n, dR1.dval, inc, R->val, inc, queue );
}
}
// construct Q matrix
if ( Q != NULL ) {
magma_zungqr_gpu( m, n, k, dA, lda, tau, dT, nb, &info );
if ( A.memory_location == Magma_DEV ) {
CHECK( magma_zvinit( Q, Magma_DEV, lda, n, c_zero, queue ) );
magma_zcopyvector( lda * n, dA, inc, Q->dval, inc, queue );
} else {
CHECK( magma_zvinit( Q, Magma_CPU, lda, n, c_zero, queue ) );
magma_zgetvector( lda * n, dA, inc, Q->val, inc, queue );
}
}
cleanup:
if( info != 0 ){
magma_zmfree( Q, queue );
magma_zmfree( R, queue );
magma_zmfree( &dR1, queue );
}
// free resources
magma_free_pinned( tau );
magma_free( dT );
if ( A.memory_location == Magma_CPU ) {
magma_free( dA );
}
return info;
}
extern "C" magma_int_t
magma_zlobpcg(
magma_z_matrix A,
magma_z_solver_par *solver_par,
magma_z_preconditioner *precond_par,
magma_queue_t queue )
{
magma_int_t info = 0;
#define residualNorms(i,iter) ( residualNorms + (i) + (iter)*n )
#define 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_z_bspmv_tuned(m, n, alpha, A, X, beta, AX, queue) { \
magma_z_matrix x={Magma_CSR}, ax={Magma_CSR}; \
x.memory_location = Magma_DEV; x.num_rows = m; x.num_cols = n; x.major = MagmaColMajor; x.nnz = m*n; x.dval = X; x.storage_type = Magma_DENSE; \
ax.memory_location= Magma_DEV; ax.num_rows = m; ax.num_cols = n; ax.major = MagmaColMajor; ax.nnz = m*n; ax.dval = AX; ax.storage_type = Magma_DENSE; \
CHECK( magma_z_spmv(alpha, A, x, beta, ax, queue )); \
}
//**************************************************************
// 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);
magmaDoubleComplex *blockX = solver_par->eigenvectors;
double *evalues = solver_par->eigenvalues;
solver_par->numiter = 0;
solver_par->spmv_count = 0;
magmaDoubleComplex *dwork=NULL, *hwork=NULL;
magmaDoubleComplex *blockP=NULL, *blockAP=NULL, *blockR=NULL, *blockAR=NULL, *blockAX=NULL, *blockW=NULL;
magmaDoubleComplex *gramA=NULL, *gramB=NULL, *gramM=NULL;
magmaDoubleComplex *gevectors=NULL, *h_gramB=NULL;
dwork = NULL;
hwork = NULL;
blockP = NULL;
blockR = NULL;
blockAP = NULL;
blockAR = NULL;
blockAX = NULL;
blockW = NULL;
gramA = NULL;
gramB = NULL;
gramM = NULL;
gevectors = NULL;
h_gramB = NULL;
magmaDoubleComplex *pointer, *origX = blockX;
double *eval_gpu=NULL;
magma_int_t iterationNumber, cBlockSize, restart = 1, iter;
//Chronometry
real_Double_t tempo1, tempo2, tempop1, tempop2;
magma_int_t lwork = max( 2*n+n*magma_get_dsytrd_nb(n),
1 + 6*3*n + 2* 3*n* 3*n);
magma_int_t *iwork={0}, liwork = 15*n+9;
magma_int_t gramDim, ldgram = 3*n, ikind = 3;
magmaDoubleComplex *hW={0};
// === Set solver parameters ===
double residualTolerance = solver_par->rtol;
magma_int_t maxIterations = solver_par->maxiter;
double tmp;
double r0=0; // set in 1st iteration
// === Set some constants & defaults ===
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
double *residualNorms={0}, *condestGhistory={0}, condestG={0};
double *gevalues={0};
magma_int_t *activeMask={0};
double *hresidualNorms={0};
#ifdef COMPLEX
double *rwork={0};
magma_int_t lrwork = 1 + 5*(3*n) + 2*(3*n)*(3*n);
CHECK( magma_dmalloc_cpu(&rwork, lrwork));
#endif
CHECK( magma_zmalloc_pinned( &hwork , lwork ));
CHECK( magma_zmalloc( &blockAX , m*n ));
CHECK( magma_zmalloc( &blockAR , m*n ));
CHECK( magma_zmalloc( &blockAP , m*n ));
CHECK( magma_zmalloc( &blockR , m*n ));
CHECK( magma_zmalloc( &blockP , m*n ));
CHECK( magma_zmalloc( &blockW , m*n ));
CHECK( magma_zmalloc( &dwork , m*n ));
CHECK( magma_dmalloc( &eval_gpu , 3*n ));
//**********************************************************+
// === Check some parameters for possible quick exit ===
solver_par->info = MAGMA_SUCCESS;
if (m < 2)
info = MAGMA_DIVERGENCE;
else if (n > m)
info = MAGMA_SLOW_CONVERGENCE;
if (solver_par->info != 0) {
magma_xerbla( __func__, -(info) );
goto cleanup;
}
solver_par->info = info; // local info variable;
// === Allocate GPU memory for the residual norms' history ===
CHECK( magma_dmalloc(&residualNorms, (maxIterations+1) * n));
CHECK( magma_malloc( (void **)&activeMask, (n+1) * sizeof(magma_int_t) ));
// === Allocate CPU work space ===
CHECK( magma_dmalloc_cpu(&condestGhistory, maxIterations+1));
CHECK( magma_dmalloc_cpu(&gevalues, 3 * n));
CHECK( magma_malloc_cpu((void **)&iwork, liwork * sizeof(magma_int_t)));
CHECK( magma_zmalloc_pinned(&hW, n*n));
CHECK( magma_zmalloc_pinned(&gevectors, 9*n*n));
CHECK( magma_zmalloc_pinned(&h_gramB , 9*n*n));
// === Allocate GPU workspace ===
CHECK( magma_zmalloc(&gramM, n * n));
CHECK( magma_zmalloc(&gramA, 9 * n * n));
CHECK( magma_zmalloc(&gramB, 9 * n * n));
// === Set activemask to one ===
for(magma_int_t k =0; k<n; k++){
iwork[k]=1;
}
magma_setmatrix(n, 1, sizeof(magma_int_t), iwork, n , activeMask, n, queue);
#if defined(PRECISION_s)
ikind = 3;
#endif
// === Make the initial vectors orthonormal ===
magma_zgegqr_gpu(ikind, m, n, blockX, m, dwork, hwork, &info );
//magma_zorthomgs( m, n, blockX, queue );
magma_z_bspmv_tuned(m, n, c_one, A, blockX, c_zero, blockAX, queue );
solver_par->spmv_count++;
// === Compute the Gram matrix = (X, AX) & its eigenstates ===
magma_zgemm( MagmaConjTrans, MagmaNoTrans, n, n, m,
c_one, blockX, m, blockAX, m, c_zero, gramM, n, queue );
magma_zheevd_gpu( MagmaVec, MagmaUpper,
n, gramM, n, evalues, hW, n, hwork, lwork,
#ifdef COMPLEX
rwork, lrwork,
#endif
iwork, liwork, &info );
// === Update X = X * evectors ===
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, blockX, m, gramM, n, c_zero, blockW, m, queue );
SWAP(blockW, blockX);
// === Update AX = AX * evectors ===
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, blockAX, m, gramM, n, c_zero, blockW, m, queue );
SWAP(blockW, blockAX);
condestGhistory[1] = 7.82;
tempo1 = magma_sync_wtime( queue );
// === Main LOBPCG loop ============================================================
for(iterationNumber = 1; iterationNumber < maxIterations; iterationNumber++)
{
// === compute the residuals (R = Ax - x evalues )
magmablas_zlacpy( MagmaFull, m, n, blockAX, m, blockR, m, queue );
/*
for(magma_int_t i=0; i<n; i++) {
magma_zaxpy( m, MAGMA_Z_MAKE(-evalues[i],0), blockX+i*m, 1, blockR+i*m, 1, queue );
}
*/
magma_dsetmatrix( 3*n, 1, evalues, 3*n, eval_gpu, 3*n, queue );
CHECK( magma_zlobpcg_res( m, n, eval_gpu, blockX, blockR, eval_gpu, queue ));
magmablas_dznrm2_cols( m, n, blockR, m, residualNorms(0, iterationNumber), queue );
// === remove the residuals corresponding to already converged evectors
CHECK( magma_zcompact(m, n, blockR, m,
residualNorms(0, iterationNumber), residualTolerance,
activeMask, &cBlockSize, queue ));
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_zlacpy( MagmaFull, m, cBlockSize, blockR, m, blockW, m, queue );
//SWAP(blockW, blockR);
// preconditioner
magma_z_matrix bWv={Magma_CSR}, bRv={Magma_CSR};
bWv.memory_location = Magma_DEV; bWv.num_rows = m; bWv.num_cols = cBlockSize; bWv.major = MagmaColMajor; bWv.nnz = m*cBlockSize; bWv.dval = blockW;
bRv.memory_location = Magma_DEV; bRv.num_rows = m; bRv.num_cols = cBlockSize; bRv.major = MagmaColMajor; bRv.nnz = m*cBlockSize; bRv.dval = blockR;
tempop1 = magma_sync_wtime( queue );
CHECK( magma_z_applyprecond_left( MagmaNoTrans, A, bRv, &bWv, precond_par, queue ));
CHECK( magma_z_applyprecond_right( MagmaNoTrans, A, bWv, &bRv, precond_par, queue ));
tempop2 = magma_sync_wtime( queue );
precond_par->runtime += tempop2-tempop1;
// === make the preconditioned residuals orthogonal to X
if( precond_par->solver != Magma_NONE){
magma_zgemm( MagmaConjTrans, MagmaNoTrans, n, cBlockSize, m,
c_one, blockX, m, blockR, m, c_zero, gramB(0,0), ldgram, queue );
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, cBlockSize, n,
c_neg_one, blockX, m, gramB(0,0), ldgram, c_one, blockR, m, queue );
}
// === make the active preconditioned residuals orthonormal
magma_zgegqr_gpu(ikind, m, cBlockSize, blockR, m, dwork, hwork, &info );
#if defined(PRECISION_s)
// re-orthogonalization
SWAP(blockX, dwork);
magma_zgegqr_gpu(ikind, m, cBlockSize, blockR, m, dwork, hwork, &info );
#endif
//magma_zorthomgs( m, cBlockSize, blockR, queue );
// === compute AR
magma_z_bspmv_tuned(m, cBlockSize, c_one, A, blockR, c_zero, blockAR, queue );
solver_par->spmv_count++;
if (!restart) {
// === compact P & AP as well
CHECK( magma_zcompactActive(m, n, blockP, m, activeMask, queue ));
CHECK( magma_zcompactActive(m, n, blockAP, m, activeMask, queue ));
/*
// === make P orthogonal to X ?
magma_zgemm( MagmaConjTrans, MagmaNoTrans, n, cBlockSize, m,
c_one, blockX, m, blockP, m, c_zero, gramB(0,0), ldgram, queue );
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, cBlockSize, n,
c_neg_one, blockX, m, gramB(0,0), ldgram, c_one, blockP, m, queue );
// === make P orthogonal to R ?
magma_zgemm( MagmaConjTrans, MagmaNoTrans, cBlockSize, cBlockSize, m,
c_one, blockR, m, blockP, m, c_zero, gramB(0,0), ldgram, queue );
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, cBlockSize, cBlockSize,
c_neg_one, blockR, m, gramB(0,0), ldgram, c_one, blockP, m, queue );
*/
// === Make P orthonormal & properly change AP (without multiplication by A)
magma_zgegqr_gpu(ikind, m, cBlockSize, blockP, m, dwork, hwork, &info );
#if defined(PRECISION_s)
// re-orthogonalization
SWAP(blockX, dwork);
magma_zgegqr_gpu(ikind, m, cBlockSize, blockP, m, dwork, hwork, &info );
#endif
//magma_zorthomgs( m, cBlockSize, blockP, queue );
//magma_z_bspmv_tuned(m, cBlockSize, c_one, A, blockP, c_zero, blockAP, queue );
magma_zsetmatrix( cBlockSize, cBlockSize, hwork, cBlockSize, dwork, cBlockSize, queue );
// replacement according to Stan
#if defined(PRECISION_s) || defined(PRECISION_d)
magmablas_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
m, cBlockSize, c_one, dwork, cBlockSize, blockAP, m, queue );
#else
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
m, cBlockSize, c_one, dwork, cBlockSize, blockAP, m, queue );
#endif
}
iter = max( 1, iterationNumber - 10 - int(log(1.*cBlockSize)) );
double condestGmean = 0.;
for(magma_int_t 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_zlaset( MagmaFull, ldgram, ldgram, c_zero, c_one, gramB, ldgram, queue ); // identity
if (!restart) {
magma_zgemm( MagmaConjTrans, MagmaNoTrans, cBlockSize, n, m,
c_one, blockP, m, blockX, m, c_zero, gramB(n+cBlockSize,0), ldgram, queue );
magma_zgemm( MagmaConjTrans, MagmaNoTrans, cBlockSize, cBlockSize, m,
c_one, blockP, m, blockR, m, c_zero, gramB(n+cBlockSize,n), ldgram, queue );
}
magma_zgemm( MagmaConjTrans, MagmaNoTrans, cBlockSize, n, m,
c_one, blockR, m, blockX, m, c_zero, gramB(n,0), ldgram, queue );
// === get GramB from the GPU to the CPU and compute its eigenvalues only
magma_zgetmatrix( gramDim, gramDim, gramB, ldgram, h_gramB, ldgram, queue );
lapackf77_zheev("N", "L", &gramDim, h_gramB, &ldgram, gevalues,
hwork, &lwork,
#ifdef COMPLEX
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_zlaset( MagmaFull, ldgram, ldgram, c_zero, c_one, gramA, ldgram, queue ); // identity
magma_zgemm( MagmaConjTrans, MagmaNoTrans, cBlockSize, n, m,
c_one, blockR, m, blockAX, m, c_zero, gramA(n,0), ldgram, queue );
magma_zgemm( MagmaConjTrans, MagmaNoTrans, cBlockSize, cBlockSize, m,
c_one, blockR, m, blockAR, m, c_zero, gramA(n,n), ldgram, queue );
if (!restart) {
magma_zgemm( MagmaConjTrans, MagmaNoTrans, cBlockSize, n, m,
c_one, blockP, m, blockAX, m, c_zero,
gramA(n+cBlockSize,0), ldgram, queue );
magma_zgemm( MagmaConjTrans, MagmaNoTrans, cBlockSize, cBlockSize, m,
c_one, blockP, m, blockAR, m, c_zero,
gramA(n+cBlockSize,n), ldgram, queue );
magma_zgemm( MagmaConjTrans, MagmaNoTrans, cBlockSize, cBlockSize, m,
c_one, blockP, m, blockAP, m, c_zero,
gramA(n+cBlockSize,n+cBlockSize), ldgram, queue );
}
/*
// === Compute X' AX or just use the eigenvalues below ?
magma_zgemm( MagmaConjTrans, MagmaNoTrans, n, n, m,
c_one, blockX, m, blockAX, m, c_zero,
gramA(0,0), ldgram, queue );
*/
if (restart==0) {
magma_zgetmatrix( gramDim, gramDim, gramA, ldgram, gevectors, ldgram, queue );
}
else {
gramDim = n+cBlockSize;
magma_zgetmatrix( gramDim, gramDim, gramA, ldgram, gevectors, ldgram, queue );
}
for(magma_int_t k=0; k<n; k++)
*gevectors(k,k) = MAGMA_Z_MAKE(evalues[k], 0);
// === the previous eigensolver destroyed what is in h_gramB => must copy it again
magma_zgetmatrix( gramDim, gramDim, gramB, ldgram, h_gramB, ldgram, queue );
magma_int_t itype = 1;
lapackf77_zhegvd(&itype, "V", "L", &gramDim,
gevectors, &ldgram, h_gramB, &ldgram,
gevalues, hwork, &lwork,
#ifdef COMPLEX
rwork, &lrwork,
#endif
iwork, &liwork, &info);
for(magma_int_t 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_zsetmatrix( gramDim, gramDim, gevectors, ldgram, gramA, ldgram, queue );
if (restart == 0) {
// === contribution from P to the new X (in new search direction P)
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, n, cBlockSize,
c_one, blockP, m, gramA(n+cBlockSize,0), ldgram, c_zero, dwork, m, queue );
SWAP(dwork, blockP);
// === contribution from R to the new X (in new search direction P)
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, n, cBlockSize,
c_one, blockR, m, gramA(n,0), ldgram, c_one, blockP, m, queue );
// === corresponding contribution from AP to the new AX (in AP)
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, n, cBlockSize,
c_one, blockAP, m, gramA(n+cBlockSize,0), ldgram, c_zero, dwork, m, queue );
SWAP(dwork, blockAP);
// === corresponding contribution from AR to the new AX (in AP)
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, n, cBlockSize,
c_one, blockAR, m, gramA(n,0), ldgram, c_one, blockAP, m, queue );
}
else {
// === contribution from R (only) to the new X
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, n, cBlockSize,
c_one, blockR, m, gramA(n,0), ldgram, c_zero, blockP, m, queue );
// === corresponding contribution from AR (only) to the new AX
magma_zgemm( MagmaNoTrans, MagmaNoTrans,m, n, cBlockSize,
c_one, blockAR, m, gramA(n,0), ldgram, c_zero, blockAP, m, queue );
}
// === contribution from old X to the new X + the new search direction P
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, blockX, m, gramA, ldgram, c_zero, dwork, m, queue );
SWAP(dwork, blockX);
//magma_zaxpy( m*n, c_one, blockP, 1, blockX, 1, queue );
CHECK( magma_zlobpcg_maxpy( m, n, blockP, blockX, queue ));
// === corresponding contribution from old AX to new AX + AP
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, blockAX, m, gramA, ldgram, c_zero, dwork, m, queue );
SWAP(dwork, blockAX);
//magma_zaxpy( m*n, c_one, blockAP, 1, blockAX, 1, queue );
CHECK( magma_zlobpcg_maxpy( m, n, blockAP, blockAX, queue ));
condestGhistory[iterationNumber+1]=condestG;
magma_dgetmatrix( 1, 1, residualNorms(0, iterationNumber), 1, &tmp, 1, queue );
if ( iterationNumber == 1 ) {
solver_par->init_res = tmp;
r0 = tmp * solver_par->rtol;
if ( r0 < ATOLERANCE )
r0 = ATOLERANCE;
}
solver_par->final_res = tmp;
if ( tmp < r0 ) {
break;
}
if (cBlockSize == 0) {
break;
}
if ( solver_par->verbose!=0 ) {
if ( iterationNumber%solver_par->verbose == 0 ) {
// double res;
// magma_zgetmatrix( 1, 1,
// (magmaDoubleComplex*)residualNorms(0, iterationNumber), 1,
// (magmaDoubleComplex*)&res, 1, queue );
//
// 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
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t) tempo2-tempo1;
solver_par->numiter = iterationNumber;
if ( solver_par->numiter < solver_par->maxiter) {
info = MAGMA_SUCCESS;
} else if ( solver_par->init_res > solver_par->final_res )
info = MAGMA_SLOW_CONVERGENCE;
else
info = MAGMA_DIVERGENCE;
// =============================================================================
// === postprocessing;
// =============================================================================
// === compute the real AX and corresponding eigenvalues
magma_z_bspmv_tuned(m, n, c_one, A, blockX, c_zero, blockAX, queue );
magma_zgemm( MagmaConjTrans, MagmaNoTrans, n, n, m,
c_one, blockX, m, blockAX, m, c_zero, gramM, n, queue );
magma_zheevd_gpu( MagmaVec, MagmaUpper,
n, gramM, n, gevalues, dwork, n, hwork, lwork,
#ifdef COMPLEX
rwork, lrwork,
#endif
iwork, liwork, &info );
for(magma_int_t k =0; k<n; k++)
evalues[k] = gevalues[k];
// === update X = X * evectors
SWAP(blockX, dwork);
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, dwork, m, gramM, n, c_zero, blockX, m, queue );
// === update AX = AX * evectors to compute the final residual
SWAP(blockAX, dwork);
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, dwork, m, gramM, n, c_zero, blockAX, m, queue );
// === compute R = AX - evalues X
magmablas_zlacpy( MagmaFull, m, n, blockAX, m, blockR, m, queue );
for(magma_int_t i=0; i<n; i++)
magma_zaxpy( m, MAGMA_Z_MAKE(-evalues[i], 0), blockX+i*m, 1, blockR+i*m, 1, queue );
// === residualNorms[iterationNumber] = || R ||
magmablas_dznrm2_cols( m, n, blockR, m, residualNorms(0, iterationNumber), queue );
// === restore blockX if needed
if (blockX != origX)
magmablas_zlacpy( MagmaFull, m, n, blockX, m, origX, m, queue );
printf("Eigenvalues:\n");
for(magma_int_t 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");
//=== Prmagma_int_t residual history in a file for plotting ====
CHECK( magma_dmalloc_cpu(&hresidualNorms, (iterationNumber+1) * n));
magma_dgetmatrix( n, iterationNumber,
residualNorms, n,
hresidualNorms, n, queue );
solver_par->iter_res = *hresidualNorms(0, iterationNumber-1);
printf("Residuals are stored in file residualNorms\n");
printf("Plot the residuals using: myplot \n");
FILE *residuals_file;
residuals_file = fopen("residualNorms", "w");
for(magma_int_t i =1; i<iterationNumber; i++) {
for(magma_int_t j = 0; j<n; j++)
fprintf(residuals_file, "%f ", *hresidualNorms(j,i));
fprintf(residuals_file, "\n");
}
fclose(residuals_file);
cleanup:
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 );
if (blockX != (solver_par->eigenvectors))
magma_free( blockX );
if (blockAX != (solver_par->eigenvectors))
magma_free( blockAX );
if (blockAR != (solver_par->eigenvectors))
magma_free( blockAR );
if (blockAP != (solver_par->eigenvectors))
magma_free( blockAP );
if (blockR != (solver_par->eigenvectors))
magma_free( blockR );
if (blockP != (solver_par->eigenvectors))
magma_free( blockP );
if (blockW != (solver_par->eigenvectors))
magma_free( blockW );
if (dwork != (solver_par->eigenvectors))
magma_free( dwork );
magma_free( eval_gpu );
magma_free_pinned( hwork );
#ifdef COMPLEX
magma_free_cpu( rwork );
rwork = NULL;
#endif
return info;
}
extern "C" magma_int_t
magma_zlahr2_m(
magma_int_t n, magma_int_t k, magma_int_t nb,
magmaDoubleComplex *A, magma_int_t lda,
magmaDoubleComplex *tau,
magmaDoubleComplex *T, magma_int_t ldt,
magmaDoubleComplex *Y, magma_int_t ldy,
struct zgehrd_data* data )
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
ZLAHR2 reduces the first NB columns of a complex 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 ZGEHRD.
Arguments
=========
N (input) INTEGER
The order of the matrix A.
K (input) INTEGER
The offset for the reduction. Elements below the k-th
subdiagonal in the first NB columns are reduced to zero.
K < N.
NB (input) INTEGER
The number of columns to be reduced.
A (input/output) COMPLEX_16 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.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
TAU (output) COMPLEX_16 array, dimension (NB)
The scalar factors of the elementary reflectors. See Further
Details.
T (output) COMPLEX_16 array, dimension (LDT,NB)
The upper triangular matrix T.
LDT (input) INTEGER
The leading dimension of the array T. LDT >= NB.
Y (output) COMPLEX_16 array, dimension (LDY,NB)
The n-by-nb matrix Y.
LDY (input) INTEGER
The leading dimension of the array Y. LDY >= N.
dA (input/output) COMPLEX_16 array on the GPU, dimension (LDA,N-K+1)
On entry, the n-by-(n-k+1) general matrix A.
On exit, the elements in rows K:N of the first NB columns are
overwritten with the matrix Y.
DV (output) COMPLEX_16 array on the GPU, dimension (N, NB)
On exit this contains the Householder vectors of the transformation.
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 complex scalar, and v is a complex 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:
( 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 )
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.
===================================================================== */
#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)
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex 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;
magmaDoubleComplex scale;
magma_int_t i;
magmaDoubleComplex ei = MAGMA_Z_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_zlaset( MagmaUpperLower, nb, nb, dV(d,k,0), ldv );
}
// set all Y=0
lapackf77_zlaset( "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_Z_NEGATE( tau[i-1] );
blasf77_zaxpy( &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_zcopy( &i,
A(k+1,i), &ione,
T(0,nb-1), &ione );
// w := V1' * b1 = VA(k+1:k+i, 0:i-1)' * w
blasf77_ztrmv( "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_zgemv( "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_ztrmv( "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_zgemv( "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_ztrmv( "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_zaxpy( &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_zlarfg( &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_zsetvector_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_zlacpy( 'F', 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_zgemv( 'N', 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_zgetvector_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_Z_NEGATE( tau[i] );
blasf77_zgemv( "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_ztrmv( "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 complex, conjugate row of V, and undo afterwards
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv( &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_zgemv( "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_zlacgv( &i1, A(k+i1,0), &lda );
#endif
// A(k:n, i+1) -= Y(k:n, 0:i) * w
blasf77_zgemv( "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_zaxpy( &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_zgemm( 'N', 'N', 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_zgetmatrix_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_zaxpy( &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_zsetmatrix_async( n, nb, Y, ldy, dY(d, 0, 0), ldda, data->streams[d] );
magma_zsetmatrix_async( nb, nb, T, nb, dTi(d), nb, data->streams[d] );
}
return 0;
} // magma_zlahr2
extern "C" magma_int_t
magma_zcgeqrsv_gpu(magma_int_t m, magma_int_t n, magma_int_t nrhs,
magmaDoubleComplex *dA, magma_int_t ldda,
magmaDoubleComplex *dB, magma_int_t lddb,
magmaDoubleComplex *dX, magma_int_t lddx,
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
=======
ZCGEQRSV 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.
ZCGEQRSV first attempts to factorize the matrix in complex SINGLE PRECISION
and use this factorization within an iterative refinement procedure
to produce a solution with complex DOUBLE PRECISION norm-wise backward error
quality (see below). If the approach fails the method switches to a
complex DOUBLE PRECISION factorization and solve.
The iterative refinement is not going to be a winning strategy if
the ratio complex SINGLE PRECISION performance over complex 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
=========
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. M >= 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) COMPLEX_16 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.
LDDA (input) INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
dB (input or input/output) COMPLEX_16 array on the GPU, dimension (LDDB,NRHS)
The M-by-NRHS right hand side matrix B.
May be overwritten (e.g., if refinement fails).
LDDB (input) INTEGER
The leading dimension of the array dB. LDDB >= max(1,M).
dX (output) COMPLEX_16 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).
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 SGEQRF
-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
===================================================================== */
#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)
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magma_int_t ione = 1;
magmaDoubleComplex *dworkd, *hworkd;
magmaFloatComplex *dworks, *hworks;
magmaDoubleComplex *dR, *tau, *dT;
magmaFloatComplex *dSA, *dSX, *dST, *stau;
magmaDoubleComplex 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_cgeqrf_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_cmalloc( &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_zmalloc( &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_cmalloc_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_zlange('I', m, n, dA, ldda, (double*)dworkd );
cte = Anrm * eps * pow((double)n, 0.5) * BWDMAX;
/*
* Convert to single precision
*/
magmablas_zlag2c( m, nrhs, dB, lddb, dSX, lddsx, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
magmablas_zlag2c( m, n, dA, ldda, dSA, lddsa, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
// factor dSA in single precision
magma_cgeqrf_gpu( m, n, dSA, lddsa, stau, dST, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// solve dSA*dSX = dB in single precision
magma_cgeqrs_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_clag2z( n, nrhs, dSX, lddsx, dX, lddx, info );
magmablas_zlacpy( MagmaUpperLower, m, nrhs, dB, lddb, dR, lddr );
if ( nrhs == 1 ) {
magma_zgemv( MagmaNoTrans, m, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, nrhs, n,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
// TODO: use MAGMA_Z_ABS( dX(i,j) ) instead of zlange?
for( j=0; j < nrhs; j++ ) {
i = magma_izamax( n, dX(0,j), 1) - 1;
magma_zgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_zlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_izamax ( m, dR(0,j), 1 ) - 1;
magma_zgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_zlange( "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_zlag2c( m, nrhs, dR, lddr, dSX, lddsx, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
// solve dSA*dSX = R in single precision
magma_cgeqrs_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_zcaxpycp( n, dSX(0,j), dX(0,j), dB(0,j), dR(0,j) );
}
// dR = dB (whole m rows)
magmablas_zlacpy( MagmaUpperLower, m, nrhs, dB, lddb, dR, lddr );
// residual dR = dB - dA*dX in double precision
if ( nrhs == 1 ) {
magma_zgemv( MagmaNoTrans, m, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_zgemm( 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_izamax( n, dX(0,j), 1) - 1;
magma_zgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_zlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_izamax ( m, dR(0,j), 1 ) - 1;
magma_zgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_zlange( "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 zgeqrf */
nb = magma_get_zgeqrf_nb( m );
size = (2*min(m, n) + (n+31)/32*32 )*nb;
if ( size > ldworkd ) {
magma_free( dworkd );
if (MAGMA_SUCCESS != magma_zmalloc( &dworkd, size )) {
fprintf(stderr, "Allocation of dworkd2 failed\n");
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
}
dT = dworkd;
/* hworkd(dtau + workspace for zgeqrs) = min(m,n) + lhwork */
size = lhwork + minmn;
magma_zmalloc_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_zgeqrf_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_zgeqrs_gpu( m, n, nrhs, dA, ldda, tau, dT, dB, lddb, hworkd, lhwork, info );
magmablas_zlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dX, lddx );
}
magma_free( dworkd );
magma_free_cpu( hworkd );
return *info;
}
void magmablas_zhemm_mgpu_spec(
char side, char uplo, magma_int_t m, magma_int_t n,
magmaDoubleComplex alpha,
magmaDoubleComplex *dA[], magma_int_t ldda, magma_int_t offset,
magmaDoubleComplex *dB[], magma_int_t lddb,
magmaDoubleComplex beta,
magmaDoubleComplex *dC[], magma_int_t lddc,
magmaDoubleComplex *dwork[], magma_int_t dworksiz,
magmaDoubleComplex *C, magma_int_t ldc,
magmaDoubleComplex *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 );
magmaDoubleComplex *dwork1[MagmaMaxGPUs];
magmaDoubleComplex *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_zgemm( MagmaConjTrans, 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_zgemm( MagmaConjTrans, 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 complex and they should
// send it out to their actives slaves.
magma_setdevice( dev );
//==============================================
// sending to the master of the active complex
//==============================================
//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){ //complex 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(magmaDoubleComplex),
&dwork[dev][maxgsize*dev], mycolsize*sizeof(magmaDoubleComplex),
mycolsize*sizeof(magmaDoubleComplex), n,
cudaMemcpyDeviceToDevice, streams[dev][gmaster]);
magma_event_record(redevents[dev][gmaster*ngpu+dev], streams[dev][gmaster]);
}
}
}
//==============================================
//
//==============================================
// sending to the active GPUs of my complex
//==============================================
//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(magmaDoubleComplex),
&dwork[dev][maxgsize*dev], mycolsize*sizeof(magmaDoubleComplex),
mycolsize*sizeof(magmaDoubleComplex), 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 complex 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(magmaDoubleComplex),
&dwork[masterdev][maxgsize*gdev], gcolsize*sizeof(magmaDoubleComplex),
gcolsize*sizeof(magmaDoubleComplex), 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 zlacpy\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_zlacpy( '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 zlacpy\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_zlacpy( '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 complex 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
-------
ZGERFS improve the computed solution to a system of linear
equations.
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]
dA COMPLEX_16 array on the GPU, dimension (ldda,N)
the N-by-N coefficient matrix A.
@param[in]
ldda INTEGER
The leading dimension of the array dA. ldda >= max(1,N).
@param[in]
dB COMPLEX_16 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[in, out]
dX COMPLEX_16 array on the GPU, dimension (lddx,NRHS)
On entry, the solution matrix X, as computed by
ZGETRS_NOPIV. On exit, the improved solution matrix X.
@param[in]
lddx INTEGER
The leading dimension of the array dX. lddx >= max(1,N).
@param
dworkd (workspace) COMPLEX_16 array on the GPU, dimension (N*NRHS)
This array is used to hold the residual vectors.
@param
dAF COMPLEX*16 array on the GPU, dimension (ldda,n)
The factors L and U from the factorization A = L*U
as computed by ZGETRF_NOPIV.
@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_zgesv_driver
********************************************************************/
extern "C" magma_int_t
magma_zgerfs_nopiv_gpu(
magma_trans_t trans, magma_int_t n, magma_int_t nrhs,
magmaDoubleComplex_ptr dA, magma_int_t ldda,
magmaDoubleComplex_ptr dB, magma_int_t lddb,
magmaDoubleComplex_ptr dX, magma_int_t lddx,
magmaDoubleComplex_ptr dworkd, magmaDoubleComplex_ptr dAF,
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)
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magma_int_t ione = 1;
magmaDoubleComplex_ptr dR;
magmaDoubleComplex 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;
dR = dworkd;
eps = lapackf77_dlamch("Epsilon");
Anrm = magmablas_zlange(MagmaInfNorm, n, n, dA, ldda, (double*)dworkd );
cte = Anrm * eps * pow( (double)n, (double)0.5 ) * BWDMAX;
// residual dR = dB - dA*dX in double precision
magmablas_zlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dR, lddr );
if ( nrhs == 1 ) {
magma_zgemv( trans, n, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_zgemm( trans, MagmaNoTrans, n, nrhs, n,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
// TODO: use MAGMA_Z_ABS( dX(i,j) ) instead of zlange?
for( j=0; j < nrhs; j++ ) {
i = magma_izamax( n, dX(0,j), 1) - 1;
magma_zgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_zlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_izamax ( n, dR(0,j), 1 ) - 1;
magma_zgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_zlange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
// printf("Rnrm : %e, Xnrm*cte : %e\n", Rnrm, Xnrm*cte);
if ( Rnrm > Xnrm*cte ) {
goto REFINEMENT;
}
}
*iter = 0;
return *info;
REFINEMENT:
for( iiter=1; iiter < ITERMAX; ) {
*info = 0;
// solve dAF*dX = dR
// it's okay that dR is used for both dB input and dX output.
magma_zgetrs_nopiv_gpu( trans, n, nrhs, dAF, lddsa, dR, lddr, 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_zaxpycp2( n, dR(0,j), dX(0,j), dB(0,j) );
}
// residual dR = dB - dA*dX in double precision
if ( nrhs == 1 ) {
magma_zgemv( trans, n, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_zgemm( 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_izamax( n, dX(0,j), 1) - 1;
magma_zgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_zlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_izamax ( n, dR(0,j), 1 ) - 1;
magma_zgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_zlange( "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. */
*iter = -ITERMAX - 1;
FALLBACK:
/* Iterative refinement failed to converge to a
* satisfactory solution. */
return *info;
}
/**
Purpose
-------
ZGETRI computes the inverse of a matrix using the LU factorization
computed by ZGETRF. 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 ZGESV,
or ZGETRF and ZGETRS, 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 COMPLEX_16 array on the GPU, dimension (LDDA,N)
On entry, the factors L and U from the factorization
A = P*L*U as computed by ZGETRF_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 ZGETRF; for 1 <= i <= N, row i of the
matrix was interchanged with row IPIV(i).
@param[out]
dwork (workspace) COMPLEX_16 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_zgetri_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_zgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_zgetri_gpu( magma_int_t n, magmaDoubleComplex *dA, magma_int_t ldda,
magma_int_t *ipiv, magmaDoubleComplex *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 */
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex *dL = dwork;
magma_int_t lddl = n;
magma_int_t nb = magma_get_zgetri_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_ztrtri_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_zlacpy( MagmaFull, n-j, jb,
dA(j,j), ldda,
dL(j,0), lddl );
magmablas_zlaset( 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_zgemm( 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_ztrsm( 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_zswap( n, dA(0,j), 1, dA(0,jp), 1 );
}
}
return *info;
}
extern "C" magma_int_t
magma_zcgeqrsv_gpu(magma_int_t M, magma_int_t N, magma_int_t NRHS,
cuDoubleComplex *dA, magma_int_t ldda,
cuDoubleComplex *dB, magma_int_t lddb,
cuDoubleComplex *dX, magma_int_t lddx,
magma_int_t *iter, magma_int_t *info)
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
Purpose
=======
ZCGEQRSV 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.
ZCGEQRSV first attempts to factorize the matrix in SINGLE PRECISION
and use this factorization within an iterative refinement procedure
to produce a solution with DOUBLE PRECISION norm-wise backward error
quality (see below). If the approach fails the method switches to a
DOUBLE PRECISION factorization and solve.
The iterative refinement is not going to be a winning strategy if
the ratio SINGLE PRECISION performance over 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
=========
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. M >= N >= 0.
NRHS (input) INTEGER
The number of right hand sides, i.e., the number of columns
of the matrix B. NRHS >= 0.
A (input or input/output) DOUBLE PRECISION array, 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 A contains the QR factorization of A as returned by
function DGEQRF_GPU.
ldda (input) INTEGER
The leading dimension of the array A. ldda >= max(1,M).
B (input) DOUBLE PRECISION array, dimension (LDB,NRHS)
The M-by-NRHS right hand side matrix B.
LDB (input) INTEGER
The leading dimension of the array B. LDB >= max(1,M).
X (output) DOUBLE PRECISION array, dimension (LDX,NRHS)
If info = 0, the N-by-NRHS solution matrix X.
LDX (input) INTEGER
The leading dimension of the array X. LDX >= max(1,N).
WORK (workspace) DOUBLE PRECISION array, dimension (N*NRHS)
This array is used to hold the residual vectors.
SWORK (workspace) REAL array, dimension (M*(N+NRHS))
This array is used to store the 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
iterations
> 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
TAU (output) REAL array, dimension (N)
On exit, TAU(i) contains the scalar factor of the elementary
reflector H(i), as returned by magma_cgeqrf_gpu.
LWORK (input) INTEGER
The dimension of the array H_WORK. LWORK >= (M+N+NB)*NB,
where NB can be obtained through magma_get_sgeqrf_nb(M).
H_WORK (workspace/output) REAL array, dimension (MAX(1,LWORK))
Higher performance is achieved if H_WORK is in pinned memory, e.g.
allocated using magma_malloc_pinned.
D_WORK (workspace/output) REAL array on the GPU, dimension 2*N*NB,
where NB can be obtained through magma_get_sgeqrf_nb(M).
It starts with NB*NB blocks that store the triangular T
matrices, followed by the NB*NB blocks of the diagonal
inverses for the R matrix.
TAU_D (output) DOUBLE REAL array, dimension (N)
On exit, if the matrix had to be factored in double precision,
TAU(i) contains the scalar factor of the elementary
reflector H(i), as returned by magma_zgeqrf_gpu.
LWORK_D (input) INTEGER
The dimension of the array H_WORK_D. LWORK_D >= (M+N+NB)*NB,
where NB can be obtained through magma_get_dgeqrf_nb(M).
H_WORK_D (workspace/output) DOUBLE REAL array, dimension (MAX(1,LWORK_D))
This memory is unattached if the iterative refinement worked,
otherwise it is used as workspace to factor the matrix in
double precision. Higher performance is achieved if H_WORK_D is
in pinned memory, e.g. allocated using magma_malloc_pinned.
D_WORK_D (workspace/output) DOUBLE REAL array on the GPU, dimension 2*N*NB,
where NB can be obtained through magma_get_dgeqrf_nb(M).
This memory is unattached if the iterative refinement worked,
otherwise it is used as workspace to factor the matrix in
double precision. It starts with NB*NB blocks that store the
triangular T matrices, followed by the NB*NB blocks of the
diagonal inverses for the R matrix.
===================================================================== */
cuDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
cuDoubleComplex c_one = MAGMA_Z_ONE;
magma_int_t ione = 1;
cuDoubleComplex *dworkd, *hworkd;
cuFloatComplex *dworks, *hworks;
cuDoubleComplex *dR, *tau, *dT;
cuFloatComplex *dSA, *dSX, *dST, *stau;
cuDoubleComplex Xnrmv, Rnrmv;
double Anrm, Xnrm, Rnrm, cte, eps;
magma_int_t i, j, iiter, nb, lhwork, minmn, size;
/*
Check The Parameters.
*/
*iter = 0 ;
*info = 0 ;
if ( N < 0 )
*info = -1;
else if(NRHS<0)
*info = -3;
else if( ldda < max(1,N))
*info = -5;
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;
nb = magma_get_cgeqrf_nb(M);
minmn= min(M, N);
/*
* Allocate temporary buffers
*/
/* dworks(dSA + dSX + dST) */
size = ldda*N + N*NRHS + ( 2*minmn + ((N+31)/32)*32 )*nb;
if (MAGMA_SUCCESS != magma_cmalloc( &dworks, size )) {
fprintf(stderr, "Allocation of dworks failed (%d)\n", (int) size);
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dSA = dworks;
dSX = dSA + ldda*N;
dST = dSX + N*NRHS;
/* dworkd(dR) = N*NRHS */
size = N*NRHS;
if (MAGMA_SUCCESS != magma_zmalloc( &dworkd, size )) {
magma_free( dworks );
fprintf(stderr, "Allocation of dworkd failed\n");
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dR = dworkd;
/* hworks(stau + workspace for cgeqrs) = min(M,N) + lhworks */
lhwork = nb*max((M-N+nb+2*(NRHS)), 1);
lhwork = max(lhwork, N*nb); /* We hope that magma nb is bigger than lapack nb to have enough memory in workspace */
size = minmn + lhwork;
magma_cmalloc_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_zlange('I', M, N, dA, ldda, (double*)dworkd );
cte = Anrm * eps * pow((double)N, 0.5) * BWDMAX ;
/*
* Convert to single precision
*/
magmablas_zlag2c(N, NRHS, dB, lddb, dSX, N, info );
if( *info != 0 ) {
*iter = -2; goto L40;
}
magmablas_zlag2c(N, N, dA, ldda, dSA, ldda, info );
if(*info !=0){
*iter = -2; goto L40;
}
// In an ideal version these variables should come from user.
magma_cgeqrf_gpu(M, N, dSA, ldda, stau, dST, info);
if( *info != 0 ) {
*iter = -3; goto L40;
}
magma_cgeqrs_gpu(M, N, NRHS, dSA, ldda, stau, dST, dSX, N, hworks, lhwork, info);
// dX = dSX
magmablas_clag2z(N, NRHS, dSX, N, dX, lddx, info);
// dR = dB
magmablas_zlacpy(MagmaUpperLower, N, NRHS, dB, lddb, dR, N);
// dR = dB - dA * dX
if( NRHS == 1 )
magma_zgemv( MagmaNoTrans, N, N,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1);
else
magma_zgemm( MagmaNoTrans, MagmaNoTrans, N, NRHS, N,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, N );
for(i=0; i<NRHS; i++){
j = magma_izamax( N, dX+i*N, 1);
magma_zgetmatrix( 1, 1, dX+i*N+j-1, 1, &Xnrmv, 1 );
Xnrm = lapackf77_zlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
j = magma_izamax ( N, dR+i*N, 1 );
magma_zgetmatrix( 1, 1, dR+i*N+j-1, 1, &Rnrmv, 1 );
Rnrm = lapackf77_zlange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
if( Rnrm > Xnrm *cte ) goto L10;
}
*iter = 0;
/* Free workspaces */
magma_free( dworks );
magma_free( dworkd );
magma_free_cpu( hworks );
return *info;
L10:
for(iiter=1; iiter<ITERMAX; ) {
*info = 0 ;
/* Convert R from double precision to single precision
and store the result in SX.
Solve the system SA*SX = SR.
-- These two Tasks are merged here. */
// make SWORK = WORK ... residuals...
magmablas_zlag2c( N, NRHS, dR, N, dSX, N, info );
magma_cgeqrs_gpu( M, N, NRHS, dSA, ldda, stau, dST, dSX, N, hworks, lhwork, info);
if( *info != 0 ){
*iter = -3; goto L40;
}
for(i=0; i<NRHS; i++) {
magmablas_zcaxpycp( dSX+i*N, dX+i*lddx, N, dB+i*lddb, dR+i*N );
}
/* unnecessary may be */
magmablas_zlacpy(MagmaUpperLower, N, NRHS, dB, lddb, dR, N);
if( NRHS == 1 )
magma_zgemv( MagmaNoTrans, N, N,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1);
else
magma_zgemm( MagmaNoTrans, MagmaNoTrans, N, NRHS, N,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, N);
/* Check whether the NRHS normwise backward errors satisfy the
stopping criterion. If yes, set ITER=IITER>0 and return. */
for(i=0;i<NRHS;i++)
{
j = magma_izamax( N, dX+i*N, 1);
magma_zgetmatrix( 1, 1, dX+i*N+j-1, 1, &Xnrmv, 1 );
Xnrm = lapackf77_zlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
j = magma_izamax ( N, dR+i*N, 1 );
magma_zgetmatrix( 1, 1, dR+i*N+j-1, 1, &Rnrmv, 1 );
Rnrm = lapackf77_zlange( "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 ;
L40:
magma_free( dworks );
/*
* Allocate temporary buffers
*/
/* dworkd(dT + tau) = min_mn + min_mn*nb*3 */
nb = magma_get_zgeqrf_nb(M);
size = minmn * (3 * nb + 1);
if ( size > (N*NRHS) ) {
magma_free( dworkd );
if (MAGMA_SUCCESS != magma_zmalloc( &dworkd, size )) {
fprintf(stderr, "Allocation of dworkd2 failed\n");
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
}
tau = dworkd;
dT = tau + minmn;
/* hworks(stau + workspace for cgeqrs) = min(M,N) + lhworks */
/* re-use hworks memory for hworkd if possible, else re-allocate. */
if ( (2*lhwork) <= (minmn+lhwork) ) {
hworkd = (cuDoubleComplex*) hworks;
}
else {
magma_free_cpu( hworks );
magma_zmalloc_cpu( &hworkd, lhwork );
if ( hworkd == NULL ) {
magma_free( dworkd );
fprintf(stderr, "Allocation of hworkd2 failed\n");
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
}
/* Single-precision iterative refinement failed to converge to a
satisfactory solution, so we resort to double precision. */
magma_zgeqrf_gpu(M, N, dA, ldda, tau, dT, info);
if ( *info == 0 ) {
magmablas_zlacpy(MagmaUpperLower, N, NRHS, dB, lddb, dX, lddx);
magma_zgeqrs_gpu(M, N, NRHS, dA, ldda, tau, dT, dX, lddx, hworkd, lhwork, info);
}
magma_free( dworkd );
magma_free_cpu( hworkd );
return *info;
}