extern "C" magma_int_t
magma_zlarfb2_gpu(
magma_int_t m, magma_int_t n, magma_int_t k,
magmaDoubleComplex_const_ptr dV, magma_int_t lddv,
magmaDoubleComplex_const_ptr dT, magma_int_t lddt,
magmaDoubleComplex_ptr dC, magma_int_t lddc,
magmaDoubleComplex_ptr dwork, magma_int_t ldwork,
magma_queue_t queue )
{
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
if (m <= 0 || n <= 0)
return MAGMA_SUCCESS;
// W = C^H V
magma_zgemm( MagmaConjTrans, MagmaNoTrans,
//magmablas_zgemm_reduce(
n, k, m,
c_one, dC, lddc,
dV, lddv,
c_zero, dwork, ldwork, queue );
// W = W T^H = C^H V T^H
magma_ztrmm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
n, k,
c_one, dT, lddt,
dwork, ldwork, queue );
// C = C - V W^H = C - V T V^H C = (I - V T V^H) C = H C
magma_zgemm( MagmaNoTrans, MagmaConjTrans,
m, n, k,
c_neg_one, dV, lddv,
dwork, ldwork,
c_one, dC, lddc, queue );
return MAGMA_SUCCESS;
}
/**
Purpose
-------
ZSSSSM applies the LU factorization update from a complex
matrix formed by a lower triangular IB-by-K tile L1 on top of a
M2-by-K tile L2 to a second complex matrix formed by a M1-by-N1
tile A1 on top of a M2-by-N2 tile A2 (N1 == N2).
This is the right-looking Level 2.5 BLAS version of the algorithm.
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]
ib INTEGER
The inner-blocking size. IB >= 0.
@param[in]
NB INTEGER
The blocking size. NB >= 0.
@param[in,out]
hU COMPLEX_16 array, dimension(LDHU, N), on cpu.
On entry, the NB-by-N upper triangular tile hU.
On exit, the content is incomplete. Shouldn't be used.
@param[in]
ldhu INTEGER
The leading dimension of the array hU. LDHU >= max(1,NB).
@param[in,out]
dU COMPLEX_16 array, dimension(LDDU, N), on gpu.
On entry, the NB-by-N upper triangular tile dU identical to hU.
On exit, the new factor U from the factorization.
@param[in]
lddu INTEGER
The leading dimension of the array dU. LDDU >= max(1,NB).
@param[in,out]
hA COMPLEX_16 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.
@param[in]
ldha INTEGER
The leading dimension of the array hA. LDHA >= max(1,M).
@param[in,out]
dA COMPLEX_16 array, dimension(LDDA, N), on gpu.
On entry, the M-by-N tile to be factored.
On exit, the factor L from the factorization
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
@param[out]
hL COMPLEX_16 array, dimension(LDHL, K), on vpu.
On exit, contains in the upper part the IB-by-K lower triangular tile,
and in the lower part IB-by-K the inverse of the top part.
@param[in]
ldhl INTEGER
The leading dimension of the array hL. LDHL >= max(1,2*IB).
@param[out]
dL COMPLEX_16 array, dimension(LDDL, K), on gpu.
On exit, contains in the upper part the IB-by-K lower triangular tile,
and in the lower part IB-by-K the inverse of the top part.
@param[in]
lddl INTEGER
The leading dimension of the array dL. LDDL >= max(1,2*IB).
@param[out]
hWORK COMPLEX_16 array, dimension(LDHWORK, 2*IB), on cpu.
Workspace.
@param[in]
ldhwork INTEGER
The leading dimension of the array hWORK. LDHWORK >= max(NB, 1).
@param[out]
dWORK COMPLEX_16 array, dimension(LDDWORK, 2*IB), on gpu.
Workspace.
@param[in]
lddwork INTEGER
The leading dimension of the array dWORK. LDDWORK >= max(NB, 1).
@param[out]
ipiv INTEGER array on the cpu.
The pivot indices array of size K as returned by ZTSTRF
@param[out]
info 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.
@ingroup magma_zgesv_tile
********************************************************************/
extern "C" magma_int_t
magma_ztstrf_gpu(
magma_order_t order, magma_int_t m, magma_int_t n,
magma_int_t ib, magma_int_t nb,
magmaDoubleComplex *hU, magma_int_t ldhu,
magmaDoubleComplex_ptr dU, magma_int_t lddu,
magmaDoubleComplex *hA, magma_int_t ldha,
magmaDoubleComplex_ptr dA, magma_int_t ldda,
magmaDoubleComplex *hL, magma_int_t ldhl,
magmaDoubleComplex_ptr dL, magma_int_t lddl,
magma_int_t *ipiv,
magmaDoubleComplex *hwork, magma_int_t ldhwork,
magmaDoubleComplex_ptr dwork, magma_int_t lddwork,
magma_int_t *info)
{
#define UT(i,j) (dUT + (i)*ib*lddu + (j)*ib )
#define AT(i,j) (dAT + (i)*ib*ldda + (j)*ib )
#define L(i) (dL + (i)*ib*lddl )
#define L2(i) (dL2 + (i)*ib*lddl )
#define hU(i,j) (hU + (j)*ib*ldhu + (i)*ib )
#define hA(i,j) (hA + (j)*ib*ldha + (i)*ib )
#define hL(i) (hL + (i)*ib*ldhl )
#define hL2(i) (hL2 + (i)*ib*ldhl )
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
int iinfo = 0;
int maxm, mindim;
int i, j, im, s, ip, ii, sb, p = 1;
magmaDoubleComplex_ptr dAT, dUT;
magmaDoubleComplex_ptr dAp, dUp;
#ifndef WITHOUTTRTRI
magmaDoubleComplex_ptr dL2 = dL + ib;
magmaDoubleComplex *hL2 = hL + ib;
p = 2;
#endif
/* Check input arguments */
*info = 0;
if (m < 0) {
*info = -1;
}
else if (n < 0) {
*info = -2;
}
else if (ib < 0) {
*info = -3;
}
else if ((lddu < max(1,m)) && (m > 0)) {
*info = -6;
}
else if ((ldda < max(1,m)) && (m > 0)) {
*info = -8;
}
else if ((lddl < max(1,ib)) && (ib > 0)) {
*info = -10;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* quick return */
if ((m == 0) || (n == 0) || (ib == 0))
return *info;
ip = 0;
/* Function Body */
mindim = min(m, n);
s = mindim / ib;
if ( ib >= mindim ) {
/* Use CPU code. */
CORE_ztstrf(m, n, ib, nb,
(PLASMA_Complex64_t*)hU, ldhu,
(PLASMA_Complex64_t*)hA, ldha,
(PLASMA_Complex64_t*)hL, ldhl,
ipiv,
(PLASMA_Complex64_t*)hwork, ldhwork,
info);
#ifndef WITHOUTTRTRI
CORE_zlacpy( PlasmaUpperLower, mindim, mindim,
(PLASMA_Complex64_t*)hL, ldhl,
(PLASMA_Complex64_t*)hL2, ldhl );
CORE_ztrtri( PlasmaLower, PlasmaUnit, mindim,
(PLASMA_Complex64_t*)hL2, ldhl, info );
if (*info != 0 ) {
fprintf(stderr, "ERROR, trtri returned with info = %d\n", *info);
}
#endif
if ( order == MagmaRowMajor ) {
magma_zsetmatrix( m, n, hU, ldhu, dwork, lddwork );
magmablas_ztranspose( m, n, dwork, lddwork, dU, lddu );
magma_zsetmatrix( m, n, hA, ldha, dwork, lddwork );
magmablas_ztranspose( m, n, dwork, lddwork, dA, ldda );
} else {
magma_zsetmatrix( m, n, hU, ldhu, dU, lddu );
magma_zsetmatrix( m, n, hA, ldha, dA, ldda );
}
magma_zsetmatrix( p*ib, n, hL, ldhl, dL, lddl );
}
else {
/* Use hybrid blocked code. */
maxm = magma_roundup( m, 32 );
if ( order == MagmaColMajor ) {
magmablas_zgetmo_in( dU, dUT, lddu, m, n );
magmablas_zgetmo_in( dA, dAT, ldda, m, n );
} else {
dUT = dU; dAT = dA;
}
dAp = dwork;
dUp = dAp + ib*lddwork;
ip = 0;
for( i=0; i < s; i++ ) {
ii = i * ib;
sb = min(mindim-ii, ib);
if ( i > 0 ) {
// download i-th panel
magmablas_ztranspose( sb, ii, UT(0,i), lddu, dUp, lddu );
magmablas_ztranspose( sb, m, AT(0,i), ldda, dAp, ldda );
magma_zgetmatrix( ii, sb, dUp, lddu, hU(0, i), ldhu );
magma_zgetmatrix( m, sb, dAp, ldda, hA(0, i), ldha );
// make sure that gpu queue is empty
//magma_device_sync();
#ifndef WITHOUTTRTRI
magma_ztrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n-(ii+sb), ib,
c_one, L2(i-1), lddl,
UT(i-1, i+1), lddu);
#else
magma_ztrsm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n-(ii+sb), ib,
c_one, L(i-1), lddl,
UT(i-1, i+1), lddu);
#endif
magma_zgemm( MagmaNoTrans, MagmaNoTrans,
n-(ii+sb), m, ib,
c_neg_one, UT(i-1, i+1), lddu,
AT(0, i-1), ldda,
c_one, AT(0, i+1), ldda );
}
// do the cpu part
CORE_ztstrf(m, sb, ib, nb,
(PLASMA_Complex64_t*)hU(i, i), ldhu,
(PLASMA_Complex64_t*)hA(0, i), ldha,
(PLASMA_Complex64_t*)hL(i), ldhl,
ipiv+ii,
(PLASMA_Complex64_t*)hwork, ldhwork,
info);
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + ii;
// Need to swap betw U and A
#ifndef NOSWAPBLK
magmablas_zswapblk( MagmaRowMajor, n-(ii+sb),
UT(i, i+1), lddu,
AT(0, i+1), ldda,
1, sb, ipiv+ii, 1, nb );
for (j=0; j < ib; j++) {
im = ipiv[ip]-1;
if ( im == j ) {
ipiv[ip] += ii;
}
ip++;
}
#else
for (j=0; j < ib; j++) {
im = ipiv[ip]-1;
if ( im != (j) ) {
im = im - nb;
assert( (im >= 0) && (im < m) );
magmablas_zswap( n-(ii+sb), UT(i, i+1)+j*lddu, 1, AT(0, i+1)+im*ldda, 1 );
} else {
ipiv[ip] += ii;
}
ip++;
}
#endif
#ifndef WITHOUTTRTRI
CORE_zlacpy( PlasmaUpperLower, sb, sb,
(PLASMA_Complex64_t*)hL(i), ldhl,
(PLASMA_Complex64_t*)hL2(i), ldhl );
CORE_ztrtri( PlasmaLower, PlasmaUnit, sb,
(PLASMA_Complex64_t*)hL2(i), ldhl, info );
if (*info != 0 ) {
fprintf(stderr, "ERROR, trtri returned with info = %d\n", *info);
}
#endif
// upload i-th panel
magma_zsetmatrix( sb, sb, hU(i, i), ldhu, dUp, lddu );
magma_zsetmatrix( m, sb, hA(0, i), ldha, dAp, ldda );
magma_zsetmatrix( p*ib, sb, hL(i), ldhl, L(i), lddl );
magmablas_ztranspose( sb, sb, dUp, lddu, UT(i,i), lddu );
magmablas_ztranspose( m, sb, dAp, ldda, AT(0,i), ldda );
// make sure that gpu queue is empty
//magma_device_sync();
// do the small non-parallel computations
if ( s > (i+1) ) {
#ifndef WITHOUTTRTRI
magma_ztrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
sb, sb,
c_one, L2(i), lddl,
UT(i, i+1), lddu);
#else
magma_ztrsm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
sb, sb,
c_one, L(i), lddl,
UT(i, i+1), lddu);
#endif
magma_zgemm( MagmaNoTrans, MagmaNoTrans,
sb, m, sb,
c_neg_one, UT(i, i+1), lddu,
AT(0, i ), ldda,
c_one, AT(0, i+1), ldda );
}
else {
#ifndef WITHOUTTRTRI
magma_ztrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n-mindim, sb,
c_one, L2(i), lddl,
UT(i, i+1), lddu);
#else
magma_ztrsm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n-mindim, sb,
c_one, L(i), lddl,
UT(i, i+1), lddu);
#endif
magma_zgemm( MagmaNoTrans, MagmaNoTrans,
n-mindim, m, sb,
c_neg_one, UT(i, i+1), lddu,
AT(0, i ), ldda,
c_one, AT(0, i+1), ldda );
}
}
if ( order == MagmaColMajor ) {
magmablas_zgetmo_out( dU, dUT, lddu, m, n );
magmablas_zgetmo_out( dA, dAT, ldda, m, n );
}
}
return *info;
}
extern "C" magma_int_t
magma_ztrtri_gpu(
magma_uplo_t uplo, magma_diag_t diag, magma_int_t n,
magmaDoubleComplex_ptr dA, size_t dA_offset, magma_int_t ldda,
magma_queue_t queues[2],
magma_int_t *info)
{
/* -- clMAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date November 2014
Purpose
=======
ZTRTRI computes the inverse of a real upper or lower triangular
matrix dA.
This is the Level 3 BLAS version of the algorithm.
Arguments
=========
UPLO (input) CHARACTER*1
= 'U': A is upper triangular;
= 'L': A is lower triangular.
DIAG (input) CHARACTER*1
= 'N': A is non-unit triangular;
= 'U': A is unit triangular.
N (input) INTEGER
The order of the matrix A. N >= 0.
dA (input/output) COMPLEX_16 array ON THE GPU, dimension (LDDA,N)
On entry, the triangular matrix A. If UPLO = 'U', the
leading N-by-N upper triangular part of the array dA contains
the upper triangular matrix, and the strictly lower
triangular part of A is not referenced. If UPLO = 'L', the
leading N-by-N lower triangular part of the array dA contains
the lower triangular matrix, and the strictly upper
triangular part of A is not referenced. If DIAG = 'U', 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.
LDDA (input) INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
> 0: if INFO = i, dA(i,i) is exactly zero. The triangular
matrix is singular and its inverse cannot be computed.
(Singularity check is currently disabled.)
===================================================================== */
/* Local variables */
magma_int_t nb, nn, j, jb;
//magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex *work;
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 (ldda < max(1,n))
*info = -5;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Check for singularity if non-unit */
/* cannot do here with matrix dA on GPU -- need kernel */
/*
if (nounit) {
for (j=0; j < n; ++j) {
if ( MAGMA_Z_EQUAL( *dA(j,j), c_zero )) {
*info = j+1; // Fortran index
return *info;
}
}
}
*/
/* Determine the block size for this environment */
nb = magma_get_zpotrf_nb(n);
/* Create Queues */
//magma_queue_t queues[2];
//magma_device_t device[MagmaMaxGPUs];
//magma_int_t num = 0;
//magma_int_t err;
//
//err = magma_getdevices( device, MagmaMaxGPUs, &num );
//if ( err != 0 || num < 1 ) {
// fprintf( stderr, "magma_getdevices failed: %d\n", err );
// exit(-1);
//}
//err = magma_queue_create( device[0], &queues[0] );
//if ( err != 0 ) {
// fprintf( stderr, "magma_queue_create 0 failed: %d\n", err );
// exit(-1);
//}
//err = magma_queue_create( device[0], &queues[1] );
//if ( err != 0 ) {
// fprintf( stderr, "magma_queue_create 1 failed: %d\n", err );
// exit(-1);
//}
if (MAGMA_SUCCESS != magma_zmalloc_cpu( &work, nb*nb )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
if (nb <= 1 || nb >= n) {
magma_zgetmatrix( n, n, dA, dA_offset, ldda, work, n, queues[0] );
lapackf77_ztrtri( lapack_const(uplo), lapack_const(diag), &n, work, &n, info );
magma_zsetmatrix( n, n, work, n, dA, dA_offset, ldda, queues[0] );
}
else {
if (upper) {
/* Compute inverse of upper triangular matrix */
for (j=0; j < n; j += nb){
jb = min(nb, (n-j));
/* Compute rows 1:j-1 of current block column */
magma_ztrmm( MagmaLeft, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_one, dA(0,0), ldda, dA(0, j), ldda,
queues[0] );
magma_ztrsm( MagmaRight, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_neg_one, dA(j,j), ldda, dA(0, j), ldda,
queues[0] );
magma_zgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, jb, queues[1], NULL );
magma_queue_sync( queues[1] );
/* Compute inverse of current diagonal block */
lapackf77_ztrtri( MagmaUpperStr, lapack_const(diag), &jb, work, &jb, info );
/*
magma_zsetmatrix_async( jb, jb,
work, 0, jb,
dA(j, j), ldda, queues[0], NULL );
*/
magma_zsetmatrix( jb, jb,
work, jb,
dA(j, j), ldda, queues[0] );
}
}
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){
/* Compute rows j+jb:n of current block column */
magma_ztrmm( MagmaLeft, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_one, dA(j+jb,j+jb), ldda, dA(j+jb, j), ldda,
queues[0] );
magma_ztrsm( MagmaRight, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_neg_one, dA(j,j), ldda, dA(j+jb, j), ldda,
queues[0] );
}
magma_zgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, jb, queues[1], NULL );
magma_queue_sync( queues[1] );
/* Compute inverse of current diagonal block */
lapackf77_ztrtri( MagmaLowerStr, lapack_const(diag), &jb, work, &jb, info );
/*
magma_zsetmatrix_async( jb, jb,
work, 0, jb,
dA(j, j), ldda, queues[0], NULL );
*/
magma_zsetmatrix( jb, jb,
work, jb,
dA(j, j), ldda, queues[0] );
}
}
}
//magma_queue_destroy( queues[0] );
//magma_queue_destroy( queues[1] );
magma_free_cpu( work );
return *info;
}
/**
Purpose
-------
ZHEGVDX_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
---------
@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]
range magma_range_t
- = MagmaRangeAll: all eigenvalues will be found.
- = MagmaRangeV: all eigenvalues in the half-open interval (VL,VU]
will be found.
- = MagmaRangeI: the IL-th through IU-th eigenvalues will be found.
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper 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 COMPLEX_16 array, dimension (LDA, N)
On entry, the Hermitian matrix A. If UPLO = MagmaUpper, the
leading N-by-N upper triangular part of A contains the
upper triangular part of the matrix A. If UPLO = MagmaLower,
the leading N-by-N lower triangular part of A contains
the lower triangular part of the matrix A.
\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**H*B*Z = I;
if ITYPE = 3, Z**H*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 COMPLEX_16 array, dimension (LDB, N)
On entry, the Hermitian 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**H*U or B = L*L**H.
@param[in]
ldb INTEGER
The leading dimension of the array B. LDB >= max(1,N).
@param[in]
vl DOUBLE PRECISION
@param[in]
vu DOUBLE PRECISION
If RANGE=MagmaRangeV, the lower and upper bounds of the interval to
be searched for eigenvalues. VL < VU.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeI.
@param[in]
il INTEGER
@param[in]
iu INTEGER
If RANGE=MagmaRangeI, the indices (in ascending order) of the
smallest and largest eigenvalues to be returned.
1 <= IL <= IU <= N, if N > 0; IL = 1 and IU = 0 if N = 0.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeV.
@param[out]
m INTEGER
The total number of eigenvalues found. 0 <= M <= N.
If RANGE = MagmaRangeAll, M = N, and if RANGE = MagmaRangeI, M = IU-IL+1.
@param[out]
w DOUBLE PRECISION array, dimension (N)
If INFO = 0, the eigenvalues in ascending order.
@param[out]
work (workspace) COMPLEX_16 array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) 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 >= LQ2 + N * (NB + 1).
If JOBZ = MagmaVec and N > 1, LWORK >= LQ2 + 2*N + N**2.
where LQ2 is the size needed to store the Q2 matrix
and is returned by magma_bulge_get_lq2.
\n
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.
@param[out]
rwork (workspace) DOUBLE PRECISION array, dimension (MAX(1,LRWORK))
On exit, if INFO = 0, RWORK(1) returns the optimal LRWORK.
@param[in]
lrwork INTEGER
The dimension of the array RWORK.
If N <= 1, LRWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LRWORK >= N.
If JOBZ = MagmaVec and N > 1, LRWORK >= 1 + 5*N + 2*N**2.
\n
If LRWORK = -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.
@param[out]
iwork (workspace) INTEGER array, dimension (MAX(1,LIWORK))
On exit, if INFO = 0, IWORK(1) 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, 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.
@param[out]
info 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 = 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 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.
@ingroup magma_zhegv_driver
********************************************************************/
extern "C" magma_int_t
magma_zhegvdx_2stage(magma_int_t itype, magma_vec_t jobz, magma_range_t range, magma_uplo_t uplo, magma_int_t n,
magmaDoubleComplex *A, magma_int_t lda,
magmaDoubleComplex *B, magma_int_t ldb,
double vl, double vu, magma_int_t il, magma_int_t iu,
magma_int_t *m, double *w,
magmaDoubleComplex *work, magma_int_t lwork,
double *rwork, magma_int_t lrwork,
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 );
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex *dA;
magmaDoubleComplex *dB;
magma_int_t ldda = n;
magma_int_t lddb = n;
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;
magma_int_t lrwmin;
magma_queue_t stream;
magma_queue_create( &stream );
/* determine the number of threads */
magma_int_t parallel_threads = magma_get_parallel_numthreads();
wantz = (jobz == MagmaVec);
lower = (uplo == MagmaLower);
alleig = (range == MagmaRangeAll);
valeig = (range == MagmaRangeV);
indeig = (range == MagmaRangeI);
lquery = (lwork == -1 || lrwork == -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_zbulge_nb(n, parallel_threads);
magma_int_t lq2 = magma_zbulge_get_lq2(n, parallel_threads);
if (wantz) {
lwmin = lq2 + 2 * n + n * n;
lrwmin = 1 + 5 * n + 2 * n * n;
liwmin = 5 * n + 3;
} else {
lwmin = lq2 + n * (nb + 1);
lrwmin = n;
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_dlamch("Epsilon");
work[0] = MAGMA_Z_MAKE( lwmin * one_eps, 0.); // round up
rwork[0] = lrwmin * one_eps;
iwork[0] = liwmin;
if (lwork < lwmin && ! lquery) {
*info = -17;
} else if (lrwork < lrwmin && ! lquery) {
*info = -19;
} else if (liwork < liwmin && ! lquery) {
*info = -21;
}
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_zhegvd(&itype, jobz_, uplo_,
&n, A, &lda, B, &ldb,
w, work, &lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, &lrwork,
#endif
iwork, &liwork, info);
*m = n;
return *info;
}
// TODO: fix memory leak
if (MAGMA_SUCCESS != magma_zmalloc( &dA, n*ldda ) ||
MAGMA_SUCCESS != magma_zmalloc( &dB, n*lddb )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/* Form a Cholesky factorization of B. */
magma_zsetmatrix( n, n, B, ldb, dB, lddb );
magma_zsetmatrix_async( n, n,
A, lda,
dA, ldda, stream );
magma_timer_t time=0;
timer_start( time );
magma_zpotrf_gpu(uplo, n, dB, lddb, info);
if (*info != 0) {
*info = n + *info;
return *info;
}
timer_stop( time );
timer_printf( "time zpotrf_gpu = %6.2f\n", time );
magma_queue_sync( stream );
magma_zgetmatrix_async( n, n,
dB, lddb,
B, ldb, stream );
/* Transform problem to standard eigenvalue problem and solve. */
timer_start( time );
magma_zhegst_gpu(itype, uplo, n, dA, ldda, dB, lddb, info);
timer_stop( time );
timer_printf( "time zhegst_gpu = %6.2f\n", time );
magma_zgetmatrix( n, n, dA, ldda, A, lda );
magma_queue_sync( stream );
magma_free( dA );
magma_free( dB );
timer_start( time );
magma_zheevdx_2stage(jobz, range, uplo, n, A, lda, vl, vu, il, iu, m, w, work, lwork, rwork, lrwork, iwork, liwork, info);
timer_stop( time );
timer_printf( "time zheevdx_2stage = %6.2f\n", time );
if (wantz && *info == 0) {
// TODO fix memory leak
if (MAGMA_SUCCESS != magma_zmalloc( &dA, n*ldda ) ||
MAGMA_SUCCESS != magma_zmalloc( &dB, n*lddb )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
timer_start( time );
magma_zsetmatrix( n, *m, A, lda, dA, ldda );
magma_zsetmatrix( 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 = MagmaConjTrans;
} else {
trans = MagmaNoTrans;
}
magma_ztrsm(MagmaLeft, uplo, trans, MagmaNonUnit, n, *m, c_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 = MagmaConjTrans;
}
magma_ztrmm(MagmaLeft, uplo, trans, MagmaNonUnit, n, *m, c_one, dB, lddb, dA, ldda);
}
magma_zgetmatrix( n, *m, dA, ldda, A, lda );
timer_stop( time );
timer_printf( "time trsm/mm + getmatrix = %6.2f\n", time );
magma_free( dA );
magma_free( dB );
}
magma_queue_destroy( stream );
work[0] = MAGMA_Z_MAKE( lwmin * one_eps, 0.); // round up
rwork[0] = lrwmin * one_eps;
iwork[0] = liwmin;
return *info;
} /* magma_zhegvdx_2stage */
/**
Purpose
-------
ZTRTRI 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 COMPLEX_16 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_zgesv_aux
********************************************************************/
extern "C" magma_int_t
magma_ztrtri(
magma_uplo_t uplo, magma_diag_t diag, magma_int_t n,
magmaDoubleComplex *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;
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex *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_Z_EQUAL( *A(j,j), c_zero )) {
*info = j+1; // Fortran index
return *info;
}
}
}
/* Determine the block size for this environment */
nb = magma_get_zpotrf_nb(n);
ldda = ((n+31)/32)*32;
if (MAGMA_SUCCESS != magma_zmalloc( &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_ztrtri(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_zsetmatrix( jb, (n-j),
A(j, j), lda,
dA(j, j), ldda );
/* Compute rows 1:j-1 of current block column */
magma_ztrmm( MagmaLeft, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_one, dA(0,0), ldda, dA(0, j),ldda);
magma_ztrsm( MagmaRight, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_neg_one, dA(j,j), ldda, dA(0, j),ldda);
magma_zgetmatrix_async( jb, jb,
dA(j, j), ldda,
A(j, j), lda, stream[1] );
magma_zgetmatrix_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_ztrtri(MagmaUpperStr, diag_, &jb, A(j,j), &lda, info);
magma_zsetmatrix( 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_zsetmatrix( (n-j), jb,
A(j, j), lda,
dA(j, j), ldda );
/* Compute rows j+jb:n of current block column */
magma_ztrmm( MagmaLeft, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_one, dA(j+jb,j+jb), ldda, dA(j+jb, j), ldda );
magma_ztrsm( MagmaRight, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_neg_one, dA(j,j), ldda, dA(j+jb, j), ldda );
magma_zgetmatrix_async( n-j-jb, jb,
dA(j+jb, j), ldda,
A(j+jb, j), lda, stream[1] );
magma_zgetmatrix_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_ztrtri(MagmaLowerStr, diag_, &jb, A(j,j), &lda, info);
magma_zsetmatrix( 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;
}
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, t1, t2;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t ione = 1;
magma_trans_t trans[] = { MagmaNoTrans, MagmaConjTrans, MagmaTrans };
magma_uplo_t uplo [] = { MagmaLower, MagmaUpper };
magma_diag_t diag [] = { MagmaUnit, MagmaNonUnit };
magma_side_t side [] = { MagmaLeft, MagmaRight };
magmaDoubleComplex *A, *B, *C, *C2, *LU;
magmaDoubleComplex *dA, *dB, *dC1, *dC2;
magmaDoubleComplex alpha = MAGMA_Z_MAKE( 0.5, 0.1 );
magmaDoubleComplex beta = MAGMA_Z_MAKE( 0.7, 0.2 );
double dalpha = 0.6;
double dbeta = 0.8;
double 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_int_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 itest = 0; itest < opts.ntest; ++itest ) {
m = opts.msize[itest];
n = opts.nsize[itest];
k = opts.ksize[itest];
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 = max( 1, maxn );
size = ld*maxn;
err = magma_malloc_cpu( (void**) &piv, maxn*sizeof(magma_int_t) ); assert( err == 0 );
err = magma_zmalloc_pinned( &A, size ); assert( err == 0 );
err = magma_zmalloc_pinned( &B, size ); assert( err == 0 );
err = magma_zmalloc_pinned( &C, size ); assert( err == 0 );
err = magma_zmalloc_pinned( &C2, size ); assert( err == 0 );
err = magma_zmalloc_pinned( &LU, size ); assert( err == 0 );
err = magma_zmalloc( &dA, size ); assert( err == 0 );
err = magma_zmalloc( &dB, size ); assert( err == 0 );
err = magma_zmalloc( &dC1, size ); assert( err == 0 );
err = magma_zmalloc( &dC2, size ); assert( err == 0 );
// initialize matrices
size = maxn*maxn;
lapackf77_zlarnv( &ione, ISEED, &size, A );
lapackf77_zlarnv( &ione, ISEED, &size, B );
lapackf77_zlarnv( &ione, ISEED, &size, C );
printf( "========== Level 1 BLAS ==========\n" );
// ----- test ZSWAP
// swap columns 2 and 3 of dA, then copy to C2 and compare with A
if ( n >= 3 ) {
magma_zsetmatrix( m, n, A, ld, dA, ld );
magma_zsetmatrix( m, n, A, ld, dB, ld );
magma_zswap( m, dA(0,1), 1, dA(0,2), 1 );
magma_zswap( m, dB(0,1), 1, dB(0,2), 1 );
// check results, storing diff between magma and cuda calls in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dA, 1, dB, 1 );
magma_zgetmatrix( m, n, dB, ld, C2, ld );
error = lapackf77_zlange( "F", &m, &k, C2, &ld, work );
total_error += error;
printf( "zswap diff %.2g\n", error );
}
else {
printf( "zswap skipped for n < 3\n" );
}
// ----- test IZAMAX
// get argmax of column of A
magma_zsetmatrix( m, k, A, ld, dA, ld );
error = 0;
for( int j = 0; j < k; ++j ) {
magma_int_t i1 = magma_izamax( m, dA(0,j), 1 );
int i2; // NOT magma_int_t, for cublas
cublasIzamax( handle, m, dA(0,j), 1, &i2 );
// todo need sync here?
assert( i1 == i2 );
error += abs( i1 - i2 );
}
total_error += error;
gflops = (double)m * k / 1e9;
printf( "izamax diff %.2g\n", error );
printf( "\n" );
printf( "========== Level 2 BLAS ==========\n" );
// ----- test ZGEMV
// 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_zsetmatrix( m, n, A, ld, dA, ld );
magma_zsetvector( maxn, B, 1, dB, 1 );
magma_zsetvector( maxn, C, 1, dC1, 1 );
magma_zsetvector( maxn, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_zgemv( trans[ia], m, n, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZgemv( handle, cublas_trans_const(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] == MagmaNoTrans ? m : n);
cublasZaxpy( handle, size, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetvector( size, dC2, 1, C2, 1 );
error = lapackf77_zlange( "F", &size, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZGEMV( m, n ) / 1e9;
printf( "zgemv( %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_trans_const(trans[ia]), error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test ZHEMV
// 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_zsetmatrix( m, m, A, ld, dA, ld );
magma_zsetvector( m, B, 1, dB, 1 );
magma_zsetvector( m, C, 1, dC1, 1 );
magma_zsetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_zhemv( uplo[iu], m, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZhemv( handle, cublas_uplo_const(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
cublasZaxpy( handle, m, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_zlange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZHEMV( m ) / 1e9;
printf( "zhemv( %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test ZTRSV
// 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_zlacpy( "Full", &maxn, &maxn, A, &ld, LU, &ld );
lapackf77_zgetrf( &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_zsetmatrix( m, m, LU, ld, dA, ld );
magma_zsetvector( m, C, 1, dC1, 1 );
magma_zsetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_ztrsv( uplo[iu], trans[it], diag[id], m, dA, ld, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZtrsv( handle, cublas_uplo_const(uplo[iu]), cublas_trans_const(trans[it]),
cublas_diag_const(diag[id]), m, dA, ld, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, m, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_zlange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZTRSM( MagmaLeft, m, 1 ) / 1e9;
printf( "ztrsv( %c, %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]), lapacke_diag_const(diag[id]),
error, gflops/t1, gflops/t2 );
}}}
printf( "\n" );
printf( "========== Level 3 BLAS ==========\n" );
// ----- test ZGEMM
// 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] == MagmaNoTrans);
bool ntb = (trans[ib] == MagmaNoTrans);
magma_zsetmatrix( (nta ? m : k), (nta ? m : k), A, ld, dA, ld );
magma_zsetmatrix( (ntb ? k : n), (ntb ? n : k), B, ld, dB, ld );
magma_zsetmatrix( m, n, C, ld, dC1, ld );
magma_zsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_zgemm( 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 );
cublasZgemm( handle, cublas_trans_const(trans[ia]), cublas_trans_const(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
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZGEMM( m, n, k ) / 1e9;
printf( "zgemm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_trans_const(trans[ia]), lapacke_trans_const(trans[ib]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test ZHEMM
// 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_zsetmatrix( m, m, A, ld, dA, ld );
magma_zsetmatrix( m, n, B, ld, dB, ld );
magma_zsetmatrix( m, n, C, ld, dC1, ld );
magma_zsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_zhemm( 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 );
cublasZhemm( handle, cublas_side_const(side[is]), cublas_uplo_const(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
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZHEMM( side[is], m, n ) / 1e9;
printf( "zhemm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_side_const(side[is]), lapacke_uplo_const(uplo[iu]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test ZHERK
// 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_zsetmatrix( n, k, A, ld, dA, ld );
magma_zsetmatrix( n, n, C, ld, dC1, ld );
magma_zsetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_zherk( uplo[iu], trans[it], n, k, dalpha, dA, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZherk( handle, cublas_uplo_const(uplo[iu]), cublas_trans_const(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
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZHERK( k, n ) / 1e9;
printf( "zherk( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test ZHER2K
// 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] == MagmaNoTrans);
magma_zsetmatrix( (nt ? n : k), (nt ? n : k), A, ld, dA, ld );
magma_zsetmatrix( n, n, C, ld, dC1, ld );
magma_zsetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_zher2k( 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 );
cublasZher2k( handle, cublas_uplo_const(uplo[iu]), cublas_trans_const(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
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZHER2K( k, n ) / 1e9;
printf( "zher2k( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test ZTRMM
// 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] == MagmaLeft);
magma_zsetmatrix( (left ? m : n), (left ? m : n), A, ld, dA, ld );
magma_zsetmatrix( m, n, C, ld, dC1, ld );
magma_zsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ztrmm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
// note cublas does trmm out-of-place (i.e., adds output matrix C),
// but allows C=B to do in-place.
t2 = magma_sync_wtime( 0 );
cublasZtrmm( handle, cublas_side_const(side[is]), cublas_uplo_const(uplo[iu]),
cublas_trans_const(trans[it]), cublas_diag_const(diag[id]),
m, n, &alpha, dA, ld, dC2, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZTRMM( side[is], m, n ) / 1e9;
printf( "ztrmm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// ----- test ZTRSM
// 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] == MagmaLeft);
magma_zsetmatrix( (left ? m : n), (left ? m : n), LU, ld, dA, ld );
magma_zsetmatrix( m, n, C, ld, dC1, ld );
magma_zsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ztrsm( 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 );
cublasZtrsm( handle, cublas_side_const(side[is]), cublas_uplo_const(uplo[iu]),
cublas_trans_const(trans[it]), cublas_diag_const(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
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZTRSM( side[is], m, n ) / 1e9;
printf( "ztrsm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(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 );
fflush( stdout );
}
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();
int status = (total_error != 0.);
return status;
}
/**
Purpose
-------
ZHEGVD 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.
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]
nrgpu INTEGER
Number of GPUs to use.
@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 COMPLEX_16 array, dimension (LDA, N)
On entry, the Hermitian matrix A. If UPLO = MagmaUpper, the
leading N-by-N upper triangular part of A contains the
upper triangular part of the matrix A. If UPLO = MagmaLower,
the leading N-by-N lower triangular part of A contains
the lower triangular part of the matrix A.
\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**H*B*Z = I;
if ITYPE = 3, Z**H*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 COMPLEX_16 array, dimension (LDB, N)
On entry, the Hermitian 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**H*U or B = L*L**H.
@param[in]
ldb INTEGER
The leading dimension of the array B. LDB >= max(1,N).
@param[out]
w DOUBLE PRECISION array, dimension (N)
If INFO = 0, the eigenvalues in ascending order.
@param[out]
work (workspace) COMPLEX_16 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 >= N + N*NB.
If JOBZ = MagmaVec and N > 1, LWORK >= max( N + N*NB, 2*N + N**2 ).
NB can be obtained through magma_get_zhetrd_nb(N).
\n
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.
@param[out]
rwork (workspace) DOUBLE PRECISION array, dimension (MAX(1,LRWORK))
On exit, if INFO = 0, RWORK[0] returns the optimal LRWORK.
@param[in]
lrwork INTEGER
The dimension of the array RWORK.
If N <= 1, LRWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LRWORK >= N.
If JOBZ = MagmaVec and N > 1, LRWORK >= 1 + 5*N + 2*N**2.
\n
If LRWORK = -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.
@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, 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.
@param[out]
info 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 = 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 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.
@ingroup magma_zhegv_driver
********************************************************************/
extern "C" magma_int_t
magma_zhegvd_m(magma_int_t nrgpu, magma_int_t itype, magma_vec_t jobz, magma_uplo_t uplo, magma_int_t n,
magmaDoubleComplex *A, magma_int_t lda, magmaDoubleComplex *B, magma_int_t ldb,
double *w, magmaDoubleComplex *work, magma_int_t lwork,
double *rwork, magma_int_t lrwork,
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 );
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magma_int_t lower;
magma_trans_t trans;
magma_int_t wantz;
magma_int_t lquery;
magma_int_t lwmin;
magma_int_t liwmin;
magma_int_t lrwmin;
magma_queue_t stream;
magma_queue_create( &stream );
wantz = (jobz == MagmaVec);
lower = (uplo == MagmaLower);
lquery = (lwork == -1 || lrwork == -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_zhetrd_nb( n );
if ( n <= 1 ) {
lwmin = 1;
lrwmin = 1;
liwmin = 1;
}
else if ( wantz ) {
lwmin = max( n + n*nb, 2*n + n*n );
lrwmin = 1 + 5*n + 2*n*n;
liwmin = 3 + 5*n;
}
else {
lwmin = n + n*nb;
lrwmin = n;
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_dlamch("Epsilon");
work[0] = MAGMA_Z_MAKE( lwmin * one_eps, 0.); // round up
rwork[0] = lrwmin * one_eps;
iwork[0] = liwmin;
if (lwork < lwmin && ! lquery) {
*info = -11;
} else if (lrwork < lrwmin && ! lquery) {
*info = -13;
} else if (liwork < liwmin && ! lquery) {
*info = -15;
}
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_zhegvd(&itype, jobz_, uplo_,
&n, A, &lda, B, &ldb,
w, work, &lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, &lrwork,
#endif
iwork, &liwork, info);
return *info;
}
magma_timer_t time=0;
timer_start( time );
magma_zpotrf_m(nrgpu, uplo, n, B, ldb, info);
if (*info != 0) {
*info = n + *info;
return *info;
}
timer_stop( time );
timer_printf( "time zpotrf = %6.2f\n", time );
timer_start( time );
/* Transform problem to standard eigenvalue problem and solve. */
magma_zhegst_m(nrgpu, itype, uplo, n, A, lda, B, ldb, info);
timer_stop( time );
timer_printf( "time zhegst = %6.2f\n", time );
timer_start( time );
magma_zheevd_m(nrgpu, jobz, uplo, n, A, lda, w, work, lwork, rwork, lrwork, iwork, liwork, info);
timer_stop( time );
timer_printf( "time zheevd = %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 = MagmaConjTrans;
} else {
trans = MagmaNoTrans;
}
magma_ztrsm_m(nrgpu, MagmaLeft, uplo, trans, MagmaNonUnit,
n, n, 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 = MagmaConjTrans;
}
printf("--- the multi GPU version is falling back to 1 GPU to perform the last TRMM since there is no TRMM_mgpu --- \n");
magmaDoubleComplex *dA=NULL, *dB=NULL;
magma_int_t ldda = n;
magma_int_t lddb = n;
if (MAGMA_SUCCESS != magma_zmalloc( &dB, n*lddb ) ) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
if (MAGMA_SUCCESS != magma_zmalloc( &dA, n*ldda ) ) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_zsetmatrix( n, n, B, ldb, dB, lddb );
magma_zsetmatrix( n, n, A, lda, dA, ldda );
magma_ztrmm(MagmaLeft, uplo, trans, MagmaNonUnit,
n, n, c_one, dB, lddb, dA, ldda);
magma_zgetmatrix( n, n, dA, ldda, A, lda );
}
timer_stop( time );
timer_printf( "time setmatrices trsm/mm + getmatrices = %6.2f\n", time );
}
work[0] = MAGMA_Z_MAKE( lwmin * one_eps, 0.); // round up
rwork[0] = lrwmin * one_eps;
iwork[0] = liwmin;
return *info;
} /* magma_zhegvd_m */
extern "C" magma_int_t
magma_zlarfb_gpu( char side, char trans, char direct, char storev,
magma_int_t m, magma_int_t n, magma_int_t k,
const magmaDoubleComplex *dV, magma_int_t ldv,
const magmaDoubleComplex *dT, magma_int_t ldt,
magmaDoubleComplex *dC, magma_int_t ldc,
magmaDoubleComplex *dwork, magma_int_t ldwork )
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Univ. of California Berkeley
August 2013
Purpose
=======
ZLARFB applies a complex block reflector H or its transpose H^H to a
COMPLEX_16 m by n matrix C, from the left.
Arguments
=========
SIDE (input) CHARACTER
= 'L': apply H or H^H from the Left
= 'R': apply H or H^H from the Right
TRANS (input) CHARACTER
= 'N': apply H (No transpose)
= 'C': apply H^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) COMPLEX_16 array on the GPU, dimension
(LDV,K) if STOREV = 'C'
(LDV,M) if STOREV = 'R' and SIDE = 'L'
(LDV,N) if STOREV = 'R' and SIDE = 'R'
The matrix V. See further details.
LDV (input) INTEGER
The leading dimension of the array V.
If STOREV = 'C' and SIDE = 'L', LDV >= max(1,M);
if STOREV = 'C' and SIDE = 'R', LDV >= max(1,N);
if STOREV = 'R', LDV >= K.
DT (input) COMPLEX_16 array on the GPU, 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) COMPLEX_16 array on the GPU, dimension (LDC,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.
LDC (input) INTEGER
The leading dimension of the array C. LDA >= max(1,M).
WORK (workspace) COMPLEX_16 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);
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 = 'F' and STOREV = 'C': DIRECT = 'F' and STOREV = 'R':
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 = 'B' and STOREV = 'C': DIRECT = 'B' and STOREV = 'R':
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 )
=================================================================== */
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_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 == 'C' || storev == 'c') && (side == 'L' || side == 'l') && ldv < max(1,m)) ||
((storev == 'C' || storev == 'c') && (side == 'R' || side == 'r') && ldv < max(1,n)) ||
((storev == 'R' || storev == 'r') && ldv < k) ) {
info = -9;
} else if (ldt < k) {
info = -11;
} else if (ldc < max(1,m)) {
info = -13;
} else if ( ((side == 'L' || side == 'l') && ldwork < max(1,n)) ||
((side == 'R' || side == 'r') && ldwork < max(1,m)) ) {
info = -15;
}
if (info != 0) {
magma_xerbla( __func__, -(info) );
return info;
}
/* Function Body */
if (m <= 0 || n <= 0) {
return MAGMA_SUCCESS;
}
// opposite of trans
char transt;
if (trans == 'N' || trans == 'n')
transt = MagmaConjTrans;
else
transt = MagmaNoTrans;
// whether T is upper or lower triangular
char uplo;
if (direct == 'F' || direct == 'f')
uplo = MagmaUpper;
else
uplo = MagmaLower;
// whether V is stored transposed or not
char notransV, transV;
if (storev == 'C' || storev == 'c') {
notransV = MagmaNoTrans;
transV = MagmaConjTrans;
}
else {
notransV = MagmaConjTrans;
transV = MagmaNoTrans;
}
if ( side == 'l' || side == 'L' ) {
// 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_zgemm( MagmaConjTrans, notransV,
n, k, m,
c_one, dC, ldc,
dV, ldv,
c_zero, dwork, ldwork);
// W = W T^H = C^H V T^H
magma_ztrmm( MagmaRight, uplo, transt, 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_zgemm( notransV, MagmaConjTrans,
m, n, k,
c_neg_one, dV, ldv,
dwork, ldwork,
c_one, dC, ldc);
}
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_zgemm( MagmaNoTrans, notransV,
m, k, n,
c_one, dC, ldc,
dV, ldv,
c_zero, dwork, ldwork);
// W = W T = C V T
magma_ztrmm( MagmaRight, uplo, trans, MagmaNonUnit,
m, k,
c_one, dT, ldt,
dwork, ldwork);
// C = C - W V^H = C - C V T V^H = C (I - V T V^H) = C H
magma_zgemm( MagmaNoTrans, transV,
m, n, k,
c_neg_one, dwork, ldwork,
dV, ldv,
c_one, dC, ldc);
}
return MAGMA_SUCCESS;
} /* magma_zlarfb */
extern "C" magma_int_t
magma_zhegvx(magma_int_t itype, char jobz, char range, char uplo, magma_int_t n,
magmaDoubleComplex *a, magma_int_t lda, magmaDoubleComplex *b, magma_int_t ldb,
double vl, double vu, magma_int_t il, magma_int_t iu, double abstol,
magma_int_t *m, double *w, magmaDoubleComplex *z, magma_int_t ldz,
magmaDoubleComplex *work, magma_int_t lwork, double *rwork,
magma_int_t *iwork, magma_int_t *ifail, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
ZHEGVX computes selected eigenvalues, and optionally, 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.
Eigenvalues and eigenvectors can be selected by specifying either a
range of values or a range of indices for the desired eigenvalues.
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.
RANGE (input) CHARACTER*1
= 'A': all eigenvalues will be found.
= 'V': all eigenvalues in the half-open interval (VL,VU]
will be found.
= 'I': the IL-th through IU-th eigenvalues will be found.
UPLO (input) CHARACTER*1
= 'U': Upper 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 Hermitian matrix A. If UPLO = 'U', the
leading N-by-N upper triangular part of A contains the
upper triangular part of the matrix A. If UPLO = 'L',
the leading N-by-N lower triangular part of A contains
the lower triangular part of the matrix A.
On exit, the lower triangle (if UPLO='L') or the upper
triangle (if UPLO='U') of A, including the diagonal, is
destroyed.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
B (input/output) COMPLEX_16 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'.
ABSTOL (input) DOUBLE PRECISION
The absolute error tolerance for the eigenvalues.
An approximate eigenvalue is accepted as converged
when it is determined to lie in an interval [a,b]
of width less than or equal to
ABSTOL + EPS * max( |a|,|b| ) ,
where EPS is the machine precision. If ABSTOL is less than
or equal to zero, then EPS*|T| will be used in its place,
where |T| is the 1-norm of the tridiagonal matrix obtained
by reducing A to tridiagonal form.
Eigenvalues will be computed most accurately when ABSTOL is
set to twice the underflow threshold 2*DLAMCH('S'), not zero.
If this routine returns with INFO>0, indicating that some
eigenvectors did not converge, try setting ABSTOL to
2*DLAMCH('S').
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)
The first M elements contain the selected
eigenvalues in ascending order.
Z (output) COMPLEX_16 array, dimension (LDZ, max(1,M))
If JOBZ = 'N', then Z is not referenced.
If JOBZ = 'V', then if INFO = 0, the first M columns of Z
contain the orthonormal eigenvectors of the matrix A
corresponding to the selected eigenvalues, with the i-th
column of Z holding the eigenvector associated with W(i).
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 an eigenvector fails to converge, then that column of Z
contains the latest approximation to the eigenvector, and the
index of the eigenvector is returned in IFAIL.
Note: the user must ensure that at least max(1,M) columns are
supplied in the array Z; if RANGE = 'V', the exact value of M
is not known in advance and an upper bound must be used.
LDZ (input) INTEGER
The leading dimension of the array Z. LDZ >= 1, and if
JOBZ = 'V', LDZ >= max(1,N).
WORK (workspace/output) COMPLEX_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. LWORK >= max(1,2*N).
For optimal efficiency, LWORK >= (NB+1)*N,
where NB is the blocksize for ZHETRD returned by ILAENV.
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
RWORK (workspace) DOUBLE PRECISION array, dimension (7*N)
IWORK (workspace) INTEGER array, dimension (5*N)
IFAIL (output) INTEGER array, dimension (N)
If JOBZ = 'V', then if INFO = 0, the first M elements of
IFAIL are zero. If INFO > 0, then IFAIL contains the
indices of the eigenvectors that failed to converge.
If JOBZ = 'N', then IFAIL is not referenced.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
> 0: ZPOTRF or ZHEEVX returned an error code:
<= N: if INFO = i, ZHEEVX failed to converge;
i eigenvectors failed to converge. Their indices
are stored in array IFAIL.
> 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
===================================================================== */
char uplo_[2] = {uplo, 0};
char jobz_[2] = {jobz, 0};
char range_[2] = {range, 0};
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex *da;
magmaDoubleComplex *db;
magmaDoubleComplex *dz;
magma_int_t ldda = n;
magma_int_t lddb = n;
magma_int_t lddz = 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_queue_t stream;
magma_queue_create( &stream );
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;
*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 (ldz < 1 || (wantz && ldz < n)) {
*info = -18;
} 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_zhetrd_nb(n);
lwmin = n * (nb + 1);
MAGMA_Z_SET2REAL(work[0],(double)lwmin);
if (lwork < lwmin && ! lquery) {
*info = -20;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info));
return *info;
} else if (lquery) {
return *info;
}
/* Quick return if possible */
if (n == 0) {
return *info;
}
if (MAGMA_SUCCESS != magma_zmalloc( &da, n*ldda ) ||
MAGMA_SUCCESS != magma_zmalloc( &db, n*lddb ) ||
MAGMA_SUCCESS != magma_zmalloc( &dz, n*lddz )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/* Form a Cholesky factorization of B. */
magma_zsetmatrix( n, n, b, ldb, db, lddb );
magma_zsetmatrix_async( n, n,
a, lda,
da, ldda, stream );
magma_zpotrf_gpu(uplo_[0], n, db, lddb, info);
if (*info != 0) {
*info = n + *info;
return *info;
}
magma_queue_sync( stream );
magma_zgetmatrix_async( n, n,
db, lddb,
b, ldb, stream );
/* Transform problem to standard eigenvalue problem and solve. */
magma_zhegst_gpu(itype, uplo, n, da, ldda, db, lddb, info);
magma_zheevx_gpu(jobz, range, uplo, n, da, ldda, vl, vu, il, iu, abstol, m, w, dz, lddz, a, lda, z, ldz, work, lwork, rwork, iwork, ifail, info);
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 = MagmaConjTrans;
} else {
*(unsigned char *)trans = MagmaNoTrans;
}
magma_ztrsm(MagmaLeft, uplo, *trans, MagmaNonUnit, n, *m, c_one, db, lddb, dz, lddz);
}
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_ztrmm(MagmaLeft, uplo, *trans, MagmaNonUnit, n, *m, c_one, db, lddb, dz, lddz);
}
magma_zgetmatrix( n, *m, dz, lddz, z, ldz );
}
magma_queue_sync( stream );
magma_queue_destroy( stream );
magma_free( da );
magma_free( db );
magma_free( dz );
return *info;
} /* zhegvx */
/**
Purpose
-------
ZGETRF_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
---------
@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]
ib INTEGER
The inner-blocking size. IB >= 0.
@param[in,out]
hA 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.
@param[in]
ldha INTEGER
The leading dimension of the array hA. LDHA >= max(1,M).
@param[in,out]
dA 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.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
@param[out]
hL 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.
@param[in]
ldhl INTEGER
The leading dimension of the array hL. LDHL >= max(1,2*IB).
@param[out]
dL 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.
@param[in]
lddl INTEGER
The leading dimension of the array dL. LDDL >= max(1,2*IB).
@param[out]
ipiv INTEGER array, dimension min(M,N), on the cpu.
The pivot indices array.
@param[out]
dWORK DOUBLE COMPLEX array, dimension(LDDWORK, 2*IB), on gpu.
Workspace.
@param[in]
lddwork INTEGER
The leading dimension of the array dWORK. LDDWORK >= max(NB, 1).
@param[out]
info 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.
@ingroup magma_zgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_zgetrf_incpiv_gpu( magma_order_t order, magma_int_t m, magma_int_t n, magma_int_t ib,
magmaDoubleComplex *hA, magma_int_t ldha, magmaDoubleComplex *dA, magma_int_t ldda,
magmaDoubleComplex *hL, magma_int_t ldhl, magmaDoubleComplex *dL, magma_int_t lddl,
magma_int_t *ipiv,
magmaDoubleComplex *dwork, magma_int_t lddwork,
magma_int_t *info)
{
#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 )
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t iinfo;
magma_int_t maxm, mindim;
magma_int_t i, rows, cols, s, ii, sb;
magmaDoubleComplex *dAT;
#ifndef WITHOUTTRTRI
magmaDoubleComplex *dL2 = dL + ib;
magmaDoubleComplex *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_zgetrf(&m, &n, hA, &ldha, ipiv, info);
#ifndef WITHOUTTRTRI
CORE_zlacpy(PlasmaUpperLower, mindim, mindim,
(PLASMA_Complex64_t*)hA, ldha,
(PLASMA_Complex64_t*)hL2, ldhl );
CORE_ztrtri( PlasmaLower, PlasmaUnit, mindim,
(PLASMA_Complex64_t*)hL2, ldhl, info );
if (*info != 0 ) {
fprintf(stderr, "ERROR, trtri returned with info = %d\n", *info);
}
magma_zsetmatrix( mindim, mindim, hL2, ldhl, dL2, lddl );
#endif
if ( order == MagmaRowMajor ) {
magma_zsetmatrix( m, n, hA, ldha, dwork, lddwork );
magmablas_ztranspose( m, n, dwork, lddwork, dA, ldda );
} else {
magma_zsetmatrix( m, n, hA, ldha, dA, ldda );
}
}
else {
/* Use hybrid blocked code. */
maxm = ((m + 31)/32)*32;
if ( order == MagmaColMajor ) {
magmablas_zgetmo_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_ztranspose( sb, m, AT(0,i), ldda, dwork, maxm );
magma_zgetmatrix( m, sb, dwork, maxm, hA(0, i), ldha );
// make sure that gpu queue is empty
//magma_device_sync();
#ifndef WITHOUTTRTRI
magma_ztrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n - (ii+sb), ib,
c_one, dL2(i-1), lddl,
AT(i-1,i+1), ldda );
#else
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n - (ii+sb), ib,
c_one, AT(i-1,i-1), ldda,
AT(i-1,i+1), ldda );
#endif
magma_zgemm( 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_zgetrf( &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_zlaswp( n-ii, AT(0, i), ldda, ii+1, ii+sb, ipiv, 1 );
#ifndef WITHOUTTRTRI
CORE_zlacpy(PlasmaLower, sb, sb,
(PLASMA_Complex64_t*)hA(i, i), ldha,
(PLASMA_Complex64_t*)hL2(i), ldhl );
CORE_ztrtri( PlasmaLower, PlasmaUnit, sb,
(PLASMA_Complex64_t*)hL2(i), ldhl, info );
if (*info != 0 ) {
fprintf(stderr, "ERROR, trtri returned with info = %d\n", *info);
}
magma_zsetmatrix( sb, sb, hL2(i), ldhl, dL2(i), lddl );
#endif
// upload i-th panel
magma_zsetmatrix( rows, sb, hA(i, i), ldha, dwork, cols );
magmablas_ztranspose( rows, sb, dwork, cols, AT(i,i), ldda );
// do the small non-parallel computations
if ( s > (i+1) ) {
#ifndef WITHOUTTRTRI
magma_ztrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
sb, sb,
c_one, dL2(i), lddl,
AT(i, i+1), ldda);
#else
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
sb, sb,
c_one, AT(i, i ), ldda,
AT(i, i+1), ldda);
#endif
magma_zgemm( 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_ztrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n-mindim, sb,
c_one, dL2(i), lddl,
AT(i, i+1), ldda);
#else
magma_ztrsm( 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_zgemm( 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 ( order == MagmaColMajor ) {
magmablas_zgetmo_out( dA, dAT, ldda, m, n );
}
}
return *info;
}
extern "C" magma_int_t
magma_zlauum(char uplo, magma_int_t n,
cuDoubleComplex *a, magma_int_t lda, magma_int_t *info)
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
Purpose
=======
ZLAUUM 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 = 'U' or 'u' then the upper triangle of the result is stored,
overwriting the factor U in A.
If UPLO = 'L' or 'l' 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
=========
UPLO (input) CHARACTER*1
Specifies whether the triangular factor stored in the array A
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.
A (input/output) COPLEX_16 array, dimension (LDA,N)
On entry, the triangular factor U or L.
On exit, if UPLO = 'U', the upper triangle of A is
overwritten with the upper triangle of the product U * U';
if UPLO = 'L', the lower triangle of A is overwritten with
the lower triangle of the product L' * L.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= 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 ldda, nb;
magma_int_t i, ib;
cuDoubleComplex c_one = MAGMA_Z_ONE;
double d_one = MAGMA_D_ONE;
cuDoubleComplex *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 (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_zmalloc( &work, (n)*ldda )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
cudaStream_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
nb = magma_get_zpotrf_nb(n);
if (nb <= 1 || nb >= n)
lapackf77_zlauum(uplo_, &n, a, &lda, info);
else
{
if (upper)
{
/* Compute the product U * U'. */
for (i=0; i<n; i=i+nb)
{
ib=min(nb,n-i);
//cublasSetMatrix(ib, (n-i), sizeof(cuDoubleComplex), A(i, i), lda, dA(i, i), ldda);
magma_zsetmatrix_async( ib, ib,
A(i,i), lda,
dA(i, i), ldda, stream[1] );
magma_zsetmatrix_async( ib, (n-i-ib),
A(i,i+ib), lda,
dA(i,i+ib), ldda, stream[0] );
magma_queue_sync( stream[1] );
magma_ztrmm( MagmaRight, MagmaUpper,
MagmaConjTrans, MagmaNonUnit, i, ib,
c_one, dA(i,i), ldda, dA(0, i),ldda);
lapackf77_zlauum(MagmaUpperStr, &ib, A(i,i), &lda, info);
magma_zsetmatrix_async( ib, ib,
A(i, i), lda,
dA(i, i), ldda, stream[0] );
if (i+ib < n)
{
magma_zgemm( 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_zherk( MagmaUpper, MagmaNoTrans, ib,(n-i-ib),
d_one, dA(i, i+ib), ldda,
d_one, dA(i, i), ldda);
}
magma_zgetmatrix( i+ib, ib,
dA(0, i), ldda,
A(0, i), lda );
}
}
else
{
/* Compute the product L' * L. */
for(i=0; i<n; i=i+nb)
{
ib=min(nb,n-i);
//cublasSetMatrix((n-i), ib, sizeof(cuDoubleComplex),
// A(i, i), lda, dA(i, i), ldda);
magma_zsetmatrix_async( ib, ib,
A(i,i), lda,
dA(i, i), ldda, stream[1] );
magma_zsetmatrix_async( (n-i-ib), ib,
A(i+ib, i), lda,
dA(i+ib, i), ldda, stream[0] );
magma_queue_sync( stream[1] );
magma_ztrmm( MagmaLeft, MagmaLower,
MagmaConjTrans, MagmaNonUnit, ib,
i, c_one, dA(i,i), ldda,
dA(i, 0),ldda);
lapackf77_zlauum(MagmaLowerStr, &ib, A(i,i), &lda, info);
//cublasSetMatrix(ib, ib, sizeof(cuDoubleComplex),
// A(i, i), lda, dA(i, i), ldda);
magma_zsetmatrix_async( ib, ib,
A(i, i), lda,
dA(i, i), ldda, stream[0] );
if (i+ib < n)
{
magma_zgemm(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_zherk(MagmaLower, MagmaConjTrans, ib, (n-i-ib),
d_one, dA(i+ib, i), ldda,
d_one, dA(i, i), ldda);
}
magma_zgetmatrix( ib, i+ib,
dA(i, 0), ldda,
A(i, 0), lda );
}
}
}
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free( work );
return *info;
}
/**
Purpose
-------
ZHEGST_GPU reduces a complex Hermitian-definite generalized
eigenproblem to standard form.
If ITYPE = 1, the problem is A*x = lambda*B*x,
and A is overwritten by inv(U^H)*A*inv(U) or inv(L)*A*inv(L^H)
If ITYPE = 2 or 3, the problem is A*B*x = lambda*x or
B*A*x = lambda*x, and A is overwritten by U*A*U^H or L^H*A*L.
B must have been previously factorized as U^H*U or L*L^H by ZPOTRF.
Arguments
---------
@param[in]
itype INTEGER
= 1: compute inv(U^H)*A*inv(U) or inv(L)*A*inv(L^H);
= 2 or 3: compute U*A*U^H or L^H*A*L.
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored and B is factored as U^H*U;
- = MagmaLower: Lower triangle of A is stored and B is factored as L*L^H.
@param[in]
n INTEGER
The order of the matrices A and B. N >= 0.
@param[in,out]
dA COMPLEX_16 array, on the GPU device, dimension (LDDA,N)
On entry, the Hermitian matrix A. If UPLO = MagmaUpper, the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
\n
On exit, if INFO = 0, the transformed matrix, stored in the
same format as A.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,N).
@param[in]
dB COMPLEX_16 array, on the GPU device, dimension (LDDB,N)
The triangular factor from the Cholesky factorization of B,
as returned by ZPOTRF.
@param[in]
lddb INTEGER
The leading dimension of the array B. LDDB >= max(1,N).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_zheev_comp
********************************************************************/
extern "C" magma_int_t
magma_zhegst_gpu(
magma_int_t itype, magma_uplo_t uplo, magma_int_t n,
magmaDoubleComplex_ptr dA, magma_int_t ldda,
magmaDoubleComplex_const_ptr dB, magma_int_t lddb,
magma_int_t *info)
{
#define A(i_, j_) (work + (i_) + (j_)*lda )
#define B(i_, j_) (work + (i_) + (j_)*ldb + nb*ldb)
#define dA(i_, j_) (dA + (i_) + (j_)*ldda)
#define dB(i_, j_) (dB + (i_) + (j_)*lddb)
/* Constants */
const magmaDoubleComplex c_one = MAGMA_Z_ONE;
const magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
const magmaDoubleComplex c_half = MAGMA_Z_HALF;
const magmaDoubleComplex c_neg_half = MAGMA_Z_NEG_HALF;
const double d_one = 1.0;
/* Local variables */
const char* uplo_ = lapack_uplo_const( uplo );
magma_int_t k, kb, kb2, nb;
magma_int_t lda;
magma_int_t ldb;
magmaDoubleComplex *work;
bool upper = (uplo == MagmaUpper);
/* Test the input parameters. */
*info = 0;
if (itype < 1 || itype > 3) {
*info = -1;
} else if (! upper && uplo != MagmaLower) {
*info = -2;
} else if (n < 0) {
*info = -3;
} else if (ldda < max(1,n)) {
*info = -5;
} else if (lddb < max(1,n)) {
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return */
if ( n == 0 )
return *info;
nb = magma_get_zhegst_nb( n );
lda = nb;
ldb = nb;
if (MAGMA_SUCCESS != magma_zmalloc_pinned( &work, 2*nb*nb )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_queue_t queues[2];
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[0] );
magma_queue_create( cdev, &queues[1] );
/* Use hybrid blocked code */
if (itype == 1) {
if (upper) {
kb = min( n, nb );
/* Compute inv(U^H)*A*inv(U) */
magma_zgetmatrix_async( kb, kb,
dA(0, 0), ldda,
A(0, 0), lda, queues[0] );
magma_zgetmatrix_async( kb, kb,
dB(0, 0), lddb,
B(0, 0), ldb, queues[0] );
for (k = 0; k < n; k += nb) {
kb = min( n-k, nb );
kb2 = min( n-k-nb, nb );
magma_queue_sync( queues[0] ); // finish get dA(k,k) -> A(0,0) and dB(k,k) -> B(0,0)
/* Update the upper triangle of A(k:n,k:n) */
lapackf77_zhegst( &itype, uplo_, &kb, A(0,0), &lda, B(0,0), &ldb, info );
magma_zsetmatrix_async( kb, kb,
A(0, 0), lda,
dA(k, k), ldda, queues[0] );
if (k+kb < n) {
magma_ztrsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
kb, n-k-kb,
c_one, dB(k,k), lddb,
dA(k,k+kb), ldda, queues[1] );
magma_queue_sync( queues[0] ); // finish set dA(k,k)
// Start copying next B block
magma_zgetmatrix_async( kb2, kb2,
dB(k+kb, k+kb), lddb,
B(0, 0), ldb, queues[0] );
magma_zhemm( MagmaLeft, MagmaUpper,
kb, n-k-kb,
c_neg_half, dA(k,k), ldda,
dB(k,k+kb), lddb,
c_one, dA(k,k+kb), ldda, queues[1] );
magma_zher2k( MagmaUpper, MagmaConjTrans,
n-k-kb, kb,
c_neg_one, dA(k,k+kb), ldda,
dB(k,k+kb), lddb,
d_one, dA(k+kb,k+kb), ldda, queues[1] );
// Start copying next A block
magma_queue_sync( queues[1] );
magma_zgetmatrix_async( kb2, kb2,
dA(k+kb, k+kb), ldda,
A(0, 0), lda, queues[0] );
magma_zhemm( MagmaLeft, MagmaUpper,
kb, n-k-kb,
c_neg_half, dA(k,k), ldda,
dB(k,k+kb), lddb,
c_one, dA(k,k+kb), ldda, queues[1] );
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
kb, n-k-kb,
c_one, dB(k+kb,k+kb), lddb,
dA(k,k+kb), ldda, queues[1] );
}
}
}
else {
kb = min( n, nb );
/* Compute inv(L)*A*inv(L^H) */
magma_zgetmatrix_async( kb, kb,
dA(0, 0), ldda,
A(0, 0), lda, queues[0] );
magma_zgetmatrix_async( kb, kb,
dB(0, 0), lddb,
B(0, 0), ldb, queues[0] );
for (k = 0; k < n; k += nb) {
kb = min( n-k, nb );
kb2 = min( n-k-nb, nb );
magma_queue_sync( queues[0] ); // finish get dA(k,k) -> A(0,0) and dB(k,k) -> B(0,0)
/* Update the lower triangle of A(k:n,k:n) */
lapackf77_zhegst( &itype, uplo_, &kb, A(0, 0), &lda, B(0, 0), &ldb, info );
magma_zsetmatrix_async( kb, kb,
A(0, 0), lda,
dA(k, k), ldda, queues[0] );
if (k+kb < n) {
magma_ztrsm( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
n-k-kb, kb,
c_one, dB(k,k), lddb,
dA(k+kb,k), ldda, queues[1] );
magma_queue_sync( queues[0] ); // finish set dA(k,k)
// Start copying next B block
magma_zgetmatrix_async( kb2, kb2,
dB(k+kb, k+kb), lddb,
B(0, 0), ldb, queues[0] );
magma_zhemm( MagmaRight, MagmaLower,
n-k-kb, kb,
c_neg_half, dA(k,k), ldda,
dB(k+kb,k), lddb,
c_one, dA(k+kb, k), ldda, queues[1] );
magma_zher2k( MagmaLower, MagmaNoTrans,
n-k-kb, kb,
c_neg_one, dA(k+kb,k), ldda,
dB(k+kb,k), lddb,
d_one, dA(k+kb,k+kb), ldda, queues[1] );
// Start copying next A block
magma_queue_sync( queues[1] );
magma_zgetmatrix_async( kb2, kb2,
dA(k+kb, k+kb), ldda,
A(0, 0), lda, queues[0] );
magma_zhemm( MagmaRight, MagmaLower,
n-k-kb, kb,
c_neg_half, dA(k,k), ldda,
dB(k+kb,k), lddb,
c_one, dA(k+kb,k), ldda, queues[1] );
magma_ztrsm( MagmaLeft, MagmaLower, MagmaNoTrans, MagmaNonUnit,
n-k-kb, kb,
c_one, dB(k+kb,k+kb), lddb,
dA(k+kb,k), ldda, queues[1] );
}
}
}
}
else { // itype == 2 or 3
if (upper) {
/* Compute U*A*U^H */
for (k = 0; k < n; k += nb) {
kb = min( n-k, nb );
magma_zgetmatrix_async( kb, kb,
dA(k, k), ldda,
A(0, 0), lda, queues[0] );
magma_zgetmatrix_async( kb, kb,
dB(k, k), lddb,
B(0, 0), ldb, queues[0] );
/* Update the upper triangle of A(1:k+kb-1,1:k+kb-1) */
if (k > 0) {
magma_ztrmm( MagmaLeft, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
k, kb,
c_one, dB(0,0), lddb,
dA(0,k), ldda, queues[1] );
magma_zhemm( MagmaRight, MagmaUpper,
k, kb,
c_half, dA(k,k), ldda,
dB(0,k), lddb,
c_one, dA(0,k), ldda, queues[1] );
magma_zher2k( MagmaUpper, MagmaNoTrans,
k, kb,
c_one, dA(0,k), ldda,
dB(0,k), lddb,
d_one, dA(0,0), ldda, queues[1] );
magma_zhemm( MagmaRight, MagmaUpper,
k, kb,
c_half, dA(k,k), ldda,
dB(0,k), lddb,
c_one, dA(0,k), ldda, queues[1] );
magma_ztrmm( MagmaRight, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
k, kb,
c_one, dB(k,k), lddb,
dA(0,k), ldda, queues[1] );
}
magma_queue_sync( queues[0] ); // finish get dA(k,k) -> A(0,0) and dB(k,k) -> B(0,0)
lapackf77_zhegst( &itype, uplo_, &kb, A(0, 0), &lda, B(0, 0), &ldb, info );
magma_zsetmatrix_async( kb, kb,
A(0, 0), lda,
dA(k, k), ldda, queues[1] );
magma_queue_sync( queues[1] ); // wait for A(0,0) before getting next panel
}
}
else {
/* Compute L^H*A*L */
for (k = 0; k < n; k += nb) {
kb = min( n-k, nb );
magma_zgetmatrix_async( kb, kb,
dA(k, k), ldda,
A(0, 0), lda, queues[0] );
magma_zgetmatrix_async( kb, kb,
dB(k, k), lddb,
B(0, 0), ldb, queues[0] );
/* Update the lower triangle of A(1:k+kb-1,1:k+kb-1) */
if (k > 0) {
magma_ztrmm( MagmaRight, MagmaLower, MagmaNoTrans, MagmaNonUnit,
kb, k,
c_one, dB(0,0), lddb,
dA(k,0), ldda, queues[1] );
magma_zhemm( MagmaLeft, MagmaLower,
kb, k,
c_half, dA(k,k), ldda,
dB(k,0), lddb,
c_one, dA(k,0), ldda, queues[1] );
magma_queue_sync( queues[1] );
magma_zher2k( MagmaLower, MagmaConjTrans,
k, kb,
c_one, dA(k,0), ldda,
dB(k,0), lddb,
d_one, dA(0,0), ldda, queues[1] );
magma_zhemm( MagmaLeft, MagmaLower,
kb, k,
c_half, dA(k,k), ldda,
dB(k,0), lddb,
c_one, dA(k,0), ldda, queues[1] );
magma_ztrmm( MagmaLeft, MagmaLower, MagmaConjTrans, MagmaNonUnit,
kb, k,
c_one, dB(k,k), lddb,
dA(k,0), ldda, queues[1] );
}
magma_queue_sync( queues[0] ); // finish get dA(k,k) -> A(0,0) and dB(k,k) -> B(0,0)
lapackf77_zhegst( &itype, uplo_, &kb, A(0, 0), &lda, B(0, 0), &ldb, info );
magma_zsetmatrix_async( kb, kb,
A(0, 0), lda,
dA(k, k), ldda, queues[1] );
magma_queue_sync( queues[1] ); // wait for A(0,0) before getting next panel
}
}
}
magma_queue_sync( queues[0] );
magma_queue_sync( queues[1] );
magma_queue_destroy( queues[0] );
magma_queue_destroy( queues[1] );
magma_free_pinned( work );
return *info;
} /* magma_zhegst_gpu */
/**
Purpose
-------
ZGESSM applies the factors L computed by ZGETRF_INCPIV to
a complex 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
ZGETRF_INCPIV.
@param[in]
dL1 DOUBLE COMPLEX 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 DOUBLE COMPLEX 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 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.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,M).
@ingroup magma_zgesv_tile
********************************************************************/
extern "C" magma_int_t
magma_zgessm_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,
magmaDoubleComplex *dL1, magma_int_t lddl1,
magmaDoubleComplex *dL, magma_int_t lddl,
magmaDoubleComplex *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)
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
int i, s, sb;
magmaDoubleComplex *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_zgetmo_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_zlaswp( n, dAT, ldda, i+1, i+sb, ipiv, 1 );
#ifndef WITHOUTTRTRI
magma_ztrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n, sb,
c_one, dL1(i), lddl1,
AT(i, 0), ldda);
#else
magma_ztrsm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n, sb,
c_one, L( i, i), lddl,
AT(i, 0), ldda);
#endif
if ( (i+sb) < m) {
magma_zgemm( 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_zgetmo_in( dA, dAT, ldda, m, n );
}
return *info;
} /* magma_zgessm_gpu */
/**
Purpose
-------
ZLAUUM 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_zposv_aux
***************************************************************************/
extern "C" magma_int_t
magma_zlauum(
magma_uplo_t uplo, magma_int_t n,
magmaDoubleComplex *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;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
double d_one = MAGMA_D_ONE;
magmaDoubleComplex *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_zmalloc( &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_zpotrf_nb(n);
if (nb <= 1 || nb >= n)
lapackf77_zlauum(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_zsetmatrix_async( ib, ib,
A(i,i), lda,
dA(i, i), ldda, stream[1] );
magma_zsetmatrix_async( ib, (n-i-ib),
A(i,i+ib), lda,
dA(i,i+ib), ldda, stream[0] );
magma_queue_sync( stream[1] );
magma_ztrmm( MagmaRight, MagmaUpper,
MagmaConjTrans, MagmaNonUnit, i, ib,
c_one, dA(i,i), ldda, dA(0, i),ldda);
lapackf77_zlauum(MagmaUpperStr, &ib, A(i,i), &lda, info);
magma_zsetmatrix_async( ib, ib,
A(i, i), lda,
dA(i, i), ldda, stream[0] );
if (i+ib < n) {
magma_zgemm( 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_zherk( MagmaUpper, MagmaNoTrans, ib,(n-i-ib),
d_one, dA(i, i+ib), ldda,
d_one, dA(i, i), ldda);
}
magma_zgetmatrix( 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_zsetmatrix_async( ib, ib,
A(i,i), lda,
dA(i, i), ldda, stream[1] );
magma_zsetmatrix_async( (n-i-ib), ib,
A(i+ib, i), lda,
dA(i+ib, i), ldda, stream[0] );
magma_queue_sync( stream[1] );
magma_ztrmm( MagmaLeft, MagmaLower,
MagmaConjTrans, MagmaNonUnit, ib,
i, c_one, dA(i,i), ldda,
dA(i, 0),ldda);
lapackf77_zlauum(MagmaLowerStr, &ib, A(i,i), &lda, info);
magma_zsetmatrix_async( ib, ib,
A(i, i), lda,
dA(i, i), ldda, stream[0] );
if (i+ib < n) {
magma_zgemm(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_zherk(MagmaLower, MagmaConjTrans, ib, (n-i-ib),
d_one, dA(i+ib, i), ldda,
d_one, dA(i, i), ldda);
}
magma_zgetmatrix( 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;
}
extern "C" magma_int_t
magma_ztrtri(char uplo, char diag, magma_int_t n,
magmaDoubleComplex *A, magma_int_t lda, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
ZTRTRI computes the inverse of a real upper or lower triangular
matrix A.
This is the Level 3 BLAS version of the algorithm.
Arguments
=========
UPLO (input) CHARACTER*1
= 'U': A is upper triangular;
= 'L': A is lower triangular.
DIAG (input) CHARACTER*1
= 'N': A is non-unit triangular;
= 'U': A is unit triangular.
N (input) INTEGER
The order of the matrix A. N >= 0.
A (input/output) COMPLEX_16 array, dimension (LDA,N)
On entry, the triangular matrix A. If UPLO = 'U', 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 = 'L', 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 = 'U', 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.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
INFO (output) 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.
===================================================================== */
#define A(i, j) ( A + (i) + (j)*lda )
#define dA(i, j) (dA + (i) + (j)*ldda)
/* Local variables */
char uplo_[2] = {uplo, 0};
char diag_[2] = {diag, 0};
magma_int_t ldda, nb, nn, j, jb;
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex *dA;
int upper = lapackf77_lsame(uplo_, "U");
int nounit = lapackf77_lsame(diag_, "N");
*info = 0;
if ((! upper) && (! lapackf77_lsame(uplo_, "L")))
*info = -1;
else if ((! nounit) && (! lapackf77_lsame(diag_, "U")))
*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_Z_EQUAL( *A(j,j), c_zero )) {
*info = j+1; // Fortran index
return *info;
}
}
}
/* Determine the block size for this environment */
nb = magma_get_zpotrf_nb(n);
ldda = ((n+31)/32)*32;
if (MAGMA_SUCCESS != magma_zmalloc( &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_ztrtri(uplo_, diag_, &n, A, &lda, info);
else {
if (upper) {
/* Compute inverse of upper triangular matrix */
for (j=0; j<n; j=j+nb) {
jb = min(nb, (n-j));
magma_zsetmatrix( jb, (n-j),
A(j, j), lda,
dA(j, j), ldda );
/* Compute rows 1:j-1 of current block column */
magma_ztrmm( MagmaLeft, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_one, dA(0,0), ldda, dA(0, j),ldda);
magma_ztrsm( MagmaRight, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_neg_one, dA(j,j), ldda, dA(0, j),ldda);
magma_zgetmatrix_async( jb, jb,
dA(j, j), ldda,
A(j, j), lda, stream[1] );
magma_zgetmatrix_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_ztrtri(MagmaUpperStr, diag_, &jb, A(j,j), &lda, info);
magma_zsetmatrix( 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=j-nb) {
jb=min(nb,(n-j));
if((j+jb) < n) {
magma_zsetmatrix( (n-j), jb,
A(j, j), lda,
dA(j, j), ldda );
/* Compute rows j+jb:n of current block column */
magma_ztrmm( MagmaLeft, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_one, dA(j+jb,j+jb), ldda, dA(j+jb, j), ldda );
magma_ztrsm( MagmaRight, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_neg_one, dA(j,j), ldda, dA(j+jb, j), ldda );
magma_zgetmatrix_async( n-j-jb, jb,
dA(j+jb, j), ldda,
A(j+jb, j), lda, stream[1] );
magma_zgetmatrix_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_ztrtri(MagmaLowerStr, diag_, &jb, A(j,j), &lda, info);
magma_zsetmatrix( 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;
}
extern "C" magma_int_t
magma_ztrtri_gpu(magma_uplo_t uplo, magma_diag_t diag, magma_int_t n,
magmaDoubleComplex_ptr dA, size_t dA_offset, magma_int_t ldda, magma_int_t *info)
{
/* -- clMAGMA (version 1.0.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2012
Purpose
=======
ZTRTRI computes the inverse of a real upper or lower triangular
matrix dA.
This is the Level 3 BLAS version of the algorithm.
Arguments
=========
UPLO (input) CHARACTER*1
= 'U': A is upper triangular;
= 'L': A is lower triangular.
DIAG (input) CHARACTER*1
= 'N': A is non-unit triangular;
= 'U': A is unit triangular.
N (input) INTEGER
The order of the matrix A. N >= 0.
dA (input/output) DOUBLE PRECISION array ON THE GPU, dimension (LDDA,N)
On entry, the triangular matrix A. If UPLO = 'U', the
leading N-by-N upper triangular part of the array dA contains
the upper triangular matrix, and the strictly lower
triangular part of A is not referenced. If UPLO = 'L', the
leading N-by-N lower triangular part of the array dA contains
the lower triangular matrix, and the strictly upper
triangular part of A is not referenced. If DIAG = 'U', 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.
LDDA (input) INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
> 0: if INFO = i, dA(i,i) is exactly zero. The triangular
matrix is singular and its inverse can not be computed.
===================================================================== */
/* Local variables */
magma_uplo_t uplo_ = uplo;
magma_diag_t diag_ = diag;
magma_int_t nb, nn, j, jb;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex *work;
int upper = lapackf77_lsame(lapack_const(uplo_), lapack_const(MagmaUpper));
int nounit = lapackf77_lsame(lapack_const(diag_), lapack_const(MagmaNonUnit));
*info = 0;
if ((! upper) && (! lapackf77_lsame(lapack_const(uplo_), lapack_const(MagmaLower))))
*info = -1;
else if ((! nounit) && (! lapackf77_lsame(lapack_const(diag_), lapack_const(MagmaUnit))))
*info = -2;
else if (n < 0)
*info = -3;
else if (ldda < max(1,n))
*info = -5;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
nb = magma_get_zpotrf_nb(n);
/* Create Queues */
magma_queue_t queues[2];
magma_device_t device;
int num = 0;
magma_err_t err;
err = magma_get_devices( &device, 1, &num );
if ( err != 0 || num < 1 ) {
fprintf( stderr, "magma_get_devices failed: %d\n", err );
exit(-1);
}
err = magma_queue_create( device, &queues[0] );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create 0 failed: %d\n", err );
exit(-1);
}
err = magma_queue_create( device, &queues[1] );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create 1 failed: %d\n", err );
exit(-1);
}
if (MAGMA_SUCCESS != magma_malloc_host( (void**)&work, nb*nb*sizeof(magmaDoubleComplex) )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
if (nb <= 1 || nb >= n)
{
magma_zgetmatrix( n, n, dA, dA_offset, ldda, work, 0, n, queues[0] );
lapackf77_ztrtri(lapack_const(uplo_), lapack_const(diag_), &n, work, &n, info);
magma_zsetmatrix( n, n, work, 0, n, dA, dA_offset, ldda, queues[0] );
}
else
{
if (upper){
/* Compute inverse of upper triangular matrix */
for (j=0; j<n; j =j+ nb){
jb = min(nb, (n-j));
/* Compute rows 1:j-1 of current block column */
magma_ztrmm(MagmaLeft, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_one, dA(0,0), ldda, dA(0, j), ldda,
queues[0]);
magma_ztrsm(MagmaRight, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_neg_one, dA(j,j), ldda, dA(0, j), ldda,
queues[0]);
magma_zgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, 0, jb, queues[1], NULL );
magma_queue_sync( queues[1] );
/* Compute inverse of current diagonal block */
lapackf77_ztrtri(MagmaUpperStr, lapack_const(diag_), &jb, work, &jb, info);
/*
magma_zsetmatrix_async( jb, jb,
work, 0, jb,
dA(j, j), ldda, queues[0], NULL );
*/
magma_zsetmatrix( jb, jb,
work, 0, jb,
dA(j, j), ldda, queues[0]);
}
}
else{
/* Compute inverse of lower triangular matrix */
nn=((n-1)/nb)*nb+1;
for(j=nn-1; j>=0; j=j-nb){
jb=min(nb,(n-j));
if((j+jb) < n){
/* Compute rows j+jb:n of current block column */
magma_ztrmm(MagmaLeft, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_one, dA(j+jb,j+jb), ldda, dA(j+jb, j), ldda,
queues[0]);
magma_ztrsm(MagmaRight, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_neg_one, dA(j,j), ldda, dA(j+jb, j), ldda,
queues[0]);
}
magma_zgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, 0, jb, queues[1], NULL );
magma_queue_sync( queues[1] );
/* Compute inverse of current diagonal block */
lapackf77_ztrtri(MagmaLowerStr, lapack_const(diag_), &jb, work, &jb, info);
/*
magma_zsetmatrix_async( jb, jb,
work, 0, jb,
dA(j, j), ldda, queues[0], NULL );
*/
magma_zsetmatrix( jb, jb,
work, 0, jb,
dA(j, j), ldda, queues[0] );
}
}
}
magma_free_host( work );
magma_queue_destroy(queues[0]);
magma_queue_destroy(queues[1]);
return *info;
}
magma_err_t
magma_zlarfb_gpu( int side, int trans, int direct, int storev,
magma_int_t m, magma_int_t n, magma_int_t k,
magmaDoubleComplex_ptr dV, size_t dV_offset, magma_int_t ldv,
magmaDoubleComplex_ptr dT, size_t dT_offset, magma_int_t ldt,
magmaDoubleComplex_ptr dC, size_t dC_offset, magma_int_t ldc,
magmaDoubleComplex_ptr dwork, size_t dwork_offset, magma_int_t ldwork,
magma_queue_t queue)
{
/* -- clMAGMA (version 1.1.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date January 2014
Purpose
=======
ZLARFB applies a complex block reflector H or its transpose H' to a
COMPLEX_16 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) COMPLEX_16 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) COMPLEX_16 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) COMPLEX_16 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) COMPLEX_16 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);
=================================================================== */
/* TODO: replace with updated larfb_gpu from CUDA MAGMA */
#define dV(i) dV, (i)
#define dT(i) dT, (i)
#define dC(i) dC, (i)
#define dwork(i) dwork, (i)
magmaDoubleComplex c_zero = MAGMA_Z_MAKE( 0.0, 0.0 );
magmaDoubleComplex c_one = MAGMA_Z_MAKE( 1.0, 0.0 );
magmaDoubleComplex c_neg_one = MAGMA_Z_MAKE( -1.0, 0.0 );
if (m <= 0 || n <= 0) {
return MAGMA_SUCCESS;
}
magma_int_t transt;
if (trans == MagmaNoTrans)
transt = MagmaConjTrans;
else
transt = MagmaNoTrans;
if ( side == MagmaLeft ) {
if ( storev == MagmaColumnwise )
{
magma_zgemm( MagmaConjTrans, 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_ztrmm( MagmaRight, MagmaUpper, transt, MagmaNonUnit,
n, k,
c_one, dT(dT_offset), ldt,
dwork(dwork_offset), ldwork, queue);
else
magma_ztrmm( MagmaRight, MagmaLower, transt, MagmaNonUnit,
n, k,
c_one, dT(dT_offset), ldt,
dwork(dwork_offset), ldwork, queue);
magma_zgemm( MagmaNoTrans, MagmaConjTrans,
m, n, k,
c_neg_one, dV(dV_offset), ldv,
dwork(dwork_offset), ldwork,
c_one, dC(dC_offset), ldc, queue);
}
else {
magma_zgemm( MagmaNoTrans, MagmaConjTrans,
m, k, n,
c_one, dC(dC_offset), ldc,
dV(dV_offset), ldv,
c_zero, dwork(dwork_offset), ldwork, queue);
magma_ztrmm( MagmaRight, MagmaUpper, transt, MagmaNonUnit,
m, k,
c_one, dT(dT_offset), ldt,
dwork(dwork_offset), ldwork, queue);
magma_zgemm( 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_zgemm( 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_ztrmm( MagmaRight, MagmaUpper, transt, MagmaNonUnit,
m, k,
c_one, dT(dT_offset), ldt,
dwork(dwork_offset), ldwork, queue);
else
magma_ztrmm( MagmaRight, MagmaLower, transt, MagmaNonUnit,
m, k,
c_one, dT(dT_offset), ldt,
dwork(dwork_offset), ldwork, queue);
magma_zgemm( MagmaNoTrans, MagmaConjTrans,
m, n, k,
c_neg_one, dwork(dwork_offset), ldwork,
dV(dV_offset), ldv,
c_one, dC(dC_offset), ldc, queue);
}
else {
magma_zgemm( MagmaNoTrans, MagmaConjTrans,
m, k, n,
c_one, dC(dC_offset), ldc,
dV(dV_offset), ldv,
c_zero, dwork(dwork_offset), ldwork, queue);
magma_ztrmm( MagmaRight, MagmaUpper, transt, MagmaNonUnit,
m, k,
c_one, dT(dT_offset), ldt,
dwork(dwork_offset), ldwork, queue);
magma_zgemm( 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_zlarfb */
extern "C" magma_int_t
magma_zgessm_gpu( char storev, magma_int_t m, magma_int_t n, magma_int_t k, magma_int_t ib,
magma_int_t *ipiv,
magmaDoubleComplex *dL1, magma_int_t lddl1,
magmaDoubleComplex *dL, magma_int_t lddl,
magmaDoubleComplex *dA, magma_int_t ldda,
magma_int_t *info)
{
/* -- MAGMA (version 1.4.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
December 2013
Purpose
=======
ZGESSM applies the factors L computed by ZGETRF_INCPIV to
a complex 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
ZGETRF_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)
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
int i, s, sb;
magmaDoubleComplex *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_zgetmo_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_zlaswp( n, dAT, ldda, i+1, i+sb, ipiv, 1 );
#ifndef WITHOUTTRTRI
magma_ztrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n, sb,
c_one, dL1(i), lddl1,
AT(i, 0), ldda);
#else
magma_ztrsm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n, sb,
c_one, L( i, i), lddl,
AT(i, 0), ldda);
#endif
if ( (i+sb) < m) {
magma_zgemm( 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_zgetmo_in( dA, dAT, ldda, m, n );
}
return *info;
/* End of MAGMA_ZGETRF_GPU */
}
/**
Purpose
-------
ZHEGVD 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.
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 COMPLEX_16 array, dimension (LDA, N)
On entry, the Hermitian matrix A. If UPLO = MagmaUpper, the
leading N-by-N upper triangular part of A contains the
upper triangular part of the matrix A. If UPLO = MagmaLower,
the leading N-by-N lower triangular part of A contains
the lower triangular part of the matrix A.
\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**H*B*Z = I;
if ITYPE = 3, Z**H*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 COMPLEX_16 array, dimension (LDB, N)
On entry, the Hermitian 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**H*U or B = L*L**H.
@param[in]
ldb INTEGER
The leading dimension of the array B. LDB >= max(1,N).
@param[out]
w DOUBLE PRECISION array, dimension (N)
If INFO = 0, the eigenvalues in ascending order.
@param[out]
work (workspace) COMPLEX_16 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 >= N + N*NB.
If JOBZ = MagmaVec and N > 1, LWORK >= max( N + N*NB, 2*N + N**2 ).
NB can be obtained through magma_get_zhetrd_nb(N).
\n
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.
@param[out]
rwork (workspace) DOUBLE PRECISION array, dimension (LRWORK)
On exit, if INFO = 0, RWORK[0] returns the optimal LRWORK.
@param[in]
lrwork INTEGER
The dimension of the array RWORK.
If N <= 1, LRWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LRWORK >= N.
If JOBZ = MagmaVec and N > 1, LRWORK >= 1 + 5*N + 2*N**2.
\n
If LRWORK = -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.
@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, 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.
@param[out]
info 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 = 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 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.
@ingroup magma_zhegv_driver
********************************************************************/
extern "C" magma_int_t
magma_zhegvd(magma_int_t itype, magma_vec_t jobz, magma_uplo_t uplo, magma_int_t n,
magmaDoubleComplex *A, magma_int_t lda, magmaDoubleComplex *B, magma_int_t ldb,
double *w, magmaDoubleComplex *work, magma_int_t lwork,
double *rwork, magma_int_t lrwork,
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 );
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex *da;
magmaDoubleComplex *db;
magma_int_t ldda = n;
magma_int_t lddb = n;
magma_int_t lower;
magma_trans_t trans;
magma_int_t wantz;
magma_int_t lquery;
magma_int_t lwmin;
magma_int_t liwmin;
magma_int_t lrwmin;
magma_queue_t stream;
magma_queue_create( &stream );
wantz = (jobz == MagmaVec);
lower = (uplo == MagmaLower);
lquery = (lwork == -1 || lrwork == -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_zhetrd_nb( n );
if ( n <= 1 ) {
lwmin = 1;
lrwmin = 1;
liwmin = 1;
}
else if ( wantz ) {
lwmin = max( n + n*nb, 2*n + n*n );
lrwmin = 1 + 5*n + 2*n*n;
liwmin = 3 + 5*n;
}
else {
lwmin = n + n*nb;
lrwmin = n;
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_dlamch("Epsilon");
work[0] = MAGMA_Z_MAKE( lwmin * one_eps, 0.); // round up
rwork[0] = lrwmin * one_eps;
iwork[0] = liwmin;
if (lwork < lwmin && ! lquery) {
*info = -11;
} else if (lrwork < lrwmin && ! lquery) {
*info = -13;
} else if (liwork < liwmin && ! lquery) {
*info = -15;
}
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_zhegvd(&itype, jobz_, uplo_,
&n, A, &lda, B, &ldb,
w, work, &lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, &lrwork,
#endif
iwork, &liwork, info);
return *info;
}
// TODO fix memory leak
if (MAGMA_SUCCESS != magma_zmalloc( &da, n*ldda ) ||
MAGMA_SUCCESS != magma_zmalloc( &db, n*lddb )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/* Form a Cholesky factorization of B. */
magma_zsetmatrix( n, n, B, ldb, db, lddb );
magma_zsetmatrix_async( n, n,
A, lda,
da, ldda, stream );
magma_timer_t time=0;
timer_start( time );
magma_zpotrf_gpu(uplo, n, db, lddb, info);
if (*info != 0) {
*info = n + *info;
return *info;
}
timer_stop( time );
timer_printf( "time zpotrf_gpu = %6.2f\n", time );
magma_queue_sync( stream );
magma_zgetmatrix_async( n, n,
db, lddb,
B, ldb, stream );
timer_start( time );
/* Transform problem to standard eigenvalue problem and solve. */
magma_zhegst_gpu(itype, uplo, n, da, ldda, db, lddb, info);
timer_stop( time );
timer_printf( "time zhegst_gpu = %6.2f\n", time );
/* simple fix to be able to run bigger size.
* need to have a dwork here that will be used
* a db and then passed to dsyevd.
* */
if (n > 5000) {
magma_queue_sync( stream );
magma_free( db );
}
timer_start( time );
magma_zheevd_gpu(jobz, uplo, n, da, ldda, w, A, lda,
work, lwork, rwork, lrwork, iwork, liwork, info);
timer_stop( time );
timer_printf( "time zheevd_gpu = %6.2f\n", time );
if (wantz && *info == 0) {
timer_start( time );
/* allocate and copy db back */
if (n > 5000) {
if (MAGMA_SUCCESS != magma_zmalloc( &db, n*lddb ) ) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_zsetmatrix( 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 = MagmaConjTrans;
} else {
trans = MagmaNoTrans;
}
magma_ztrsm(MagmaLeft, uplo, trans, MagmaNonUnit,
n, n, c_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 = MagmaConjTrans;
}
magma_ztrmm(MagmaLeft, uplo, trans, MagmaNonUnit,
n, n, c_one, db, lddb, da, ldda);
}
magma_zgetmatrix( n, n, da, ldda, A, lda );
/* free db */
if (n > 5000) {
magma_free( db );
}
timer_stop( time );
timer_printf( "time ztrsm/mm + getmatrix = %6.2f\n", time );
}
magma_queue_sync( stream );
magma_queue_destroy( stream );
work[0] = MAGMA_Z_MAKE( lwmin * one_eps, 0.); // round up
rwork[0] = lrwmin * one_eps;
iwork[0] = liwmin;
magma_free( da );
if (n <= 5000) {
magma_free( db );
}
return *info;
} /* magma_zhegvd */
/**
Purpose
-------
ZSSSSM applies the LU factorization update from a complex
matrix formed by a lower triangular IB-by-K tile L1 on top of a
M2-by-K tile L2 to a second complex matrix formed by a M1-by-N1
tile A1 on top of a M2-by-N2 tile A2 (N1 == N2).
This is the right-looking Level 2.5 BLAS version of the algorithm.
Arguments
---------
@param[in]
m1 INTEGER
The number of rows of the matrix A1. M1 >= 0.
@param[in]
n1 INTEGER
The number of columns of the matrix A1. N1 >= 0.
@param[in]
m2 INTEGER
The number of rows of the matrix A2. M2 >= 0.
@param[in]
n2 INTEGER
The number of columns of the matrix A2. N2 >= 0.
@param[in]
k INTEGER
The number of columns of the matrix L1 and L2. K >= 0.
@param[in]
ib INTEGER
The inner-blocking size. IB >= 0.
@param[in,out]
dA1 COMPLEX_16 array, dimension(LDDA1, N), on gpu.
On entry, the M1-by-N1 tile dA1.
On exit, dA1 is updated by the application of dL (dL1 dL2).
@param[in]
ldda1 INTEGER
The leading dimension of the array dA1. LDDA1 >= max(1,M1).
@param[in,out]
dA2 COMPLEX_16 array, dimension(LDDA2, N), on gpu.
On entry, the M2-by-N2 tile dA2.
On exit, dA2 is updated by the application of dL (dL1 dL2).
@param[in]
ldda2 INTEGER
The leading dimension of the array dA2. LDDA2 >= max(1,M2).
@param[in]
dL1 COMPLEX_16 array, dimension(LDDL1, K), on gpu.
The inverse of the IB-by-K lower triangular tile as returned by
ZTSTRF.
@param[in]
lddl1 INTEGER
The leading dimension of the array L1. LDDL1 >= max(1,2*IB).
@param[in]
dL2 COMPLEX_16 array, dimension(LDDL2, K)
The M2-by-K tile as returned by ZTSTRF.
@param[in]
lddl2 INTEGER
The leading dimension of the array L2. LDDL2 >= max(1,M2).
@param[in]
ipiv INTEGER array on the cpu.
The pivot indices array of size K as returned by ZTSTRF
@ingroup magma_zgesv_tile
********************************************************************/
extern "C" magma_int_t
magma_zssssm_gpu(
magma_order_t order, magma_int_t m1, magma_int_t n1,
magma_int_t m2, magma_int_t n2, magma_int_t k, magma_int_t ib,
magmaDoubleComplex_ptr dA1, magma_int_t ldda1,
magmaDoubleComplex_ptr dA2, magma_int_t ldda2,
magmaDoubleComplex_ptr dL1, magma_int_t lddl1,
magmaDoubleComplex_ptr dL2, magma_int_t lddl2,
magma_int_t *ipiv,
magma_int_t *info)
{
#define A1T(i,j) (dA1T + (i)*ldda1 + (j))
#define A2T(i,j) (dA2T + (i)*ldda2 + (j))
#define L1(i) (dL1 + (i)*lddl1 )
#define L2(i,j) (dL2 + (i)*lddl2i + (j)*lddl2j)
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
int ip, ii, sb;
magmaDoubleComplex_ptr dA1T, dA2T;
magma_trans_t transL;
int lddl2i, lddl2j;
MAGMA_UNUSED( ip ); // used only if NOSWAPBLK
/* Check input arguments */
*info = 0;
if (m1 < 0) {
*info = -1;
}
else if (n1 < 0) {
*info = -2;
}
else if (m2 < 0) {
*info = -3;
}
else if (n2 < 0) {
*info = -4;
}
else if (k < 0) {
*info = -5;
}
else if (ib < 0) {
*info = -6;
}
else if (ldda1 < max(1,m1)) {
*info = -8;
}
else if (ldda2 < max(1,m2)) {
*info = -10;
}
else if (lddl1 < max(1,ib)) {
*info = -12;
}
else if (lddl2 < max(1,m2)) {
*info = -14;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return */
if ((m1 == 0) || (n1 == 0) || (m2 == 0) || (n2 == 0) || (k == 0) || (ib == 0))
return *info;
if ( order == MagmaColMajor ) {
magmablas_zgetmo_in( dA1, dA1T, ldda1, m1, n1 );
magmablas_zgetmo_in( dA2, dA2T, ldda2, m2, n2 );
transL = MagmaTrans;
lddl2i = 1; lddl2j = lddl2;
} else {
dA1T = dA1;
dA2T = dA2;
transL = MagmaNoTrans;
lddl2i = lddl2; lddl2j = 1;
}
ip = 0;
for( ii=0; ii < k; ii += ib ) {
sb = min( k-ii, ib);
#ifndef NOSWAPBLK
magmablas_zswapblk( MagmaRowMajor, n1,
A1T(0, 0), ldda1,
A2T(0, 0), ldda2,
ii+1, ii+ib, ipiv, 1, m1 );
#else
{
int im;
for (i=0; i < ib; i++) {
im = ipiv[ip]-1;
if (im != (ii+i)) {
im = im - m1;
assert( (im >= 0) && (im < m1) && (im < m2) );
magmablas_zswap( n1, A1T(ii+i, 0), 1, A2T(im, 0), 1 );
}
ip++;
}
}
#endif
#ifndef WITHOUTTRTRI
/* Lower, Trans, because L1 is not transposed */
magma_ztrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n1, sb,
c_one, L1( ii), lddl1,
A1T(ii, 0), ldda1);
#else
/* Lower, Trans, because L1 is not transposed */
magma_ztrsm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n1, sb,
c_one, L1( ii), lddl1,
A1T(ii, 0), ldda1);
#endif
/* Second parameter is trans because L2 is not transposed */
magma_zgemm( MagmaNoTrans, transL,
n2, m2, sb,
c_neg_one, A1T(ii, 0), ldda1,
L2( 0, ii), lddl2,
c_one, A2T(0, 0 ), ldda2 );
}
if ( order == MagmaColMajor ) {
magmablas_zgetmo_out( dA1, dA1T, ldda1, m1, n1 );
magmablas_zgetmo_out( dA2, dA2T, ldda2, m2, n2 );
}
return *info;
}
extern "C" magma_int_t
magma_zhegvdx_2stage(magma_int_t itype, char jobz, char range, char uplo, magma_int_t n,
magmaDoubleComplex *a, magma_int_t lda,
magmaDoubleComplex *b, magma_int_t ldb,
double vl, double vu, magma_int_t il, magma_int_t iu,
magma_int_t *m, double *w,
magmaDoubleComplex *work, magma_int_t lwork,
double *rwork, magma_int_t lrwork,
magma_int_t *iwork, magma_int_t liwork,
magma_int_t *info)
{
/* -- MAGMA (version 1.4.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
December 2013
Purpose
=======
ZHEGVDX_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) COMPLEX_16 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) COMPLEX_16 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) 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 >= LQ2 + N * (NB + 1).
If JOBZ = 'V' and N > 1, LWORK >= LQ2 + 2*N + 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.
RWORK (workspace/output) DOUBLE PRECISION array, dimension (MAX(1,LRWORK))
On exit, if INFO = 0, RWORK(1) returns the optimal LRWORK.
LRWORK (input) INTEGER
The dimension of the array RWORK.
If N <= 1, LRWORK >= 1.
If JOBZ = 'N' and N > 1, LRWORK >= N.
If JOBZ = 'V' and N > 1, LRWORK >= 1 + 5*N + 2*N**2.
If LRWORK = -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};
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex *da;
magmaDoubleComplex *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 lopt;
magma_int_t lwmin;
// magma_int_t liopt;
magma_int_t liwmin;
// magma_int_t lropt;
magma_int_t lrwmin;
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 || lrwork == -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_zbulge_nb(n, threads);
magma_int_t lq2 = magma_zbulge_get_lq2(n, threads);
if (wantz) {
lwmin = lq2 + 2 * n + n * n;
lrwmin = 1 + 5 * n + 2 * n * n;
liwmin = 5 * n + 3;
} else {
lwmin = lq2 + n * (nb + 1);
lrwmin = n;
liwmin = 1;
}
work[0] = MAGMA_Z_MAKE( lwmin * (1. + lapackf77_dlamch("Epsilon")), 0.); // round up
rwork[0] = lrwmin * (1. + lapackf77_dlamch("Epsilon"));
iwork[0] = liwmin;
if (lwork < lwmin && ! lquery) {
*info = -17;
} else if (lrwork < lrwmin && ! lquery) {
*info = -19;
} else if (liwork < liwmin && ! lquery) {
*info = -21;
}
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_zhegvd(&itype, jobz_, uplo_,
&n, a, &lda, b, &ldb,
w, work, &lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, &lrwork,
#endif
iwork, &liwork, info);
*m = n;
return *info;
}
if (MAGMA_SUCCESS != magma_zmalloc( &da, n*ldda ) ||
MAGMA_SUCCESS != magma_zmalloc( &db, n*lddb )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/* Form a Cholesky factorization of B. */
magma_zsetmatrix( n, n, b, ldb, db, lddb );
magma_zsetmatrix_async( n, n,
a, lda,
da, ldda, stream );
#ifdef ENABLE_TIMER
magma_timestr_t start, end;
start = get_current_time();
#endif
magma_zpotrf_gpu(uplo_[0], n, db, lddb, info);
if (*info != 0) {
*info = n + *info;
return *info;
}
#ifdef ENABLE_TIMER
end = get_current_time();
printf("time zpotrf_gpu = %6.2f\n", GetTimerValue(start,end)/1000.);
#endif
magma_queue_sync( stream );
magma_zgetmatrix_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_zhegst_gpu(itype, uplo, n, da, ldda, db, lddb, info);
#ifdef ENABLE_TIMER
end = get_current_time();
printf("time zhegst_gpu = %6.2f\n", GetTimerValue(start,end)/1000.);
#endif
magma_zgetmatrix( 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_zheevdx_2stage(jobz, range, uplo, n, a, lda, vl, vu, il, iu, m, w, work, lwork, rwork, lrwork, iwork, liwork, info);
#ifdef ENABLE_TIMER
end = get_current_time();
printf("time zheevdx_2stage = %6.2f\n", GetTimerValue(start,end)/1000.);
#endif
if (wantz && *info == 0) {
if (MAGMA_SUCCESS != magma_zmalloc( &da, n*ldda ) ||
MAGMA_SUCCESS != magma_zmalloc( &db, n*lddb )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
#ifdef ENABLE_TIMER
start = get_current_time();
#endif
magma_zsetmatrix( n, *m, a, lda, da, ldda );
magma_zsetmatrix( 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_ztrsm(MagmaLeft, uplo, *trans, MagmaNonUnit, n, *m, c_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_ztrmm(MagmaLeft, uplo, *trans, MagmaNonUnit, n, *m, c_one, db, lddb, da, ldda);
}
magma_zgetmatrix( n, *m, da, ldda, a, lda );
#ifdef ENABLE_TIMER
end = get_current_time();
printf("time trsm/mm + getmatrix = %6.2f\n", GetTimerValue(start,end)/1000.);
#endif
magma_free( da );
magma_free( db );
}
magma_queue_destroy( stream );
work[0] = MAGMA_Z_MAKE( lwmin * (1. + lapackf77_dlamch("Epsilon")), 0.); // round up
rwork[0] = lrwmin * (1. + lapackf77_dlamch("Epsilon"));
iwork[0] = liwmin;
return *info;
} /* zhegvdx_2stage */
extern "C" magma_int_t
magma_zhegst(magma_int_t itype, char uplo, magma_int_t n,
magmaDoubleComplex *a, magma_int_t lda,
magmaDoubleComplex *b, magma_int_t ldb, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
ZHEGST reduces a complex Hermitian-definite generalized
eigenproblem to standard form.
If ITYPE = 1, the problem is A*x = lambda*B*x,
and A is overwritten by inv(U**H)*A*inv(U) or inv(L)*A*inv(L**H)
If ITYPE = 2 or 3, the problem is A*B*x = lambda*x or
B*A*x = lambda*x, and A is overwritten by U*A*U**H or L**H*A*L.
B must have been previously factorized as U**H*U or L*L**H by ZPOTRF.
Arguments
=========
ITYPE (input) INTEGER
= 1: compute inv(U**H)*A*inv(U) or inv(L)*A*inv(L**H);
= 2 or 3: compute U*A*U**H or L**H*A*L.
UPLO (input) CHARACTER*1
= 'U': Upper triangle of A is stored and B is factored as
U**H*U;
= 'L': Lower triangle of A is stored and B is factored as
L*L**H.
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 Hermitian matrix A. If UPLO = 'U', the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = 'L', the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit, if INFO = 0, the transformed matrix, stored in the
same format as A.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
B (input) COMPLEX_16 array, dimension (LDB,N)
The triangular factor from the Cholesky factorization of B,
as returned by ZPOTRF.
LDB (input) INTEGER
The leading dimension of the array B. LDB >= max(1,N).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
=====================================================================*/
char uplo_[2] = {uplo, 0};
magma_int_t nb;
magma_int_t k, kb, kb2;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_half = MAGMA_Z_HALF;
magmaDoubleComplex c_neg_half = MAGMA_Z_NEG_HALF;
magmaDoubleComplex *dw;
magma_int_t ldda = n;
magma_int_t lddb = n;
double d_one = 1.0;
int upper = lapackf77_lsame(uplo_, "U");
/* Test the input parameters. */
*info = 0;
if (itype<1 || itype>3){
*info = -1;
}else if ((! upper) && (! lapackf77_lsame(uplo_, "L"))) {
*info = -2;
} else if (n < 0) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
}else if (ldb < max(1,n)) {
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return */
if ( n == 0 )
return *info;
if (MAGMA_SUCCESS != magma_zmalloc( &dw, 2*n*n )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
nb = magma_get_zhegst_nb(n);
magma_queue_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
magma_zsetmatrix( n, n, A(0, 0), lda, dA(0, 0), ldda );
magma_zsetmatrix( n, n, B(0, 0), ldb, dB(0, 0), lddb );
/* Use hybrid blocked code */
if (itype==1) {
if (upper) {
/* Compute inv(U')*A*inv(U) */
for(k = 0; k<n; k+=nb){
kb = min(n-k,nb);
kb2= min(n-k-nb,nb);
/* Update the upper triangle of A(k:n,k:n) */
lapackf77_zhegst( &itype, uplo_, &kb, A(k,k), &lda, B(k,k), &ldb, info);
magma_zsetmatrix_async( kb, kb,
A(k, k), lda,
dA(k, k), ldda, stream[0] );
if(k+kb<n){
magma_ztrsm(MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
kb, n-k-kb,
c_one, dB(k,k), lddb,
dA(k,k+kb), ldda);
magma_queue_sync( stream[0] );
magma_zhemm(MagmaLeft, MagmaUpper,
kb, n-k-kb,
c_neg_half, dA(k,k), ldda,
dB(k,k+kb), lddb,
c_one, dA(k, k+kb), ldda);
magma_zher2k(MagmaUpper, MagmaConjTrans,
n-k-kb, kb,
c_neg_one, dA(k,k+kb), ldda,
dB(k,k+kb), lddb,
d_one, dA(k+kb,k+kb), ldda);
magma_zgetmatrix_async( kb2, kb2,
dA(k+kb, k+kb), ldda,
A(k+kb, k+kb), lda, stream[1] );
magma_zhemm(MagmaLeft, MagmaUpper,
kb, n-k-kb,
c_neg_half, dA(k,k), ldda,
dB(k,k+kb), lddb,
c_one, dA(k, k+kb), ldda);
magma_ztrsm(MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
kb, n-k-kb,
c_one ,dB(k+kb,k+kb), lddb,
dA(k,k+kb), ldda);
magma_queue_sync( stream[1] );
}
}
magma_queue_sync( stream[0] );
}
else {
/* Compute inv(L)*A*inv(L') */
for(k = 0; k<n; k+=nb){
kb= min(n-k,nb);
kb2= min(n-k-nb,nb);
/* Update the lower triangle of A(k:n,k:n) */
lapackf77_zhegst( &itype, uplo_, &kb, A(k,k), &lda, B(k,k), &ldb, info);
magma_zsetmatrix_async( kb, kb,
A(k, k), lda,
dA(k, k), ldda, stream[0] );
if(k+kb<n){
magma_ztrsm(MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
n-k-kb, kb,
c_one, dB(k,k), lddb,
dA(k+kb,k), ldda);
magma_queue_sync( stream[0] );
magma_zhemm(MagmaRight, MagmaLower,
n-k-kb, kb,
c_neg_half, dA(k,k), ldda,
dB(k+kb,k), lddb,
c_one, dA(k+kb, k), ldda);
magma_zher2k(MagmaLower, MagmaNoTrans,
n-k-kb, kb,
c_neg_one, dA(k+kb,k), ldda,
dB(k+kb,k), lddb,
d_one, dA(k+kb,k+kb), ldda);
magma_zgetmatrix_async( kb2, kb2,
dA(k+kb, k+kb), ldda,
A(k+kb, k+kb), lda, stream[1] );
magma_zhemm(MagmaRight, MagmaLower,
n-k-kb, kb,
c_neg_half, dA(k,k), ldda,
dB(k+kb,k), lddb,
c_one, dA(k+kb, k), ldda);
magma_ztrsm(MagmaLeft, MagmaLower, MagmaNoTrans, MagmaNonUnit,
n-k-kb, kb,
c_one, dB(k+kb,k+kb), lddb,
dA(k+kb,k), ldda);
}
magma_queue_sync( stream[1] );
}
}
magma_queue_sync( stream[0] );
}
else {
if (upper) {
/* Compute U*A*U' */
for(k = 0; k<n; k+=nb){
kb= min(n-k,nb);
magma_zgetmatrix_async( kb, kb,
dA(k, k), ldda,
A(k, k), lda, stream[0] );
/* Update the upper triangle of A(1:k+kb-1,1:k+kb-1) */
if(k>0){
magma_ztrmm(MagmaLeft, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
k, kb,
c_one ,dB(0,0), lddb,
dA(0,k), ldda);
magma_zhemm(MagmaRight, MagmaUpper,
k, kb,
c_half, dA(k,k), ldda,
dB(0,k), lddb,
c_one, dA(0, k), ldda);
magma_queue_sync( stream[1] );
magma_zher2k(MagmaUpper, MagmaNoTrans,
k, kb,
c_one, dA(0,k), ldda,
dB(0,k), lddb,
d_one, dA(0,0), ldda);
magma_zhemm(MagmaRight, MagmaUpper,
k, kb,
c_half, dA(k,k), ldda,
dB(0,k), lddb,
c_one, dA(0, k), ldda);
magma_ztrmm(MagmaRight, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
k, kb,
c_one, dB(k,k), lddb,
dA(0,k), ldda);
}
magma_queue_sync( stream[0] );
lapackf77_zhegst( &itype, uplo_, &kb, A(k, k), &lda, B(k, k), &ldb, info);
magma_zsetmatrix_async( kb, kb,
A(k, k), lda,
dA(k, k), ldda, stream[1] );
}
magma_queue_sync( stream[1] );
}
else {
/* Compute L'*A*L */
for(k = 0; k<n; k+=nb){
kb= min(n-k,nb);
magma_zgetmatrix_async( kb, kb,
dA(k, k), ldda,
A(k, k), lda, stream[0] );
/* Update the lower triangle of A(1:k+kb-1,1:k+kb-1) */
if(k>0){
magma_ztrmm(MagmaRight, MagmaLower, MagmaNoTrans, MagmaNonUnit,
kb, k,
c_one ,dB(0,0), lddb,
dA(k,0), ldda);
magma_zhemm(MagmaLeft, MagmaLower,
kb, k,
c_half, dA(k,k), ldda,
dB(k,0), lddb,
c_one, dA(k, 0), ldda);
magma_queue_sync( stream[1] );
magma_zher2k(MagmaLower, MagmaConjTrans,
k, kb,
c_one, dA(k,0), ldda,
dB(k,0), lddb,
d_one, dA(0,0), ldda);
magma_zhemm(MagmaLeft, MagmaLower,
kb, k,
c_half, dA(k,k), ldda,
dB(k,0), lddb,
c_one, dA(k, 0), ldda);
magma_ztrmm(MagmaLeft, MagmaLower, MagmaConjTrans, MagmaNonUnit,
kb, k,
c_one, dB(k,k), lddb,
dA(k,0), ldda);
}
magma_queue_sync( stream[0] );
lapackf77_zhegst( &itype, uplo_, &kb, A(k,k), &lda, B(k,k), &ldb, info);
magma_zsetmatrix_async( kb, kb,
A(k, k), lda,
dA(k, k), ldda, stream[1] );
}
magma_queue_sync( stream[1] );
}
}
magma_zgetmatrix( n, n, dA(0, 0), ldda, A(0, 0), lda );
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free( dw );
return *info;
} /* magma_zhegst_gpu */
/**
Purpose
-------
ZLARFB applies a complex block reflector H or its transpose H^H to a
COMPLEX_16 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)
- = Magma_ConjTrans: 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 COMPLEX_16 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 COMPLEX_16 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 COMPLEX_16 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. LDDC >= max(1,M).
@param
dwork (workspace) COMPLEX_16 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);
@param[in]
queue magma_queue_t
Queue to execute in.
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_zaux3
********************************************************************/
extern "C" magma_int_t
magma_zlarfb_gpu_q(
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,
magmaDoubleComplex_const_ptr dV, magma_int_t lddv,
magmaDoubleComplex_const_ptr dT, magma_int_t lddt,
magmaDoubleComplex_ptr dC, magma_int_t lddc,
magmaDoubleComplex_ptr dwork, magma_int_t ldwork,
magma_queue_t queue )
{
#define dV(i_,j_) (dV + (i_) + (j_)*lddv)
#define dT(i_,j_) (dT + (i_) + (j_)*lddt)
#define dC(i_,j_) (dC + (i_) + (j_)*lddc)
#define dwork(i_) (dwork + (i_))
/* Constants */
const magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
const magmaDoubleComplex c_one = MAGMA_Z_ONE;
const magmaDoubleComplex c_neg_one = MAGMA_Z_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;
}
/* Local variables */
// opposite of trans
magma_trans_t transt;
if (trans == MagmaNoTrans)
transt = Magma_ConjTrans;
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 = Magma_ConjTrans;
}
else {
notransV = Magma_ConjTrans;
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_zgemm( Magma_ConjTrans, notransV,
n, k, m,
c_one, dC(0,0), lddc,
dV(0,0), lddv,
c_zero, dwork(0), ldwork, queue );
// W = W T^H = C^H V T^H
magma_ztrmm( MagmaRight, uplo, transt, MagmaNonUnit,
n, k,
c_one, dT(0,0), lddt,
dwork(0), ldwork, queue );
// C = C - V W^H = C - V T V^H C = (I - V T V^H) C = H C
magma_zgemm( notransV, Magma_ConjTrans,
m, n, k,
c_neg_one, dV(0,0), lddv,
dwork(0), ldwork,
c_one, dC(0,0), lddc, queue );
}
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_zgemm( MagmaNoTrans, notransV,
m, k, n,
c_one, dC(0,0), lddc,
dV(0,0), lddv,
c_zero, dwork(0), ldwork, queue );
// W = W T = C V T
magma_ztrmm( MagmaRight, uplo, trans, MagmaNonUnit,
m, k,
c_one, dT(0,0), lddt,
dwork(0), ldwork, queue );
// C = C - W V^H = C - C V T V^H = C (I - V T V^H) = C H
magma_zgemm( MagmaNoTrans, transV,
m, n, k,
c_neg_one, dwork(0), ldwork,
dV(0,0), lddv,
c_one, dC(0,0), lddc, queue );
}
return info;
} /* magma_zlarfb */