// On input, A and ipiv is LU factorization of A. On output, A is overwritten.
// Requires m == n.
// Uses init_matrix() to re-generate original A as needed.
// Generates random RHS b and solves Ax=b.
// Returns residual, |Ax - b| / (n |A| |x|).
double get_residual(
magma_int_t m, magma_int_t n,
double *A, magma_int_t lda,
magma_int_t *ipiv )
{
if ( m != n ) {
printf( "\nERROR: residual check defined only for square matrices\n" );
return -1;
}
const double c_one = MAGMA_D_ONE;
const double c_neg_one = MAGMA_D_NEG_ONE;
const magma_int_t ione = 1;
// this seed should be DIFFERENT than used in init_matrix
// (else x is column of A, so residual can be exactly zero)
magma_int_t ISEED[4] = {0,0,0,2};
magma_int_t info = 0;
double *x, *b;
// initialize RHS
TESTING_MALLOC_CPU( x, double, n );
TESTING_MALLOC_CPU( b, double, n );
lapackf77_dlarnv( &ione, ISEED, &n, b );
blasf77_dcopy( &n, b, &ione, x, &ione );
// solve Ax = b
lapackf77_dgetrs( "Notrans", &n, &ione, A, &lda, ipiv, x, &n, &info );
if (info != 0)
printf("lapackf77_dgetrs returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// reset to original A
init_matrix( m, n, A, lda );
// compute r = Ax - b, saved in b
blasf77_dgemv( "Notrans", &m, &n, &c_one, A, &lda, x, &ione, &c_neg_one, b, &ione );
// compute residual |Ax - b| / (n*|A|*|x|)
double norm_x, norm_A, norm_r, work[1];
norm_A = lapackf77_dlange( "F", &m, &n, A, &lda, work );
norm_r = lapackf77_dlange( "F", &n, &ione, b, &n, work );
norm_x = lapackf77_dlange( "F", &n, &ione, x, &n, work );
//printf( "r=\n" ); magma_dprint( 1, n, b, 1 );
TESTING_FREE_CPU( x );
TESTING_FREE_CPU( b );
//printf( "r=%.2e, A=%.2e, x=%.2e, n=%d\n", norm_r, norm_A, norm_x, n );
return norm_r / (n * norm_A * norm_x);
}
extern "C" magma_int_t
magma_zheevx_gpu(char jobz, char range, char uplo, magma_int_t n,
magmaDoubleComplex *da, magma_int_t ldda, double vl, double vu,
magma_int_t il, magma_int_t iu, double abstol, magma_int_t *m,
double *w, magmaDoubleComplex *dz, magma_int_t lddz,
magmaDoubleComplex *wa, magma_int_t ldwa,
magmaDoubleComplex *wz, magma_int_t ldwz,
magmaDoubleComplex *work, magma_int_t lwork,
double *rwork, magma_int_t *iwork, magma_int_t *ifail, magma_int_t *info)
{
/* -- MAGMA (version 1.4.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
December 2013
Purpose
=======
ZHEEVX computes selected eigenvalues and, optionally, eigenvectors
of a complex Hermitian matrix A. Eigenvalues and eigenvectors can
be selected by specifying either a range of values or a range of
indices for the desired eigenvalues.
Arguments
=========
JOBZ (input) CHARACTER*1
= 'N': Compute eigenvalues only;
= 'V': Compute eigenvalues and eigenvectors.
RANGE (input) CHARACTER*1
= 'A': all eigenvalues will be found.
= 'V': all eigenvalues in the half-open interval (VL,VU]
will be found.
= 'I': the IL-th through IU-th eigenvalues will be found.
UPLO (input) CHARACTER*1
= 'U': Upper triangle of A is stored;
= 'L': Lower triangle of A is stored.
N (input) INTEGER
The order of the matrix A. N >= 0.
DA (device input/output) COMPLEX_16 array, 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. If UPLO = 'L',
the leading N-by-N lower triangular part of A contains
the lower triangular part of the matrix A.
On exit, the lower triangle (if UPLO='L') or the upper
triangle (if UPLO='U') of A, including the diagonal, is
destroyed.
LDDA (input) INTEGER
The leading dimension of the array DA. LDDA >= max(1,N).
VL (input) DOUBLE PRECISION
VU (input) DOUBLE PRECISION
If RANGE='V', the lower and upper bounds of the interval to
be searched for eigenvalues. VL < VU.
Not referenced if RANGE = 'A' or 'I'.
IL (input) INTEGER
IU (input) INTEGER
If RANGE='I', the indices (in ascending order) of the
smallest and largest eigenvalues to be returned.
1 <= IL <= IU <= N, if N > 0; IL = 1 and IU = 0 if N = 0.
Not referenced if RANGE = 'A' or 'V'.
ABSTOL (input) DOUBLE PRECISION
The absolute error tolerance for the eigenvalues.
An approximate eigenvalue is accepted as converged
when it is determined to lie in an interval [a,b]
of width less than or equal to
ABSTOL + EPS * max( |a|,|b| ) ,
where EPS is the machine precision. If ABSTOL is less than
or equal to zero, then EPS*|T| will be used in its place,
where |T| is the 1-norm of the tridiagonal matrix obtained
by reducing A to tridiagonal form.
Eigenvalues will be computed most accurately when ABSTOL is
set to twice the underflow threshold 2*DLAMCH('S'), not zero.
If this routine returns with INFO>0, indicating that some
eigenvectors did not converge, try setting ABSTOL to
2*DLAMCH('S').
See "Computing Small Singular Values of Bidiagonal Matrices
with Guaranteed High Relative Accuracy," by Demmel and
Kahan, LAPACK Working Note #3.
M (output) INTEGER
The total number of eigenvalues found. 0 <= M <= N.
If RANGE = 'A', M = N, and if RANGE = 'I', M = IU-IL+1.
W (output) DOUBLE PRECISION array, dimension (N)
On normal exit, the first M elements contain the selected
eigenvalues in ascending order.
DZ (device output) COMPLEX_16 array, dimension (LDDZ, max(1,M))
If JOBZ = 'V', then if INFO = 0, the first M columns of Z
contain the orthonormal eigenvectors of the matrix A
corresponding to the selected eigenvalues, with the i-th
column of Z holding the eigenvector associated with W(i).
If an eigenvector fails to converge, then that column of Z
contains the latest approximation to the eigenvector, and the
index of the eigenvector is returned in IFAIL.
If JOBZ = 'N', then Z is not referenced.
Note: the user must ensure that at least max(1,M) columns are
supplied in the array Z; if RANGE = 'V', the exact value of M
is not known in advance and an upper bound must be used.
********* (workspace) If FAST_HEMV is defined DZ should be (LDDZ, max(1,N)) in both cases.
LDDZ (input) INTEGER
The leading dimension of the array DZ. LDDZ >= 1, and if
JOBZ = 'V', LDDZ >= max(1,N).
WA (workspace) COMPLEX_16 array, dimension (LDWA, N)
LDWA (input) INTEGER
The leading dimension of the array WA. LDWA >= max(1,N).
WZ (workspace) COMPLEX_16 array, dimension (LDWZ, max(1,M))
LDWZ (input) INTEGER
The leading dimension of the array DZ. LDWZ >= 1, and if
JOBZ = 'V', LDWZ >= max(1,N).
WORK (workspace/output) COMPLEX_16 array, dimension (LWORK)
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The length of the array WORK. LWORK >= (NB+1)*N,
where NB is the max of the blocksize for ZHETRD.
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
RWORK (workspace) DOUBLE PRECISION array, dimension (7*N)
IWORK (workspace) INTEGER array, dimension (5*N)
IFAIL (output) INTEGER array, dimension (N)
If JOBZ = 'V', then if INFO = 0, the first M elements of
IFAIL are zero. If INFO > 0, then IFAIL contains the
indices of the eigenvectors that failed to converge.
If JOBZ = 'N', then IFAIL is not referenced.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
> 0: if INFO = i, then i eigenvectors failed to converge.
Their indices are stored in array IFAIL.
===================================================================== */
char uplo_[2] = {uplo, 0};
char jobz_[2] = {jobz, 0};
char range_[2] = {range, 0};
magma_int_t ione = 1;
char order[1];
magma_int_t indd, inde;
magma_int_t imax;
magma_int_t lopt, itmp1, indee;
magma_int_t lower, wantz;
magma_int_t i, j, jj, i__1;
magma_int_t alleig, valeig, indeig;
magma_int_t iscale, indibl;
magma_int_t indiwk, indisp, indtau;
magma_int_t indrwk, indwrk;
magma_int_t llwork, nsplit;
magma_int_t lquery;
magma_int_t iinfo;
double safmin;
double bignum;
double smlnum;
double eps, tmp1;
double anrm;
double sigma, d__1;
double rmin, rmax;
double *dwork;
/* Function Body */
lower = lapackf77_lsame(uplo_, MagmaLowerStr);
wantz = lapackf77_lsame(jobz_, MagmaVecStr);
alleig = lapackf77_lsame(range_, "A");
valeig = lapackf77_lsame(range_, "V");
indeig = lapackf77_lsame(range_, "I");
lquery = lwork == -1;
*info = 0;
if (! (wantz || lapackf77_lsame(jobz_, MagmaNoVecStr))) {
*info = -1;
} else if (! (alleig || valeig || indeig)) {
*info = -2;
} else if (! (lower || lapackf77_lsame(uplo_, MagmaUpperStr))) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (ldda < max(1,n)) {
*info = -6;
} else if (lddz < 1 || (wantz && lddz < n)) {
*info = -15;
} else if (ldwa < max(1,n)) {
*info = -17;
} else if (ldwz < 1 || (wantz && ldwz < n)) {
*info = -19;
} else {
if (valeig) {
if (n > 0 && vu <= vl) {
*info = -8;
}
} else if (indeig) {
if (il < 1 || il > max(1,n)) {
*info = -9;
} else if (iu < min(n,il) || iu > n) {
*info = -10;
}
}
}
magma_int_t nb = magma_get_zhetrd_nb(n);
lopt = n * (nb + 1);
work[0] = MAGMA_Z_MAKE( lopt, 0 );
if (lwork < lopt && ! lquery) {
*info = -21;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info));
return *info;
} else if (lquery) {
return *info;
}
*m = 0;
/* Check if matrix is very small then just call LAPACK on CPU, no need for GPU */
if (n <= 128) {
#ifdef ENABLE_DEBUG
printf("--------------------------------------------------------------\n");
printf(" warning matrix too small N=%d NB=%d, calling lapack on CPU \n", (int) n, (int) nb);
printf("--------------------------------------------------------------\n");
#endif
magmaDoubleComplex *a = (magmaDoubleComplex *) malloc( n * n * sizeof(magmaDoubleComplex) );
magma_zgetmatrix(n, n, da, ldda, a, n);
lapackf77_zheevx(jobz_, range_, uplo_,
&n, a, &n, &vl, &vu, &il, &iu, &abstol, m,
w, wz, &ldwz, work, &lwork,
rwork, iwork, ifail, info);
magma_zsetmatrix( n, n, a, n, da, ldda);
magma_zsetmatrix( n, *m, wz, ldwz, dz, lddz);
free(a);
return *info;
}
if (MAGMA_SUCCESS != magma_dmalloc( &dwork, n )) {
fprintf (stderr, "!!!! device memory allocation error (magma_zheevx_gpu)\n");
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
--w;
--work;
--rwork;
--iwork;
--ifail;
/* Get machine constants. */
safmin = lapackf77_dlamch("Safe minimum");
eps = lapackf77_dlamch("Precision");
smlnum = safmin / eps;
bignum = 1. / smlnum;
rmin = magma_dsqrt(smlnum);
rmax = magma_dsqrt(bignum);
/* Scale matrix to allowable range, if necessary. */
anrm = magmablas_zlanhe('M', uplo, n, da, ldda, dwork);
iscale = 0;
sigma = 1;
if (anrm > 0. && anrm < rmin) {
iscale = 1;
sigma = rmin / anrm;
} else if (anrm > rmax) {
iscale = 1;
sigma = rmax / anrm;
}
if (iscale == 1) {
d__1 = 1.;
magmablas_zlascl(uplo, 0, 0, 1., sigma, n, n, da,
ldda, info);
if (abstol > 0.) {
abstol *= sigma;
}
if (valeig) {
vl *= sigma;
vu *= sigma;
}
}
/* Call ZHETRD to reduce Hermitian matrix to tridiagonal form. */
indd = 1;
inde = indd + n;
indrwk = inde + n;
indtau = 1;
indwrk = indtau + n;
llwork = lwork - indwrk + 1;
#ifdef FAST_HEMV
magma_zhetrd2_gpu(uplo, n, da, ldda, &rwork[indd], &rwork[inde],
&work[indtau], wa, ldwa, &work[indwrk], llwork, dz, lddz*n, &iinfo);
#else
magma_zhetrd_gpu (uplo, n, da, ldda, &rwork[indd], &rwork[inde],
&work[indtau], wa, ldwa, &work[indwrk], llwork, &iinfo);
#endif
lopt = n + (magma_int_t)MAGMA_Z_REAL(work[indwrk]);
/* If all eigenvalues are desired and ABSTOL is less than or equal to
zero, then call DSTERF or ZUNGTR and ZSTEQR. If this fails for
some eigenvalue, then try DSTEBZ. */
if ((alleig || (indeig && il == 1 && iu == n)) && abstol <= 0.) {
blasf77_dcopy(&n, &rwork[indd], &ione, &w[1], &ione);
indee = indrwk + 2*n;
if (! wantz) {
i__1 = n - 1;
blasf77_dcopy(&i__1, &rwork[inde], &ione, &rwork[indee], &ione);
lapackf77_dsterf(&n, &w[1], &rwork[indee], info);
}
else {
lapackf77_zlacpy("A", &n, &n, wa, &ldwa, wz, &ldwz);
lapackf77_zungtr(uplo_, &n, wz, &ldwz, &work[indtau], &work[indwrk], &llwork, &iinfo);
i__1 = n - 1;
blasf77_dcopy(&i__1, &rwork[inde], &ione, &rwork[indee], &ione);
lapackf77_zsteqr(jobz_, &n, &w[1], &rwork[indee], wz, &ldwz, &rwork[indrwk], info);
if (*info == 0) {
for (i = 1; i <= n; ++i) {
ifail[i] = 0;
}
magma_zsetmatrix( n, n, wz, ldwz, dz, lddz );
}
}
if (*info == 0) {
*m = n;
}
}
/* Otherwise, call DSTEBZ and, if eigenvectors are desired, ZSTEIN. */
if (*m == 0) {
*info = 0;
if (wantz) {
*(unsigned char *)order = 'B';
} else {
*(unsigned char *)order = 'E';
}
indibl = 1;
indisp = indibl + n;
indiwk = indisp + n;
lapackf77_dstebz(range_, order, &n, &vl, &vu, &il, &iu, &abstol, &rwork[indd], &rwork[inde], m,
&nsplit, &w[1], &iwork[indibl], &iwork[indisp], &rwork[indrwk], &iwork[indiwk], info);
if (wantz) {
lapackf77_zstein(&n, &rwork[indd], &rwork[inde], m, &w[1], &iwork[indibl], &iwork[indisp],
wz, &ldwz, &rwork[indrwk], &iwork[indiwk], &ifail[1], info);
magma_zsetmatrix( n, *m, wz, ldwz, dz, lddz );
/* Apply unitary matrix used in reduction to tridiagonal
form to eigenvectors returned by ZSTEIN. */
magma_zunmtr_gpu(MagmaLeft, uplo, MagmaNoTrans, n, *m, da, ldda, &work[indtau],
dz, lddz, wa, ldwa, &iinfo);
}
}
/* If matrix was scaled, then rescale eigenvalues appropriately. */
if (iscale == 1) {
if (*info == 0) {
imax = *m;
} else {
imax = *info - 1;
}
d__1 = 1. / sigma;
blasf77_dscal(&imax, &d__1, &w[1], &ione);
}
/* If eigenvalues are not in order, then sort them, along with
eigenvectors. */
if (wantz) {
for (j = 1; j <= *m-1; ++j) {
i = 0;
tmp1 = w[j];
for (jj = j + 1; jj <= *m; ++jj) {
if (w[jj] < tmp1) {
i = jj;
tmp1 = w[jj];
}
}
if (i != 0) {
itmp1 = iwork[indibl + i - 1];
w[i] = w[j];
iwork[indibl + i - 1] = iwork[indibl + j - 1];
w[j] = tmp1;
iwork[indibl + j - 1] = itmp1;
magma_zswap(n, dz + (i-1)*lddz, ione, dz + (j-1)*lddz, ione);
if (*info != 0) {
itmp1 = ifail[i];
ifail[i] = ifail[j];
ifail[j] = itmp1;
}
}
}
}
/* Set WORK(1) to optimal complex workspace size. */
work[1] = MAGMA_Z_MAKE( lopt, 0 );
return *info;
} /* magma_zheevx_gpu */
extern "C" magma_int_t
magma_dlaex1(magma_int_t n, double* d, double* q, magma_int_t ldq,
magma_int_t* indxq, double rho, magma_int_t cutpnt,
double* work, magma_int_t* iwork, magmaDouble_ptr dwork,
magma_vec_t range, double vl, double vu,
magma_int_t il, magma_int_t iu, magma_int_t* info, magma_queue_t queue)
{
/*
-- clMAGMA (version 1.0.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
September 2012
.. Scalar Arguments ..
CHARACTER RANGE
INTEGER IL, IU, CUTPNT, INFO, LDQ, N
DOUBLE PRECISION RHO, VL, VU
..
.. Array Arguments ..
INTEGER INDXQ( * ), iwork[* )
DOUBLE PRECISION D( * ), Q( LDQ, * ), WORK( * ), DWORK( * )
..
Purpose
=======
DLAEX1 computes the updated eigensystem of a diagonal
matrix after modification by a rank-one symmetric matrix.
T = Q(in) ( D(in) + RHO * Z*Z' ) Q'(in) = Q(out) * D(out) * Q'(out)
where Z = Q'u, u is a vector of length N with ones in the
CUTPNT and CUTPNT + 1 th elements and zeros elsewhere.
The eigenvectors of the original matrix are stored in Q, and the
eigenvalues are in D. The algorithm consists of three stages:
The first stage consists of deflating the size of the problem
when there are multiple eigenvalues or if there is a zero in
the Z vector. For each such occurence the dimension of the
secular equation problem is reduced by one. This stage is
performed by the routine DLAED2.
The second stage consists of calculating the updated
eigenvalues. This is done by finding the roots of the secular
equation via the routine DLAED4 (as called by DLAED3).
This routine also calculates the eigenvectors of the current
problem.
The final stage consists of computing the updated eigenvectors
directly using the updated eigenvalues. The eigenvectors for
the current problem are multiplied with the eigenvectors from
the overall problem.
Arguments
=========
N (input) INTEGER
The dimension of the symmetric tridiagonal matrix. N >= 0.
D (input/output) DOUBLE PRECISION array, dimension (N)
On entry, the eigenvalues of the rank-1-perturbed matrix.
On exit, the eigenvalues of the repaired matrix.
Q (input/output) DOUBLE PRECISION array, dimension (LDQ,N)
On entry, the eigenvectors of the rank-1-perturbed matrix.
On exit, the eigenvectors of the repaired tridiagonal matrix.
LDQ (input) INTEGER
The leading dimension of the array Q. LDQ >= max(1,N).
INDXQ (input/output) INTEGER array, dimension (N)
On entry, the permutation which separately sorts the two
subproblems in D into ascending order.
On exit, the permutation which will reintegrate the
subproblems back into sorted order,
i.e. D( INDXQ( I = 1, N ) ) will be in ascending order.
RHO (input) DOUBLE PRECISION
The subdiagonal entry used to create the rank-1 modification.
CUTPNT (input) INTEGER
The location of the last eigenvalue in the leading sub-matrix.
min(1,N) <= CUTPNT <= N/2.
WORK (workspace) DOUBLE PRECISION array, dimension (4*N + N**2)
IWORK (workspace) INTEGER array, dimension (4*N)
DWORK (device workspace) DOUBLE PRECISION array, dimension (3*N*N/2+3*N)
RANGE (input) CHARACTER*1
= 'A': all eigenvalues will be found.
= 'V': all eigenvalues in the half-open interval (VL,VU]
will be found.
= 'I': the IL-th through IU-th eigenvalues will be found.
VL (input) DOUBLE PRECISION
VU (input) DOUBLE PRECISION
if RANGE='V', the lower and upper bounds of the interval to
be searched for eigenvalues. VL < VU.
Not referenced if RANGE = 'A' or 'I'.
IL (input) INTEGER
IU (input) INTEGER
if RANGE='I', the indices (in ascending order) of the
smallest and largest eigenvalues to be returned.
1 <= IL <= IU <= N, if N > 0; IL = 1 and IU = 0 if N = 0.
Not referenced if RANGE = 'A' or 'V'.
INFO (output) INTEGER
= 0: successful exit.
< 0: if INFO = -i, the i-th argument had an illegal value.
> 0: if INFO = 1, an eigenvalue did not converge
Further Details
===============
Based on contributions by
Jeff Rutter, Computer Science Division, University of California
at Berkeley, USA
Modified by Francoise Tisseur, University of Tennessee.
=====================================================================
*/
magma_int_t coltyp, i, idlmda;
magma_int_t indx, indxc, indxp;
magma_int_t iq2, is, iw, iz, k, tmp;
magma_int_t ione = 1;
// Test the input parameters.
*info = 0;
if( n < 0 )
*info = -1;
else if( ldq < max(1, n) )
*info = -4;
else if( min( 1, n/2 ) > cutpnt || n/2 < cutpnt )
*info = -7;
if( *info != 0 ){
magma_xerbla( __func__, -*info );
return MAGMA_ERR_ILLEGAL_VALUE;
}
// Quick return if possible
if( n == 0 )
return MAGMA_SUCCESS;
// The following values are integer pointers which indicate
// the portion of the workspace
// used by a particular array in DLAED2 and DLAED3.
iz = 0;
idlmda = iz + n;
iw = idlmda + n;
iq2 = iw + n;
indx = 0;
indxc = indx + n;
coltyp = indxc + n;
indxp = coltyp + n;
// Form the z-vector which consists of the last row of Q_1 and the
// first row of Q_2.
blasf77_dcopy( &cutpnt, Q(cutpnt-1, 0), &ldq, &work[iz], &ione);
tmp = n-cutpnt;
blasf77_dcopy( &tmp, Q(cutpnt, cutpnt), &ldq, &work[iz+cutpnt], &ione);
// Deflate eigenvalues.
magma_dlaed2(&k, &n, &cutpnt, d, q, &ldq, indxq, &rho, &work[iz],
&work[idlmda], &work[iw], &work[iq2],
&iwork[indx], &iwork[indxc], &iwork[indxp],
&iwork[coltyp], info);
if( *info != 0 )
return MAGMA_SUCCESS;
// Solve Secular Equation.
if( k != 0 ){
is = (iwork[coltyp]+iwork[coltyp+1])*cutpnt + (iwork[coltyp+1]+iwork[coltyp+2])*(n-cutpnt) + iq2;
magma_dlaex3(k, n, cutpnt, d, q, ldq, rho,
&work[idlmda], &work[iq2], &iwork[indxc],
&iwork[coltyp], &work[iw], &work[is],
indxq, dwork, range, vl, vu, il, iu, info, queue );
if( *info != 0 )
return MAGMA_SUCCESS;
}
else {
for (i = 0; i<n; ++i)
indxq[i] = i+1;
}
return MAGMA_SUCCESS;
} /* magma_dlaex1 */
/**
Purpose
-------
DLAHR2 reduces the first NB columns of a real general n-BY-(n-k+1)
matrix A so that elements below the k-th subdiagonal are zero. The
reduction is performed by an orthogonal similarity transformation
Q' * A * Q. The routine returns the matrices V and T which determine
Q as a block reflector I - V*T*V', and also the matrix Y = A * V.
(Note this is different than LAPACK, which computes Y = A * V * T.)
This is an auxiliary routine called by DGEHRD.
Arguments
---------
@param[in]
n INTEGER
The order of the matrix A.
@param[in]
k INTEGER
The offset for the reduction. Elements below the k-th
subdiagonal in the first NB columns are reduced to zero.
K < N.
@param[in]
nb INTEGER
The number of columns to be reduced.
@param[in,out]
A DOUBLE_PRECISION array, dimension (LDA,N-K+1)
On entry, the n-by-(n-k+1) general matrix A.
On exit, the elements on and above the k-th subdiagonal in
the first NB columns are overwritten with the corresponding
elements of the reduced matrix; the elements below the k-th
subdiagonal, with the array TAU, represent the matrix Q as a
product of elementary reflectors. The other columns of A are
unchanged. See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
tau DOUBLE_PRECISION array, dimension (NB)
The scalar factors of the elementary reflectors. See Further
Details.
@param[out]
T DOUBLE_PRECISION array, dimension (LDT,NB)
The upper triangular matrix T.
@param[in]
ldt INTEGER
The leading dimension of the array T. LDT >= NB.
@param[out]
Y DOUBLE_PRECISION array, dimension (LDY,NB)
The n-by-nb matrix Y.
@param[in]
ldy INTEGER
The leading dimension of the array Y. LDY >= N.
@param[in,out]
data Structure with pointers to dA, dT, dV, dW, dY
which are distributed across multiple GPUs.
Further Details
---------------
The matrix Q is represented as a product of nb elementary reflectors
Q = H(1) H(2) . . . H(nb).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i+k-1) = 0, v(i+k) = 1; v(i+k+1:n) is stored on exit in
A(i+k+1:n,i), and tau in TAU(i).
The elements of the vectors v together form the (n-k+1)-by-nb matrix
V which is needed, with T and Y, to apply the transformation to the
unreduced part of the matrix, using an update of the form:
A := (I - V*T*V') * (A - Y*T*V').
The contents of A on exit are illustrated by the following example
with n = 7, k = 3 and nb = 2:
@verbatim
( a a a a a )
( a a a a a )
( a a a a a )
( h h a a a )
( v1 h a a a )
( v1 v2 a a a )
( v1 v2 a a a )
@endverbatim
where "a" denotes an element of the original matrix A, h denotes a
modified element of the upper Hessenberg matrix H, and vi denotes an
element of the vector defining H(i).
This implementation follows the hybrid algorithm and notations described in
S. Tomov and J. Dongarra, "Accelerating the reduction to upper Hessenberg
form through hybrid GPU-based computing," University of Tennessee Computer
Science Technical Report, UT-CS-09-642 (also LAPACK Working Note 219),
May 24, 2009.
@ingroup magma_dgeev_aux
********************************************************************/
extern "C" magma_int_t
magma_dlahr2_m(
magma_int_t n, magma_int_t k, magma_int_t nb,
double *A, magma_int_t lda,
double *tau,
double *T, magma_int_t ldt,
double *Y, magma_int_t ldy,
struct dgehrd_data* data )
{
#define A( i, j ) ( A + (i) + (j)*lda)
#define Y( i, j ) ( Y + (i) + (j)*ldy)
#define T( i, j ) ( T + (i) + (j)*ldt)
#define dA( d, i, j ) (data->A [d] + (i) + (j)*ldda)
#define dTi( d ) (data->Ti[d])
#define dV( d, i, j ) (data->V [d] + (i) + (j)*ldv )
#define dVd( d, i, j ) (data->Vd[d] + (i) + (j)*ldvd)
#define dY( d, i, j ) (data->Y [d] + (i) + (j)*ldda)
double c_zero = MAGMA_D_ZERO;
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
double tmp;
magma_int_t ngpu = data->ngpu;
magma_int_t ldda = data->ldda;
magma_int_t ldv = data->ldv;
magma_int_t ldvd = data->ldvd;
magma_int_t ione = 1;
magma_int_t d, dki1, dn, nblocks, gblock, lblock, lgid;
magma_int_t n_k_i_1, n_k;
double scale;
magma_int_t i;
double ei = MAGMA_D_ZERO;
magma_int_t info_data = 0;
magma_int_t *info = &info_data;
if (n < 0) {
*info = -1;
} else if (k < 0 || k >= n) {
*info = -2;
} else if (nb < 1 || nb > n) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
} else if (ldt < nb) {
*info = -8;
} else if (ldy < max(1,n)) {
*info = -10;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
// adjust from 1-based indexing
k -= 1;
// Function Body
if (n <= 1)
return 0;
// zero out current top block of V on all GPUs
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magmablasSetKernelStream( data->streams[d] );
magmablas_dlaset( MagmaFull, nb, nb, c_zero, c_zero, dV(d,k,0), ldv );
}
// set all Y=0
lapackf77_dlaset( "Full", &n, &nb, &c_zero, &c_zero, Y, &ldy );
for (i = 0; i < nb; ++i) {
n_k_i_1 = n - k - i - 1;
n_k = n - k;
if (i > 0) {
// Finish applying I - V * T * V' on right
tmp = MAGMA_D_NEGATE( tau[i-1] );
blasf77_daxpy( &n_k, &tmp, Y(k,i-1), &ione, A(k,i), &ione );
// Apply I - V * T' * V' to this column (call it b) from the
// left, using the last column of T as workspace, w.
//
// Let V = ( V1 ) and b = ( b1 ) (first i-1 rows)
// ( V2 ) ( b2 )
// where V1 is unit lower triangular
// w := b1 = A(k+1:k+i, i)
blasf77_dcopy( &i,
A(k+1,i), &ione,
T(0,nb-1), &ione );
// w := V1' * b1 = VA(k+1:k+i, 0:i-1)' * w
blasf77_dtrmv( "Lower", "Conj", "Unit", &i,
A(k+1,0), &lda,
T(0,nb-1), &ione );
// w := w + V2'*b2 = w + VA(k+i+1:n-1, 0:i-1)' * A(k+i+1:n-1, i)
blasf77_dgemv( "Conj", &n_k_i_1, &i,
&c_one, A(k+i+1,0), &lda,
A(k+i+1,i), &ione,
&c_one, T(0,nb-1), &ione );
// w := T'*w = T(0:i-1, 0:i-1)' * w
blasf77_dtrmv( "Upper", "Conj", "Non-unit", &i,
T(0,0), &ldt,
T(0,nb-1), &ione );
// b2 := b2 - V2*w = A(k+i+1:n-1, i) - VA(k+i+1:n-1, 0:i-1) * w
blasf77_dgemv( "No trans", &n_k_i_1, &i,
&c_neg_one, A(k+i+1,0), &lda,
T(0,nb-1), &ione,
&c_one, A(k+i+1,i), &ione );
// w := V1*w = VA(k+1:k+i, 0:i-1) * w
blasf77_dtrmv( "Lower", "No trans", "Unit", &i,
A(k+1,0), &lda,
T(0,nb-1), &ione );
// b1 := b1 - w = A(k+1:k+i-1, i) - w
blasf77_daxpy( &i,
&c_neg_one, T(0,nb-1), &ione,
A(k+1,i), &ione );
// Restore diagonal element, saved below during previous iteration
*A(k+i,i-1) = ei;
}
// Generate the elementary reflector H(i) to annihilate A(k+i+1:n-1,i)
lapackf77_dlarfg( &n_k_i_1,
A(k+i+1,i),
A(k+i+2,i), &ione, &tau[i] );
// Save diagonal element and set to one, to simplify multiplying by V
ei = *A(k+i+1,i);
*A(k+i+1,i) = c_one;
// compute yi = A vi = sum_g A{d} vi{d}
nblocks = (n-1) / nb / ngpu + 1;
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magmablasSetKernelStream( data->streams[d] );
// dV(k+i+1:n-1, i) = VA(k+i:n, i)
magma_dsetvector_async( n_k_i_1,
A(k+i+1,i), 1,
dV(d, k+i+1, i), 1, data->streams[d] );
// copy column of dV -> dVd, using block cyclic distribution.
// This assumes V and Vd have been padded so that
// a 2D matrix copy doesn't access them out-of-bounds
gblock = k / nb;
lblock = gblock / ngpu;
lgid = gblock % ngpu;
if ( d < lgid ) {
lblock += 1;
}
// treat V as (nb*ngpu) x nblock matrix, and Vd as nb x nblock matrix
magmablas_dlacpy( MagmaFull, nb, nblocks-lblock,
dV (d, d*nb + lblock*nb*ngpu, i), nb*ngpu,
dVd(d, 0 + lblock*nb, i), nb );
// convert global indices (k) to local indices (dk)
magma_indices_1D_bcyclic( nb, ngpu, d, k+i+1, n, &dki1, &dn );
// dY(k:n, i) = dA(k:n, k+i+1:n) * dV(k+i+1:n, i)
// skip if matrix is empty
// each GPU copies to different temporary vector in Y,
// which are summed in separate loop below
if ( dn-dki1 > 0 ) {
magma_dgemv( MagmaNoTrans, n-k, dn-dki1,
c_one, dA (d, k, dki1), ldda,
dVd(d, dki1, i), 1,
c_zero, dY (d, k, i), 1 );
// copy vector to host, storing in column nb+d of Y
// as temporary space (Y has >= nb+ngpu columns)
magma_dgetvector_async( n-k,
dY(d, k, i), 1,
Y(k, nb+d), 1, data->streams[d] );
}
}
// while GPU is doing above Ag*v...
// Compute T(0:i,i) = [ -tau T V' vi ]
// [ tau ]
// T(0:i-1, i) = -tau VA(k+i+1:n-1, 0:i-1)' VA(k+i+1:n-1, i)
scale = MAGMA_D_NEGATE( tau[i] );
blasf77_dgemv( "Conj", &n_k_i_1, &i,
&scale, A(k+i+1,0), &lda,
A(k+i+1,i), &ione,
&c_zero, T(0,i), &ione );
// T(0:i-1, i) = T(0:i-1, 0:i-1) * T(0:i-1, i)
blasf77_dtrmv( "Upper", "No trans", "Non-unit", &i,
T(0,0), &ldt,
T(0,i), &ione );
*T(i,i) = tau[i];
// apply reflectors to next column, A(i+1), on right only.
// one axpy will be required to finish this, in the next iteration above
if ( i > 0 && i+1 < nb ) {
// Update next column, A(k:n,i+1), applying Q on right.
// One axpy will be required to finish this, in the next iteration
// above, after yi is computed.
// This updates one more row than LAPACK does (row k),
// making block above panel an even multiple of nb.
// Use last column of T as workspace, w.
magma_int_t i1 = i+1;
// If real, conjugate row of V, and undo afterwards
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i1, A(k+i1,0), &lda );
#endif
// w = T(0:i, 0:i+1) * VA(k+i+1, 0:i+1)'
// T is now rectangular, so we use gemv instead of trmv as in lapack.
blasf77_dgemv( "No trans", &i, &i1,
&c_one, T(0,0), &ldt,
A(k+i1,0), &lda,
&c_zero, T(0,nb-1), &ione );
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i1, A(k+i1,0), &lda );
#endif
// A(k:n, i+1) -= Y(k:n, 0:i) * w
blasf77_dgemv( "No trans", &n_k, &i,
&c_neg_one, Y(k,0), &ldy,
T(0,nb-1), &ione,
&c_one, A(k,i1), &ione );
}
// yi = sum_g yi{d}
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magma_queue_sync( data->streams[d] );
magma_indices_1D_bcyclic( nb, ngpu, d, k+i+1, n, &dki1, &dn );
if ( dn-dki1 > 0 ) {
// yi = yi + yi{d}
blasf77_daxpy( &n_k, &c_one, Y(k,nb+d), &ione, Y(k,i), &ione );
}
}
}
// Restore diagonal element
*A(k+nb,nb-1) = ei;
// compute Y = Am V = sum_g Am{d} V{d} --- top part, Y(0:k-1,:)
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magmablasSetKernelStream( data->streams[d] );
// convert global indices (k) to local indices (dk)
magma_indices_1D_bcyclic( nb, ngpu, d, k+1, n, &dki1, &dn );
// dY(0:k, :) = dA(0:k, k+i+1:n-1) * dV(k+i+1:n-1, :)
// skip if matrix is empty
// each GPU copies to different temporary block in Y,
// which are summed in separate loop below
if ( dn-dki1 > 0 ) {
magma_dgemm( MagmaNoTrans, MagmaNoTrans, k, nb, dn-dki1,
c_one, dA (d, 0, dki1), ldda,
dVd(d, dki1, 0), ldvd,
c_zero, dY (d, 0, 0), ldda );
// copy result to host, storing in columns [nb + nb*d : nb + nb*(d+1)] of Y
// as temporary space (Y has nb + nb*ngpu columns)
magma_dgetmatrix_async( k, nb,
dY(d, 0, 0), ldda,
Y(0,nb+nb*d), ldy, data->streams[d] );
}
}
// Y = sum_g Y{d}
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magma_queue_sync( 0 );
magma_indices_1D_bcyclic( nb, ngpu, d, k+1, n, &dki1, &dn );
if ( dn-dki1 > 0 ) {
// Y = Y + Am V
for( i = 0; i < nb; ++i ) {
blasf77_daxpy( &k, &c_one, Y(0,nb+nb*d+i), &ione, Y(0,i), &ione );
}
}
}
// copy Y and T matrices to GPUs
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magma_dsetmatrix_async( n, nb, Y, ldy, dY(d, 0, 0), ldda, data->streams[d] );
magma_dsetmatrix_async( nb, nb, T, nb, dTi(d), nb, data->streams[d] );
}
return 0;
} /* magma_dlahr2 */
extern "C" magma_int_t
magma_dlaex0(magma_int_t n, double* d, double* e, double* q, magma_int_t ldq,
double* work, magma_int_t* iwork, magmaDouble_ptr dwork,
magma_vec_t range, double vl, double vu,
magma_int_t il, magma_int_t iu, magma_int_t* info, magma_queue_t queue)
{
/*
-- MAGMA (version 1.1.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date January 2014
.. Scalar Arguments ..
CHARACTER RANGE
INTEGER IL, IU, INFO, LDQ, N
DOUBLE PRECISION VL, VU
..
.. Array Arguments ..
INTEGER IWORK( * )
DOUBLE PRECISION D( * ), E( * ), Q( LDQ, * ),
$ WORK( * ), DWORK( * )
..
Purpose
=======
DLAEX0 computes all eigenvalues and the choosen eigenvectors of a
symmetric tridiagonal matrix using the divide and conquer method.
Arguments
=========
N (input) INTEGER
The dimension of the symmetric tridiagonal matrix. N >= 0.
D (input/output) DOUBLE PRECISION array, dimension (N)
On entry, the main diagonal of the tridiagonal matrix.
On exit, its eigenvalues.
E (input) DOUBLE PRECISION array, dimension (N-1)
The off-diagonal elements of the tridiagonal matrix.
On exit, E has been destroyed.
Q (input/output) DOUBLE PRECISION array, dimension (LDQ, N)
On entry, Q will be the identity matrix.
On exit, Q contains the eigenvectors of the
tridiagonal matrix.
LDQ (input) INTEGER
The leading dimension of the array Q. If eigenvectors are
desired, then LDQ >= max(1,N). In any case, LDQ >= 1.
WORK (workspace) DOUBLE PRECISION array,
the dimension of WORK must be at least 4*N + N**2.
IWORK (workspace) INTEGER array,
the dimension of IWORK must be at least 3 + 5*N.
DWORK (device workspace) DOUBLE PRECISION array, dimension (3*N*N/2+3*N)
RANGE (input) CHARACTER*1
= 'A': all eigenvalues will be found.
= 'V': all eigenvalues in the half-open interval (VL,VU]
will be found.
= 'I': the IL-th through IU-th eigenvalues will be found.
VL (input) DOUBLE PRECISION
VU (input) DOUBLE PRECISION
If RANGE='V', the lower and upper bounds of the interval to
be searched for eigenvalues. VL < VU.
Not referenced if RANGE = 'A' or 'I'.
IL (input) INTEGER
IU (input) INTEGER
If RANGE='I', the indices (in ascending order) of the
smallest and largest eigenvalues to be returned.
1 <= IL <= IU <= N, if N > 0; IL = 1 and IU = 0 if N = 0.
Not referenced if RANGE = 'A' or 'V'.
INFO (output) INTEGER
= 0: successful exit.
< 0: if INFO = -i, the i-th argument had an illegal value.
> 0: The algorithm failed to compute an eigenvalue while
working on the submatrix lying in rows and columns
INFO/(N+1) through mod(INFO,N+1).
Further Details
===============
Based on contributions by
Jeff Rutter, Computer Science Division, University of California
at Berkeley, USA
=====================================================================
*/
magma_int_t ione = 1;
magma_vec_t range_ = range;
magma_int_t curlvl, curprb, i, indxq;
magma_int_t j, k, matsiz, msd2, smlsiz;
magma_int_t submat, subpbs, tlvls;
// Test the input parameters.
*info = 0;
if( n < 0 )
*info = -1;
else if( ldq < max(1, n) )
*info = -5;
if( *info != 0 ){
magma_xerbla( __func__, -*info );
return MAGMA_ERR_ILLEGAL_VALUE;
}
// Quick return if possible
if(n == 0)
return MAGMA_SUCCESS;
smlsiz = get_dlaex0_smlsize();
// Determine the size and placement of the submatrices, and save in
// the leading elements of IWORK.
iwork[0] = n;
subpbs= 1;
tlvls = 0;
while (iwork[subpbs - 1] > smlsiz) {
for (j = subpbs; j > 0; --j){
iwork[2*j - 1] = (iwork[j-1]+1)/2;
iwork[2*j - 2] = iwork[j-1]/2;
}
++tlvls;
subpbs *= 2;
}
for (j=1; j<subpbs; ++j)
iwork[j] += iwork[j-1];
// Divide the matrix into SUBPBS submatrices of size at most SMLSIZ+1
// using rank-1 modifications (cuts).
for(i=0; i < subpbs-1; ++i){
submat = iwork[i];
d[submat-1] -= MAGMA_D_ABS(e[submat-1]);
d[submat] -= MAGMA_D_ABS(e[submat-1]);
}
indxq = 4*n + 3;
// Solve each submatrix eigenproblem at the bottom of the divide and
// conquer tree.
char char_I[] = {'I', 0};
//#define ENABLE_TIMER
#ifdef ENABLE_TIMER
magma_timestr_t start, end;
start = get_current_time();
#endif
for (i = 0; i < subpbs; ++i){
if(i == 0){
submat = 0;
matsiz = iwork[0];
} else {
submat = iwork[i-1];
matsiz = iwork[i] - iwork[i-1];
}
lapackf77_dsteqr(char_I , &matsiz, &d[submat], &e[submat],
Q(submat, submat), &ldq, work, info); // change to edc?
if(*info != 0){
printf("info: %d\n, submat: %d\n", (int) *info, (int) submat);
*info = (submat+1)*(n+1) + submat + matsiz;
printf("info: %d\n", (int) *info);
return MAGMA_SUCCESS;
}
k = 1;
for(j = submat; j < iwork[i]; ++j){
iwork[indxq+j] = k;
++k;
}
}
#ifdef ENABLE_TIMER
end = get_current_time();
printf("for: dsteqr = %6.2f\n", GetTimerValue(start,end)/1000.);
#endif
// Successively merge eigensystems of adjacent submatrices
// into eigensystem for the corresponding larger matrix.
curlvl = 1;
while (subpbs > 1){
#ifdef ENABLE_TIMER
magma_timestr_t start, end;
start = get_current_time();
#endif
for (i=0; i<subpbs-1; i+=2){
if(i == 0){
submat = 0;
matsiz = iwork[1];
msd2 = iwork[0];
} else {
submat = iwork[i-1];
matsiz = iwork[i+1] - iwork[i-1];
msd2 = matsiz / 2;
}
// Merge lower order eigensystems (of size MSD2 and MATSIZ - MSD2)
// into an eigensystem of size MATSIZ.
// DLAEX1 is used only for the full eigensystem of a tridiagonal
// matrix.
if (matsiz == n)
range_=range;
else
// We need all the eigenvectors if it is not last step
range_= MagmaAllVec;
magma_dlaex1(matsiz, &d[submat], Q(submat, submat), ldq,
&iwork[indxq+submat], e[submat+msd2-1], msd2,
work, &iwork[subpbs], dwork,
range_, vl, vu, il, iu, info, queue);
if(*info != 0){
*info = (submat+1)*(n+1) + submat + matsiz;
return MAGMA_SUCCESS;
}
iwork[i/2]= iwork[i+1];
}
subpbs /= 2;
++curlvl;
#ifdef ENABLE_TIMER
end = get_current_time();
printf("%d: time: %6.2f\n", curlvl, GetTimerValue(start,end)/1000.);
#endif
}
// Re-merge the eigenvalues/vectors which were deflated at the final
// merge step.
for(i = 0; i<n; ++i){
j = iwork[indxq+i] - 1;
work[i] = d[j];
blasf77_dcopy(&n, Q(0, j), &ione, &work[ n*(i+1) ], &ione);
}
blasf77_dcopy(&n, work, &ione, d, &ione);
char char_A[] = {'A',0};
lapackf77_dlacpy ( char_A, &n, &n, &work[n], &n, q, &ldq );
return MAGMA_SUCCESS;
} /* magma_dlaex0 */
/**
Purpose
-------
DLAEX1 computes the updated eigensystem of a diagonal
matrix after modification by a rank-one symmetric matrix.
T = Q(in) ( D(in) + RHO * Z*Z' ) Q'(in) = Q(out) * D(out) * Q'(out)
where Z = Q'u, u is a vector of length N with ones in the
CUTPNT and CUTPNT + 1 th elements and zeros elsewhere.
The eigenvectors of the original matrix are stored in Q, and the
eigenvalues are in D. The algorithm consists of three stages:
The first stage consists of deflating the size of the problem
when there are multiple eigenvalues or if there is a zero in
the Z vector. For each such occurence the dimension of the
secular equation problem is reduced by one. This stage is
performed by the routine DLAED2.
The second stage consists of calculating the updated
eigenvalues. This is done by finding the roots of the secular
equation via the routine DLAED4 (as called by DLAED3).
This routine also calculates the eigenvectors of the current
problem.
The final stage consists of computing the updated eigenvectors
directly using the updated eigenvalues. The eigenvectors for
the current problem are multiplied with the eigenvectors from
the overall problem.
Arguments
---------
@param[in]
nrgpu INTEGER
Number of GPUs to use.
@param[in]
n INTEGER
The dimension of the symmetric tridiagonal matrix. N >= 0.
@param[in,out]
d DOUBLE PRECISION array, dimension (N)
On entry, the eigenvalues of the rank-1-perturbed matrix.
On exit, the eigenvalues of the repaired matrix.
@param[in,out]
Q DOUBLE PRECISION array, dimension (LDQ,N)
On entry, the eigenvectors of the rank-1-perturbed matrix.
On exit, the eigenvectors of the repaired tridiagonal matrix.
@param[in]
ldq INTEGER
The leading dimension of the array Q. LDQ >= max(1,N).
@param[in,out]
indxq INTEGER array, dimension (N)
On entry, the permutation which separately sorts the two
subproblems in D into ascending order.
On exit, the permutation which will reintegrate the
subproblems back into sorted order,
i.e. D( INDXQ( I = 1, N ) ) will be in ascending order.
@param[in]
rho DOUBLE PRECISION
The subdiagonal entry used to create the rank-1 modification.
@param[in]
cutpnt INTEGER
The location of the last eigenvalue in the leading sub-matrix.
min(1,N) <= CUTPNT <= N/2.
@param
work (workspace) DOUBLE PRECISION array, dimension (4*N + N**2)
@param
iwork (workspace) INTEGER array, dimension (4*N)
@param
dwork (devices workspaces) DOUBLE PRECISION array of arrays,
dimension NRGPU.
if NRGPU = 1 the dimension of the first workspace
should be (3*N*N/2+3*N)
otherwise the NRGPU workspaces should have the size
ceil((N-N1) * (N-N1) / floor(nrgpu/2)) +
NB * ((N-N1) + (N-N1) / floor(nrgpu/2))
@param
stream (device stream) magma_queue_t array,
dimension (MagmaMaxGPUs,2)
@param[in]
range magma_range_t
- = MagmaRangeAll: all eigenvalues will be found.
- = MagmaRangeV: all eigenvalues in the half-open interval (VL,VU]
will be found.
- = MagmaRangeI: the IL-th through IU-th eigenvalues will be found.
@param[in]
vl DOUBLE PRECISION
@param[in]
vu DOUBLE PRECISION
if RANGE=MagmaRangeV, the lower and upper bounds of the interval to
be searched for eigenvalues. VL < VU.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeI.
@param[in]
il INTEGER
@param[in]
iu INTEGER
if RANGE=MagmaRangeI, the indices (in ascending order) of the
smallest and largest eigenvalues to be returned.
1 <= IL <= IU <= N, if N > 0; IL = 1 and IU = 0 if N = 0.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeV.
@param[out]
info INTEGER
- = 0: successful exit.
- < 0: if INFO = -i, the i-th argument had an illegal value.
- > 0: if INFO = 1, an eigenvalue did not converge
Further Details
---------------
Based on contributions by
Jeff Rutter, Computer Science Division, University of California
at Berkeley, USA
Modified by Francoise Tisseur, University of Tennessee.
@ingroup magma_dsyev_aux
********************************************************************/
extern "C" magma_int_t
magma_dlaex1_m(magma_int_t nrgpu, magma_int_t n, double* d, double* Q, magma_int_t ldq,
magma_int_t* indxq, double rho, magma_int_t cutpnt,
double* work, magma_int_t* iwork, double** dwork,
magma_queue_t stream[MagmaMaxGPUs][2],
magma_range_t range, double vl, double vu,
magma_int_t il, magma_int_t iu, magma_int_t* info)
{
#define Q(i_,j_) (Q + (i_) + (j_)*ldq)
magma_int_t coltyp, i, idlmda;
magma_int_t indx, indxc, indxp;
magma_int_t iq2, is, iw, iz, k, tmp;
magma_int_t ione = 1;
// Test the input parameters.
*info = 0;
if ( n < 0 )
*info = -1;
else if ( ldq < max(1, n) )
*info = -4;
else if ( min( 1, n/2 ) > cutpnt || n/2 < cutpnt )
*info = -7;
if ( *info != 0 ) {
magma_xerbla( __func__, -(*info) );
return *info;
}
// Quick return if possible
if ( n == 0 )
return MAGMA_SUCCESS;
// The following values are integer pointers which indicate
// the portion of the workspace
// used by a particular array in DLAED2 and DLAED3.
iz = 0;
idlmda = iz + n;
iw = idlmda + n;
iq2 = iw + n;
indx = 0;
indxc = indx + n;
coltyp = indxc + n;
indxp = coltyp + n;
// Form the z-vector which consists of the last row of Q_1 and the
// first row of Q_2.
blasf77_dcopy( &cutpnt, Q(cutpnt-1, 0), &ldq, &work[iz], &ione);
tmp = n-cutpnt;
blasf77_dcopy( &tmp, Q(cutpnt, cutpnt), &ldq, &work[iz+cutpnt], &ione);
// Deflate eigenvalues.
lapackf77_dlaed2(&k, &n, &cutpnt, d, Q, &ldq, indxq, &rho, &work[iz],
&work[idlmda], &work[iw], &work[iq2],
&iwork[indx], &iwork[indxc], &iwork[indxp],
&iwork[coltyp], info);
if ( *info != 0 )
return MAGMA_SUCCESS;
// Solve Secular Equation.
if ( k != 0 ) {
is = (iwork[coltyp]+iwork[coltyp+1])*cutpnt + (iwork[coltyp+1]+iwork[coltyp+2])*(n-cutpnt) + iq2;
magma_dlaex3_m(nrgpu, k, n, cutpnt, d, Q, ldq, rho,
&work[idlmda], &work[iq2], &iwork[indxc],
&iwork[coltyp], &work[iw], &work[is],
indxq, dwork, stream, range, vl, vu, il, iu, info );
if ( *info != 0 )
return MAGMA_SUCCESS;
}
else {
for (i = 0; i < n; ++i)
indxq[i] = i+1;
}
return MAGMA_SUCCESS;
} /* magma_dlaex1_m */
/**
Purpose
-------
DLAHR2 reduces the first NB columns of a real general n-BY-(n-k+1)
matrix A so that elements below the k-th subdiagonal are zero. The
reduction is performed by an orthogonal similarity transformation
Q' * A * Q. The routine returns the matrices V and T which determine
Q as a block reflector I - V*T*V', and also the matrix Y = A * V.
(Note this is different than LAPACK, which computes Y = A * V * T.)
This is an auxiliary routine called by DGEHRD.
Arguments
---------
@param[in]
n INTEGER
The order of the matrix A.
@param[in]
k INTEGER
The offset for the reduction. Elements below the k-th
subdiagonal in the first NB columns are reduced to zero.
K < N.
@param[in]
nb INTEGER
The number of columns to be reduced.
@param[in,out]
dA DOUBLE PRECISION array on the GPU, dimension (LDDA,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.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
@param[out]
dV DOUBLE PRECISION array on the GPU, dimension (LDDV, NB)
On exit this n-by-nb array contains the Householder vectors of the transformation.
@param[in]
lddv INTEGER
The leading dimension of the array dV. LDDV >= max(1,N).
@param[in,out]
A DOUBLE PRECISION array, dimension (LDA,N-K+1)
On entry, the n-by-(n-k+1) general matrix A.
On exit, the elements on and above the k-th subdiagonal in
the first NB columns are overwritten with the corresponding
elements of the reduced matrix; the elements below the k-th
subdiagonal, with the array TAU, represent the matrix Q as a
product of elementary reflectors. The other columns of A are
unchanged. See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
tau DOUBLE PRECISION array, dimension (NB)
The scalar factors of the elementary reflectors. See Further
Details.
@param[out]
T DOUBLE PRECISION array, dimension (LDT,NB)
The upper triangular matrix T.
@param[in]
ldt INTEGER
The leading dimension of the array T. LDT >= NB.
@param[out]
Y DOUBLE PRECISION array, dimension (LDY,NB)
The n-by-nb matrix Y.
@param[in]
ldy INTEGER
The leading dimension of the array Y. LDY >= N.
@param[in]
queue magma_queue_t
Queue to execute in.
Further Details
---------------
The matrix Q is represented as a product of nb elementary reflectors
Q = H(1) H(2) . . . H(nb).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i+k-1) = 0, v(i+k) = 1; v(i+k+1:n) is stored on exit in
A(i+k+1:n,i), and tau in TAU(i).
The elements of the vectors v together form the (n-k+1)-by-nb matrix
V which is needed, with T and Y, to apply the transformation to the
unreduced part of the matrix, using an update of the form:
A := (I - V*T*V') * (A - Y*T*V').
The contents of A on exit are illustrated by the following example
with n = 7, k = 3 and nb = 2:
@verbatim
( a a a a a )
( a a a a a )
( a a a a a )
( h h a a a )
( v1 h a a a )
( v1 v2 a a a )
( v1 v2 a a a )
@endverbatim
where "a" denotes an element of the original matrix A, h denotes a
modified element of the upper Hessenberg matrix H, and vi denotes an
element of the vector defining H(i).
This implementation follows the hybrid algorithm and notations described in
S. Tomov and J. Dongarra, "Accelerating the reduction to upper Hessenberg
form through hybrid GPU-based computing," University of Tennessee Computer
Science Technical Report, UT-CS-09-642 (also LAPACK Working Note 219),
May 24, 2009.
@ingroup magma_dgeev_aux
********************************************************************/
extern "C" magma_int_t
magma_dlahr2(
magma_int_t n, magma_int_t k, magma_int_t nb,
magmaDouble_ptr dA, magma_int_t ldda,
magmaDouble_ptr dV, magma_int_t lddv,
double *A, magma_int_t lda,
double *tau,
double *T, magma_int_t ldt,
double *Y, magma_int_t ldy,
magma_queue_t queue )
{
#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(i_,j_) (dA + (i_) + (j_)*ldda)
#define dV(i_,j_) (dV + (i_) + (j_)*lddv)
double c_zero = MAGMA_D_ZERO;
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
magma_int_t ione = 1;
magma_int_t n_k_i_1, n_k;
double scale;
magma_int_t i;
double ei = MAGMA_D_ZERO;
magma_int_t info = 0;
if (n < 0) {
info = -1;
} else if (k < 0 || k > n) {
info = -2;
} else if (nb < 1 || nb > n) {
info = -3;
} else if (ldda < max(1,n)) {
info = -5;
} else if (lddv < max(1,n)) {
info = -7;
} else if (lda < max(1,n)) {
info = -9;
} else if (ldt < max(1,nb)) {
info = -12;
} else if (ldy < max(1,n)) {
info = -13;
}
if (info != 0) {
magma_xerbla( __func__, -(info) );
return info;
}
// adjust from 1-based indexing
k -= 1;
if (n <= 1)
return info;
for (i = 0; i < nb; ++i) {
n_k_i_1 = n - k - i - 1;
n_k = n - k;
if (i > 0) {
// Update A(k:n-1,i); Update i-th column of A - Y * T * V'
// This updates one more row than LAPACK does (row k),
// making the block above the panel an even multiple of nb.
// Use last column of T as workspace, w.
// w(0:i-1, nb-1) = VA(k+i, 0:i-1)'
blasf77_dcopy( &i,
A(k+i,0), &lda,
T(0,nb-1), &ione );
#ifdef COMPLEX
// If real, conjugate row of V.
lapackf77_dlacgv(&i, T(0,nb-1), &ione);
#endif
// w = T(0:i-1, 0:i-1) * w
blasf77_dtrmv( "Upper", "No trans", "No trans", &i,
T(0,0), &ldt,
T(0,nb-1), &ione );
// A(k:n-1, i) -= Y(k:n-1, 0:i-1) * w
blasf77_dgemv( "No trans", &n_k, &i,
&c_neg_one, Y(k,0), &ldy,
T(0,nb-1), &ione,
&c_one, A(k,i), &ione );
// Apply I - V * T' * V' to this column (call it b) from the
// left, using the last column of T as workspace, w.
//
// Let V = ( V1 ) and b = ( b1 ) (first i-1 rows)
// ( V2 ) ( b2 )
// where V1 is unit lower triangular
// w := b1 = A(k+1:k+i, i)
blasf77_dcopy( &i,
A(k+1,i), &ione,
T(0,nb-1), &ione );
// w := V1' * b1 = VA(k+1:k+i, 0:i-1)' * w
blasf77_dtrmv( "Lower", "Conj", "Unit", &i,
A(k+1,0), &lda,
T(0,nb-1), &ione );
// w := w + V2'*b2 = w + VA(k+i+1:n-1, 0:i-1)' * A(k+i+1:n-1, i)
blasf77_dgemv( "Conj", &n_k_i_1, &i,
&c_one, A(k+i+1,0), &lda,
A(k+i+1,i), &ione,
&c_one, T(0,nb-1), &ione );
// w := T'*w = T(0:i-1, 0:i-1)' * w
blasf77_dtrmv( "Upper", "Conj", "Non-unit", &i,
T(0,0), &ldt,
T(0,nb-1), &ione );
// b2 := b2 - V2*w = A(k+i+1:n-1, i) - VA(k+i+1:n-1, 0:i-1) * w
blasf77_dgemv( "No trans", &n_k_i_1, &i,
&c_neg_one, A(k+i+1,0), &lda,
T(0,nb-1), &ione,
&c_one, A(k+i+1,i), &ione );
// w := V1*w = VA(k+1:k+i, 0:i-1) * w
blasf77_dtrmv( "Lower", "No trans", "Unit", &i,
A(k+1,0), &lda,
T(0,nb-1), &ione );
// b1 := b1 - w = A(k+1:k+i-1, i) - w
blasf77_daxpy( &i,
&c_neg_one, T(0,nb-1), &ione,
A(k+1,i), &ione );
// Restore diagonal element, saved below during previous iteration
*A(k+i,i-1) = ei;
}
// Generate the elementary reflector H(i) to annihilate A(k+i+1:n-1,i)
lapackf77_dlarfg( &n_k_i_1,
A(k+i+1,i),
A(k+i+2,i), &ione, &tau[i] );
// Save diagonal element and set to one, to simplify multiplying by V
ei = *A(k+i+1,i);
*A(k+i+1,i) = c_one;
// dV(i+1:n-k-1, i) = VA(k+i+1:n-1, i)
magma_dsetvector( n_k_i_1,
A(k+i+1,i), 1,
dV(i+1,i), 1, queue );
// Compute Y(k+1:n,i) = A vi
// dA(k:n-1, i) = dA(k:n-1, i+1:n-k-1) * dV(i+1:n-k-1, i)
magma_dgemv( MagmaNoTrans, n_k, n_k_i_1,
c_one, dA(k,i+1), ldda,
dV(i+1,i), ione,
c_zero, dA(k,i), ione, queue );
// Compute T(0:i,i) = [ -tau T V' vi ]
// [ tau ]
// T(0:i-1, i) = -tau VA(k+i+1:n-1, 0:i-1)' VA(k+i+1:n-1, i)
scale = MAGMA_D_NEGATE( tau[i]);
blasf77_dgemv( "Conj", &n_k_i_1, &i,
&scale, A(k+i+1,0), &lda,
A(k+i+1,i), &ione,
&c_zero, T(0,i), &ione );
// T(0:i-1, i) = T(0:i-1, 0:i-1) * T(0:i-1, i)
blasf77_dtrmv( "Upper", "No trans", "Non-unit", &i,
T(0,0), &ldt,
T(0,i), &ione );
*T(i,i) = tau[i];
// Y(k:n-1, i) = dA(k:n-1, i)
magma_dgetvector( n-k,
dA(k,i), 1,
Y(k,i), 1, queue );
}
// Restore diagonal element
*A(k+nb,nb-1) = ei;
return info;
} /* magma_dlahr2 */
/**
Purpose
-------
DLAEX3 finds the roots of the secular equation, as defined by the
values in D, W, and RHO, between 1 and K. It makes the
appropriate calls to DLAED4 and then updates the eigenvectors by
multiplying the matrix of eigenvectors of the pair of eigensystems
being combined by the matrix of eigenvectors of the K-by-K system
which is solved here.
It is used in the last step when only a part of the eigenvectors
is required.
It compute only the required part of the eigenvectors and the rest
is not used.
This code makes very mild assumptions about floating point
arithmetic. It will work on machines with a guard digit in
add/subtract, or on those binary machines without guard digits
which subtract like the Cray X-MP, Cray Y-MP, Cray C-90, or Cray-2.
It could conceivably fail on hexadecimal or decimal machines
without guard digits, but we know of none.
Arguments
---------
@param[in]
k INTEGER
The number of terms in the rational function to be solved by
DLAED4. K >= 0.
@param[in]
n INTEGER
The number of rows and columns in the Q matrix.
N >= K (deflation may result in N > K).
@param[in]
n1 INTEGER
The location of the last eigenvalue in the leading submatrix.
min(1,N) <= N1 <= N/2.
@param[out]
d DOUBLE PRECISION array, dimension (N)
D(I) contains the updated eigenvalues for
1 <= I <= K.
@param[out]
Q DOUBLE PRECISION array, dimension (LDQ,N)
Initially the first K columns are used as workspace.
On output the columns ??? to ??? contain
the updated eigenvectors.
@param[in]
ldq INTEGER
The leading dimension of the array Q. LDQ >= max(1,N).
@param[in]
rho DOUBLE PRECISION
The value of the parameter in the rank one update equation.
RHO >= 0 required.
@param[in,out]
dlamda DOUBLE PRECISION array, dimension (K)
The first K elements of this array contain the old roots
of the deflated updating problem. These are the poles
of the secular equation. May be changed on output by
having lowest order bit set to zero on Cray X-MP, Cray Y-MP,
Cray-2, or Cray C-90, as described above.
@param[in]
Q2 DOUBLE PRECISION array, dimension (LDQ2, N)
The first K columns of this matrix contain the non-deflated
eigenvectors for the split problem.
TODO what is LDQ2?
@param[in]
indx INTEGER array, dimension (N)
The permutation used to arrange the columns of the deflated
Q matrix into three groups (see DLAED2).
The rows of the eigenvectors found by DLAED4 must be likewise
permuted before the matrix multiply can take place.
@param[in]
ctot INTEGER array, dimension (4)
A count of the total number of the various types of columns
in Q, as described in INDX. The fourth column type is any
column which has been deflated.
@param[in,out]
w DOUBLE PRECISION array, dimension (K)
The first K elements of this array contain the components
of the deflation-adjusted updating vector. Destroyed on
output.
@param
s (workspace) DOUBLE PRECISION array, dimension (N1 + 1)*K
Will contain the eigenvectors of the repaired matrix which
will be multiplied by the previously accumulated eigenvectors
to update the system.
@param[out]
indxq INTEGER array, dimension (N)
On exit, the permutation which will reintegrate the
subproblems back into sorted order,
i.e. D( INDXQ( I = 1, N ) ) will be in ascending order.
@param
dwork (workspace) DOUBLE PRECISION array, dimension (3*N*N/2+3*N)
@param[in]
range magma_range_t
- = MagmaRangeAll: all eigenvalues will be found.
- = MagmaRangeV: all eigenvalues in the half-open interval (VL,VU]
will be found.
- = MagmaRangeI: the IL-th through IU-th eigenvalues will be found.
TODO verify range, vl, vu, il, iu -- copied from dlaex1.
@param[in]
vl DOUBLE PRECISION
@param[in]
vu DOUBLE PRECISION
if RANGE=MagmaRangeV, the lower and upper bounds of the interval to
be searched for eigenvalues. VL < VU.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeI.
@param[in]
il INTEGER
@param[in]
iu INTEGER
if RANGE=MagmaRangeI, the indices (in ascending order) of the
smallest and largest eigenvalues to be returned.
1 <= IL <= IU <= N, if N > 0; IL = 1 and IU = 0 if N = 0.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeV.
@param[out]
info INTEGER
- = 0: successful exit.
- < 0: if INFO = -i, the i-th argument had an illegal value.
- > 0: if INFO = 1, an eigenvalue did not converge
Further Details
---------------
Based on contributions by
Jeff Rutter, Computer Science Division, University of California
at Berkeley, USA
Modified by Francoise Tisseur, University of Tennessee.
@ingroup magma_dsyev_aux
********************************************************************/
extern "C" magma_int_t
magma_dlaex3(magma_int_t k, magma_int_t n, magma_int_t n1, double* d,
double* Q, magma_int_t ldq, double rho,
double* dlamda, double* Q2, magma_int_t* indx,
magma_int_t* ctot, double* w, double* s, magma_int_t* indxq,
double* dwork,
magma_range_t range, double vl, double vu, magma_int_t il, magma_int_t iu,
magma_int_t* info )
{
#define Q(i_,j_) (Q + (i_) + (j_)*ldq)
double d_one = 1.;
double d_zero = 0.;
magma_int_t ione = 1;
magma_int_t ineg_one = -1;
magma_int_t iil, iiu, rk;
double* dq2= dwork;
double* ds = dq2 + n*(n/2+1);
double* dq = ds + n*(n/2+1);
magma_int_t lddq = n/2 + 1;
magma_int_t i, iq2, j, n12, n2, n23, tmp, lq2;
double temp;
magma_int_t alleig, valeig, indeig;
alleig = (range == MagmaRangeAll);
valeig = (range == MagmaRangeV);
indeig = (range == MagmaRangeI);
*info = 0;
if (k < 0)
*info=-1;
else if (n < k)
*info=-2;
else if (ldq < max(1,n))
*info=-6;
else if (! (alleig || valeig || indeig))
*info = -15;
else {
if (valeig) {
if (n > 0 && vu <= vl)
*info = -17;
}
else if (indeig) {
if (il < 1 || il > max(1,n))
*info = -18;
else if (iu < min(n,il) || iu > n)
*info = -19;
}
}
if (*info != 0) {
magma_xerbla(__func__, -(*info));
return *info;
}
// Quick return if possible
if (k == 0)
return *info;
/*
Modify values DLAMDA(i) to make sure all DLAMDA(i)-DLAMDA(j) can
be computed with high relative accuracy (barring over/underflow).
This is a problem on machines without a guard digit in
add/subtract (Cray XMP, Cray YMP, Cray C 90 and Cray 2).
The following code replaces DLAMDA(I) by 2*DLAMDA(I)-DLAMDA(I),
which on any of these machines zeros out the bottommost
bit of DLAMDA(I) if it is 1; this makes the subsequent
subtractions DLAMDA(I)-DLAMDA(J) unproblematic when cancellation
occurs. On binary machines with a guard digit (almost all
machines) it does not change DLAMDA(I) at all. On hexadecimal
and decimal machines with a guard digit, it slightly
changes the bottommost bits of DLAMDA(I). It does not account
for hexadecimal or decimal machines without guard digits
(we know of none). We use a subroutine call to compute
2*DLAMBDA(I) to prevent optimizing compilers from eliminating
this code.*/
n2 = n - n1;
n12 = ctot[0] + ctot[1];
n23 = ctot[1] + ctot[2];
iq2 = n1 * n12;
lq2 = iq2 + n2 * n23;
magma_dsetvector_async( lq2, Q2, 1, dq2, 1, NULL );
#ifdef _OPENMP
/////////////////////////////////////////////////////////////////////////////////
//openmp implementation
/////////////////////////////////////////////////////////////////////////////////
magma_timer_t time=0;
timer_start( time );
#pragma omp parallel private(i, j, tmp, temp)
{
magma_int_t id = omp_get_thread_num();
magma_int_t tot = omp_get_num_threads();
magma_int_t ib = ( id * k) / tot; //start index of local loop
magma_int_t ie = ((id+1) * k) / tot; //end index of local loop
magma_int_t ik = ie - ib; //number of local indices
for (i = ib; i < ie; ++i)
dlamda[i]=lapackf77_dlamc3(&dlamda[i], &dlamda[i]) - dlamda[i];
for (j = ib; j < ie; ++j) {
magma_int_t tmpp=j+1;
magma_int_t iinfo = 0;
lapackf77_dlaed4(&k, &tmpp, dlamda, w, Q(0,j), &rho, &d[j], &iinfo);
// If the zero finder fails, the computation is terminated.
if (iinfo != 0) {
#pragma omp critical (info)
*info=iinfo;
break;
}
}
#pragma omp barrier
if (*info == 0) {
#pragma omp single
{
//Prepare the INDXQ sorting permutation.
magma_int_t nk = n - k;
lapackf77_dlamrg( &k, &nk, d, &ione, &ineg_one, indxq);
//compute the lower and upper bound of the non-deflated eigenvectors
if (valeig)
magma_dvrange(k, d, &iil, &iiu, vl, vu);
else if (indeig)
magma_dirange(k, indxq, &iil, &iiu, il, iu);
else {
iil = 1;
iiu = k;
}
rk = iiu - iil + 1;
}
if (k == 2) {
#pragma omp single
{
for (j = 0; j < k; ++j) {
w[0] = *Q(0,j);
w[1] = *Q(1,j);
i = indx[0] - 1;
*Q(0,j) = w[i];
i = indx[1] - 1;
*Q(1,j) = w[i];
}
}
}
else if (k != 1) {
// Compute updated W.
blasf77_dcopy( &ik, &w[ib], &ione, &s[ib], &ione);
// Initialize W(I) = Q(I,I)
tmp = ldq + 1;
blasf77_dcopy( &ik, Q(ib,ib), &tmp, &w[ib], &ione);
for (j = 0; j < k; ++j) {
magma_int_t i_tmp = min(j, ie);
for (i = ib; i < i_tmp; ++i)
w[i] = w[i] * ( *Q(i, j) / ( dlamda[i] - dlamda[j] ) );
i_tmp = max(j+1, ib);
for (i = i_tmp; i < ie; ++i)
w[i] = w[i] * ( *Q(i, j) / ( dlamda[i] - dlamda[j] ) );
}
for (i = ib; i < ie; ++i)
w[i] = copysign( sqrt( -w[i] ), s[i]);
#pragma omp barrier
//reduce the number of used threads to have enough S workspace
tot = min(n1, omp_get_num_threads());
if (id < tot) {
ib = ( id * rk) / tot + iil - 1;
ie = ((id+1) * rk) / tot + iil - 1;
ik = ie - ib;
}
else {
ib = -1;
ie = -1;
ik = -1;
}
// Compute eigenvectors of the modified rank-1 modification.
for (j = ib; j < ie; ++j) {
for (i = 0; i < k; ++i)
s[id*k + i] = w[i] / *Q(i,j);
temp = magma_cblas_dnrm2( k, s+id*k, 1 );
for (i = 0; i < k; ++i) {
magma_int_t iii = indx[i] - 1;
*Q(i,j) = s[id*k + iii] / temp;
}
}
}
}
}
if (*info != 0)
return *info;
timer_stop( time );
timer_printf( "eigenvalues/vector D+zzT = %6.2f\n", time );
#else
/////////////////////////////////////////////////////////////////////////////////
// Non openmp implementation
/////////////////////////////////////////////////////////////////////////////////
magma_timer_t time=0;
timer_start( time );
for (i = 0; i < k; ++i)
dlamda[i]=lapackf77_dlamc3(&dlamda[i], &dlamda[i]) - dlamda[i];
for (j = 0; j < k; ++j) {
magma_int_t tmpp=j+1;
magma_int_t iinfo = 0;
lapackf77_dlaed4(&k, &tmpp, dlamda, w, Q(0,j), &rho, &d[j], &iinfo);
// If the zero finder fails, the computation is terminated.
if (iinfo != 0)
*info=iinfo;
}
if (*info != 0)
return *info;
//Prepare the INDXQ sorting permutation.
magma_int_t nk = n - k;
lapackf77_dlamrg( &k, &nk, d, &ione, &ineg_one, indxq);
//compute the lower and upper bound of the non-deflated eigenvectors
if (valeig)
magma_dvrange(k, d, &iil, &iiu, vl, vu);
else if (indeig)
magma_dirange(k, indxq, &iil, &iiu, il, iu);
else {
iil = 1;
iiu = k;
}
rk = iiu - iil + 1;
if (k == 2) {
for (j = 0; j < k; ++j) {
w[0] = *Q(0,j);
w[1] = *Q(1,j);
i = indx[0] - 1;
*Q(0,j) = w[i];
i = indx[1] - 1;
*Q(1,j) = w[i];
}
}
else if (k != 1) {
// Compute updated W.
blasf77_dcopy( &k, w, &ione, s, &ione);
// Initialize W(I) = Q(I,I)
tmp = ldq + 1;
blasf77_dcopy( &k, Q, &tmp, w, &ione);
for (j = 0; j < k; ++j) {
for (i = 0; i < j; ++i)
w[i] = w[i] * ( *Q(i, j) / ( dlamda[i] - dlamda[j] ) );
for (i = j+1; i < k; ++i)
w[i] = w[i] * ( *Q(i, j) / ( dlamda[i] - dlamda[j] ) );
}
for (i = 0; i < k; ++i)
w[i] = copysign( sqrt( -w[i] ), s[i]);
// Compute eigenvectors of the modified rank-1 modification.
for (j = iil-1; j < iiu; ++j) {
for (i = 0; i < k; ++i)
s[i] = w[i] / *Q(i,j);
temp = magma_cblas_dnrm2( k, s, 1 );
for (i = 0; i < k; ++i) {
magma_int_t iii = indx[i] - 1;
*Q(i,j) = s[iii] / temp;
}
}
}
timer_stop( time );
timer_printf( "eigenvalues/vector D+zzT = %6.2f\n", time );
#endif //_OPENMP
// Compute the updated eigenvectors.
timer_start( time );
magma_queue_sync( NULL );
if (rk != 0) {
if ( n23 != 0 ) {
if (rk < magma_get_dlaed3_k()) {
lapackf77_dlacpy("A", &n23, &rk, Q(ctot[0],iil-1), &ldq, s, &n23);
blasf77_dgemm("N", "N", &n2, &rk, &n23, &d_one, &Q2[iq2], &n2,
s, &n23, &d_zero, Q(n1,iil-1), &ldq );
} else {
magma_dsetmatrix( n23, rk, Q(ctot[0],iil-1), ldq, ds, n23 );
magma_dgemm( MagmaNoTrans, MagmaNoTrans, n2, rk, n23, d_one, &dq2[iq2], n2, ds, n23, d_zero, dq, lddq);
magma_dgetmatrix( n2, rk, dq, lddq, Q(n1,iil-1), ldq );
}
} else
lapackf77_dlaset("A", &n2, &rk, &d_zero, &d_zero, Q(n1,iil-1), &ldq);
if ( n12 != 0 ) {
if (rk < magma_get_dlaed3_k()) {
lapackf77_dlacpy("A", &n12, &rk, Q(0,iil-1), &ldq, s, &n12);
blasf77_dgemm("N", "N", &n1, &rk, &n12, &d_one, Q2, &n1,
s, &n12, &d_zero, Q(0,iil-1), &ldq);
} else {
magma_dsetmatrix( n12, rk, Q(0,iil-1), ldq, ds, n12 );
magma_dgemm( MagmaNoTrans, MagmaNoTrans, n1, rk, n12, d_one, dq2, n1, ds, n12, d_zero, dq, lddq);
magma_dgetmatrix( n1, rk, dq, lddq, Q(0,iil-1), ldq );
}
} else
lapackf77_dlaset("A", &n1, &rk, &d_zero, &d_zero, Q(0,iil-1), &ldq);
}
timer_stop( time );
timer_printf( "gemms = %6.2f\n", time );
return *info;
} /* magma_dlaex3 */
/**
Purpose
-------
ZHEEVR computes selected eigenvalues and, optionally, eigenvectors
of a complex Hermitian matrix T. Eigenvalues and eigenvectors can
be selected by specifying either a range of values or a range of
indices for the desired eigenvalues.
Whenever possible, ZHEEVR calls ZSTEGR to compute the
eigenspectrum using Relatively Robust Representations. ZSTEGR
computes eigenvalues by the dqds algorithm, while orthogonal
eigenvectors are computed from various "good" L D L^T representations
(also known as Relatively Robust Representations). Gram-Schmidt
orthogonalization is avoided as far as possible. More specifically,
the various steps of the algorithm are as follows. For the i-th
unreduced block of T,
1. Compute T - sigma_i = L_i D_i L_i^T, such that L_i D_i L_i^T
is a relatively robust representation,
2. Compute the eigenvalues, lambda_j, of L_i D_i L_i^T to high
relative accuracy by the dqds algorithm,
3. If there is a cluster of close eigenvalues, "choose" sigma_i
close to the cluster, and go to step (a),
4. Given the approximate eigenvalue lambda_j of L_i D_i L_i^T,
compute the corresponding eigenvector by forming a
rank-revealing twisted factorization.
The desired accuracy of the output can be specified by the input
parameter ABSTOL.
For more details, see "A new O(n^2) algorithm for the symmetric
tridiagonal eigenvalue/eigenvector problem", by Inderjit Dhillon,
Computer Science Division Technical Report No. UCB//CSD-97-971,
UC Berkeley, May 1997.
Note 1 : ZHEEVR calls ZSTEGR when the full spectrum is requested
on machines which conform to the ieee-754 floating point standard.
ZHEEVR calls DSTEBZ and ZSTEIN on non-ieee machines and
when partial spectrum requests are made.
Normal execution of ZSTEGR may create NaNs and infinities and
hence may abort due to a floating point exception in environments
which do not handle NaNs and infinities in the ieee standard default
manner.
Arguments
---------
@param[in]
jobz magma_vec_t
- = MagmaNoVec: Compute eigenvalues only;
- = MagmaVec: Compute eigenvalues and eigenvectors.
@param[in]
range magma_range_t
- = MagmaRangeAll: all eigenvalues will be found.
- = MagmaRangeV: all eigenvalues in the half-open interval (VL,VU]
will be found.
- = MagmaRangeI: the IL-th through IU-th eigenvalues will be found.
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored;
- = MagmaLower: Lower triangle of A is stored.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
A COMPLEX_16 array, dimension (LDA, N)
On entry, the Hermitian matrix A. If UPLO = MagmaUpper, the
leading N-by-N upper triangular part of A contains the
upper triangular part of the matrix A. If UPLO = MagmaLower,
the leading N-by-N lower triangular part of A contains
the lower triangular part of the matrix A.
On exit, the lower triangle (if UPLO=MagmaLower) or the upper
triangle (if UPLO=MagmaUpper) of A, including the diagonal, is
destroyed.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[in]
vl DOUBLE PRECISION
@param[in]
vu DOUBLE PRECISION
If RANGE=MagmaRangeV, the lower and upper bounds of the interval to
be searched for eigenvalues. VL < VU.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeI.
@param[in]
il INTEGER
@param[in]
iu INTEGER
If RANGE=MagmaRangeI, the indices (in ascending order) of the
smallest and largest eigenvalues to be returned.
1 <= IL <= IU <= N, if N > 0; IL = 1 and IU = 0 if N = 0.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeV.
@param[in]
abstol DOUBLE PRECISION
The absolute error tolerance for the eigenvalues.
An approximate eigenvalue is accepted as converged
when it is determined to lie in an interval [a,b]
of width less than or equal to
ABSTOL + EPS * max( |a|,|b| ),
\n
where EPS is the machine precision. If ABSTOL is less than
or equal to zero, then EPS*|T| will be used in its place,
where |T| is the 1-norm of the tridiagonal matrix obtained
by reducing A to tridiagonal form.
\n
See "Computing Small Singular Values of Bidiagonal Matrices
with Guaranteed High Relative Accuracy," by Demmel and
Kahan, LAPACK Working Note #3.
\n
If high relative accuracy is important, set ABSTOL to
DLAMCH( 'Safe minimum' ). Doing so will guarantee that
eigenvalues are computed to high relative accuracy when
possible in future releases. The current code does not
make any guarantees about high relative accuracy, but
furutre releases will. See J. Barlow and J. Demmel,
"Computing Accurate Eigensystems of Scaled Diagonally
Dominant Matrices", LAPACK Working Note #7, for a discussion
of which matrices define their eigenvalues to high relative
accuracy.
@param[out]
m INTEGER
The total number of eigenvalues found. 0 <= M <= N.
If RANGE = MagmaRangeAll, M = N, and if RANGE = MagmaRangeI, M = IU-IL+1.
@param[out]
w DOUBLE PRECISION array, dimension (N)
The first M elements contain the selected eigenvalues in
ascending order.
@param[out]
Z COMPLEX_16 array, dimension (LDZ, max(1,M))
If JOBZ = MagmaVec, then if INFO = 0, the first M columns of Z
contain the orthonormal eigenvectors of the matrix A
corresponding to the selected eigenvalues, with the i-th
column of Z holding the eigenvector associated with W(i).
If JOBZ = MagmaNoVec, then Z is not referenced.
Note: the user must ensure that at least max(1,M) columns are
supplied in the array Z; if RANGE = MagmaRangeV, the exact value of M
is not known in advance and an upper bound must be used.
@param[in]
ldz INTEGER
The leading dimension of the array Z. LDZ >= 1, and if
JOBZ = MagmaVec, LDZ >= max(1,N).
@param[out]
isuppz INTEGER ARRAY, dimension ( 2*max(1,M) )
The support of the eigenvectors in Z, i.e., the indices
indicating the nonzero elements in Z. The i-th eigenvector
is nonzero only in elements ISUPPZ( 2*i-1 ) through
ISUPPZ( 2*i ).
__Implemented only for__ RANGE = MagmaRangeAll or MagmaRangeI and IU - IL = N - 1
@param[out]
work (workspace) COMPLEX_16 array, dimension (LWORK)
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The length of the array WORK. LWORK >= max(1,2*N).
For optimal efficiency, LWORK >= (NB+1)*N,
where NB is the max of the blocksize for ZHETRD and for
ZUNMTR as returned by ILAENV.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
rwork (workspace) DOUBLE PRECISION array, dimension (LRWORK)
On exit, if INFO = 0, RWORK[0] returns the optimal
(and minimal) LRWORK.
@param[in]
lrwork INTEGER
The length of the array RWORK. LRWORK >= max(1,24*N).
\n
If LRWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the RWORK array, returns
this value as the first entry of the RWORK array, and no error
message related to LRWORK is issued by XERBLA.
@param[out]
iwork (workspace) INTEGER array, dimension (LIWORK)
On exit, if INFO = 0, IWORK[0] returns the optimal
(and minimal) LIWORK.
@param[in]
liwork INTEGER
The dimension of the array IWORK. LIWORK >= max(1,10*N).
\n
If LIWORK = -1, then a workspace query is assumed; the
routine only calculates the optimal size of the IWORK array,
returns this value as the first entry of the IWORK array, and
no error message related to LIWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: Internal error
Further Details
---------------
Based on contributions by
Inderjit Dhillon, IBM Almaden, USA
Osni Marques, LBNL/NERSC, USA
Ken Stanley, Computer Science Division, University of
California at Berkeley, USA
@ingroup magma_zheev_driver
********************************************************************/
extern "C" magma_int_t
magma_zheevr(
magma_vec_t jobz, magma_range_t range, magma_uplo_t uplo, magma_int_t n,
magmaDoubleComplex *A, magma_int_t lda,
double vl, double vu,
magma_int_t il, magma_int_t iu, double abstol, magma_int_t *m,
double *w,
magmaDoubleComplex *Z, magma_int_t ldz,
magma_int_t *isuppz,
magmaDoubleComplex *work, magma_int_t lwork,
double *rwork, magma_int_t lrwork,
magma_int_t *iwork, magma_int_t liwork,
magma_int_t *info)
{
/* Constants */
const magma_int_t izero = 0;
const magma_int_t ione = 1;
const float szero = 0.;
const float sone = 1.;
/* Local variables */
const char* uplo_ = lapack_uplo_const( uplo );
const char* jobz_ = lapack_vec_const( jobz );
const char* range_ = lapack_range_const( range );
magma_int_t indrd, indre;
magma_int_t imax;
magma_int_t lopt, itmp1, indree, indrdd;
magma_int_t tryrac;
magma_int_t i, j, jj, i__1;
magma_int_t iscale, indibl, indifl;
magma_int_t indiwo, indisp, indtau;
magma_int_t indrwk, indwk;
magma_int_t llwork, llrwork, nsplit;
magma_int_t ieeeok;
magma_int_t iinfo;
magma_int_t lwmin, lrwmin, liwmin;
double safmin;
double bignum;
double smlnum;
double eps, tmp1;
double anrm;
double sigma, d__1;
double rmin, rmax;
bool lower = (uplo == MagmaLower);
bool wantz = (jobz == MagmaVec);
bool alleig = (range == MagmaRangeAll);
bool valeig = (range == MagmaRangeV);
bool indeig = (range == MagmaRangeI);
bool lquery = (lwork == -1 || lrwork == -1 || liwork == -1);
*info = 0;
if (! (wantz || (jobz == MagmaNoVec))) {
*info = -1;
} else if (! (alleig || valeig || indeig)) {
*info = -2;
} else if (! (lower || (uplo == MagmaUpper))) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (lda < max(1,n)) {
*info = -6;
} else if (ldz < 1 || (wantz && ldz < n)) {
*info = -15;
} else {
if (valeig) {
if (n > 0 && vu <= vl) {
*info = -8;
}
} else if (indeig) {
if (il < 1 || il > max(1,n)) {
*info = -9;
} else if (iu < min(n,il) || iu > n) {
*info = -10;
}
}
}
magma_int_t nb = magma_get_zhetrd_nb(n);
lwmin = n * (nb + 1);
lrwmin = 24 * n;
liwmin = 10 * n;
work[0] = magma_zmake_lwork( lwmin );
rwork[0] = magma_dmake_lwork( lrwmin );
iwork[0] = liwmin;
if (lwork < lwmin && ! lquery) {
*info = -18;
} else if ((lrwork < lrwmin) && ! lquery) {
*info = -20;
} else if ((liwork < liwmin) && ! lquery) {
*info = -22;
}
if (*info != 0) {
magma_xerbla(__func__, -(*info));
return *info;
} else if (lquery) {
return *info;
}
*m = 0;
/* Check if matrix is very small then just call LAPACK on CPU, no need for GPU */
if (n <= 128) {
#ifdef ENABLE_DEBUG
printf("--------------------------------------------------------------\n");
printf(" warning matrix too small N=%d NB=%d, calling lapack on CPU \n", (int) n, (int) nb);
printf("--------------------------------------------------------------\n");
#endif
lapackf77_zheevr(jobz_, range_, uplo_,
&n, A, &lda, &vl, &vu, &il, &iu, &abstol, m,
w, Z, &ldz, isuppz, work, &lwork,
rwork, &lrwork, iwork, &liwork, info);
return *info;
}
--w;
--work;
--rwork;
--iwork;
--isuppz;
/* Get machine constants. */
safmin = lapackf77_dlamch("Safe minimum");
eps = lapackf77_dlamch("Precision");
smlnum = safmin / eps;
bignum = 1. / smlnum;
rmin = magma_dsqrt(smlnum);
rmax = magma_dsqrt(bignum);
/* Scale matrix to allowable range, if necessary. */
anrm = lapackf77_zlanhe("M", uplo_, &n, A, &lda, &rwork[1]);
iscale = 0;
if (anrm > 0. && anrm < rmin) {
iscale = 1;
sigma = rmin / anrm;
} else if (anrm > rmax) {
iscale = 1;
sigma = rmax / anrm;
}
if (iscale == 1) {
d__1 = 1.;
lapackf77_zlascl(uplo_, &izero, &izero, &d__1, &sigma, &n, &n, A,
&lda, info);
if (abstol > 0.) {
abstol *= sigma;
}
if (valeig) {
vl *= sigma;
vu *= sigma;
}
}
/* Call ZHETRD to reduce Hermitian matrix to tridiagonal form. */
indtau = 1;
indwk = indtau + n;
indre = 1;
indrd = indre + n;
indree = indrd + n;
indrdd = indree + n;
indrwk = indrdd + n;
llwork = lwork - indwk + 1;
llrwork = lrwork - indrwk + 1;
indifl = 1;
indibl = indifl + n;
indisp = indibl + n;
indiwo = indisp + n;
magma_zhetrd(uplo, n, A, lda, &rwork[indrd], &rwork[indre], &work[indtau], &work[indwk], llwork, &iinfo);
lopt = n + (magma_int_t)MAGMA_Z_REAL(work[indwk]);
/* If all eigenvalues are desired and ABSTOL is less than or equal to
zero, then call DSTERF
or ZUNGTR and ZSTEQR. If this fails for
some eigenvalue, then try DSTEBZ. */
ieeeok = lapackf77_ieeeck( &ione, &szero, &sone);
/* If only the eigenvalues are required call DSTERF for all or DSTEBZ for a part */
if (! wantz) {
blasf77_dcopy(&n, &rwork[indrd], &ione, &w[1], &ione);
i__1 = n - 1;
if (alleig || (indeig && il == 1 && iu == n)) {
lapackf77_dsterf(&n, &w[1], &rwork[indre], info);
*m = n;
} else {
lapackf77_dstebz(range_, "E", &n, &vl, &vu, &il, &iu, &abstol,
&rwork[indrd], &rwork[indre], m,
&nsplit, &w[1], &iwork[indibl], &iwork[indisp],
&rwork[indrwk], &iwork[indiwo], info);
}
/* Otherwise call ZSTEMR if infinite and NaN arithmetic is supported */
}
else if (ieeeok == 1) {
i__1 = n - 1;
blasf77_dcopy(&i__1, &rwork[indre], &ione, &rwork[indree], &ione);
blasf77_dcopy(&n, &rwork[indrd], &ione, &rwork[indrdd], &ione);
if (abstol < 2*n*eps)
tryrac = 1;
else
tryrac = 0;
lapackf77_zstemr(jobz_, range_, &n, &rwork[indrdd], &rwork[indree], &vl, &vu, &il,
&iu, m, &w[1], Z, &ldz, &n, &isuppz[1], &tryrac, &rwork[indrwk],
&llrwork, &iwork[1], &liwork, info);
if (*info == 0 && wantz) {
magma_zunmtr(MagmaLeft, uplo, MagmaNoTrans, n, *m, A, lda, &work[indtau],
Z, ldz, &work[indwk], llwork, &iinfo);
}
}
/* Call DSTEBZ and ZSTEIN if infinite and NaN arithmetic is not supported or ZSTEMR didn't converge. */
if (wantz && (ieeeok == 0 || *info != 0)) {
*info = 0;
lapackf77_dstebz(range_, "B", &n, &vl, &vu, &il, &iu, &abstol, &rwork[indrd], &rwork[indre], m,
&nsplit, &w[1], &iwork[indibl], &iwork[indisp], &rwork[indrwk], &iwork[indiwo], info);
lapackf77_zstein(&n, &rwork[indrd], &rwork[indre], m, &w[1], &iwork[indibl], &iwork[indisp],
Z, &ldz, &rwork[indrwk], &iwork[indiwo], &iwork[indifl], info);
/* Apply unitary matrix used in reduction to tridiagonal
form to eigenvectors returned by ZSTEIN. */
magma_zunmtr(MagmaLeft, uplo, MagmaNoTrans, n, *m, A, lda, &work[indtau],
Z, ldz, &work[indwk], llwork, &iinfo);
}
/* If matrix was scaled, then rescale eigenvalues appropriately. */
if (iscale == 1) {
if (*info == 0) {
imax = *m;
} else {
imax = *info - 1;
}
d__1 = 1. / sigma;
blasf77_dscal(&imax, &d__1, &w[1], &ione);
}
/* If eigenvalues are not in order, then sort them, along with
eigenvectors. */
if (wantz) {
for (j = 1; j <= *m-1; ++j) {
i = 0;
tmp1 = w[j];
for (jj = j + 1; jj <= *m; ++jj) {
if (w[jj] < tmp1) {
i = jj;
tmp1 = w[jj];
}
}
if (i != 0) {
itmp1 = iwork[indibl + i - 1];
w[i] = w[j];
iwork[indibl + i - 1] = iwork[indibl + j - 1];
w[j] = tmp1;
iwork[indibl + j - 1] = itmp1;
blasf77_zswap(&n, Z + (i-1)*ldz, &ione, Z + (j-1)*ldz, &ione);
}
}
}
/* Set WORK[0] to optimal complex workspace size. */
work[1] = magma_zmake_lwork( lopt );
rwork[1] = magma_dmake_lwork( lrwmin );
iwork[1] = liwmin;
return *info;
} /* magma_zheevr */
extern "C" magma_int_t
magma_dgegqr_gpu( magma_int_t m, magma_int_t n,
double *dA, magma_int_t ldda,
double *dwork, double *work,
magma_int_t *info )
{
/* -- MAGMA (version 1.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2011
Purpose
=======
ZGEGQR orthogonalizes the N vectors given by a real M-by-N matrix A:
A = Q * R.
On exit, if successful, the orthogonal vectors Q overwrite A
and R is given in work (on the CPU memory).
This version uses normal equations and SVD in an iterative process that
makes the computation numerically accurate.
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. N >= 0.
dA (input/output) DOUBLE_PRECISION array on the GPU, dimension (LDDA,N)
On entry, the M-by-N matrix A.
On exit, the M-by-N matrix Q with orthogonal columns.
LDDA (input) INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
To benefit from coalescent memory accesses LDDA must be
dividable by 16.
dwork (GPU workspace) DOUBLE_PRECISION array, dimension (N,N)
work (CPU workspace/output) DOUBLE_PRECISION array, dimension 3n^2.
On exit, work(1:n^2) holds the rectangular matrix R.
Preferably, for higher performance, work must be in pinned memory.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
Further Details
===============
===================================================================== */
magma_int_t i = 0, j, k, n2 = n*n, ione = 1;
double zero = MAGMA_D_ZERO, one = MAGMA_D_ONE;
double cn = 200., mins, maxs;
/* check arguments */
*info = 0;
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (ldda < max(1,m)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
double *U, *VT, *vt, *R, *G, *hwork, *tau;
double *S;
R = work; // Size n * n
G = R + n*n; // Size n * n
VT = G + n*n; // Size n * n
magma_dmalloc_cpu( &hwork, 2*n*n + 2*n);
if ( hwork == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_int_t lwork = n*n; // First part f hwork; used as workspace in svd
U = hwork + n*n; // Size n*n
S = (double *)(U+n*n);// Size n
tau = U + n*n + n ; // Size n
do {
i++;
magma_dgemm(MagmaTrans, MagmaNoTrans, ??, ??, ??, one, dA, ldda, dA, ldda, zero, dwork, n );
magma_dgetmatrix(??, ??, dwork, n, G, n);
lapackf77_dgesvd("n", "a", &??, &??, G, &n, S, U, &n, VT, &n,
hwork, &lwork, info);
mins = 100.f, maxs = 0.f;
for(k=0; k<n; k++){
S[k] = magma_dsqrt( S[k] );
if (S[k] < mins) mins = S[k];
if (S[k] > maxs) maxs = S[k];
}
for(k=0; k<n;k++){
vt = VT + k*n;
for(j=0; j<n; j++)
vt[j]*=S[j];
}
lapackf77_dgeqrf(&??, &??, ??, &n, tau, hwork, &lwork, info);
if (i==1)
blasf77_dcopy(&n2, VT, &ione, R, &ione);
else
blasf77_dtrmm("l", "u", "n", "n", &n, &n, &one, VT, &n, R, &n);
magma_dsetmatrix(n, n, VT, n, G, n);
magma_dtrsm('r', 'u', 'n', 'n', ??, ??, one, ??, n, ??, ldda);
if (mins > 0.00001f)
cn = maxs/mins;
//fprintf(stderr, "Iteration %d, cond num = %f \n", i, cn);
} while (cn > 10.f);
magma_free_cpu( hwork );
return *info;
} /* magma_dgegqr_gpu */
extern "C" magma_int_t
magma_dlahr2(magma_int_t n, magma_int_t k, magma_int_t nb,
double *da, double *dv,
double *a, magma_int_t lda,
double *tau, double *t, magma_int_t ldt,
double *y, magma_int_t ldy)
{
/* -- MAGMA auxiliary routine (version 1.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
Purpose
=======
DLAHR2 reduces the first NB columns of a real general n-BY-(n-k+1)
matrix A so that elements below the k-th subdiagonal are zero. The
reduction is performed by an orthogonal similarity transformation
Q' * A * Q. The routine returns the matrices V and T which determine
Q as a block reflector I - V*T*V', and also the matrix Y = A * V.
This is an auxiliary routine called by DGEHRD.
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.
DA (input/output) DOUBLE_PRECISION array on the GPU, 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.
DV (output) DOUBLE_PRECISION array on the GPU, dimension (N, NB)
On exit this contains the Householder vectors of the transformation.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
TAU (output) DOUBLE_PRECISION array, dimension (NB)
The scalar factors of the elementary reflectors. See Further
Details.
T (output) DOUBLE_PRECISION array, dimension (LDT,NB)
The upper triangular matrix T.
LDT (input) INTEGER
The leading dimension of the array T. LDT >= NB.
Y (output) DOUBLE_PRECISION array, dimension (LDY,NB)
The n-by-nb matrix Y.
LDY (input) INTEGER
The leading dimension of the array Y. LDY >= N.
Further Details
===============
The matrix Q is represented as a product of nb elementary reflectors
Q = H(1) H(2) . . . H(nb).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i+k-1) = 0, v(i+k) = 1; v(i+k+1:n) is stored on exit in
A(i+k+1:n,i), and tau in TAU(i).
The elements of the vectors v together form the (n-k+1)-by-nb matrix
V which is needed, with T and Y, to apply the transformation to the
unreduced part of the matrix, using an update of the form:
A := (I - V*T*V') * (A - Y*T*V').
The contents of A on exit are illustrated by the following example
with n = 7, k = 3 and nb = 2:
( 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.
===================================================================== */
double c_zero = MAGMA_D_ZERO;
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
magma_int_t ldda = lda;
magma_int_t c__1 = 1;
magma_int_t a_dim1, a_offset, t_dim1, t_offset, y_dim1, y_offset, i__2, i__3;
double d__1;
magma_int_t i__;
double ei;
--tau;
a_dim1 = lda;
a_offset = 1 + a_dim1;
a -= a_offset;
t_dim1 = ldt;
t_offset = 1 + t_dim1;
t -= t_offset;
y_dim1 = ldy;
y_offset = 1 + y_dim1;
y -= y_offset;
/* Function Body */
if (n <= 1)
return 0;
for (i__ = 1; i__ <= nb; ++i__) {
if (i__ > 1) {
/* Update A(K+1:N,I); Update I-th column of A - Y * V' */
i__2 = n - k + 1;
i__3 = i__ - 1;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv(&i__3, &a[k+i__-1+a_dim1], &lda);
#endif
blasf77_dcopy(&i__3, &a[k+i__-1+a_dim1], &lda, &t[nb*t_dim1+1], &c__1);
blasf77_dtrmv("u","n","n",&i__3,&t[t_offset], &ldt, &t[nb*t_dim1+1], &c__1);
blasf77_dgemv("NO TRANSPOSE", &i__2, &i__3, &c_neg_one, &y[k + y_dim1],
&ldy, &t[nb*t_dim1+1], &c__1, &c_one, &a[k+i__*a_dim1],&c__1);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv(&i__3, &a[k+i__-1+a_dim1], &lda);
#endif
/* Apply I - V * T' * V' to this column (call it b) from the
left, using the last column of T as workspace
Let V = ( V1 ) and b = ( b1 ) (first I-1 rows)
( V2 ) ( b2 )
where V1 is unit lower triangular
w := V1' * b1 */
i__2 = i__ - 1;
blasf77_dcopy(&i__2, &a[k+1+i__*a_dim1], &c__1, &t[nb*t_dim1+1], &c__1);
blasf77_dtrmv("Lower", MagmaTransStr, "UNIT", &i__2,
&a[k + 1 + a_dim1], &lda, &t[nb * t_dim1 + 1], &c__1);
/* w := w + V2'*b2 */
i__2 = n - k - i__ + 1;
i__3 = i__ - 1;
blasf77_dgemv(MagmaTransStr, &i__2, &i__3, &c_one,
&a[k + i__ + a_dim1], &lda, &a[k+i__+i__*a_dim1], &c__1,
&c_one, &t[nb*t_dim1+1], &c__1);
/* w := T'*w */
i__2 = i__ - 1;
blasf77_dtrmv("U", MagmaTransStr, "N", &i__2, &t[t_offset], &ldt,
&t[nb*t_dim1+1], &c__1);
/* b2 := b2 - V2*w */
i__2 = n - k - i__ + 1;
i__3 = i__ - 1;
blasf77_dgemv("N", &i__2, &i__3, &c_neg_one, &a[k + i__ + a_dim1], &lda,
&t[nb*t_dim1+1], &c__1, &c_one, &a[k+i__+i__*a_dim1], &c__1);
/* b1 := b1 - V1*w */
i__2 = i__ - 1;
blasf77_dtrmv("L","N","U",&i__2,&a[k+1+a_dim1],&lda,&t[nb*t_dim1+1],&c__1);
blasf77_daxpy(&i__2, &c_neg_one, &t[nb * t_dim1 + 1], &c__1,
&a[k + 1 + i__ * a_dim1], &c__1);
a[k + i__ - 1 + (i__ - 1) * a_dim1] = ei;
}
/* Generate the elementary reflector H(I) to annihilate A(K+I+1:N,I) */
i__2 = n - k - i__ + 1;
i__3 = k + i__ + 1;
lapackf77_dlarfg(&i__2, &a[k + i__ + i__ * a_dim1],
&a[min(i__3,n) + i__ * a_dim1], &c__1, &tau[i__]);
ei = a[k + i__ + i__ * a_dim1];
a[k + i__ + i__ * a_dim1] = c_one;
/* Compute Y(K+1:N,I) */
i__2 = n - k;
i__3 = n - k - i__ + 1;
magma_dsetvector( i__3,
&a[k + i__ + i__*a_dim1], 1,
dv+(i__-1)*(ldda+1), 1 );
magma_dgemv(MagmaNoTrans, i__2+1, i__3, c_one,
da -1 + k + i__ * ldda, ldda,
dv+(i__-1)*(ldda+1), c__1, c_zero,
da-1 + k + (i__-1)*ldda, c__1);
i__2 = n - k - i__ + 1;
i__3 = i__ - 1;
blasf77_dgemv(MagmaTransStr, &i__2, &i__3, &c_one,
&a[k + i__ + a_dim1], &lda, &a[k+i__+i__*a_dim1], &c__1,
&c_zero, &t[i__*t_dim1+1], &c__1);
/* Compute T(1:I,I) */
i__2 = i__ - 1;
d__1 = MAGMA_D_NEGATE( tau[i__] );
blasf77_dscal(&i__2, &d__1, &t[i__ * t_dim1 + 1], &c__1);
blasf77_dtrmv("U","N","N", &i__2, &t[t_offset], &ldt, &t[i__*t_dim1+1], &c__1);
t[i__ + i__ * t_dim1] = tau[i__];
magma_dgetvector( n - k + 1,
da-1+ k+(i__-1)*ldda, 1,
y+ k + i__*y_dim1, 1 );
}
a[k + nb + nb * a_dim1] = ei;
return 0;
} /* magma_dlahr2 */
/**
Purpose
-------
ZHEEVX computes selected eigenvalues and, optionally, eigenvectors
of a complex Hermitian matrix A. Eigenvalues and eigenvectors can
be selected by specifying either a range of values or a range of
indices for the desired eigenvalues.
Arguments
---------
@param[in]
jobz magma_vec_t
- = MagmaNoVec: Compute eigenvalues only;
- = MagmaVec: Compute eigenvalues and eigenvectors.
@param[in]
range magma_range_t
- = MagmaRangeAll: all eigenvalues will be found.
- = MagmaRangeV: all eigenvalues in the half-open interval (VL,VU]
will be found.
- = MagmaRangeI: the IL-th through IU-th eigenvalues will be found.
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored;
- = MagmaLower: Lower triangle of A is stored.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
dA COMPLEX_16 array, dimension (LDDA, N)
On entry, the Hermitian matrix A. If UPLO = MagmaUpper, the
leading N-by-N upper triangular part of A contains the
upper triangular part of the matrix A. If UPLO = MagmaLower,
the leading N-by-N lower triangular part of A contains
the lower triangular part of the matrix A.
On exit, the lower triangle (if UPLO=MagmaLower) or the upper
triangle (if UPLO=MagmaUpper) of A, including the diagonal, is
destroyed.
@param[in]
ldda INTEGER
The leading dimension of the array DA. LDDA >= max(1,N).
@param[in]
vl DOUBLE PRECISION
@param[in]
vu DOUBLE PRECISION
If RANGE=MagmaRangeV, the lower and upper bounds of the interval to
be searched for eigenvalues. VL < VU.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeI.
@param[in]
il INTEGER
@param[in]
iu INTEGER
If RANGE=MagmaRangeI, the indices (in ascending order) of the
smallest and largest eigenvalues to be returned.
1 <= IL <= IU <= N, if N > 0; IL = 1 and IU = 0 if N = 0.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeV.
@param[in]
abstol DOUBLE PRECISION
The absolute error tolerance for the eigenvalues.
An approximate eigenvalue is accepted as converged
when it is determined to lie in an interval [a,b]
of width less than or equal to
ABSTOL + EPS * max( |a|,|b| ),
\n
where EPS is the machine precision. If ABSTOL is less than
or equal to zero, then EPS*|T| will be used in its place,
where |T| is the 1-norm of the tridiagonal matrix obtained
by reducing A to tridiagonal form.
\n
Eigenvalues will be computed most accurately when ABSTOL is
set to twice the underflow threshold 2*DLAMCH('S'), not zero.
If this routine returns with INFO > 0, indicating that some
eigenvectors did not converge, try setting ABSTOL to
2*DLAMCH('S').
\n
See "Computing Small Singular Values of Bidiagonal Matrices
with Guaranteed High Relative Accuracy," by Demmel and
Kahan, LAPACK Working Note #3.
@param[out]
m INTEGER
The total number of eigenvalues found. 0 <= M <= N.
If RANGE = MagmaRangeAll, M = N, and if RANGE = MagmaRangeI, M = IU-IL+1.
@param[out]
w DOUBLE PRECISION array, dimension (N)
On normal exit, the first M elements contain the selected
eigenvalues in ascending order.
@param[out]
dZ COMPLEX_16 array, dimension (LDDZ, max(1,M))
If JOBZ = MagmaVec, then if INFO = 0, the first M columns of Z
contain the orthonormal eigenvectors of the matrix A
corresponding to the selected eigenvalues, with the i-th
column of Z holding the eigenvector associated with W(i).
If an eigenvector fails to converge, then that column of Z
contains the latest approximation to the eigenvector, and the
index of the eigenvector is returned in IFAIL.
If JOBZ = MagmaNoVec, then Z is not referenced.
Note: the user must ensure that at least max(1,M) columns are
supplied in the array Z; if RANGE = MagmaRangeV, the exact value of M
is not known in advance and an upper bound must be used.
********* (workspace) If FAST_HEMV is defined DZ should be (LDDZ, max(1,N)) in both cases.
@param[in]
lddz INTEGER
The leading dimension of the array DZ. LDDZ >= 1, and if
JOBZ = MagmaVec, LDDZ >= max(1,N).
@param
wA (workspace) COMPLEX_16 array, dimension (LDWA, N)
@param[in]
ldwa INTEGER
The leading dimension of the array wA. LDWA >= max(1,N).
@param
wZ (workspace) COMPLEX_16 array, dimension (LDWZ, max(1,M))
@param[in]
ldwz INTEGER
The leading dimension of the array wZ. LDWZ >= 1, and if
JOBZ = MagmaVec, LDWZ >= max(1,N).
@param[out]
work (workspace) COMPLEX_16 array, dimension (LWORK)
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The length of the array WORK. LWORK >= (NB+1)*N,
where NB is the max of the blocksize for ZHETRD.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param
rwork (workspace) DOUBLE PRECISION array, dimension (7*N)
@param
iwork (workspace) INTEGER array, dimension (5*N)
@param[out]
ifail INTEGER array, dimension (N)
If JOBZ = MagmaVec, then if INFO = 0, the first M elements of
IFAIL are zero. If INFO > 0, then IFAIL contains the
indices of the eigenvectors that failed to converge.
If JOBZ = MagmaNoVec, then IFAIL is not referenced.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: if INFO = i, then i eigenvectors failed to converge.
Their indices are stored in array IFAIL.
@ingroup magma_zheev_driver
********************************************************************/
extern "C" magma_int_t
magma_zheevx_gpu(magma_vec_t jobz, magma_range_t range, magma_uplo_t uplo, magma_int_t n,
magmaDoubleComplex *dA, magma_int_t ldda, double vl, double vu,
magma_int_t il, magma_int_t iu, double abstol, magma_int_t *m,
double *w, magmaDoubleComplex *dZ, magma_int_t lddz,
magmaDoubleComplex *wA, magma_int_t ldwa,
magmaDoubleComplex *wZ, magma_int_t ldwz,
magmaDoubleComplex *work, magma_int_t lwork,
double *rwork, magma_int_t *iwork, magma_int_t *ifail, magma_int_t *info)
{
const char* uplo_ = lapack_uplo_const( uplo );
const char* jobz_ = lapack_vec_const( jobz );
const char* range_ = lapack_range_const( range );
magma_int_t ione = 1;
const char* order_;
magma_int_t indd, inde;
magma_int_t imax;
magma_int_t lopt, itmp1, indee;
magma_int_t lower, wantz;
magma_int_t i, j, jj, i__1;
magma_int_t alleig, valeig, indeig;
magma_int_t iscale, indibl;
magma_int_t indiwk, indisp, indtau;
magma_int_t indrwk, indwrk;
magma_int_t llwork, nsplit;
magma_int_t lquery;
magma_int_t iinfo;
double safmin;
double bignum;
double smlnum;
double eps, tmp1;
double anrm;
double sigma, d__1;
double rmin, rmax;
double *dwork;
/* Function Body */
lower = (uplo == MagmaLower);
wantz = (jobz == MagmaVec);
alleig = (range == MagmaRangeAll);
valeig = (range == MagmaRangeV);
indeig = (range == MagmaRangeI);
lquery = (lwork == -1);
*info = 0;
if (! (wantz || (jobz == MagmaNoVec))) {
*info = -1;
} else if (! (alleig || valeig || indeig)) {
*info = -2;
} else if (! (lower || (uplo == MagmaUpper))) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (ldda < max(1,n)) {
*info = -6;
} else if (lddz < 1 || (wantz && lddz < n)) {
*info = -15;
} else if (ldwa < max(1,n)) {
*info = -17;
} else if (ldwz < 1 || (wantz && ldwz < n)) {
*info = -19;
} else {
if (valeig) {
if (n > 0 && vu <= vl) {
*info = -8;
}
} else if (indeig) {
if (il < 1 || il > max(1,n)) {
*info = -9;
} else if (iu < min(n,il) || iu > n) {
*info = -10;
}
}
}
magma_int_t nb = magma_get_zhetrd_nb(n);
lopt = n * (nb + 1);
work[0] = MAGMA_Z_MAKE( lopt, 0 );
if (lwork < lopt && ! lquery) {
*info = -21;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info));
return *info;
} else if (lquery) {
return *info;
}
*m = 0;
/* Check if matrix is very small then just call LAPACK on CPU, no need for GPU */
if (n <= 128) {
#ifdef ENABLE_DEBUG
printf("--------------------------------------------------------------\n");
printf(" warning matrix too small N=%d NB=%d, calling lapack on CPU \n", (int) n, (int) nb);
printf("--------------------------------------------------------------\n");
#endif
magmaDoubleComplex *a;
magma_zmalloc_cpu( &a, n*n );
magma_zgetmatrix(n, n, dA, ldda, a, n);
lapackf77_zheevx(jobz_, range_, uplo_,
&n, a, &n, &vl, &vu, &il, &iu, &abstol, m,
w, wZ, &ldwz, work, &lwork,
rwork, iwork, ifail, info);
magma_zsetmatrix( n, n, a, n, dA, ldda);
magma_zsetmatrix( n, *m, wZ, ldwz, dZ, lddz);
magma_free_cpu(a);
return *info;
}
if (MAGMA_SUCCESS != magma_dmalloc( &dwork, n )) {
fprintf (stderr, "!!!! device memory allocation error (magma_zheevx_gpu)\n");
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
--w;
--work;
--rwork;
--iwork;
--ifail;
/* Get machine constants. */
safmin = lapackf77_dlamch("Safe minimum");
eps = lapackf77_dlamch("Precision");
smlnum = safmin / eps;
bignum = 1. / smlnum;
rmin = magma_dsqrt(smlnum);
rmax = magma_dsqrt(bignum);
/* Scale matrix to allowable range, if necessary. */
anrm = magmablas_zlanhe(MagmaMaxNorm, uplo, n, dA, ldda, dwork);
iscale = 0;
sigma = 1;
if (anrm > 0. && anrm < rmin) {
iscale = 1;
sigma = rmin / anrm;
} else if (anrm > rmax) {
iscale = 1;
sigma = rmax / anrm;
}
if (iscale == 1) {
d__1 = 1.;
magmablas_zlascl(uplo, 0, 0, 1., sigma, n, n, dA,
ldda, info);
if (abstol > 0.) {
abstol *= sigma;
}
if (valeig) {
vl *= sigma;
vu *= sigma;
}
}
/* Call ZHETRD to reduce Hermitian matrix to tridiagonal form. */
indd = 1;
inde = indd + n;
indrwk = inde + n;
indtau = 1;
indwrk = indtau + n;
llwork = lwork - indwrk + 1;
#ifdef FAST_HEMV
magma_zhetrd2_gpu(uplo, n, dA, ldda, &rwork[indd], &rwork[inde],
&work[indtau], wA, ldwa, &work[indwrk], llwork, dZ, lddz*n, &iinfo);
#else
magma_zhetrd_gpu (uplo, n, dA, ldda, &rwork[indd], &rwork[inde],
&work[indtau], wA, ldwa, &work[indwrk], llwork, &iinfo);
#endif
lopt = n + (magma_int_t)MAGMA_Z_REAL(work[indwrk]);
/* If all eigenvalues are desired and ABSTOL is less than or equal to
zero, then call DSTERF or ZUNGTR and ZSTEQR. If this fails for
some eigenvalue, then try DSTEBZ. */
if ((alleig || (indeig && il == 1 && iu == n)) && abstol <= 0.) {
blasf77_dcopy(&n, &rwork[indd], &ione, &w[1], &ione);
indee = indrwk + 2*n;
if (! wantz) {
i__1 = n - 1;
blasf77_dcopy(&i__1, &rwork[inde], &ione, &rwork[indee], &ione);
lapackf77_dsterf(&n, &w[1], &rwork[indee], info);
}
else {
lapackf77_zlacpy("A", &n, &n, wA, &ldwa, wZ, &ldwz);
lapackf77_zungtr(uplo_, &n, wZ, &ldwz, &work[indtau], &work[indwrk], &llwork, &iinfo);
i__1 = n - 1;
blasf77_dcopy(&i__1, &rwork[inde], &ione, &rwork[indee], &ione);
lapackf77_zsteqr(jobz_, &n, &w[1], &rwork[indee], wZ, &ldwz, &rwork[indrwk], info);
if (*info == 0) {
for (i = 1; i <= n; ++i) {
ifail[i] = 0;
}
magma_zsetmatrix( n, n, wZ, ldwz, dZ, lddz );
}
}
if (*info == 0) {
*m = n;
}
}
/* Otherwise, call DSTEBZ and, if eigenvectors are desired, ZSTEIN. */
if (*m == 0) {
*info = 0;
if (wantz) {
order_ = "B";
} else {
order_ = "E";
}
indibl = 1;
indisp = indibl + n;
indiwk = indisp + n;
lapackf77_dstebz(range_, order_, &n, &vl, &vu, &il, &iu, &abstol, &rwork[indd], &rwork[inde], m,
&nsplit, &w[1], &iwork[indibl], &iwork[indisp], &rwork[indrwk], &iwork[indiwk], info);
if (wantz) {
lapackf77_zstein(&n, &rwork[indd], &rwork[inde], m, &w[1], &iwork[indibl], &iwork[indisp],
wZ, &ldwz, &rwork[indrwk], &iwork[indiwk], &ifail[1], info);
magma_zsetmatrix( n, *m, wZ, ldwz, dZ, lddz );
/* Apply unitary matrix used in reduction to tridiagonal
form to eigenvectors returned by ZSTEIN. */
magma_zunmtr_gpu(MagmaLeft, uplo, MagmaNoTrans, n, *m, dA, ldda, &work[indtau],
dZ, lddz, wA, ldwa, &iinfo);
}
}
/* If matrix was scaled, then rescale eigenvalues appropriately. */
if (iscale == 1) {
if (*info == 0) {
imax = *m;
} else {
imax = *info - 1;
}
d__1 = 1. / sigma;
blasf77_dscal(&imax, &d__1, &w[1], &ione);
}
/* If eigenvalues are not in order, then sort them, along with
eigenvectors. */
if (wantz) {
for (j = 1; j <= *m-1; ++j) {
i = 0;
tmp1 = w[j];
for (jj = j + 1; jj <= *m; ++jj) {
if (w[jj] < tmp1) {
i = jj;
tmp1 = w[jj];
}
}
if (i != 0) {
itmp1 = iwork[indibl + i - 1];
w[i] = w[j];
iwork[indibl + i - 1] = iwork[indibl + j - 1];
w[j] = tmp1;
iwork[indibl + j - 1] = itmp1;
magma_zswap(n, dZ + (i-1)*lddz, ione, dZ + (j-1)*lddz, ione);
if (*info != 0) {
itmp1 = ifail[i];
ifail[i] = ifail[j];
ifail[j] = itmp1;
}
}
}
}
/* Set WORK[0] to optimal complex workspace size. */
work[1] = MAGMA_Z_MAKE( lopt, 0 );
return *info;
} /* magma_zheevx_gpu */
extern "C" magma_int_t
magma_dlaex3_m(magma_int_t nrgpu,
magma_int_t k, magma_int_t n, magma_int_t n1, double* d,
double* q, magma_int_t ldq, double rho,
double* dlamda, double* q2, magma_int_t* indx,
magma_int_t* ctot, double* w, double* s, magma_int_t* indxq,
double** dwork, magma_queue_t stream[MagmaMaxGPUs][2],
char range, double vl, double vu, magma_int_t il, magma_int_t iu,
magma_int_t* info )
{
/*
Purpose
=======
DLAEX3 finds the roots of the secular equation, as defined by the
values in D, W, and RHO, between 1 and K. It makes the
appropriate calls to DLAED4 and then updates the eigenvectors by
multiplying the matrix of eigenvectors of the pair of eigensystems
being combined by the matrix of eigenvectors of the K-by-K system
which is solved here.
It is used in the last step when only a part of the eigenvectors
is required.
It compute only the required part of the eigenvectors and the rest
is not used.
This code makes very mild assumptions about floating point
arithmetic. It will work on machines with a guard digit in
add/subtract, or on those binary machines without guard digits
which subtract like the Cray X-MP, Cray Y-MP, Cray C-90, or Cray-2.
It could conceivably fail on hexadecimal or decimal machines
without guard digits, but we know of none.
Arguments
=========
K (input) INTEGER
The number of terms in the rational function to be solved by
DLAED4. K >= 0.
N (input) INTEGER
The number of rows and columns in the Q matrix.
N >= K (deflation may result in N>K).
N1 (input) INTEGER
The location of the last eigenvalue in the leading submatrix.
min(1,N) <= N1 <= N/2.
D (output) DOUBLE PRECISION array, dimension (N)
D(I) contains the updated eigenvalues for
1 <= I <= K.
Q (output) DOUBLE PRECISION array, dimension (LDQ,N)
Initially the first K columns are used as workspace.
On output the columns ??? to ??? contain
the updated eigenvectors.
LDQ (input) INTEGER
The leading dimension of the array Q. LDQ >= max(1,N).
RHO (input) DOUBLE PRECISION
The value of the parameter in the rank one update equation.
RHO >= 0 required.
DLAMDA (input/output) DOUBLE PRECISION array, dimension (K)
The first K elements of this array contain the old roots
of the deflated updating problem. These are the poles
of the secular equation. May be changed on output by
having lowest order bit set to zero on Cray X-MP, Cray Y-MP,
Cray-2, or Cray C-90, as described above.
Q2 (input) DOUBLE PRECISION array, dimension (LDQ2, N)
The first K columns of this matrix contain the non-deflated
eigenvectors for the split problem.
INDX (input) INTEGER array, dimension (N)
The permutation used to arrange the columns of the deflated
Q matrix into three groups (see DLAED2).
The rows of the eigenvectors found by DLAED4 must be likewise
permuted before the matrix multiply can take place.
CTOT (input) INTEGER array, dimension (4)
A count of the total number of the various types of columns
in Q, as described in INDX. The fourth column type is any
column which has been deflated.
W (input/output) DOUBLE PRECISION array, dimension (K)
The first K elements of this array contain the components
of the deflation-adjusted updating vector. Destroyed on
output.
S (workspace) DOUBLE PRECISION array, dimension (N1 + 1)*K
Will contain the eigenvectors of the repaired matrix which
will be multiplied by the previously accumulated eigenvectors
to update the system.
INDXQ (output) INTEGER array, dimension (N)
On exit, the permutation which will reintegrate the
subproblems back into sorted order,
i.e. D( INDXQ( I = 1, N ) ) will be in ascending order.
DWORK (devices workspaces) DOUBLE PRECISION array of arrays,
dimension NRGPU.
if NRGPU = 1 the dimension of the first workspace
should be (3*N*N/2+3*N)
otherwise the NRGPU workspaces should have the size
ceil((N-N1) * (N-N1) / floor(nrgpu/2)) +
NB * ((N-N1) + (N-N1) / floor(nrgpu/2))
STREAM (device stream) magma_queue_t array,
dimension (MagmaMaxGPUs,2)
INFO (output) INTEGER
= 0: successful exit.
< 0: if INFO = -i, the i-th argument had an illegal value.
> 0: if INFO = 1, an eigenvalue did not converge
Further Details
===============
Based on contributions by
Jeff Rutter, Computer Science Division, University of California
at Berkeley, USA
Modified by Francoise Tisseur, University of Tennessee.
===================================================================== */
if (nrgpu==1){
magma_setdevice(0);
magma_dlaex3(k, n, n1, d, q, ldq, rho,
dlamda, q2, indx, ctot, w, s, indxq,
*dwork, range, vl, vu, il, iu, info );
return MAGMA_SUCCESS;
}
double d_one = 1.;
double d_zero = 0.;
magma_int_t ione = 1;
magma_int_t ineg_one = -1;
char range_[] = {range, 0};
magma_int_t iil, iiu, rk;
magma_int_t n1_loc, n2_loc, ib, nb, ib2, igpu;
magma_int_t ni_loc[MagmaMaxGPUs];
magma_int_t i,ind,iq2,j,n12,n2,n23,tmp,lq2;
double temp;
magma_int_t alleig, valeig, indeig;
alleig = lapackf77_lsame(range_, "A");
valeig = lapackf77_lsame(range_, "V");
indeig = lapackf77_lsame(range_, "I");
*info = 0;
if(k < 0)
*info=-1;
else if(n < k)
*info=-2;
else if(ldq < max(1,n))
*info=-6;
else if (! (alleig || valeig || indeig))
*info = -15;
else {
if (valeig) {
if (n > 0 && vu <= vl)
*info = -17;
}
else if (indeig) {
if (il < 1 || il > max(1,n))
*info = -18;
else if (iu < min(n,il) || iu > n)
*info = -19;
}
}
if(*info != 0){
magma_xerbla(__func__, -(*info));
return MAGMA_ERR_ILLEGAL_VALUE;
}
// Quick return if possible
if(k == 0)
return MAGMA_SUCCESS;
/*
Modify values DLAMDA(i) to make sure all DLAMDA(i)-DLAMDA(j) can
be computed with high relative accuracy (barring over/underflow).
This is a problem on machines without a guard digit in
add/subtract (Cray XMP, Cray YMP, Cray C 90 and Cray 2).
The following code replaces DLAMDA(I) by 2*DLAMDA(I)-DLAMDA(I),
which on any of these machines zeros out the bottommost
bit of DLAMDA(I) if it is 1; this makes the subsequent
subtractions DLAMDA(I)-DLAMDA(J) unproblematic when cancellation
occurs. On binary machines with a guard digit (almost all
machines) it does not change DLAMDA(I) at all. On hexadecimal
and decimal machines with a guard digit, it slightly
changes the bottommost bits of DLAMDA(I). It does not account
for hexadecimal or decimal machines without guard digits
(we know of none). We use a subroutine call to compute
2*DLAMBDA(I) to prevent optimizing compilers from eliminating
this code.*/
//#define CHECK_CPU
#ifdef CHECK_CPU
double *hwS[2][MagmaMaxGPUs], *hwQ[2][MagmaMaxGPUs], *hwQ2[MagmaMaxGPUs];
#define hQ2(id) (hwQ2[id])
#define hS(id, ii) (hwS[ii][id])
#define hQ(id, ii) (hwQ[ii][id])
#endif
n2 = n - n1;
n12 = ctot[0] + ctot[1];
n23 = ctot[1] + ctot[2];
iq2 = n1 * n12;
lq2 = iq2 + n2 * n23;
n1_loc = (n1-1) / (nrgpu/2) + 1;
n2_loc = (n2-1) / (nrgpu/2) + 1;
nb = magma_get_dlaex3_m_nb();
if (n1 >= magma_get_dlaex3_m_k()){
#ifdef CHECK_CPU
for (igpu = 0; igpu < nrgpu; ++igpu){
magma_dmalloc_pinned( &(hwS[0][igpu]), n2*nb );
magma_dmalloc_pinned( &(hwS[1][igpu]), n2*nb );
magma_dmalloc_pinned( &(hwQ2[igpu]), n2*n2_loc );
magma_dmalloc_pinned( &(hwQ[0][igpu]), n2_loc*nb );
magma_dmalloc_pinned( &(hwQ[1][igpu]), n2_loc*nb );
}
#endif
for (igpu = 0; igpu < nrgpu-1; igpu += 2){
ni_loc[igpu] = min(n1_loc, n1 - igpu/2 * n1_loc);
#ifdef CHECK_CPU
lapackf77_dlacpy("A", &ni_loc[igpu], &n12, q2+n1_loc*(igpu/2), &n1, hQ2(igpu), &n1_loc);
#endif
magma_setdevice(igpu);
magma_dsetmatrix_async( ni_loc[igpu], n12,
q2+n1_loc*(igpu/2), n1,
dQ2(igpu), n1_loc, stream[igpu][0] );
ni_loc[igpu+1] = min(n2_loc, n2 - igpu/2 * n2_loc);
#ifdef CHECK_CPU
lapackf77_dlacpy("A", &ni_loc[igpu+1], &n23, q2+iq2+n2_loc*(igpu/2), &n2, hQ2(igpu+1), &n2_loc);
#endif
magma_setdevice(igpu+1);
magma_dsetmatrix_async( ni_loc[igpu+1], n23,
q2+iq2+n2_loc*(igpu/2), n2,
dQ2(igpu+1), n2_loc, stream[igpu+1][0] );
}
}
//
#ifdef _OPENMP
/////////////////////////////////////////////////////////////////////////////////
//openmp implementation
/////////////////////////////////////////////////////////////////////////////////
#ifdef ENABLE_TIMER_DIVIDE_AND_CONQUER
magma_timestr_t start, end;
start = get_current_time();
#endif
#pragma omp parallel private(i, j, tmp, temp)
{
magma_int_t id = omp_get_thread_num();
magma_int_t tot = omp_get_num_threads();
magma_int_t ib = ( id * k) / tot; //start index of local loop
magma_int_t ie = ((id+1) * k) / tot; //end index of local loop
magma_int_t ik = ie - ib; //number of local indices
for(i = ib; i < ie; ++i)
dlamda[i]=lapackf77_dlamc3(&dlamda[i], &dlamda[i]) - dlamda[i];
for(j = ib; j < ie; ++j){
magma_int_t tmpp=j+1;
magma_int_t iinfo = 0;
lapackf77_dlaed4(&k, &tmpp, dlamda, w, Q(0,j), &rho, &d[j], &iinfo);
// If the zero finder fails, the computation is terminated.
if(iinfo != 0){
#pragma omp critical (info)
*info=iinfo;
break;
}
}
#pragma omp barrier
if(*info == 0){
#pragma omp single
{
//Prepare the INDXQ sorting permutation.
magma_int_t nk = n - k;
lapackf77_dlamrg( &k, &nk, d, &ione , &ineg_one, indxq);
//compute the lower and upper bound of the non-deflated eigenvectors
if (valeig)
magma_dvrange(k, d, &iil, &iiu, vl, vu);
else if (indeig)
magma_dirange(k, indxq, &iil, &iiu, il, iu);
else {
iil = 1;
iiu = k;
}
rk = iiu - iil + 1;
}
if (k == 2){
#pragma omp single
{
for(j = 0; j < k; ++j){
w[0] = *Q(0,j);
w[1] = *Q(1,j);
i = indx[0] - 1;
*Q(0,j) = w[i];
i = indx[1] - 1;
*Q(1,j) = w[i];
}
}
}
else if(k != 1){
// Compute updated W.
blasf77_dcopy( &ik, &w[ib], &ione, &s[ib], &ione);
// Initialize W(I) = Q(I,I)
tmp = ldq + 1;
blasf77_dcopy( &ik, Q(ib,ib), &tmp, &w[ib], &ione);
for(j = 0; j < k; ++j){
magma_int_t i_tmp = min(j, ie);
for(i = ib; i < i_tmp; ++i)
w[i] = w[i] * ( *Q(i, j) / ( dlamda[i] - dlamda[j] ) );
i_tmp = max(j+1, ib);
for(i = i_tmp; i < ie; ++i)
w[i] = w[i] * ( *Q(i, j) / ( dlamda[i] - dlamda[j] ) );
}
for(i = ib; i < ie; ++i)
w[i] = copysign( sqrt( -w[i] ), s[i]);
#pragma omp barrier
//reduce the number of used threads to have enough S workspace
tot = min(n1, omp_get_num_threads());
if(id < tot){
ib = ( id * rk) / tot + iil - 1;
ie = ((id+1) * rk) / tot + iil - 1;
ik = ie - ib;
}
else{
ib = -1;
ie = -1;
ik = -1;
}
// Compute eigenvectors of the modified rank-1 modification.
for(j = ib; j < ie; ++j){
for(i = 0; i < k; ++i)
s[id*k + i] = w[i] / *Q(i,j);
temp = cblas_dnrm2( k, s+id*k, 1);
for(i = 0; i < k; ++i){
magma_int_t iii = indx[i] - 1;
*Q(i,j) = s[id*k + iii] / temp;
}
}
}
}
}
if (*info != 0)
return MAGMA_SUCCESS; //??????
#ifdef ENABLE_TIMER_DIVIDE_AND_CONQUER
end = get_current_time();
printf("eigenvalues/vector D+zzT = %6.2f\n", GetTimerValue(start,end)/1000.);
#endif
#else
/////////////////////////////////////////////////////////////////////////////////
// Non openmp implementation
/////////////////////////////////////////////////////////////////////////////////
#ifdef ENABLE_TIMER_DIVIDE_AND_CONQUER
magma_timestr_t start, end;
start = get_current_time();
#endif
for(i = 0; i < k; ++i)
dlamda[i]=lapackf77_dlamc3(&dlamda[i], &dlamda[i]) - dlamda[i];
for(j = 0; j < k; ++j){
magma_int_t tmpp=j+1;
magma_int_t iinfo = 0;
lapackf77_dlaed4(&k, &tmpp, dlamda, w, Q(0,j), &rho, &d[j], &iinfo);
// If the zero finder fails, the computation is terminated.
if(iinfo != 0)
*info=iinfo;
}
if(*info != 0)
return MAGMA_SUCCESS;
//Prepare the INDXQ sorting permutation.
magma_int_t nk = n - k;
lapackf77_dlamrg( &k, &nk, d, &ione , &ineg_one, indxq);
//compute the lower and upper bound of the non-deflated eigenvectors
if (valeig)
magma_dvrange(k, d, &iil, &iiu, vl, vu);
else if (indeig)
magma_dirange(k, indxq, &iil, &iiu, il, iu);
else {
iil = 1;
iiu = k;
}
rk = iiu - iil + 1;
if (k == 2){
for(j = 0; j < k; ++j){
w[0] = *Q(0,j);
w[1] = *Q(1,j);
i = indx[0] - 1;
*Q(0,j) = w[i];
i = indx[1] - 1;
*Q(1,j) = w[i];
}
}
else if(k != 1){
// Compute updated W.
blasf77_dcopy( &k, w, &ione, s, &ione);
// Initialize W(I) = Q(I,I)
tmp = ldq + 1;
blasf77_dcopy( &k, q, &tmp, w, &ione);
for(j = 0; j < k; ++j){
for(i = 0; i < j; ++i)
w[i] = w[i] * ( *Q(i, j) / ( dlamda[i] - dlamda[j] ) );
for(i = j+1; i < k; ++i)
w[i] = w[i] * ( *Q(i, j) / ( dlamda[i] - dlamda[j] ) );
}
for(i = 0; i < k; ++i)
w[i] = copysign( sqrt( -w[i] ), s[i]);
// Compute eigenvectors of the modified rank-1 modification.
for(j = iil-1; j < iiu; ++j){
for(i = 0; i < k; ++i)
s[i] = w[i] / *Q(i,j);
temp = cblas_dnrm2( k, s, 1);
for(i = 0; i < k; ++i){
magma_int_t iii = indx[i] - 1;
*Q(i,j) = s[iii] / temp;
}
}
}
#ifdef ENABLE_TIMER_DIVIDE_AND_CONQUER
end = get_current_time();
printf("eigenvalues/vector D+zzT = %6.2f\n", GetTimerValue(start,end)/1000.);
#endif
#endif //_OPENMP
// Compute the updated eigenvectors.
#ifdef ENABLE_TIMER_DIVIDE_AND_CONQUER
start = get_current_time();
#endif
if(rk > 0){
if (n1 < magma_get_dlaex3_m_k()){
// stay on the CPU
if( n23 != 0 ){
lapackf77_dlacpy("A", &n23, &rk, Q(ctot[0],iil-1), &ldq, s, &n23);
blasf77_dgemm("N", "N", &n2, &rk, &n23, &d_one, &q2[iq2], &n2,
s, &n23, &d_zero, Q(n1,iil-1), &ldq );
}
else
lapackf77_dlaset("A", &n2, &rk, &d_zero, &d_zero, Q(n1,iil-1), &ldq);
if( n12 != 0 ) {
lapackf77_dlacpy("A", &n12, &rk, Q(0,iil-1), &ldq, s, &n12);
blasf77_dgemm("N", "N", &n1, &rk, &n12, &d_one, q2, &n1,
s, &n12, &d_zero, Q(0,iil-1), &ldq);
}
else
lapackf77_dlaset("A", &n1, &rk, &d_zero, &d_zero, Q(0,iil-1), &ldq);
}
else {
//use the gpus
ib = min(nb, rk);
for (igpu = 0; igpu < nrgpu-1; igpu += 2){
if (n23 != 0) {
magma_setdevice(igpu+1);
magma_dsetmatrix_async( n23, ib,
Q(ctot[0],iil-1), ldq,
dS(igpu+1,0), n23, stream[igpu+1][0] );
}
if (n12 != 0) {
magma_setdevice(igpu);
magma_dsetmatrix_async( n12, ib,
Q(0,iil-1), ldq,
dS(igpu,0), n12, stream[igpu][0] );
}
}
for (i = 0; i<rk; i+=nb){
ib = min(nb, rk - i);
ind = (i/nb)%2;
if (i+nb<rk){
ib2 = min(nb, rk - i - nb);
for (igpu = 0; igpu < nrgpu-1; igpu += 2){
if (n23 != 0) {
magma_setdevice(igpu+1);
magma_dsetmatrix_async( n23, ib2,
Q(ctot[0],iil-1+i+nb), ldq,
dS(igpu+1,(ind+1)%2), n23, stream[igpu+1][(ind+1)%2] );
}
if (n12 != 0) {
magma_setdevice(igpu);
magma_dsetmatrix_async( n12, ib2,
Q(0,iil-1+i+nb), ldq,
dS(igpu,(ind+1)%2), n12, stream[igpu][(ind+1)%2] );
}
}
}
// Ensure that the data is copied on gpu since we will overwrite it.
for (igpu = 0; igpu < nrgpu-1; igpu += 2){
if (n23 != 0) {
#ifdef CHECK_CPU
lapackf77_dlacpy("A", &n23, &ib, Q(ctot[0],iil-1+i), &ldq, hS(igpu+1,ind), &n23);
#endif
magma_setdevice(igpu+1);
magma_queue_sync( stream[igpu+1][ind] );
}
if (n12 != 0) {
#ifdef CHECK_CPU
lapackf77_dlacpy("A", &n12, &ib, Q(0,iil-1+i), &ldq, hS(igpu,ind), &n12);
#endif
magma_setdevice(igpu);
magma_queue_sync( stream[igpu][ind] );
}
}
for (igpu = 0; igpu < nrgpu-1; igpu += 2){
if (n23 != 0) {
#ifdef CHECK_CPU
blasf77_dgemm("N", "N", &ni_loc[igpu+1], &ib, &n23, &d_one, hQ2(igpu+1), &n2_loc,
hS(igpu+1,ind), &n23, &d_zero, hQ(igpu+1, ind), &n2_loc);
#endif
magma_setdevice(igpu+1);
magmablasSetKernelStream(stream[igpu+1][ind]);
magma_dgemm(MagmaNoTrans, MagmaNoTrans, ni_loc[igpu+1], ib, n23, d_one, dQ2(igpu+1), n2_loc,
dS(igpu+1, ind), n23, d_zero, dQ(igpu+1, ind), n2_loc);
#ifdef CHECK_CPU
printf("norm Q %d: %f\n", igpu+1, cpu_gpu_ddiff(ni_loc[igpu+1], ib, hQ(igpu+1, ind), n2_loc, dQ(igpu+1, ind), n2_loc));
#endif
}
if (n12 != 0) {
#ifdef CHECK_CPU
blasf77_dgemm("N", "N", &ni_loc[igpu], &ib, &n12, &d_one, hQ2(igpu), &n1_loc,
hS(igpu,ind%2), &n12, &d_zero, hQ(igpu, ind%2), &n1_loc);
#endif
magma_setdevice(igpu);
magmablasSetKernelStream(stream[igpu][ind]);
magma_dgemm(MagmaNoTrans, MagmaNoTrans, ni_loc[igpu], ib, n12, d_one, dQ2(igpu), n1_loc,
dS(igpu, ind), n12, d_zero, dQ(igpu, ind), n1_loc);
#ifdef CHECK_CPU
printf("norm Q %d: %f\n", igpu, cpu_gpu_ddiff(ni_loc[igpu], ib, hQ(igpu, ind), n1_loc, dQ(igpu, ind), n1_loc));
#endif
}
}
for (igpu = 0; igpu < nrgpu-1; igpu += 2){
if (n23 != 0) {
magma_setdevice(igpu+1);
magma_dgetmatrix( ni_loc[igpu+1], ib, dQ(igpu+1, ind), n2_loc,
Q(n1+n2_loc*(igpu/2),iil-1+i), ldq );
// magma_dgetmatrix_async( ni_loc[igpu+1], ib, dQ(igpu+1, ind), n2_loc,
// Q(n1+n2_loc*(igpu/2),iil-1+i), ldq, stream[igpu+1][ind] );
}
if (n12 != 0) {
magma_setdevice(igpu);
magma_dgetmatrix( ni_loc[igpu], ib, dQ(igpu, ind), n1_loc,
Q(n1_loc*(igpu/2),iil-1+i), ldq );
// magma_dgetmatrix_async( ni_loc[igpu], ib, dQ(igpu, ind), n1_loc,
// Q(n1_loc*(igpu/2),iil-1+i), ldq, stream[igpu][ind] );
}
}
}
for (igpu = 0; igpu < nrgpu; ++igpu){
#ifdef CHECK_CPU
magma_free_pinned( hwS[1][igpu] );
magma_free_pinned( hwS[0][igpu] );
magma_free_pinned( hwQ2[igpu] );
magma_free_pinned( hwQ[1][igpu] );
magma_free_pinned( hwQ[0][igpu] );
#endif
magma_setdevice(igpu);
magmablasSetKernelStream(NULL);
magma_queue_sync( stream[igpu][0] );
magma_queue_sync( stream[igpu][1] );
}
if( n23 == 0 )
lapackf77_dlaset("A", &n2, &rk, &d_zero, &d_zero, Q(n1,iil-1), &ldq);
if( n12 == 0 )
lapackf77_dlaset("A", &n1, &rk, &d_zero, &d_zero, Q(0,iil-1), &ldq);
}
}
#ifdef ENABLE_TIMER_DIVIDE_AND_CONQUER
end = get_current_time();
printf("gemms = %6.2f\n", GetTimerValue(start,end)/1000.);
#endif
return MAGMA_SUCCESS;
} /*magma_dlaed3_m*/
