extern "C" magma_int_t
magma_slarfb2_gpu( magma_int_t m, magma_int_t n, magma_int_t k,
const float *dV, magma_int_t ldv,
const float *dT, magma_int_t ldt,
float *dC, magma_int_t ldc,
float *dwork, magma_int_t ldwork )
{
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
if (m <= 0 || n <= 0)
return MAGMA_SUCCESS;
// W = C^H V
// magma_sgemm( MagmaTrans, MagmaNoTrans,
magmablas_sgemm_reduce(
n, k, m,
c_one, dC, ldc,
dV, ldv,
c_zero, dwork, ldwork);
// W = W T^H = C^H V T^H
magma_strmm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
n, k,
c_one, dT, ldt,
dwork, ldwork);
// C = C - V W^H = C - V T V^H C = (I - V T V^H) C = H C
magma_sgemm( MagmaNoTrans, MagmaTrans,
m, n, k,
c_neg_one, dV, ldv,
dwork, ldwork,
c_one, dC, ldc);
return MAGMA_SUCCESS;
}
extern "C" magma_int_t
magma_ssygvd(magma_int_t itype, char jobz, char uplo, magma_int_t n,
float *a, magma_int_t lda, float *b, magma_int_t ldb,
float *w, float *work, magma_int_t lwork,
magma_int_t *iwork, magma_int_t liwork, magma_int_t *info)
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
Purpose
=======
SSYGVD computes all the eigenvalues, and optionally, the eigenvectors
of a real generalized symmetric-definite eigenproblem, of the form
A*x=(lambda)*B*x, A*Bx=(lambda)*x, or B*A*x=(lambda)*x. Here A and
B are assumed to be symmetric and B is also positive definite.
If eigenvectors are desired, it uses a divide and conquer algorithm.
The divide and conquer algorithm 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
=========
ITYPE (input) INTEGER
Specifies the problem type to be solved:
= 1: A*x = (lambda)*B*x
= 2: A*B*x = (lambda)*x
= 3: B*A*x = (lambda)*x
JOBZ (input) CHARACTER*1
= 'N': Compute eigenvalues only;
= 'V': Compute eigenvalues and eigenvectors.
UPLO (input) CHARACTER*1
= 'U': Upper triangles of A and B are stored;
= 'L': Lower triangles of A and B are stored.
N (input) INTEGER
The order of the matrices A and B. N >= 0.
A (input/output) COMPLEX*16 array, dimension (LDA, N)
On entry, the symmetric 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, if JOBZ = 'V', then if INFO = 0, A contains the
matrix Z of eigenvectors. The eigenvectors are normalized
as follows:
if ITYPE = 1 or 2, Z**T * B * Z = I;
if ITYPE = 3, Z**T * inv(B) * Z = I.
If JOBZ = 'N', then on exit the upper triangle (if UPLO='U')
or the lower triangle (if UPLO='L') of A, including the
diagonal, is destroyed.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
B (input/output) COMPLEX*16 array, dimension (LDB, N)
On entry, the symmetric matrix B. If UPLO = 'U', the
leading N-by-N upper triangular part of B contains the
upper triangular part of the matrix B. If UPLO = 'L',
the leading N-by-N lower triangular part of B contains
the lower triangular part of the matrix B.
On exit, if INFO <= N, the part of B containing the matrix is
overwritten by the triangular factor U or L from the Cholesky
factorization B = U**T * U or B = L * L**T.
LDB (input) INTEGER
The leading dimension of the array B. LDB >= max(1,N).
W (output) DOUBLE PRECISION array, dimension (N)
If INFO = 0, the eigenvalues in ascending order.
WORK (workspace/output) COMPLEX*16 array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The length of the array WORK.
If N <= 1, LWORK >= 1.
If JOBZ = 'N' and N > 1, LWORK >= 2*N*nb + 1.
If JOBZ = 'V' and N > 1, LWORK >= 1 + 6*N*nb + 2*N**2.
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal sizes of the WORK and
IWORK arrays, returns these values as the first entries of
the WORK and IWORK arrays, and no error message
related to LWORK or LIWORK is issued by XERBLA.
IWORK (workspace/output) INTEGER array, dimension (MAX(1,LIWORK))
On exit, if INFO = 0, IWORK(1) returns the optimal LIWORK.
LIWORK (input) INTEGER
The dimension of the array IWORK.
If N <= 1, LIWORK >= 1.
If JOBZ = 'N' and N > 1, LIWORK >= 1.
If JOBZ = 'V' and N > 1, LIWORK >= 3 + 5*N.
If LIWORK = -1, then a workspace query is assumed; the
routine only calculates the optimal sizes of the WORK
and IWORK arrays, returns these values as the first entries
of the WORK and IWORK arrays, and no error message
related to LWORK or LIWORK is issued by XERBLA.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
> 0: SPOTRF or SSYEVD returned an error code:
<= N: if INFO = i and JOBZ = 'N', then the algorithm
failed to converge; i off-diagonal elements of an
intermediate tridiagonal form did not converge to
zero;
if INFO = i and JOBZ = 'V', then 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);
> N: if INFO = N + i, for 1 <= i <= N, then the leading
minor of order i of B is not positive definite.
The factorization of B could not be completed and
no eigenvalues or eigenvectors were computed.
Further Details
===============
Based on contributions by
Mark Fahey, Department of Mathematics, Univ. of Kentucky, USA
Modified so that no backsubstitution is performed if SSYEVD fails to
converge (NEIG in old code could be greater than N causing out of
bounds reference to A - reported by Ralf Meyer). Also corrected the
description of INFO and the test on ITYPE. Sven, 16 Feb 05.
===================================================================== */
char uplo_[2] = {uplo, 0};
char jobz_[2] = {jobz, 0};
float d_one = MAGMA_S_ONE;
float *da;
float *db;
magma_int_t ldda = n;
magma_int_t lddb = n;
magma_int_t lower;
char trans[1];
magma_int_t wantz, lquery;
magma_int_t lopt, lwmin, liopt, liwmin;
cudaStream_t stream;
magma_queue_create( &stream );
wantz = lapackf77_lsame(jobz_, MagmaVectorsStr);
lower = lapackf77_lsame(uplo_, MagmaLowerStr);
lquery = lwork == -1 || liwork == -1;
*info = 0;
if (itype < 1 || itype > 3) {
*info = -1;
} else if (! (wantz || lapackf77_lsame(jobz_, MagmaNoVectorsStr))) {
*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 (ldb < max(1,n)) {
*info = -8;
}
magma_int_t nb = magma_get_ssytrd_nb(n);
if (n < 1) {
liwmin = 1;
lwmin = 1;
} else if (wantz) {
lwmin = 1 + 6 * n * nb + 2* n * n;
liwmin = 5 * n + 3;
} else {
lwmin = 2 * n * nb + 1;
liwmin = 1;
}
lopt = lwmin;
liopt = liwmin;
work[ 0] = lopt;
iwork[0] = liopt;
if (lwork < lwmin && ! lquery) {
*info = -11;
} else if (liwork < liwmin && ! lquery) {
*info = -13;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return MAGMA_ERR_ILLEGAL_VALUE;
}
else if (lquery) {
return MAGMA_SUCCESS;
}
/* Quick return if possible */
if (n == 0) {
return 0;
}
if (MAGMA_SUCCESS != magma_smalloc( &da, n*ldda ) ||
MAGMA_SUCCESS != magma_smalloc( &db, n*lddb )) {
*info = -17;
return MAGMA_ERR_DEVICE_ALLOC;
}
/* Form a Cholesky factorization of B. */
magma_ssetmatrix( n, n, b, ldb, db, lddb );
magma_ssetmatrix_async( n, n,
a, lda,
da, ldda, stream );
magma_spotrf_gpu(uplo_[0], n, db, lddb, info);
if (*info != 0) {
*info = n + *info;
return 0;
}
magma_queue_sync( stream );
magma_sgetmatrix_async( n, n,
db, lddb,
b, ldb, stream );
/* Transform problem to standard eigenvalue problem and solve. */
magma_ssygst_gpu(itype, uplo_[0], n, da, ldda, db, lddb, info);
magma_ssyevd_gpu(jobz_[0], uplo_[0], n, da, ldda, w, a, lda,
work, lwork, iwork, liwork, info);
lopt = max( lopt, (magma_int_t) work[0]);
liopt = max(liopt, iwork[0]);
if (wantz && *info == 0)
{
/* Backtransform eigenvectors to the original problem. */
if (itype == 1 || itype == 2)
{
/* For A*x=(lambda)*B*x and A*B*x=(lambda)*x;
backtransform eigenvectors: x = inv(L)'*y or inv(U)*y */
if (lower) {
*(unsigned char *)trans = MagmaTrans;
} else {
*(unsigned char *)trans = MagmaNoTrans;
}
magma_strsm(MagmaLeft, uplo_[0], *trans, MagmaNonUnit,
n, n, d_one, db, lddb, da, ldda);
} else if (itype == 3)
{
/* For B*A*x=(lambda)*x;
backtransform eigenvectors: x = L*y or U'*y */
if (lower) {
*(unsigned char *)trans = MagmaNoTrans;
} else {
*(unsigned char *)trans = MagmaTrans;
}
magma_strmm(MagmaLeft, uplo_[0], *trans, MagmaNonUnit,
n, n, d_one, db, lddb, da, ldda);
}
magma_sgetmatrix( n, n, da, ldda, a, lda );
}
magma_queue_sync( stream );
magma_queue_destroy( stream );
work[0] = (float) lopt;
iwork[0] = liopt;
magma_free( da );
magma_free( db );
return MAGMA_SUCCESS;
} /* magma_ssygvd */
/**
Purpose
-------
SSYGVD computes all the eigenvalues, and optionally, the eigenvectors
of a real generalized symmetric-definite eigenproblem, of the form
A*x=(lambda)*B*x, A*Bx=(lambda)*x, or B*A*x=(lambda)*x. Here A and
B are assumed to be symmetric and B is also positive definite.
If eigenvectors are desired, it uses a divide and conquer algorithm.
The divide and conquer algorithm 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]
itype INTEGER
Specifies the problem type to be solved:
= 1: A*x = (lambda)*B*x
= 2: A*B*x = (lambda)*x
= 3: B*A*x = (lambda)*x
@param[in]
jobz magma_vec_t
- = MagmaNoVec: Compute eigenvalues only;
- = MagmaVec: Compute eigenvalues and eigenvectors.
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangles of A and B are stored;
- = MagmaLower: Lower triangles of A and B are stored.
@param[in]
n INTEGER
The order of the matrices A and B. N >= 0.
@param[in,out]
A REAL array, dimension (LDA, N)
On entry, the symmetric matrix A. If UPLO = MagmaUpper, the
leading N-by-N upper triangular part of A contains the
upper triangular part of the matrix A. If UPLO = MagmaLower,
the leading N-by-N lower triangular part of A contains
the lower triangular part of the matrix A.
\n
On exit, if JOBZ = MagmaVec, then if INFO = 0, A contains the
matrix Z of eigenvectors. The eigenvectors are normalized
as follows:
if ITYPE = 1 or 2, Z**T * B * Z = I;
if ITYPE = 3, Z**T * inv(B) * Z = I.
If JOBZ = MagmaNoVec, then on exit the upper triangle (if UPLO=MagmaUpper)
or the lower triangle (if UPLO=MagmaLower) of A, including the
diagonal, is destroyed.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[in,out]
B REAL array, dimension (LDB, N)
On entry, the symmetric matrix B. If UPLO = MagmaUpper, the
leading N-by-N upper triangular part of B contains the
upper triangular part of the matrix B. If UPLO = MagmaLower,
the leading N-by-N lower triangular part of B contains
the lower triangular part of the matrix B.
\n
On exit, if INFO <= N, the part of B containing the matrix is
overwritten by the triangular factor U or L from the Cholesky
factorization B = U**T * U or B = L * L**T.
@param[in]
ldb INTEGER
The leading dimension of the array B. LDB >= max(1,N).
@param[out]
w REAL array, dimension (N)
If INFO = 0, the eigenvalues in ascending order.
@param[out]
work (workspace) REAL array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The length of the array WORK.
If N <= 1, LWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LWORK >= 2*N + N*NB.
If JOBZ = MagmaVec and N > 1, LWORK >= max( 2*N + N*NB, 1 + 6*N + 2*N**2 ).
NB can be obtained through magma_get_ssytrd_nb(N).
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal sizes of the WORK and IWORK
arrays, returns these values as the first entries of the WORK
and IWORK arrays, and no error message related to LWORK or
LIWORK is issued by XERBLA.
@param[out]
iwork (workspace) INTEGER array, dimension (MAX(1,LIWORK))
On exit, if INFO = 0, IWORK[0] returns the optimal LIWORK.
@param[in]
liwork INTEGER
The dimension of the array IWORK.
If N <= 1, LIWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LIWORK >= 1.
If JOBZ = MagmaVec and N > 1, LIWORK >= 3 + 5*N.
\n
If LIWORK = -1, then a workspace query is assumed; the
routine only calculates the optimal sizes of the WORK and
IWORK arrays, returns these values as the first entries of
the WORK and IWORK arrays, and no error message related to
LWORK or LIWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: SPOTRF or SSYEVD returned an error code:
<= N: if INFO = i and JOBZ = MagmaNoVec, then the algorithm
failed to converge; i off-diagonal elements of an
intermediate tridiagonal form did not converge to
zero;
if INFO = i and JOBZ = MagmaVec, then 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);
> N: if INFO = N + i, for 1 <= i <= N, then the leading
minor of order i of B is not positive definite.
The factorization of B could not be completed and
no eigenvalues or eigenvectors were computed.
Further Details
---------------
Based on contributions by
Mark Fahey, Department of Mathematics, Univ. of Kentucky, USA
Modified so that no backsubstitution is performed if SSYEVD fails to
converge (NEIG in old code could be greater than N causing out of
bounds reference to A - reported by Ralf Meyer). Also corrected the
description of INFO and the test on ITYPE. Sven, 16 Feb 05.
@ingroup magma_ssygv_driver
********************************************************************/
extern "C" magma_int_t
magma_ssygvd_m(
magma_int_t ngpu,
magma_int_t itype, magma_vec_t jobz, magma_uplo_t uplo, magma_int_t n,
float *A, magma_int_t lda,
float *B, magma_int_t ldb,
float *w, float *work, magma_int_t lwork,
#ifdef COMPLEX
float *rwork, magma_int_t lrwork,
#endif
magma_int_t *iwork, magma_int_t liwork,
magma_int_t *info)
{
const char* uplo_ = lapack_uplo_const( uplo );
const char* jobz_ = lapack_vec_const( jobz );
float d_one = MAGMA_S_ONE;
magma_int_t lower;
magma_trans_t trans;
magma_int_t wantz, lquery;
magma_int_t lwmin, liwmin;
magma_queue_t stream;
magma_queue_create( &stream );
wantz = (jobz == MagmaVec);
lower = (uplo == MagmaLower);
lquery = (lwork == -1 || liwork == -1);
*info = 0;
if (itype < 1 || itype > 3) {
*info = -1;
} else if (! (wantz || (jobz == MagmaNoVec))) {
*info = -2;
} else if (! (lower || (uplo == MagmaUpper))) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (lda < max(1,n)) {
*info = -6;
} else if (ldb < max(1,n)) {
*info = -8;
}
magma_int_t nb = magma_get_ssytrd_nb( n );
if ( n <= 1 ) {
lwmin = 1;
liwmin = 1;
}
else if ( wantz ) {
lwmin = max( 2*n + n*nb, 1 + 6*n + 2*n*n );
liwmin = 3 + 5*n;
}
else {
lwmin = 2*n + n*nb;
liwmin = 1;
}
// multiply by 1+eps (in Double!) to ensure length gets rounded up,
// if it cannot be exactly represented in floating point.
real_Double_t one_eps = 1. + lapackf77_slamch("Epsilon");
work[0] = lwmin * one_eps;
iwork[0] = liwmin;
if (lwork < lwmin && ! lquery) {
*info = -11;
} else if (liwork < liwmin && ! lquery) {
*info = -13;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (n == 0) {
return *info;
}
/* If matrix is very small, then just call LAPACK on CPU, no need for GPU */
if (n <= 128) {
lapackf77_ssygvd( &itype, jobz_, uplo_,
&n, A, &lda, B, &ldb,
w, work, &lwork,
iwork, &liwork, info );
return *info;
}
magma_timer_t time=0;
timer_start( time );
magma_spotrf_m( ngpu, uplo, n, B, ldb, info );
if (*info != 0) {
*info = n + *info;
return *info;
}
timer_stop( time );
timer_printf( "time spotrf = %6.2f\n", time );
timer_start( time );
/* Transform problem to standard eigenvalue problem and solve. */
magma_ssygst_m( ngpu, itype, uplo, n, A, lda, B, ldb, info );
timer_stop( time );
timer_printf( "time ssygst = %6.2f\n", time );
timer_start( time );
magma_ssyevd_m( ngpu, jobz, uplo, n, A, lda, w, work, lwork, iwork, liwork, info );
timer_stop( time );
timer_printf( "time ssyevd = %6.2f\n", time );
if (wantz && *info == 0) {
timer_start( time );
/* Backtransform eigenvectors to the original problem. */
if (itype == 1 || itype == 2) {
/* For A*x=(lambda)*B*x and A*B*x=(lambda)*x;
backtransform eigenvectors: x = inv(L)'*y or inv(U)*y */
if (lower) {
trans = MagmaTrans;
} else {
trans = MagmaNoTrans;
}
magma_strsm_m( ngpu, MagmaLeft, uplo, trans, MagmaNonUnit,
n, n, d_one, B, ldb, A, lda );
}
else if (itype == 3) {
/* For B*A*x=(lambda)*x;
backtransform eigenvectors: x = L*y or U'*y */
if (lower) {
trans = MagmaNoTrans;
} else {
trans = MagmaTrans;
}
printf("--- the multi GPU version is falling back to 1 GPU to perform the last TRMM since there is no TRMM_mgpu --- \n");
float *dA=NULL, *dB=NULL;
magma_int_t ldda = roundup( n, 32 );
magma_int_t lddb = ldda;
if (MAGMA_SUCCESS != magma_smalloc( &dA, n*ldda ) ||
MAGMA_SUCCESS != magma_smalloc( &dB, n*lddb ) ) {
magma_free( dA );
magma_free( dB );
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_ssetmatrix( n, n, B, ldb, dB, lddb );
magma_ssetmatrix( n, n, A, lda, dA, ldda );
magma_strmm( MagmaLeft, uplo, trans, MagmaNonUnit,
n, n, d_one, dB, lddb, dA, ldda );
magma_sgetmatrix( n, n, dA, ldda, A, lda );
magma_free( dA );
magma_free( dB );
}
timer_stop( time );
timer_printf( "time setmatrices trsm/mm + getmatrices = %6.2f\n", time );
}
work[0] = lwmin * one_eps; // round up
iwork[0] = liwmin;
return *info;
} /* magma_ssygvd_m */
/**
Purpose
-------
SLARFB applies a real block reflector H or its transpose H^H to a
REAL m by n matrix C, from the left.
Arguments
---------
@param[in]
side magma_side_t
- = MagmaLeft: apply H or H^H from the Left
- = MagmaRight: apply H or H^H from the Right
@param[in]
trans magma_trans_t
- = MagmaNoTrans: apply H (No transpose)
- = MagmaTrans: apply H^H (Conjugate transpose)
@param[in]
direct magma_direct_t
Indicates how H is formed from a product of elementary
reflectors
- = MagmaForward: H = H(1) H(2) . . . H(k) (Forward)
- = MagmaBackward: H = H(k) . . . H(2) H(1) (Backward)
@param[in]
storev magma_storev_t
Indicates how the vectors which define the elementary
reflectors are stored:
- = MagmaColumnwise: Columnwise
- = MagmaRowwise: Rowwise
@param[in]
m INTEGER
The number of rows of the matrix C.
@param[in]
n INTEGER
The number of columns of the matrix C.
@param[in]
k INTEGER
The order of the matrix T (= the number of elementary
reflectors whose product defines the block reflector).
@param[in]
dV REAL array on the GPU, dimension
(LDDV,K) if STOREV = MagmaColumnwise
(LDDV,M) if STOREV = MagmaRowwise and SIDE = MagmaLeft
(LDDV,N) if STOREV = MagmaRowwise and SIDE = MagmaRight
The matrix V. See further details.
@param[in]
lddv INTEGER
The leading dimension of the array V.
If STOREV = MagmaColumnwise and SIDE = MagmaLeft, LDDV >= max(1,M);
if STOREV = MagmaColumnwise and SIDE = MagmaRight, LDDV >= max(1,N);
if STOREV = MagmaRowwise, LDDV >= K.
@param[in]
dT REAL array on the GPU, dimension (LDDT,K)
The triangular k by k matrix T in the representation of the
block reflector.
@param[in]
lddt INTEGER
The leading dimension of the array T. LDDT >= K.
@param[in,out]
dC REAL array on the GPU, dimension (LDDC,N)
On entry, the m by n matrix C.
On exit, C is overwritten by H*C, or H^H*C, or C*H, or C*H^H.
@param[in]
lddc INTEGER
The leading dimension of the array C. LDA >= max(1,M).
@param
dwork (workspace) REAL array, dimension (LDWORK,K)
@param[in]
ldwork INTEGER
The leading dimension of the array WORK.
If SIDE = MagmaLeft, LDWORK >= max(1,N);
if SIDE = MagmaRight, LDWORK >= max(1,M);
Further Details
---------------
The shape of the matrix V and the storage of the vectors which define
the H(i) is best illustrated by the following example with n = 5 and
k = 3.
All elements including 0's and 1's are stored, unlike LAPACK.
DIRECT = MagmaForward and DIRECT = MagmaForward and
STOREV = MagmaColumnwise: STOREV = MagmaRowwise:
V = ( 1 0 0 ) V = ( 1 v1 v1 v1 v1 )
( v1 1 0 ) ( 0 1 v2 v2 v2 )
( v1 v2 1 ) ( 0 0 1 v3 v3 )
( v1 v2 v3 )
( v1 v2 v3 )
DIRECT = MagmaBackward and DIRECT = MagmaBackward and
STOREV = MagmaColumnwise: STOREV = MagmaRowwise:
V = ( v1 v2 v3 ) V = ( v1 v1 1 0 0 )
( v1 v2 v3 ) ( v2 v2 v2 1 0 )
( 1 v2 v3 ) ( v3 v3 v3 v3 1 )
( 0 1 v3 )
( 0 0 1 )
@ingroup magma_saux3
********************************************************************/
extern "C" magma_int_t
magma_slarfb_gpu(
magma_side_t side, magma_trans_t trans, magma_direct_t direct, magma_storev_t storev,
magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloat_const_ptr dV, magma_int_t lddv,
magmaFloat_const_ptr dT, magma_int_t lddt,
magmaFloat_ptr dC, magma_int_t lddc,
magmaFloat_ptr dwork, magma_int_t ldwork )
{
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
/* Check input arguments */
magma_int_t info = 0;
if (m < 0) {
info = -5;
} else if (n < 0) {
info = -6;
} else if (k < 0) {
info = -7;
} else if ( ((storev == MagmaColumnwise) && (side == MagmaLeft) && lddv < max(1,m)) ||
((storev == MagmaColumnwise) && (side == MagmaRight) && lddv < max(1,n)) ||
((storev == MagmaRowwise) && lddv < k) ) {
info = -9;
} else if (lddt < k) {
info = -11;
} else if (lddc < max(1,m)) {
info = -13;
} else if ( ((side == MagmaLeft) && ldwork < max(1,n)) ||
((side == MagmaRight) && ldwork < max(1,m)) ) {
info = -15;
}
if (info != 0) {
magma_xerbla( __func__, -(info) );
return info;
}
/* Function Body */
if (m <= 0 || n <= 0) {
return info;
}
// opposite of trans
magma_trans_t transt;
if (trans == MagmaNoTrans)
transt = MagmaTrans;
else
transt = MagmaNoTrans;
// whether T is upper or lower triangular
magma_uplo_t uplo;
if (direct == MagmaForward)
uplo = MagmaUpper;
else
uplo = MagmaLower;
// whether V is stored transposed or not
magma_trans_t notransV, transV;
if (storev == MagmaColumnwise) {
notransV = MagmaNoTrans;
transV = MagmaTrans;
}
else {
notransV = MagmaTrans;
transV = MagmaNoTrans;
}
if ( side == MagmaLeft ) {
// Form H C or H^H C
// Comments assume H C. When forming H^H C, T gets transposed via transt.
// W = C^H V
magma_sgemm( MagmaTrans, notransV,
n, k, m,
c_one, dC, lddc,
dV, lddv,
c_zero, dwork, ldwork);
// W = W T^H = C^H V T^H
magma_strmm( MagmaRight, uplo, transt, MagmaNonUnit,
n, k,
c_one, dT, lddt,
dwork, ldwork);
// C = C - V W^H = C - V T V^H C = (I - V T V^H) C = H C
magma_sgemm( notransV, MagmaTrans,
m, n, k,
c_neg_one, dV, lddv,
dwork, ldwork,
c_one, dC, lddc);
}
else {
// Form C H or C H^H
// Comments assume C H. When forming C H^H, T gets transposed via trans.
// W = C V
magma_sgemm( MagmaNoTrans, notransV,
m, k, n,
c_one, dC, lddc,
dV, lddv,
c_zero, dwork, ldwork);
// W = W T = C V T
magma_strmm( MagmaRight, uplo, trans, MagmaNonUnit,
m, k,
c_one, dT, lddt,
dwork, ldwork);
// C = C - W V^H = C - C V T V^H = C (I - V T V^H) = C H
magma_sgemm( MagmaNoTrans, transV,
m, n, k,
c_neg_one, dwork, ldwork,
dV, lddv,
c_one, dC, lddc);
}
return info;
} /* magma_slarfb */
/**
Purpose
-------
SSYGVD computes all the eigenvalues, and optionally, the eigenvectors
of a real generalized symmetric-definite eigenproblem, of the form
A*x=(lambda)*B*x, A*Bx=(lambda)*x, or B*A*x=(lambda)*x. Here A and
B are assumed to be symmetric and B is also positive definite.
If eigenvectors are desired, it uses a divide and conquer algorithm.
The divide and conquer algorithm 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]
itype INTEGER
Specifies the problem type to be solved:
= 1: A*x = (lambda)*B*x
= 2: A*B*x = (lambda)*x
= 3: B*A*x = (lambda)*x
@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]
jobz magma_vec_t
- = MagmaNoVec: Compute eigenvalues only;
- = MagmaVec: Compute eigenvalues and eigenvectors.
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangles of A and B are stored;
- = MagmaLower: Lower triangles of A and B are stored.
@param[in]
n INTEGER
The order of the matrices A and B. N >= 0.
@param[in,out]
A REAL array, dimension (LDA, N)
On entry, the symmetric matrix A. If UPLO = MagmaUpper, the
leading N-by-N upper triangular part of A contains the
upper triangular part of the matrix A. If UPLO = MagmaLower,
the leading N-by-N lower triangular part of A contains
the lower triangular part of the matrix A.
\n
On exit, if JOBZ = MagmaVec, then if INFO = 0, A contains the
matrix Z of eigenvectors. The eigenvectors are normalized
as follows:
if ITYPE = 1 or 2, Z**T*B*Z = I;
if ITYPE = 3, Z**T*inv(B)*Z = I.
If JOBZ = MagmaNoVec, then on exit the upper triangle (if UPLO=MagmaUpper)
or the lower triangle (if UPLO=MagmaLower) of A, including the
diagonal, is destroyed.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[in,out]
B REAL array, dimension (LDB, N)
On entry, the symmetric matrix B. If UPLO = MagmaUpper, the
leading N-by-N upper triangular part of B contains the
upper triangular part of the matrix B. If UPLO = MagmaLower,
the leading N-by-N lower triangular part of B contains
the lower triangular part of the matrix B.
\n
On exit, if INFO <= N, the part of B containing the matrix is
overwritten by the triangular factor U or L from the Cholesky
factorization B = U**T*U or B = L*L**T.
@param[in]
ldb INTEGER
The leading dimension of the array B. LDB >= 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[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)
If INFO = 0, the eigenvalues in ascending order.
@param[out]
work (workspace) REAL array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The length of the array WORK.
If N <= 1, LWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LWORK >= 2*N + N*NB.
If JOBZ = MagmaVec and N > 1, LWORK >= max( 2*N + N*NB, 1 + 6*N + 2*N**2 ).
NB can be obtained through magma_get_ssytrd_nb(N).
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal sizes of the WORK and
IWORK arrays, returns these values as the first entries of
the WORK and IWORK arrays, and no error message
related to LWORK or LIWORK is issued by XERBLA.
@param[out]
iwork (workspace) INTEGER array, dimension (MAX(1,LIWORK))
On exit, if INFO = 0, IWORK[0] returns the optimal LIWORK.
@param[in]
liwork INTEGER
The dimension of the array IWORK.
If N <= 1, LIWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LIWORK >= 1.
If JOBZ = MagmaVec and N > 1, LIWORK >= 3 + 5*N.
\n
If LIWORK = -1, then a workspace query is assumed; the
routine only calculates the optimal sizes of the WORK
and IWORK arrays, returns these values as the first entries
of the WORK and IWORK arrays, and no error message
related to LWORK or LIWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: SPOTRF or SSYEVD returned an error code:
<= N: if INFO = i and JOBZ = MagmaNoVec, then the algorithm
failed to converge; i off-diagonal elements of an
intermediate tridiagonal form did not converge to
zero;
if INFO = i and JOBZ = MagmaVec, then 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);
> N: if INFO = N + i, for 1 <= i <= N, then the leading
minor of order i of B is not positive definite.
The factorization of B could not be completed and
no eigenvalues or eigenvectors were computed.
Further Details
---------------
Based on contributions by
Mark Fahey, Department of Mathematics, Univ. of Kentucky, USA
Modified so that no backsubstitution is performed if SSYEVD fails to
converge (NEIG in old code could be greater than N causing out of
bounds reference to A - reported by Ralf Meyer). Also corrected the
description of INFO and the test on ITYPE. Sven, 16 Feb 05.
@ingroup magma_ssygv_driver
********************************************************************/
extern "C" magma_int_t
magma_ssygvdx_m(
magma_int_t ngpu,
magma_int_t itype, magma_vec_t jobz, magma_range_t range, magma_uplo_t uplo, magma_int_t n,
float *A, magma_int_t lda,
float *B, magma_int_t ldb,
float vl, float vu, magma_int_t il, magma_int_t iu,
magma_int_t *m, float *w,
float *work, magma_int_t lwork,
#ifdef COMPLEX
float *rwork, magma_int_t lrwork,
#endif
magma_int_t *iwork, magma_int_t liwork,
magma_int_t *info)
{
/* Constants */
float c_one = MAGMA_S_ONE;
/* Local variables */
const char* uplo_ = lapack_uplo_const( uplo );
const char* jobz_ = lapack_vec_const( jobz );
magma_int_t lower;
magma_trans_t trans;
magma_int_t wantz;
magma_int_t lquery;
magma_int_t alleig, valeig, indeig;
magma_int_t lwmin;
magma_int_t liwmin;
wantz = (jobz == MagmaVec);
lower = (uplo == MagmaLower);
alleig = (range == MagmaRangeAll);
valeig = (range == MagmaRangeV);
indeig = (range == MagmaRangeI);
lquery = (lwork == -1 || liwork == -1);
*info = 0;
if (itype < 1 || itype > 3) {
*info = -1;
} else if (! (alleig || valeig || indeig)) {
*info = -2;
} else if (! (wantz || (jobz == MagmaNoVec))) {
*info = -3;
} else if (! (lower || (uplo == MagmaUpper))) {
*info = -4;
} else if (n < 0) {
*info = -5;
} else if (lda < max(1,n)) {
*info = -7;
} else if (ldb < max(1,n)) {
*info = -9;
} else {
if (valeig) {
if (n > 0 && vu <= vl) {
*info = -11;
}
} else if (indeig) {
if (il < 1 || il > max(1,n)) {
*info = -12;
} else if (iu < min(n,il) || iu > n) {
*info = -13;
}
}
}
magma_int_t nb = magma_get_ssytrd_nb( n );
if ( n <= 1 ) {
lwmin = 1;
liwmin = 1;
}
else if ( wantz ) {
lwmin = max( 2*n + n*nb, 1 + 6*n + 2*n*n );
liwmin = 3 + 5*n;
}
else {
lwmin = 2*n + n*nb;
liwmin = 1;
}
work[0] = magma_smake_lwork( lwmin );
iwork[0] = liwmin;
if (lwork < lwmin && ! lquery) {
*info = -17;
} else if (liwork < liwmin && ! lquery) {
*info = -19;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info));
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (n == 0) {
return *info;
}
/* 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_ssygvd(&itype, jobz_, uplo_,
&n, A, &lda, B, &ldb,
w, work, &lwork,
iwork, &liwork, info);
*m = n;
return *info;
}
magma_timer_t time=0;
timer_start( time );
magma_spotrf_m(ngpu, uplo, n, B, ldb, info);
if (*info != 0) {
*info = n + *info;
return *info;
}
timer_stop( time );
timer_printf( "time spotrf = %6.2f\n", time );
timer_start( time );
/* Transform problem to standard eigenvalue problem and solve. */
magma_ssygst_m(ngpu, itype, uplo, n, A, lda, B, ldb, info);
timer_stop( time );
timer_printf( "time ssygst = %6.2f\n", time );
timer_start( time );
magma_ssyevdx_m(ngpu, jobz, range, uplo, n, A, lda, vl, vu, il, iu, m, w, work, lwork, iwork, liwork, info);
timer_stop( time );
timer_printf( "time ssyevd = %6.2f\n", time );
if (wantz && *info == 0) {
timer_start( time );
/* Backtransform eigenvectors to the original problem. */
if (itype == 1 || itype == 2) {
/* For A*x=(lambda)*B*x and A*B*x=(lambda)*x;
backtransform eigenvectors: x = inv(L)'*y or inv(U)*y */
if (lower) {
trans = MagmaTrans;
} else {
trans = MagmaNoTrans;
}
magma_strsm_m( ngpu, MagmaLeft, uplo, trans, MagmaNonUnit,
n, *m, c_one, B, ldb, A, lda );
}
else if (itype == 3) {
/* For B*A*x=(lambda)*x;
backtransform eigenvectors: x = L*y or U'*y */
if (lower) {
trans = MagmaNoTrans;
} else {
trans = MagmaTrans;
}
#ifdef ENABLE_DEBUG
printf("--- the multi GPU version is falling back to 1 GPU to perform the last TRMM since there is no TRMM_mgpu --- \n");
#endif
float *dA=NULL, *dB=NULL;
magma_int_t ldda = magma_roundup( n, 32 );
magma_int_t lddb = ldda;
if (MAGMA_SUCCESS != magma_smalloc( &dA, ldda*(*m) ) ||
MAGMA_SUCCESS != magma_smalloc( &dB, lddb*n ) ) {
magma_free( dA );
magma_free( dB );
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_queue_t queue;
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queue );
magma_ssetmatrix( n, n, B, ldb, dB, lddb, queue );
magma_ssetmatrix( n, (*m), A, lda, dA, ldda, queue );
magma_strmm( MagmaLeft, uplo, trans, MagmaNonUnit,
n, (*m), c_one, dB, lddb, dA, ldda, queue );
magma_sgetmatrix( n, (*m), dA, ldda, A, lda, queue );
magma_queue_destroy( queue );
magma_free( dA );
magma_free( dB );
}
timer_stop( time );
timer_printf( "time setmatrices trsm/mm + getmatrices = %6.2f\n", time );
}
work[0] = magma_smake_lwork( lwmin );
iwork[0] = liwmin;
return *info;
} /* magma_ssygvd_m */
magma_err_t
magma_slarfb_gpu( int side, int trans, int direct, int storev,
magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloat_ptr dV, size_t dV_offset, magma_int_t ldv,
magmaFloat_ptr dT, size_t dT_offset, magma_int_t ldt,
magmaFloat_ptr dC, size_t dC_offset, magma_int_t ldc,
magmaFloat_ptr dwork, size_t dwork_offset, magma_int_t ldwork,
magma_queue_t queue)
{
/* -- clMAGMA (version 1.0.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
April 2012
Purpose
=======
SLARFB applies a real block reflector H or its transpose H' to a
REAL m by n matrix C, from the left.
Arguments
=========
SIDE (input) CHARACTER
= 'L': apply H or H' from the Left
= 'R': apply H or H' from the Right
TRANS (input) CHARACTER
= 'N': apply H (No transpose)
= 'C': apply H' (Conjugate transpose)
DIRECT (input) CHARACTER
Indicates how H is formed from a product of elementary
reflectors
= 'F': H = H(1) H(2) . . . H(k) (Forward)
= 'B': H = H(k) . . . H(2) H(1) (Backward)
STOREV (input) CHARACTER
Indicates how the vectors which define the elementary
reflectors are stored:
= 'C': Columnwise
= 'R': Rowwise
M (input) INTEGER
The number of rows of the matrix C.
N (input) INTEGER
The number of columns of the matrix C.
K (input) INTEGER
The order of the matrix T (= the number of elementary
reflectors whose product defines the block reflector).
DV (input) REAL array, dimension (LDV,K)
The matrix V. See further details.
LDV (input) INTEGER
The leading dimension of the array V. LDV >= max(1,M);
DT (input) REAL array, dimension (LDT,K)
The triangular k by k matrix T in the representation of the
block reflector.
LDT (input) INTEGER
The leading dimension of the array T. LDT >= K.
DC (input/output) REAL array, dimension (LDC,N)
On entry, the m by n matrix C.
On exit, C is overwritten by H*C.
LDC (input) INTEGER
The leading dimension of the array C. LDA >= max(1,M).
WORK (workspace) REAL array, dimension (LDWORK,K)
LDWORK (input) INTEGER
The leading dimension of the array WORK.
If SIDE == 'L', LDWORK >= max(1,N);
if SIDE == 'R', LDWORK >= max(1,M);
=================================================================== */
#define dV(i) dV, (i)
#define dT(i) dT, (i)
#define dC(i) dC, (i)
#define dwork(i) dwork, (i)
float c_zero = MAGMA_S_MAKE( 0.0, 0.0 );
float c_one = MAGMA_S_MAKE( 1.0, 0.0 );
float c_neg_one = MAGMA_S_MAKE( -1.0, 0.0 );
if (m <= 0 || n <= 0) {
return MAGMA_SUCCESS;
}
magma_int_t transt;
if (trans == MagmaNoTrans)
transt = MagmaTrans;
else
transt = MagmaNoTrans;
if ( side == MagmaLeft ) {
if ( storev == MagmaColumnwise )
{
magma_sgemm( MagmaTrans, MagmaNoTrans,
n, k, m,
c_one, dC(dC_offset), ldc,
dV(dV_offset), ldv,
c_zero, dwork(dwork_offset), ldwork, queue);
if (direct == MagmaForward)
magma_strmm( MagmaRight, MagmaUpper, transt, MagmaNonUnit,
n, k,
c_one, dT(dT_offset), ldt,
dwork(dwork_offset), ldwork, queue);
else
magma_strmm( MagmaRight, MagmaLower, transt, MagmaNonUnit,
n, k,
c_one, dT(dT_offset), ldt,
dwork(dwork_offset), ldwork, queue);
magma_sgemm( MagmaNoTrans, MagmaTrans,
m, n, k,
c_neg_one, dV(dV_offset), ldv,
dwork(dwork_offset), ldwork,
c_one, dC(dC_offset), ldc, queue);
}
else {
magma_sgemm( MagmaNoTrans, MagmaTrans,
m, k, n,
c_one, dC(dC_offset), ldc,
dV(dV_offset), ldv,
c_zero, dwork(dwork_offset), ldwork, queue);
magma_strmm( MagmaRight, MagmaUpper, transt, MagmaNonUnit,
m, k,
c_one, dT(dT_offset), ldt,
dwork(dwork_offset), ldwork, queue);
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
m, n, k,
c_neg_one, dwork(dwork_offset), ldwork,
dV(dV_offset), ldv,
c_one, dC(dC_offset), ldc, queue);
}
}
else {
/* Case side == 'R' */
if ( storev == MagmaColumnwise ) {
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
m, k, n,
c_one, dC(dC_offset), ldc,
dV(dV_offset), ldv,
c_zero, dwork(dwork_offset), ldwork, queue);
// ??? ldwork replaced by k for case n < k
if (direct == MagmaForward)
magma_strmm( MagmaRight, MagmaUpper, transt, MagmaNonUnit,
m, k,
c_one, dT(dT_offset), ldt,
dwork(dwork_offset), ldwork, queue);
else
magma_strmm( MagmaRight, MagmaLower, transt, MagmaNonUnit,
m, k,
c_one, dT(dT_offset), ldt,
dwork(dwork_offset), ldwork, queue);
magma_sgemm( MagmaNoTrans, MagmaTrans,
m, n, k,
c_neg_one, dwork(dwork_offset), ldwork,
dV(dV_offset), ldv,
c_one, dC(dC_offset), ldc, queue);
}
else {
magma_sgemm( MagmaNoTrans, MagmaTrans,
m, k, n,
c_one, dC(dC_offset), ldc,
dV(dV_offset), ldv,
c_zero, dwork(dwork_offset), ldwork, queue);
magma_strmm( MagmaRight, MagmaUpper, transt, MagmaNonUnit,
m, k,
c_one, dT(dT_offset), ldt,
dwork(dwork_offset), ldwork, queue);
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
m, n, k,
c_neg_one, dwork(dwork_offset), ldwork,
dV(dV_offset), ldv,
c_one, dC(dC_offset), ldc, queue);
}
}
return MAGMA_SUCCESS;
} /* magma_slarfb */
/**
Purpose
-------
SLAUUM computes the product U * U' or L' * L, where the triangular
factor U or L is stored in the upper or lower triangular part of
the array A.
If UPLO = MagmaUpper then the upper triangle of the result is stored,
overwriting the factor U in A.
If UPLO = MagmaLower then the lower triangle of the result is stored,
overwriting the factor L in A.
This is the blocked form of the algorithm, calling Level 3 BLAS.
Arguments
---------
@param[in]
uplo magma_uplo_t
Specifies whether the triangular factor stored in the array A
is upper or lower triangular:
- = MagmaUpper: Upper triangular
- = MagmaLower: Lower triangular
@param[in]
n INTEGER
The order of the triangular factor U or L. N >= 0.
@param[in,out]
A COPLEX_16 array, dimension (LDA,N)
On entry, the triangular factor U or L.
On exit, if UPLO = MagmaUpper, the upper triangle of A is
overwritten with the upper triangle of the product U * U';
if UPLO = MagmaLower, the lower triangle of A is overwritten with
the lower triangle of the product L' * L.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -k, the k-th argument had an illegal value
@ingroup magma_sposv_aux
***************************************************************************/
extern "C" magma_int_t
magma_slauum(
magma_uplo_t uplo, magma_int_t n,
float *A, magma_int_t lda,
magma_int_t *info)
{
#define A(i, j) (A + (j)*lda + (i))
#define dA(i, j) (dA + (j)*ldda + (i))
/* Local variables */
const char* uplo_ = lapack_uplo_const( uplo );
magma_int_t ldda, nb;
magma_int_t i, ib;
float c_one = MAGMA_S_ONE;
float d_one = MAGMA_D_ONE;
float *dA;
int upper = (uplo == MagmaUpper);
*info = 0;
if (! upper && uplo != MagmaLower)
*info = -1;
else if (n < 0)
*info = -2;
else if (lda < max(1,n))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return */
if ( n == 0 )
return *info;
ldda = ((n+31)/32)*32;
if (MAGMA_SUCCESS != magma_smalloc( &dA, (n)*ldda )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_queue_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
nb = magma_get_spotrf_nb(n);
if (nb <= 1 || nb >= n)
lapackf77_slauum(uplo_, &n, A, &lda, info);
else {
if (upper) {
/* Compute the product U * U'. */
for (i=0; i < n; i += nb) {
ib=min(nb,n-i);
magma_ssetmatrix_async( ib, ib,
A(i,i), lda,
dA(i, i), ldda, stream[1] );
magma_ssetmatrix_async( ib, (n-i-ib),
A(i,i+ib), lda,
dA(i,i+ib), ldda, stream[0] );
magma_queue_sync( stream[1] );
magma_strmm( MagmaRight, MagmaUpper,
MagmaConjTrans, MagmaNonUnit, i, ib,
c_one, dA(i,i), ldda, dA(0, i),ldda);
lapackf77_slauum(MagmaUpperStr, &ib, A(i,i), &lda, info);
magma_ssetmatrix_async( ib, ib,
A(i, i), lda,
dA(i, i), ldda, stream[0] );
if (i+ib < n) {
magma_sgemm( MagmaNoTrans, MagmaConjTrans,
i, ib, (n-i-ib), c_one, dA(0,i+ib),
ldda, dA(i, i+ib),ldda, c_one,
dA(0,i), ldda);
magma_queue_sync( stream[0] );
magma_ssyrk( MagmaUpper, MagmaNoTrans, ib,(n-i-ib),
d_one, dA(i, i+ib), ldda,
d_one, dA(i, i), ldda);
}
magma_sgetmatrix( i+ib, ib,
dA(0, i), ldda,
A(0, i), lda );
}
}
else {
/* Compute the product L' * L. */
for (i=0; i < n; i += nb) {
ib=min(nb,n-i);
magma_ssetmatrix_async( ib, ib,
A(i,i), lda,
dA(i, i), ldda, stream[1] );
magma_ssetmatrix_async( (n-i-ib), ib,
A(i+ib, i), lda,
dA(i+ib, i), ldda, stream[0] );
magma_queue_sync( stream[1] );
magma_strmm( MagmaLeft, MagmaLower,
MagmaConjTrans, MagmaNonUnit, ib,
i, c_one, dA(i,i), ldda,
dA(i, 0),ldda);
lapackf77_slauum(MagmaLowerStr, &ib, A(i,i), &lda, info);
magma_ssetmatrix_async( ib, ib,
A(i, i), lda,
dA(i, i), ldda, stream[0] );
if (i+ib < n) {
magma_sgemm(MagmaConjTrans, MagmaNoTrans,
ib, i, (n-i-ib), c_one, dA( i+ib,i),
ldda, dA(i+ib, 0),ldda, c_one,
dA(i,0), ldda);
magma_queue_sync( stream[0] );
magma_ssyrk(MagmaLower, MagmaConjTrans, ib, (n-i-ib),
d_one, dA(i+ib, i), ldda,
d_one, dA(i, i), ldda);
}
magma_sgetmatrix( ib, i+ib,
dA(i, 0), ldda,
A(i, 0), lda );
}
}
}
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free( dA );
return *info;
}
/**
Purpose
-------
SSYGVDX computes selected eigenvalues and, optionally, eigenvectors
of a real generalized symmetric-definite eigenproblem, of the form
A*x=(lambda)*B*x, A*Bx=(lambda)*x, or B*A*x=(lambda)*x. Here A and
B are assumed to be symmetric and B is also positive definite.
Eigenvalues and eigenvectors can be selected by specifying either a
range of values or a range of indices for the desired eigenvalues.
If eigenvectors are desired, it uses a divide and conquer algorithm.
The divide and conquer algorithm 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]
itype INTEGER
Specifies the problem type to be solved:
= 1: A*x = (lambda)*B*x
= 2: A*B*x = (lambda)*x
= 3: B*A*x = (lambda)*x
@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]
jobz magma_vec_t
- = MagmaNoVec: Compute eigenvalues only;
- = MagmaVec: Compute eigenvalues and eigenvectors.
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangles of A and B are stored;
- = MagmaLower: Lower triangles of A and B are stored.
@param[in]
n INTEGER
The order of the matrices A and B. N >= 0.
@param[in,out]
A REAL array, dimension (LDA, N)
On entry, the symmetric matrix A. If UPLO = MagmaUpper, the
leading N-by-N upper triangular part of A contains the
upper triangular part of the matrix A. If UPLO = MagmaLower,
the leading N-by-N lower triangular part of A contains
the lower triangular part of the matrix A.
\n
On exit, if JOBZ = MagmaVec, then if INFO = 0, A contains the
matrix Z of eigenvectors. The eigenvectors are normalized
as follows:
if ITYPE = 1 or 2, Z**T * B * Z = I;
if ITYPE = 3, Z**T * inv(B) * Z = I.
If JOBZ = MagmaNoVec, then on exit the upper triangle (if UPLO=MagmaUpper)
or the lower triangle (if UPLO=MagmaLower) of A, including the
diagonal, is destroyed.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[in,out]
B REAL array, dimension (LDB, N)
On entry, the symmetric matrix B. If UPLO = MagmaUpper, the
leading N-by-N upper triangular part of B contains the
upper triangular part of the matrix B. If UPLO = MagmaLower,
the leading N-by-N lower triangular part of B contains
the lower triangular part of the matrix B.
\n
On exit, if INFO <= N, the part of B containing the matrix is
overwritten by the triangular factor U or L from the Cholesky
factorization B = U**T * U or B = L * L**T.
@param[in]
ldb INTEGER
The leading dimension of the array B. LDB >= 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[out]
mout INTEGER
The total number of eigenvalues found. 0 <= MOUT <= N.
If RANGE = MagmaRangeAll, MOUT = N, and if RANGE = MagmaRangeI, MOUT = IU-IL+1.
@param[out]
w REAL array, dimension (N)
If INFO = 0, the eigenvalues in ascending order.
@param[out]
work (workspace) REAL array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[out]
work (workspace) REAL array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The length of the array WORK.
If N <= 1, LWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LWORK >= 2*N + N*NB.
If JOBZ = MagmaVec and N > 1, LWORK >= max( 2*N + N*NB, 1 + 6*N + 2*N**2 ).
NB can be obtained through magma_get_ssytrd_nb(N).
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal sizes of the WORK and IWORK
arrays, returns these values as the first entries of the WORK
and IWORK arrays, and no error message related to LWORK or
LIWORK is issued by XERBLA.
@param[out]
iwork (workspace) INTEGER array, dimension (MAX(1,LIWORK))
On exit, if INFO = 0, IWORK[0] returns the optimal LIWORK.
@param[in]
liwork INTEGER
The dimension of the array IWORK.
If N <= 1, LIWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LIWORK >= 1.
If JOBZ = MagmaVec and N > 1, LIWORK >= 3 + 5*N.
\n
If LIWORK = -1, then a workspace query is assumed; the
routine only calculates the optimal sizes of the WORK and
IWORK arrays, returns these values as the first entries of
the WORK and IWORK arrays, and no error message related to
LWORK or LIWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: SPOTRF or SSYEVD returned an error code:
<= N: if INFO = i and JOBZ = MagmaNoVec, then the algorithm
failed to converge; i off-diagonal elements of an
intermediate tridiagonal form did not converge to
zero;
if INFO = i and JOBZ = MagmaVec, then 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);
> N: if INFO = N + i, for 1 <= i <= N, then the leading
minor of order i of B is not positive definite.
The factorization of B could not be completed and
no eigenvalues or eigenvectors were computed.
Further Details
---------------
Based on contributions by
Mark Fahey, Department of Mathematics, Univ. of Kentucky, USA
Modified so that no backsubstitution is performed if SSYEVD fails to
converge (NEIG in old code could be greater than N causing out of
bounds reference to A - reported by Ralf Meyer). Also corrected the
description of INFO and the test on ITYPE. Sven, 16 Feb 05.
@ingroup magma_ssygv_driver
********************************************************************/
extern "C" magma_int_t
magma_ssygvdx(
magma_int_t itype, magma_vec_t jobz, magma_range_t range, magma_uplo_t uplo, magma_int_t n,
float *A, magma_int_t lda,
float *B, magma_int_t ldb,
float vl, float vu, magma_int_t il, magma_int_t iu,
magma_int_t *mout, float *w,
float *work, magma_int_t lwork,
#ifdef COMPLEX
float *rwork, magma_int_t lrwork,
#endif
magma_int_t *iwork, magma_int_t liwork,
magma_int_t *info)
{
const char* uplo_ = lapack_uplo_const( uplo );
const char* jobz_ = lapack_vec_const( jobz );
float d_one = MAGMA_S_ONE;
float *dA=NULL, *dB=NULL;
magma_int_t ldda = roundup( n, 32 );
magma_int_t lddb = ldda;
magma_int_t lower;
magma_trans_t trans;
magma_int_t wantz, lquery;
magma_int_t alleig, valeig, indeig;
magma_int_t lwmin, liwmin;
magma_queue_t stream;
magma_queue_create( &stream );
wantz = (jobz == MagmaVec);
lower = (uplo == MagmaLower);
alleig = (range == MagmaRangeAll);
valeig = (range == MagmaRangeV);
indeig = (range == MagmaRangeI);
lquery = (lwork == -1 || liwork == -1);
*info = 0;
if (itype < 1 || itype > 3) {
*info = -1;
} else if (! (alleig || valeig || indeig)) {
*info = -2;
} else if (! (wantz || (jobz == MagmaNoVec))) {
*info = -3;
} else if (! (lower || (uplo == MagmaUpper))) {
*info = -4;
} else if (n < 0) {
*info = -5;
} else if (lda < max(1,n)) {
*info = -7;
} else if (ldb < max(1,n)) {
*info = -9;
} else {
if (valeig) {
if (n > 0 && vu <= vl) {
*info = -11;
}
} else if (indeig) {
if (il < 1 || il > max(1,n)) {
*info = -12;
} else if (iu < min(n,il) || iu > n) {
*info = -13;
}
}
}
magma_int_t nb = magma_get_ssytrd_nb( n );
if ( n <= 1 ) {
lwmin = 1;
liwmin = 1;
}
else if ( wantz ) {
lwmin = max( 2*n + n*nb, 1 + 6*n + 2*n*n );
liwmin = 3 + 5*n;
}
else {
lwmin = 2*n + n*nb;
liwmin = 1;
}
// multiply by 1+eps (in Double!) to ensure length gets rounded up,
// if it cannot be exactly represented in floating point.
real_Double_t one_eps = 1. + lapackf77_slamch("Epsilon");
work[0] = lwmin * one_eps;
iwork[0] = liwmin;
if (lwork < lwmin && ! lquery) {
*info = -17;
} else if (liwork < liwmin && ! lquery) {
*info = -19;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (n == 0) {
return *info;
}
/* If matrix is very small, then just call LAPACK on CPU, no need for GPU */
if (n <= 128) {
lapackf77_ssygvd( &itype, jobz_, uplo_,
&n, A, &lda, B, &ldb,
w, work, &lwork,
iwork, &liwork, info );
*mout = n;
return *info;
}
if (MAGMA_SUCCESS != magma_smalloc( &dA, n*ldda ) ||
MAGMA_SUCCESS != magma_smalloc( &dB, n*lddb )) {
magma_free( dA );
magma_free( dB );
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/* Form a Cholesky factorization of B. */
magma_ssetmatrix( n, n, B, ldb, dB, lddb );
magma_ssetmatrix_async( n, n,
A, lda,
dA, ldda, stream );
magma_timer_t time=0;
timer_start( time );
magma_spotrf_gpu( uplo, n, dB, lddb, info );
if (*info != 0) {
*info = n + *info;
return *info;
}
timer_stop( time );
timer_printf( "time spotrf_gpu = %6.2f\n", time );
magma_queue_sync( stream );
magma_sgetmatrix_async( n, n,
dB, lddb,
B, ldb, stream );
timer_start( time );
/* Transform problem to standard eigenvalue problem and solve. */
magma_ssygst_gpu( itype, uplo, n, dA, ldda, dB, lddb, info );
timer_stop( time );
timer_printf( "time ssygst_gpu = %6.2f\n", time );
/* simple fix to be able to run bigger size.
* set dB=NULL so we know to re-allocate below
* TODO: have dwork here that will be used as dB and then passed to ssyevd.
*/
if (n > 5000) {
magma_queue_sync( stream );
magma_free( dB ); dB=NULL;
}
timer_start( time );
magma_ssyevdx_gpu( jobz, range, uplo, n, dA, ldda, vl, vu, il, iu, mout, w, A, lda,
work, lwork, iwork, liwork, info );
timer_stop( time );
timer_printf( "time ssyevdx_gpu = %6.2f\n", time );
if (wantz && *info == 0) {
timer_start( time );
/* allocate and copy dB back */
if (dB == NULL) {
if (MAGMA_SUCCESS != magma_smalloc( &dB, n*lddb ) ) {
magma_free( dA ); dA=NULL;
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_ssetmatrix( n, n, B, ldb, dB, lddb );
}
/* Backtransform eigenvectors to the original problem. */
if (itype == 1 || itype == 2) {
/* For A*x=(lambda)*B*x and A*B*x=(lambda)*x;
backtransform eigenvectors: x = inv(L)'*y or inv(U)*y */
if (lower) {
trans = MagmaTrans;
} else {
trans = MagmaNoTrans;
}
magma_strsm( MagmaLeft, uplo, trans, MagmaNonUnit,
n, *mout, d_one, dB, lddb, dA, ldda );
}
else if (itype == 3) {
/* For B*A*x=(lambda)*x;
backtransform eigenvectors: x = L*y or U'*y */
if (lower) {
trans = MagmaNoTrans;
} else {
trans = MagmaTrans;
}
magma_strmm( MagmaLeft, uplo, trans, MagmaNonUnit,
n, *mout, d_one, dB, lddb, dA, ldda );
}
magma_sgetmatrix( n, *mout, dA, ldda, A, lda );
timer_stop( time );
timer_printf( "time strsm/mm + getmatrix = %6.2f\n", time );
}
magma_queue_sync( stream );
magma_queue_destroy( stream );
work[0] = lwmin * one_eps; // round up
iwork[0] = liwmin;
magma_free( dA ); dA=NULL;
magma_free( dB ); dB=NULL;
return *info;
} /* magma_ssygvd */
extern "C" magma_int_t
magma_ssygvdx_2stage(magma_int_t itype, char jobz, char range, char uplo, magma_int_t n,
float *a, magma_int_t lda, float *b, magma_int_t ldb,
float vl, float vu, magma_int_t il, magma_int_t iu,
magma_int_t *m, float *w, float *work, magma_int_t lwork,
magma_int_t *iwork, magma_int_t liwork, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
SSYGVDX_2STAGE computes all the eigenvalues, and optionally, the eigenvectors
of a complex generalized Hermitian-definite eigenproblem, of the form
A*x=(lambda)*B*x, A*Bx=(lambda)*x, or B*A*x=(lambda)*x. Here A and
B are assumed to be Hermitian and B is also positive definite.
It uses a two-stage algorithm for the tridiagonalization.
If eigenvectors are desired, it uses a divide and conquer algorithm.
The divide and conquer algorithm 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
=========
ITYPE (input) INTEGER
Specifies the problem type to be solved:
= 1: A*x = (lambda)*B*x
= 2: A*B*x = (lambda)*x
= 3: B*A*x = (lambda)*x
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.
JOBZ (input) CHARACTER*1
= 'N': Compute eigenvalues only;
= 'V': Compute eigenvalues and eigenvectors.
UPLO (input) CHARACTER*1
= 'U': Upper triangles of A and B are stored;
= 'L': Lower triangles of A and B are stored.
N (input) INTEGER
The order of the matrices A and B. N >= 0.
A (input/output) DOUBLE PRECISION 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, if JOBZ = 'V', then if INFO = 0, A contains the
matrix Z of eigenvectors. The eigenvectors are normalized
as follows:
if ITYPE = 1 or 2, Z**H*B*Z = I;
if ITYPE = 3, Z**H*inv(B)*Z = I.
If JOBZ = 'N', then on exit the upper triangle (if UPLO='U')
or the lower triangle (if UPLO='L') of A, including the
diagonal, is destroyed.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
B (input/output) DOUBLE PRECISION array, dimension (LDB, N)
On entry, the Hermitian matrix B. If UPLO = 'U', the
leading N-by-N upper triangular part of B contains the
upper triangular part of the matrix B. If UPLO = 'L',
the leading N-by-N lower triangular part of B contains
the lower triangular part of the matrix B.
On exit, if INFO <= N, the part of B containing the matrix is
overwritten by the triangular factor U or L from the Cholesky
factorization B = U**H*U or B = L*L**H.
LDB (input) INTEGER
The leading dimension of the array B. LDB >= max(1,N).
VL (input) DOUBLE PRECISION
VU (input) DOUBLE PRECISION
If RANGE='V', the lower and upper bounds of the interval to
be searched for eigenvalues. VL < VU.
Not referenced if RANGE = 'A' or 'I'.
IL (input) INTEGER
IU (input) INTEGER
If RANGE='I', the indices (in ascending order) of the
smallest and largest eigenvalues to be returned.
1 <= IL <= IU <= N, if N > 0; IL = 1 and IU = 0 if N = 0.
Not referenced if RANGE = 'A' or 'V'.
M (output) INTEGER
The total number of eigenvalues found. 0 <= M <= N.
If RANGE = 'A', M = N, and if RANGE = 'I', M = IU-IL+1.
W (output) DOUBLE PRECISION array, dimension (N)
If INFO = 0, the eigenvalues in ascending order.
WORK (workspace/output) DOUBLE PRECISION array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The length of the array WORK.
If N <= 1, LWORK >= 1.
If JOBZ = 'N' and N > 1, LWORK >= LQ2 + N * (NB + 2).
If JOBZ = 'V' and N > 1, LWORK >= LQ2 + 1 + 6*N + 2*N**2.
where LQ2 is the size needed to store
the Q2 matrix and is returned by
MAGMA_BULGE_GET_LQ2.
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal sizes of the WORK, RWORK and
IWORK arrays, returns these values as the first entries of
the WORK, RWORK and IWORK arrays, and no error message
related to LWORK or LRWORK or LIWORK is issued by XERBLA.
IWORK (workspace/output) INTEGER array, dimension (MAX(1,LIWORK))
On exit, if INFO = 0, IWORK(1) returns the optimal LIWORK.
LIWORK (input) INTEGER
The dimension of the array IWORK.
If N <= 1, LIWORK >= 1.
If JOBZ = 'N' and N > 1, LIWORK >= 1.
If JOBZ = 'V' and N > 1, LIWORK >= 3 + 5*N.
If LIWORK = -1, then a workspace query is assumed; the
routine only calculates the optimal sizes of the WORK, RWORK
and IWORK arrays, returns these values as the first entries
of the WORK, RWORK and IWORK arrays, and no error message
related to LWORK or LRWORK or LIWORK is issued by XERBLA.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
> 0: ZPOTRF or ZHEEVD returned an error code:
<= N: if INFO = i and JOBZ = 'N', then the algorithm
failed to converge; i off-diagonal elements of an
intermediate tridiagonal form did not converge to
zero;
if INFO = i and JOBZ = 'V', then 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);
> N: if INFO = N + i, for 1 <= i <= N, then the leading
minor of order i of B is not positive definite.
The factorization of B could not be completed and
no eigenvalues or eigenvectors were computed.
Further Details
===============
Based on contributions by
Mark Fahey, Department of Mathematics, Univ. of Kentucky, USA
Modified so that no backsubstitution is performed if ZHEEVD fails to
converge (NEIG in old code could be greater than N causing out of
bounds reference to A - reported by Ralf Meyer). Also corrected the
description of INFO and the test on ITYPE. Sven, 16 Feb 05.
===================================================================== */
char uplo_[2] = {uplo, 0};
char jobz_[2] = {jobz, 0};
char range_[2] = {range, 0};
float d_one = MAGMA_S_ONE;
float *da;
float *db;
magma_int_t ldda = n;
magma_int_t lddb = n;
magma_int_t lower;
char trans[1];
magma_int_t wantz;
magma_int_t lquery;
magma_int_t alleig, valeig, indeig;
magma_int_t lwmin;
magma_int_t liwmin;
magma_queue_t stream;
magma_queue_create( &stream );
/* determine the number of threads */
magma_int_t threads = magma_get_numthreads();
magma_setlapack_numthreads(threads);
wantz = lapackf77_lsame(jobz_, MagmaVecStr);
lower = lapackf77_lsame(uplo_, MagmaLowerStr);
alleig = lapackf77_lsame(range_, "A");
valeig = lapackf77_lsame(range_, "V");
indeig = lapackf77_lsame(range_, "I");
lquery = lwork == -1 || liwork == -1;
*info = 0;
if (itype < 1 || itype > 3) {
*info = -1;
} else if (! (alleig || valeig || indeig)) {
*info = -2;
} else if (! (wantz || lapackf77_lsame(jobz_, MagmaNoVecStr))) {
*info = -3;
} else if (! (lower || lapackf77_lsame(uplo_, MagmaUpperStr))) {
*info = -4;
} else if (n < 0) {
*info = -5;
} else if (lda < max(1,n)) {
*info = -7;
} else if (ldb < max(1,n)) {
*info = -9;
} else {
if (valeig) {
if (n > 0 && vu <= vl) {
*info = -11;
}
} else if (indeig) {
if (il < 1 || il > max(1,n)) {
*info = -12;
} else if (iu < min(n,il) || iu > n) {
*info = -13;
}
}
}
magma_int_t nb = magma_get_sbulge_nb(n, threads);
magma_int_t lq2 = magma_sbulge_get_lq2(n, threads);
if (wantz) {
lwmin = lq2 + 1 + 6*n + 2*n*n;
liwmin = 3 + 5*n;
} else {
lwmin = n * (nb + 2);
liwmin = 1;
}
work[0] = lwmin * (1. + lapackf77_slamch("Epsilon"));
iwork[0] = liwmin;
if (lwork < lwmin && ! lquery) {
*info = -17;
} else if (liwork < liwmin && ! lquery) {
*info = -19;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info));
return *info;
} else if (lquery) {
return *info;
}
/* Quick return if possible */
if (n == 0) {
return *info;
}
/* 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_ssygvd(&itype, jobz_, uplo_,
&n, a, &lda, b, &ldb,
w, work, &lwork,
iwork, &liwork, info);
*m = n;
return *info;
}
if (MAGMA_SUCCESS != magma_smalloc( &da, n*ldda ) ||
MAGMA_SUCCESS != magma_smalloc( &db, n*lddb )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/* Form a Cholesky factorization of B. */
magma_ssetmatrix( n, n, b, ldb, db, lddb );
magma_ssetmatrix_async( n, n,
a, lda,
da, ldda, stream );
#ifdef ENABLE_TIMER
magma_timestr_t start, end;
start = get_current_time();
#endif
magma_spotrf_gpu(uplo_[0], n, db, lddb, info);
if (*info != 0) {
*info = n + *info;
return *info;
}
#ifdef ENABLE_TIMER
end = get_current_time();
printf("time spotrf_gpu = %6.2f\n", GetTimerValue(start,end)/1000.);
#endif
magma_queue_sync( stream );
magma_sgetmatrix_async( n, n,
db, lddb,
b, ldb, stream );
#ifdef ENABLE_TIMER
start = get_current_time();
#endif
/* Transform problem to standard eigenvalue problem and solve. */
magma_ssygst_gpu(itype, uplo, n, da, ldda, db, lddb, info);
#ifdef ENABLE_TIMER
end = get_current_time();
printf("time ssygst_gpu = %6.2f\n", GetTimerValue(start,end)/1000.);
#endif
magma_sgetmatrix( n, n, da, ldda, a, lda );
magma_queue_sync( stream );
magma_free( da );
magma_free( db );
#ifdef ENABLE_TIMER
start = get_current_time();
#endif
magma_ssyevdx_2stage(jobz, range, uplo, n, a, lda, vl, vu, il, iu, m, w, work, lwork, iwork, liwork, info);
#ifdef ENABLE_TIMER
end = get_current_time();
printf("time ssyevdx_2stage = %6.2f\n", GetTimerValue(start,end)/1000.);
#endif
if (wantz && *info == 0) {
if (MAGMA_SUCCESS != magma_smalloc( &da, n*ldda ) ||
MAGMA_SUCCESS != magma_smalloc( &db, n*lddb )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
#ifdef ENABLE_TIMER
start = get_current_time();
#endif
magma_ssetmatrix( n, *m, a, lda, da, ldda );
magma_ssetmatrix( n, n, b, ldb, db, lddb );
/* Backtransform eigenvectors to the original problem. */
if (itype == 1 || itype == 2) {
/* For A*x=(lambda)*B*x and A*B*x=(lambda)*x;
backtransform eigenvectors: x = inv(L)'*y or inv(U)*y */
if (lower) {
*(unsigned char *)trans = MagmaConjTrans;
} else {
*(unsigned char *)trans = MagmaNoTrans;
}
magma_strsm(MagmaLeft, uplo, *trans, MagmaNonUnit, n, *m, d_one, db, lddb, da, ldda);
}
else if (itype == 3) {
/* For B*A*x=(lambda)*x;
backtransform eigenvectors: x = L*y or U'*y */
if (lower) {
*(unsigned char *)trans = MagmaNoTrans;
} else {
*(unsigned char *)trans = MagmaConjTrans;
}
magma_strmm(MagmaLeft, uplo, *trans, MagmaNonUnit, n, *m, d_one, db, lddb, da, ldda);
}
magma_sgetmatrix( n, *m, da, ldda, a, lda );
#ifdef ENABLE_TIMER
end = get_current_time();
printf("time strsm/mm + getmatrix = %6.2f\n", GetTimerValue(start,end)/1000.);
#endif
magma_free( da );
magma_free( db );
}
magma_queue_destroy( stream );
work[0] = lwmin * (1. + lapackf77_slamch("Epsilon"));
iwork[0] = liwmin;
return *info;
} /* ssygvdx_2stage */
/**
Purpose
-------
SSYGVD computes all the eigenvalues, and optionally, the eigenvectors
of a real generalized symmetric-definite eigenproblem, of the form
A*x=(lambda)*B*x, A*Bx=(lambda)*x, or B*A*x=(lambda)*x. Here A and
B are assumed to be symmetric and B is also positive definite.
If eigenvectors are desired, it uses a divide and conquer algorithm.
The divide and conquer algorithm 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]
itype INTEGER
Specifies the problem type to be solved:
= 1: A*x = (lambda)*B*x
= 2: A*B*x = (lambda)*x
= 3: B*A*x = (lambda)*x
@param[in]
jobz magma_vec_t
- = MagmaNoVec: Compute eigenvalues only;
- = MagmaVec: Compute eigenvalues and eigenvectors.
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangles of A and B are stored;
- = MagmaLower: Lower triangles of A and B are stored.
@param[in]
n INTEGER
The order of the matrices A and B. N >= 0.
@param[in,out]
A REAL array, dimension (LDA, N)
On entry, the symmetric matrix A. If UPLO = MagmaUpper, the
leading N-by-N upper triangular part of A contains the
upper triangular part of the matrix A. If UPLO = MagmaLower,
the leading N-by-N lower triangular part of A contains
the lower triangular part of the matrix A.
\n
On exit, if JOBZ = MagmaVec, then if INFO = 0, A contains the
matrix Z of eigenvectors. The eigenvectors are normalized
as follows:
if ITYPE = 1 or 2, Z**T * B * Z = I;
if ITYPE = 3, Z**T * inv(B) * Z = I.
If JOBZ = MagmaNoVec, then on exit the upper triangle (if UPLO=MagmaUpper)
or the lower triangle (if UPLO=MagmaLower) of A, including the
diagonal, is destroyed.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[in,out]
B REAL array, dimension (LDB, N)
On entry, the symmetric matrix B. If UPLO = MagmaUpper, the
leading N-by-N upper triangular part of B contains the
upper triangular part of the matrix B. If UPLO = MagmaLower,
the leading N-by-N lower triangular part of B contains
the lower triangular part of the matrix B.
\n
On exit, if INFO <= N, the part of B containing the matrix is
overwritten by the triangular factor U or L from the Cholesky
factorization B = U**T * U or B = L * L**T.
@param[in]
ldb INTEGER
The leading dimension of the array B. LDB >= max(1,N).
@param[out]
w REAL array, dimension (N)
If INFO = 0, the eigenvalues in ascending order.
@param[out]
work (workspace) REAL array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The length of the array WORK.
If N <= 1, LWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LWORK >= 2*N + N*NB.
If JOBZ = MagmaVec and N > 1, LWORK >= max( 2*N + N*NB, 1 + 6*N + 2*N**2 ).
NB can be obtained through magma_get_ssytrd_nb(N).
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal sizes of the WORK and IWORK
arrays, returns these values as the first entries of the WORK
and IWORK arrays, and no error message related to LWORK or
LIWORK is issued by XERBLA.
@param[out]
iwork (workspace) INTEGER array, dimension (MAX(1,LIWORK))
On exit, if INFO = 0, IWORK[0] returns the optimal LIWORK.
@param[in]
liwork INTEGER
The dimension of the array IWORK.
If N <= 1, LIWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LIWORK >= 1.
If JOBZ = MagmaVec and N > 1, LIWORK >= 3 + 5*N.
\n
If LIWORK = -1, then a workspace query is assumed; the
routine only calculates the optimal sizes of the WORK and
IWORK arrays, returns these values as the first entries of
the WORK and IWORK arrays, and no error message related to
LWORK or LIWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: SPOTRF or SSYEVD returned an error code:
<= N: if INFO = i and JOBZ = MagmaNoVec, then the algorithm
failed to converge; i off-diagonal elements of an
intermediate tridiagonal form did not converge to
zero;
if INFO = i and JOBZ = MagmaVec, then 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);
> N: if INFO = N + i, for 1 <= i <= N, then the leading
minor of order i of B is not positive definite.
The factorization of B could not be completed and
no eigenvalues or eigenvectors were computed.
Further Details
---------------
Based on contributions by
Mark Fahey, Department of Mathematics, Univ. of Kentucky, USA
Modified so that no backsubstitution is performed if SSYEVD fails to
converge (NEIG in old code could be greater than N causing out of
bounds reference to A - reported by Ralf Meyer). Also corrected the
description of INFO and the test on ITYPE. Sven, 16 Feb 05.
@ingroup magma_ssygv_driver
********************************************************************/
extern "C" magma_int_t
magma_ssygvd(
magma_int_t itype, magma_vec_t jobz, magma_uplo_t uplo, magma_int_t n,
float *A, magma_int_t lda,
float *B, magma_int_t ldb,
float *w,
float *work, magma_int_t lwork,
#ifdef COMPLEX
float *rwork, magma_int_t lrwork,
#endif
magma_int_t *iwork, magma_int_t liwork,
magma_int_t *info)
{
const char* uplo_ = lapack_uplo_const( uplo );
const char* jobz_ = lapack_vec_const( jobz );
float d_one = MAGMA_S_ONE;
float *dA=NULL, *dB=NULL;
magma_int_t ldda = magma_roundup( n, 32 );
magma_int_t lddb = ldda;
magma_int_t lower;
magma_trans_t trans;
magma_int_t wantz, lquery;
magma_int_t lwmin, liwmin;
wantz = (jobz == MagmaVec);
lower = (uplo == MagmaLower);
lquery = (lwork == -1 || liwork == -1);
*info = 0;
if (itype < 1 || itype > 3) {
*info = -1;
} else if (! (wantz || (jobz == MagmaNoVec))) {
*info = -2;
} else if (! (lower || (uplo == MagmaUpper))) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (lda < max(1,n)) {
*info = -6;
} else if (ldb < max(1,n)) {
*info = -8;
}
magma_int_t nb = magma_get_ssytrd_nb( n );
if ( n <= 1 ) {
lwmin = 1;
liwmin = 1;
}
else if ( wantz ) {
lwmin = max( 2*n + n*nb, 1 + 6*n + 2*n*n );
liwmin = 3 + 5*n;
}
else {
lwmin = 2*n + n*nb;
liwmin = 1;
}
work[0] = magma_smake_lwork( lwmin );
iwork[0] = liwmin;
if (lwork < lwmin && ! lquery) {
*info = -11;
} else if (liwork < liwmin && ! lquery) {
*info = -13;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (n == 0) {
return *info;
}
/* If matrix is very small, then just call LAPACK on CPU, no need for GPU */
if (n <= 128) {
lapackf77_ssygvd( &itype, jobz_, uplo_,
&n, A, &lda, B, &ldb,
w, work, &lwork,
iwork, &liwork, info );
return *info;
}
if (MAGMA_SUCCESS != magma_smalloc( &dA, n*ldda ) ||
MAGMA_SUCCESS != magma_smalloc( &dB, n*lddb )) {
magma_free( dA );
magma_free( dB );
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_queue_t queue;
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queue );
/* Form a Cholesky factorization of B. */
magma_ssetmatrix( n, n, B, ldb, dB, lddb, queue );
magma_ssetmatrix_async( n, n,
A, lda,
dA, ldda, queue );
magma_timer_t time=0;
timer_start( time );
magma_spotrf_gpu( uplo, n, dB, lddb, info );
if (*info != 0) {
*info = n + *info;
return *info;
}
timer_stop( time );
timer_printf( "time spotrf_gpu = %6.2f\n", time );
magma_queue_sync( queue );
magma_sgetmatrix_async( n, n,
dB, lddb,
B, ldb, queue );
timer_start( time );
/* Transform problem to standard eigenvalue problem and solve. */
magma_ssygst_gpu( itype, uplo, n, dA, ldda, dB, lddb, info );
timer_stop( time );
timer_printf( "time ssygst_gpu = %6.2f\n", time );
/* simple fix to be able to run bigger size.
* set dB=NULL so we know to re-allocate below
* TODO: have dwork here that will be used as dB and then passed to ssyevd.
*/
if (n > 5000) {
magma_queue_sync( queue );
magma_free( dB ); dB=NULL;
}
timer_start( time );
magma_ssyevd_gpu( jobz, uplo, n, dA, ldda, w, A, lda,
work, lwork, iwork, liwork, info );
timer_stop( time );
timer_printf( "time ssyevd_gpu = %6.2f\n", time );
if (wantz && *info == 0) {
timer_start( time );
/* allocate and copy dB back */
if (dB == NULL) {
if (MAGMA_SUCCESS != magma_smalloc( &dB, n*lddb ) ) {
magma_free( dA );
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_ssetmatrix( n, n, B, ldb, dB, lddb, queue );
}
/* Backtransform eigenvectors to the original problem. */
if (itype == 1 || itype == 2) {
/* For A*x=(lambda)*B*x and A*B*x=(lambda)*x;
backtransform eigenvectors: x = inv(L)'*y or inv(U)*y */
if (lower) {
trans = MagmaTrans;
} else {
trans = MagmaNoTrans;
}
magma_strsm( MagmaLeft, uplo, trans, MagmaNonUnit,
n, n, d_one, dB, lddb, dA, ldda, queue );
}
else if (itype == 3) {
/* For B*A*x=(lambda)*x;
backtransform eigenvectors: x = L*y or U'*y */
if (lower) {
trans = MagmaNoTrans;
} else {
trans = MagmaTrans;
}
magma_strmm( MagmaLeft, uplo, trans, MagmaNonUnit,
n, n, d_one, dB, lddb, dA, ldda, queue );
}
magma_sgetmatrix( n, n, dA, ldda, A, lda, queue );
timer_stop( time );
timer_printf( "time strsm/mm + getmatrix = %6.2f\n", time );
}
magma_queue_sync( queue );
magma_queue_destroy( queue );
work[0] = magma_smake_lwork( lwmin );
iwork[0] = liwmin;
magma_free( dA ); dA=NULL;
magma_free( dB ); dB=NULL;
return *info;
} /* magma_ssygvd */
extern "C" magma_int_t
magma_slauum_gpu(char uplo, magma_int_t n,
float *dA, magma_int_t ldda, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
SLAUUM computes the product U * U' or L' * L, where the triangular
factor U or L is stored in the upper or lower triangular part of
the array dA.
If UPLO = 'U' or 'u' then the upper triangle of the result is stored,
overwriting the factor U in dA.
If UPLO = 'L' or 'l' then the lower triangle of the result is stored,
overwriting the factor L in dA.
This is the blocked form of the algorithm, calling Level 3 BLAS.
Arguments
=========
UPLO (input) CHARACTER*1
Specifies whether the triangular factor stored in the array dA
is upper or lower triangular:
= 'U': Upper triangular
= 'L': Lower triangular
N (input) INTEGER
The order of the triangular factor U or L. N >= 0.
dA (input/output) DOUBLE PRECISION array on the GPU, dimension (LDDA,N)
On entry, the triangular factor U or L.
On exit, if UPLO = 'U', the upper triangle of dA is
overwritten with the upper triangle of the product U * U';
if UPLO = 'L', the lower triangle of dA is overwritten with
the lower triangle of the product L' * L.
LDDA (input) INTEGER
The leading dimension of the array A. LDDA >= max(1,N).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -k, the k-th argument had an illegal value
===================================================================== */
/* Local variables */
char uplo_[2] = {uplo, 0};
magma_int_t nb, i, ib;
float d_one = MAGMA_D_ONE;
float c_one = MAGMA_S_ONE;
float *work;
int upper = lapackf77_lsame(uplo_, "U");
*info = 0;
if ((! upper) && (! lapackf77_lsame(uplo_, "L")))
*info = -1;
else if (n < 0)
*info = -2;
else if (ldda < max(1,n))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
nb = magma_get_spotrf_nb(n);
if (MAGMA_SUCCESS != magma_smalloc_pinned( &work, nb*nb )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_queue_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
if (nb <= 1 || nb >= n)
{
magma_sgetmatrix( n, n, dA, ldda, work, n );
lapackf77_slauum(uplo_, &n, work, &n, info);
magma_ssetmatrix( n, n, work, n, dA, ldda );
}
else
{
if (upper)
{
/* Compute inverse of upper triangular matrix */
for (i=0; i < n; i += nb)
{
ib = min(nb, (n-i));
/* Compute the product U * U'. */
magma_strmm( MagmaRight, MagmaUpper,
MagmaTrans, MagmaNonUnit, i, ib,
c_one, dA(i,i), ldda, dA(0, i),ldda);
magma_sgetmatrix( ib, ib,
dA(i, i), ldda,
work, ib );
lapackf77_slauum(MagmaUpperStr, &ib, work, &ib, info);
magma_ssetmatrix( ib, ib,
work, ib,
dA(i, i), ldda );
if(i+ib < n)
{
magma_sgemm( MagmaNoTrans, MagmaTrans,
i, ib, (n-i-ib), c_one, dA(0,i+ib),
ldda, dA(i, i+ib), ldda, c_one,
dA(0,i), ldda);
magma_ssyrk( MagmaUpper, MagmaNoTrans, ib,(n-i-ib),
d_one, dA(i, i+ib), ldda,
d_one, dA(i, i), ldda);
}
}
}
else
{
/* Compute the product L' * L. */
for(i=0; i<n; i=i+nb)
{
ib=min(nb,(n-i));
magma_strmm( MagmaLeft, MagmaLower,
MagmaTrans, MagmaNonUnit, ib,
i, c_one, dA(i,i), ldda,
dA(i, 0),ldda);
magma_sgetmatrix( ib, ib,
dA(i, i), ldda,
work, ib );
lapackf77_slauum(MagmaLowerStr, &ib, work, &ib, info);
magma_ssetmatrix( ib, ib,
work, ib,
dA(i, i), ldda );
if((i+ib) < n)
{
magma_sgemm( MagmaTrans, MagmaNoTrans,
ib, i, (n-i-ib), c_one, dA( i+ib,i),
ldda, dA(i+ib, 0),ldda, c_one,
dA(i,0), ldda);
magma_ssyrk( MagmaLower, MagmaTrans, ib, (n-i-ib),
d_one, dA(i+ib, i), ldda,
d_one, dA(i, i), ldda);
}
}
}
}
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free_pinned( work );
return *info;
}
extern "C" magma_int_t
magma_sgessm_gpu( char storev, magma_int_t m, magma_int_t n, magma_int_t k, magma_int_t ib,
magma_int_t *ipiv,
float *dL1, magma_int_t lddl1,
float *dL, magma_int_t lddl,
float *dA, magma_int_t ldda,
magma_int_t *info)
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
Purpose
=======
SGESSM applies the factors L computed by SGETRF_INCPIV to
a real M-by-N tile A.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. N >= 0.
K (input) INTEGER
The number of columns of the matrix L. K >= 0.
IB (input) INTEGER
The inner-blocking size. IB >= 0.
IPIV (input) INTEGER array on the cpu.
The pivot indices array of size K as returned by
SGETRF_INCPIV.
dL1 (input) DOUBLE COMPLEX array, dimension(LDDL1, N)
The IB-by-K matrix in which is stored L^(-1) as returned by GETRF_INCPIV
LDDL1 (input) INTEGER
The leading dimension of the array L1. LDDL1 >= max(1,2*IB).
dL (input) DOUBLE COMPLEX array, dimension(LDDL, N)
The M-by-K lower triangular tile on the gpu.
LDDL (input) INTEGER
The leading dimension of the array L. LDDL >= max(1,M).
dA (input/output) DOUBLE COMPLEX array, dimension (LDDA, N)
On entry, the M-by-N tile A on the gpu.
On exit, updated by the application of L on the gpu.
===================================================================== */
#define AT(i,j) (dAT + (i)*ldda + (j) )
#define L(i,j) (dL + (i) + (j)*lddl )
#define dL1(j) (dL1 + (j)*lddl1)
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
int i, s, sb;
float *dAT;
/* Check arguments */
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (ldda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
if ( (storev == 'C') || (storev == 'c') ) {
magmablas_sgetmo_in( dA, dAT, ldda, m, n );
} else {
dAT = dA;
}
s = k / ib;
for(i = 0; i < k; i += ib) {
sb = min(ib, k-i);
magmablas_slaswp( n, dAT, ldda, i+1, i+sb, ipiv, 1 );
#ifndef WITHOUTTRTRI
magma_strmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n, sb,
c_one, dL1(i), lddl1,
AT(i, 0), ldda);
#else
magma_strsm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n, sb,
c_one, L( i, i), lddl,
AT(i, 0), ldda);
#endif
if ( (i+sb) < m) {
magma_sgemm( MagmaNoTrans, MagmaTrans,
n, m-(i+sb), sb,
c_neg_one, AT(i, 0), ldda,
L( i+sb, i), lddl,
c_one, AT(i+sb, 0), ldda );
}
}
if ( (storev == 'C') || (storev == 'c') ) {
magmablas_sgetmo_in( dA, dAT, ldda, m, n );
}
return *info;
/* End of MAGMA_SGETRF_GPU */
}
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, t1, t2;
float c_neg_one = MAGMA_S_NEG_ONE;
magma_int_t ione = 1;
const char trans[] = { 'N', 'C', 'T' };
const char uplo[] = { 'L', 'U' };
const char diag[] = { 'U', 'N' };
const char side[] = { 'L', 'R' };
float *A, *B, *C, *C2, *LU;
float *dA, *dB, *dC1, *dC2;
float alpha = MAGMA_S_MAKE( 0.5, 0.1 );
float beta = MAGMA_S_MAKE( 0.7, 0.2 );
float dalpha = 0.6;
float dbeta = 0.8;
float work[1], error, total_error;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t m, n, k, size, maxn, ld, info;
magma_int_t *piv;
magma_err_t err;
magma_opts opts;
parse_opts( argc, argv, &opts );
printf( "Compares magma wrapper function to cublas function; all diffs should be exactly 0.\n\n" );
total_error = 0.;
for( int i = 0; i < opts.ntest; ++i ) {
m = opts.msize[i];
n = opts.nsize[i];
k = opts.ksize[i];
printf("=========================================================================\n");
printf( "M %d, N %d, K %d\n", (int) m, (int) n, (int) k );
// allocate matrices
// over-allocate so they can be any combination of {m,n,k} x {m,n,k}.
maxn = max( max( m, n ), k );
ld = maxn;
size = maxn*maxn;
err = magma_malloc_cpu( (void**) &piv, maxn*sizeof(magma_int_t) ); assert( err == 0 );
err = magma_smalloc_pinned( &A, size ); assert( err == 0 );
err = magma_smalloc_pinned( &B, size ); assert( err == 0 );
err = magma_smalloc_pinned( &C, size ); assert( err == 0 );
err = magma_smalloc_pinned( &C2, size ); assert( err == 0 );
err = magma_smalloc_pinned( &LU, size ); assert( err == 0 );
err = magma_smalloc( &dA, size ); assert( err == 0 );
err = magma_smalloc( &dB, size ); assert( err == 0 );
err = magma_smalloc( &dC1, size ); assert( err == 0 );
err = magma_smalloc( &dC2, size ); assert( err == 0 );
// initialize matrices
size = maxn*maxn;
lapackf77_slarnv( &ione, ISEED, &size, A );
lapackf77_slarnv( &ione, ISEED, &size, B );
lapackf77_slarnv( &ione, ISEED, &size, C );
printf( "========== Level 1 BLAS ==========\n" );
// ----- test SSWAP
// swap 2nd and 3rd columns of dA, then copy to C2 and compare with A
assert( n >= 4 );
magma_ssetmatrix( m, n, A, ld, dA, ld );
magma_ssetmatrix( m, n, A, ld, dB, ld );
magma_sswap( m, dA(0,1), 1, dA(0,2), 1 );
magma_sswap( m, dB(0,1), 1, dB(0,2), 1 );
// check results, storing diff between magma and cuda calls in C2
cublasSaxpy( ld*n, c_neg_one, dA, 1, dB, 1 );
magma_sgetmatrix( m, n, dB, ld, C2, ld );
error = lapackf77_slange( "F", &m, &k, C2, &ld, work );
total_error += error;
printf( "sswap diff %.2g\n", error );
// ----- test ISAMAX
// get argmax of column of A
magma_ssetmatrix( m, k, A, ld, dA, ld );
error = 0;
for( int j = 0; j < k; ++j ) {
magma_int_t i1 = magma_isamax( m, dA(0,j), 1 );
magma_int_t i2 = cublasIsamax( m, dA(0,j), 1 );
assert( i1 == i2 );
error += abs( i1 - i2 );
}
total_error += error;
gflops = (float)m * k / 1e9;
printf( "isamax diff %.2g\n", error );
printf( "\n" );
printf( "========== Level 2 BLAS ==========\n" );
// ----- test SGEMV
// c = alpha*A*b + beta*c, with A m*n; b,c m or n-vectors
// try no-trans/trans
for( int ia = 0; ia < 3; ++ia ) {
magma_ssetmatrix( m, n, A, ld, dA, ld );
magma_ssetvector( maxn, B, 1, dB, 1 );
magma_ssetvector( maxn, C, 1, dC1, 1 );
magma_ssetvector( maxn, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_sgemv( trans[ia], m, n, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasSgemv( trans[ia], m, n, alpha, dA, ld, dB, 1, beta, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
size = (trans[ia] == 'N' ? m : n);
cublasSaxpy( size, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetvector( size, dC2, 1, C2, 1 );
error = lapackf77_slange( "F", &size, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_SGEMV( m, n ) / 1e9;
printf( "sgemv( %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
trans[ia], error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test SSYMV
// c = alpha*A*b + beta*c, with A m*m symmetric; b,c m-vectors
// try upper/lower
for( int iu = 0; iu < 2; ++iu ) {
magma_ssetmatrix( m, m, A, ld, dA, ld );
magma_ssetvector( m, B, 1, dB, 1 );
magma_ssetvector( m, C, 1, dC1, 1 );
magma_ssetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_ssymv( uplo[iu], m, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasSsymv( uplo[iu], m, alpha, dA, ld, dB, 1, beta, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasSaxpy( m, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_slange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_SSYMV( m ) / 1e9;
printf( "ssymv( %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test STRSV
// solve A*c = c, with A m*m triangular; c m-vector
// try upper/lower, no-trans/trans, unit/non-unit diag
// Factor A into LU to get well-conditioned triangles, else solve yields garbage.
// Still can give garbage if solves aren't consistent with LU factors,
// e.g., using unit diag for U, so copy lower triangle to upper triangle.
// Also used for trsm later.
lapackf77_slacpy( "Full", &maxn, &maxn, A, &ld, LU, &ld );
lapackf77_sgetrf( &maxn, &maxn, LU, &ld, piv, &info );
for( int j = 0; j < maxn; ++j ) {
for( int i = 0; i < j; ++i ) {
*LU(i,j) = *LU(j,i);
}
}
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
magma_ssetmatrix( m, m, LU, ld, dA, ld );
magma_ssetvector( m, C, 1, dC1, 1 );
magma_ssetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_strsv( uplo[iu], trans[it], diag[id], m, dA, ld, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasStrsv( uplo[iu], trans[it], diag[id], m, dA, ld, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasSaxpy( m, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_slange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_STRSM( MagmaLeft, m, 1 ) / 1e9;
printf( "strsv( %c, %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], diag[id], error, gflops/t1, gflops/t2 );
}}}
printf( "\n" );
printf( "========== Level 3 BLAS ==========\n" );
// ----- test SGEMM
// C = alpha*A*B + beta*C, with A m*k or k*m; B k*n or n*k; C m*n
// try combinations of no-trans/trans
for( int ia = 0; ia < 3; ++ia ) {
for( int ib = 0; ib < 3; ++ib ) {
bool nta = (trans[ia] == 'N');
bool ntb = (trans[ib] == 'N');
magma_ssetmatrix( (nta ? m : k), (nta ? m : k), A, ld, dA, ld );
magma_ssetmatrix( (ntb ? k : n), (ntb ? n : k), B, ld, dB, ld );
magma_ssetmatrix( m, n, C, ld, dC1, ld );
magma_ssetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_sgemm( trans[ia], trans[ib], m, n, k, alpha, dA, ld, dB, ld, beta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasSgemm( trans[ia], trans[ib], m, n, k, alpha, dA, ld, dB, ld, beta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasSaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_slange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_SGEMM( m, n, k ) / 1e9;
printf( "sgemm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
trans[ia], trans[ib], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test SSYMM
// C = alpha*A*B + beta*C (left) with A m*m symmetric; B,C m*n; or
// C = alpha*B*A + beta*C (right) with A n*n symmetric; B,C m*n
// try left/right, upper/lower
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
magma_ssetmatrix( m, m, A, ld, dA, ld );
magma_ssetmatrix( m, n, B, ld, dB, ld );
magma_ssetmatrix( m, n, C, ld, dC1, ld );
magma_ssetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ssymm( side[is], uplo[iu], m, n, alpha, dA, ld, dB, ld, beta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasSsymm( side[is], uplo[iu], m, n, alpha, dA, ld, dB, ld, beta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasSaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_slange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_SSYMM( side[is], m, n ) / 1e9;
printf( "ssymm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
side[is], uplo[iu], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test SSYRK
// C = alpha*A*A^H + beta*C (no-trans) with A m*k and C m*m symmetric; or
// C = alpha*A^H*A + beta*C (trans) with A k*m and C m*m symmetric
// try upper/lower, no-trans/trans
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
magma_ssetmatrix( n, k, A, ld, dA, ld );
magma_ssetmatrix( n, n, C, ld, dC1, ld );
magma_ssetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ssyrk( uplo[iu], trans[it], n, k, dalpha, dA, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasSsyrk( uplo[iu], trans[it], n, k, dalpha, dA, ld, dbeta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasSaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_slange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_SSYRK( k, n ) / 1e9;
printf( "ssyrk( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test SSYR2K
// C = alpha*A*B^H + ^alpha*B*A^H + beta*C (no-trans) with A,B n*k; C n*n symmetric; or
// C = alpha*A^H*B + ^alpha*B^H*A + beta*C (trans) with A,B k*n; C n*n symmetric
// try upper/lower, no-trans/trans
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
bool nt = (trans[it] == 'N');
magma_ssetmatrix( (nt ? n : k), (nt ? n : k), A, ld, dA, ld );
magma_ssetmatrix( n, n, C, ld, dC1, ld );
magma_ssetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ssyr2k( uplo[iu], trans[it], n, k, alpha, dA, ld, dB, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasSsyr2k( uplo[iu], trans[it], n, k, alpha, dA, ld, dB, ld, dbeta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasSaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_slange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_SSYR2K( k, n ) / 1e9;
printf( "ssyr2k( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test STRMM
// C = alpha*A*C (left) with A m*m triangular; C m*n; or
// C = alpha*C*A (right) with A n*n triangular; C m*n
// try left/right, upper/lower, no-trans/trans, unit/non-unit
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
bool left = (side[is] == 'L');
magma_ssetmatrix( (left ? m : n), (left ? m : n), A, ld, dA, ld );
magma_ssetmatrix( m, n, C, ld, dC1, ld );
magma_ssetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_strmm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasStrmm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasSaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_slange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_STRMM( side[is], m, n ) / 1e9;
printf( "strmm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// ----- test STRSM
// solve A*X = alpha*B (left) with A m*m triangular; B m*n; or
// solve X*A = alpha*B (right) with A n*n triangular; B m*n
// try left/right, upper/lower, no-trans/trans, unit/non-unit
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
bool left = (side[is] == 'L');
magma_ssetmatrix( (left ? m : n), (left ? m : n), LU, ld, dA, ld );
magma_ssetmatrix( m, n, C, ld, dC1, ld );
magma_ssetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_strsm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasStrsm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasSaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_slange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_STRSM( side[is], m, n ) / 1e9;
printf( "strsm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// cleanup
magma_free_cpu( piv );
magma_free_pinned( A );
magma_free_pinned( B );
magma_free_pinned( C );
magma_free_pinned( C2 );
magma_free_pinned( LU );
magma_free( dA );
magma_free( dB );
magma_free( dC1 );
magma_free( dC2 );
}
if ( total_error != 0. ) {
printf( "total error %.2g -- ought to be 0 -- some test failed (see above).\n",
total_error );
}
else {
printf( "all tests passed\n" );
}
TESTING_FINALIZE();
return 0;
}
/**
Purpose
-------
STRTRI computes the inverse of a real upper or lower triangular
matrix A.
This is the Level 3 BLAS version of the algorithm.
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: A is upper triangular;
- = MagmaLower: A is lower triangular.
@param[in]
diag magma_diag_t
- = MagmaNonUnit: A is non-unit triangular;
- = MagmaUnit: A is unit triangular.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
A REAL array, dimension (LDA,N)
On entry, the triangular matrix A. If UPLO = MagmaUpper, the
leading N-by-N upper triangular part of the array A contains
the upper triangular matrix, and the strictly lower
triangular part of A is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of the array A contains
the lower triangular matrix, and the strictly upper
triangular part of A is not referenced. If DIAG = MagmaUnit, the
diagonal elements of A are also not referenced and are
assumed to be 1.
On exit, the (triangular) inverse of the original matrix, in
the same storage format.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: if INFO = i, A(i,i) is exactly zero. The triangular
matrix is singular and its inverse cannot be computed.
@ingroup magma_sgesv_aux
********************************************************************/
extern "C" magma_int_t
magma_strtri(
magma_uplo_t uplo, magma_diag_t diag, magma_int_t n,
float *A, magma_int_t lda,
magma_int_t *info)
{
#define A(i, j) ( A + (i) + (j)*lda )
#define dA(i, j) (dA + (i) + (j)*ldda)
/* Local variables */
const char* uplo_ = lapack_uplo_const( uplo );
const char* diag_ = lapack_diag_const( diag );
magma_int_t ldda, nb, nn, j, jb;
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
float *dA;
int upper = (uplo == MagmaUpper);
int nounit = (diag == MagmaNonUnit);
*info = 0;
if (! upper && uplo != MagmaLower)
*info = -1;
else if (! nounit && diag != MagmaUnit)
*info = -2;
else if (n < 0)
*info = -3;
else if (lda < max(1,n))
*info = -5;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return */
if ( n == 0 )
return *info;
/* Check for singularity if non-unit */
if (nounit) {
for (j=0; j < n; ++j) {
if ( MAGMA_S_EQUAL( *A(j,j), c_zero )) {
*info = j+1; // Fortran index
return *info;
}
}
}
/* Determine the block size for this environment */
nb = magma_get_spotrf_nb(n);
ldda = ((n+31)/32)*32;
if (MAGMA_SUCCESS != magma_smalloc( &dA, (n)*ldda )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_queue_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
if (nb <= 1 || nb >= n)
lapackf77_strtri(uplo_, diag_, &n, A, &lda, info);
else {
if (upper) {
/* Compute inverse of upper triangular matrix */
for (j=0; j < n; j += nb) {
jb = min(nb, (n-j));
magma_ssetmatrix( jb, (n-j),
A(j, j), lda,
dA(j, j), ldda );
/* Compute rows 1:j-1 of current block column */
magma_strmm( MagmaLeft, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_one, dA(0,0), ldda, dA(0, j),ldda);
magma_strsm( MagmaRight, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_neg_one, dA(j,j), ldda, dA(0, j),ldda);
magma_sgetmatrix_async( jb, jb,
dA(j, j), ldda,
A(j, j), lda, stream[1] );
magma_sgetmatrix_async( j, jb,
dA(0, j), ldda,
A(0, j), lda, stream[0] );
magma_queue_sync( stream[1] );
/* Compute inverse of current diagonal block */
lapackf77_strtri(MagmaUpperStr, diag_, &jb, A(j,j), &lda, info);
magma_ssetmatrix( jb, jb,
A(j, j), lda,
dA(j, j), ldda );
}
}
else {
/* Compute inverse of lower triangular matrix */
nn=((n-1)/nb)*nb+1;
for (j=nn-1; j >= 0; j -= nb) {
jb=min(nb,(n-j));
if ((j+jb) < n) {
magma_ssetmatrix( (n-j), jb,
A(j, j), lda,
dA(j, j), ldda );
/* Compute rows j+jb:n of current block column */
magma_strmm( MagmaLeft, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_one, dA(j+jb,j+jb), ldda, dA(j+jb, j), ldda );
magma_strsm( MagmaRight, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_neg_one, dA(j,j), ldda, dA(j+jb, j), ldda );
magma_sgetmatrix_async( n-j-jb, jb,
dA(j+jb, j), ldda,
A(j+jb, j), lda, stream[1] );
magma_sgetmatrix_async( jb, jb,
dA(j,j), ldda,
A(j,j), lda, stream[0] );
magma_queue_sync( stream[0] );
}
/* Compute inverse of current diagonal block */
lapackf77_strtri(MagmaLowerStr, diag_, &jb, A(j,j), &lda, info);
magma_ssetmatrix( jb, jb,
A(j, j), lda,
dA(j, j), ldda );
}
}
}
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free( dA );
return *info;
}
/**
Purpose
-------
SGESSM applies the factors L computed by SGETRF_INCPIV to
a real M-by-N tile A.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in]
k INTEGER
The number of columns of the matrix L. K >= 0.
@param[in]
ib INTEGER
The inner-blocking size. IB >= 0.
@param[in]
ipiv INTEGER array on the cpu.
The pivot indices array of size K as returned by
SGETRF_INCPIV.
@param[in]
dL1 REAL array, dimension(LDDL1, N)
The IB-by-K matrix in which is stored L^(-1) as returned by GETRF_INCPIV
@param[in]
lddl1 INTEGER
The leading dimension of the array L1. LDDL1 >= max(1,2*IB).
@param[in]
dL REAL array, dimension(LDDL, N)
The M-by-K lower triangular tile on the gpu.
@param[in]
lddl INTEGER
The leading dimension of the array L. LDDL >= max(1,M).
@param[in,out]
dA REAL array, dimension (LDDA, N)
On entry, the M-by-N tile A on the gpu.
On exit, updated by the application of L on the gpu.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,M).
@ingroup magma_sgesv_tile
********************************************************************/
extern "C" magma_int_t
magma_sgessm_gpu( magma_order_t order, magma_int_t m, magma_int_t n, magma_int_t k, magma_int_t ib,
magma_int_t *ipiv,
float *dL1, magma_int_t lddl1,
float *dL, magma_int_t lddl,
float *dA, magma_int_t ldda,
magma_int_t *info)
{
#define AT(i,j) (dAT + (i)*ldda + (j) )
#define L(i,j) (dL + (i) + (j)*lddl )
#define dL1(j) (dL1 + (j)*lddl1)
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
int i, s, sb;
float *dAT;
/* Check arguments */
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (ldda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
if ( order == MagmaColMajor ) {
magmablas_sgetmo_in( dA, dAT, ldda, m, n );
} else {
dAT = dA;
}
s = k / ib;
for (i = 0; i < k; i += ib) {
sb = min(ib, k-i);
magmablas_slaswp( n, dAT, ldda, i+1, i+sb, ipiv, 1 );
#ifndef WITHOUTTRTRI
magma_strmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n, sb,
c_one, dL1(i), lddl1,
AT(i, 0), ldda);
#else
magma_strsm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n, sb,
c_one, L( i, i), lddl,
AT(i, 0), ldda);
#endif
if ( (i+sb) < m) {
magma_sgemm( MagmaNoTrans, MagmaTrans,
n, m-(i+sb), sb,
c_neg_one, AT(i, 0), ldda,
L( i+sb, i), lddl,
c_one, AT(i+sb, 0), ldda );
}
}
if ( order == MagmaColMajor ) {
magmablas_sgetmo_in( dA, dAT, ldda, m, n );
}
return *info;
} /* magma_sgessm_gpu */
extern "C" magma_int_t
magma_sgetrf_incpiv_gpu( char storev, magma_int_t m, magma_int_t n, magma_int_t ib,
float *hA, magma_int_t ldha, float *dA, magma_int_t ldda,
float *hL, magma_int_t ldhl, float *dL, magma_int_t lddl,
magma_int_t *ipiv,
float *dwork, magma_int_t lddwork,
magma_int_t *info)
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
Purpose
=======
SGETRF_INCPIV computes an LU factorization of a general M-by-N tile A
using partial pivoting with row interchanges.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 2.5 BLAS version of the algorithm.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. N >= 0.
IB (input) INTEGER
The inner-blocking size. IB >= 0.
hA (input,output) DOUBLE COMPLEX array, dimension(LDHA, N), on cpu.
On entry, only the M-by-IB first panel needs to be identical to dA(1..M, 1..IB).
On exit, the content is incomplete. Shouldn't be used.
LDHA (input) INTEGER
The leading dimension of the array hA. LDHA >= max(1,M).
dA (input,output) DOUBLE COMPLEX array, dimension(LDDA, N) , on gpu.
On entry, the M-by-N tile to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
LDDA (input) INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
hL (output) DOUBLE COMPLEX array, dimension(LDHL, min(M,N)), on vpu.
On exit, contains in the upper part the IB-by-K lower triangular tile,
and in the lower part IB-by-min(M,N) the inverse of the top part.
LDHL (input) INTEGER
The leading dimension of the array hL. LDHL >= max(1,2*IB).
dL (output) DOUBLE COMPLEX array, dimension(LDDL, K), on gpu.
On exit, contains in the upper part the IB-by-min(M,N) lower triangular tile,
and in the lower part IB-by-min(M,N) the inverse of the top part.
LDDL (input) INTEGER
The leading dimension of the array dL. LDDL >= max(1,2*IB).
IPIV (output) INTEGER array, dimension min(M,N), on the cpu.
The pivot indices array.
dWORK (output) DOUBLE COMPLEX array, dimension(LDDWORK, 2*IB), on gpu.
Workspace.
LDDWORK (input) INTEGER
The leading dimension of the array dWORK. LDDWORK >= max(NB, 1).
INFO (output) INTEGER
- PLASMA_SUCCESS successful exit
- < 0 if INFO = -k, the k-th argument had an illegal value
- > 0 if INFO = k, U(k,k) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
===================================================================== */
#define AT(i,j) (dAT + (i)*ib*ldda + (j)*ib)
#define hA(i,j) (hA + (i)*ib + (j)*ib*ldha)
#define hL(j) (hL + (j)*ib*ldhl )
#define hL2(j) (hL2 + (j)*ib*ldhl )
#define dL(j) (dL + (j)*ib*lddl )
#define dL2(j) (dL2 + (j)*ib*lddl )
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
magma_int_t iinfo;
magma_int_t maxm, mindim;
magma_int_t i, rows, cols, s, ii, sb;
float *dAT;
#ifndef WITHOUTTRTRI
float *dL2 = dL + ib;
float *hL2 = hL + ib;
#endif
/* Check arguments */
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (ldda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
/* Function Body */
mindim = min(m, n);
s = mindim / ib;
if ( ib >= mindim ) {
/* Use CPU code. */
lapackf77_sgetrf(&m, &n, hA, &ldha, ipiv, info);
#ifndef WITHOUTTRTRI
CORE_slacpy(PlasmaUpperLower, mindim, mindim,
(float*)hA, ldha,
(float*)hL2, ldhl );
CORE_strtri( PlasmaLower, PlasmaUnit, mindim,
(float*)hL2, ldhl, info );
if (*info != 0 ) {
fprintf(stderr, "ERROR, trtri returned with info = %d\n", *info);
}
magma_ssetmatrix( mindim, mindim, hL2, ldhl, dL2, lddl );
#endif
if ( (storev == 'R') || (storev == 'r') ) {
magma_ssetmatrix( m, n, hA, ldha, dwork, lddwork );
magmablas_stranspose( dA, ldda, dwork, lddwork, m, n );
} else {
magma_ssetmatrix( m, n, hA, ldha, dA, ldda );
}
}
else {
/* Use hybrid blocked code. */
maxm = ((m + 31)/32)*32;
if ( (storev == 'C') || (storev == 'c') ) {
magmablas_sgetmo_in( dA, dAT, ldda, m, n );
} else {
dAT = dA;
}
for( i=0; i<s; i++ )
{
ii = i * ib;
sb = min(ib, mindim-ii);
cols = maxm - ii;
if ( i>0 ){
// download i-th panel
magmablas_stranspose( dwork, maxm, AT(0, i), ldda, sb, m );
magma_sgetmatrix( m, sb, dwork, maxm, hA(0, i), ldha );
// make sure that gpu queue is empty
//magma_device_sync();
#ifndef WITHOUTTRTRI
magma_strmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n - (ii+sb), ib,
c_one, dL2(i-1), lddl,
AT(i-1,i+1), ldda );
#else
magma_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n - (ii+sb), ib,
c_one, AT(i-1,i-1), ldda,
AT(i-1,i+1), ldda );
#endif
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
n-(ii+sb), m-ii, ib,
c_neg_one, AT(i-1,i+1), ldda,
AT(i, i-1), ldda,
c_one, AT(i, i+1), ldda );
}
// do the cpu part
rows = m - ii;
lapackf77_sgetrf( &rows, &sb, hA(i, i), &ldha, ipiv+ii, &iinfo);
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + ii;
{
int j;
int fin = ii + sb;
for(j=ii ; j <fin; j++) {
ipiv[j] = ii + ipiv[j];
}
}
magmablas_slaswp( n-ii, AT(0, i), ldda, ii+1, ii+sb, ipiv, 1 );
#ifndef WITHOUTTRTRI
CORE_slacpy(PlasmaLower, sb, sb,
(float*)hA(i, i), ldha,
(float*)hL2(i), ldhl );
CORE_strtri( PlasmaLower, PlasmaUnit, sb,
(float*)hL2(i), ldhl, info );
if (*info != 0 ) {
fprintf(stderr, "ERROR, trtri returned with info = %d\n", *info);
}
magma_ssetmatrix( sb, sb, hL2(i), ldhl, dL2(i), lddl );
#endif
// upload i-th panel
magma_ssetmatrix( rows, sb, hA(i, i), ldha, dwork, cols );
magmablas_stranspose( AT(i,i), ldda, dwork, cols, rows, sb);
// do the small non-parallel computations
if ( s > (i+1) ) {
#ifndef WITHOUTTRTRI
magma_strmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
sb, sb,
c_one, dL2(i), lddl,
AT(i, i+1), ldda);
#else
magma_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
sb, sb,
c_one, AT(i, i ), ldda,
AT(i, i+1), ldda);
#endif
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
sb, m-(ii+sb), sb,
c_neg_one, AT(i, i+1), ldda,
AT(i+1, i ), ldda,
c_one, AT(i+1, i+1), ldda );
}
else {
/* Update of the last panel */
#ifndef WITHOUTTRTRI
magma_strmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n-mindim, sb,
c_one, dL2(i), lddl,
AT(i, i+1), ldda);
#else
magma_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-mindim, sb,
c_one, AT(i, i ), ldda,
AT(i, i+1), ldda);
#endif
/* m-(ii+sb) should be always 0 */
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
n-mindim, m-(ii+sb), sb,
c_neg_one, AT(i, i+1), ldda,
AT(i+1, i ), ldda,
c_one, AT(i+1, i+1), ldda );
}
}
if ( (storev == 'C') || (storev == 'c') ) {
magmablas_sgetmo_out( dA, dAT, ldda, m, n );
}
}
return *info;
}
