extern "C" magma_err_t
magma_clarfb2_gpu( magma_int_t m, magma_int_t n, magma_int_t k,
const magmaFloatComplex_ptr dV, size_t dV_offset, magma_int_t ldv,
const magmaFloatComplex_ptr dT, size_t dT_offset, magma_int_t ldt,
magmaFloatComplex_ptr dC, size_t dC_offset, magma_int_t ldc,
magmaFloatComplex_ptr dwork, size_t dwork_offset, magma_int_t ldwork,
magma_queue_t queue )
{
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
if (m <= 0 || n <= 0)
return MAGMA_SUCCESS;
// W = C^H V
// magma_cgemm( MagmaConjTrans, MagmaNoTrans,
magmablas_cgemm_reduce(
n, k, m,
c_one, dC, dC_offset, ldc,
dV, dV_offset, ldv,
c_zero, dwork, dwork_offset, ldwork, queue);
// W = W T^H = C^H V T^H
magma_ctrmm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
n, k,
c_one, dT, dT_offset, ldt,
dwork, dwork_offset, ldwork,
queue);
// C = C - V W^H = C - V T V^H C = (I - V T V^H) C = H C
magma_cgemm( MagmaNoTrans, MagmaConjTrans,
m, n, k,
c_neg_one, dV, dV_offset, ldv,
dwork, dwork_offset, ldwork,
c_one, dC, dC_offset, ldc, queue );
}
/**
Purpose
-------
CGETRF_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 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 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 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 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 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_cgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_cgetrf_incpiv_gpu( magma_order_t order, magma_int_t m, magma_int_t n, magma_int_t ib,
magmaFloatComplex *hA, magma_int_t ldha, magmaFloatComplex *dA, magma_int_t ldda,
magmaFloatComplex *hL, magma_int_t ldhl, magmaFloatComplex *dL, magma_int_t lddl,
magma_int_t *ipiv,
magmaFloatComplex *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 )
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t iinfo;
magma_int_t maxm, mindim;
magma_int_t i, rows, cols, s, ii, sb;
magmaFloatComplex *dAT;
#ifndef WITHOUTTRTRI
magmaFloatComplex *dL2 = dL + ib;
magmaFloatComplex *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_cgetrf(&m, &n, hA, &ldha, ipiv, info);
#ifndef WITHOUTTRTRI
CORE_clacpy(PlasmaUpperLower, mindim, mindim,
(PLASMA_Complex32_t*)hA, ldha,
(PLASMA_Complex32_t*)hL2, ldhl );
CORE_ctrtri( PlasmaLower, PlasmaUnit, mindim,
(PLASMA_Complex32_t*)hL2, ldhl, info );
if (*info != 0 ) {
fprintf(stderr, "ERROR, trtri returned with info = %d\n", *info);
}
magma_csetmatrix( mindim, mindim, hL2, ldhl, dL2, lddl );
#endif
if ( order == MagmaRowMajor ) {
magma_csetmatrix( m, n, hA, ldha, dwork, lddwork );
magmablas_ctranspose( m, n, dwork, lddwork, dA, ldda );
} else {
magma_csetmatrix( m, n, hA, ldha, dA, ldda );
}
}
else {
/* Use hybrid blocked code. */
maxm = ((m + 31)/32)*32;
if ( order == MagmaColMajor ) {
magmablas_cgetmo_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_ctranspose( sb, m, AT(0,i), ldda, dwork, maxm );
magma_cgetmatrix( m, sb, dwork, maxm, hA(0, i), ldha );
// make sure that gpu queue is empty
//magma_device_sync();
#ifndef WITHOUTTRTRI
magma_ctrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n - (ii+sb), ib,
c_one, dL2(i-1), lddl,
AT(i-1,i+1), ldda );
#else
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n - (ii+sb), ib,
c_one, AT(i-1,i-1), ldda,
AT(i-1,i+1), ldda );
#endif
magma_cgemm( 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_cgetrf( &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_claswp( n-ii, AT(0, i), ldda, ii+1, ii+sb, ipiv, 1 );
#ifndef WITHOUTTRTRI
CORE_clacpy(PlasmaLower, sb, sb,
(PLASMA_Complex32_t*)hA(i, i), ldha,
(PLASMA_Complex32_t*)hL2(i), ldhl );
CORE_ctrtri( PlasmaLower, PlasmaUnit, sb,
(PLASMA_Complex32_t*)hL2(i), ldhl, info );
if (*info != 0 ) {
fprintf(stderr, "ERROR, trtri returned with info = %d\n", *info);
}
magma_csetmatrix( sb, sb, hL2(i), ldhl, dL2(i), lddl );
#endif
// upload i-th panel
magma_csetmatrix( rows, sb, hA(i, i), ldha, dwork, cols );
magmablas_ctranspose( rows, sb, dwork, cols, AT(i,i), ldda );
// do the small non-parallel computations
if ( s > (i+1) ) {
#ifndef WITHOUTTRTRI
magma_ctrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
sb, sb,
c_one, dL2(i), lddl,
AT(i, i+1), ldda);
#else
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
sb, sb,
c_one, AT(i, i ), ldda,
AT(i, i+1), ldda);
#endif
magma_cgemm( 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_ctrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n-mindim, sb,
c_one, dL2(i), lddl,
AT(i, i+1), ldda);
#else
magma_ctrsm( 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_cgemm( 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_cgetmo_out( dA, dAT, ldda, m, n );
}
}
return *info;
}
/**
Purpose
-------
CLARFB applies a complex block reflector H or its transpose H^H to a
COMPLEX 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 array on the GPU, dimension
(LDV,K) if STOREV = MagmaColumnwise
(LDV,M) if STOREV = MagmaRowwise and SIDE = MagmaLeft
(LDV,N) if STOREV = MagmaRowwise and SIDE = MagmaRight
The matrix V. See further details.
@param[in]
ldv INTEGER
The leading dimension of the array V.
If STOREV = MagmaColumnwise and SIDE = MagmaLeft, LDV >= max(1,M);
if STOREV = MagmaColumnwise and SIDE = MagmaRight, LDV >= max(1,N);
if STOREV = MagmaRowwise, LDV >= K.
@param[in]
dT COMPLEX array on the GPU, dimension (LDT,K)
The triangular k by k matrix T in the representation of the
block reflector.
@param[in]
ldt INTEGER
The leading dimension of the array T. LDT >= K.
@param[in,out]
dC COMPLEX 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.
@param[in]
ldc INTEGER
The leading dimension of the array C. LDA >= max(1,M).
@param
dwork (workspace) COMPLEX array, dimension (LDWORK,K)
@param[in]
ldwork INTEGER
The leading dimension of the array WORK.
If SIDE = MagmaLeft, LDWORK >= max(1,N);
if SIDE = MagmaRight, LDWORK >= max(1,M);
Further Details
---------------
The shape of the matrix V and the storage of the vectors which define
the H(i) is best illustrated by the following example with n = 5 and
k = 3.
All elements including 0's and 1's are stored, unlike LAPACK.
DIRECT = MagmaForward and DIRECT = MagmaForward and
STOREV = MagmaColumnwise: STOREV = MagmaRowwise:
V = ( 1 0 0 ) V = ( 1 v1 v1 v1 v1 )
( v1 1 0 ) ( 0 1 v2 v2 v2 )
( v1 v2 1 ) ( 0 0 1 v3 v3 )
( v1 v2 v3 )
( v1 v2 v3 )
DIRECT = MagmaBackward and DIRECT = MagmaBackward and
STOREV = MagmaColumnwise: STOREV = MagmaRowwise:
V = ( v1 v2 v3 ) V = ( v1 v1 1 0 0 )
( v1 v2 v3 ) ( v2 v2 v2 1 0 )
( 1 v2 v3 ) ( v3 v3 v3 v3 1 )
( 0 1 v3 )
( 0 0 1 )
@ingroup magma_caux3
********************************************************************/
extern "C" magma_int_t
magma_clarfb_gpu( magma_side_t side, magma_trans_t trans, magma_direct_t direct, magma_storev_t storev,
magma_int_t m, magma_int_t n, magma_int_t k,
const magmaFloatComplex *dV, magma_int_t ldv,
const magmaFloatComplex *dT, magma_int_t ldt,
magmaFloatComplex *dC, magma_int_t ldc,
magmaFloatComplex *dwork, magma_int_t ldwork )
{
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_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) && ldv < max(1,m)) ||
((storev == MagmaColumnwise) && (side == MagmaRight) && ldv < max(1,n)) ||
((storev == MagmaRowwise) && ldv < k) ) {
info = -9;
} else if (ldt < k) {
info = -11;
} else if (ldc < max(1,m)) {
info = -13;
} else if ( ((side == MagmaLeft) && ldwork < max(1,n)) ||
((side == MagmaRight) && ldwork < max(1,m)) ) {
info = -15;
}
if (info != 0) {
magma_xerbla( __func__, -(info) );
return info;
}
/* Function Body */
if (m <= 0 || n <= 0) {
return info;
}
// opposite of trans
magma_trans_t transt;
if (trans == MagmaNoTrans)
transt = 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_cgemm( Magma_ConjTrans, notransV,
n, k, m,
c_one, dC, ldc,
dV, ldv,
c_zero, dwork, ldwork);
// W = W T^H = C^H V T^H
magma_ctrmm( 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_cgemm( notransV, Magma_ConjTrans,
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_cgemm( MagmaNoTrans, notransV,
m, k, n,
c_one, dC, ldc,
dV, ldv,
c_zero, dwork, ldwork);
// W = W T = C V T
magma_ctrmm( 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_cgemm( MagmaNoTrans, transV,
m, n, k,
c_neg_one, dwork, ldwork,
dV, ldv,
c_one, dC, ldc);
}
return info;
} /* magma_clarfb */
/**
Purpose
-------
CHEGST_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 CPOTRF.
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 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, 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. LDA >= max(1,N).
@param[in]
dB COMPLEX array, dimension (LDB,N)
The triangular factor from the Cholesky factorization of B,
as returned by CPOTRF.
@param[in]
lddb INTEGER
The leading dimension of the array B. LDB >= max(1,N).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_cheev_comp
********************************************************************/
extern "C" magma_int_t
magma_chegst_gpu(magma_int_t itype, magma_uplo_t uplo, magma_int_t n,
magmaFloatComplex *dA, magma_int_t ldda,
magmaFloatComplex *dB, magma_int_t lddb, magma_int_t *info)
{
#define A(i, j) (w + (j)*lda + (i))
#define B(i, j) (w + nb*lda + (j)*ldb + (i))
#define dA(i, j) (dA + (j)*ldda + (i))
#define dB(i, j) (dB + (j)*lddb + (i))
const char* uplo_ = lapack_uplo_const( uplo );
magma_int_t nb;
magma_int_t k, kb, kb2;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex c_half = MAGMA_C_HALF;
magmaFloatComplex c_neg_half = MAGMA_C_NEG_HALF;
magmaFloatComplex *w;
magma_int_t lda;
magma_int_t ldb;
float d_one = 1.0;
int 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_chegst_nb(n);
lda = nb;
ldb = nb;
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &w, 2*nb*nb )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_queue_t stream[3];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
magma_queue_create( &stream[2] );
/* Use hybrid blocked code */
if (itype == 1) {
if (upper) {
kb = min(n,nb);
/* Compute inv(U')*A*inv(U) */
magma_cgetmatrix_async( kb, kb,
dB(0, 0), lddb,
B(0, 0), nb, stream[2] );
magma_cgetmatrix_async( kb, kb,
dA(0, 0), ldda,
A(0, 0), nb, stream[1] );
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) */
magma_queue_sync( stream[2] );
magma_queue_sync( stream[1] );
lapackf77_chegst( &itype, uplo_, &kb, A(0,0), &lda, B(0,0), &ldb, info);
magma_csetmatrix_async( kb, kb,
A(0, 0), lda,
dA(k, k), ldda, stream[0] );
if (k+kb < n) {
// Start copying the new B block
magma_cgetmatrix_async( kb2, kb2,
dB(k+kb, k+kb), lddb,
B(0, 0), nb, stream[2] );
magma_ctrsm(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_chemm(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_cher2k(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_cgetmatrix_async( kb2, kb2,
dA(k+kb, k+kb), ldda,
A(0, 0), lda, stream[1] );
magma_chemm(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_ctrsm(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[0] );
}
else {
kb = min(n,nb);
/* Compute inv(L)*A*inv(L') */
magma_cgetmatrix_async( kb, kb,
dB(0, 0), lddb,
B(0, 0), nb, stream[2] );
magma_cgetmatrix_async( kb, kb,
dA(0, 0), ldda,
A(0, 0), nb, stream[1] );
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) */
magma_queue_sync( stream[2] );
magma_queue_sync( stream[1] );
lapackf77_chegst( &itype, uplo_, &kb, A(0, 0), &lda, B(0, 0), &ldb, info);
magma_csetmatrix_async( kb, kb,
A(0, 0), lda,
dA(k, k), ldda, stream[0] );
if (k+kb < n) {
// Start copying the new B block
magma_cgetmatrix_async( kb2, kb2,
dB(k+kb, k+kb), lddb,
B(0, 0), nb, stream[2] );
magma_ctrsm(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_chemm(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_cher2k(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_cgetmatrix_async( kb2, kb2,
dA(k+kb, k+kb), ldda,
A(0, 0), lda, stream[1] );
magma_chemm(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_ctrsm(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[0] );
}
else {
if (upper) {
/* Compute U*A*U' */
for (k = 0; k < n; k += nb) {
kb= min(n-k,nb);
magma_cgetmatrix_async( kb, kb,
dB(k, k), lddb,
B(0, 0), nb, stream[2] );
/* Update the upper triangle of A(1:k+kb-1,1:k+kb-1) */
if (k > 0) {
magma_ctrmm(MagmaLeft, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
k, kb,
c_one, dB(0,0), lddb,
dA(0,k), ldda);
magma_chemm(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_cgetmatrix_async( kb, kb,
dA(k, k), ldda,
A(0, 0), lda, stream[0] );
if (k > 0) {
magma_cher2k(MagmaUpper, MagmaNoTrans,
k, kb,
c_one, dA(0,k), ldda,
dB(0,k), lddb,
d_one, dA(0,0), ldda);
magma_chemm(MagmaRight, MagmaUpper,
k, kb,
c_half, dA(k,k), ldda,
dB(0,k), lddb,
c_one, dA(0, k), ldda);
magma_ctrmm(MagmaRight, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
k, kb,
c_one, dB(k,k), lddb,
dA(0,k), ldda);
}
magma_queue_sync( stream[2] );
magma_queue_sync( stream[0] );
lapackf77_chegst( &itype, uplo_, &kb, A(0, 0), &lda, B(0, 0), &ldb, info);
magma_csetmatrix_async( kb, kb,
A(0, 0), 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_cgetmatrix_async( kb, kb,
dB(k, k), lddb,
B(0, 0), nb, stream[2] );
/* Update the lower triangle of A(1:k+kb-1,1:k+kb-1) */
if (k > 0) {
magma_ctrmm(MagmaRight, MagmaLower, MagmaNoTrans, MagmaNonUnit,
kb, k,
c_one, dB(0,0), lddb,
dA(k,0), ldda);
magma_chemm(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_cgetmatrix_async( kb, kb,
dA(k, k), ldda,
A(0, 0), lda, stream[0] );
if (k > 0) {
magma_cher2k(MagmaLower, MagmaConjTrans,
k, kb,
c_one, dA(k,0), ldda,
dB(k,0), lddb,
d_one, dA(0,0), ldda);
magma_chemm(MagmaLeft, MagmaLower,
kb, k,
c_half, dA(k,k), ldda,
dB(k,0), lddb,
c_one, dA(k, 0), ldda);
magma_ctrmm(MagmaLeft, MagmaLower, MagmaConjTrans, MagmaNonUnit,
kb, k,
c_one, dB(k,k), lddb,
dA(k,0), ldda);
}
magma_queue_sync( stream[2] );
magma_queue_sync( stream[0] );
lapackf77_chegst( &itype, uplo_, &kb, A(0, 0), &lda, B(0, 0), &ldb, info);
magma_csetmatrix_async( kb, kb,
A(0, 0), lda,
dA(k, k), ldda, stream[1] );
}
magma_queue_sync( stream[1] );
}
}
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_queue_destroy( stream[2] );
magma_free_pinned( w );
return *info;
} /* magma_chegst_gpu */
extern "C" magma_int_t
magma_cgessm_gpu( char storev, magma_int_t m, magma_int_t n, magma_int_t k, magma_int_t ib,
magma_int_t *ipiv,
cuFloatComplex *dL1, magma_int_t lddl1,
cuFloatComplex *dL, magma_int_t lddl,
cuFloatComplex *dA, magma_int_t ldda,
magma_int_t *info)
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
Purpose
=======
CGESSM applies the factors L computed by CGETRF_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
CGETRF_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)
cuFloatComplex c_one = MAGMA_C_ONE;
cuFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
int i, s, sb;
cuFloatComplex *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_cgetmo_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_claswp( n, dAT, ldda, i+1, i+sb, ipiv, 1 );
#ifndef WITHOUTTRTRI
magma_ctrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n, sb,
c_one, dL1(i), lddl1,
AT(i, 0), ldda);
#else
magma_ctrsm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n, sb,
c_one, L( i, i), lddl,
AT(i, 0), ldda);
#endif
if ( (i+sb) < m) {
magma_cgemm( 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_cgetmo_in( dA, dAT, ldda, m, n );
}
return *info;
/* End of MAGMA_CGETRF_GPU */
}
/**
Purpose
-------
CTRTRI computes the inverse of a real upper or lower triangular
matrix dA.
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]
dA COMPLEX array ON THE GPU, dimension (LDDA,N)
On entry, the triangular matrix A. If UPLO = MagmaUpper, 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 = MagmaLower, 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 = 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]
ldda INTEGER
The leading dimension of the array dA. LDDA >= 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, dA(i,i) is exactly zero. The triangular
matrix is singular and its inverse cannot be computed.
(Singularity check is currently disabled.)
@ingroup magma_cgesv_aux
********************************************************************/
extern "C" magma_int_t
magma_ctrtri_gpu(
magma_uplo_t uplo, magma_diag_t diag, magma_int_t n,
magmaFloatComplex_ptr dA, magma_int_t ldda,
magma_int_t *info)
{
#define dA(i, j) (dA+(j)*ldda + (i))
/* Local variables */
const char* uplo_ = lapack_uplo_const( uplo );
const char* diag_ = lapack_diag_const( diag );
magma_int_t nb, nn, j, jb;
//magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *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_C_EQUAL( *dA(j,j), c_zero )) {
*info = j+1; // Fortran index
return *info;
}
}
}
*/
/* Determine the block size for this environment */
nb = magma_get_cpotrf_nb(n);
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &work, nb*nb )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_queue_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
if (nb <= 1 || nb >= n) {
magma_cgetmatrix( n, n, dA, ldda, work, n );
lapackf77_ctrtri( uplo_, diag_, &n, work, &n, info );
magma_csetmatrix( n, n, work, n, dA, ldda );
}
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_ctrmm( MagmaLeft, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_one, dA(0,0), ldda, dA(0, j), ldda );
magma_ctrsm( MagmaRight, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_neg_one, dA(j,j), ldda, dA(0, j), ldda );
magma_cgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, jb, stream[1] );
magma_queue_sync( stream[1] );
/* Compute inverse of current diagonal block */
lapackf77_ctrtri( MagmaUpperStr, diag_, &jb, work, &jb, info );
magma_csetmatrix_async( jb, jb,
work, jb,
dA(j, j), ldda, stream[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_ctrmm( MagmaLeft, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_one, dA(j+jb,j+jb), ldda, dA(j+jb, j), ldda );
magma_ctrsm( MagmaRight, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_neg_one, dA(j,j), ldda, dA(j+jb, j), ldda );
}
magma_cgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, jb, stream[1] );
magma_queue_sync( stream[1] );
/* Compute inverse of current diagonal block */
lapackf77_ctrtri( MagmaLowerStr, diag_, &jb, work, &jb, info );
magma_csetmatrix_async( jb, jb,
work, jb,
dA(j, j), ldda, stream[0] );
}
}
}
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free_pinned( work );
return *info;
}
/**
Purpose
-------
CHEGVX 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
---------
@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 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, the lower triangle (if UPLO=MagmaLower) or the upper
triangle (if UPLO=MagmaUpper) of A, including the diagonal, is
destroyed.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[in,out]
B COMPLEX 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 REAL
@param[in]
vu REAL
If RANGE=MagmaRangeV, the lower and upper bounds of the interval to
be searched for eigenvalues. VL < VU.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeI.
@param[in]
il INTEGER
@param[in]
iu INTEGER
If RANGE=MagmaRangeI, the indices (in ascending order) of the
smallest and largest eigenvalues to be returned.
1 <= IL <= IU <= N, if N > 0; IL = 1 and IU = 0 if N = 0.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeV.
@param[in]
abstol REAL
The absolute error tolerance for the eigenvalues.
An approximate eigenvalue is accepted as converged
when it is determined to lie in an interval [a,b]
of width less than or equal to
\n
ABSTOL + EPS * max( |a|,|b| ),
\n
where EPS is the machine precision. If ABSTOL is less than
or equal to zero, then EPS*|T| will be used in its place,
where |T| is the 1-norm of the tridiagonal matrix obtained
by reducing A to tridiagonal form.
\n
Eigenvalues will be computed most accurately when ABSTOL is
set to twice the underflow threshold 2*SLAMCH('S'), not zero.
If this routine returns with INFO > 0, indicating that some
eigenvectors did not converge, try setting ABSTOL to
2*SLAMCH('S').
@param[out]
m INTEGER
The total number of eigenvalues found. 0 <= M <= N.
If RANGE = MagmaRangeAll, M = N, and if RANGE = MagmaRangeI, M = IU-IL+1.
@param[out]
w REAL array, dimension (N)
The first M elements contain the selected
eigenvalues in ascending order.
@param[out]
Z COMPLEX array, dimension (LDZ, max(1,M))
If JOBZ = MagmaNoVec, then Z is not referenced.
If JOBZ = MagmaVec, then if INFO = 0, the first M columns of Z
contain the orthonormal eigenvectors of the matrix A
corresponding to the selected eigenvalues, with the i-th
column of Z holding the eigenvector associated with W(i).
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.
\n
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 = MagmaRangeV, the exact value of M
is not known in advance and an upper bound must be used.
@param[in]
ldz INTEGER
The leading dimension of the array Z. LDZ >= 1, and if
JOBZ = MagmaVec, LDZ >= max(1,N).
@param[out]
work (workspace) COMPLEX 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. LWORK >= max(1,2*N).
For optimal efficiency, LWORK >= (NB+1)*N,
where NB is the blocksize for CHETRD returned by ILAENV.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param
rwork (workspace) REAL array, dimension (7*N)
@param
iwork (workspace) INTEGER array, dimension (5*N)
@param[out]
ifail INTEGER array, dimension (N)
If JOBZ = MagmaVec, then if INFO = 0, the first M elements of
IFAIL are zero. If INFO > 0, then IFAIL contains the
indices of the eigenvectors that failed to converge.
If JOBZ = MagmaNoVec, then IFAIL is not referenced.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: CPOTRF or CHEEVX returned an error code:
<= N: if INFO = i, CHEEVX 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
@ingroup magma_chegv_driver
********************************************************************/
extern "C" magma_int_t
magma_chegvx(
magma_int_t itype, magma_vec_t jobz, magma_range_t range, magma_uplo_t uplo, magma_int_t n,
magmaFloatComplex *A, magma_int_t lda, magmaFloatComplex *B, magma_int_t ldb,
float vl, float vu, magma_int_t il, magma_int_t iu, float abstol,
magma_int_t *m, float *w, magmaFloatComplex *Z, magma_int_t ldz,
magmaFloatComplex *work, magma_int_t lwork, float *rwork,
magma_int_t *iwork, magma_int_t *ifail,
magma_int_t *info)
{
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex *dA;
magmaFloatComplex *dB;
magmaFloatComplex *dZ;
magma_int_t ldda = n;
magma_int_t lddb = n;
magma_int_t lddz = 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_queue_t stream;
magma_queue_create( &stream );
wantz = (jobz == MagmaVec);
lower = (uplo == MagmaLower);
alleig = (range == MagmaRangeAll);
valeig = (range == MagmaRangeV);
indeig = (range == MagmaRangeI);
lquery = (lwork == -1);
*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 (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_chetrd_nb(n);
lwmin = n * (nb + 1);
work[0] = MAGMA_C_MAKE( lwmin, 0 );
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_cmalloc( &dA, n*ldda ) ||
MAGMA_SUCCESS != magma_cmalloc( &dB, n*lddb ) ||
MAGMA_SUCCESS != magma_cmalloc( &dZ, n*lddz )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/* Form a Cholesky factorization of B. */
magma_csetmatrix( n, n, B, ldb, dB, lddb );
magma_csetmatrix_async( n, n,
A, lda,
dA, ldda, stream );
magma_cpotrf_gpu(uplo, n, dB, lddb, info);
if (*info != 0) {
*info = n + *info;
return *info;
}
magma_queue_sync( stream );
magma_cgetmatrix_async( n, n,
dB, lddb,
B, ldb, stream );
/* Transform problem to standard eigenvalue problem and solve. */
magma_chegst_gpu(itype, uplo, n, dA, ldda, dB, lddb, info);
magma_cheevx_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) {
trans = MagmaConjTrans;
} else {
trans = MagmaNoTrans;
}
magma_ctrsm(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) {
trans = MagmaNoTrans;
} else {
trans = MagmaConjTrans;
}
magma_ctrmm(MagmaLeft, uplo, trans, MagmaNonUnit, n, *m, c_one, dB, lddb, dZ, lddz);
}
magma_cgetmatrix( 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;
} /* magma_chegvx */
/**
Purpose
-------
CLAUUM 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_cposv_aux
***************************************************************************/
extern "C" magma_int_t
magma_clauum(magma_uplo_t uplo, magma_int_t n,
magmaFloatComplex *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;
magmaFloatComplex c_one = MAGMA_C_ONE;
float d_one = MAGMA_D_ONE;
magmaFloatComplex *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_cmalloc( &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_cpotrf_nb(n);
if (nb <= 1 || nb >= n)
lapackf77_clauum(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_csetmatrix_async( ib, ib,
A(i,i), lda,
dA(i, i), ldda, stream[1] );
magma_csetmatrix_async( ib, (n-i-ib),
A(i,i+ib), lda,
dA(i,i+ib), ldda, stream[0] );
magma_queue_sync( stream[1] );
magma_ctrmm( MagmaRight, MagmaUpper,
MagmaConjTrans, MagmaNonUnit, i, ib,
c_one, dA(i,i), ldda, dA(0, i),ldda);
lapackf77_clauum(MagmaUpperStr, &ib, A(i,i), &lda, info);
magma_csetmatrix_async( ib, ib,
A(i, i), lda,
dA(i, i), ldda, stream[0] );
if (i+ib < n) {
magma_cgemm( 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_cherk( MagmaUpper, MagmaNoTrans, ib,(n-i-ib),
d_one, dA(i, i+ib), ldda,
d_one, dA(i, i), ldda);
}
magma_cgetmatrix( 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_csetmatrix_async( ib, ib,
A(i,i), lda,
dA(i, i), ldda, stream[1] );
magma_csetmatrix_async( (n-i-ib), ib,
A(i+ib, i), lda,
dA(i+ib, i), ldda, stream[0] );
magma_queue_sync( stream[1] );
magma_ctrmm( MagmaLeft, MagmaLower,
MagmaConjTrans, MagmaNonUnit, ib,
i, c_one, dA(i,i), ldda,
dA(i, 0),ldda);
lapackf77_clauum(MagmaLowerStr, &ib, A(i,i), &lda, info);
magma_csetmatrix_async( ib, ib,
A(i, i), lda,
dA(i, i), ldda, stream[0] );
if (i+ib < n) {
magma_cgemm(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_cherk(MagmaLower, MagmaConjTrans, ib, (n-i-ib),
d_one, dA(i+ib, i), ldda,
d_one, dA(i, i), ldda);
}
magma_cgetmatrix( 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;
}
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, t1, t2;
magmaFloatComplex c_neg_one = MAGMA_C_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 };
magmaFloatComplex *A, *B, *C, *C2, *LU;
magmaFloatComplex_ptr dA, dB, dC1, dC2;
magmaFloatComplex alpha = MAGMA_C_MAKE( 0.5, 0.1 );
magmaFloatComplex beta = MAGMA_C_MAKE( 0.7, 0.2 );
float dalpha = 0.6;
float dbeta = 0.8;
float work[1], error, total_error;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t m, n, k, size, maxn, ld, info;
magma_int_t *piv;
magma_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_cmalloc_pinned( &A, size ); assert( err == 0 );
err = magma_cmalloc_pinned( &B, size ); assert( err == 0 );
err = magma_cmalloc_pinned( &C, size ); assert( err == 0 );
err = magma_cmalloc_pinned( &C2, size ); assert( err == 0 );
err = magma_cmalloc_pinned( &LU, size ); assert( err == 0 );
err = magma_cmalloc( &dA, size ); assert( err == 0 );
err = magma_cmalloc( &dB, size ); assert( err == 0 );
err = magma_cmalloc( &dC1, size ); assert( err == 0 );
err = magma_cmalloc( &dC2, size ); assert( err == 0 );
// initialize matrices
size = maxn*maxn;
lapackf77_clarnv( &ione, ISEED, &size, A );
lapackf77_clarnv( &ione, ISEED, &size, B );
lapackf77_clarnv( &ione, ISEED, &size, C );
printf( "========== Level 1 BLAS ==========\n" );
// ----- test CSWAP
// swap columns 2 and 3 of dA, then copy to C2 and compare with A
if ( n >= 3 ) {
magma_csetmatrix( m, n, A, ld, dA, ld );
magma_csetmatrix( m, n, A, ld, dB, ld );
magma_cswap( m, dA(0,1), 1, dA(0,2), 1 );
magma_cswap( m, dB(0,1), 1, dB(0,2), 1 );
// check results, storing diff between magma and cuda calls in C2
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dA, 1, dB, 1 );
magma_cgetmatrix( m, n, dB, ld, C2, ld );
error = lapackf77_clange( "F", &m, &k, C2, &ld, work );
total_error += error;
printf( "cswap diff %.2g\n", error );
}
else {
printf( "cswap skipped for n < 3\n" );
}
// ----- test ICAMAX
// get argmax of column of A
magma_csetmatrix( m, k, A, ld, dA, ld );
error = 0;
for( int j = 0; j < k; ++j ) {
magma_int_t i1 = magma_icamax( m, dA(0,j), 1 );
int i2; // NOT magma_int_t, for cublas
cublasIcamax( opts.handle, m, dA(0,j), 1, &i2 );
// todo need sync here?
assert( i1 == i2 );
error += abs( i1 - i2 );
}
total_error += error;
gflops = (float)m * k / 1e9;
printf( "icamax diff %.2g\n", error );
printf( "\n" );
printf( "========== Level 2 BLAS ==========\n" );
// ----- test CGEMV
// 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_csetmatrix( m, n, A, ld, dA, ld );
magma_csetvector( maxn, B, 1, dB, 1 );
magma_csetvector( maxn, C, 1, dC1, 1 );
magma_csetvector( maxn, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_cgemv( trans[ia], m, n, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasCgemv( opts.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);
cublasCaxpy( opts.handle, size, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetvector( size, dC2, 1, C2, 1 );
error = lapackf77_clange( "F", &size, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_CGEMV( m, n ) / 1e9;
printf( "cgemv( %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_trans_const(trans[ia]), error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test CHEMV
// 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_csetmatrix( m, m, A, ld, dA, ld );
magma_csetvector( m, B, 1, dB, 1 );
magma_csetvector( m, C, 1, dC1, 1 );
magma_csetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_chemv( uplo[iu], m, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasChemv( opts.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
cublasCaxpy( opts.handle, m, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_clange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_CHEMV( m ) / 1e9;
printf( "chemv( %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test CTRSV
// 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_clacpy( "Full", &maxn, &maxn, A, &ld, LU, &ld );
lapackf77_cgetrf( &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_csetmatrix( m, m, LU, ld, dA, ld );
magma_csetvector( m, C, 1, dC1, 1 );
magma_csetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_ctrsv( uplo[iu], trans[it], diag[id], m, dA, ld, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasCtrsv( opts.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
cublasCaxpy( opts.handle, m, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_clange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_CTRSM( MagmaLeft, m, 1 ) / 1e9;
printf( "ctrsv( %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 CGEMM
// 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_csetmatrix( (nta ? m : k), (nta ? m : k), A, ld, dA, ld );
magma_csetmatrix( (ntb ? k : n), (ntb ? n : k), B, ld, dB, ld );
magma_csetmatrix( m, n, C, ld, dC1, ld );
magma_csetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_cgemm( 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 );
cublasCgemm( opts.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
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_clange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_CGEMM( m, n, k ) / 1e9;
printf( "cgemm( %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 CHEMM
// 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_csetmatrix( m, m, A, ld, dA, ld );
magma_csetmatrix( m, n, B, ld, dB, ld );
magma_csetmatrix( m, n, C, ld, dC1, ld );
magma_csetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_chemm( 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 );
cublasChemm( opts.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
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_clange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_CHEMM( side[is], m, n ) / 1e9;
printf( "chemm( %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 CHERK
// 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_csetmatrix( n, k, A, ld, dA, ld );
magma_csetmatrix( n, n, C, ld, dC1, ld );
magma_csetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_cherk( uplo[iu], trans[it], n, k, dalpha, dA, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasCherk( opts.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
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_clange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_CHERK( k, n ) / 1e9;
printf( "cherk( %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 CHER2K
// 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_csetmatrix( (nt ? n : k), (nt ? n : k), A, ld, dA, ld );
magma_csetmatrix( n, n, C, ld, dC1, ld );
magma_csetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_cher2k( 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 );
cublasCher2k( opts.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
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_clange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_CHER2K( k, n ) / 1e9;
printf( "cher2k( %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 CTRMM
// 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_csetmatrix( (left ? m : n), (left ? m : n), A, ld, dA, ld );
magma_csetmatrix( m, n, C, ld, dC1, ld );
magma_csetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ctrmm( 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 );
cublasCtrmm( opts.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
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_clange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_CTRMM( side[is], m, n ) / 1e9;
printf( "ctrmm( %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 CTRSM
// 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_csetmatrix( (left ? m : n), (left ? m : n), LU, ld, dA, ld );
magma_csetmatrix( m, n, C, ld, dC1, ld );
magma_csetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ctrsm( 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 );
cublasCtrsm( opts.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
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_clange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_CTRSM( side[is], m, n ) / 1e9;
printf( "ctrsm( %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;
}
extern "C" magma_int_t
magma_ctrtri_gpu(
magma_uplo_t uplo, magma_diag_t diag, magma_int_t n,
magmaFloatComplex_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
=======
CTRTRI 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 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;
//magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *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_C_EQUAL( *dA(j,j), c_zero )) {
*info = j+1; // Fortran index
return *info;
}
}
}
*/
/* Determine the block size for this environment */
nb = magma_get_cpotrf_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_cmalloc_cpu( &work, nb*nb )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
if (nb <= 1 || nb >= n) {
magma_cgetmatrix( n, n, dA, dA_offset, ldda, work, n, queues[0] );
lapackf77_ctrtri( lapack_const(uplo), lapack_const(diag), &n, work, &n, info );
magma_csetmatrix( 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_ctrmm( MagmaLeft, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_one, dA(0,0), ldda, dA(0, j), ldda,
queues[0] );
magma_ctrsm( MagmaRight, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_neg_one, dA(j,j), ldda, dA(0, j), ldda,
queues[0] );
magma_cgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, jb, queues[1], NULL );
magma_queue_sync( queues[1] );
/* Compute inverse of current diagonal block */
lapackf77_ctrtri( MagmaUpperStr, lapack_const(diag), &jb, work, &jb, info );
/*
magma_csetmatrix_async( jb, jb,
work, 0, jb,
dA(j, j), ldda, queues[0], NULL );
*/
magma_csetmatrix( 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_ctrmm( MagmaLeft, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_one, dA(j+jb,j+jb), ldda, dA(j+jb, j), ldda,
queues[0] );
magma_ctrsm( MagmaRight, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_neg_one, dA(j,j), ldda, dA(j+jb, j), ldda,
queues[0] );
}
magma_cgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, jb, queues[1], NULL );
magma_queue_sync( queues[1] );
/* Compute inverse of current diagonal block */
lapackf77_ctrtri( MagmaLowerStr, lapack_const(diag), &jb, work, &jb, info );
/*
magma_csetmatrix_async( jb, jb,
work, 0, jb,
dA(j, j), ldda, queues[0], NULL );
*/
magma_csetmatrix( 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
-------
CTRTRI 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 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_cgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_ctrtri(
magma_uplo_t uplo, magma_diag_t diag, magma_int_t n,
magmaFloatComplex *A, magma_int_t lda,
magma_int_t *info)
{
#define A(i_, j_) ( A + (i_) + (j_)*lda )
#ifdef HAVE_clBLAS
#define dA(i_, j_) dA, ((i_) + (j_)*ldda)
#else
#define dA(i_, j_) (dA + (i_) + (j_)*ldda)
#endif
// Constants
const magmaFloatComplex c_zero = MAGMA_C_ZERO;
const magmaFloatComplex c_one = MAGMA_C_ONE;
const magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
const char* uplo_ = lapack_uplo_const( uplo );
const char* diag_ = lapack_diag_const( diag );
// Local variables
magma_int_t ldda, nb, nn, j, jb;
magmaFloatComplex_ptr dA;
bool upper = (uplo == MagmaUpper);
bool 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_C_EQUAL( *A(j,j), c_zero )) {
*info = j+1; // Fortran index
return *info;
}
}
}
// Determine the block size for this environment
nb = magma_get_cpotrf_nb( n );
ldda = magma_roundup( n, 32 );
if (MAGMA_SUCCESS != magma_cmalloc( &dA, (n)*ldda )) {
*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] ); // unused
if (nb <= 1 || nb >= n) {
lapackf77_ctrtri( 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 );
if (j > 0) {
// Send current block column (with diagonal) to device
// This must finish before trtri below
magma_csetmatrix( j+jb, jb,
A(0,j), lda,
dA(0,j), ldda, queues[0] );
// Compute rows 0:j of current block column
magma_ctrmm( MagmaLeft, MagmaUpper,
MagmaNoTrans, diag, j, jb, c_one,
dA(0,0), ldda,
dA(0,j), ldda, queues[0] );
magma_ctrsm( MagmaRight, MagmaUpper,
MagmaNoTrans, diag, j, jb, c_neg_one,
dA(j,j), ldda,
dA(0,j), ldda, queues[0] );
// Get above diagonal from device
// TODO: could be on another queue, after trmm/trsm finish
magma_cgetmatrix_async( j, jb,
dA(0,j), ldda,
A(0,j), lda, queues[0] );
}
// Compute inverse of current diagonal block
// TODO: problem if diagonal has not finished sending yet?
lapackf77_ctrtri( MagmaUpperStr, diag_, &jb, A(j,j), &lda, info );
if (j+jb < n) {
// Send inverted diagonal block to device
magma_csetmatrix( jb, jb,
A(j,j), lda,
dA(j,j), ldda, queues[0] );
}
}
}
else {
// Compute inverse of lower triangular matrix
nn = ((n-1)/nb)*nb;
for (j=nn; j >= 0; j -= nb) {
jb = min( nb, n-j );
if (j+jb < n) {
// Send current block row (with diagonal) to device
// This must finish before trtri below
magma_csetmatrix( n-j, jb,
A(j,j), lda,
dA(j,j), ldda, queues[0] );
// Compute rows j+jb:n of current block column
magma_ctrmm( MagmaLeft, MagmaLower,
MagmaNoTrans, diag, n-j-jb, jb, c_one,
dA(j+jb,j+jb), ldda,
dA(j+jb,j), ldda, queues[0] );
magma_ctrsm( MagmaRight, MagmaLower,
MagmaNoTrans, diag, n-j-jb, jb, c_neg_one,
dA(j,j), ldda,
dA(j+jb,j), ldda, queues[0] );
// Get below diagonal block from device
magma_cgetmatrix_async( n-j-jb, jb,
dA(j+jb,j), ldda,
A(j+jb,j), lda, queues[0] );
}
// Compute inverse of current diagonal block
lapackf77_ctrtri( MagmaLowerStr, diag_, &jb, A(j,j), &lda, info );
if (j > 0) {
// Send inverted diagonal block to device
magma_csetmatrix( jb, jb,
A(j,j), lda,
dA(j,j), ldda, queues[0] );
}
}
}
magma_queue_destroy( queues[0] );
//magma_queue_destroy( queues[1] ); // unused
magma_free( dA );
return *info;
}
extern "C" magma_int_t
magma_ctrtri(char uplo, char diag, magma_int_t n,
cuFloatComplex *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
=======
CTRTRI 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 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.
===================================================================== */
/* Local variables */
char uplo_[2] = {uplo, 0};
char diag_[2] = {diag, 0};
magma_int_t ldda, nb, nn, j, jb;
cuFloatComplex c_zero = MAGMA_C_ZERO;
cuFloatComplex c_one = MAGMA_C_ONE;
cuFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
cuFloatComplex *work;
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_C_EQUAL( *A(j,j), c_zero )) {
*info = j+1; // Fortran index
return *info;
}
}
}
/* Determine the block size for this environment */
nb = magma_get_cpotrf_nb(n);
ldda = ((n+31)/32)*32;
if (MAGMA_SUCCESS != magma_cmalloc( &work, (n)*ldda )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
cudaStream_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
if (nb <= 1 || nb >= n)
lapackf77_ctrtri(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_csetmatrix( jb, (n-j),
A(j, j), lda,
dA(j, j), ldda );
/* Compute rows 1:j-1 of current block column */
magma_ctrmm( MagmaLeft, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_one, dA(0,0), ldda, dA(0, j),ldda);
magma_ctrsm( MagmaRight, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_neg_one, dA(j,j), ldda, dA(0, j),ldda);
//cublasGetMatrix(j ,jb, sizeof( cuFloatComplex),
//dA(0, j), ldda, A(0, j), lda);
magma_cgetmatrix_async( jb, jb,
dA(j, j), ldda,
A(j, j), lda, stream[1] );
magma_cgetmatrix_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_ctrtri(MagmaUpperStr, diag_, &jb, A(j,j), &lda, info);
magma_csetmatrix( 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_csetmatrix( (n-j), jb,
A(j, j), lda,
dA(j, j), ldda );
/* Compute rows j+jb:n of current block column */
magma_ctrmm( MagmaLeft, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_one, dA(j+jb,j+jb), ldda, dA(j+jb, j), ldda );
magma_ctrsm( MagmaRight, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_neg_one, dA(j,j), ldda, dA(j+jb, j), ldda );
//cublasGetMatrix((n-j), jb, sizeof( cuFloatComplex),dA(j, j), ldda, A(j, j), lda);
magma_cgetmatrix_async( n-j-jb, jb,
dA(j+jb, j), ldda,
A(j+jb, j), lda, stream[1] );
magma_cgetmatrix_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_ctrtri(MagmaLowerStr, diag_, &jb, A(j,j), &lda, info);
magma_csetmatrix( jb, jb,
A(j, j), lda,
dA(j, j), ldda );
}
}
}
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free( work );
return *info;
}
extern "C" magma_int_t
magma_chegvr(magma_int_t itype, char jobz, char range, char uplo, magma_int_t n,
magmaFloatComplex *a, magma_int_t lda, magmaFloatComplex *b, magma_int_t ldb,
float vl, float vu, magma_int_t il, magma_int_t iu, float abstol,
magma_int_t *m, float *w, magmaFloatComplex *z, magma_int_t ldz,
magma_int_t *isuppz, magmaFloatComplex *work, magma_int_t lwork,
float *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
=======
CHEGVR 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.
Whenever possible, CHEEVR calls CSTEGR to compute the
eigenspectrum using Relatively Robust Representations. CSTEGR
computes eigenvalues by the dqds algorithm, while orthogonal
eigenvectors are computed from various "good" L D L^T representations
(also known as Relatively Robust Representations). Gram-Schmidt
orthogonalization is avoided as far as possible. More specifically,
the various steps of the algorithm are as follows. For the i-th
unreduced block of T,
(a) Compute T - sigma_i = L_i D_i L_i^T, such that L_i D_i L_i^T
is a relatively robust representation,
(b) Compute the eigenvalues, lambda_j, of L_i D_i L_i^T to high
relative accuracy by the dqds algorithm,
(c) If there is a cluster of close eigenvalues, "choose" sigma_i
close to the cluster, and go to step (a),
(d) Given the approximate eigenvalue lambda_j of L_i D_i L_i^T,
compute the corresponding eigenvector by forming a
rank-revealing twisted factorization.
The desired accuracy of the output can be specified by the input
parameter ABSTOL.
For more details, see "A new O(n^2) algorithm for the symmetric
tridiagonal eigenvalue/eigenvector problem", by Inderjit Dhillon,
Computer Science Division Technical Report No. UCB//CSD-97-971,
UC Berkeley, May 1997.
Note 1 : CHEEVR calls CSTEGR when the full spectrum is requested
on machines which conform to the ieee-754 floating point standard.
CHEEVR calls SSTEBZ and CSTEIN on non-ieee machines and
when partial spectrum requests are made.
Normal execution of CSTEGR may create NaNs and infinities and
hence may abort due to a floating point exception in environments
which do not handle NaNs and infinities in the ieee standard default
manner.
Arguments
=========
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 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 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) REAL
VU (input) REAL
If RANGE='V', the lower and upper bounds of the interval to
be searched for eigenvalues. VL < VU.
Not referenced if RANGE = 'A' or 'I'.
IL (input) INTEGER
IU (input) INTEGER
If RANGE='I', the indices (in ascending order) of the
smallest and largest eigenvalues to be returned.
1 <= IL <= IU <= N, if N > 0; IL = 1 and IU = 0 if N = 0.
Not referenced if RANGE = 'A' or 'V'.
ABSTOL (input) REAL
The absolute error tolerance for the eigenvalues.
An approximate eigenvalue is accepted as converged
when it is determined to lie in an interval [a,b]
of width less than or equal to
ABSTOL + EPS * max( |a|,|b| ) ,
where EPS is the machine precision. If ABSTOL is less than
or equal to zero, then EPS*|T| will be used in its place,
where |T| is the 1-norm of the tridiagonal matrix obtained
by reducing A to tridiagonal form.
See "Computing Small Singular Values of Bidiagonal Matrices
with Guaranteed High Relative Accuracy," by Demmel and
Kahan, LAPACK Working Note #3.
If high relative accuracy is important, set ABSTOL to
SLAMCH( 'Safe minimum' ). Doing so will guarantee that
eigenvalues are computed to high relative accuracy when
possible in future releases. The current code does not
make any guarantees about high relative accuracy, but
furutre releases will. See J. Barlow and J. Demmel,
"Computing Accurate Eigensystems of Scaled Diagonally
Dominant Matrices", LAPACK Working Note #7, for a discussion
of which matrices define their eigenvalues to high relative
accuracy.
M (output) INTEGER
The total number of eigenvalues found. 0 <= M <= N.
If RANGE = 'A', M = N, and if RANGE = 'I', M = IU-IL+1.
W (output) REAL array, dimension (N)
If INFO = 0, the eigenvalues in ascending order.
Z (output) COMPLEX array, dimension (LDZ, max(1,M))
If JOBZ = 'V', then if INFO = 0, the first M columns of Z
contain the orthonormal eigenvectors of the matrix A
corresponding to the selected eigenvalues, with the i-th
column of Z holding the eigenvector associated with W(i).
If JOBZ = 'N', then Z is not referenced.
Note: the user must ensure that at least max(1,M) columns are
supplied in the array Z; if RANGE = 'V', the exact value of M
is not known in advance and an upper bound must be used.
LDZ (input) INTEGER
The leading dimension of the array Z. LDZ >= 1, and if
JOBZ = 'V', LDZ >= max(1,N).
ISUPPZ (output) INTEGER ARRAY, dimension ( 2*max(1,M) )
The support of the eigenvectors in Z, i.e., the indices
indicating the nonzero elements in Z. The i-th eigenvector
is nonzero only in elements ISUPPZ( 2*i-1 ) through
ISUPPZ( 2*i ).
********* Implemented only for RANGE = 'A' or 'I' and IU - IL = N - 1
WORK (workspace/output) COMPLEX array, dimension (LWORK)
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The length of the array WORK. LWORK >= max(1,2*N).
For optimal efficiency, LWORK >= (NB+1)*N,
where NB is the max of the blocksize for CHETRD and for
CUNMTR as returned by ILAENV.
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
RWORK (workspace/output) REAL array, dimension (LRWORK)
On exit, if INFO = 0, RWORK(1) returns the optimal
(and minimal) LRWORK.
LRWORK (input) INTEGER
The length of the array RWORK. LRWORK >= max(1,24*N).
If LRWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the RWORK array, returns
this value as the first entry of the RWORK array, and no error
message related to LRWORK is issued by XERBLA.
IWORK (workspace/output) INTEGER array, dimension (LIWORK)
On exit, if INFO = 0, IWORK(1) returns the optimal
(and minimal) LIWORK.
LIWORK (input) INTEGER
The dimension of the array IWORK. LIWORK >= max(1,10*N).
If LIWORK = -1, then a workspace query is assumed; the
routine only calculates the optimal size of the IWORK array,
returns this value as the first entry of the IWORK array, and
no error message related to LIWORK is issued by XERBLA.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
> 0: Internal error
Further Details
===============
Based on contributions by
Inderjit Dhillon, IBM Almaden, USA
Osni Marques, LBNL/NERSC, USA
Ken Stanley, Computer Science Division, University of
California at Berkeley, USA
===================================================================== */
char uplo_[2] = {uplo, 0};
char jobz_[2] = {jobz, 0};
char range_[2] = {range, 0};
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex *da;
magmaFloatComplex *db;
magmaFloatComplex *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, lrwmin, liwmin;
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_chetrd_nb(n);
lwmin = n * (nb + 1);
lrwmin = 24 * n;
liwmin = 10 * n;
work[0] = MAGMA_C_MAKE( lwmin, 0 );
rwork[0] = lrwmin;
iwork[0] = liwmin;
if (lwork < lwmin && ! lquery) {
*info = -21;
} else if ((lrwork < lrwmin) && ! lquery) {
*info = -23;
} else if ((liwork < liwmin) && ! lquery) {
*info = -25;
}
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_cmalloc( &da, n*ldda ) ||
MAGMA_SUCCESS != magma_cmalloc( &db, n*lddb ) ||
MAGMA_SUCCESS != magma_cmalloc( &dz, n*lddz )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/* Form a Cholesky factorization of B. */
magma_csetmatrix( n, n, b, ldb, db, lddb );
magma_csetmatrix_async( n, n,
a, lda,
da, ldda, stream );
magma_cpotrf_gpu(uplo_[0], n, db, lddb, info);
if (*info != 0) {
*info = n + *info;
return *info;
}
magma_queue_sync( stream );
magma_cgetmatrix_async( n, n,
db, lddb,
b, ldb, stream );
/* Transform problem to standard eigenvalue problem and solve. */
magma_chegst_gpu(itype, uplo, n, da, ldda, db, lddb, info);
magma_cheevr_gpu(jobz, range, uplo, n, da, ldda, vl, vu, il, iu, abstol,
m, w, dz, lddz, isuppz, a, lda, z, ldz, work, lwork,
rwork, lrwork, iwork, liwork, 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_ctrsm(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_ctrmm(MagmaLeft, uplo, *trans, MagmaNonUnit, n, *m, c_one,
db, lddb, dz, lddz);
}
magma_cgetmatrix( 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;
} /* chegvr */
extern "C" magma_int_t
magma_ctstrf_gpu( char storev, magma_int_t m, magma_int_t n, magma_int_t ib, magma_int_t nb,
magmaFloatComplex *hU, magma_int_t ldhu, magmaFloatComplex *dU, magma_int_t lddu,
magmaFloatComplex *hA, magma_int_t ldha, magmaFloatComplex *dA, magma_int_t ldda,
magmaFloatComplex *hL, magma_int_t ldhl, magmaFloatComplex *dL, magma_int_t lddl,
magma_int_t *ipiv,
magmaFloatComplex *hwork, magma_int_t ldhwork, magmaFloatComplex *dwork, magma_int_t lddwork,
magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CSSSSM 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
=========
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. N >= 0.
IB (input) INTEGER
The inner-blocking size. IB >= 0.
NB (input) INTEGER
The blocking size. NB >= 0.
hU (input,output) COMPLEX 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.
LDHU (input) INTEGER
The leading dimension of the array hU. LDHU >= max(1,NB).
dU (input,output) COMPLEX 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.
LDDU (input) INTEGER
The leading dimension of the array dU. LDDU >= max(1,NB).
hA (input,output) COMPLEX array, dimension(LDHA, N), on cpu.
On entry, only the M-by-IB first panel needs to be identical to dA(1..M, 1..IB).
On exit, the content is incomplete. Shouldn't be used.
LDHA (input) INTEGER
The leading dimension of the array hA. LDHA >= max(1,M).
dA (input,output) COMPLEX array, dimension(LDDA, N) , on gpu.
On entry, the M-by-N tile to be factored.
On exit, the factor L from the factorization
LDDA (input) INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
hL (output) COMPLEX 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.
LDHL (input) INTEGER
The leading dimension of the array hL. LDHL >= max(1,2*IB).
dL (output) COMPLEX 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.
LDDL (input) INTEGER
The leading dimension of the array dL. LDDL >= max(1,2*IB).
hWORK (output) COMPLEX array, dimension(LDHWORK, 2*IB), on cpu.
Workspace.
LDHWORK (input) INTEGER
The leading dimension of the array hWORK. LDHWORK >= max(NB, 1).
dWORK (output) COMPLEX array, dimension(LDDWORK, 2*IB), on gpu.
Workspace.
LDDWORK (input) INTEGER
The leading dimension of the array dWORK. LDDWORK >= max(NB, 1).
IPIV (output) INTEGER array on the cpu.
The pivot indices array of size K as returned by CTSTRF
INFO (output) INTEGER
- PLASMA_SUCCESS successful exit
- < 0 if INFO = -k, the k-th argument had an illegal value
- > 0 if INFO = k, U(k,k) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
===================================================================== */
#define 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 )
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
int iinfo = 0;
int maxm, mindim;
int i, j, im, s, ip, ii, sb, p = 1;
magmaFloatComplex *dAT, *dUT;
magmaFloatComplex *dAp, *dUp;
#ifndef WITHOUTTRTRI
magmaFloatComplex *dL2 = dL + ib;
magmaFloatComplex *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_ctstrf(m, n, ib, nb,
(PLASMA_Complex32_t*)hU, ldhu,
(PLASMA_Complex32_t*)hA, ldha,
(PLASMA_Complex32_t*)hL, ldhl,
ipiv,
(PLASMA_Complex32_t*)hwork, ldhwork,
info);
#ifndef WITHOUTTRTRI
CORE_clacpy( PlasmaUpperLower, mindim, mindim,
(PLASMA_Complex32_t*)hL, ldhl,
(PLASMA_Complex32_t*)hL2, ldhl );
CORE_ctrtri( PlasmaLower, PlasmaUnit, mindim,
(PLASMA_Complex32_t*)hL2, ldhl, info );
if (*info != 0 ) {
fprintf(stderr, "ERROR, trtri returned with info = %d\n", *info);
}
#endif
if ( (storev == 'R') || (storev == 'r') ) {
magma_csetmatrix( m, n, hU, ldhu, dwork, lddwork );
magmablas_ctranspose( dU, lddu, dwork, lddwork, m, n );
magma_csetmatrix( m, n, hA, ldha, dwork, lddwork );
magmablas_ctranspose( dA, ldda, dwork, lddwork, m, n );
} else {
magma_csetmatrix( m, n, hU, ldhu, dU, lddu );
magma_csetmatrix( m, n, hA, ldha, dA, ldda );
}
magma_csetmatrix( p*ib, n, hL, ldhl, dL, lddl );
}
else {
/* Use hybrid blocked code. */
maxm = ((m + 31)/32)*32;
if ( (storev == 'C') || (storev == 'c') ) {
magmablas_cgetmo_in( dU, dUT, lddu, m, n );
magmablas_cgetmo_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_ctranspose( dUp, lddu, UT(0, i), lddu, sb, ii );
magmablas_ctranspose( dAp, ldda, AT(0, i), ldda, sb, m );
magma_cgetmatrix( ii, sb, dUp, lddu, hU(0, i), ldhu );
magma_cgetmatrix( m, sb, dAp, ldda, hA(0, i), ldha );
// make sure that gpu queue is empty
//magma_device_sync();
#ifndef WITHOUTTRTRI
magma_ctrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n-(ii+sb), ib,
c_one, L2(i-1), lddl,
UT(i-1, i+1), lddu);
#else
magma_ctrsm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n-(ii+sb), ib,
c_one, L(i-1), lddl,
UT(i-1, i+1), lddu);
#endif
magma_cgemm( 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_ctstrf(m, sb, ib, nb,
(PLASMA_Complex32_t*)hU(i, i), ldhu,
(PLASMA_Complex32_t*)hA(0, i), ldha,
(PLASMA_Complex32_t*)hL(i), ldhl,
ipiv+ii,
(PLASMA_Complex32_t*)hwork, ldhwork,
info);
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + ii;
// Need to swap betw U and A
#ifndef NOSWAPBLK
magmablas_cswapblk( 'R', 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_cswap( 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_clacpy( PlasmaUpperLower, sb, sb,
(PLASMA_Complex32_t*)hL(i), ldhl,
(PLASMA_Complex32_t*)hL2(i), ldhl );
CORE_ctrtri( PlasmaLower, PlasmaUnit, sb,
(PLASMA_Complex32_t*)hL2(i), ldhl, info );
if (*info != 0 ) {
fprintf(stderr, "ERROR, trtri returned with info = %d\n", *info);
}
#endif
// upload i-th panel
magma_csetmatrix( sb, sb, hU(i, i), ldhu, dUp, lddu );
magma_csetmatrix( m, sb, hA(0, i), ldha, dAp, ldda );
magma_csetmatrix( p*ib, sb, hL(i), ldhl, L(i), lddl );
magmablas_ctranspose( UT(i, i), lddu, dUp, lddu, sb, sb);
magmablas_ctranspose( AT(0, i), ldda, dAp, ldda, m, sb);
// make sure that gpu queue is empty
//magma_device_sync();
// do the small non-parallel computations
if ( s > (i+1) ) {
#ifndef WITHOUTTRTRI
magma_ctrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
sb, sb,
c_one, L2(i), lddl,
UT(i, i+1), lddu);
#else
magma_ctrsm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
sb, sb,
c_one, L(i), lddl,
UT(i, i+1), lddu);
#endif
magma_cgemm( 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_ctrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n-mindim, sb,
c_one, L2(i), lddl,
UT(i, i+1), lddu);
#else
magma_ctrsm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n-mindim, sb,
c_one, L(i), lddl,
UT(i, i+1), lddu);
#endif
magma_cgemm( 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 ( (storev == 'C') || (storev == 'c') ) {
magmablas_cgetmo_out( dU, dUT, lddu, m, n );
magmablas_cgetmo_out( dA, dAT, ldda, m, n );
}
}
return *info;
}
/**
Purpose
-------
CLAUUM computes the product U * U' or L' * L, where the triangular
factor U or L is stored in the upper or lower triangular part of
the array dA.
If UPLO = MagmaUpper then the upper triangle of the result is stored,
overwriting the factor U in dA.
If UPLO = MagmaLower then the lower triangle of the result is stored,
overwriting the factor L in dA.
This is the blocked form of the algorithm, calling Level 3 BLAS.
Arguments
---------
@param[in]
uplo magma_uplo_t
Specifies whether the triangular factor stored in the array dA
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]
dA REAL array on the GPU, dimension (LDDA,N)
On entry, the triangular factor U or L.
On exit, if UPLO = MagmaUpper, the upper triangle of dA is
overwritten with the upper triangle of the product U * U';
if UPLO = MagmaLower, the lower triangle of dA is overwritten with
the lower triangle of the product L' * L.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,N).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -k, the k-th argument had an illegal value
@ingroup magma_cposv_aux
***************************************************************************/
extern "C" magma_int_t
magma_clauum_gpu(magma_uplo_t uplo, magma_int_t n,
magmaFloatComplex *dA, magma_int_t ldda, magma_int_t *info)
{
#define dA(i, j) (dA + (j)*ldda + (i))
/* Local variables */
const char* uplo_ = lapack_uplo_const( uplo );
magma_int_t nb, i, ib;
float d_one = MAGMA_D_ONE;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex *work;
int upper = (uplo == MagmaUpper);
*info = 0;
if (! upper && uplo != MagmaLower)
*info = -1;
else if (n < 0)
*info = -2;
else if (ldda < max(1,n))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
nb = magma_get_cpotrf_nb(n);
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &work, nb*nb )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_queue_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
if (nb <= 1 || nb >= n) {
magma_cgetmatrix( n, n, dA, ldda, work, n );
lapackf77_clauum(uplo_, &n, work, &n, info);
magma_csetmatrix( n, n, work, n, dA, ldda );
}
else {
if (upper) {
/* Compute inverse of upper triangular matrix */
for (i=0; i < n; i += nb) {
ib = min(nb, (n-i));
/* Compute the product U * U'. */
magma_ctrmm( MagmaRight, MagmaUpper,
MagmaConjTrans, MagmaNonUnit, i, ib,
c_one, dA(i,i), ldda, dA(0, i),ldda);
magma_cgetmatrix( ib, ib,
dA(i, i), ldda,
work, ib );
lapackf77_clauum(MagmaUpperStr, &ib, work, &ib, info);
magma_csetmatrix( ib, ib,
work, ib,
dA(i, i), ldda );
if (i+ib < n) {
magma_cgemm( 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_cherk( MagmaUpper, MagmaNoTrans, ib,(n-i-ib),
d_one, dA(i, i+ib), ldda,
d_one, dA(i, i), ldda);
}
}
}
else {
/* Compute the product L' * L. */
for (i=0; i < n; i += nb) {
ib=min(nb,(n-i));
magma_ctrmm( MagmaLeft, MagmaLower,
MagmaConjTrans, MagmaNonUnit, ib,
i, c_one, dA(i,i), ldda,
dA(i, 0),ldda);
magma_cgetmatrix( ib, ib,
dA(i, i), ldda,
work, ib );
lapackf77_clauum(MagmaLowerStr, &ib, work, &ib, info);
magma_csetmatrix( ib, ib,
work, ib,
dA(i, i), ldda );
if (i+ib < n) {
magma_cgemm( 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_cherk( MagmaLower, MagmaConjTrans, ib, (n-i-ib),
d_one, dA(i+ib, i), ldda,
d_one, dA(i, i), ldda);
}
}
}
}
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free_pinned( work );
return *info;
}
/**
Purpose
-------
CGESSM applies the factors L computed by CGETRF_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
CGETRF_INCPIV.
@param[in]
dL1 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 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 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_cgesv_tile
********************************************************************/
extern "C" magma_int_t
magma_cgessm_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,
magmaFloatComplex_ptr dL1, magma_int_t lddl1,
magmaFloatComplex_ptr dL, magma_int_t lddl,
magmaFloatComplex_ptr 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)
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
int i, sb;
magmaFloatComplex_ptr 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_cgetmo_in( dA, dAT, ldda, m, n );
} else {
dAT = dA;
}
for (i = 0; i < k; i += ib) {
sb = min(ib, k-i);
magmablas_claswp( n, dAT, ldda, i+1, i+sb, ipiv, 1 );
#ifndef WITHOUTTRTRI
magma_ctrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n, sb,
c_one, dL1(i), lddl1,
AT(i, 0), ldda);
#else
magma_ctrsm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n, sb,
c_one, L( i, i), lddl,
AT(i, 0), ldda);
#endif
if ( (i+sb) < m) {
magma_cgemm( 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_cgetmo_in( dA, dAT, ldda, m, n );
}
return *info;
} /* magma_cgessm_gpu */
/**
Purpose
-------
CHEGVD 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]
ngpu INTEGER
Number of GPUs to use. ngpu > 0.
@param[in]
itype INTEGER
Specifies the problem type to be solved:
= 1: A*x = (lambda)*B*x
= 2: A*B*x = (lambda)*x
= 3: B*A*x = (lambda)*x
@param[in]
jobz magma_vec_t
- = MagmaNoVec: Compute eigenvalues only;
- = MagmaVec: Compute eigenvalues and eigenvectors.
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangles of A and B are stored;
- = MagmaLower: Lower triangles of A and B are stored.
@param[in]
n INTEGER
The order of the matrices A and B. N >= 0.
@param[in,out]
A COMPLEX 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 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 REAL array, dimension (N)
If INFO = 0, the eigenvalues in ascending order.
@param[out]
work (workspace) COMPLEX 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_chetrd_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) REAL 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: CPOTRF or CHEEVD 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 CHEEVD 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_chegv_driver
********************************************************************/
extern "C" magma_int_t
magma_chegvd_m(
magma_int_t ngpu,
magma_int_t itype, magma_vec_t jobz, magma_uplo_t uplo, magma_int_t n,
magmaFloatComplex *A, magma_int_t lda,
magmaFloatComplex *B, magma_int_t ldb,
float *w, magmaFloatComplex *work, magma_int_t lwork,
#ifdef COMPLEX
float *rwork, magma_int_t lrwork,
#endif
magma_int_t *iwork, magma_int_t liwork,
magma_int_t *info)
{
const char* uplo_ = lapack_uplo_const( uplo );
const char* jobz_ = lapack_vec_const( jobz );
magmaFloatComplex c_one = MAGMA_C_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;
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_chetrd_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;
}
work[0] = magma_cmake_lwork( lwmin );
rwork[0] = magma_smake_lwork( lrwmin );
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;
}
/* If matrix is very small, then just call LAPACK on CPU, no need for GPU */
if (n <= 128) {
lapackf77_chegvd( &itype, jobz_, uplo_,
&n, A, &lda, B, &ldb,
w, work, &lwork,
#ifdef COMPLEX
rwork, &lrwork,
#endif
iwork, &liwork, info);
return *info;
}
magma_timer_t time=0;
timer_start( time );
magma_cpotrf_m( ngpu, uplo, n, B, ldb, info );
if (*info != 0) {
*info = n + *info;
return *info;
}
timer_stop( time );
timer_printf( "time cpotrf = %6.2f\n", time );
timer_start( time );
/* Transform problem to standard eigenvalue problem and solve. */
magma_chegst_m( ngpu, itype, uplo, n, A, lda, B, ldb, info );
timer_stop( time );
timer_printf( "time chegst = %6.2f\n", time );
timer_start( time );
magma_cheevd_m( ngpu, jobz, uplo, n, A, lda, w, work, lwork, rwork, lrwork, iwork, liwork, info );
timer_stop( time );
timer_printf( "time cheevd = %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_ctrsm_m( ngpu, 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;
}
#ifdef ENABLE_DEBUG
printf("--- the multi GPU version is falling back to 1 GPU to perform the last TRMM since there is no TRMM_mgpu --- \n");
#endif
magmaFloatComplex *dA=NULL, *dB=NULL;
magma_int_t ldda = magma_roundup( n, 32 );
magma_int_t lddb = ldda;
if (MAGMA_SUCCESS != magma_cmalloc( &dA, n*ldda ) ||
MAGMA_SUCCESS != magma_cmalloc( &dB, n*lddb )) {
magma_free( dA );
magma_free( dB );
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_queue_t queue;
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queue );
magma_csetmatrix( n, n, B, ldb, dB, lddb, queue );
magma_csetmatrix( n, n, A, lda, dA, ldda, queue );
magma_ctrmm( MagmaLeft, uplo, trans, MagmaNonUnit,
n, n, c_one, dB, lddb, dA, ldda, queue );
magma_cgetmatrix( n, n, dA, ldda, A, lda, queue );
magma_queue_destroy( queue );
magma_free( dA );
magma_free( dB );
}
timer_stop( time );
timer_printf( "time setmatrices trsm/mm + getmatrices = %6.2f\n", time );
}
work[0] = magma_cmake_lwork( lwmin );
rwork[0] = magma_smake_lwork( lrwmin );
iwork[0] = liwmin;
return *info;
} /* magma_chegvd_m */