extern "C" void
magma_strdtype1cbHLsym_withQ_v2(magma_int_t n, magma_int_t nb,
float *A, magma_int_t lda,
float *V, magma_int_t ldv,
float *TAU,
magma_int_t st, magma_int_t ed,
magma_int_t sweep, magma_int_t Vblksiz,
float *work)
{
/*
WORK (workspace) float real array, dimension N
*/
magma_int_t ione = 1;
magma_int_t vpos, taupos, len, len2;
float c_one = MAGMA_S_ONE;
magma_bulge_findVTAUpos(n, nb, Vblksiz, sweep-1, st-1, ldv, &vpos, &taupos);
//printf("voici vpos %d taupos %d tpos %d blkid %d \n", vpos, taupos, tpos, blkid);
len = ed-st+1;
*V(vpos) = c_one;
len2 = len-1;
blasf77_scopy( &len2, A(st+1, st-1), &ione, V(vpos+1), &ione );
//memcpy(V(vpos+1), A(st+1, st-1), (len-1)*sizeof(float));
memset(A(st+1, st-1), 0, (len-1)*sizeof(float));
/* Eliminate the col at st-1 */
lapackf77_slarfg( &len, A(st, st-1), V(vpos+1), &ione, TAU(taupos) );
/* apply left and right on A(st:ed,st:ed)*/
magma_slarfxsym_v2(len, A(st,st), lda-1, V(vpos), TAU(taupos), work);
}
// 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|).
float get_residual(
magma_opts &opts,
magma_int_t m, magma_int_t n,
float *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 float c_one = MAGMA_S_ONE;
const float c_neg_one = MAGMA_S_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;
float *x, *b;
// initialize RHS
TESTING_MALLOC_CPU( x, float, n );
TESTING_MALLOC_CPU( b, float, n );
lapackf77_slarnv( &ione, ISEED, &n, b );
blasf77_scopy( &n, b, &ione, x, &ione );
// solve Ax = b
lapackf77_sgetrs( "Notrans", &n, &ione, A, &lda, ipiv, x, &n, &info );
if (info != 0) {
printf("lapackf77_sgetrs returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
// reset to original A
init_matrix( opts, m, n, A, lda );
// compute r = Ax - b, saved in b
blasf77_sgemv( "Notrans", &m, &n, &c_one, A, &lda, x, &ione, &c_neg_one, b, &ione );
// compute residual |Ax - b| / (n*|A|*|x|)
float norm_x, norm_A, norm_r, work[1];
norm_A = lapackf77_slange( "F", &m, &n, A, &lda, work );
norm_r = lapackf77_slange( "F", &n, &ione, b, &n, work );
norm_x = lapackf77_slange( "F", &n, &ione, x, &n, work );
//printf( "r=\n" ); magma_sprint( 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" void
magma_strdtype2cbHLsym_withQ_v2(
magma_int_t n, magma_int_t nb,
float *A, magma_int_t lda,
float *V, magma_int_t ldv,
float *TAU,
magma_int_t st, magma_int_t ed,
magma_int_t sweep, magma_int_t Vblksiz,
float *work)
{
/*
WORK (workspace) float real array, dimension NB
*/
magma_int_t ione = 1;
magma_int_t vpos, taupos;
float conjtmp;
float c_one = MAGMA_S_ONE;
magma_int_t ldx = lda-1;
magma_int_t len = ed - st + 1;
magma_int_t lem = min(ed+nb, n) - ed;
magma_int_t lem2;
if (lem > 0) {
magma_bulge_findVTAUpos(n, nb, Vblksiz, sweep-1, st-1, ldv, &vpos, &taupos);
/* apply remaining right coming from the top block */
lapackf77_slarfx("R", &lem, &len, V(vpos), TAU(taupos), A(ed+1, st), &ldx, work);
}
if (lem > 1) {
magma_bulge_findVTAUpos(n, nb, Vblksiz, sweep-1, ed, ldv, &vpos, &taupos);
/* remove the first column of the created bulge */
*V(vpos) = c_one;
//memcpy(V(vpos+1), A(ed+2, st), (lem-1)*sizeof(float));
lem2 = lem-1;
blasf77_scopy( &lem2, A(ed+2, st), &ione, V(vpos+1), &ione );
memset(A(ed+2, st),0,(lem-1)*sizeof(float));
/* Eliminate the col at st */
lapackf77_slarfg( &lem, A(ed+1, st), V(vpos+1), &ione, TAU(taupos) );
/* apply left on A(J1:J2,st+1:ed) */
len = len-1; /* because we start at col st+1 instead of st. col st is the col that has been removed; */
conjtmp = MAGMA_S_CNJG(*TAU(taupos));
lapackf77_slarfx("L", &lem, &len, V(vpos), &conjtmp, A(ed+1, st+1), &ldx, work);
}
}
/**
Purpose
-------
SLAHR2 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 SGEHRD.
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 REAL 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 REAL 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 REAL 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 REAL array, dimension (NB)
The scalar factors of the elementary reflectors. See Further
Details.
@param[out]
T REAL 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 REAL array, dimension (LDY,NB)
The n-by-nb matrix Y.
@param[in]
ldy 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:
@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_sgeev_aux
********************************************************************/
extern "C" magma_int_t
magma_slahr2(
magma_int_t n, magma_int_t k, magma_int_t nb,
magmaFloat_ptr dA, magma_int_t ldda,
magmaFloat_ptr dV, magma_int_t lddv,
float *A, magma_int_t lda,
float *tau,
float *T, magma_int_t ldt,
float *Y, magma_int_t ldy )
{
#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)
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
magma_int_t ione = 1;
magma_int_t n_k_i_1, n_k;
float scale;
magma_int_t i;
float ei = MAGMA_S_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_scopy( &i,
A(k+i,0), &lda,
T(0,nb-1), &ione );
#if defined(PRECISION_z) || defined(PRECISION_c)
// If real, conjugate row of V.
lapackf77_slacgv(&i, T(0,nb-1), &ione);
#endif
// w = T(0:i-1, 0:i-1) * w
blasf77_strmv( "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_sgemv( "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_scopy( &i,
A(k+1,i), &ione,
T(0,nb-1), &ione );
// w := V1' * b1 = VA(k+1:k+i, 0:i-1)' * w
blasf77_strmv( "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_sgemv( "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_strmv( "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_sgemv( "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_strmv( "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_saxpy( &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_slarfg( &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_ssetvector( n_k_i_1,
A(k+i+1,i), 1,
dV(i+1,i), 1 );
// 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_sgemv( 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 );
// 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_S_NEGATE( tau[i]);
blasf77_sgemv( "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_strmv( "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_sgetvector( n-k,
dA(k,i), 1,
Y(k,i), 1 );
}
// Restore diagonal element
*A(k+nb,nb-1) = ei;
return info;
} /* magma_slahr2 */
/**
Purpose
-------
CHEEVX 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 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 REAL
@param[in]
vu REAL
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 REAL
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*SLAMCH('S'), not zero.
If this routine returns with INFO > 0, indicating that some
eigenvectors did not converge, try setting ABSTOL to
2*SLAMCH('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 REAL array, dimension (N)
On normal exit, the first M elements contain the selected
eigenvalues in ascending order.
@param[out]
dZ COMPLEX 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 array, dimension (LDWA, N)
@param[in]
ldwa INTEGER
The leading dimension of the array wA. LDWA >= max(1,N).
@param
wZ (workspace) COMPLEX 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 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 CHETRD.
\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) REAL 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_cheev_driver
********************************************************************/
extern "C" magma_int_t
magma_cheevx_gpu(
magma_vec_t jobz, magma_range_t range, magma_uplo_t uplo, magma_int_t n,
magmaFloatComplex_ptr dA, magma_int_t ldda,
float vl, float vu,
magma_int_t il, magma_int_t iu, float abstol, magma_int_t *m,
float *w,
magmaFloatComplex_ptr dZ, magma_int_t lddz,
magmaFloatComplex *wA, magma_int_t ldwa,
magmaFloatComplex *wZ, magma_int_t ldwz,
magmaFloatComplex *work, magma_int_t lwork,
float *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;
float safmin;
float bignum;
float smlnum;
float eps, tmp1;
float anrm;
float sigma, d__1;
float rmin, rmax;
magmaFloat_ptr 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_chetrd_nb(n);
lopt = n * (nb + 1);
work[0] = MAGMA_C_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
magmaFloatComplex *a;
magma_cmalloc_cpu( &a, n*n );
magma_cgetmatrix(n, n, dA, ldda, a, n);
lapackf77_cheevx(jobz_, range_, uplo_,
&n, a, &n, &vl, &vu, &il, &iu, &abstol, m,
w, wZ, &ldwz, work, &lwork,
rwork, iwork, ifail, info);
magma_csetmatrix( n, n, a, n, dA, ldda);
magma_csetmatrix( n, *m, wZ, ldwz, dZ, lddz);
magma_free_cpu(a);
return *info;
}
if (MAGMA_SUCCESS != magma_smalloc( &dwork, n )) {
fprintf (stderr, "!!!! device memory allocation error (magma_cheevx_gpu)\n");
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
--w;
--work;
--rwork;
--iwork;
--ifail;
/* Get machine constants. */
safmin = lapackf77_slamch("Safe minimum");
eps = lapackf77_slamch("Precision");
smlnum = safmin / eps;
bignum = 1. / smlnum;
rmin = magma_ssqrt(smlnum);
rmax = magma_ssqrt(bignum);
/* Scale matrix to allowable range, if necessary. */
anrm = magmablas_clanhe(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_clascl(uplo, 0, 0, 1., sigma, n, n, dA,
ldda, info);
if (abstol > 0.) {
abstol *= sigma;
}
if (valeig) {
vl *= sigma;
vu *= sigma;
}
}
/* Call CHETRD 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_chetrd2_gpu(uplo, n, dA, ldda, &rwork[indd], &rwork[inde],
&work[indtau], wA, ldwa, &work[indwrk], llwork, dZ, lddz*n, &iinfo);
#else
magma_chetrd_gpu (uplo, n, dA, ldda, &rwork[indd], &rwork[inde],
&work[indtau], wA, ldwa, &work[indwrk], llwork, &iinfo);
#endif
lopt = n + (magma_int_t)MAGMA_C_REAL(work[indwrk]);
/* If all eigenvalues are desired and ABSTOL is less than or equal to
zero, then call SSTERF or CUNGTR and CSTEQR. If this fails for
some eigenvalue, then try SSTEBZ. */
if ((alleig || (indeig && il == 1 && iu == n)) && abstol <= 0.) {
blasf77_scopy(&n, &rwork[indd], &ione, &w[1], &ione);
indee = indrwk + 2*n;
if (! wantz) {
i__1 = n - 1;
blasf77_scopy(&i__1, &rwork[inde], &ione, &rwork[indee], &ione);
lapackf77_ssterf(&n, &w[1], &rwork[indee], info);
}
else {
lapackf77_clacpy("A", &n, &n, wA, &ldwa, wZ, &ldwz);
lapackf77_cungtr(uplo_, &n, wZ, &ldwz, &work[indtau], &work[indwrk], &llwork, &iinfo);
i__1 = n - 1;
blasf77_scopy(&i__1, &rwork[inde], &ione, &rwork[indee], &ione);
lapackf77_csteqr(jobz_, &n, &w[1], &rwork[indee], wZ, &ldwz, &rwork[indrwk], info);
if (*info == 0) {
for (i = 1; i <= n; ++i) {
ifail[i] = 0;
}
magma_csetmatrix( n, n, wZ, ldwz, dZ, lddz );
}
}
if (*info == 0) {
*m = n;
}
}
/* Otherwise, call SSTEBZ and, if eigenvectors are desired, CSTEIN. */
if (*m == 0) {
*info = 0;
if (wantz) {
order_ = "B";
} else {
order_ = "E";
}
indibl = 1;
indisp = indibl + n;
indiwk = indisp + n;
lapackf77_sstebz(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_cstein(&n, &rwork[indd], &rwork[inde], m, &w[1], &iwork[indibl], &iwork[indisp],
wZ, &ldwz, &rwork[indrwk], &iwork[indiwk], &ifail[1], info);
magma_csetmatrix( n, *m, wZ, ldwz, dZ, lddz );
/* Apply unitary matrix used in reduction to tridiagonal
form to eigenvectors returned by CSTEIN. */
magma_cunmtr_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_sscal(&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_cswap(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_C_MAKE( lopt, 0 );
return *info;
} /* magma_cheevx_gpu */
/**
Purpose
-------
SLAEX3 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 SLAED4 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
SLAED4. 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 REAL array, dimension (N)
D(I) contains the updated eigenvalues for
1 <= I <= K.
@param[out]
Q REAL 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 REAL
The value of the parameter in the rank one update equation.
RHO >= 0 required.
@param[in,out]
dlamda REAL 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 REAL 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 SLAED2).
The rows of the eigenvectors found by SLAED4 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 REAL 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) REAL 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) REAL 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 slaex1.
@param[in]
vl REAL
@param[in]
vu REAL
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_ssyev_aux
********************************************************************/
extern "C" magma_int_t
magma_slaex3(magma_int_t k, magma_int_t n, magma_int_t n1, float* d,
float* Q, magma_int_t ldq, float rho,
float* dlamda, float* Q2, magma_int_t* indx,
magma_int_t* ctot, float* w, float* s, magma_int_t* indxq,
float* dwork,
magma_range_t range, float vl, float vu, magma_int_t il, magma_int_t iu,
magma_int_t* info )
{
#define Q(i_,j_) (Q + (i_) + (j_)*ldq)
float d_one = 1.;
float d_zero = 0.;
magma_int_t ione = 1;
magma_int_t ineg_one = -1;
magma_int_t iil, iiu, rk;
float* dq2= dwork;
float* ds = dq2 + n*(n/2+1);
float* dq = ds + n*(n/2+1);
magma_int_t lddq = n/2 + 1;
magma_int_t i, iq2, j, n12, n2, n23, tmp, lq2;
float 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_ssetvector_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_slamc3(&dlamda[i], &dlamda[i]) - dlamda[i];
for (j = ib; j < ie; ++j) {
magma_int_t tmpp=j+1;
magma_int_t iinfo = 0;
lapackf77_slaed4(&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_slamrg( &k, &nk, d, &ione, &ineg_one, indxq);
//compute the lower and upper bound of the non-deflated eigenvectors
if (valeig)
magma_svrange(k, d, &iil, &iiu, vl, vu);
else if (indeig)
magma_sirange(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_scopy( &ik, &w[ib], &ione, &s[ib], &ione);
// Initialize W(I) = Q(I,I)
tmp = ldq + 1;
blasf77_scopy( &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_snrm2( 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_slamc3(&dlamda[i], &dlamda[i]) - dlamda[i];
for (j = 0; j < k; ++j) {
magma_int_t tmpp=j+1;
magma_int_t iinfo = 0;
lapackf77_slaed4(&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_slamrg( &k, &nk, d, &ione, &ineg_one, indxq);
//compute the lower and upper bound of the non-deflated eigenvectors
if (valeig)
magma_svrange(k, d, &iil, &iiu, vl, vu);
else if (indeig)
magma_sirange(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_scopy( &k, w, &ione, s, &ione);
// Initialize W(I) = Q(I,I)
tmp = ldq + 1;
blasf77_scopy( &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_snrm2( 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_slaed3_k()) {
lapackf77_slacpy("A", &n23, &rk, Q(ctot[0],iil-1), &ldq, s, &n23);
blasf77_sgemm("N", "N", &n2, &rk, &n23, &d_one, &Q2[iq2], &n2,
s, &n23, &d_zero, Q(n1,iil-1), &ldq );
} else {
magma_ssetmatrix( n23, rk, Q(ctot[0],iil-1), ldq, ds, n23 );
magma_sgemm( MagmaNoTrans, MagmaNoTrans, n2, rk, n23, d_one, &dq2[iq2], n2, ds, n23, d_zero, dq, lddq);
magma_sgetmatrix( n2, rk, dq, lddq, Q(n1,iil-1), ldq );
}
} else
lapackf77_slaset("A", &n2, &rk, &d_zero, &d_zero, Q(n1,iil-1), &ldq);
if ( n12 != 0 ) {
if (rk < magma_get_slaed3_k()) {
lapackf77_slacpy("A", &n12, &rk, Q(0,iil-1), &ldq, s, &n12);
blasf77_sgemm("N", "N", &n1, &rk, &n12, &d_one, Q2, &n1,
s, &n12, &d_zero, Q(0,iil-1), &ldq);
} else {
magma_ssetmatrix( n12, rk, Q(0,iil-1), ldq, ds, n12 );
magma_sgemm( MagmaNoTrans, MagmaNoTrans, n1, rk, n12, d_one, dq2, n1, ds, n12, d_zero, dq, lddq);
magma_sgetmatrix( n1, rk, dq, lddq, Q(0,iil-1), ldq );
}
} else
lapackf77_slaset("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_slaex3 */
/**
Purpose
-------
SGEGQR 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).
The routine is designed for tall-and-skinny matrices: M >> N, N <= 128.
This version uses normal equations and SVD in an iterative process that
makes the computation numerically accurate.
Arguments
---------
@param[in]
ikind INTEGER
Several versions are implemented indiceted by the ikind value:
1: This version uses normal equations and SVD in an iterative process
that makes the computation numerically accurate.
2: This version uses a standard LAPACK-based orthogonalization through
MAGMA's QR panel factorization (magma_sgeqr2x3_gpu) and magma_sorgqr
3: MGS
4. Cholesky QR [ Note: this method uses the normal equations which
squares the condition number of A, therefore
||I - Q'Q|| < O(eps cond(A)^2) ]
@param[in]
m INTEGER
The number of rows of the matrix A. m >= n >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. 128 >= n >= 0.
@param[in,out]
dA REAL 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.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,m).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
@param
dwork (GPU workspace) REAL array, dimension:
n^2 for ikind = 1
3 n^2 + min(m, n) + 2 for ikind = 2
0 (not used) for ikind = 3
n^2 for ikind = 4
@param[out]
work (CPU workspace) REAL array, dimension 3 n^2.
On exit, work(1:n^2) holds the rectangular matrix R.
Preferably, for higher performance, work should be in pinned memory.
@param[out]
info 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.
@ingroup magma_sgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_sgegqr_gpu( magma_int_t ikind, magma_int_t m, magma_int_t n,
float *dA, magma_int_t ldda,
float *dwork, float *work,
magma_int_t *info )
{
#define work(i_,j_) (work + (i_) + (j_)*n)
#define dA(i_,j_) (dA + (i_) + (j_)*ldda)
magma_int_t i = 0, j, k, n2 = n*n;
magma_int_t ione = 1;
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float cn = 200., mins, maxs;
/* check arguments */
*info = 0;
if (ikind < 1 || ikind > 4) {
*info = -1;
} else if (m < 0 || m < n) {
*info = -2;
} else if (n < 0 || n > 128) {
*info = -3;
} else if (ldda < max(1,m)) {
*info = -5;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (ikind == 1) {
// === Iterative, based on SVD ============================================================
float *U, *VT, *vt, *R, *G, *hwork, *tau;
float *S;
R = work; // Size n * n
G = R + n*n; // Size n * n
VT = G + n*n; // Size n * n
magma_smalloc_cpu( &hwork, 32 + 2*n*n + 2*n);
if ( hwork == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_int_t lwork=n*n+32; // First part f hwork; used as workspace in svd
U = hwork + n*n + 32; // Size n*n
S = (float *)(U+n*n); // Size n
tau = U + n*n + n; // Size n
#if defined(PRECISION_c) || defined(PRECISION_z)
float *rwork;
magma_smalloc_cpu( &rwork, 5*n);
if ( rwork == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
#endif
do {
i++;
magma_sgemm(MagmaConjTrans, MagmaNoTrans, n, n, m, c_one, dA, ldda, dA, ldda, c_zero, dwork, n );
magma_sgetmatrix(n, n, dwork, n, G, n);
#if defined(PRECISION_s) || defined(PRECISION_d)
lapackf77_sgesvd("n", "a", &n, &n, G, &n, S, U, &n, VT, &n,
hwork, &lwork, info);
#else
lapackf77_sgesvd("n", "a", &n, &n, G, &n, S, U, &n, VT, &n,
hwork, &lwork, rwork, info);
#endif
mins = 100.f, maxs = 0.f;
for (k=0; k < n; k++) {
S[k] = magma_ssqrt( 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_sgeqrf(&n, &n, VT, &n, tau, hwork, &lwork, info);
if (i == 1)
blasf77_scopy(&n2, VT, &ione, R, &ione);
else
blasf77_strmm("l", "u", "n", "n", &n, &n, &c_one, VT, &n, R, &n);
magma_ssetmatrix(n, n, VT, n, dwork, n);
magma_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit, m, n, c_one, dwork, n, dA, 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 );
#if defined(PRECISION_c) || defined(PRECISION_z)
magma_free_cpu( rwork );
#endif
// ================== end of ikind == 1 ===================================================
}
else if (ikind == 2) {
// ================== LAPACK based ===================================================
magma_int_t min_mn = min(m, n);
magma_int_t nb = n;
float *dtau = dwork + 2*n*n, *d_T = dwork, *ddA = dwork + n*n;
float *tau = work+n*n;
magmablas_slaset( MagmaFull, n, n, c_zero, c_zero, d_T, n );
magma_sgeqr2x3_gpu(m, n, dA, ldda, dtau, d_T, ddA,
(float *)(dwork+min_mn+2*n*n), info);
magma_sgetmatrix( min_mn, 1, dtau, min_mn, tau, min_mn);
magma_sgetmatrix( n, n, ddA, n, work, n);
magma_sorgqr_gpu( m, n, n, dA, ldda, tau, d_T, nb, info );
// ================== end of ikind == 2 ===================================================
}
else if (ikind == 3) {
// ================== MGS ===================================================
for(magma_int_t j = 0; j<n; j++){
for(magma_int_t i = 0; i<j; i++){
*work(i, j) = magma_sdot(m, dA(0,i), 1, dA(0,j), 1);
magma_saxpy(m, -(*work(i,j)), dA(0,i), 1, dA(0,j), 1);
}
for(magma_int_t i = j; i<n; i++)
*work(i, j) = MAGMA_S_ZERO;
//*work(j,j) = MAGMA_S_MAKE( magma_snrm2(m, dA(0,j), 1), 0. );
*work(j,j) = magma_sdot(m, dA(0,j), 1, dA(0,j), 1);
*work(j,j) = MAGMA_S_MAKE( sqrt(MAGMA_S_REAL( *work(j,j) )), 0.);
magma_sscal(m, 1./ *work(j,j), dA(0,j), 1);
}
// ================== end of ikind == 3 ===================================================
}
else if (ikind == 4) {
// ================== Cholesky QR ===================================================
magma_sgemm(MagmaConjTrans, MagmaNoTrans, n, n, m, c_one, dA, ldda, dA, ldda, c_zero, dwork, n );
magma_sgetmatrix(n, n, dwork, n, work, n);
lapackf77_spotrf("u", &n, work, &n, info);
magma_ssetmatrix(n, n, work, n, dwork, n);
magma_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit, m, n, c_one, dwork, n, dA, ldda);
// ================== end of ikind == 4 ===================================================
}
return *info;
} /* magma_sgegqr_gpu */
magma_int_t magma_strevc3(
magma_side_t side, magma_vec_t howmany,
magma_int_t *select, // logical in fortran
magma_int_t n,
float *T, magma_int_t ldt,
float *VL, magma_int_t ldvl,
float *VR, magma_int_t ldvr,
magma_int_t mm, magma_int_t *mout,
float *work, magma_int_t lwork,
#ifdef COMPLEX
float *rwork,
#endif
magma_int_t *info )
{
#define T(i,j) (T + (i) + (j)*ldt)
#define VL(i,j) (VL + (i) + (j)*ldvl)
#define VR(i,j) (VR + (i) + (j)*ldvr)
#define X(i,j) (X + (i)-1 + ((j)-1)*2) // still as 1-based indices
#define work(i,j) (work + (i) + (j)*n)
// constants
const magma_int_t ione = 1;
const float c_zero = 0;
const float c_one = 1;
const magma_int_t nbmin = 16, nbmax = 256;
// .. Local Scalars ..
magma_int_t allv, bothv, leftv, over, pair, rightv, somev;
magma_int_t i, ierr, ii, ip, is, j, k, ki, ki2,
iv, n2, nb, nb2, version;
float emax, remax;
// .. Local Arrays ..
// since iv is a 1-based index, allocate one extra here
magma_int_t iscomplex[ nbmax+1 ];
// Decode and test the input parameters
bothv = (side == MagmaBothSides);
rightv = (side == MagmaRight) || bothv;
leftv = (side == MagmaLeft ) || bothv;
allv = (howmany == MagmaAllVec);
over = (howmany == MagmaBacktransVec);
somev = (howmany == MagmaSomeVec);
*info = 0;
if ( ! rightv && ! leftv )
*info = -1;
else if ( ! allv && ! over && ! somev )
*info = -2;
else if ( n < 0 )
*info = -4;
else if ( ldt < max( 1, n ) )
*info = -6;
else if ( ldvl < 1 || ( leftv && ldvl < n ) )
*info = -8;
else if ( ldvr < 1 || ( rightv && ldvr < n ) )
*info = -10;
else if ( lwork < max( 1, 3*n ) )
*info = -14;
else {
// Set mout to the number of columns required to store the selected
// eigenvectors, standardize the array select if necessary, and
// test mm.
if ( somev ) {
*mout = 0;
pair = false;
for( j=0; j < n; ++j ) {
if ( pair ) {
pair = false;
select[j] = false;
}
else {
if ( j < n-1 ) {
if ( *T(j+1,j) == c_zero ) {
if ( select[j] ) {
*mout += 1;
}
}
else {
pair = true;
if ( select[j] || select[j+1] ) {
select[j] = true;
*mout += 2;
}
}
}
else if ( select[n-1] ) {
*mout += 1;
}
}
}
}
else {
*mout = n;
}
if ( mm < *mout ) {
*info = -11;
}
}
if ( *info != 0 ) {
magma_xerbla( __func__, -(*info) );
return *info;
}
// Quick return if possible.
if ( n == 0 ) {
return *info;
}
// Use blocked version (2) if sufficient workspace.
// Requires 1 vector for 1-norms, and 2*nb vectors for x and Q*x.
// Zero-out the workspace to avoid potential NaN propagation.
nb = 2;
if ( lwork >= n + 2*n*nbmin ) {
version = 2;
nb = (lwork - n) / (2*n);
nb = min( nb, nbmax );
nb2 = 1 + 2*nb;
lapackf77_slaset( "F", &n, &nb2, &c_zero, &c_zero, work, &n );
}
else {
version = 1;
}
// Compute 1-norm of each column of strictly upper triangular
// part of T to control overflow in triangular solver.
*work(0,0) = c_zero;
for( j=1; j < n; ++j ) {
*work(j,0) = c_zero;
for( i=0; i < j; ++i ) {
*work(j,0) += fabsf( *T(i,j) );
}
}
magma_timer_t time_total=0, time_trsv=0, time_gemm=0, time_gemv=0, time_trsv_sum=0, time_gemm_sum=0, time_gemv_sum=0;
timer_start( time_total );
// Index ip is used to specify the real or complex eigenvalue:
// ip = 0, real eigenvalue (wr),
// = 1, first of conjugate complex pair: (wr,wi)
// = -1, second of conjugate complex pair: (wr,wi)
// iscomplex array stores ip for each column in current block.
if ( rightv ) {
// ============================================================
// Compute right eigenvectors.
// iv is index of column in current block (1-based).
// For complex right vector, uses iv-1 for real part and iv for complex part.
// Non-blocked version always uses iv=2;
// blocked version starts with iv=nb, goes down to 1 or 2.
// (Note the "0-th" column is used for 1-norms computed above.)
iv = 2;
if ( version == 2 ) {
iv = nb;
}
timer_start( time_trsv );
ip = 0;
is = *mout - 1;
for( ki=n-1; ki >= 0; --ki ) {
if ( ip == -1 ) {
// previous iteration (ki+1) was second of conjugate pair,
// so this ki is first of conjugate pair; skip to end of loop
ip = 1;
continue;
}
else if ( ki == 0 ) {
// last column, so this ki must be real eigenvalue
ip = 0;
}
else if ( *T(ki,ki-1) == c_zero ) {
// zero on sub-diagonal, so this ki is real eigenvalue
ip = 0;
}
else {
// non-zero on sub-diagonal, so this ki is second of conjugate pair
ip = -1;
}
if ( somev ) {
if ( ip == 0 ) {
if ( ! select[ki] ) {
continue;
}
}
else {
if ( ! select[ki-1] ) {
continue;
}
}
}
if ( ip == 0 ) {
// ------------------------------------------------------------
// Real right eigenvector
// Solve upper quasi-triangular system:
// [ T(0:ki-1,0:ki-1) - wr ]*X = -T(0:ki-1,ki)
magma_slaqtrsd( MagmaNoTrans, ki+1, T(0,0), ldt,
work(0,iv), n, work(0,0), &ierr );
// Copy the vector x or Q*x to VR and normalize.
if ( ! over ) {
// ------------------------------
// no back-transform: copy x to VR and normalize.
n2 = ki+1;
blasf77_scopy( &n2, work(0,iv), &ione, VR(0,is), &ione );
ii = blasf77_isamax( &n2, VR(0,is), &ione ) - 1; // subtract 1; ii is 0-based
remax = c_one / fabsf( *VR(ii,is) );
blasf77_sscal( &n2, &remax, VR(0,is), &ione );
for( k=ki + 1; k < n; ++k ) {
*VR(k,is) = c_zero;
}
}
else if ( version == 1 ) {
// ------------------------------
// version 1: back-transform each vector with GEMV, Q*x.
time_trsv_sum += timer_stop( time_trsv );
timer_start( time_gemv );
if ( ki > 0 ) {
n2 = ki;
blasf77_sgemv( "n", &n, &n2, &c_one,
VR, &ldvr,
work(0, iv), &ione,
work(ki,iv), VR(0,ki), &ione );
}
time_gemv_sum += timer_stop( time_gemv );
ii = blasf77_isamax( &n, VR(0,ki), &ione ) - 1; // subtract 1; ii is 0-based
remax = c_one / fabsf( *VR(ii,ki) );
blasf77_sscal( &n, &remax, VR(0,ki), &ione );
timer_start( time_trsv );
}
else if ( version == 2 ) {
// ------------------------------
// version 2: back-transform block of vectors with GEMM
// zero out below vector
for( k=ki + 1; k < n; ++k ) {
*work(k,iv) = c_zero;
}
iscomplex[ iv ] = ip;
// back-transform and normalization is done below
}
} // end real eigenvector
else {
// ------------------------------------------------------------
// Complex right eigenvector
// Solve upper quasi-triangular system:
// [ T(0:ki-2,0:ki-2) - (wr+i*wi) ]*x = u
magma_slaqtrsd( MagmaNoTrans, ki+1, T(0,0), ldt,
work(0,iv-1), n, work(0,0), &ierr );
// Copy the vector x or Q*x to VR and normalize.
if ( ! over ) {
// ------------------------------
// no back-transform: copy x to VR and normalize.
n2 = ki+1;
blasf77_scopy( &n2, work(0,iv-1), &ione, VR(0,is-1), &ione );
blasf77_scopy( &n2, work(0,iv ), &ione, VR(0,is ), &ione );
emax = c_zero;
for( k=0; k <= ki; ++k ) {
emax = max( emax, fabsf(*VR(k,is-1)) + fabsf(*VR(k,is)) );
}
remax = c_one / emax;
blasf77_sscal( &n2, &remax, VR(0,is-1), &ione );
blasf77_sscal( &n2, &remax, VR(0,is ), &ione );
for( k=ki + 1; k < n; ++k ) {
*VR(k,is-1) = c_zero;
*VR(k,is ) = c_zero;
}
}
else if ( version == 1 ) {
// ------------------------------
// version 1: back-transform each vector with GEMV, Q*x.
time_trsv_sum += timer_stop( time_trsv );
timer_start( time_gemv );
if ( ki > 1 ) {
n2 = ki-1;
blasf77_sgemv( "n", &n, &n2, &c_one,
VR, &ldvr,
work(0, iv-1), &ione,
work(ki-1,iv-1), VR(0,ki-1), &ione );
blasf77_sgemv( "n", &n, &n2, &c_one,
VR, &ldvr,
work(0, iv), &ione,
work(ki,iv), VR(0,ki), &ione );
}
else {
blasf77_sscal( &n, work(ki-1,iv-1), VR(0,ki-1), &ione );
blasf77_sscal( &n, work(ki, iv ), VR(0,ki ), &ione );
}
time_gemv_sum += timer_stop( time_gemv );
emax = c_zero;
for( k=0; k < n; ++k ) {
emax = max( emax, fabsf(*VR(k,ki-1)) + fabsf(*VR(k,ki)) );
}
remax = c_one / emax;
blasf77_sscal( &n, &remax, VR(0,ki-1), &ione );
blasf77_sscal( &n, &remax, VR(0,ki ), &ione );
timer_start( time_trsv );
}
else if ( version == 2 ) {
// ------------------------------
// version 2: back-transform block of vectors with GEMM
// zero out below vector
for( k=ki + 1; k < n; ++k ) {
*work(k,iv-1) = c_zero;
*work(k,iv ) = c_zero;
}
iscomplex[ iv-1 ] = -ip;
iscomplex[ iv ] = ip;
iv -= 1;
// back-transform and normalization is done below
}
} // end real or complex vector
if ( version == 2 ) {
// ------------------------------------------------------------
// Blocked version of back-transform
// For complex case, ki2 includes both vectors (ki-1 and ki)
if ( ip == 0 ) {
ki2 = ki;
}
else {
ki2 = ki - 1;
}
// Columns iv:nb of work are valid vectors.
// When the number of vectors stored reaches nb-1 or nb,
// or if this was last vector, do the GEMM
if ( (iv <= 2) || (ki2 == 0) ) {
time_trsv_sum += timer_stop( time_trsv );
timer_start( time_gemm );
nb2 = nb-iv+1;
n2 = ki2+nb-iv+1;
blasf77_sgemm( "n", "n", &n, &nb2, &n2, &c_one,
VR, &ldvr,
work(0,iv), &n,
&c_zero,
work(0,nb+iv), &n );
time_gemm_sum += timer_stop( time_gemm );
// normalize vectors
// TODO if somev, should copy vectors individually to correct location.
for( k=iv; k <= nb; ++k ) {
if ( iscomplex[k] == 0 ) {
// real eigenvector
ii = blasf77_isamax( &n, work(0,nb+k), &ione ) - 1; // subtract 1; ii is 0-based
remax = c_one / fabsf( *work(ii,nb+k) );
}
else if ( iscomplex[k] == 1 ) {
// first eigenvector of conjugate pair
emax = c_zero;
for( ii=0; ii < n; ++ii ) {
emax = max( emax, fabsf( *work(ii,nb+k ) )
+ fabsf( *work(ii,nb+k+1) ) );
}
remax = c_one / emax;
// else if iscomplex[k] == -1
// second eigenvector of conjugate pair
// reuse same remax as previous k
}
blasf77_sscal( &n, &remax, work(0,nb+k), &ione );
}
nb2 = nb-iv+1;
lapackf77_slacpy( "F", &n, &nb2,
work(0,nb+iv), &n,
VR(0,ki2), &ldvr );
iv = nb;
timer_start( time_trsv );
}
else {
iv -= 1;
}
} // end blocked back-transform
is -= 1;
if ( ip != 0 ) {
is -= 1;
}
}
}
timer_stop( time_trsv );
timer_stop( time_total );
timer_printf( "trevc trsv %.4f, gemm %.4f, gemv %.4f, total %.4f\n",
time_trsv_sum, time_gemm_sum, time_gemv_sum, time_total );
if ( leftv ) {
// ============================================================
// Compute left eigenvectors.
// iv is index of column in current block (1-based).
// For complex left vector, uses iv for real part and iv+1 for complex part.
// Non-blocked version always uses iv=1;
// blocked version starts with iv=1, goes up to nb-1 or nb.
// (Note the "0-th" column is used for 1-norms computed above.)
iv = 1;
ip = 0;
is = 0;
for( ki=0; ki < n; ++ki ) {
if ( ip == 1 ) {
// previous iteration (ki-1) was first of conjugate pair,
// so this ki is second of conjugate pair; skip to end of loop
ip = -1;
continue;
}
else if ( ki == n-1 ) {
// last column, so this ki must be real eigenvalue
ip = 0;
}
else if ( *T(ki+1,ki) == c_zero ) {
// zero on sub-diagonal, so this ki is real eigenvalue
ip = 0;
}
else {
// non-zero on sub-diagonal, so this ki is first of conjugate pair
ip = 1;
}
if ( somev ) {
if ( ! select[ki] ) {
continue;
}
}
if ( ip == 0 ) {
// ------------------------------------------------------------
// Real left eigenvector
// Solve transposed quasi-triangular system:
// [ T(ki+1:n,ki+1:n) - wr ]**T * X = -T(ki+1:n,ki)
magma_slaqtrsd( MagmaTrans, n-ki, T(ki,ki), ldt,
work(ki,iv), n, work(ki,0), &ierr );
// Copy the vector x or Q*x to VL and normalize.
if ( ! over ) {
// ------------------------------
// no back-transform: copy x to VL and normalize.
n2 = n-ki;
blasf77_scopy( &n2, work(ki,iv), &ione, VL(ki,is), &ione );
ii = blasf77_isamax( &n2, VL(ki,is), &ione ) + ki - 1; // subtract 1; ii is 0-based
remax = c_one / fabsf( *VL(ii,is) );
blasf77_sscal( &n2, &remax, VL(ki,is), &ione );
for( k=0; k < ki; ++k ) {
*VL(k,is) = c_zero;
}
}
else if ( version == 1 ) {
// ------------------------------
// version 1: back-transform each vector with GEMV, Q*x.
if ( ki < n-1 ) {
n2 = n-ki-1;
blasf77_sgemv( "n", &n, &n2, &c_one,
VL(0,ki+1), &ldvl,
work(ki+1,iv), &ione,
work(ki, iv), VL(0,ki), &ione );
}
ii = blasf77_isamax( &n, VL(0,ki), &ione ) - 1; // subtract 1; ii is 0-based
remax = c_one / fabsf( *VL(ii,ki) );
blasf77_sscal( &n, &remax, VL(0,ki), &ione );
}
else if ( version == 2 ) {
// ------------------------------
// version 2: back-transform block of vectors with GEMM
// zero out above vector
// could go from (ki+1)-NV+1 to ki
for( k=0; k < ki; ++k ) {
*work(k,iv) = c_zero;
}
iscomplex[ iv ] = ip;
// back-transform and normalization is done below
}
} // end real eigenvector
else {
// ------------------------------------------------------------
// Complex left eigenvector
// Solve transposed quasi-triangular system:
// [ T(ki+2:n,ki+2:n)**T - (wr-i*wi) ]*X = V
magma_slaqtrsd( MagmaTrans, n-ki, T(ki,ki), ldt,
work(ki,iv), n, work(ki,0), &ierr );
// Copy the vector x or Q*x to VL and normalize.
if ( ! over ) {
// ------------------------------
// no back-transform: copy x to VL and normalize.
n2 = n-ki;
blasf77_scopy( &n2, work(ki,iv ), &ione, VL(ki,is ), &ione );
blasf77_scopy( &n2, work(ki,iv+1), &ione, VL(ki,is+1), &ione );
emax = c_zero;
for( k=ki; k < n; ++k ) {
emax = max( emax, fabsf(*VL(k,is))+ fabsf(*VL(k,is+1)) );
}
remax = c_one / emax;
blasf77_sscal( &n2, &remax, VL(ki,is ), &ione );
blasf77_sscal( &n2, &remax, VL(ki,is+1), &ione );
for( k=0; k < ki; ++k ) {
*VL(k,is ) = c_zero;
*VL(k,is+1) = c_zero;
}
}
else if ( version == 1 ) {
// ------------------------------
// version 1: back-transform each vector with GEMV, Q*x.
if ( ki < n-2 ) {
n2 = n-ki-2;
blasf77_sgemv( "n", &n, &n2, &c_one,
VL(0,ki+2), &ldvl,
work(ki+2,iv), &ione,
work(ki, iv), VL(0,ki), &ione );
blasf77_sgemv( "n", &n, &n2, &c_one,
VL(0,ki+2), &ldvl,
work(ki+2,iv+1), &ione,
work(ki+1,iv+1), VL(0,ki+1), &ione );
}
else {
blasf77_sscal( &n, work(ki, iv ), VL(0, ki ), &ione );
blasf77_sscal( &n, work(ki+1,iv+1), VL(0, ki+1), &ione );
}
emax = c_zero;
for( k=0; k < n; ++k ) {
emax = max( emax, fabsf(*VL(k,ki))+ fabsf(*VL(k,ki+1)) );
}
remax = c_one / emax;
blasf77_sscal( &n, &remax, VL(0,ki ), &ione );
blasf77_sscal( &n, &remax, VL(0,ki+1), &ione );
}
else if ( version == 2 ) {
// ------------------------------
// version 2: back-transform block of vectors with GEMM
// zero out above vector
// could go from (ki+1)-NV+1 to ki
for( k=0; k < ki; ++k ) {
*work(k,iv ) = c_zero;
*work(k,iv+1) = c_zero;
}
iscomplex[ iv ] = ip;
iscomplex[ iv+1 ] = -ip;
iv += 1;
// back-transform and normalization is done below
}
} // end real or complex eigenvector
if ( version == 2 ) {
// -------------------------------------------------
// Blocked version of back-transform
// For complex case, (ki2+1) includes both vectors (ki+1) and (ki+2)
if ( ip == 0 ) {
ki2 = ki;
}
else {
ki2 = ki + 1;
}
// Columns 1:iv of work are valid vectors.
// When the number of vectors stored reaches nb-1 or nb,
// or if this was last vector, do the GEMM
if ( (iv >= nb-1) || (ki2 == n-1) ) {
n2 = n-(ki2+1)+iv;
blasf77_sgemm( "n", "n", &n, &iv, &n2, &c_one,
VL(0,ki2-iv+1), &ldvl,
work(ki2-iv+1,1), &n,
&c_zero,
work(0,nb+1), &n );
// normalize vectors
for( k=1; k <= iv; ++k ) {
if ( iscomplex[k] == 0 ) {
// real eigenvector
ii = blasf77_isamax( &n, work(0,nb+k), &ione ) - 1; // subtract 1; ii is 0-based
remax = c_one / fabsf( *work(ii,nb+k) );
}
else if ( iscomplex[k] == 1) {
// first eigenvector of conjugate pair
emax = c_zero;
for( ii=0; ii < n; ++ii ) {
emax = max( emax, fabsf( *work(ii,nb+k ) )
+ fabsf( *work(ii,nb+k+1) ) );
}
remax = c_one / emax;
// else if iscomplex[k] == -1
// second eigenvector of conjugate pair
// reuse same remax as previous k
}
blasf77_sscal( &n, &remax, work(0,nb+k), &ione );
}
lapackf77_slacpy( "F", &n, &iv,
work(0,nb+1), &n,
VL(0,ki2-iv+1), &ldvl );
iv = 1;
}
else {
iv += 1;
}
} // blocked back-transform
is += 1;
if ( ip != 0 ) {
is += 1;
}
}
}
return *info;
} // end of STREVC3
/**
Purpose
-------
SLAEX1 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 SLAED2.
The second stage consists of calculating the updated
eigenvalues. This is done by finding the roots of the secular
equation via the routine SLAED4 (as called by SLAED3).
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 REAL 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 REAL 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 REAL
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) REAL array, dimension (4*N + N**2)
@param
iwork (workspace) INTEGER array, dimension (4*N)
@param
dwork (devices workspaces) REAL 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 REAL
@param[in]
vu REAL
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_ssyev_aux
********************************************************************/
extern "C" magma_int_t
magma_slaex1_m(magma_int_t nrgpu, magma_int_t n, float* d, float* Q, magma_int_t ldq,
magma_int_t* indxq, float rho, magma_int_t cutpnt,
float* work, magma_int_t* iwork, float** dwork,
magma_queue_t stream[MagmaMaxGPUs][2],
magma_range_t range, float vl, float 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 *info;
// The following values are integer pointers which indicate
// the portion of the workspace
// used by a particular array in SLAED2 and SLAED3.
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_scopy( &cutpnt, Q(cutpnt-1, 0), &ldq, &work[iz], &ione);
tmp = n-cutpnt;
blasf77_scopy( &tmp, Q(cutpnt, cutpnt), &ldq, &work[iz+cutpnt], &ione);
// Deflate eigenvalues.
lapackf77_slaed2(&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 *info;
// Solve Secular Equation.
if ( k != 0 ) {
is = (iwork[coltyp]+iwork[coltyp+1])*cutpnt + (iwork[coltyp+1]+iwork[coltyp+2])*(n-cutpnt) + iq2;
magma_slaex3_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 *info;
}
else {
for (i = 0; i < n; ++i)
indxq[i] = i+1;
}
return *info;
} /* magma_slaex1_m */
extern "C" magma_int_t
magma_slaex1(
magma_int_t n, float* d, float* q, magma_int_t ldq,
magma_int_t* indxq, float rho, magma_int_t cutpnt,
float* work, magma_int_t* iwork, magmaFloat_ptr dwork,
magma_range_t range, float vl, float vu,
magma_int_t il, magma_int_t iu,
magma_queue_t queue,
magma_int_t* info)
{
/*
-- clMAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date November 2014
.. Scalar Arguments ..
CHARACTER RANGE
INTEGER IL, IU, CUTPNT, INFO, LDQ, N
REAL RHO, VL, VU
..
.. Array Arguments ..
INTEGER INDXQ( * ), iwork[* )
REAL D( * ), Q( LDQ, * ), WORK( * ), DWORK( * )
..
Purpose
=======
SLAEX1 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 SLAED2.
The second stage consists of calculating the updated
eigenvalues. This is done by finding the roots of the secular
equation via the routine SLAED4 (as called by SLAED3).
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) REAL array, dimension (N)
On entry, the eigenvalues of the rank-1-perturbed matrix.
On exit, the eigenvalues of the repaired matrix.
Q (input/output) REAL 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) REAL
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) REAL array, dimension (4*N + N**2)
IWORK (workspace) INTEGER array, dimension (4*N)
DWORK (device workspace) REAL 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) REAL
VU (input) REAL
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 SLAED2 and SLAED3.
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_scopy( &cutpnt, Q(cutpnt-1, 0), &ldq, &work[iz], &ione);
tmp = n-cutpnt;
blasf77_scopy( &tmp, Q(cutpnt, cutpnt), &ldq, &work[iz+cutpnt], &ione);
// Deflate eigenvalues.
magma_slaed2(&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_slaex3(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, queue, info );
if( *info != 0 )
return MAGMA_SUCCESS;
}
else {
for (i = 0; i<n; ++i)
indxq[i] = i+1;
}
return MAGMA_SUCCESS;
} /* magma_slaex1 */
extern "C" magma_int_t
magma_slaex3(magma_int_t k, magma_int_t n, magma_int_t n1, float* d,
float* q, magma_int_t ldq, float rho,
float* dlamda, float* q2, magma_int_t* indx,
magma_int_t* ctot, float* w, float* s, magma_int_t* indxq,
float* dwork,
char range, float vl, float vu, magma_int_t il, magma_int_t iu,
magma_int_t* info )
{
/*
Purpose
=======
SLAEX3 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 SLAED4 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
SLAED4. 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) REAL array, dimension (N)
D(I) contains the updated eigenvalues for
1 <= I <= K.
Q (output) REAL 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) REAL
The value of the parameter in the rank one update equation.
RHO >= 0 required.
DLAMDA (input/output) REAL 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) REAL 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 SLAED2).
The rows of the eigenvectors found by SLAED4 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) REAL array, dimension (K)
The first K elements of this array contain the components
of the deflation-adjusted updating vector. Destroyed on
output.
S (workspace) REAL 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 (device workspace) REAL array, dimension (3*N*N/2+3*N)
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.
===================================================================== */
float d_one = 1.;
float d_zero = 0.;
magma_int_t ione = 1;
magma_int_t ineg_one = -1;
char range_[] = {range, 0};
magma_int_t iil, iiu, rk;
float* dq2= dwork;
float* ds = dq2 + n*(n/2+1);
float* dq = ds + n*(n/2+1);
magma_int_t lddq = n/2 + 1;
magma_int_t i,iq2,j,n12,n2,n23,tmp,lq2;
float 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.*/
n2 = n - n1;
n12 = ctot[0] + ctot[1];
n23 = ctot[1] + ctot[2];
iq2 = n1 * n12;
lq2 = iq2 + n2 * n23;
magma_ssetvector_async( lq2, q2, 1, dq2, 1, NULL );
#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_slamc3(&dlamda[i], &dlamda[i]) - dlamda[i];
for(j = ib; j < ie; ++j){
magma_int_t tmpp=j+1;
magma_int_t iinfo = 0;
lapackf77_slaed4(&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_slamrg( &k, &nk, d, &ione , &ineg_one, indxq);
//compute the lower and upper bound of the non-deflated eigenvectors
if (valeig)
magma_svrange(k, d, &iil, &iiu, vl, vu);
else if (indeig)
magma_sirange(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_scopy( &ik, &w[ib], &ione, &s[ib], &ione);
// Initialize W(I) = Q(I,I)
tmp = ldq + 1;
blasf77_scopy( &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_snrm2( 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_slamc3(&dlamda[i], &dlamda[i]) - dlamda[i];
for(j = 0; j < k; ++j){
magma_int_t tmpp=j+1;
magma_int_t iinfo = 0;
lapackf77_slaed4(&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_slamrg( &k, &nk, d, &ione , &ineg_one, indxq);
//compute the lower and upper bound of the non-deflated eigenvectors
if (valeig)
magma_svrange(k, d, &iil, &iiu, vl, vu);
else if (indeig)
magma_sirange(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_scopy( &k, w, &ione, s, &ione);
// Initialize W(I) = Q(I,I)
tmp = ldq + 1;
blasf77_scopy( &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_snrm2( 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
magma_queue_sync( NULL );
if (rk != 0){
if( n23 != 0 ){
if (rk < magma_get_slaed3_k()){
lapackf77_slacpy("A", &n23, &rk, Q(ctot[0],iil-1), &ldq, s, &n23);
blasf77_sgemm("N", "N", &n2, &rk, &n23, &d_one, &q2[iq2], &n2,
s, &n23, &d_zero, Q(n1,iil-1), &ldq );
} else {
magma_ssetmatrix( n23, rk, Q(ctot[0],iil-1), ldq, ds, n23 );
magma_sgemm('N', 'N', n2, rk, n23, d_one, &dq2[iq2], n2, ds, n23, d_zero, dq, lddq);
magma_sgetmatrix( n2, rk, dq, lddq, Q(n1,iil-1), ldq );
}
} else
lapackf77_slaset("A", &n2, &rk, &d_zero, &d_zero, Q(n1,iil-1), &ldq);
if( n12 != 0 ) {
if (rk < magma_get_slaed3_k()){
lapackf77_slacpy("A", &n12, &rk, Q(0,iil-1), &ldq, s, &n12);
blasf77_sgemm("N", "N", &n1, &rk, &n12, &d_one, q2, &n1,
s, &n12, &d_zero, Q(0,iil-1), &ldq);
} else {
magma_ssetmatrix( n12, rk, Q(0,iil-1), ldq, ds, n12 );
magma_sgemm('N', 'N', n1, rk, n12, d_one, dq2, n1, ds, n12, d_zero, dq, lddq);
magma_sgetmatrix( n1, rk, dq, lddq, Q(0,iil-1), ldq );
}
} else
lapackf77_slaset("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_slaed3*/
/**
Purpose
-------
SLAEX0 computes all eigenvalues and the choosen eigenvectors of a
symmetric tridiagonal matrix using the divide and conquer method.
Arguments
---------
@param[in]
n INTEGER
The dimension of the symmetric tridiagonal matrix. N >= 0.
@param[in,out]
d REAL array, dimension (N)
On entry, the main diagonal of the tridiagonal matrix.
On exit, its eigenvalues.
@param[in]
e REAL array, dimension (N-1)
The off-diagonal elements of the tridiagonal matrix.
On exit, E has been destroyed.
@param[in,out]
Q REAL array, dimension (LDQ, N)
On entry, Q will be the identity matrix.
On exit, Q contains the eigenvectors of the
tridiagonal matrix.
@param[in]
ldq INTEGER
The leading dimension of the array Q. If eigenvectors are
desired, then LDQ >= max(1,N). In any case, LDQ >= 1.
@param
work (workspace) REAL array,
the dimension of WORK >= 4*N + N**2.
@param
iwork (workspace) INTEGER array,
the dimension of IWORK >= 3 + 5*N.
@param
dwork (workspace) REAL 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.
@param[in]
vl REAL
@param[in]
vu REAL
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: 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
@ingroup magma_ssyev_aux
********************************************************************/
extern "C" magma_int_t
magma_slaex0(
magma_int_t n,
float *d, float *e,
float *Q, magma_int_t ldq,
float *work, magma_int_t *iwork,
magmaFloat_ptr dwork,
magma_range_t range, float vl, float vu,
magma_int_t il, magma_int_t iu,
magma_int_t *info)
{
#define Q(i_,j_) (Q + (i_) + (j_)*ldq)
magma_int_t ione = 1;
magma_range_t range2;
magma_int_t curlvl, 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 *info;
}
// Quick return if possible
if (n == 0)
return *info;
smlsiz = magma_get_smlsize_divideconquer();
// 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_S_ABS(e[submat-1]);
d[submat] -= MAGMA_S_ABS(e[submat-1]);
}
indxq = 4*n + 3;
// Solve each submatrix eigenproblem at the bottom of the divide and
// conquer tree.
magma_timer_t time=0;
timer_start( time );
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_ssteqr("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 *info;
}
k = 1;
for (j = submat; j < iwork[i]; ++j) {
iwork[indxq+j] = k;
++k;
}
}
timer_stop( time );
timer_printf( " for: ssteqr = %6.2f\n", time );
// Successively merge eigensystems of adjacent submatrices
// into eigensystem for the corresponding larger matrix.
curlvl = 1;
while (subpbs > 1) {
timer_start( time );
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.
// SLAEX1 is used only for the full eigensystem of a tridiagonal
// matrix.
if (matsiz == n)
range2 = range;
else
// We need all the eigenvectors if it is not last step
range2 = MagmaRangeAll;
magma_slaex1(matsiz, &d[submat], Q(submat, submat), ldq,
&iwork[indxq+submat], e[submat+msd2-1], msd2,
work, &iwork[subpbs], dwork,
range2, vl, vu, il, iu, info);
if (*info != 0) {
*info = (submat+1)*(n+1) + submat + matsiz;
return *info;
}
iwork[i/2]= iwork[i+1];
}
subpbs /= 2;
++curlvl;
timer_stop( time );
timer_printf("%d: time: %6.2f\n", (int) curlvl, time );
}
// 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_scopy(&n, Q(0, j), &ione, &work[ n*(i+1) ], &ione);
}
blasf77_scopy(&n, work, &ione, d, &ione);
lapackf77_slacpy( "A", &n, &n, &work[n], &n, Q, &ldq );
return *info;
} /* magma_slaex0 */
/**
Purpose
-------
SLAHR2 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 SGEHRD.
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 REAL 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 REAL array, dimension (NB)
The scalar factors of the elementary reflectors. See Further
Details.
@param[out]
T REAL 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 REAL 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_sgeev_aux
********************************************************************/
extern "C" magma_int_t
magma_slahr2_m(
magma_int_t n, magma_int_t k, magma_int_t nb,
float *A, magma_int_t lda,
float *tau,
float *T, magma_int_t ldt,
float *Y, magma_int_t ldy,
struct sgehrd_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)
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
float 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;
float scale;
magma_int_t i;
float ei = MAGMA_S_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 *info;
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
// zero out current top block of V on all GPUs
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magmablas_slaset( MagmaFull, nb, nb, c_zero, c_zero, dV(d,k,0), ldv, data->queues[d] );
}
// set all Y=0
lapackf77_slaset( "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_S_NEGATE( tau[i-1] );
blasf77_saxpy( &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_scopy( &i,
A(k+1,i), &ione,
T(0,nb-1), &ione );
// w := V1' * b1 = VA(k+1:k+i, 0:i-1)' * w
blasf77_strmv( "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_sgemv( "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_strmv( "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_sgemv( "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_strmv( "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_saxpy( &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_slarfg( &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 );
// dV(k+i+1:n-1, i) = VA(k+i:n, i)
magma_ssetvector_async( n_k_i_1,
A(k+i+1,i), 1,
dV(d, k+i+1, i), 1, data->queues[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_slacpy( MagmaFull, nb, nblocks-lblock,
dV (d, d*nb + lblock*nb*ngpu, i), nb*ngpu,
dVd(d, 0 + lblock*nb, i), nb, data->queues[d] );
// 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_sgemv( MagmaNoTrans, n-k, dn-dki1,
c_one, dA (d, k, dki1), ldda,
dVd(d, dki1, i), 1,
c_zero, dY (d, k, i), 1, data->queues[d] );
// copy vector to host, storing in column nb+d of Y
// as temporary space (Y has >= nb+ngpu columns)
magma_sgetvector_async( n-k,
dY(d, k, i), 1,
Y(k, nb+d), 1, data->queues[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_S_NEGATE( tau[i] );
blasf77_sgemv( "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_strmv( "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
#ifdef COMPLEX
lapackf77_slacgv( &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_sgemv( "No trans", &i, &i1,
&c_one, T(0,0), &ldt,
A(k+i1,0), &lda,
&c_zero, T(0,nb-1), &ione );
#ifdef COMPLEX
lapackf77_slacgv( &i1, A(k+i1,0), &lda );
#endif
// A(k:n, i+1) -= Y(k:n, 0:i) * w
blasf77_sgemv( "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->queues[d] );
magma_indices_1D_bcyclic( nb, ngpu, d, k+i+1, n, &dki1, &dn );
if ( dn-dki1 > 0 ) {
// yi = yi + yi{d}
blasf77_saxpy( &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 );
// 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_sgemm( 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, data->queues[d] );
// 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_sgetmatrix_async( k, nb,
dY(d, 0, 0), ldda,
Y(0,nb+nb*d), ldy, data->queues[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_saxpy( &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_ssetmatrix_async( n, nb, Y, ldy, dY(d, 0, 0), ldda, data->queues[d] );
magma_ssetmatrix_async( nb, nb, T, nb, dTi(d), nb, data->queues[d] );
}
magma_setdevice( orig_dev );
return *info;
} /* magma_slahr2 */
extern "C" magma_int_t
magma_cheevx(char jobz, char range, char uplo, magma_int_t n,
magmaFloatComplex *a, magma_int_t lda, float vl, float vu,
magma_int_t il, magma_int_t iu, float abstol, magma_int_t *m,
float *w, magmaFloatComplex *z, magma_int_t ldz, magmaFloatComplex *work, magma_int_t lwork,
float *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
=======
CHEEVX 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.
A (input/output) COMPLEX array, dimension (LDA, 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.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
VL (input) REAL
VU (input) REAL
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) REAL
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*SLAMCH('S'), not zero.
If this routine returns with INFO>0, indicating that some
eigenvectors did not converge, try setting ABSTOL to
2*SLAMCH('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) REAL array, dimension (N)
On normal exit, the first M elements contain the selected
eigenvalues in ascending order.
Z (output) COMPLEX array, dimension (LDZ, 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.
LDZ (input) INTEGER
The leading dimension of the array Z. LDZ >= 1, and if
JOBZ = 'V', LDZ >= max(1,N).
WORK (workspace/output) COMPLEX array, dimension (LWORK)
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The length of the array WORK. LWORK >= max(1,2*N-1).
For optimal efficiency, LWORK >= (NB+1)*N,
where NB is the max of the blocksize for CHETRD and for
CUNMTR as returned by ILAENV.
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) REAL 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 izero = 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;
float safmin;
float bignum;
float smlnum;
float eps, tmp1;
float anrm;
float sigma, d__1;
float rmin, rmax;
/* 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 (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_chetrd_nb(n);
lopt = n * (nb + 1);
work[0] = MAGMA_C_MAKE( lopt, 0 );
if (lwork < lopt && ! lquery) {
*info = -17;
}
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_cheevx(jobz_, range_, uplo_,
&n, a, &lda, &vl, &vu, &il, &iu, &abstol, m,
w, z, &ldz, work, &lwork,
rwork, iwork, ifail, info);
return *info;
}
--w;
--work;
--rwork;
--iwork;
--ifail;
/* Get machine constants. */
safmin = lapackf77_slamch("Safe minimum");
eps = lapackf77_slamch("Precision");
smlnum = safmin / eps;
bignum = 1. / smlnum;
rmin = magma_ssqrt(smlnum);
rmax = magma_ssqrt(bignum);
/* Scale matrix to allowable range, if necessary. */
anrm = lapackf77_clanhe("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_clascl(uplo_, &izero, &izero, &d__1, &sigma, &n, &n, a,
&lda, info);
if (abstol > 0.) {
abstol *= sigma;
}
if (valeig) {
vl *= sigma;
vu *= sigma;
}
}
/* Call CHETRD to reduce Hermitian matrix to tridiagonal form. */
indd = 1;
inde = indd + n;
indrwk = inde + n;
indtau = 1;
indwrk = indtau + n;
llwork = lwork - indwrk + 1;
magma_chetrd(uplo, n, a, lda, &rwork[indd], &rwork[inde], &work[indtau], &work[indwrk], llwork, &iinfo);
lopt = n + (magma_int_t)MAGMA_C_REAL(work[indwrk]);
/* If all eigenvalues are desired and ABSTOL is less than or equal to
zero, then call SSTERF or CUNGTR and CSTEQR. If this fails for
some eigenvalue, then try SSTEBZ. */
if ((alleig || (indeig && il == 1 && iu == n)) && abstol <= 0.) {
blasf77_scopy(&n, &rwork[indd], &ione, &w[1], &ione);
indee = indrwk + 2*n;
if (! wantz) {
i__1 = n - 1;
blasf77_scopy(&i__1, &rwork[inde], &ione, &rwork[indee], &ione);
lapackf77_ssterf(&n, &w[1], &rwork[indee], info);
}
else {
lapackf77_clacpy("A", &n, &n, a, &lda, z, &ldz);
lapackf77_cungtr(uplo_, &n, z, &ldz, &work[indtau], &work[indwrk], &llwork, &iinfo);
i__1 = n - 1;
blasf77_scopy(&i__1, &rwork[inde], &ione, &rwork[indee], &ione);
lapackf77_csteqr(jobz_, &n, &w[1], &rwork[indee], z, &ldz, &rwork[indrwk], info);
if (*info == 0) {
for (i = 1; i <= n; ++i) {
ifail[i] = 0;
}
}
}
if (*info == 0) {
*m = n;
}
}
/* Otherwise, call SSTEBZ and, if eigenvectors are desired, CSTEIN. */
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_sstebz(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_cstein(&n, &rwork[indd], &rwork[inde], m, &w[1], &iwork[indibl], &iwork[indisp],
z, &ldz, &rwork[indrwk], &iwork[indiwk], &ifail[1], info);
/* Apply unitary matrix used in reduction to tridiagonal
form to eigenvectors returned by CSTEIN. */
magma_cunmtr(MagmaLeft, uplo, MagmaNoTrans, n, *m, a, lda, &work[indtau],
z, ldz, &work[indwrk], 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_sscal(&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_cswap(&n, z + (i-1)*ldz, &ione, z + (j-1)*ldz, &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_C_MAKE( lopt, 0 );
return *info;
} /* magma_cheevx */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing strsm
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, cublas_perf, cublas_time, cpu_perf=0, cpu_time=0;
float cublas_error, normA, normx, normr, work[1];
magma_int_t N, info;
magma_int_t sizeA;
magma_int_t lda, ldda;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t *ipiv;
float *h_A, *h_b, *h_x, *h_xcublas;
float *d_A, *d_x;
float c_neg_one = MAGMA_S_NEG_ONE;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
printf("uplo = %s, transA = %s, diag = %s\n",
lapack_uplo_const(opts.uplo), lapack_trans_const(opts.transA), lapack_diag_const(opts.diag) );
printf(" N CUBLAS Gflop/s (ms) CPU Gflop/s (ms) CUBLAS error\n");
printf("============================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
gflops = FLOPS_STRSM(opts.side, N, 1) / 1e9;
lda = N;
ldda = ((lda+31)/32)*32;
sizeA = lda*N;
TESTING_MALLOC_CPU( ipiv, magma_int_t, N );
TESTING_MALLOC_CPU( h_A, float, lda*N );
TESTING_MALLOC_CPU( h_b, float, N );
TESTING_MALLOC_CPU( h_x, float, N );
TESTING_MALLOC_CPU( h_xcublas, float, N );
TESTING_MALLOC_DEV( d_A, float, ldda*N );
TESTING_MALLOC_DEV( d_x, float, N );
/* Initialize the matrices */
/* Factor A into LU to get well-conditioned triangular matrix.
* Copy L to U, since L seems okay when used with non-unit diagonal
* (i.e., from U), while U fails when used with unit diagonal. */
lapackf77_slarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_sgetrf( &N, &N, h_A, &lda, ipiv, &info );
for( int j = 0; j < N; ++j ) {
for( int i = 0; i < j; ++i ) {
*h_A(i,j) = *h_A(j,i);
}
}
lapackf77_slarnv( &ione, ISEED, &N, h_b );
blasf77_scopy( &N, h_b, &ione, h_x, &ione );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_ssetmatrix( N, N, h_A, lda, d_A, ldda );
magma_ssetvector( N, h_x, 1, d_x, 1 );
cublas_time = magma_sync_wtime( NULL );
cublasStrsv( handle, cublas_uplo_const(opts.uplo),
cublas_trans_const(opts.transA), cublas_diag_const(opts.diag),
N,
d_A, ldda,
d_x, 1 );
cublas_time = magma_sync_wtime( NULL ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_sgetvector( N, d_x, 1, h_xcublas, 1 );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
blasf77_strsv( lapack_uplo_const(opts.uplo), lapack_trans_const(opts.transA), lapack_diag_const(opts.diag),
&N,
h_A, &lda,
h_x, &ione );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
}
/* =====================================================================
Check the result
=================================================================== */
// ||b - Ax|| / (||A||*||x||)
// error for CUBLAS
normA = lapackf77_slange( "F", &N, &N, h_A, &lda, work );
normx = lapackf77_slange( "F", &N, &ione, h_xcublas, &ione, work );
blasf77_strmv( lapack_uplo_const(opts.uplo), lapack_trans_const(opts.transA), lapack_diag_const(opts.diag),
&N,
h_A, &lda,
h_xcublas, &ione );
blasf77_saxpy( &N, &c_neg_one, h_b, &ione, h_xcublas, &ione );
normr = lapackf77_slange( "F", &N, &ione, h_xcublas, &N, work );
cublas_error = normr / (normA*normx);
if ( opts.lapack ) {
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) N,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
cublas_error, (cublas_error < tol ? "ok" : "failed"));
status += ! (cublas_error < tol);
}
else {
printf("%5d %7.2f (%7.2f) --- ( --- ) %8.2e %s\n",
(int) N,
cublas_perf, 1000.*cublas_time,
cublas_error, (cublas_error < tol ? "ok" : "failed"));
status += ! (cublas_error < tol);
}
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_b );
TESTING_FREE_CPU( h_x );
TESTING_FREE_CPU( h_xcublas );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_x );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/**
Purpose
-------
SLAEX3 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 SLAED4 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]
ngpu INTEGER
Number of GPUs to use. ngpu > 0.
@param[in]
k INTEGER
The number of terms in the rational function to be solved by
SLAED4. 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 REAL array, dimension (N)
D(I) contains the updated eigenvalues for
1 <= I <= K.
@param[out]
Q REAL 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 REAL
The value of the parameter in the rank one update equation.
RHO >= 0 required.
@param[in,out]
dlamda REAL 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 REAL array, dimension (LDQ2, N)
The first K columns of this matrix contain the non-deflated
eigenvectors for the split problem.
@param[in]
indx INTEGER array, dimension (N)
The permutation used to arrange the columns of the deflated
Q matrix into three groups (see SLAED2).
The rows of the eigenvectors found by SLAED4 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 REAL 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) REAL 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 (devices workspaces) REAL 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(ngpu/2)) +
NB * ((N-N1) + (N-N1) / floor(ngpu/2))
@param
queues (device queues) 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.
TODO verify range, vl, vu, il, iu -- copied from slaex1.
@param[in]
vl REAL
@param[in]
vu REAL
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_ssyev_aux
********************************************************************/
extern "C" magma_int_t
magma_slaex3_m(
magma_int_t ngpu,
magma_int_t k, magma_int_t n, magma_int_t n1, float *d,
float *Q, magma_int_t ldq, float rho,
float *dlamda, float *Q2, magma_int_t *indx,
magma_int_t *ctot, float *w, float *s, magma_int_t *indxq,
magmaFloat_ptr dwork[],
magma_queue_t queues[MagmaMaxGPUs][2],
magma_range_t range, float vl, float vu, magma_int_t il, magma_int_t iu,
magma_int_t *info )
{
#define Q(i_,j_) (Q + (i_) + (j_)*ldq)
#define dQ2(id) (dwork[id])
#define dS(id, ii) (dwork[id] + n2*n2_loc + (ii)*(n2*nb))
#define dQ(id, ii) (dwork[id] + n2*n2_loc + 2*(n2*nb) + (ii)*(n2_loc*nb))
if (ngpu == 1) {
magma_setdevice(0);
magma_slaex3(k, n, n1, d, Q, ldq, rho,
dlamda, Q2, indx, ctot, w, s, indxq,
*dwork, range, vl, vu, il, iu, info );
return *info;
}
float d_one = 1.;
float d_zero = 0.;
magma_int_t ione = 1;
magma_int_t ineg_one = -1;
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;
float 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;
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
/*
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
float *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) / (ngpu/2) + 1;
n2_loc = (n2-1) / (ngpu/2) + 1;
nb = magma_get_slaex3_m_nb();
if (n1 >= magma_get_slaex3_m_k()) {
#ifdef CHECK_CPU
for (igpu = 0; igpu < ngpu; ++igpu) {
magma_smalloc_pinned( &(hwS[0][igpu]), n2*nb );
magma_smalloc_pinned( &(hwS[1][igpu]), n2*nb );
magma_smalloc_pinned( &(hwQ2[igpu]), n2*n2_loc );
magma_smalloc_pinned( &(hwQ[0][igpu]), n2_loc*nb );
magma_smalloc_pinned( &(hwQ[1][igpu]), n2_loc*nb );
}
#endif
for (igpu = 0; igpu < ngpu-1; igpu += 2) {
ni_loc[igpu] = min(n1_loc, n1 - igpu/2 * n1_loc);
#ifdef CHECK_CPU
lapackf77_slacpy("A", &ni_loc[igpu], &n12, Q2+n1_loc*(igpu/2), &n1, hQ2(igpu), &n1_loc);
#endif
magma_setdevice(igpu);
magma_ssetmatrix_async( ni_loc[igpu], n12,
Q2+n1_loc*(igpu/2), n1,
dQ2(igpu), n1_loc, queues[igpu][0] );
ni_loc[igpu+1] = min(n2_loc, n2 - igpu/2 * n2_loc);
#ifdef CHECK_CPU
lapackf77_slacpy("A", &ni_loc[igpu+1], &n23, Q2+iq2+n2_loc*(igpu/2), &n2, hQ2(igpu+1), &n2_loc);
#endif
magma_setdevice(igpu+1);
magma_ssetmatrix_async( ni_loc[igpu+1], n23,
Q2+iq2+n2_loc*(igpu/2), n2,
dQ2(igpu+1), n2_loc, queues[igpu+1][0] );
}
}
//
#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_slamc3(&dlamda[i], &dlamda[i]) - dlamda[i];
for (j = ib; j < ie; ++j) {
magma_int_t tmpp=j+1;
magma_int_t iinfo = 0;
lapackf77_slaed4(&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_slamrg( &k, &nk, d, &ione, &ineg_one, indxq);
//compute the lower and upper bound of the non-deflated eigenvectors
if (valeig)
magma_svrange(k, d, &iil, &iiu, vl, vu);
else if (indeig)
magma_sirange(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_scopy( &ik, &w[ib], &ione, &s[ib], &ione);
// Initialize W(I) = Q(I,I)
tmp = ldq + 1;
blasf77_scopy( &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_snrm2( 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_slamc3(&dlamda[i], &dlamda[i]) - dlamda[i];
for (j = 0; j < k; ++j) {
magma_int_t tmpp=j+1;
magma_int_t iinfo = 0;
lapackf77_slaed4(&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_slamrg( &k, &nk, d, &ione, &ineg_one, indxq);
//compute the lower and upper bound of the non-deflated eigenvectors
if (valeig)
magma_svrange(k, d, &iil, &iiu, vl, vu);
else if (indeig)
magma_sirange(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_scopy( &k, w, &ione, s, &ione);
// Initialize W(I) = Q(I,I)
tmp = ldq + 1;
blasf77_scopy( &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_snrm2( 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 );
if (rk > 0) {
if (n1 < magma_get_slaex3_m_k()) {
// stay on the CPU
if ( n23 != 0 ) {
lapackf77_slacpy("A", &n23, &rk, Q(ctot[0],iil-1), &ldq, s, &n23);
blasf77_sgemm("N", "N", &n2, &rk, &n23, &d_one, &Q2[iq2], &n2,
s, &n23, &d_zero, Q(n1,iil-1), &ldq );
}
else
lapackf77_slaset("A", &n2, &rk, &d_zero, &d_zero, Q(n1,iil-1), &ldq);
if ( n12 != 0 ) {
lapackf77_slacpy("A", &n12, &rk, Q(0,iil-1), &ldq, s, &n12);
blasf77_sgemm("N", "N", &n1, &rk, &n12, &d_one, Q2, &n1,
s, &n12, &d_zero, Q(0,iil-1), &ldq);
}
else
lapackf77_slaset("A", &n1, &rk, &d_zero, &d_zero, Q(0,iil-1), &ldq);
}
else {
//use the gpus
ib = min(nb, rk);
for (igpu = 0; igpu < ngpu-1; igpu += 2) {
if (n23 != 0) {
magma_setdevice(igpu+1);
magma_ssetmatrix_async( n23, ib,
Q(ctot[0],iil-1), ldq,
dS(igpu+1,0), n23, queues[igpu+1][0] );
}
if (n12 != 0) {
magma_setdevice(igpu);
magma_ssetmatrix_async( n12, ib,
Q(0,iil-1), ldq,
dS(igpu,0), n12, queues[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 < ngpu-1; igpu += 2) {
if (n23 != 0) {
magma_setdevice(igpu+1);
magma_ssetmatrix_async( n23, ib2,
Q(ctot[0],iil-1+i+nb), ldq,
dS(igpu+1,(ind+1)%2), n23, queues[igpu+1][(ind+1)%2] );
}
if (n12 != 0) {
magma_setdevice(igpu);
magma_ssetmatrix_async( n12, ib2,
Q(0,iil-1+i+nb), ldq,
dS(igpu,(ind+1)%2), n12, queues[igpu][(ind+1)%2] );
}
}
}
// Ensure that the data is copied on gpu since we will overwrite it.
for (igpu = 0; igpu < ngpu-1; igpu += 2) {
if (n23 != 0) {
#ifdef CHECK_CPU
lapackf77_slacpy("A", &n23, &ib, Q(ctot[0],iil-1+i), &ldq, hS(igpu+1,ind), &n23);
#endif
magma_setdevice(igpu+1);
magma_queue_sync( queues[igpu+1][ind] );
}
if (n12 != 0) {
#ifdef CHECK_CPU
lapackf77_slacpy("A", &n12, &ib, Q(0,iil-1+i), &ldq, hS(igpu,ind), &n12);
#endif
magma_setdevice(igpu);
magma_queue_sync( queues[igpu][ind] );
}
}
for (igpu = 0; igpu < ngpu-1; igpu += 2) {
if (n23 != 0) {
#ifdef CHECK_CPU
blasf77_sgemm("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(queues[igpu+1][ind]);
magma_sgemm(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_sdiff(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_sgemm("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(queues[igpu][ind]);
magma_sgemm(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_sdiff(ni_loc[igpu], ib, hQ(igpu, ind), n1_loc, dQ(igpu, ind), n1_loc));
#endif
}
}
for (igpu = 0; igpu < ngpu-1; igpu += 2) {
if (n23 != 0) {
magma_setdevice(igpu+1);
magma_sgetmatrix( ni_loc[igpu+1], ib, dQ(igpu+1, ind), n2_loc,
Q(n1+n2_loc*(igpu/2),iil-1+i), ldq );
// magma_sgetmatrix_async( ni_loc[igpu+1], ib, dQ(igpu+1, ind), n2_loc,
// Q(n1+n2_loc*(igpu/2),iil-1+i), ldq, queues[igpu+1][ind] );
}
if (n12 != 0) {
magma_setdevice(igpu);
magma_sgetmatrix( ni_loc[igpu], ib, dQ(igpu, ind), n1_loc,
Q(n1_loc*(igpu/2),iil-1+i), ldq );
// magma_sgetmatrix_async( ni_loc[igpu], ib, dQ(igpu, ind), n1_loc,
// Q(n1_loc*(igpu/2),iil-1+i), ldq, queues[igpu][ind] );
}
}
}
for (igpu = 0; igpu < ngpu; ++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);
magma_queue_sync( queues[igpu][0] );
magma_queue_sync( queues[igpu][1] );
}
if ( n23 == 0 )
lapackf77_slaset("A", &n2, &rk, &d_zero, &d_zero, Q(n1,iil-1), &ldq);
if ( n12 == 0 )
lapackf77_slaset("A", &n1, &rk, &d_zero, &d_zero, Q(0,iil-1), &ldq);
}
}
timer_stop( time );
timer_printf( "gemms = %6.2f\n", time );
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
return *info;
} /* magma_slaed3_m */
extern "C" magma_int_t
magma_slahr2(
magma_int_t n, magma_int_t k, magma_int_t nb,
magmaFloat_ptr da, size_t da_offset, magma_int_t ldda,
magmaFloat_ptr dv, size_t dv_offset, magma_int_t lddv,
float *a, magma_int_t lda,
float *tau,
float *t, magma_int_t ldt,
float *y, magma_int_t ldy,
magma_queue_t queue)
{
/* -- clMAGMA auxiliary routine (version 0.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date November 2014
Purpose
=======
SLAHR2 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 SGEHRD.
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) REAL 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) REAL 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) REAL array, dimension (NB)
The scalar factors of the elementary reflectors. See Further
Details.
T (output) REAL array, dimension (LDT,NB)
The upper triangular matrix T.
LDT (input) INTEGER
The leading dimension of the array T. LDT >= NB.
Y (output) REAL 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.
===================================================================== */
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
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;
float d__1;
magma_int_t i__;
float 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;
if (n <= 1)
return MAGMA_SUCCESS;
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_slacgv(&i__3, &a[k+i__-1+a_dim1], &lda);
#endif
blasf77_scopy(&i__3, &a[k+i__-1+a_dim1], &lda, &t[nb*t_dim1+1], &c__1);
blasf77_strmv("u","n","n",&i__3,&t[t_offset], &ldt, &t[nb*t_dim1+1], &c__1);
blasf77_sgemv("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_slacgv(&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_scopy(&i__2, &a[k+1+i__*a_dim1], &c__1, &t[nb*t_dim1+1], &c__1);
blasf77_strmv("Lower", MagmaConjTransStr, "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_sgemv(MagmaConjTransStr, &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_strmv("U", MagmaConjTransStr, "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_sgemv("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_strmv("L","N","U",&i__2,&a[k+1+a_dim1],&lda,&t[nb*t_dim1+1],&c__1);
blasf77_saxpy(&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_slarfg(&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_ssetvector( i__3, &a[k + i__ + i__*a_dim1], 1, dv, dv_offset+(i__-1)*(lddv+1), 1, queue );
magma_sgemv(MagmaNoTrans, i__2+1, i__3, c_one,
da, da_offset + (-1 + k + i__ * ldda), ldda,
dv, dv_offset + (i__-1)*(lddv+1), c__1, c_zero,
da, da_offset + (-1 + k + (i__-1)*ldda), c__1, queue);
i__2 = n - k - i__ + 1;
i__3 = i__ - 1;
blasf77_sgemv(MagmaConjTransStr, &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_S_NEGATE( tau[i__] );
blasf77_sscal(&i__2, &d__1, &t[i__ * t_dim1 + 1], &c__1);
blasf77_strmv("U","N","N", &i__2, &t[t_offset], &ldt, &t[i__*t_dim1+1], &c__1);
t[i__ + i__ * t_dim1] = tau[i__];
magma_sgetvector( n - k + 1, da, da_offset+(-1+ k+(i__-1)*ldda), 1, y+ k + i__*y_dim1, 1, queue );
}
a[k + nb + nb * a_dim1] = ei;
return MAGMA_SUCCESS;
} /* magma_slahr2 */
extern "C" magma_int_t
magma_slaex0(magma_int_t n, float* d, float* e, float* q, magma_int_t ldq,
float* work, magma_int_t* iwork, magmaFloat_ptr dwork,
magma_vec_t range, float vl, float 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
REAL VL, VU
..
.. Array Arguments ..
INTEGER IWORK( * )
REAL D( * ), E( * ), Q( LDQ, * ),
$ WORK( * ), DWORK( * )
..
Purpose
=======
SLAEX0 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) REAL array, dimension (N)
On entry, the main diagonal of the tridiagonal matrix.
On exit, its eigenvalues.
E (input) REAL array, dimension (N-1)
The off-diagonal elements of the tridiagonal matrix.
On exit, E has been destroyed.
Q (input/output) REAL 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) REAL 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) REAL 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) REAL
VU (input) REAL
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_slaex0_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_S_ABS(e[submat-1]);
d[submat] -= MAGMA_S_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_ssteqr(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: ssteqr = %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.
// SLAEX1 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_slaex1(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_scopy(&n, Q(0, j), &ione, &work[ n*(i+1) ], &ione);
}
blasf77_scopy(&n, work, &ione, d, &ione);
char char_A[] = {'A',0};
lapackf77_slacpy ( char_A, &n, &n, &work[n], &n, q, &ldq );
return MAGMA_SUCCESS;
} /* magma_slaex0 */
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float *h_x, *h_x1, *h_x2, *h_tau;
float *d_x, *d_tau;
float c_neg_one = MAGMA_S_NEG_ONE;
float error, work[1];
magma_int_t N, size, nb;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
parse_opts( argc, argv, &opts );
// does larfg on nb columns, one after another
nb = (opts.nb > 0 ? opts.nb : 64);
magma_queue_t queue = 0;
printf(" N nb CPU GFLop/s (ms) GPU GFlop/s (ms) error \n");
printf("==============================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[i];
gflops = FLOPS_SLARFG( N ) / 1e9 * nb;
TESTING_MALLOC( h_x, float, N*nb );
TESTING_MALLOC( h_x1, float, N*nb );
TESTING_MALLOC( h_x2, float, N*nb );
TESTING_MALLOC( h_tau, float, nb );
TESTING_DEVALLOC( d_x, float, N*nb );
TESTING_DEVALLOC( d_tau, float, nb );
/* Initialize the vector */
size = N*nb;
lapackf77_slarnv( &ione, ISEED, &size, h_x );
blasf77_scopy( &size, h_x, &ione, h_x1, &ione );
/* =====================================================================
Performs operation using MAGMA-BLAS
=================================================================== */
magma_ssetvector( size, h_x, ione, d_x, ione );
gpu_time = magma_sync_wtime( queue );
for( int j = 0; j < nb; ++j ) {
magma_slarfg( N, &d_x[0+j*N], &d_x[1+j*N], ione, &d_tau[j] );
}
gpu_time = magma_sync_wtime( queue ) - gpu_time;
gpu_perf = gflops / gpu_time;
magma_sgetvector( size, d_x, ione, h_x2, ione );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
for( int j = 0; j < nb; ++j ) {
lapackf77_slarfg( &N, &h_x1[0+j*N], &h_x1[1+j*N], &ione, &h_tau[j] );
}
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/* =====================================================================
Error Computation and Performance Compariosn
=================================================================== */
blasf77_saxpy( &size, &c_neg_one, h_x1, &ione, h_x2, &ione);
error = lapackf77_slange( "F", &N, &nb, h_x2, &N, work );
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2g\n",
(int) N, (int) nb, cpu_perf, 1000.*cpu_time, gpu_perf, 1000.*gpu_time, error );
TESTING_FREE( h_x );
TESTING_FREE( h_x1 );
TESTING_FREE( h_x2 );
TESTING_FREE( h_tau );
TESTING_DEVFREE( d_x );
TESTING_DEVFREE( d_tau );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
