/**
Purpose
-------
DORMQR overwrites the general real M-by-N matrix C with
@verbatim
SIDE = MagmaLeft SIDE = MagmaRight
TRANS = MagmaNoTrans: Q * C C * Q
TRANS = MagmaTrans: Q**H * C C * Q**H
@endverbatim
where Q is a real unitary matrix defined as the product of k
elementary reflectors
Q = H(1) H(2) . . . H(k)
as returned by DGEQRF. Q is of order M if SIDE = MagmaLeft and of order N
if SIDE = MagmaRight.
Arguments
---------
@param[in]
ngpu INTEGER
Number of GPUs to use. ngpu > 0.
@param[in]
side magma_side_t
- = MagmaLeft: apply Q or Q**H from the Left;
- = MagmaRight: apply Q or Q**H from the Right.
@param[in]
trans magma_trans_t
- = MagmaNoTrans: No transpose, apply Q;
- = MagmaTrans: Conjugate transpose, apply Q**H.
@param[in]
m INTEGER
The number of rows of the matrix C. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix C. N >= 0.
@param[in]
k INTEGER
The number of elementary reflectors whose product defines
the matrix Q.
If SIDE = MagmaLeft, M >= K >= 0;
if SIDE = MagmaRight, N >= K >= 0.
@param[in]
A DOUBLE_PRECISION array, dimension (LDA,K)
The i-th column must contain the vector which defines the
elementary reflector H(i), for i = 1,2,...,k, as returned by
DGEQRF in the first k columns of its array argument A.
@param[in]
lda INTEGER
The leading dimension of the array A.
If SIDE = MagmaLeft, LDA >= max(1,M);
if SIDE = MagmaRight, LDA >= max(1,N).
@param[in]
tau DOUBLE_PRECISION array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by DGEQRF.
@param[in,out]
C DOUBLE_PRECISION array, dimension (LDC,N)
On entry, the M-by-N matrix C.
On exit, C is overwritten by Q*C or Q**H*C or C*Q**H or C*Q.
@param[in]
ldc INTEGER
The leading dimension of the array C. LDC >= max(1,M).
@param[out]
work (workspace) DOUBLE_PRECISION array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK.
If SIDE = MagmaLeft, LWORK >= max(1,N);
if SIDE = MagmaRight, LWORK >= max(1,M).
For optimum performance LWORK >= N*NB if SIDE = MagmaLeft, and
LWORK >= M*NB if SIDE = MagmaRight, where NB is the optimal
blocksize.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_dgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_dormqr_m(
magma_int_t ngpu,
magma_side_t side, magma_trans_t trans,
magma_int_t m, magma_int_t n, magma_int_t k,
double *A, magma_int_t lda,
double *tau,
double *C, magma_int_t ldc,
double *work, magma_int_t lwork,
magma_int_t *info)
{
#define A(i, j) (A + (j)*lda + (i))
#define C(i, j) (C + (j)*ldc + (i))
#define dC(gpui, i, j) (dw[gpui] + (j)*lddc + (i))
#define dA_c(gpui, ind, i, j) (dw[gpui] + maxnlocal*lddc + (ind)*lddar*lddac + (i) + (j)*lddac)
#define dA_r(gpui, ind, i, j) (dw[gpui] + maxnlocal*lddc + (ind)*lddar*lddac + (i) + (j)*lddar)
#define dT(gpui, ind) (dw[gpui] + maxnlocal*lddc + 2*lddac*lddar + (ind)*((nb+1)*nb))
#define dwork(gpui, ind) (dw[gpui] + maxnlocal*lddc + 2*lddac*lddar + 2*((nb+1)*nb) + (ind)*(lddwork*nb))
double c_zero = MAGMA_D_ZERO;
double c_one = MAGMA_D_ONE;
const char* side_ = lapack_side_const( side );
const char* trans_ = lapack_trans_const( trans );
// TODO fix memory leak (alloc after argument checks)
magma_int_t nb = 128;
double *T;
magma_dmalloc_pinned(&T, nb*nb);
//printf("calling dormqr_m with nb=%d\n", (int) nb);
double* dw[MagmaMaxGPUs];
magma_queue_t stream [MagmaMaxGPUs][2];
magma_event_t event [MagmaMaxGPUs][2];
magma_int_t ind_c;
magma_device_t igpu;
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
*info = 0;
magma_int_t left = (side == MagmaLeft);
magma_int_t notran = (trans == MagmaNoTrans);
magma_int_t lquery = (lwork == -1);
/* NQ is the order of Q and NW is the minimum dimension of WORK */
magma_int_t nq, nw;
if (left) {
nq = m;
nw = n;
} else {
nq = n;
nw = m;
}
if (! left && side != MagmaRight) {
*info = -1;
} else if (! notran && trans != MagmaTrans) {
*info = -2;
} else if (m < 0) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (k < 0 || k > nq) {
*info = -5;
} else if (lda < max(1,nq)) {
*info = -7;
} else if (ldc < max(1,m)) {
*info = -10;
} else if (lwork < max(1,nw) && ! lquery) {
*info = -12;
}
magma_int_t lwkopt = max(1,nw) * nb;
if (*info == 0) {
work[0] = MAGMA_D_MAKE( lwkopt, 0 );
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0 || k == 0) {
work[0] = c_one;
return *info;
}
if (nb >= k) {
/* Use CPU code */
lapackf77_dormqr(side_, trans_, &m, &n, &k, A, &lda, tau,
C, &ldc, work, &lwork, info);
return *info;
}
magma_int_t lddc = (m+63)/64*64;
magma_int_t lddac = nq;
magma_int_t lddar = nb;
magma_int_t lddwork = nw;
magma_int_t nlocal[ MagmaMaxGPUs ] = { 0 };
magma_int_t nb_l=256;
magma_int_t nbl = (n-1)/nb_l+1; // number of blocks
magma_int_t maxnlocal = (nbl+ngpu-1)/ngpu*nb_l;
ngpu = min(ngpu, (n+nb_l-1)/nb_l); // Don't use GPU that will not have data.
magma_int_t ldw = maxnlocal*lddc // dC
+ 2*lddac*lddar // 2*dA
+ 2*(nb + 1 + lddwork)*nb; // 2*(dT and dwork)
for (igpu = 0; igpu < ngpu; ++igpu) {
magma_setdevice(igpu);
if (MAGMA_SUCCESS != magma_dmalloc( &dw[igpu], ldw )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
magma_xerbla( __func__, -(*info) );
return *info;
}
magma_queue_create( &stream[igpu][0] );
magma_queue_create( &stream[igpu][1] );
magma_event_create( &event[igpu][0] );
magma_event_create( &event[igpu][1] );
}
/* Use hybrid CPU-MGPU code */
if (left) {
//copy C to mgpus
for (magma_int_t i = 0; i < nbl; ++i) {
magma_int_t igpu = i%ngpu;
magma_setdevice(igpu);
magma_int_t kb = min(nb_l, n-i*nb_l);
magma_dsetmatrix_async( m, kb,
C(0, i*nb_l), ldc,
dC(igpu, 0, i/ngpu*nb_l), lddc, stream[igpu][0] );
nlocal[igpu] += kb;
}
magma_int_t i1, i2, i3;
if ( !notran ) {
i1 = 0;
i2 = k;
i3 = nb;
} else {
i1 = (k - 1) / nb * nb;
i2 = 0;
i3 = -nb;
}
ind_c = 0;
for (magma_int_t i = i1; (i3 < 0 ? i >= i2 : i < i2); i += i3) {
// start the copy of A panel
magma_int_t kb = min(nb, k - i);
for (igpu = 0; igpu < ngpu; ++igpu) {
magma_setdevice(igpu);
magma_event_sync(event[igpu][ind_c]); // check if the new data can be copied
magma_dsetmatrix_async(nq-i, kb,
A(i, i), lda,
dA_c(igpu, ind_c, i, 0), lddac, stream[igpu][0] );
// set upper triangular part of dA to identity
magmablas_dlaset_band_q( MagmaUpper, kb, kb, kb, c_zero, c_one, dA_c(igpu, ind_c, i, 0), lddac, stream[igpu][0] );
}
/* Form the triangular factor of the block reflector
H = H(i) H(i+1) . . . H(i+ib-1) */
magma_int_t nqi = nq - i;
lapackf77_dlarft("F", "C", &nqi, &kb, A(i, i), &lda,
&tau[i], T, &kb);
/* H or H' is applied to C(1:m,i:n) */
/* Apply H or H'; First copy T to the GPU */
for (igpu = 0; igpu < ngpu; ++igpu) {
magma_setdevice(igpu);
magma_dsetmatrix_async(kb, kb,
T, kb,
dT(igpu, ind_c), kb, stream[igpu][0] );
}
for (igpu = 0; igpu < ngpu; ++igpu) {
magma_setdevice(igpu);
magma_queue_sync( stream[igpu][0] ); // check if the data was copied
magmablasSetKernelStream(stream[igpu][1]);
magma_dlarfb_gpu( side, trans, MagmaForward, MagmaColumnwise,
m-i, nlocal[igpu], kb,
dA_c(igpu, ind_c, i, 0), lddac, dT(igpu, ind_c), kb,
dC(igpu, i, 0), lddc,
dwork(igpu, ind_c), lddwork);
magma_event_record(event[igpu][ind_c], stream[igpu][1] );
}
ind_c = (ind_c+1)%2;
}
for (igpu = 0; igpu < ngpu; ++igpu) {
magma_setdevice(igpu);
magma_queue_sync( stream[igpu][1] );
}
//copy C from mgpus
for (magma_int_t i = 0; i < nbl; ++i) {
magma_int_t igpu = i%ngpu;
magma_setdevice(igpu);
magma_int_t kb = min(nb_l, n-i*nb_l);
magma_dgetmatrix( m, kb,
dC(igpu, 0, i/ngpu*nb_l), lddc,
C(0, i*nb_l), ldc );
// magma_dgetmatrix_async( m, kb,
// dC(igpu, 0, i/ngpu*nb_l), lddc,
// C(0, i*nb_l), ldc, stream[igpu][0] );
}
} else {
// TODO fix memory leak T, dw, event, stream
fprintf(stderr, "The case (side == right) is not implemented\n");
*info = MAGMA_ERR_NOT_IMPLEMENTED;
magma_xerbla( __func__, -(*info) );
return *info;
/*
if ( notran ) {
i1 = 0;
i2 = k;
i3 = nb;
} else {
i1 = (k - 1) / nb * nb;
i2 = 0;
i3 = -nb;
}
mi = m;
ic = 0;
for (i = i1; (i3 < 0 ? i >= i2 : i < i2); i += i3) {
ib = min(nb, k - i);
// Form the triangular factor of the block reflector
// H = H(i) H(i+1) . . . H(i+ib-1)
i__4 = nq - i;
lapackf77_dlarft("F", "C", &i__4, &ib, A(i, i), &lda,
&tau[i], T, &ib);
// 1) copy the panel from A to the GPU, and
// 2) set upper triangular part of dA to identity
magma_dsetmatrix( i__4, ib, A(i, i), lda, dA(i, 0), ldda );
magmablas_dlaset_band( MagmaUpper, ib, ib, ib, c_zero, c_one, dA(i, 0), ldda );
// H or H' is applied to C(1:m,i:n)
ni = n - i;
jc = i;
// Apply H or H'; First copy T to the GPU
magma_dsetmatrix( ib, ib, T, ib, dT, ib );
magma_dlarfb_gpu( side, trans, MagmaForward, MagmaColumnwise,
mi, ni, ib,
dA(i, 0), ldda, dT, ib,
dC(ic, jc), lddc,
dwork, lddwork);
}
*/
}
work[0] = MAGMA_D_MAKE( lwkopt, 0 );
for (igpu = 0; igpu < ngpu; ++igpu) {
magma_setdevice(igpu);
magma_event_destroy( event[igpu][0] );
magma_event_destroy( event[igpu][1] );
magma_queue_destroy( stream[igpu][0] );
magma_queue_destroy( stream[igpu][1] );
magma_free( dw[igpu] );
}
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
return *info;
} /* magma_dormqr */
/**
Purpose
-------
If VECT = MagmaQ, DORMBR overwrites the general real M-by-N matrix C with
SIDE = MagmaLeft SIDE = MagmaRight
TRANS = MagmaNoTrans: Q*C C*Q
TRANS = MagmaTrans: Q**H*C C*Q**H
If VECT = MagmaP, DORMBR overwrites the general real M-by-N matrix C with
SIDE = MagmaLeft SIDE = MagmaRight
TRANS = MagmaNoTrans: P*C C*P
TRANS = MagmaTrans: P**H*C C*P**H
Here Q and P**H are the unitary matrices determined by DGEBRD when
reducing A real matrix A to bidiagonal form: A = Q*B * P**H. Q
and P**H are defined as products of elementary reflectors H(i) and
G(i) respectively.
Let nq = m if SIDE = MagmaLeft and nq = n if SIDE = MagmaRight. Thus nq is the
order of the unitary matrix Q or P**H that is applied.
If VECT = MagmaQ, A is assumed to have been an NQ-by-K matrix:
if nq >= k, Q = H(1) H(2) . . . H(k);
if nq < k, Q = H(1) H(2) . . . H(nq-1).
If VECT = MagmaP, A is assumed to have been A K-by-NQ matrix:
if k < nq, P = G(1) G(2) . . . G(k);
if k >= nq, P = G(1) G(2) . . . G(nq-1).
Arguments
---------
@param[in]
vect magma_vect_t
- = MagmaQ: apply Q or Q**H;
- = MagmaP: apply P or P**H.
@param[in]
side magma_side_t
- = MagmaLeft: apply Q, Q**H, P or P**H from the Left;
- = MagmaRight: apply Q, Q**H, P or P**H from the Right.
@param[in]
trans magma_trans_t
- = MagmaNoTrans: No transpose, apply Q or P;
- = MagmaTrans: Conjugate transpose, apply Q**H or P**H.
@param[in]
m INTEGER
The number of rows of the matrix C. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix C. N >= 0.
@param[in]
k INTEGER
If VECT = MagmaQ, the number of columns in the original
matrix reduced by DGEBRD.
If VECT = MagmaP, the number of rows in the original
matrix reduced by DGEBRD.
K >= 0.
@param[in]
A DOUBLE_PRECISION array, dimension
(LDA,min(nq,K)) if VECT = MagmaQ
(LDA,nq) if VECT = MagmaP
The vectors which define the elementary reflectors H(i) and
G(i), whose products determine the matrices Q and P, as
returned by DGEBRD.
@param[in]
lda INTEGER
The leading dimension of the array A.
If VECT = MagmaQ, LDA >= max(1,nq);
if VECT = MagmaP, LDA >= max(1,min(nq,K)).
@param[in]
tau DOUBLE_PRECISION array, dimension (min(nq,K))
TAU(i) must contain the scalar factor of the elementary
reflector H(i) or G(i) which determines Q or P, as returned
by DGEBRD in the array argument TAUQ or TAUP.
@param[in,out]
C DOUBLE_PRECISION array, dimension (LDC,N)
On entry, the M-by-N matrix C.
On exit, C is overwritten by Q*C or Q**H*C or C*Q**H or C*Q
or P*C or P**H*C or C*P or C*P**H.
@param[in]
ldc INTEGER
The leading dimension of the array C. LDC >= max(1,M).
@param[out]
work (workspace) DOUBLE_PRECISION array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK.
If SIDE = MagmaLeft, LWORK >= max(1,N);
if SIDE = MagmaRight, LWORK >= max(1,M);
if N = 0 or M = 0, LWORK >= 1.
For optimum performance
if SIDE = MagmaLeft, LWORK >= max(1,N*NB);
if SIDE = MagmaRight, LWORK >= max(1,M*NB),
where NB is the optimal blocksize. (NB = 0 if M = 0 or N = 0.)
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_dgesvd_comp
********************************************************************/
extern "C" magma_int_t
magma_dormbr(
magma_vect_t vect, magma_side_t side, magma_trans_t trans,
magma_int_t m, magma_int_t n, magma_int_t k,
double *A, magma_int_t lda,
double *tau,
double *C, magma_int_t ldc,
double *work, magma_int_t lwork,
magma_int_t *info)
{
#define A(i,j) (A + (i) + (j)*lda)
#define C(i,j) (C + (i) + (j)*ldc)
magma_int_t i1, i2, nb, mi, ni, nq, nq_1, nw, iinfo, lwkopt;
magma_int_t left, notran, applyq, lquery;
magma_trans_t transt;
MAGMA_UNUSED( nq_1 ); // used only in version 1
*info = 0;
applyq = (vect == MagmaQ);
left = (side == MagmaLeft);
notran = (trans == MagmaNoTrans);
lquery = (lwork == -1);
/* NQ is the order of Q or P and NW is the minimum dimension of WORK */
if (left) {
nq = m;
nw = n;
}
else {
nq = n;
nw = m;
}
if (m == 0 || n == 0) {
nw = 0;
}
/* check arguments */
if (! applyq && vect != MagmaP) {
*info = -1;
}
else if (! left && side != MagmaRight) {
*info = -2;
}
else if (! notran && trans != MagmaTrans) {
*info = -3;
}
else if (m < 0) {
*info = -4;
}
else if (n < 0) {
*info = -5;
}
else if (k < 0) {
*info = -6;
}
else if ( ( applyq && lda < max(1,nq) ) ||
( ! applyq && lda < max(1,min(nq,k)) ) ) {
*info = -8;
}
else if (ldc < max(1,m)) {
*info = -11;
}
else if (lwork < max(1,nw) && ! lquery) {
*info = -13;
}
if (*info == 0) {
if (nw > 0) {
// TODO have get_dormqr_nb and get_dormlq_nb routines? see original LAPACK dormbr.
// TODO make them dependent on m, n, and k?
nb = magma_get_dgebrd_nb( min( m, n ));
lwkopt = max(1, nw*nb);
}
else {
lwkopt = 1;
}
work[0] = MAGMA_D_MAKE( lwkopt, 0 );
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0) {
return *info;
}
if (applyq) {
/* Apply Q */
if (nq >= k) {
/* Q was determined by a call to DGEBRD with nq >= k */
#if VERSION == 1
lapackf77_dormqr( lapack_side_const(side), lapack_trans_const(trans),
&m, &n, &k, A, &lda, tau, C, &ldc, work, &lwork, &iinfo);
#else
magma_dormqr( side, trans,
m, n, k, A, lda, tau, C, ldc, work, lwork, &iinfo);
#endif
}
else if (nq > 1) {
/* Q was determined by a call to DGEBRD with nq < k */
if (left) {
mi = m - 1;
ni = n;
i1 = 1;
i2 = 0;
}
else {
mi = m;
ni = n - 1;
i1 = 0;
i2 = 1;
}
#if VERSION == 1
nq_1 = nq - 1;
lapackf77_dormqr( lapack_side_const(side), lapack_trans_const(trans),
&mi, &ni, &nq_1, A(1,0), &lda, tau, C(i1,i2), &ldc, work, &lwork, &iinfo);
#else
magma_dormqr( side, trans,
mi, ni, nq-1, A(1,0), lda, tau, C(i1,i2), ldc, work, lwork, &iinfo);
#endif
}
}
else {
/* Apply P */
if (notran) {
transt = MagmaTrans;
}
else {
transt = MagmaNoTrans;
}
if (nq > k) {
/* P was determined by a call to DGEBRD with nq > k */
#if VERSION == 1
lapackf77_dormlq( lapack_side_const(side), lapack_trans_const(transt),
&m, &n, &k, A, &lda, tau, C, &ldc, work, &lwork, &iinfo);
#else
magma_dormlq( side, transt,
m, n, k, A, lda, tau, C, ldc, work, lwork, &iinfo);
#endif
}
else if (nq > 1) {
/* P was determined by a call to DGEBRD with nq <= k */
if (left) {
mi = m - 1;
ni = n;
i1 = 1;
i2 = 0;
}
else {
mi = m;
ni = n - 1;
i1 = 0;
i2 = 1;
}
#if VERSION == 1
nq_1 = nq - 1;
lapackf77_dormlq( lapack_side_const(side), lapack_trans_const(transt),
&mi, &ni, &nq_1, A(0,1), &lda, tau, C(i1,i2), &ldc, work, &lwork, &iinfo);
#else
magma_dormlq( side, transt,
mi, ni, nq-1, A(0,1), lda, tau, C(i1,i2), ldc, work, lwork, &iinfo);
#endif
}
}
work[0] = MAGMA_D_MAKE( lwkopt, 0 );
return *info;
} /* magma_dormbr */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing clarfb_gpu
*/
int main( int argc, char** argv )
{
TESTING_INIT();
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t M, N, K, size, ldc, ldv, ldt, ldw, nv;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
float error, work[1];
magma_int_t status = 0;
// test all combinations of input parameters
magma_side_t side [] = { MagmaLeft, MagmaRight };
magma_trans_t trans [] = { MagmaConjTrans, MagmaNoTrans };
magma_direct_t direct[] = { MagmaForward, MagmaBackward };
magma_storev_t storev[] = { MagmaColumnwise, MagmaRowwise };
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
printf(" M N K storev side direct trans ||R||_F / ||HC||_F\n");
printf("========================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
M = opts.msize[itest];
N = opts.nsize[itest];
K = opts.ksize[itest];
if ( M < K || N < K || K <= 0 ) {
printf( "%5d %5d %5d skipping because clarfb requires M >= K, N >= K, K >= 0\n",
(int) M, (int) N, (int) K );
continue;
}
for( int istor = 0; istor < 2; ++istor ) {
for( int iside = 0; iside < 2; ++iside ) {
for( int idir = 0; idir < 2; ++idir ) {
for( int itran = 0; itran < 2; ++itran ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
ldc = ((M+31)/32)*32;
ldt = ((K+31)/32)*32;
ldw = (side[iside] == MagmaLeft ? N : M);
// (ldv, nv) get swapped later if rowwise
ldv = (side[iside] == MagmaLeft ? M : N);
nv = K;
// Allocate memory for matrices
magmaFloatComplex *C, *R, *V, *T, *W;
TESTING_MALLOC_CPU( C, magmaFloatComplex, ldc*N );
TESTING_MALLOC_CPU( R, magmaFloatComplex, ldc*N );
TESTING_MALLOC_CPU( V, magmaFloatComplex, ldv*K );
TESTING_MALLOC_CPU( T, magmaFloatComplex, ldt*K );
TESTING_MALLOC_CPU( W, magmaFloatComplex, ldw*K );
magmaFloatComplex_ptr dC, dV, dT, dW;
TESTING_MALLOC_DEV( dC, magmaFloatComplex, ldc*N );
TESTING_MALLOC_DEV( dV, magmaFloatComplex, ldv*K );
TESTING_MALLOC_DEV( dT, magmaFloatComplex, ldt*K );
TESTING_MALLOC_DEV( dW, magmaFloatComplex, ldw*K );
// C is M x N.
size = ldc*N;
lapackf77_clarnv( &ione, ISEED, &size, C );
//printf( "C=" ); magma_cprint( M, N, C, ldc );
// V is ldv x nv. See larfb docs for description.
// if column-wise and left, M x K
// if column-wise and right, N x K
// if row-wise and left, K x M
// if row-wise and right, K x N
size = ldv*nv;
lapackf77_clarnv( &ione, ISEED, &size, V );
if ( storev[istor] == MagmaColumnwise ) {
if ( direct[idir] == MagmaForward ) {
lapackf77_claset( MagmaUpperStr, &K, &K, &c_zero, &c_one, V, &ldv );
}
else {
lapackf77_claset( MagmaLowerStr, &K, &K, &c_zero, &c_one, &V[(ldv-K)], &ldv );
}
}
else {
// rowwise, swap V's dimensions
std::swap( ldv, nv );
if ( direct[idir] == MagmaForward ) {
lapackf77_claset( MagmaLowerStr, &K, &K, &c_zero, &c_one, V, &ldv );
}
else {
lapackf77_claset( MagmaUpperStr, &K, &K, &c_zero, &c_one, &V[(nv-K)*ldv], &ldv );
}
}
//printf( "# ldv %d, nv %d\n", ldv, nv );
//printf( "V=" ); magma_cprint( ldv, nv, V, ldv );
// T is K x K, upper triangular for forward, and lower triangular for backward
magma_int_t k1 = K-1;
size = ldt*K;
lapackf77_clarnv( &ione, ISEED, &size, T );
if ( direct[idir] == MagmaForward ) {
lapackf77_claset( MagmaLowerStr, &k1, &k1, &c_zero, &c_zero, &T[1], &ldt );
}
else {
lapackf77_claset( MagmaUpperStr, &k1, &k1, &c_zero, &c_zero, &T[1*ldt], &ldt );
}
//printf( "T=" ); magma_cprint( K, K, T, ldt );
magma_csetmatrix( M, N, C, ldc, dC, ldc );
magma_csetmatrix( ldv, nv, V, ldv, dV, ldv );
magma_csetmatrix( K, K, T, ldt, dT, ldt );
lapackf77_clarfb( lapack_side_const( side[iside] ), lapack_trans_const( trans[itran] ),
lapack_direct_const( direct[idir] ), lapack_storev_const( storev[istor] ),
&M, &N, &K,
V, &ldv, T, &ldt, C, &ldc, W, &ldw );
//printf( "HC=" ); magma_cprint( M, N, C, ldc );
magma_clarfb_gpu( side[iside], trans[itran], direct[idir], storev[istor],
M, N, K,
dV, ldv, dT, ldt, dC, ldc, dW, ldw );
magma_cgetmatrix( M, N, dC, ldc, R, ldc );
//printf( "dHC=" ); magma_cprint( M, N, R, ldc );
// compute relative error |HC_magma - HC_lapack| / |HC_lapack|
error = lapackf77_clange( "Fro", &M, &N, C, &ldc, work );
size = ldc*N;
blasf77_caxpy( &size, &c_neg_one, C, &ione, R, &ione );
error = lapackf77_clange( "Fro", &M, &N, R, &ldc, work ) / error;
printf( "%5d %5d %5d %c %c %c %c %8.2e %s\n",
(int) M, (int) N, (int) K,
lapacke_storev_const(storev[istor]), lapacke_side_const(side[iside]),
lapacke_direct_const(direct[idir]), lapacke_trans_const(trans[itran]),
error, (error < tol ? "ok" : "failed") );
status += ! (error < tol);
TESTING_FREE_CPU( C );
TESTING_FREE_CPU( R );
TESTING_FREE_CPU( V );
TESTING_FREE_CPU( T );
TESTING_FREE_CPU( W );
TESTING_FREE_DEV( dC );
TESTING_FREE_DEV( dV );
TESTING_FREE_DEV( dT );
TESTING_FREE_DEV( dW );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}}}}
printf( "\n" );
}
TESTING_FINALIZE();
return status;
}
/**
Purpose
-------
SGEEV computes for an N-by-N real nonsymmetric matrix A, the
eigenvalues and, optionally, the left and/or right eigenvectors.
The right eigenvector v(j) of A satisfies
A * v(j) = lambda(j) * v(j)
where lambda(j) is its eigenvalue.
The left eigenvector u(j) of A satisfies
u(j)**T * A = lambda(j) * u(j)**T
where u(j)**T denotes the transpose of u(j).
The computed eigenvectors are normalized to have Euclidean norm
equal to 1 and largest component real.
Arguments
---------
@param[in]
jobvl magma_vec_t
- = MagmaNoVec: left eigenvectors of A are not computed;
- = MagmaVec: left eigenvectors of are computed.
@param[in]
jobvr magma_vec_t
- = MagmaNoVec: right eigenvectors of A are not computed;
- = MagmaVec: right eigenvectors of A are computed.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
A REAL array, dimension (LDA,N)
On entry, the N-by-N matrix A.
On exit, A has been overwritten.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
wr REAL array, dimension (N)
@param[out]
wi REAL array, dimension (N)
WR and WI contain the real and imaginary parts,
respectively, of the computed eigenvalues. Complex
conjugate pairs of eigenvalues appear consecutively
with the eigenvalue having the positive imaginary part
first.
@param[out]
VL REAL array, dimension (LDVL,N)
If JOBVL = MagmaVec, the left eigenvectors u(j) are stored one
after another in the columns of VL, in the same order
as their eigenvalues.
If JOBVL = MagmaNoVec, VL is not referenced.
u(j) = VL(:,j), the j-th column of VL.
@param[in]
ldvl INTEGER
The leading dimension of the array VL. LDVL >= 1; if
JOBVL = MagmaVec, LDVL >= N.
@param[out]
VR REAL array, dimension (LDVR,N)
If JOBVR = MagmaVec, the right eigenvectors v(j) are stored one
after another in the columns of VR, in the same order
as their eigenvalues.
If JOBVR = MagmaNoVec, VR is not referenced.
v(j) = VR(:,j), the j-th column of VR.
@param[in]
ldvr INTEGER
The leading dimension of the array VR. LDVR >= 1; if
JOBVR = MagmaVec, LDVR >= N.
@param[out]
work (workspace) REAL array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK. LWORK >= (2 + nb + nb*ngpu)*N.
For optimal performance, LWORK >= (2 + 2*nb + nb*ngpu)*N.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value.
- > 0: if INFO = i, the QR algorithm failed to compute all the
eigenvalues, and no eigenvectors have been computed;
elements and i+1:N of W contain eigenvalues which have
converged.
@ingroup magma_sgeev_driver
********************************************************************/
extern "C" magma_int_t
magma_sgeev_m(
magma_vec_t jobvl, magma_vec_t jobvr, magma_int_t n,
float *A, magma_int_t lda,
#ifdef COMPLEX
float *w,
#else
float *wr, float *wi,
#endif
float *VL, magma_int_t ldvl,
float *VR, magma_int_t ldvr,
float *work, magma_int_t lwork,
#ifdef COMPLEX
float *rwork,
#endif
magma_int_t *info )
{
#define VL(i,j) (VL + (i) + (j)*ldvl)
#define VR(i,j) (VR + (i) + (j)*ldvr)
const magma_int_t ione = 1;
const magma_int_t izero = 0;
float d__1, d__2;
float r, cs, sn, scl;
float dum[1], eps;
float anrm, cscale, bignum, smlnum;
magma_int_t i, k, ilo, ihi;
magma_int_t ibal, ierr, itau, iwrk, nout, liwrk, nb;
magma_int_t scalea, minwrk, optwrk, lquery, wantvl, wantvr, select[1];
magma_side_t side = MagmaRight;
magma_int_t ngpu = magma_num_gpus();
magma_timer_t time_total=0, time_gehrd=0, time_unghr=0, time_hseqr=0, time_trevc=0, time_sum=0;
magma_flops_t flop_total=0, flop_gehrd=0, flop_unghr=0, flop_hseqr=0, flop_trevc=0, flop_sum=0;
timer_start( time_total );
flops_start( flop_total );
*info = 0;
lquery = (lwork == -1);
wantvl = (jobvl == MagmaVec);
wantvr = (jobvr == MagmaVec);
if (! wantvl && jobvl != MagmaNoVec) {
*info = -1;
} else if (! wantvr && jobvr != MagmaNoVec) {
*info = -2;
} else if (n < 0) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
} else if ( (ldvl < 1) || (wantvl && (ldvl < n))) {
*info = -9;
} else if ( (ldvr < 1) || (wantvr && (ldvr < n))) {
*info = -11;
}
/* Compute workspace */
nb = magma_get_sgehrd_nb( n );
if (*info == 0) {
minwrk = (2 + nb + nb*ngpu)*n;
optwrk = (2 + 2*nb + nb*ngpu)*n;
work[0] = magma_smake_lwork( optwrk );
if (lwork < minwrk && ! lquery) {
*info = -13;
}
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (n == 0) {
return *info;
}
#if defined(Version3)
float *dT;
if (MAGMA_SUCCESS != magma_smalloc( &dT, nb*n )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
#endif
#if defined(Version5)
float *T;
if (MAGMA_SUCCESS != magma_smalloc_cpu( &T, nb*n )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
#endif
/* Get machine constants */
eps = lapackf77_slamch( "P" );
smlnum = lapackf77_slamch( "S" );
bignum = 1. / smlnum;
lapackf77_slabad( &smlnum, &bignum );
smlnum = magma_ssqrt( smlnum ) / eps;
bignum = 1. / smlnum;
/* Scale A if max element outside range [SMLNUM,BIGNUM] */
anrm = lapackf77_slange( "M", &n, &n, A, &lda, dum );
scalea = 0;
if (anrm > 0. && anrm < smlnum) {
scalea = 1;
cscale = smlnum;
} else if (anrm > bignum) {
scalea = 1;
cscale = bignum;
}
if (scalea) {
lapackf77_slascl( "G", &izero, &izero, &anrm, &cscale, &n, &n, A, &lda, &ierr );
}
/* Balance the matrix
* (Workspace: need N)
* - this space is reserved until after gebak */
ibal = 0;
lapackf77_sgebal( "B", &n, A, &lda, &ilo, &ihi, &work[ibal], &ierr );
/* Reduce to upper Hessenberg form
* (Workspace: need 3*N, prefer 2*N + N*NB + NB*NGPU)
* - added NB*NGPU needed for multi-GPU magma_sgehrd_m
* - including N reserved for gebal/gebak, unused by sgehrd */
itau = ibal + n;
iwrk = itau + n;
liwrk = lwork - iwrk;
timer_start( time_gehrd );
flops_start( flop_gehrd );
#if defined(Version1)
// Version 1 - LAPACK
lapackf77_sgehrd( &n, &ilo, &ihi, A, &lda,
&work[itau], &work[iwrk], &liwrk, &ierr );
#elif defined(Version2)
// Version 2 - LAPACK consistent HRD
magma_sgehrd2( n, ilo, ihi, A, lda,
&work[itau], &work[iwrk], liwrk, &ierr );
#elif defined(Version3)
// Version 3 - LAPACK consistent MAGMA HRD + T matrices stored,
magma_sgehrd( n, ilo, ihi, A, lda,
&work[itau], &work[iwrk], liwrk, dT, &ierr );
#elif defined(Version5)
// Version 4 - Multi-GPU, T on host
magma_sgehrd_m( n, ilo, ihi, A, lda,
&work[itau], &work[iwrk], liwrk, T, &ierr );
#endif
time_sum += timer_stop( time_gehrd );
flop_sum += flops_stop( flop_gehrd );
if (wantvl) {
/* Want left eigenvectors
* Copy Householder vectors to VL */
side = MagmaLeft;
lapackf77_slacpy( MagmaLowerStr, &n, &n, A, &lda, VL, &ldvl );
/* Generate orthogonal matrix in VL
* (Workspace: need 3*N-1, prefer 2*N + (N-1)*NB)
* - including N reserved for gebal/gebak, unused by sorghr */
timer_start( time_unghr );
flops_start( flop_unghr );
#if defined(Version1) || defined(Version2)
// Version 1 & 2 - LAPACK
lapackf77_sorghr( &n, &ilo, &ihi, VL, &ldvl, &work[itau],
&work[iwrk], &liwrk, &ierr );
#elif defined(Version3)
// Version 3 - LAPACK consistent MAGMA HRD + T matrices stored
magma_sorghr( n, ilo, ihi, VL, ldvl, &work[itau], dT, nb, &ierr );
#elif defined(Version5)
// Version 5 - Multi-GPU, T on host
magma_sorghr_m( n, ilo, ihi, VL, ldvl, &work[itau], T, nb, &ierr );
#endif
time_sum += timer_stop( time_unghr );
flop_sum += flops_stop( flop_unghr );
timer_start( time_hseqr );
flops_start( flop_hseqr );
/* Perform QR iteration, accumulating Schur vectors in VL
* (Workspace: need N+1, prefer N+HSWORK (see comments) )
* - including N reserved for gebal/gebak, unused by shseqr */
iwrk = itau;
liwrk = lwork - iwrk;
lapackf77_shseqr( "S", "V", &n, &ilo, &ihi, A, &lda, wr, wi,
VL, &ldvl, &work[iwrk], &liwrk, info );
time_sum += timer_stop( time_hseqr );
flop_sum += flops_stop( flop_hseqr );
if (wantvr) {
/* Want left and right eigenvectors
* Copy Schur vectors to VR */
side = MagmaBothSides;
lapackf77_slacpy( "F", &n, &n, VL, &ldvl, VR, &ldvr );
}
}
else if (wantvr) {
/* Want right eigenvectors
* Copy Householder vectors to VR */
side = MagmaRight;
lapackf77_slacpy( "L", &n, &n, A, &lda, VR, &ldvr );
/* Generate orthogonal matrix in VR
* (Workspace: need 3*N-1, prefer 2*N + (N-1)*NB)
* - including N reserved for gebal/gebak, unused by sorghr */
timer_start( time_unghr );
flops_start( flop_unghr );
#if defined(Version1) || defined(Version2)
// Version 1 & 2 - LAPACK
lapackf77_sorghr( &n, &ilo, &ihi, VR, &ldvr, &work[itau],
&work[iwrk], &liwrk, &ierr );
#elif defined(Version3)
// Version 3 - LAPACK consistent MAGMA HRD + T matrices stored
magma_sorghr( n, ilo, ihi, VR, ldvr, &work[itau], dT, nb, &ierr );
#elif defined(Version5)
// Version 5 - Multi-GPU, T on host
magma_sorghr_m( n, ilo, ihi, VR, ldvr, &work[itau], T, nb, &ierr );
#endif
time_sum += timer_stop( time_unghr );
flop_sum += flops_stop( flop_unghr );
/* Perform QR iteration, accumulating Schur vectors in VR
* (Workspace: need N+1, prefer N+HSWORK (see comments) )
* - including N reserved for gebal/gebak, unused by shseqr */
timer_start( time_hseqr );
flops_start( flop_hseqr );
iwrk = itau;
liwrk = lwork - iwrk;
lapackf77_shseqr( "S", "V", &n, &ilo, &ihi, A, &lda, wr, wi,
VR, &ldvr, &work[iwrk], &liwrk, info );
time_sum += timer_stop( time_hseqr );
flop_sum += flops_stop( flop_hseqr );
}
else {
/* Compute eigenvalues only
* (Workspace: need N+1, prefer N+HSWORK (see comments) )
* - including N reserved for gebal/gebak, unused by shseqr */
timer_start( time_hseqr );
flops_start( flop_hseqr );
iwrk = itau;
liwrk = lwork - iwrk;
lapackf77_shseqr( "E", "N", &n, &ilo, &ihi, A, &lda, wr, wi,
VR, &ldvr, &work[iwrk], &liwrk, info );
time_sum += timer_stop( time_hseqr );
flop_sum += flops_stop( flop_hseqr );
}
/* If INFO > 0 from SHSEQR, then quit */
if (*info > 0) {
goto CLEANUP;
}
timer_start( time_trevc );
flops_start( flop_trevc );
if (wantvl || wantvr) {
/* Compute left and/or right eigenvectors
* (Workspace: need 4*N, prefer (2 + 2*nb)*N)
* - including N reserved for gebal/gebak, unused by strevc */
liwrk = lwork - iwrk;
#if TREVC_VERSION == 1
lapackf77_strevc( lapack_side_const(side), "B", select, &n, A, &lda, VL, &ldvl,
VR, &ldvr, &n, &nout, &work[iwrk], &ierr );
#elif TREVC_VERSION == 2
lapackf77_strevc3( lapack_side_const(side), "B", select, &n, A, &lda, VL, &ldvl,
VR, &ldvr, &n, &nout, &work[iwrk], &liwrk, &ierr );
#elif TREVC_VERSION == 3
magma_strevc3( side, MagmaBacktransVec, select, n, A, lda, VL, ldvl,
VR, ldvr, n, &nout, &work[iwrk], liwrk, &ierr );
#elif TREVC_VERSION == 4
magma_strevc3_mt( side, MagmaBacktransVec, select, n, A, lda, VL, ldvl,
VR, ldvr, n, &nout, &work[iwrk], liwrk, &ierr );
#elif TREVC_VERSION == 5
magma_strevc3_mt_gpu( side, MagmaBacktransVec, select, n, A, lda, VL, ldvl,
VR, ldvr, n, &nout, &work[iwrk], liwrk, &ierr );
#else
#error Unknown TREVC_VERSION
#endif
}
time_sum += timer_stop( time_trevc );
flop_sum += flops_stop( flop_trevc );
if (wantvl) {
/* Undo balancing of left eigenvectors
* (Workspace: need N) */
lapackf77_sgebak( "B", "L", &n, &ilo, &ihi, &work[ibal], &n,
VL, &ldvl, &ierr );
/* Normalize left eigenvectors and make largest component real */
for (i = 0; i < n; ++i) {
if ( wi[i] == 0. ) {
scl = 1. / magma_cblas_snrm2( n, VL(0,i), 1 );
blasf77_sscal( &n, &scl, VL(0,i), &ione );
}
else if ( wi[i] > 0. ) {
d__1 = magma_cblas_snrm2( n, VL(0,i), 1 );
d__2 = magma_cblas_snrm2( n, VL(0,i+1), 1 );
scl = 1. / lapackf77_slapy2( &d__1, &d__2 );
blasf77_sscal( &n, &scl, VL(0,i), &ione );
blasf77_sscal( &n, &scl, VL(0,i+1), &ione );
for (k = 0; k < n; ++k) {
/* Computing 2nd power */
d__1 = *VL(k,i);
d__2 = *VL(k,i+1);
work[iwrk + k] = d__1*d__1 + d__2*d__2;
}
k = blasf77_isamax( &n, &work[iwrk], &ione ) - 1; // subtract 1; k is 0-based
lapackf77_slartg( VL(k,i), VL(k,i+1), &cs, &sn, &r );
blasf77_srot( &n, VL(0,i), &ione, VL(0,i+1), &ione, &cs, &sn );
*VL(k,i+1) = 0.;
}
}
}
if (wantvr) {
/* Undo balancing of right eigenvectors
* (Workspace: need N) */
lapackf77_sgebak( "B", "R", &n, &ilo, &ihi, &work[ibal], &n,
VR, &ldvr, &ierr );
/* Normalize right eigenvectors and make largest component real */
for (i = 0; i < n; ++i) {
if ( wi[i] == 0. ) {
scl = 1. / magma_cblas_snrm2( n, VR(0,i), 1 );
blasf77_sscal( &n, &scl, VR(0,i), &ione );
}
else if ( wi[i] > 0. ) {
d__1 = magma_cblas_snrm2( n, VR(0,i), 1 );
d__2 = magma_cblas_snrm2( n, VR(0,i+1), 1 );
scl = 1. / lapackf77_slapy2( &d__1, &d__2 );
blasf77_sscal( &n, &scl, VR(0,i), &ione );
blasf77_sscal( &n, &scl, VR(0,i+1), &ione );
for (k = 0; k < n; ++k) {
/* Computing 2nd power */
d__1 = *VR(k,i);
d__2 = *VR(k,i+1);
work[iwrk + k] = d__1*d__1 + d__2*d__2;
}
k = blasf77_isamax( &n, &work[iwrk], &ione ) - 1; // subtract 1; k is 0-based
lapackf77_slartg( VR(k,i), VR(k,i+1), &cs, &sn, &r );
blasf77_srot( &n, VR(0,i), &ione, VR(0,i+1), &ione, &cs, &sn );
*VR(k,i+1) = 0.;
}
}
}
CLEANUP:
/* Undo scaling if necessary */
if (scalea) {
// converged eigenvalues, stored in wr[i+1:n] and wi[i+1:n] for i = INFO
magma_int_t nval = n - (*info);
magma_int_t ld = max( nval, 1 );
lapackf77_slascl( "G", &izero, &izero, &cscale, &anrm, &nval, &ione, wr + (*info), &ld, &ierr );
lapackf77_slascl( "G", &izero, &izero, &cscale, &anrm, &nval, &ione, wi + (*info), &ld, &ierr );
if (*info > 0) {
// first ilo columns were already upper triangular,
// so the corresponding eigenvalues are also valid.
nval = ilo - 1;
lapackf77_slascl( "G", &izero, &izero, &cscale, &anrm, &nval, &ione, wr, &n, &ierr );
lapackf77_slascl( "G", &izero, &izero, &cscale, &anrm, &nval, &ione, wi, &n, &ierr );
}
}
#if defined(Version3)
magma_free( dT );
#endif
#if defined(Version5)
magma_free_cpu( T );
#endif
timer_stop( time_total );
flops_stop( flop_total );
timer_printf( "sgeev times n %5d, gehrd %7.3f, unghr %7.3f, hseqr %7.3f, trevc %7.3f, total %7.3f, sum %7.3f\n",
(int) n, time_gehrd, time_unghr, time_hseqr, time_trevc, time_total, time_sum );
timer_printf( "sgeev flops n %5d, gehrd %7lld, unghr %7lld, hseqr %7lld, trevc %7lld, total %7lld, sum %7lld\n",
(int) n, flop_gehrd, flop_unghr, flop_hseqr, flop_trevc, flop_total, flop_sum );
work[0] = magma_smake_lwork( optwrk );
return *info;
} /* magma_sgeev */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing ctrsm
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time=0, cublas_perf, cublas_time, cpu_perf=0, cpu_time=0;
float magma_error, cublas_error, work[1];
magma_int_t M, N, info;
magma_int_t Ak;
magma_int_t sizeA, sizeB;
magma_int_t lda, ldb, ldda, lddb;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t *ipiv;
magmaFloatComplex *h_A, *h_B, *h_Bcublas, *h_Bmagma, *h_B1, *h_X1, *h_X2;
magmaFloatComplex *d_A, *d_B;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex alpha = MAGMA_C_MAKE( 0.29, -0.86 );
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
printf("side = %s, uplo = %s, transA = %s, diag = %s \n",
lapack_side_const(opts.side), lapack_uplo_const(opts.uplo),
lapack_trans_const(opts.transA), lapack_diag_const(opts.diag) );
printf(" M N MAGMA Gflop/s (ms) CUBLAS Gflop/s (ms) CPU Gflop/s (ms) MAGMA error CUBLAS error\n");
printf("==================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
gflops = FLOPS_CTRSM(opts.side, M, N) / 1e9;
if ( opts.side == MagmaLeft ) {
lda = M;
Ak = M;
} else {
lda = N;
Ak = N;
}
ldb = M;
ldda = ((lda+31)/32)*32;
lddb = ((ldb+31)/32)*32;
sizeA = lda*Ak;
sizeB = ldb*N;
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, lda*Ak );
TESTING_MALLOC_CPU( h_B, magmaFloatComplex, ldb*N );
TESTING_MALLOC_CPU( h_B1, magmaFloatComplex, ldb*N );
TESTING_MALLOC_CPU( h_X1, magmaFloatComplex, ldb*N );
TESTING_MALLOC_CPU( h_X2, magmaFloatComplex, ldb*N );
TESTING_MALLOC_CPU( h_Bcublas, magmaFloatComplex, ldb*N );
TESTING_MALLOC_CPU( h_Bmagma, magmaFloatComplex, ldb*N );
TESTING_MALLOC_CPU( ipiv, magma_int_t, Ak );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*Ak );
TESTING_MALLOC_DEV( d_B, magmaFloatComplex, lddb*N );
/* Initialize the matrices */
/* Factor A into LU to get well-conditioned triangular matrix.
* Copy L to U, since L seems okay when used with non-unit diagonal
* (i.e., from U), while U fails when used with unit diagonal. */
lapackf77_clarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_cgetrf( &Ak, &Ak, h_A, &lda, ipiv, &info );
for( int j = 0; j < Ak; ++j ) {
for( int i = 0; i < j; ++i ) {
*h_A(i,j) = *h_A(j,i);
}
}
lapackf77_clarnv( &ione, ISEED, &sizeB, h_B );
memcpy(h_B1, h_B, sizeB*sizeof(magmaFloatComplex));
/* =====================================================================
Performs operation using MAGMABLAS
=================================================================== */
magma_csetmatrix( Ak, Ak, h_A, lda, d_A, ldda );
magma_csetmatrix( M, N, h_B, ldb, d_B, lddb );
magma_time = magma_sync_wtime( NULL );
magmablas_ctrsm( opts.side, opts.uplo, opts.transA, opts.diag,
M, N,
alpha, d_A, ldda,
d_B, lddb );
magma_time = magma_sync_wtime( NULL ) - magma_time;
magma_perf = gflops / magma_time;
magma_cgetmatrix( M, N, d_B, lddb, h_Bmagma, ldb );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_csetmatrix( M, N, h_B, ldb, d_B, lddb );
cublas_time = magma_sync_wtime( NULL );
cublasCtrsm( handle, cublas_side_const(opts.side), cublas_uplo_const(opts.uplo),
cublas_trans_const(opts.transA), cublas_diag_const(opts.diag),
M, N,
&alpha, d_A, ldda,
d_B, lddb );
cublas_time = magma_sync_wtime( NULL ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_cgetmatrix( M, N, d_B, lddb, h_Bcublas, ldb );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
blasf77_ctrsm( lapack_side_const(opts.side), lapack_uplo_const(opts.uplo), lapack_trans_const(opts.transA), lapack_diag_const(opts.diag),
&M, &N,
&alpha, h_A, &lda,
h_B, &ldb );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
}
/* =====================================================================
Check the result
=================================================================== */
// ||b - Ax|| / (||A||*||x||)
memcpy(h_X1, h_Bmagma, sizeB*sizeof(magmaFloatComplex));
magmaFloatComplex alpha2 = MAGMA_C_DIV( c_one, alpha );
blasf77_ctrmm( lapack_side_const(opts.side), lapack_uplo_const(opts.uplo), lapack_trans_const(opts.transA), lapack_diag_const(opts.diag),
&M, &N,
&alpha2, h_A, &lda,
h_X1, &ldb );
blasf77_caxpy( &sizeB, &c_neg_one, h_B1, &ione, h_X1, &ione );
float norm1 = lapackf77_clange( "M", &M, &N, h_X1, &ldb, work );
float normx = lapackf77_clange( "M", &M, &N, h_Bmagma, &ldb, work );
float normA = lapackf77_clange( "M", &Ak, &Ak, h_A, &lda, work );
magma_error = norm1/(normx*normA);
memcpy(h_X2, h_Bcublas, sizeB*sizeof(magmaFloatComplex));
blasf77_ctrmm( lapack_side_const(opts.side), lapack_uplo_const(opts.uplo), lapack_trans_const(opts.transA), lapack_diag_const(opts.diag),
&M, &N,
&alpha2, h_A, &lda,
h_X2, &ldb );
blasf77_caxpy( &sizeB, &c_neg_one, h_B1, &ione, h_X2, &ione );
norm1 = lapackf77_clange( "M", &M, &N, h_X2, &ldb, work );
normx = lapackf77_clange( "M", &M, &N, h_Bcublas, &ldb, work );
normA = lapackf77_clange( "M", &Ak, &Ak, h_A, &lda, work );
cublas_error = norm1/(normx*normA);
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %s\n",
(int) M, (int) N,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
magma_error, cublas_error,
(magma_error < tol && cublas_error < tol? "ok" : "failed"));
status += ! (magma_error < tol && cublas_error < tol);
}
else {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) --- ( --- ) %8.2e %8.2e %s\n",
(int) M, (int) N,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
magma_error, cublas_error,
(magma_error < tol && cublas_error < tol? "ok" : "failed"));
status += ! (magma_error < tol && cublas_error < tol);
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_B1 );
TESTING_FREE_CPU( h_X1 );
TESTING_FREE_CPU( h_X2 );
TESTING_FREE_CPU( h_Bcublas );
TESTING_FREE_CPU( h_Bmagma );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/**
Purpose
-------
CUNMQL overwrites the general complex M-by-N matrix C with
@verbatim
SIDE = MagmaLeft SIDE = MagmaRight
TRANS = MagmaNoTrans: Q * C C * Q
TRANS = Magma_ConjTrans: Q**H * C C * Q**H
@endverbatim
where Q is a complex unitary matrix defined as the product of k
elementary reflectors
Q = H(k) . . . H(2) H(1)
as returned by CGEQLF.
Q is of order M if SIDE = MagmaLeft
and of order N if SIDE = MagmaRight.
Arguments
---------
@param[in]
side magma_side_t
- = MagmaLeft: apply Q or Q**H from the Left;
- = MagmaRight: apply Q or Q**H from the Right.
@param[in]
trans magma_trans_t
- = MagmaNoTrans: No transpose, apply Q;
- = Magma_ConjTrans: Conjugate transpose, apply Q**H.
@param[in]
m INTEGER
The number of rows of the matrix C. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix C. N >= 0.
@param[in]
k INTEGER
The number of elementary reflectors whose product defines
the matrix Q.
If SIDE = MagmaLeft, M >= K >= 0;
if SIDE = MagmaRight, N >= K >= 0.
@param[in]
A COMPLEX array, dimension (LDA,K)
The i-th column must contain the vector which defines the
elementary reflector H(i), for i = 1,2,...,k, as returned by
CGEQLF in the last k columns of its array argument A.
A is modified by the routine but restored on exit.
@param[in]
lda INTEGER
The leading dimension of the array A.
If SIDE = MagmaLeft, LDA >= max(1,M);
if SIDE = MagmaRight, LDA >= max(1,N).
@param[in]
tau COMPLEX array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by CGEQLF.
@param[in,out]
C COMPLEX array, dimension (LDC,N)
On entry, the M-by-N matrix C.
On exit, C is overwritten by Q*C or Q**H*C or C*Q**H or C*Q.
@param[in]
ldc INTEGER
The leading dimension of the array C. LDC >= max(1,M).
@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 dimension of the array WORK.
If SIDE = MagmaLeft, LWORK >= max(1,N);
if SIDE = MagmaRight, LWORK >= max(1,M).
For optimum performance
if SIDE = MagmaLeft, LWORK >= N*NB;
if SIDE = MagmaRight, LWORK >= M*NB,
where NB is the optimal blocksize.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_cgeqlf_comp
********************************************************************/
extern "C" magma_int_t
magma_cunmql(
magma_side_t side, magma_trans_t trans,
magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloatComplex *A, magma_int_t lda,
magmaFloatComplex *tau,
magmaFloatComplex *C, magma_int_t ldc,
magmaFloatComplex *work, magma_int_t lwork,
magma_int_t *info)
{
#define A(i_,j_) ( A + (i_) + (j_)*lda)
#define dC(i_,j_) (dC + (i_) + (j_)*lddc)
magmaFloatComplex *T, *T2;
magma_int_t i, i1, i2, ib, nb, mi, ni, nq, nq_i, nw, step;
magma_int_t iinfo, ldwork, lwkopt;
*info = 0;
bool left = (side == MagmaLeft);
bool notran = (trans == MagmaNoTrans);
bool lquery = (lwork == -1);
/* NQ is the order of Q and NW is the minimum dimension of WORK */
if (left) {
nq = m;
nw = n;
} else {
nq = n;
nw = m;
}
/* Test the input arguments */
if (! left && side != MagmaRight) {
*info = -1;
} else if (! notran && trans != Magma_ConjTrans) {
*info = -2;
} else if (m < 0) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (k < 0 || k > nq) {
*info = -5;
} else if (lda < max(1,nq)) {
*info = -7;
} else if (ldc < max(1,m)) {
*info = -10;
} else if (lwork < max(1,nw) && ! lquery) {
*info = -12;
}
if (*info == 0) {
nb = magma_get_cgelqf_nb( m, n );
lwkopt = max(1,nw)*nb;
work[0] = magma_cmake_lwork( lwkopt );
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0 || k == 0) {
work[0] = MAGMA_C_ONE;
return *info;
}
ldwork = nw;
if ( nb >= k ) {
/* Use CPU code */
lapackf77_cunmql( lapack_side_const(side), lapack_trans_const(trans),
&m, &n, &k, A, &lda, tau, C, &ldc, work, &lwork, &iinfo );
}
else {
/* Use hybrid CPU-GPU code */
/* Allocate work space on the GPU.
* nw*nb for dwork (m or n) by nb
* nq*nb for dV (n or m) by nb
* nb*nb for dT
* lddc*n for dC.
*/
magma_int_t lddc = magma_roundup( m, 32 );
magmaFloatComplex *dwork, *dV, *dT, *dC;
magma_cmalloc( &dwork, (nw + nq + nb)*nb + lddc*n );
if ( dwork == NULL ) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dV = dwork + nw*nb;
dT = dV + nq*nb;
dC = dT + nb*nb;
/* work space on CPU.
* nb*nb for T
* nb*nb for T2, used to save and restore diagonal block of panel */
magma_cmalloc_pinned( &T, 2*nb*nb );
if ( T == NULL ) {
magma_free( dwork );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
T2 = T + nb*nb;
magma_queue_t queue;
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queue );
/* Copy matrix C from the CPU to the GPU */
magma_csetmatrix( m, n, C, ldc, dC, lddc, queue );
if ( (left && notran) || (! left && ! notran) ) {
i1 = 0;
i2 = k;
step = nb;
} else {
i1 = ((k - 1) / nb) * nb;
i2 = 0;
step = -nb;
}
// silence "uninitialized" warnings
mi = 0;
ni = 0;
if (left) {
ni = n;
} else {
mi = m;
}
for (i = i1; (step < 0 ? i >= i2 : i < i2); i += step) {
ib = min(nb, k - i);
/* Form the triangular factor of the block reflector
H = H(i+ib-1) . . . H(i+1) H(i) */
nq_i = nq - k + i + ib;
lapackf77_clarft("Backward", "Columnwise", &nq_i, &ib,
A(0,i), &lda, &tau[i], T, &ib);
/* 1) set lower triangle of panel in A to identity,
2) copy the panel from A to the GPU, and
3) restore A */
magma_cpanel_to_q( MagmaLower, ib, A(nq_i-ib,i), lda, T2 );
magma_csetmatrix( nq_i, ib, A(0,i), lda, dV, nq_i, queue );
magma_cq_to_panel( MagmaLower, ib, A(nq_i-ib,i), lda, T2 );
if (left) {
/* H or H**H is applied to C(1:m-k+i+ib-1,1:n) */
mi = m - k + i + ib;
}
else {
/* H or H**H is applied to C(1:m,1:n-k+i+ib-1) */
ni = n - k + i + ib;
}
/* Apply H or H**H; First copy T to the GPU */
magma_csetmatrix( ib, ib, T, ib, dT, ib, queue );
magma_clarfb_gpu( side, trans, MagmaBackward, MagmaColumnwise,
mi, ni, ib,
dV, nq_i,
dT, ib,
dC, lddc,
dwork, ldwork, queue );
}
magma_cgetmatrix( m, n, dC, lddc, C, ldc, queue );
magma_queue_destroy( queue );
magma_free( dwork );
magma_free_pinned( T );
}
work[0] = magma_cmake_lwork( lwkopt );
return *info;
} /* magma_cunmql */
/**
Purpose
-------
DORMLQ overwrites the general real M-by-N matrix C with
@verbatim
SIDE = MagmaLeft SIDE = MagmaRight
TRANS = MagmaNoTrans: Q * C C * Q
TRANS = MagmaTrans: Q**H * C C * Q**H
@endverbatim
where Q is a realunitary matrix defined as the product of k
elementary reflectors
Q = H(k)**H . . . H(2)**H H(1)**H
as returned by DGELQF. Q is of order M if SIDE = MagmaLeft and of order N
if SIDE = MagmaRight.
Arguments
---------
@param[in]
side magma_side_t
- = MagmaLeft: apply Q or Q**H from the Left;
- = MagmaRight: apply Q or Q**H from the Right.
@param[in]
trans magma_trans_t
- = MagmaNoTrans: No transpose, apply Q;
- = MagmaTrans: Conjugate transpose, apply Q**H.
@param[in]
m INTEGER
The number of rows of the matrix C. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix C. N >= 0.
@param[in]
k INTEGER
The number of elementary reflectors whose product defines
the matrix Q.
If SIDE = MagmaLeft, M >= K >= 0;
if SIDE = MagmaRight, N >= K >= 0.
@param[in]
A DOUBLE_PRECISION array, dimension
(LDA,M) if SIDE = MagmaLeft,
(LDA,N) if SIDE = MagmaRight.
The i-th row must contain the vector which defines the
elementary reflector H(i), for i = 1,2,...,k, as returned by
DGELQF in the first k rows of its array argument A.
A is modified by the routine but restored on exit.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,K).
@param[in]
tau DOUBLE_PRECISION array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by DGELQF.
@param[in,out]
C DOUBLE_PRECISION array, dimension (LDC,N)
On entry, the M-by-N matrix C.
On exit, C is overwritten by Q*C or Q**H*C or C*Q**H or C*Q.
@param[in]
ldc INTEGER
The leading dimension of the array C. LDC >= max(1,M).
@param[out]
work (workspace) DOUBLE_PRECISION array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK.
If SIDE = MagmaLeft, LWORK >= max(1,N);
if SIDE = MagmaRight, LWORK >= max(1,M).
For optimum performance
if SIDE = MagmaLeft, LWORK >= N*NB;
if SIDE = MagmaRight, LWORK >= M*NB,
where NB is the optimal blocksize.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_dgelqf_comp
********************************************************************/
extern "C" magma_int_t
magma_dormlq(
magma_side_t side, magma_trans_t trans,
magma_int_t m, magma_int_t n, magma_int_t k,
double *A, magma_int_t lda,
double *tau,
double *C, magma_int_t ldc,
double *work, magma_int_t lwork,
magma_int_t *info)
{
#define A(i_,j_) ( A + (i_) + (j_)*lda)
#define dC(i_,j_) (dC + (i_) + (j_)*lddc)
double *T, *T2;
magma_int_t i, i1, i2, ib, ic, jc, nb, mi, ni, nq, nq_i, nw, step;
magma_int_t iinfo, ldwork, lwkopt;
magma_int_t left, notran, lquery;
magma_trans_t transt;
*info = 0;
left = (side == MagmaLeft);
notran = (trans == MagmaNoTrans);
lquery = (lwork == -1);
/* NQ is the order of Q and NW is the minimum dimension of WORK */
if (left) {
nq = m;
nw = n;
} else {
nq = n;
nw = m;
}
/* Test the input arguments */
if (! left && side != MagmaRight) {
*info = -1;
} else if (! notran && trans != MagmaTrans) {
*info = -2;
} else if (m < 0) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (k < 0 || k > nq) {
*info = -5;
} else if (lda < max(1,k)) {
*info = -7;
} else if (ldc < max(1,m)) {
*info = -10;
} else if (lwork < max(1,nw) && ! lquery) {
*info = -12;
}
if (*info == 0) {
nb = magma_get_dgelqf_nb( min( m, n ));
lwkopt = max(1,nw)*nb;
work[0] = MAGMA_D_MAKE( lwkopt, 0 );
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0 || k == 0) {
work[0] = MAGMA_D_ONE;
return *info;
}
ldwork = nw;
if (nb >= k) {
/* Use CPU code */
lapackf77_dormlq( lapack_side_const(side), lapack_trans_const(trans),
&m, &n, &k, A, &lda, tau, C, &ldc, work, &lwork, &iinfo);
}
else {
/* Use hybrid CPU-GPU code */
/* Allocate work space on the GPU.
* nw*nb for dwork (m or n) by nb
* nq*nb for dV (n or m) by nb
* nb*nb for dT
* lddc*n for dC.
*/
magma_int_t lddc = ((m+31)/32)*32;
double *dwork, *dV, *dT, *dC;
magma_dmalloc( &dwork, (nw + nq + nb)*nb + lddc*n );
if ( dwork == NULL ) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dV = dwork + nw*nb;
dT = dV + nq*nb;
dC = dT + nb*nb;
/* work space on CPU.
* nb*nb for T
* nb*nb for T2, used to save and restore diagonal block of panel */
magma_dmalloc_cpu( &T, 2*nb*nb );
if ( T == NULL ) {
magma_free( dwork );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
T2 = T + nb*nb;
/* Copy matrix C from the CPU to the GPU */
magma_dsetmatrix( m, n, C, ldc, dC, lddc );
if ( (left && notran) || (! left && ! notran) ) {
i1 = 0;
i2 = k;
step = nb;
} else {
i1 = ((k - 1) / nb)*nb;
i2 = 0;
step = -nb;
}
// silence "uninitialized" warnings
mi = 0;
ni = 0;
if (left) {
ni = n;
jc = 0;
} else {
mi = m;
ic = 0;
}
if (notran) {
transt = MagmaTrans;
} else {
transt = MagmaNoTrans;
}
for (i = i1; (step < 0 ? i >= i2 : i < i2); i += step) {
ib = min(nb, k - i);
/* Form the triangular factor of the block reflector
H = H(i) H(i + 1) . . . H(i + ib-1) */
nq_i = nq - i;
lapackf77_dlarft("Forward", "Rowwise", &nq_i, &ib,
A(i,i), &lda, &tau[i], T, &ib);
/* 1) set upper triangle of panel in A to identity,
2) copy the panel from A to the GPU, and
3) restore A */
dpanel_to_q( MagmaLower, ib, A(i,i), lda, T2 );
magma_dsetmatrix( ib, nq_i, A(i,i), lda, dV, ib );
dq_to_panel( MagmaLower, ib, A(i,i), lda, T2 );
if (left) {
/* H or H**H is applied to C(i:m,1:n) */
mi = m - i;
ic = i;
}
else {
/* H or H**H is applied to C(1:m,i:n) */
ni = n - i;
jc = i;
}
/* Apply H or H**H; First copy T to the GPU */
magma_dsetmatrix( ib, ib, T, ib, dT, ib );
magma_dlarfb_gpu( side, transt, MagmaForward, MagmaRowwise,
mi, ni, ib,
dV, ib,
dT, ib,
dC(ic,jc), lddc,
dwork, ldwork );
}
magma_dgetmatrix( m, n, dC, lddc, C, ldc );
magma_free( dwork );
magma_free_cpu( T );
}
work[0] = MAGMA_D_MAKE( lwkopt, 0 );
return *info;
} /* magma_dormlq */
/**
Purpose
-------
SORMQR_GPU overwrites the general real M-by-N matrix C with
@verbatim
SIDE = MagmaLeft SIDE = MagmaRight
TRANS = MagmaNoTrans: Q * C C * Q
TRANS = MagmaTrans: Q**T * C C * Q**T
@endverbatim
where Q is a real unitary matrix defined as the product of k
elementary reflectors
Q = H(1) H(2) . . . H(k)
as returned by SGEQRF. Q is of order M if SIDE = MagmaLeft and of order N
if SIDE = MagmaRight.
Arguments
---------
@param[in]
side magma_side_t
- = MagmaLeft: apply Q or Q**T from the Left;
- = MagmaRight: apply Q or Q**T from the Right.
@param[in]
trans magma_trans_t
- = MagmaNoTrans: No transpose, apply Q;
- = MagmaTrans: Transpose, apply Q**T.
@param[in]
m INTEGER
The number of rows of the matrix C. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix C. N >= 0.
@param[in]
k INTEGER
The number of elementary reflectors whose product defines
the matrix Q.
If SIDE = MagmaLeft, M >= K >= 0;
if SIDE = MagmaRight, N >= K >= 0.
@param[in]
dA REAL array on the GPU, dimension (LDDA,K)
The i-th column must contain the vector which defines the
elementary reflector H(i), for i = 1,2,...,k, as returned by
SGEQRF in the first k columns of its array argument DA.
DA is modified by the routine but restored on exit.
@param[in]
ldda INTEGER
The leading dimension of the array DA.
If SIDE = MagmaLeft, LDDA >= max(1,M);
if SIDE = MagmaRight, LDDA >= max(1,N).
@param[in]
tau REAL array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by SGEQRF.
@param[in,out]
dC REAL array on the GPU, dimension (LDDC,N)
On entry, the M-by-N matrix C.
On exit, C is overwritten by Q*C or Q**T * C or C * Q**T or C*Q.
@param[in]
lddc INTEGER
The leading dimension of the array DC. LDDC >= max(1,M).
@param[out]
hwork (workspace) REAL array, dimension (MAX(1,LWORK))
\n
Currently, sgetrs_gpu assumes that on exit, hwork contains the last
block of A and C. This will change and *should not be relied on*!
@param[in]
lwork INTEGER
The dimension of the array HWORK.
LWORK >= (M-K+NB)*(N+NB) + N*NB if SIDE = MagmaLeft, and
LWORK >= (N-K+NB)*(M+NB) + M*NB if SIDE = MagmaRight,
where NB is the given blocksize.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the HWORK array, returns
this value as the first entry of the HWORK array, and no error
message related to LWORK is issued by XERBLA.
@param[in]
dT REAL array on the GPU that is the output
(the 9th argument) of magma_sgeqrf_gpu.
@param[in]
nb INTEGER
This is the blocking size that was used in pre-computing DT, e.g.,
the blocking size used in magma_sgeqrf_gpu.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_sgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_sormqr_gpu(magma_side_t side, magma_trans_t trans,
magma_int_t m, magma_int_t n, magma_int_t k,
float *dA, magma_int_t ldda,
float *tau,
float *dC, magma_int_t lddc,
float *hwork, magma_int_t lwork,
float *dT, magma_int_t nb,
magma_int_t *info)
{
#define dA(a_1,a_2) (dA + (a_1) + (a_2)*ldda)
#define dC(a_1,a_2) (dC + (a_1) + (a_2)*lddc)
#define dT(a_1) (dT + (a_1)*nb)
float c_one = MAGMA_S_ONE;
const char* side_ = lapack_side_const( side );
const char* trans_ = lapack_trans_const( trans );
float *dwork;
magma_int_t i, lddwork;
magma_int_t i1, i2, step, ib, ic, jc, ma, mi, ni, nq, nw;
int left, notran, lquery;
magma_int_t lwkopt;
*info = 0;
left = (side == MagmaLeft);
notran = (trans == MagmaNoTrans);
lquery = (lwork == -1);
/* NQ is the order of Q and NW is the minimum dimension of WORK */
if (left) {
nq = m;
nw = n;
} else {
nq = n;
nw = m;
}
lwkopt = (nq - k + nb)*(nw + nb) + nw*nb;
hwork[0] = MAGMA_S_MAKE( lwkopt, 0 );
if ( ! left && side != MagmaRight ) {
*info = -1;
} else if ( ! notran && trans != MagmaTrans ) {
*info = -2;
} else if (m < 0) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (k < 0 || k > nq) {
*info = -5;
} else if (ldda < max(1,nq)) {
*info = -7;
} else if (lddc < max(1,m)) {
*info = -10;
} else if (lwork < lwkopt && ! lquery) {
*info = -12;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0 || k == 0) {
hwork[0] = c_one;
return *info;
}
lddwork = k;
dwork = dT(2*lddwork);
if ( (left && (! notran)) || ((! left) && notran) ) {
// left trans: Q^T C
// right notrans: C Q
// multiply from first block, i = 0, to next-to-last block, i < k-nb
i1 = 0;
i2 = k-nb;
step = nb;
} else {
// left notrans: Q C
// right trans: C Q^T
// multiply from next-to-last block, i = floor((k-1-nb)/nb)*nb, to first block, i = 0
i1 = ((k - 1 - nb) / nb) * nb;
i2 = 0;
step = -nb;
}
if (left) {
ni = n;
jc = 0;
} else {
mi = m;
ic = 0;
}
/* Use unblocked code to multiply last or only block (cases Q*C or C*Q^T). */
// workspace left: A(mi*nb) + C(mi*ni) + work(ni*nb_la) = (m-k-nb)*nb + (m-k-nb)*n + n*nb
// workspace right: A(ni*nb) + C(mi*ni) + work(mi*nb_la) = (n-k-nb)*nb + m*(n-k-nb) + m*nb
if ( step < 0 ) {
// i is beginning of last block
i = i1 - step;
if ( i >= k ) {
i = i1;
}
ib = k - i;
if (left) {
// ni=n, jc=0, H or H^T is applied to C(i:m-1,0:n-1)
mi = m - i;
ma = mi;
ic = i;
}
else {
// mi=m, ic=0, H or H^T is applied to C(0:m-1,i:n-1)
ni = n - i;
ma = ni;
jc = i;
}
float* hA = hwork;
float* hC = hwork + ma*ib;
float* hW = hwork + ma*ib + mi*ni;
magma_int_t lhwork = lwork - (ma*ib + mi*ni);
magma_sgetmatrix( ma, ib, dA(i, i ), ldda, hA, ma );
magma_sgetmatrix( mi, ni, dC(ic, jc), lddc, hC, mi );
lapackf77_sormqr( side_, trans_,
&mi, &ni, &ib,
hA, &ma, tau+i,
hC, &mi,
hW, &lhwork, info );
// send the updated part of C back to the GPU
magma_ssetmatrix( mi, ni, hC, mi, dC(ic, jc), lddc );
}
/* Use blocked code to multiply blocks */
if (nb < k) {
for( i=i1; (step < 0 ? i >= i2 : i < i2); i += step ) {
ib = min(nb, k - i);
if (left) {
// ni=n, jc=0, H or H^T is applied to C(i:m-1,0:n-1)
mi = m - i;
ic = i;
}
else {
// mi=m, ic=0, H or H^T is applied to C(0:m-1,i:n-1)
ni = n - i;
jc = i;
}
magma_slarfb_gpu( side, trans, MagmaForward, MagmaColumnwise,
mi, ni, ib,
dA(i, i ), ldda, dT(i), nb,
dC(ic, jc), lddc, dwork, nw );
}
}
else {
i = i1;
}
/* Use unblocked code to multiply the last or only block (cases Q^T*C or C*Q). */
if ( step > 0 ) {
ib = k-i;
if (left) {
// ni=n, jc=0, H or H^T is applied to C(i:m-1,0:n-1)
mi = m - i;
ma = mi;
ic = i;
}
else {
// mi=m, ic=0, H or H^T is applied to C(0:m-1,i:n-1)
ni = n - i;
ma = ni;
jc = i;
}
float* hA = hwork;
float* hC = hwork + ma*ib;
float* hW = hwork + ma*ib + mi*ni;
magma_int_t lhwork = lwork - (ma*ib + mi*ni);
magma_sgetmatrix( ma, ib, dA(i, i ), ldda, hA, ma );
magma_sgetmatrix( mi, ni, dC(ic, jc), lddc, hC, mi );
lapackf77_sormqr( side_, trans_,
&mi, &ni, &ib,
hA, &ma, tau+i,
hC, &mi,
hW, &lhwork, info );
// send the updated part of C back to the GPU
magma_ssetmatrix( mi, ni, hC, mi, dC(ic, jc), lddc );
}
// TODO sync. For cases Q*C and C*Q^T, last call is magma_slarfb_gpu,
// which is async magma_gemm calls, so sormqr can be unfinished.
// TODO: sgeqrs_gpu ASSUMES that hwork contains the last block of A and C.
// That needs to be fixed, but until then, don't modify hwork[0] here.
// In LAPACK: On exit, if INFO = 0, HWORK(1) returns the optimal LWORK.
//hwork[0] = MAGMA_S_MAKE( lwkopt, 0 );
return *info;
} /* magma_sormqr_gpu */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dtrmm
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, cublas_perf, cublas_time, cpu_perf, cpu_time;
double cublas_error, Cnorm, work[1];
magma_int_t M, N;
magma_int_t Ak;
magma_int_t sizeA, sizeB;
magma_int_t lda, ldb, ldda, lddb;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double *h_A, *h_B, *h_Bcublas;
double *d_A, *d_B;
double c_neg_one = MAGMA_D_NEG_ONE;
double alpha = MAGMA_D_MAKE( 0.29, -0.86 );
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("If running lapack (option --lapack), CUBLAS error is computed\n"
"relative to CPU BLAS result.\n\n");
printf("side = %s, uplo = %s, transA = %s, diag = %s \n",
lapack_side_const(opts.side), lapack_uplo_const(opts.uplo),
lapack_trans_const(opts.transA), lapack_diag_const(opts.diag) );
printf(" M N CUBLAS Gflop/s (ms) CPU Gflop/s (ms) CUBLAS error\n");
printf("==================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
gflops = FLOPS_DTRMM(opts.side, M, N) / 1e9;
if ( opts.side == MagmaLeft ) {
lda = M;
Ak = M;
} else {
lda = N;
Ak = N;
}
ldb = M;
ldda = ((lda+31)/32)*32;
lddb = ((ldb+31)/32)*32;
sizeA = lda*Ak;
sizeB = ldb*N;
TESTING_MALLOC_CPU( h_A, double, lda*Ak );
TESTING_MALLOC_CPU( h_B, double, ldb*N );
TESTING_MALLOC_CPU( h_Bcublas, double, ldb*N );
TESTING_MALLOC_DEV( d_A, double, ldda*Ak );
TESTING_MALLOC_DEV( d_B, double, lddb*N );
/* Initialize the matrices */
lapackf77_dlarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_dlarnv( &ione, ISEED, &sizeB, h_B );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_dsetmatrix( Ak, Ak, h_A, lda, d_A, ldda );
magma_dsetmatrix( M, N, h_B, ldb, d_B, lddb );
// note cublas does trmm out-of-place (i.e., adds output matrix C),
// but allows C=B to do in-place.
cublas_time = magma_sync_wtime( NULL );
cublasDtrmm( handle, cublas_side_const(opts.side), cublas_uplo_const(opts.uplo),
cublas_trans_const(opts.transA), cublas_diag_const(opts.diag),
M, N,
&alpha, d_A, ldda,
d_B, lddb,
d_B, lddb );
cublas_time = magma_sync_wtime( NULL ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_dgetmatrix( M, N, d_B, lddb, h_Bcublas, ldb );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
blasf77_dtrmm( lapack_side_const(opts.side), lapack_uplo_const(opts.uplo),
lapack_trans_const(opts.transA), lapack_diag_const(opts.diag),
&M, &N,
&alpha, h_A, &lda,
h_B, &ldb );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
}
/* =====================================================================
Check the result
=================================================================== */
if ( opts.lapack ) {
// compute relative error for both magma & cublas, relative to lapack,
// |C_magma - C_lapack| / |C_lapack|
Cnorm = lapackf77_dlange( "M", &M, &N, h_B, &ldb, work );
blasf77_daxpy( &sizeB, &c_neg_one, h_B, &ione, h_Bcublas, &ione );
cublas_error = lapackf77_dlange( "M", &M, &N, h_Bcublas, &ldb, work ) / Cnorm;
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) M, (int) N,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
cublas_error, (cublas_error < tol ? "ok" : "failed"));
status += ! (cublas_error < tol);
}
else {
printf("%5d %5d %7.2f (%7.2f) --- ( --- ) --- ---\n",
(int) M, (int) N,
cublas_perf, 1000.*cublas_time);
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_Bcublas );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/***************************************************************************//**
Purpose
-------
CGEEV computes for an N-by-N complex nonsymmetric matrix A, the
eigenvalues and, optionally, the left and/or right eigenvectors.
The right eigenvector v(j) of A satisfies
A * v(j) = lambda(j) * v(j)
where lambda(j) is its eigenvalue.
The left eigenvector u(j) of A satisfies
u(j)**H * A = lambda(j) * u(j)**H
where u(j)**H denotes the conjugate transpose of u(j).
The computed eigenvectors are normalized to have Euclidean norm
equal to 1 and largest component real.
Arguments
---------
@param[in]
jobvl magma_vec_t
- = MagmaNoVec: left eigenvectors of A are not computed;
- = MagmaVec: left eigenvectors of are computed.
@param[in]
jobvr magma_vec_t
- = MagmaNoVec: right eigenvectors of A are not computed;
- = MagmaVec: right eigenvectors of A are computed.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
A COMPLEX array, dimension (LDA,N)
On entry, the N-by-N matrix A.
On exit, A has been overwritten.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
w COMPLEX array, dimension (N)
W contains the computed eigenvalues.
@param[out]
VL COMPLEX array, dimension (LDVL,N)
If JOBVL = MagmaVec, the left eigenvectors u(j) are stored one
after another in the columns of VL, in the same order
as their eigenvalues.
If JOBVL = MagmaNoVec, VL is not referenced.
u(j) = VL(:,j), the j-th column of VL.
@param[in]
ldvl INTEGER
The leading dimension of the array VL. LDVL >= 1; if
JOBVL = MagmaVec, LDVL >= N.
@param[out]
VR COMPLEX array, dimension (LDVR,N)
If JOBVR = MagmaVec, the right eigenvectors v(j) are stored one
after another in the columns of VR, in the same order
as their eigenvalues.
If JOBVR = MagmaNoVec, VR is not referenced.
v(j) = VR(:,j), the j-th column of VR.
@param[in]
ldvr INTEGER
The leading dimension of the array VR. LDVR >= 1; if
JOBVR = MagmaVec, LDVR >= 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 dimension of the array WORK. LWORK >= (1 + nb + nb*ngpu)*N.
For optimal performance, LWORK >= (1 + 2*nb + nb*ngpu)*N.
\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 (2*N)
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value.
- > 0: if INFO = i, the QR algorithm failed to compute all the
eigenvalues, and no eigenvectors have been computed;
elements and i+1:N of W contain eigenvalues which have
converged.
@ingroup magma_geev
*******************************************************************************/
extern "C" magma_int_t
magma_cgeev_m(
magma_vec_t jobvl, magma_vec_t jobvr, magma_int_t n,
magmaFloatComplex *A, magma_int_t lda,
#ifdef COMPLEX
magmaFloatComplex *w,
#else
float *wr, float *wi,
#endif
magmaFloatComplex *VL, magma_int_t ldvl,
magmaFloatComplex *VR, magma_int_t ldvr,
magmaFloatComplex *work, magma_int_t lwork,
#ifdef COMPLEX
float *rwork,
#endif
magma_int_t *info )
{
#define VL(i,j) (VL + (i) + (j)*ldvl)
#define VR(i,j) (VR + (i) + (j)*ldvr)
const magma_int_t ione = 1;
const magma_int_t izero = 0;
float d__1, d__2;
magmaFloatComplex tmp;
float scl;
float dum[1], eps;
float anrm, cscale, bignum, smlnum;
magma_int_t i, k, ilo, ihi;
magma_int_t ibal, ierr, itau, iwrk, nout, liwrk, nb;
magma_int_t scalea, minwrk, optwrk, irwork, lquery, wantvl, wantvr, select[1];
magma_side_t side = MagmaRight;
magma_int_t ngpu = magma_num_gpus();
irwork = 0;
*info = 0;
lquery = (lwork == -1);
wantvl = (jobvl == MagmaVec);
wantvr = (jobvr == MagmaVec);
if (! wantvl && jobvl != MagmaNoVec) {
*info = -1;
} else if (! wantvr && jobvr != MagmaNoVec) {
*info = -2;
} else if (n < 0) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
} else if ( (ldvl < 1) || (wantvl && (ldvl < n))) {
*info = -8;
} else if ( (ldvr < 1) || (wantvr && (ldvr < n))) {
*info = -10;
}
/* Compute workspace */
nb = magma_get_cgehrd_nb( n );
if (*info == 0) {
minwrk = (1 + nb + nb*ngpu)*n;
optwrk = (1 + 2*nb + nb*ngpu)*n;
work[0] = magma_cmake_lwork( optwrk );
if (lwork < minwrk && ! lquery) {
*info = -12;
}
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (n == 0) {
return *info;
}
#if defined(Version3)
magmaFloatComplex *dT;
if (MAGMA_SUCCESS != magma_cmalloc( &dT, nb*n )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
#endif
#if defined(Version5)
magmaFloatComplex *T;
if (MAGMA_SUCCESS != magma_cmalloc_cpu( &T, nb*n )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
#endif
/* Get machine constants */
eps = lapackf77_slamch( "P" );
smlnum = lapackf77_slamch( "S" );
bignum = 1. / smlnum;
lapackf77_slabad( &smlnum, &bignum );
smlnum = magma_ssqrt( smlnum ) / eps;
bignum = 1. / smlnum;
/* Scale A if max element outside range [SMLNUM,BIGNUM] */
anrm = lapackf77_clange( "M", &n, &n, A, &lda, dum );
scalea = 0;
if (anrm > 0. && anrm < smlnum) {
scalea = 1;
cscale = smlnum;
} else if (anrm > bignum) {
scalea = 1;
cscale = bignum;
}
if (scalea) {
lapackf77_clascl( "G", &izero, &izero, &anrm, &cscale, &n, &n, A, &lda, &ierr );
}
/* Balance the matrix
* (CWorkspace: none)
* (RWorkspace: need N)
* - this space is reserved until after gebak */
ibal = 0;
lapackf77_cgebal( "B", &n, A, &lda, &ilo, &ihi, &rwork[ibal], &ierr );
/* Reduce to upper Hessenberg form
* (CWorkspace: need 2*N, prefer N + N*NB + NB*NGPU)
* (RWorkspace: N)
* - added NB*NGPU needed for multi-GPU magma_cgehrd_m
* - including N reserved for gebal/gebak, unused by cgehrd */
itau = 0;
iwrk = itau + n;
liwrk = lwork - iwrk;
#if defined(Version1)
// Version 1 - LAPACK
lapackf77_cgehrd( &n, &ilo, &ihi, A, &lda,
&work[itau], &work[iwrk], &liwrk, &ierr );
#elif defined(Version2)
// Version 2 - LAPACK consistent HRD
magma_cgehrd2( n, ilo, ihi, A, lda,
&work[itau], &work[iwrk], liwrk, &ierr );
#elif defined(Version3)
// Version 3 - LAPACK consistent MAGMA HRD + T matrices stored,
magma_cgehrd( n, ilo, ihi, A, lda,
&work[itau], &work[iwrk], liwrk, dT, &ierr );
#elif defined(Version5)
// Version 4 - Multi-GPU, T on host
magma_cgehrd_m( n, ilo, ihi, A, lda,
&work[itau], &work[iwrk], liwrk, T, &ierr );
#endif
if (wantvl) {
/* Want left eigenvectors
* Copy Householder vectors to VL */
side = MagmaLeft;
lapackf77_clacpy( MagmaLowerStr, &n, &n, A, &lda, VL, &ldvl );
/* Generate unitary matrix in VL
* (CWorkspace: need 2*N-1, prefer N + (N-1)*NB)
* (RWorkspace: N)
* - including N reserved for gebal/gebak, unused by cunghr */
#if defined(Version1) || defined(Version2)
// Version 1 & 2 - LAPACK
lapackf77_cunghr( &n, &ilo, &ihi, VL, &ldvl, &work[itau],
&work[iwrk], &liwrk, &ierr );
#elif defined(Version3)
// Version 3 - LAPACK consistent MAGMA HRD + T matrices stored
magma_cunghr( n, ilo, ihi, VL, ldvl, &work[itau], dT, nb, &ierr );
#elif defined(Version5)
// Version 5 - Multi-GPU, T on host
magma_cunghr_m( n, ilo, ihi, VL, ldvl, &work[itau], T, nb, &ierr );
#endif
/* Perform QR iteration, accumulating Schur vectors in VL
* (CWorkspace: need 1, prefer HSWORK (see comments) )
* (RWorkspace: N)
* - including N reserved for gebal/gebak, unused by chseqr */
iwrk = itau;
liwrk = lwork - iwrk;
lapackf77_chseqr( "S", "V", &n, &ilo, &ihi, A, &lda, w,
VL, &ldvl, &work[iwrk], &liwrk, info );
if (wantvr) {
/* Want left and right eigenvectors
* Copy Schur vectors to VR */
side = MagmaBothSides;
lapackf77_clacpy( "F", &n, &n, VL, &ldvl, VR, &ldvr );
}
}
else if (wantvr) {
/* Want right eigenvectors
* Copy Householder vectors to VR */
side = MagmaRight;
lapackf77_clacpy( "L", &n, &n, A, &lda, VR, &ldvr );
/* Generate unitary matrix in VR
* (CWorkspace: need 2*N-1, prefer N + (N-1)*NB)
* (RWorkspace: N)
* - including N reserved for gebal/gebak, unused by cunghr */
#if defined(Version1) || defined(Version2)
// Version 1 & 2 - LAPACK
lapackf77_cunghr( &n, &ilo, &ihi, VR, &ldvr, &work[itau],
&work[iwrk], &liwrk, &ierr );
#elif defined(Version3)
// Version 3 - LAPACK consistent MAGMA HRD + T matrices stored
magma_cunghr( n, ilo, ihi, VR, ldvr, &work[itau], dT, nb, &ierr );
#elif defined(Version5)
// Version 5 - Multi-GPU, T on host
magma_cunghr_m( n, ilo, ihi, VR, ldvr, &work[itau], T, nb, &ierr );
#endif
/* Perform QR iteration, accumulating Schur vectors in VR
* (CWorkspace: need 1, prefer HSWORK (see comments) )
* (RWorkspace: N)
* - including N reserved for gebal/gebak, unused by chseqr */
iwrk = itau;
liwrk = lwork - iwrk;
lapackf77_chseqr( "S", "V", &n, &ilo, &ihi, A, &lda, w,
VR, &ldvr, &work[iwrk], &liwrk, info );
}
else {
/* Compute eigenvalues only
* (CWorkspace: need 1, prefer HSWORK (see comments) )
* (RWorkspace: N)
* - including N reserved for gebal/gebak, unused by chseqr */
iwrk = itau;
liwrk = lwork - iwrk;
lapackf77_chseqr( "E", "N", &n, &ilo, &ihi, A, &lda, w,
VR, &ldvr, &work[iwrk], &liwrk, info );
}
/* If INFO > 0 from CHSEQR, then quit */
if (*info > 0) {
goto CLEANUP;
}
if (wantvl || wantvr) {
/* Compute left and/or right eigenvectors
* (CWorkspace: need 2*N)
* (RWorkspace: need 2*N)
* - including N reserved for gebal/gebak, unused by ctrevc */
irwork = ibal + n;
#if TREVC_VERSION == 1
lapackf77_ctrevc( lapack_side_const(side), "B", select, &n, A, &lda, VL, &ldvl,
VR, &ldvr, &n, &nout, &work[iwrk], &rwork[irwork], &ierr );
#elif TREVC_VERSION == 2
liwrk = lwork - iwrk;
lapackf77_ctrevc3( lapack_side_const(side), "B", select, &n, A, &lda, VL, &ldvl,
VR, &ldvr, &n, &nout, &work[iwrk], &liwrk, &rwork[irwork], &ierr );
#elif TREVC_VERSION == 3
magma_ctrevc3( side, MagmaBacktransVec, select, n, A, lda, VL, ldvl,
VR, ldvr, n, &nout, &work[iwrk], liwrk, &rwork[irwork], &ierr );
#elif TREVC_VERSION == 4
magma_ctrevc3_mt( side, MagmaBacktransVec, select, n, A, lda, VL, ldvl,
VR, ldvr, n, &nout, &work[iwrk], liwrk, &rwork[irwork], &ierr );
#elif TREVC_VERSION == 5
magma_ctrevc3_mt_gpu( side, MagmaBacktransVec, select, n, A, lda, VL, ldvl,
VR, ldvr, n, &nout, &work[iwrk], liwrk, &rwork[irwork], &ierr );
#else
#error Unknown TREVC_VERSION
#endif
}
if (wantvl) {
/* Undo balancing of left eigenvectors
* (CWorkspace: none)
* (RWorkspace: need N) */
lapackf77_cgebak( "B", "L", &n, &ilo, &ihi, &rwork[ibal], &n,
VL, &ldvl, &ierr );
/* Normalize left eigenvectors and make largest component real */
for (i = 0; i < n; ++i) {
scl = 1. / magma_cblas_scnrm2( n, VL(0,i), 1 );
blasf77_csscal( &n, &scl, VL(0,i), &ione );
for (k = 0; k < n; ++k) {
/* Computing 2nd power */
d__1 = MAGMA_C_REAL( *VL(k,i) );
d__2 = MAGMA_C_IMAG( *VL(k,i) );
rwork[irwork + k] = d__1*d__1 + d__2*d__2;
}
k = blasf77_isamax( &n, &rwork[irwork], &ione ) - 1; // subtract 1; k is 0-based
tmp = MAGMA_C_CONJ( *VL(k,i) ) / magma_ssqrt( rwork[irwork + k] );
blasf77_cscal( &n, &tmp, VL(0,i), &ione );
*VL(k,i) = MAGMA_C_MAKE( MAGMA_C_REAL( *VL(k,i) ), 0 );
}
}
if (wantvr) {
/* Undo balancing of right eigenvectors
* (CWorkspace: none)
* (RWorkspace: need N) */
lapackf77_cgebak( "B", "R", &n, &ilo, &ihi, &rwork[ibal], &n,
VR, &ldvr, &ierr );
/* Normalize right eigenvectors and make largest component real */
for (i = 0; i < n; ++i) {
scl = 1. / magma_cblas_scnrm2( n, VR(0,i), 1 );
blasf77_csscal( &n, &scl, VR(0,i), &ione );
for (k = 0; k < n; ++k) {
/* Computing 2nd power */
d__1 = MAGMA_C_REAL( *VR(k,i) );
d__2 = MAGMA_C_IMAG( *VR(k,i) );
rwork[irwork + k] = d__1*d__1 + d__2*d__2;
}
k = blasf77_isamax( &n, &rwork[irwork], &ione ) - 1; // subtract 1; k is 0-based
tmp = MAGMA_C_CONJ( *VR(k,i) ) / magma_ssqrt( rwork[irwork + k] );
blasf77_cscal( &n, &tmp, VR(0,i), &ione );
*VR(k,i) = MAGMA_C_MAKE( MAGMA_C_REAL( *VR(k,i) ), 0 );
}
}
CLEANUP:
/* Undo scaling if necessary */
if (scalea) {
// converged eigenvalues, stored in WR[i+1:n] and WI[i+1:n] for i = INFO
magma_int_t nval = n - (*info);
magma_int_t ld = max( nval, 1 );
lapackf77_clascl( "G", &izero, &izero, &cscale, &anrm, &nval, &ione, w + (*info), &ld, &ierr );
if (*info > 0) {
// first ilo columns were already upper triangular,
// so the corresponding eigenvalues are also valid.
nval = ilo - 1;
lapackf77_clascl( "G", &izero, &izero, &cscale, &anrm, &nval, &ione, w, &n, &ierr );
}
}
#if defined(Version3)
magma_free( dT );
#endif
#if defined(Version5)
magma_free_cpu( T );
#endif
work[0] = magma_cmake_lwork( minwrk ); // TODO use optwrk as in dgeev
return *info;
} /* magma_cgeev */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dormqr_gpu
*/
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
double Cnorm, error, work[1];
double c_neg_one = MAGMA_D_NEG_ONE;
magma_int_t ione = 1;
magma_int_t mm, m, n, k, size, info;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t nb, ldc, lda, lwork, lwork_max, dt_size;
double *C, *R, *A, *hwork, *tau;
magmaDouble_ptr dC, dA, dT;
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
// need slightly looser bound (60*eps instead of 30*eps) for some tests
opts.tolerance = max( 60., opts.tolerance );
double tol = opts.tolerance * lapackf77_dlamch("E");
// test all combinations of input parameters
magma_side_t side [] = { MagmaLeft, MagmaRight };
magma_trans_t trans[] = { MagmaTrans, MagmaNoTrans };
printf("%% M N K side trans CPU Gflop/s (sec) GPU Gflop/s (sec) ||R||_F / ||QC||_F\n");
printf("%%==============================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iside = 0; iside < 2; ++iside ) {
for( int itran = 0; itran < 2; ++itran ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
m = opts.msize[itest];
n = opts.nsize[itest];
k = opts.ksize[itest];
ldc = magma_roundup( m, opts.align ); // multiple of 32 by default
// A is m x k (left) or n x k (right)
mm = (side[iside] == MagmaLeft ? m : n);
nb = magma_get_dgeqrf_nb( mm, k );
lda = magma_roundup( mm, opts.align ); // multiple of 32 by default
gflops = FLOPS_DORMQR( m, n, k, side[iside] ) / 1e9;
if ( side[iside] == MagmaLeft && m < k ) {
printf( "%5d %5d %5d %4c %5c skipping because side=left and m < k\n",
(int) m, (int) n, (int) k,
lapacke_side_const( side[iside] ),
lapacke_trans_const( trans[itran] ) );
continue;
}
if ( side[iside] == MagmaRight && n < k ) {
printf( "%5d %5d %5d %4c %5c skipping because side=right and n < k\n",
(int) m, (int) n, (int) k,
lapacke_side_const( side[iside] ),
lapacke_trans_const( trans[itran] ) );
continue;
}
if ( side[iside] == MagmaLeft ) {
// side = left
lwork_max = (m - k + nb)*(n + nb) + n*nb;
dt_size = ( 2*min(m,k) + magma_roundup( max(m,n), 32) )*nb;
}
else {
// side = right
lwork_max = (n - k + nb)*(m + nb) + m*nb;
dt_size = ( 2*min(n,k) + magma_roundup( max(m,n), 32 ) )*nb;
}
// this rounds it up slightly if needed to agree with lwork query below
lwork_max = int( real( magma_dmake_lwork( lwork_max )));
TESTING_MALLOC_CPU( C, double, ldc*n );
TESTING_MALLOC_CPU( R, double, ldc*n );
TESTING_MALLOC_CPU( A, double, lda*k );
TESTING_MALLOC_CPU( hwork, double, lwork_max );
TESTING_MALLOC_CPU( tau, double, k );
TESTING_MALLOC_DEV( dC, double, ldc*n );
TESTING_MALLOC_DEV( dA, double, lda*k );
TESTING_MALLOC_DEV( dT, double, dt_size );
// C is full, m x n
size = ldc*n;
lapackf77_dlarnv( &ione, ISEED, &size, C );
magma_dsetmatrix( m, n, C, ldc, dC, ldc );
// A is m x k (left) or n x k (right)
size = lda*k;
lapackf77_dlarnv( &ione, ISEED, &size, A );
// compute QR factorization to get Householder vectors in dA, tau, dT
magma_dsetmatrix( mm, k, A, lda, dA, lda );
magma_dgeqrf_gpu( mm, k, dA, lda, tau, dT, &info );
magma_dgetmatrix( mm, k, dA, lda, A, lda );
if (info != 0) {
printf("magma_dgeqrf_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_dormqr( lapack_side_const( side[iside] ), lapack_trans_const( trans[itran] ),
&m, &n, &k,
A, &lda, tau, C, &ldc, hwork, &lwork_max, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0) {
printf("lapackf77_dormqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
// query for workspace size
lwork = -1;
magma_dormqr_gpu( side[iside], trans[itran],
m, n, k,
dA, lda, tau, dC, ldc, hwork, lwork, dT, nb, &info );
if (info != 0) {
printf("magma_dormqr_gpu (lwork query) returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
lwork = (magma_int_t) MAGMA_D_REAL( hwork[0] );
if ( lwork < 0 || lwork > lwork_max ) {
printf("Warning: optimal lwork %d > allocated lwork_max %d\n", (int) lwork, (int) lwork_max );
lwork = lwork_max;
}
// dormqr2 takes a copy of dA in CPU memory
if ( opts.version == 2 ) {
magma_dgetmatrix( mm, k, dA, lda, A, lda );
}
magmablasSetKernelStream( opts.queue );
gpu_time = magma_sync_wtime( opts.queue ); // sync needed for L,N and R,T cases
if ( opts.version == 1 ) {
magma_dormqr_gpu( side[iside], trans[itran],
m, n, k,
dA, lda, tau, dC, ldc, hwork, lwork, dT, nb, &info );
}
else if ( opts.version == 2 ) {
magma_dormqr2_gpu( side[iside], trans[itran],
m, n, k,
dA, lda, tau, dC, ldc, A, lda, &info );
}
gpu_time = magma_sync_wtime( opts.queue ) - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0) {
printf("magma_dormqr_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
magma_dgetmatrix( m, n, dC, ldc, R, ldc );
/* =====================================================================
compute relative error |QC_magma - QC_lapack| / |QC_lapack|
=================================================================== */
size = ldc*n;
blasf77_daxpy( &size, &c_neg_one, C, &ione, R, &ione );
Cnorm = lapackf77_dlange( "Fro", &m, &n, C, &ldc, work );
error = lapackf77_dlange( "Fro", &m, &n, R, &ldc, work ) / (magma_dsqrt(m*n) * Cnorm);
printf( "%5d %5d %5d %4c %5c %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) m, (int) n, (int) k,
lapacke_side_const( side[iside] ),
lapacke_trans_const( trans[itran] ),
cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed") );
status += ! (error < tol);
TESTING_FREE_CPU( C );
TESTING_FREE_CPU( R );
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( hwork );
TESTING_FREE_CPU( tau );
TESTING_FREE_DEV( dC );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dT );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}} // end iside, itran
printf( "\n" );
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/**
Purpose
-------
ZUNMQR overwrites the general complex M-by-N matrix C with
@verbatim
SIDE = MagmaLeft SIDE = MagmaRight
TRANS = MagmaNoTrans: Q * C C * Q
TRANS = Magma_ConjTrans: Q**H * C C * Q**H
@endverbatim
where Q is a complex unitary matrix defined as the product of k
elementary reflectors
Q = H(1) H(2) . . . H(k)
as returned by ZGEQRF. Q is of order M if SIDE = MagmaLeft and of order N
if SIDE = MagmaRight.
Arguments
---------
@param[in]
ngpu INTEGER
Number of GPUs to use. ngpu > 0.
@param[in]
side magma_side_t
- = MagmaLeft: apply Q or Q**H from the Left;
- = MagmaRight: apply Q or Q**H from the Right.
@param[in]
trans magma_trans_t
- = MagmaNoTrans: No transpose, apply Q;
- = Magma_ConjTrans: Conjugate transpose, apply Q**H.
@param[in]
m INTEGER
The number of rows of the matrix C. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix C. N >= 0.
@param[in]
k INTEGER
The number of elementary reflectors whose product defines
the matrix Q.
If SIDE = MagmaLeft, M >= K >= 0;
if SIDE = MagmaRight, N >= K >= 0.
@param[in]
A COMPLEX_16 array, dimension (LDA,K)
The i-th column must contain the vector which defines the
elementary reflector H(i), for i = 1,2,...,k, as returned by
ZGEQRF in the first k columns of its array argument A.
@param[in]
lda INTEGER
The leading dimension of the array A.
If SIDE = MagmaLeft, LDA >= max(1,M);
if SIDE = MagmaRight, LDA >= max(1,N).
@param[in]
tau COMPLEX_16 array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by ZGEQRF.
@param[in,out]
C COMPLEX_16 array, dimension (LDC,N)
On entry, the M-by-N matrix C.
On exit, C is overwritten by Q*C or Q**H*C or C*Q**H or C*Q.
@param[in]
ldc INTEGER
The leading dimension of the array C. LDC >= max(1,M).
@param[out]
work (workspace) COMPLEX_16 array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK.
If SIDE = MagmaLeft, LWORK >= max(1,N);
if SIDE = MagmaRight, LWORK >= max(1,M).
For optimum performance LWORK >= N*NB if SIDE = MagmaLeft, and
LWORK >= M*NB if SIDE = MagmaRight, where NB is the optimal
blocksize.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_zgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_zunmqr_m(
magma_int_t ngpu,
magma_side_t side, magma_trans_t trans,
magma_int_t m, magma_int_t n, magma_int_t k,
magmaDoubleComplex *A, magma_int_t lda,
magmaDoubleComplex *tau,
magmaDoubleComplex *C, magma_int_t ldc,
magmaDoubleComplex *work, magma_int_t lwork,
magma_int_t *info)
{
#define A(i, j) (A + (j)*lda + (i))
#define C(i, j) (C + (j)*ldc + (i))
#define dC(gpui, i, j) (dw[gpui] + (j)*lddc + (i))
#define dA_c(gpui, ind, i, j) (dw[gpui] + maxnlocal*lddc + (ind)*lddar*lddac + (i) + (j)*lddac)
#define dA_r(gpui, ind, i, j) (dw[gpui] + maxnlocal*lddc + (ind)*lddar*lddac + (i) + (j)*lddar)
#define dT(gpui, ind) (dw[gpui] + maxnlocal*lddc + 2*lddac*lddar + (ind)*((nb+1)*nb))
#define dwork(gpui, ind) (dw[gpui] + maxnlocal*lddc + 2*lddac*lddar + 2*((nb+1)*nb) + (ind)*(lddwork*nb))
/* Constants */
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
/* Local variables */
const char* side_ = lapack_side_const( side );
const char* trans_ = lapack_trans_const( trans );
magma_int_t nb = 128;
magmaDoubleComplex *T = NULL;
magmaDoubleComplex_ptr dw[MagmaMaxGPUs] = { NULL };
magma_queue_t queues[MagmaMaxGPUs][2] = {{ NULL }};
magma_event_t events[MagmaMaxGPUs][2] = {{ NULL }};
magma_int_t ind_c;
magma_device_t dev;
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
*info = 0;
magma_int_t left = (side == MagmaLeft);
magma_int_t notran = (trans == MagmaNoTrans);
magma_int_t lquery = (lwork == -1);
/* NQ is the order of Q and NW is the minimum dimension of WORK */
magma_int_t nq, nw;
if (left) {
nq = m;
nw = n;
} else {
nq = n;
nw = m;
}
if (! left && side != MagmaRight) {
*info = -1;
} else if (! notran && trans != Magma_ConjTrans) {
*info = -2;
} else if (m < 0) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (k < 0 || k > nq) {
*info = -5;
} else if (lda < max(1,nq)) {
*info = -7;
} else if (ldc < max(1,m)) {
*info = -10;
} else if (lwork < max(1,nw) && ! lquery) {
*info = -12;
}
magma_int_t lwkopt = max(1,nw) * nb;
if (*info == 0) {
work[0] = magma_zmake_lwork( lwkopt );
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0 || k == 0) {
work[0] = c_one;
return *info;
}
if (nb >= k) {
/* Use CPU code */
lapackf77_zunmqr(side_, trans_, &m, &n, &k, A, &lda, tau,
C, &ldc, work, &lwork, info);
return *info;
}
magma_int_t lddc = magma_roundup( m, 64 ); // TODO why 64 instead of 32 ?
magma_int_t lddac = nq;
magma_int_t lddar = nb;
magma_int_t lddwork = nw;
magma_int_t nlocal[ MagmaMaxGPUs ] = { 0 };
magma_int_t nb_l=256;
magma_int_t nbl = magma_ceildiv( n, nb_l ); // number of blocks
magma_int_t maxnlocal = magma_ceildiv( nbl, ngpu )*nb_l;
ngpu = min( ngpu, magma_ceildiv( n, nb_l )); // Don't use GPU that will not have data.
magma_int_t ldw = maxnlocal*lddc // dC
+ 2*lddac*lddar // 2*dA
+ 2*(nb + 1 + lddwork)*nb; // 2*(dT and dwork)
if (MAGMA_SUCCESS != magma_zmalloc_pinned( &T, nb*nb )) {
*info = MAGMA_ERR_HOST_ALLOC;
goto cleanup;
}
for (dev = 0; dev < ngpu; ++dev) {
magma_setdevice( dev );
if (MAGMA_SUCCESS != magma_zmalloc( &dw[dev], ldw )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
goto cleanup;
}
magma_queue_create( dev, &queues[dev][0] );
magma_queue_create( dev, &queues[dev][1] );
magma_event_create( &events[dev][0] );
magma_event_create( &events[dev][1] );
}
/* Use hybrid CPU-MGPU code */
if (left) {
//copy C to mgpus
for (magma_int_t i = 0; i < nbl; ++i) {
dev = i % ngpu;
magma_setdevice( dev );
magma_int_t kb = min(nb_l, n-i*nb_l);
magma_zsetmatrix_async( m, kb,
C(0, i*nb_l), ldc,
dC(dev, 0, i/ngpu*nb_l), lddc, queues[dev][0] );
nlocal[dev] += kb;
}
magma_int_t i1, i2, i3;
if ( !notran ) {
i1 = 0;
i2 = k;
i3 = nb;
} else {
i1 = (k - 1) / nb * nb;
i2 = 0;
i3 = -nb;
}
ind_c = 0;
for (magma_int_t i = i1; (i3 < 0 ? i >= i2 : i < i2); i += i3) {
// start the copy of A panel
magma_int_t kb = min(nb, k - i);
for (dev = 0; dev < ngpu; ++dev) {
magma_setdevice( dev );
magma_event_sync( events[dev][ind_c] ); // check if the new data can be copied
magma_zsetmatrix_async(nq-i, kb,
A(i, i), lda,
dA_c(dev, ind_c, i, 0), lddac, queues[dev][0] );
// set upper triangular part of dA to identity
magmablas_zlaset_band( MagmaUpper, kb, kb, kb, c_zero, c_one, dA_c(dev, ind_c, i, 0), lddac, queues[dev][0] );
}
/* Form the triangular factor of the block reflector
H = H(i) H(i+1) . . . H(i+ib-1) */
magma_int_t nqi = nq - i;
lapackf77_zlarft("F", "C", &nqi, &kb, A(i, i), &lda,
&tau[i], T, &kb);
/* H or H' is applied to C(1:m,i:n) */
/* Apply H or H'; First copy T to the GPU */
for (dev = 0; dev < ngpu; ++dev) {
magma_setdevice( dev );
magma_zsetmatrix_async(kb, kb,
T, kb,
dT(dev, ind_c), kb, queues[dev][0] );
}
for (dev = 0; dev < ngpu; ++dev) {
magma_setdevice( dev );
magma_queue_sync( queues[dev][0] ); // check if the data was copied
magma_zlarfb_gpu( side, trans, MagmaForward, MagmaColumnwise,
m-i, nlocal[dev], kb,
dA_c(dev, ind_c, i, 0), lddac, dT(dev, ind_c), kb,
dC(dev, i, 0), lddc,
dwork(dev, ind_c), lddwork, queues[dev][1] );
magma_event_record(events[dev][ind_c], queues[dev][1] );
}
ind_c = (ind_c+1)%2;
}
for (dev = 0; dev < ngpu; ++dev) {
magma_setdevice( dev );
magma_queue_sync( queues[dev][1] );
}
//copy C from mgpus
for (magma_int_t i = 0; i < nbl; ++i) {
dev = i % ngpu;
magma_setdevice( dev );
magma_int_t kb = min(nb_l, n-i*nb_l);
magma_zgetmatrix( m, kb,
dC(dev, 0, i/ngpu*nb_l), lddc,
C(0, i*nb_l), ldc, queues[dev][1] );
// magma_zgetmatrix_async( m, kb,
// dC(dev, 0, i/ngpu*nb_l), lddc,
// C(0, i*nb_l), ldc, queues[dev][0] );
}
} else {
*info = MAGMA_ERR_NOT_IMPLEMENTED;
magma_xerbla( __func__, -(*info) );
goto cleanup;
/*
if ( notran ) {
i1 = 0;
i2 = k;
i3 = nb;
} else {
i1 = (k - 1) / nb * nb;
i2 = 0;
i3 = -nb;
}
mi = m;
ic = 0;
for (i = i1; (i3 < 0 ? i >= i2 : i < i2); i += i3) {
ib = min(nb, k - i);
// Form the triangular factor of the block reflector
// H = H(i) H(i+1) . . . H(i+ib-1)
i__4 = nq - i;
lapackf77_zlarft("F", "C", &i__4, &ib, A(i, i), &lda,
&tau[i], T, &ib);
// 1) copy the panel from A to the GPU, and
// 2) set upper triangular part of dA to identity
magma_zsetmatrix( i__4, ib, A(i, i), lda, dA(i, 0), ldda, queues[dev][1] );
magmablas_zlaset_band( MagmaUpper, ib, ib, ib, c_zero, c_one, dA(i, 0), ldda, queues[dev][1] );
// H or H' is applied to C(1:m,i:n)
ni = n - i;
jc = i;
// Apply H or H'; First copy T to the GPU
magma_zsetmatrix( ib, ib, T, ib, dT, ib, queues[dev][1] );
magma_zlarfb_gpu( side, trans, MagmaForward, MagmaColumnwise,
mi, ni, ib,
dA(i, 0), ldda, dT, ib,
dC(ic, jc), lddc,
dwork, lddwork, queues[dev][1] );
}
*/
}
cleanup:
work[0] = magma_zmake_lwork( lwkopt );
for (dev = 0; dev < ngpu; ++dev) {
magma_setdevice( dev );
magma_event_destroy( events[dev][0] );
magma_event_destroy( events[dev][1] );
magma_queue_destroy( queues[dev][0] );
magma_queue_destroy( queues[dev][1] );
magma_free( dw[dev] );
}
magma_setdevice( orig_dev );
magma_free_pinned( T );
return *info;
} /* magma_zunmqr */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dormql
*/
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
double error, work[1];
double c_neg_one = MAGMA_D_NEG_ONE;
magma_int_t ione = 1;
magma_int_t mm, m, n, k, size, info;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t nb, ldc, lda, lwork, lwork_max;
double *C, *R, *A, *W, *tau;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
// need slightly looser bound (60*eps instead of 30*eps) for some tests
opts.tolerance = max( 60., opts.tolerance );
double tol = opts.tolerance * lapackf77_dlamch("E");
// test all combinations of input parameters
magma_side_t side [] = { MagmaLeft, MagmaRight };
magma_trans_t trans[] = { MagmaTrans, MagmaNoTrans };
printf(" M N K side trans CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / ||QC||_F\n");
printf("===============================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iside = 0; iside < 2; ++iside ) {
for( int itran = 0; itran < 2; ++itran ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
m = opts.msize[itest];
n = opts.nsize[itest];
k = opts.ksize[itest];
nb = magma_get_dgeqlf_nb( m );
ldc = m;
// A is m x k (left) or n x k (right)
mm = (side[iside] == MagmaLeft ? m : n);
lda = mm;
gflops = FLOPS_DORMQL( m, n, k, side[iside] ) / 1e9;
if ( side[iside] == MagmaLeft && m < k ) {
printf( "%5d %5d %5d %4c %5c skipping because side=left and m < k\n",
(int) m, (int) n, (int) k,
lapacke_side_const( side[iside] ),
lapacke_trans_const( trans[itran] ) );
continue;
}
if ( side[iside] == MagmaRight && n < k ) {
printf( "%5d %5d %5d %4c %5c skipping because side=right and n < k\n",
(int) m, (int) n, (int) k,
lapacke_side_const( side[iside] ),
lapacke_trans_const( trans[itran] ) );
continue;
}
// need at least 2*nb*nb for geqlf
lwork_max = max( max( m*nb, n*nb ), 2*nb*nb );
TESTING_MALLOC_CPU( C, double, ldc*n );
TESTING_MALLOC_CPU( R, double, ldc*n );
TESTING_MALLOC_CPU( A, double, lda*k );
TESTING_MALLOC_CPU( W, double, lwork_max );
TESTING_MALLOC_CPU( tau, double, k );
// C is full, m x n
size = ldc*n;
lapackf77_dlarnv( &ione, ISEED, &size, C );
lapackf77_dlacpy( "Full", &m, &n, C, &ldc, R, &ldc );
size = lda*k;
lapackf77_dlarnv( &ione, ISEED, &size, A );
// compute QL factorization to get Householder vectors in A, tau
magma_dgeqlf( mm, k, A, lda, tau, W, lwork_max, &info );
if (info != 0)
printf("magma_dgeqlf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_dormql( lapack_side_const( side[iside] ), lapack_trans_const( trans[itran] ),
&m, &n, &k,
A, &lda, tau, C, &ldc, W, &lwork_max, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_dormql returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
// query for workspace size
lwork = -1;
magma_dormql( side[iside], trans[itran],
m, n, k,
A, lda, tau, R, ldc, W, lwork, &info );
if (info != 0)
printf("magma_dormql (lwork query) returned error %d: %s.\n",
(int) info, magma_strerror( info ));
lwork = (magma_int_t) MAGMA_D_REAL( W[0] );
if ( lwork < 0 || lwork > lwork_max ) {
printf("optimal lwork %d > lwork_max %d\n", (int) lwork, (int) lwork_max );
lwork = lwork_max;
}
gpu_time = magma_wtime();
magma_dormql( side[iside], trans[itran],
m, n, k,
A, lda, tau, R, ldc, W, lwork, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_dormql returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
compute relative error |QC_magma - QC_lapack| / |QC_lapack|
=================================================================== */
error = lapackf77_dlange( "Fro", &m, &n, C, &ldc, work );
size = ldc*n;
blasf77_daxpy( &size, &c_neg_one, C, &ione, R, &ione );
error = lapackf77_dlange( "Fro", &m, &n, R, &ldc, work ) / error;
printf( "%5d %5d %5d %4c %5c %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) m, (int) n, (int) k,
lapacke_side_const( side[iside] ),
lapacke_trans_const( trans[itran] ),
cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed") );
status += ! (error < tol);
TESTING_FREE_CPU( C );
TESTING_FREE_CPU( R );
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( W );
TESTING_FREE_CPU( tau );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}} // end iside, itran
printf( "\n" );
}
TESTING_FINALIZE();
return status;
}
/**
Purpose
-------
ZGEEV computes for an N-by-N complex nonsymmetric matrix A, the
eigenvalues and, optionally, the left and/or right eigenvectors.
The right eigenvector v(j) of A satisfies
A * v(j) = lambda(j) * v(j)
where lambda(j) is its eigenvalue.
The left eigenvector u(j) of A satisfies
u(j)**H * A = lambda(j) * u(j)**H
where u(j)**H denotes the conjugate transpose of u(j).
The computed eigenvectors are normalized to have Euclidean norm
equal to 1 and largest component real.
Arguments
---------
@param[in]
jobvl magma_vec_t
- = MagmaNoVec: left eigenvectors of A are not computed;
- = MagmaVec: left eigenvectors of are computed.
@param[in]
jobvr magma_vec_t
- = MagmaNoVec: right eigenvectors of A are not computed;
- = MagmaVec: right eigenvectors of A are computed.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
A COMPLEX_16 array, dimension (LDA,N)
On entry, the N-by-N matrix A.
On exit, A has been overwritten.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
w COMPLEX_16 array, dimension (N)
w contains the computed eigenvalues.
@param[out]
VL COMPLEX_16 array, dimension (LDVL,N)
If JOBVL = MagmaVec, the left eigenvectors u(j) are stored one
after another in the columns of VL, in the same order
as their eigenvalues.
If JOBVL = MagmaNoVec, VL is not referenced.
u(j) = VL(:,j), the j-th column of VL.
@param[in]
ldvl INTEGER
The leading dimension of the array VL. LDVL >= 1; if
JOBVL = MagmaVec, LDVL >= N.
@param[out]
VR COMPLEX_16 array, dimension (LDVR,N)
If JOBVR = MagmaVec, the right eigenvectors v(j) are stored one
after another in the columns of VR, in the same order
as their eigenvalues.
If JOBVR = MagmaNoVec, VR is not referenced.
v(j) = VR(:,j), the j-th column of VR.
@param[in]
ldvr INTEGER
The leading dimension of the array VR. LDVR >= 1; if
JOBVR = MagmaVec, LDVR >= N.
@param[out]
work (workspace) COMPLEX_16 array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK. LWORK >= (1+nb)*N.
For optimal performance, LWORK >= (1+2*nb)*N.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param
rwork (workspace) DOUBLE PRECISION array, dimension (2*N)
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value.
- > 0: if INFO = i, the QR algorithm failed to compute all the
eigenvalues, and no eigenvectors have been computed;
elements and i+1:N of w contain eigenvalues which have
converged.
@ingroup magma_zgeev_driver
********************************************************************/
extern "C" magma_int_t
magma_zgeev(
magma_vec_t jobvl, magma_vec_t jobvr, magma_int_t n,
magmaDoubleComplex *A, magma_int_t lda,
#ifdef COMPLEX
magmaDoubleComplex *w,
#else
double *wr, double *wi,
#endif
magmaDoubleComplex *VL, magma_int_t ldvl,
magmaDoubleComplex *VR, magma_int_t ldvr,
magmaDoubleComplex *work, magma_int_t lwork,
#ifdef COMPLEX
double *rwork,
#endif
magma_int_t *info )
{
#define VL(i,j) (VL + (i) + (j)*ldvl)
#define VR(i,j) (VR + (i) + (j)*ldvr)
const magma_int_t ione = 1;
const magma_int_t izero = 0;
double d__1, d__2;
magmaDoubleComplex tmp;
double scl;
double dum[1], eps;
double anrm, cscale, bignum, smlnum;
magma_int_t i, k, ilo, ihi;
magma_int_t ibal, ierr, itau, iwrk, nout, liwrk, nb;
magma_int_t scalea, minwrk, optwrk, irwork, lquery, wantvl, wantvr, select[1];
magma_side_t side = MagmaRight;
magma_timer_t time_total=0, time_gehrd=0, time_unghr=0, time_hseqr=0, time_trevc=0, time_sum=0;
magma_flops_t flop_total=0, flop_gehrd=0, flop_unghr=0, flop_hseqr=0, flop_trevc=0, flop_sum=0;
timer_start( time_total );
flops_start( flop_total );
irwork = 0;
*info = 0;
lquery = (lwork == -1);
wantvl = (jobvl == MagmaVec);
wantvr = (jobvr == MagmaVec);
if (! wantvl && jobvl != MagmaNoVec) {
*info = -1;
} else if (! wantvr && jobvr != MagmaNoVec) {
*info = -2;
} else if (n < 0) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
} else if ( (ldvl < 1) || (wantvl && (ldvl < n))) {
*info = -8;
} else if ( (ldvr < 1) || (wantvr && (ldvr < n))) {
*info = -10;
}
/* Compute workspace */
nb = magma_get_zgehrd_nb( n );
if (*info == 0) {
minwrk = (1+ nb)*n;
optwrk = (1+2*nb)*n;
work[0] = MAGMA_Z_MAKE( optwrk, 0 );
if (lwork < minwrk && ! lquery) {
*info = -12;
}
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (n == 0) {
return *info;
}
#if defined(VERSION3)
magmaDoubleComplex_ptr dT;
if (MAGMA_SUCCESS != magma_zmalloc( &dT, nb*n )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
#endif
/* Get machine constants */
eps = lapackf77_dlamch( "P" );
smlnum = lapackf77_dlamch( "S" );
bignum = 1. / smlnum;
lapackf77_dlabad( &smlnum, &bignum );
smlnum = magma_dsqrt( smlnum ) / eps;
bignum = 1. / smlnum;
/* Scale A if max element outside range [SMLNUM,BIGNUM] */
anrm = lapackf77_zlange( "M", &n, &n, A, &lda, dum );
scalea = 0;
if (anrm > 0. && anrm < smlnum) {
scalea = 1;
cscale = smlnum;
} else if (anrm > bignum) {
scalea = 1;
cscale = bignum;
}
if (scalea) {
lapackf77_zlascl( "G", &izero, &izero, &anrm, &cscale, &n, &n, A, &lda, &ierr );
}
/* Balance the matrix
* (CWorkspace: none)
* (RWorkspace: need N)
* - this space is reserved until after gebak */
ibal = 0;
lapackf77_zgebal( "B", &n, A, &lda, &ilo, &ihi, &rwork[ibal], &ierr );
/* Reduce to upper Hessenberg form
* (CWorkspace: need 2*N, prefer N + N*NB)
* (RWorkspace: N)
* - including N reserved for gebal/gebak, unused by zgehrd */
itau = 0;
iwrk = itau + n;
liwrk = lwork - iwrk;
timer_start( time_gehrd );
flops_start( flop_gehrd );
#if defined(VERSION1)
// Version 1 - LAPACK
lapackf77_zgehrd( &n, &ilo, &ihi, A, &lda,
&work[itau], &work[iwrk], &liwrk, &ierr );
#elif defined(VERSION2)
// Version 2 - LAPACK consistent HRD
magma_zgehrd2( n, ilo, ihi, A, lda,
&work[itau], &work[iwrk], liwrk, &ierr );
#elif defined(VERSION3)
// Version 3 - LAPACK consistent MAGMA HRD + T matrices stored,
magma_zgehrd( n, ilo, ihi, A, lda,
&work[itau], &work[iwrk], liwrk, dT, &ierr );
#endif
time_sum += timer_stop( time_gehrd );
flop_sum += flops_stop( flop_gehrd );
if (wantvl) {
/* Want left eigenvectors
* Copy Householder vectors to VL */
side = MagmaLeft;
lapackf77_zlacpy( MagmaLowerStr, &n, &n, A, &lda, VL, &ldvl );
/* Generate unitary matrix in VL
* (CWorkspace: need 2*N-1, prefer N + (N-1)*NB)
* (RWorkspace: N)
* - including N reserved for gebal/gebak, unused by zunghr */
timer_start( time_unghr );
flops_start( flop_unghr );
#if defined(VERSION1) || defined(VERSION2)
// Version 1 & 2 - LAPACK
lapackf77_zunghr( &n, &ilo, &ihi, VL, &ldvl, &work[itau],
&work[iwrk], &liwrk, &ierr );
#elif defined(VERSION3)
// Version 3 - LAPACK consistent MAGMA HRD + T matrices stored
magma_zunghr( n, ilo, ihi, VL, ldvl, &work[itau], dT, nb, &ierr );
#endif
time_sum += timer_stop( time_unghr );
flop_sum += flops_stop( flop_unghr );
timer_start( time_hseqr );
flops_start( flop_hseqr );
/* Perform QR iteration, accumulating Schur vectors in VL
* (CWorkspace: need 1, prefer HSWORK (see comments) )
* (RWorkspace: N)
* - including N reserved for gebal/gebak, unused by zhseqr */
iwrk = itau;
liwrk = lwork - iwrk;
lapackf77_zhseqr( "S", "V", &n, &ilo, &ihi, A, &lda, w,
VL, &ldvl, &work[iwrk], &liwrk, info );
time_sum += timer_stop( time_hseqr );
flop_sum += flops_stop( flop_hseqr );
if (wantvr) {
/* Want left and right eigenvectors
* Copy Schur vectors to VR */
side = MagmaBothSides;
lapackf77_zlacpy( "F", &n, &n, VL, &ldvl, VR, &ldvr );
}
}
else if (wantvr) {
/* Want right eigenvectors
* Copy Householder vectors to VR */
side = MagmaRight;
lapackf77_zlacpy( "L", &n, &n, A, &lda, VR, &ldvr );
/* Generate unitary matrix in VR
* (CWorkspace: need 2*N-1, prefer N + (N-1)*NB)
* (RWorkspace: N)
* - including N reserved for gebal/gebak, unused by zunghr */
timer_start( time_unghr );
flops_start( flop_unghr );
#if defined(VERSION1) || defined(VERSION2)
// Version 1 & 2 - LAPACK
lapackf77_zunghr( &n, &ilo, &ihi, VR, &ldvr, &work[itau],
&work[iwrk], &liwrk, &ierr );
#elif defined(VERSION3)
// Version 3 - LAPACK consistent MAGMA HRD + T matrices stored
magma_zunghr( n, ilo, ihi, VR, ldvr, &work[itau], dT, nb, &ierr );
#endif
time_sum += timer_stop( time_unghr );
flop_sum += flops_stop( flop_unghr );
/* Perform QR iteration, accumulating Schur vectors in VR
* (CWorkspace: need 1, prefer HSWORK (see comments) )
* (RWorkspace: N)
* - including N reserved for gebal/gebak, unused by zhseqr */
timer_start( time_hseqr );
flops_start( flop_hseqr );
iwrk = itau;
liwrk = lwork - iwrk;
lapackf77_zhseqr( "S", "V", &n, &ilo, &ihi, A, &lda, w,
VR, &ldvr, &work[iwrk], &liwrk, info );
time_sum += timer_stop( time_hseqr );
flop_sum += flops_stop( flop_hseqr );
}
else {
/* Compute eigenvalues only
* (CWorkspace: need 1, prefer HSWORK (see comments) )
* (RWorkspace: N)
* - including N reserved for gebal/gebak, unused by zhseqr */
timer_start( time_hseqr );
flops_start( flop_hseqr );
iwrk = itau;
liwrk = lwork - iwrk;
lapackf77_zhseqr( "E", "N", &n, &ilo, &ihi, A, &lda, w,
VR, &ldvr, &work[iwrk], &liwrk, info );
time_sum += timer_stop( time_hseqr );
flop_sum += flops_stop( flop_hseqr );
}
/* If INFO > 0 from ZHSEQR, then quit */
if (*info > 0) {
goto CLEANUP;
}
timer_start( time_trevc );
flops_start( flop_trevc );
if (wantvl || wantvr) {
/* Compute left and/or right eigenvectors
* (CWorkspace: need 2*N)
* (RWorkspace: need 2*N)
* - including N reserved for gebal/gebak, unused by ztrevc */
irwork = ibal + n;
#if TREVC_VERSION == 1
lapackf77_ztrevc( lapack_side_const(side), "B", select, &n, A, &lda, VL, &ldvl,
VR, &ldvr, &n, &nout, &work[iwrk], &rwork[irwork], &ierr );
#elif TREVC_VERSION == 2
liwrk = lwork - iwrk;
lapackf77_ztrevc3( lapack_side_const(side), "B", select, &n, A, &lda, VL, &ldvl,
VR, &ldvr, &n, &nout, &work[iwrk], &liwrk, &rwork[irwork], &ierr );
#elif TREVC_VERSION == 3
magma_ztrevc3( side, MagmaBacktransVec, select, n, A, lda, VL, ldvl,
VR, ldvr, n, &nout, &work[iwrk], liwrk, &rwork[irwork], &ierr );
#elif TREVC_VERSION == 4
magma_ztrevc3_mt( side, MagmaBacktransVec, select, n, A, lda, VL, ldvl,
VR, ldvr, n, &nout, &work[iwrk], liwrk, &rwork[irwork], &ierr );
#elif TREVC_VERSION == 5
magma_ztrevc3_mt_gpu( side, MagmaBacktransVec, select, n, A, lda, VL, ldvl,
VR, ldvr, n, &nout, &work[iwrk], liwrk, &rwork[irwork], &ierr );
#else
#error Unknown TREVC_VERSION
#endif
}
time_sum += timer_stop( time_trevc );
flop_sum += flops_stop( flop_trevc );
if (wantvl) {
/* Undo balancing of left eigenvectors
* (CWorkspace: none)
* (RWorkspace: need N) */
lapackf77_zgebak( "B", "L", &n, &ilo, &ihi, &rwork[ibal], &n,
VL, &ldvl, &ierr );
/* Normalize left eigenvectors and make largest component real */
for (i = 0; i < n; ++i) {
scl = 1. / magma_cblas_dznrm2( n, VL(0,i), 1 );
blasf77_zdscal( &n, &scl, VL(0,i), &ione );
for (k = 0; k < n; ++k) {
/* Computing 2nd power */
d__1 = MAGMA_Z_REAL( *VL(k,i) );
d__2 = MAGMA_Z_IMAG( *VL(k,i) );
rwork[irwork + k] = d__1*d__1 + d__2*d__2;
}
k = blasf77_idamax( &n, &rwork[irwork], &ione ) - 1; // subtract 1; k is 0-based
tmp = MAGMA_Z_CNJG( *VL(k,i) ) / magma_dsqrt( rwork[irwork + k] );
blasf77_zscal( &n, &tmp, VL(0,i), &ione );
*VL(k,i) = MAGMA_Z_MAKE( MAGMA_Z_REAL( *VL(k,i) ), 0 );
}
}
if (wantvr) {
/* Undo balancing of right eigenvectors
* (CWorkspace: none)
* (RWorkspace: need N) */
lapackf77_zgebak( "B", "R", &n, &ilo, &ihi, &rwork[ibal], &n,
VR, &ldvr, &ierr );
/* Normalize right eigenvectors and make largest component real */
for (i = 0; i < n; ++i) {
scl = 1. / magma_cblas_dznrm2( n, VR(0,i), 1 );
blasf77_zdscal( &n, &scl, VR(0,i), &ione );
for (k = 0; k < n; ++k) {
/* Computing 2nd power */
d__1 = MAGMA_Z_REAL( *VR(k,i) );
d__2 = MAGMA_Z_IMAG( *VR(k,i) );
rwork[irwork + k] = d__1*d__1 + d__2*d__2;
}
k = blasf77_idamax( &n, &rwork[irwork], &ione ) - 1; // subtract 1; k is 0-based
tmp = MAGMA_Z_CNJG( *VR(k,i) ) / magma_dsqrt( rwork[irwork + k] );
blasf77_zscal( &n, &tmp, VR(0,i), &ione );
*VR(k,i) = MAGMA_Z_MAKE( MAGMA_Z_REAL( *VR(k,i) ), 0 );
}
}
CLEANUP:
/* Undo scaling if necessary */
if (scalea) {
// converged eigenvalues, stored in WR[i+1:n] and WI[i+1:n] for i = INFO
magma_int_t nval = n - (*info);
magma_int_t ld = max( nval, 1 );
lapackf77_zlascl( "G", &izero, &izero, &cscale, &anrm, &nval, &ione, w + (*info), &ld, &ierr );
if (*info > 0) {
// first ilo columns were already upper triangular,
// so the corresponding eigenvalues are also valid.
nval = ilo - 1;
lapackf77_zlascl( "G", &izero, &izero, &cscale, &anrm, &nval, &ione, w, &n, &ierr );
}
}
#if defined(VERSION3)
magma_free( dT );
#endif
timer_stop( time_total );
flops_stop( flop_total );
timer_printf( "dgeev times n %5d, gehrd %7.3f, unghr %7.3f, hseqr %7.3f, trevc %7.3f, total %7.3f, sum %7.3f\n",
(int) n, time_gehrd, time_unghr, time_hseqr, time_trevc, time_total, time_sum );
timer_printf( "dgeev flops n %5d, gehrd %7lld, unghr %7lld, hseqr %7lld, trevc %7lld, total %7lld, sum %7lld\n",
(int) n, flop_gehrd, flop_unghr, flop_hseqr, flop_trevc, flop_total, flop_sum );
work[0] = MAGMA_Z_MAKE( (double) optwrk, 0. );
return *info;
} /* magma_zgeev */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing sormbr
*/
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float Cnorm, error, dwork[1];
float c_neg_one = MAGMA_S_NEG_ONE;
magma_int_t ione = 1;
magma_int_t m, n, k, mi, ni, mm, nn, nq, size, info;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t nb, ldc, lda, lwork, lwork_max;
float *C, *R, *A, *work, *tau, *tauq, *taup;
float *d, *e;
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
// need slightly looser bound (60*eps instead of 30*eps) for some tests
opts.tolerance = max( 60., opts.tolerance );
float tol = opts.tolerance * lapackf77_slamch("E");
// test all combinations of input parameters
magma_vect_t vect [] = { MagmaQ, MagmaP };
magma_side_t side [] = { MagmaLeft, MagmaRight };
magma_trans_t trans[] = { MagmaTrans, MagmaNoTrans };
printf("%% M N K vect side trans CPU Gflop/s (sec) GPU Gflop/s (sec) ||R||_F / ||QC||_F\n");
printf("%%==============================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int ivect = 0; ivect < 2; ++ivect ) {
for( int iside = 0; iside < 2; ++iside ) {
for( int itran = 0; itran < 2; ++itran ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
m = opts.msize[itest];
n = opts.nsize[itest];
k = opts.ksize[itest];
nb = magma_get_sgebrd_nb( m, n );
ldc = m;
// A is nq x k (vect=Q) or k x nq (vect=P)
// where nq=m (left) or nq=n (right)
nq = (side[iside] == MagmaLeft ? m : n );
mm = (vect[ivect] == MagmaQ ? nq : k );
nn = (vect[ivect] == MagmaQ ? k : nq);
lda = mm;
// MBR calls either MQR or MLQ in various ways
if ( vect[ivect] == MagmaQ ) {
if ( nq >= k ) {
gflops = FLOPS_SORMQR( m, n, k, side[iside] ) / 1e9;
}
else {
if ( side[iside] == MagmaLeft ) {
mi = m - 1;
ni = n;
}
else {
mi = m;
ni = n - 1;
}
gflops = FLOPS_SORMQR( mi, ni, nq-1, side[iside] ) / 1e9;
}
}
else {
if ( nq > k ) {
gflops = FLOPS_SORMLQ( m, n, k, side[iside] ) / 1e9;
}
else {
if ( side[iside] == MagmaLeft ) {
mi = m - 1;
ni = n;
}
else {
mi = m;
ni = n - 1;
}
gflops = FLOPS_SORMLQ( mi, ni, nq-1, side[iside] ) / 1e9;
}
}
// workspace for gebrd is (mm + nn)*nb
// workspace for unmbr is m*nb or n*nb, depending on side
lwork_max = max( (mm + nn)*nb, max( m*nb, n*nb ));
// this rounds it up slightly if needed to agree with lwork query below
lwork_max = int( real( magma_smake_lwork( lwork_max )));
TESTING_MALLOC_CPU( C, float, ldc*n );
TESTING_MALLOC_CPU( R, float, ldc*n );
TESTING_MALLOC_CPU( A, float, lda*nn );
TESTING_MALLOC_CPU( work, float, lwork_max );
TESTING_MALLOC_CPU( d, float, min(mm,nn) );
TESTING_MALLOC_CPU( e, float, min(mm,nn) );
TESTING_MALLOC_CPU( tauq, float, min(mm,nn) );
TESTING_MALLOC_CPU( taup, float, min(mm,nn) );
// C is full, m x n
size = ldc*n;
lapackf77_slarnv( &ione, ISEED, &size, C );
lapackf77_slacpy( "Full", &m, &n, C, &ldc, R, &ldc );
size = lda*nn;
lapackf77_slarnv( &ione, ISEED, &size, A );
// compute BRD factorization to get Householder vectors in A, tauq, taup
//lapackf77_sgebrd( &mm, &nn, A, &lda, d, e, tauq, taup, work, &lwork_max, &info );
magma_sgebrd( mm, nn, A, lda, d, e, tauq, taup, work, lwork_max, &info );
if (info != 0) {
printf("magma_sgebrd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
if ( vect[ivect] == MagmaQ ) {
tau = tauq;
} else {
tau = taup;
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_sormbr( lapack_vect_const( vect[ivect] ),
lapack_side_const( side[iside] ),
lapack_trans_const( trans[itran] ),
&m, &n, &k,
A, &lda, tau, C, &ldc, work, &lwork_max, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0) {
printf("lapackf77_sormbr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
// query for workspace size
lwork = -1;
magma_sormbr( vect[ivect], side[iside], trans[itran],
m, n, k,
A, lda, tau, R, ldc, work, lwork, &info );
if (info != 0) {
printf("magma_sormbr (lwork query) returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
lwork = (magma_int_t) MAGMA_S_REAL( work[0] );
if ( lwork < 0 || lwork > lwork_max ) {
printf("Warning: optimal lwork %d > allocated lwork_max %d\n", (int) lwork, (int) lwork_max );
lwork = lwork_max;
}
gpu_time = magma_wtime();
magma_sormbr( vect[ivect], side[iside], trans[itran],
m, n, k,
A, lda, tau, R, ldc, work, lwork, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0) {
printf("magma_sormbr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
compute relative error |QC_magma - QC_lapack| / |QC_lapack|
=================================================================== */
size = ldc*n;
blasf77_saxpy( &size, &c_neg_one, C, &ione, R, &ione );
Cnorm = lapackf77_slange( "Fro", &m, &n, C, &ldc, dwork );
error = lapackf77_slange( "Fro", &m, &n, R, &ldc, dwork ) / (magma_ssqrt(m*n) * Cnorm);
printf( "%5d %5d %5d %c %4c %5c %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) m, (int) n, (int) k,
lapacke_vect_const( vect[ivect] ),
lapacke_side_const( side[iside] ),
lapacke_trans_const( trans[itran] ),
cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed") );
status += ! (error < tol);
TESTING_FREE_CPU( C );
TESTING_FREE_CPU( R );
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( work );
TESTING_FREE_CPU( d );
TESTING_FREE_CPU( e );
TESTING_FREE_CPU( taup );
TESTING_FREE_CPU( tauq );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}}} // end ivect, iside, itran
printf( "\n" );
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cunmqr_gpu
*/
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float error, work[1];
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t ione = 1;
magma_int_t m, n, k, size, info;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t nb, ldc, lda, lwork, lwork_max, dt_size;
magmaFloatComplex *C, *R, *A, *W, *tau;
magmaFloatComplex_ptr dC, dA, dT;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = 2. * opts.tolerance * lapackf77_slamch("E");
// test all combinations of input parameters
magma_side_t side [] = { MagmaLeft, MagmaRight };
magma_trans_t trans[] = { Magma_ConjTrans, MagmaNoTrans };
printf(" M N K side trans CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / ||QC||_F\n");
printf("===============================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iside = 0; iside < 2; ++iside ) {
for( int itran = 0; itran < 2; ++itran ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
m = opts.msize[itest];
n = opts.nsize[itest];
k = opts.ksize[itest];
nb = magma_get_cgeqrf_nb( m );
ldc = ((m + 31)/32)*32;
lda = ((max(m,n) + 31)/32)*32;
gflops = FLOPS_CUNMQR( m, n, k, side[iside] ) / 1e9;
if ( side[iside] == MagmaLeft && m < k ) {
printf( "%5d %5d %5d %4c %5c skipping because side=left and m < k\n",
(int) m, (int) n, (int) k,
lapacke_side_const( side[iside] ),
lapacke_trans_const( trans[itran] ) );
continue;
}
if ( side[iside] == MagmaRight && n < k ) {
printf( "%5d %5d %5d %4c %5c skipping because side=right and n < k\n",
(int) m, (int) n, (int) k,
lapacke_side_const( side[iside] ),
lapacke_trans_const( trans[itran] ) );
continue;
}
if ( side[iside] == MagmaLeft ) {
// side = left
lwork_max = (m - k + nb)*(n + nb) + n*nb;
dt_size = ( 2*min(m,k) + ((max(m,n) + 31)/32)*32 )*nb;
}
else {
// side = right
lwork_max = (n - k + nb)*(m + nb) + m*nb;
dt_size = ( 2*min(n,k) + ((max(m,n) + 31)/32)*32 )*nb;
}
TESTING_MALLOC_CPU( C, magmaFloatComplex, ldc*n );
TESTING_MALLOC_CPU( R, magmaFloatComplex, ldc*n );
TESTING_MALLOC_CPU( A, magmaFloatComplex, lda*k );
TESTING_MALLOC_CPU( W, magmaFloatComplex, lwork_max );
TESTING_MALLOC_CPU( tau, magmaFloatComplex, k );
TESTING_MALLOC_DEV( dC, magmaFloatComplex, ldc*n );
TESTING_MALLOC_DEV( dA, magmaFloatComplex, lda*k );
TESTING_MALLOC_DEV( dT, magmaFloatComplex, dt_size );
// C is full, m x n
size = ldc*n;
lapackf77_clarnv( &ione, ISEED, &size, C );
magma_csetmatrix( m, n, C, ldc, dC, ldc );
// A is m x k (left) or n x k (right)
lda = (side[iside] == MagmaLeft ? m : n);
size = lda*k;
lapackf77_clarnv( &ione, ISEED, &size, A );
// compute QR factorization to get Householder vectors in dA, tau, dT
magma_csetmatrix( lda, k, A, lda, dA, lda );
magma_cgeqrf_gpu( lda, k, dA, lda, tau, dT, &info );
magma_cgetmatrix( lda, k, dA, lda, A, lda );
if (info != 0)
printf("magma_cgeqrf_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_cunmqr( lapack_side_const( side[iside] ), lapack_trans_const( trans[itran] ),
&m, &n, &k,
A, &lda, tau, C, &ldc, W, &lwork_max, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cunmqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
// query for workspace size
lwork = -1;
magma_cunmqr_gpu( side[iside], trans[itran],
m, n, k,
dA, lda, tau, dC, ldc, W, lwork, dT, nb, &info );
if (info != 0)
printf("magma_cunmqr_gpu (lwork query) returned error %d: %s.\n",
(int) info, magma_strerror( info ));
lwork = (magma_int_t) MAGMA_C_REAL( W[0] );
if ( lwork < 0 || lwork > lwork_max )
printf("invalid lwork %d, lwork_max %d\n", (int) lwork, (int) lwork_max );
gpu_time = magma_sync_wtime( 0 ); // sync needed for L,N and R,T cases
magma_cunmqr_gpu( side[iside], trans[itran],
m, n, k,
dA, lda, tau, dC, ldc, W, lwork, dT, nb, &info );
gpu_time = magma_sync_wtime( 0 ) - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cunmqr_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_cgetmatrix( m, n, dC, ldc, R, ldc );
/* =====================================================================
compute relative error |QC_magma - QC_lapack| / |QC_lapack|
=================================================================== */
error = lapackf77_clange( "Fro", &m, &n, C, &ldc, work );
size = ldc*n;
blasf77_caxpy( &size, &c_neg_one, C, &ione, R, &ione );
error = lapackf77_clange( "Fro", &m, &n, R, &ldc, work ) / error;
printf( "%5d %5d %5d %4c %5c %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) m, (int) n, (int) k,
lapacke_side_const( side[iside] ),
lapacke_trans_const( trans[itran] ),
cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed") );
status += ! (error < tol);
TESTING_FREE_CPU( C );
TESTING_FREE_CPU( R );
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( W );
TESTING_FREE_CPU( tau );
TESTING_FREE_DEV( dC );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dT );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}} // end iside, itran
printf( "\n" );
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cunmlq
*/
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float Cnorm, error, work[1];
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t ione = 1;
magma_int_t mm, m, n, k, size, info;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t nb, ldc, lda, lwork, lwork_max;
magmaFloatComplex *C, *R, *A, *W, *tau;
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
// need slightly looser bound (60*eps instead of 30*eps) for some tests
opts.tolerance = max( 60., opts.tolerance );
float tol = opts.tolerance * lapackf77_slamch("E");
// test all combinations of input parameters
magma_side_t side [] = { MagmaLeft, MagmaRight };
magma_trans_t trans[] = { Magma_ConjTrans, MagmaNoTrans };
printf("%% M N K side trans CPU Gflop/s (sec) GPU Gflop/s (sec) ||R||_F / ||QC||_F\n");
printf("%%==============================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iside = 0; iside < 2; ++iside ) {
for( int itran = 0; itran < 2; ++itran ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
m = opts.msize[itest];
n = opts.nsize[itest];
k = opts.ksize[itest];
nb = magma_get_cgelqf_nb( m, n );
ldc = m;
// A is k x m (left) or k x n (right)
mm = (side[iside] == MagmaLeft ? m : n);
lda = k;
gflops = FLOPS_CUNMLQ( m, n, k, side[iside] ) / 1e9;
if ( side[iside] == MagmaLeft && m < k ) {
printf( "%5d %5d %5d %4c %5c skipping because side=left and m < k\n",
(int) m, (int) n, (int) k,
lapacke_side_const( side[iside] ),
lapacke_trans_const( trans[itran] ) );
continue;
}
if ( side[iside] == MagmaRight && n < k ) {
printf( "%5d %5d %5d %4c %5c skipping because side=right and n < k\n",
(int) m, (int) n, (int) k,
lapacke_side_const( side[iside] ),
lapacke_trans_const( trans[itran] ) );
continue;
}
// need at least 2*nb*nb for gelqf
lwork_max = max( max( m*nb, n*nb ), 2*nb*nb );
// this rounds it up slightly if needed to agree with lwork query
lwork_max = int( real( magma_cmake_lwork( lwork_max )));
TESTING_MALLOC_CPU( C, magmaFloatComplex, ldc*n );
TESTING_MALLOC_CPU( R, magmaFloatComplex, ldc*n );
TESTING_MALLOC_CPU( A, magmaFloatComplex, lda*mm );
TESTING_MALLOC_CPU( W, magmaFloatComplex, lwork_max );
TESTING_MALLOC_CPU( tau, magmaFloatComplex, k );
// C is full, m x n
size = ldc*n;
lapackf77_clarnv( &ione, ISEED, &size, C );
lapackf77_clacpy( "Full", &m, &n, C, &ldc, R, &ldc );
size = lda*mm;
lapackf77_clarnv( &ione, ISEED, &size, A );
// compute LQ factorization to get Householder vectors in A, tau
magma_cgelqf( k, mm, A, lda, tau, W, lwork_max, &info );
if (info != 0) {
printf("magma_cgelqf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_cunmlq( lapack_side_const( side[iside] ), lapack_trans_const( trans[itran] ),
&m, &n, &k,
A, &lda, tau, C, &ldc, W, &lwork_max, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0) {
printf("lapackf77_cunmlq returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
// query for workspace size
lwork = -1;
magma_cunmlq( side[iside], trans[itran],
m, n, k,
A, lda, tau, R, ldc, W, lwork, &info );
if (info != 0) {
printf("magma_cunmlq (lwork query) returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
lwork = (magma_int_t) MAGMA_C_REAL( W[0] );
if ( lwork < 0 || lwork > lwork_max ) {
printf("Warning: optimal lwork %d > allocated lwork_max %d\n", (int) lwork, (int) lwork_max );
lwork = lwork_max;
}
gpu_time = magma_wtime();
magma_cunmlq( side[iside], trans[itran],
m, n, k,
A, lda, tau, R, ldc, W, lwork, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0) {
printf("magma_cunmlq returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
compute relative error |QC_magma - QC_lapack| / |QC_lapack|
=================================================================== */
size = ldc*n;
blasf77_caxpy( &size, &c_neg_one, C, &ione, R, &ione );
Cnorm = lapackf77_clange( "Fro", &m, &n, C, &ldc, work );
error = lapackf77_clange( "Fro", &m, &n, R, &ldc, work ) / (magma_ssqrt(m*n) * Cnorm);
printf( "%5d %5d %5d %4c %5c %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) m, (int) n, (int) k,
lapacke_side_const( side[iside] ),
lapacke_trans_const( trans[itran] ),
cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed") );
status += ! (error < tol);
TESTING_FREE_CPU( C );
TESTING_FREE_CPU( R );
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( W );
TESTING_FREE_CPU( tau );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}} // end iside, itran
printf( "\n" );
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing ctrsm
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf=0, magma_time=0, cublas_perf, cublas_time, cpu_perf=0, cpu_time=0;
float magma_error=0, cublas_error, lapack_error, work[1];
magma_int_t M, N, info;
magma_int_t Ak;
magma_int_t sizeA, sizeB;
magma_int_t lda, ldb, ldda, lddb;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t *ipiv;
magmaFloatComplex *h_A, *h_B, *h_Bcublas, *h_Bmagma, *h_Blapack, *h_X;
magmaFloatComplex_ptr d_A, d_B;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex alpha = MAGMA_C_MAKE( 0.29, -0.86 );
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
float tol = opts.tolerance * lapackf77_slamch("E");
// pass ngpu = -1 to test multi-GPU code using 1 gpu
magma_int_t abs_ngpu = abs( opts.ngpu );
printf("%% side = %s, uplo = %s, transA = %s, diag = %s, ngpu = %d\n",
lapack_side_const(opts.side), lapack_uplo_const(opts.uplo),
lapack_trans_const(opts.transA), lapack_diag_const(opts.diag), int(abs_ngpu) );
printf("%% M N MAGMA Gflop/s (ms) CUBLAS Gflop/s (ms) CPU Gflop/s (ms) MAGMA CUBLAS LAPACK error\n");
printf("%%============================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
gflops = FLOPS_CTRSM(opts.side, M, N) / 1e9;
if ( opts.side == MagmaLeft ) {
lda = M;
Ak = M;
} else {
lda = N;
Ak = N;
}
ldb = M;
ldda = magma_roundup( lda, opts.align ); // multiple of 32 by default
lddb = magma_roundup( ldb, opts.align ); // multiple of 32 by default
sizeA = lda*Ak;
sizeB = ldb*N;
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, lda*Ak );
TESTING_MALLOC_CPU( h_B, magmaFloatComplex, ldb*N );
TESTING_MALLOC_CPU( h_X, magmaFloatComplex, ldb*N );
TESTING_MALLOC_CPU( h_Blapack, magmaFloatComplex, ldb*N );
TESTING_MALLOC_CPU( h_Bcublas, magmaFloatComplex, ldb*N );
TESTING_MALLOC_CPU( h_Bmagma, magmaFloatComplex, ldb*N );
TESTING_MALLOC_CPU( ipiv, magma_int_t, Ak );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*Ak );
TESTING_MALLOC_DEV( d_B, magmaFloatComplex, lddb*N );
/* Initialize the matrices */
/* Factor A into LU to get well-conditioned triangular matrix.
* Copy L to U, since L seems okay when used with non-unit diagonal
* (i.e., from U), while U fails when used with unit diagonal. */
lapackf77_clarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_cgetrf( &Ak, &Ak, h_A, &lda, ipiv, &info );
for( int j = 0; j < Ak; ++j ) {
for( int i = 0; i < j; ++i ) {
*h_A(i,j) = *h_A(j,i);
}
}
lapackf77_clarnv( &ione, ISEED, &sizeB, h_B );
memcpy( h_Blapack, h_B, sizeB*sizeof(magmaFloatComplex) );
magma_csetmatrix( Ak, Ak, h_A, lda, d_A, ldda, opts.queue );
/* =====================================================================
Performs operation using MAGMABLAS
=================================================================== */
#if defined(HAVE_CUBLAS)
magma_csetmatrix( M, N, h_B, ldb, d_B, lddb, opts.queue );
magma_time = magma_sync_wtime( opts.queue );
if (opts.ngpu == 1) {
magmablas_ctrsm( opts.side, opts.uplo, opts.transA, opts.diag,
M, N,
alpha, d_A, ldda,
d_B, lddb, opts.queue );
}
else {
magma_ctrsm_m( abs_ngpu, opts.side, opts.uplo, opts.transA, opts.diag,
M, N,
alpha, d_A, ldda,
d_B, lddb );
}
magma_time = magma_sync_wtime( opts.queue ) - magma_time;
magma_perf = gflops / magma_time;
magma_cgetmatrix( M, N, d_B, lddb, h_Bmagma, ldb, opts.queue );
#endif
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_csetmatrix( M, N, h_B, ldb, d_B, lddb, opts.queue );
cublas_time = magma_sync_wtime( opts.queue );
#if defined(HAVE_CUBLAS)
// opts.handle also uses opts.queue
cublasCtrsm( opts.handle,
cublas_side_const(opts.side), cublas_uplo_const(opts.uplo),
cublas_trans_const(opts.transA), cublas_diag_const(opts.diag),
M, N,
&alpha, d_A, ldda,
d_B, lddb );
#elif defined(HAVE_clBLAS)
clblasCtrsm( clblasColumnMajor,
clblas_side_const(opts.side), clblas_uplo_const(opts.uplo),
clblas_trans_const(opts.transA), clblas_diag_const(opts.diag),
M, N,
alpha, d_A, 0, ldda,
d_B, 0, lddb,
1, &opts.queue, 0, NULL, NULL );
#endif
cublas_time = magma_sync_wtime( opts.queue ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_cgetmatrix( M, N, d_B, lddb, h_Bcublas, ldb, opts.queue );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
blasf77_ctrsm( lapack_side_const(opts.side), lapack_uplo_const(opts.uplo),
lapack_trans_const(opts.transA), lapack_diag_const(opts.diag),
&M, &N,
&alpha, h_A, &lda,
h_Blapack, &ldb );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
}
/* =====================================================================
Check the result
=================================================================== */
// ||b - 1/alpha*A*x|| / (||A||*||x||)
magmaFloatComplex inv_alpha = MAGMA_C_DIV( c_one, alpha );
float normR, normX, normA;
normA = lapackf77_clange( "M", &Ak, &Ak, h_A, &lda, work );
#if defined(HAVE_CUBLAS)
// check magma
memcpy( h_X, h_Bmagma, sizeB*sizeof(magmaFloatComplex) );
blasf77_ctrmm( lapack_side_const(opts.side), lapack_uplo_const(opts.uplo),
lapack_trans_const(opts.transA), lapack_diag_const(opts.diag),
&M, &N,
&inv_alpha, h_A, &lda,
h_X, &ldb );
blasf77_caxpy( &sizeB, &c_neg_one, h_B, &ione, h_X, &ione );
normR = lapackf77_clange( "M", &M, &N, h_X, &ldb, work );
normX = lapackf77_clange( "M", &M, &N, h_Bmagma, &ldb, work );
magma_error = normR/(normX*normA);
#endif
// check cublas
memcpy( h_X, h_Bcublas, sizeB*sizeof(magmaFloatComplex) );
blasf77_ctrmm( lapack_side_const(opts.side), lapack_uplo_const(opts.uplo),
lapack_trans_const(opts.transA), lapack_diag_const(opts.diag),
&M, &N,
&inv_alpha, h_A, &lda,
h_X, &ldb );
blasf77_caxpy( &sizeB, &c_neg_one, h_B, &ione, h_X, &ione );
normR = lapackf77_clange( "M", &M, &N, h_X, &ldb, work );
normX = lapackf77_clange( "M", &M, &N, h_Bcublas, &ldb, work );
cublas_error = normR/(normX*normA);
if ( opts.lapack ) {
// check lapack
// this verifies that the matrix wasn't so bad that it couldn't be solved accurately.
memcpy( h_X, h_Blapack, sizeB*sizeof(magmaFloatComplex) );
blasf77_ctrmm( lapack_side_const(opts.side), lapack_uplo_const(opts.uplo),
lapack_trans_const(opts.transA), lapack_diag_const(opts.diag),
&M, &N,
&inv_alpha, h_A, &lda,
h_X, &ldb );
blasf77_caxpy( &sizeB, &c_neg_one, h_B, &ione, h_X, &ione );
normR = lapackf77_clange( "M", &M, &N, h_X, &ldb, work );
normX = lapackf77_clange( "M", &M, &N, h_Blapack, &ldb, work );
lapack_error = normR/(normX*normA);
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %8.2e %s\n",
(int) M, (int) N,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
magma_error, cublas_error, lapack_error,
(magma_error < tol && cublas_error < tol? "ok" : "failed"));
status += ! (magma_error < tol && cublas_error < tol);
}
else {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) --- ( --- ) %8.2e %8.2e --- %s\n",
(int) M, (int) N,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
magma_error, cublas_error,
(magma_error < tol && cublas_error < tol ? "ok" : "failed"));
status += ! (magma_error < tol && cublas_error < tol);
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_X );
TESTING_FREE_CPU( h_Blapack );
TESTING_FREE_CPU( h_Bcublas );
TESTING_FREE_CPU( h_Bmagma );
TESTING_FREE_CPU( ipiv );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
