/***************************************************************************//**
Purpose
-------
CGEQRF computes a QR factorization of a COMPLEX M-by-N matrix A:
A = Q * R using multiple GPUs. This version does not require work space on the GPU
passed as input. GPU memory is allocated in the routine.
Arguments
---------
@param[in]
ngpu INTEGER
Number of GPUs to use. ngpu > 0.
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
A COMPLEX array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, the elements on and above the diagonal of the array
contain the min(M,N)-by-N upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of min(m,n) elementary reflectors (see Further
Details).
\n
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[out]
tau COMPLEX array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
@param[out]
work (workspace) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
\n
Higher performance is achieved if WORK is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lwork INTEGER
The dimension of the array WORK. LWORK >= N*NB,
where NB can be obtained through magma_get_cgeqrf_nb( M, 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.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
Further Details
---------------
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
@ingroup magma_geqrf
*******************************************************************************/
extern "C" magma_int_t
magma_cgeqrf_m(
magma_int_t ngpu,
magma_int_t m, magma_int_t n,
magmaFloatComplex *A, magma_int_t lda, magmaFloatComplex *tau,
magmaFloatComplex *work, magma_int_t lwork,
magma_int_t *info )
{
magmaFloatComplex *da[MagmaMaxGPUs];
magmaFloatComplex c_one = MAGMA_C_ONE;
magma_int_t i, min_mn, ldda;
*info = 0;
magma_int_t nb = magma_get_cgeqrf_nb( m, n );
magma_int_t lwkopt = n * nb;
work[0] = magma_cmake_lwork( lwkopt );
bool lquery = (lwork == -1);
if (ngpu < 0 || ngpu > MagmaMaxGPUs) {
*info = -1;
} else if (m < 0) {
*info = -2;
} else if (n < 0) {
*info = -3;
} else if (lda < max(1,m)) {
*info = -5;
} else if (lwork < max(1,n) && ! lquery) {
*info = -8;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
min_mn = min(m,n);
if (min_mn == 0) {
work[0] = c_one;
return *info;
}
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
ldda = magma_roundup( m, 32 );
magma_int_t n_local[MagmaMaxGPUs];
for (i=0; i < ngpu; i++) {
n_local[i] = ((n/nb)/ngpu)*nb;
if (i < (n/nb)%ngpu)
n_local[i] += nb;
else if (i == (n/nb)%ngpu)
n_local[i] += n%nb;
magma_setdevice(i);
// TODO on failure, free previously allocated memory
if (MAGMA_SUCCESS != magma_cmalloc( &da[i], ldda*n_local[i] )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
}
if (m > nb && n > nb) {
magma_queue_t queues[MagmaMaxGPUs];
for( magma_int_t dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
magma_queue_create( dev, &queues[dev] );
}
/* Copy the matrix to the GPUs in 1D block cyclic distribution */
magma_csetmatrix_1D_col_bcyclic( ngpu, m, n, nb, A, lda, da, ldda, queues );
for( magma_int_t dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
magma_queue_sync( queues[dev] );
}
/* Factor using the GPU interface */
magma_cgeqrf2_mgpu( ngpu, m, n, da, ldda, tau, info);
/* Copy the matrix back from the GPUs to the CPU */
magma_cgetmatrix_1D_col_bcyclic( ngpu, m, n, nb, da, ldda, A, lda, queues );
for( magma_int_t dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
magma_queue_sync( queues[dev] );
magma_queue_destroy( queues[dev] );
}
}
else {
lapackf77_cgeqrf(&m, &n, A, &lda, tau, work, &lwork, info);
}
/* Free the allocated GPU memory */
for (i=0; i < ngpu; i++) {
magma_setdevice(i);
magma_free( da[i] );
}
magma_setdevice( orig_dev );
return *info;
} /* magma_cgeqrf4 */
extern "C" magma_int_t
magma_chetrd_he2hb_mgpu( char uplo, magma_int_t n, magma_int_t nb,
magmaFloatComplex *a, magma_int_t lda,
magmaFloatComplex *tau,
magmaFloatComplex *work, magma_int_t lwork,
magmaFloatComplex *dAmgpu[], magma_int_t ldda,
magmaFloatComplex *dTmgpu[], magma_int_t lddt,
magma_int_t ngpu, magma_int_t distblk,
magma_queue_t streams[][20], magma_int_t nstream,
magma_int_t threads, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CHETRD_HE2HB reduces a complex Hermitian matrix A to real symmetric
band-diagonal form T by an orthogonal similarity transformation:
Q**H * A * Q = T.
This version stores the triangular matrices T used in the accumulated
Householder transformations (I - V T V').
Arguments
=========
UPLO (input) CHARACTER*1
= 'U': Upper triangle of A is stored;
= 'L': Lower triangle of A is stored.
N (input) INTEGER
The order of the matrix A. N >= 0.
A (input/output) COMPLEX array, dimension (LDA,N)
On entry, the Hermitian matrix A. If UPLO = 'U', the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = 'L', the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit, if UPLO = 'U', the Upper band-diagonal of A is
overwritten by the corresponding elements of the
band-diagonal matrix T, and the elements above the band
diagonal, with the array TAU, represent the orthogonal
matrix Q as a product of elementary reflectors; if UPLO
= 'L', the the Lower band-diagonal of A is overwritten by
the corresponding elements of the band-diagonal
matrix T, and the elements below the band-diagonal, with
the array TAU, represent the orthogonal matrix Q as a product
of elementary reflectors. See Further Details.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
TAU (output) COMPLEX array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details).
WORK (workspace/output) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The dimension of the array WORK. LWORK >= 1.
For optimum performance LWORK >= N*NB, where NB is the
optimal blocksize.
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.
dT (output) COMPLEX array on the GPU, dimension N*NB,
where NB is the optimal blocksize.
On exit dT holds the upper triangular matrices T from the
accumulated Householder transformations (I - V T V') used
in the factorization. The nb x nb matrices T are ordered
consecutively in memory one after another.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
Further Details
===============
If UPLO = 'U', the matrix Q is represented as a product of elementary
reflectors
Q = H(n-1) . . . H(2) H(1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(i+1:n) = 0 and v(i) = 1; v(1:i-1) is stored on exit in
A(1:i-1,i+1), and tau in TAU(i).
If UPLO = 'L', the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(n-1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i) = 0 and v(i+1) = 1; v(i+2:n) is stored on exit in A(i+2:n,i),
and tau in TAU(i).
The contents of A on exit are illustrated by the following examples
with n = 5:
if UPLO = 'U': if UPLO = 'L':
( d e v2 v3 v4 ) ( d )
( d e v3 v4 ) ( e d )
( d e v4 ) ( v1 e d )
( d e ) ( v1 v2 e d )
( d ) ( v1 v2 v3 e d )
where d and e denote diagonal and off-diagonal elements of T, and vi
denotes an element of the vector defining H(i).
===================================================================== */
#define a_ref(a_1,a_2) ( a + ((a_2)-1)*( lda) + (a_1)-1)
#define da_ref(a_1,a_2) (da + ((a_2)-1)*(ldda) + (a_1)-1)
#define tau_ref(a_1) (tau + (a_1)-1)
#define t_ref(a_1) (dT + ((a_1)-1)*(lddt))
#define Atest(a_1,a_2) ( Atest + ((a_2)-1)*( lda) + (a_1)-1)
#define dttest(a_0, a_1, a_2) (dTmgpu[a_0] + ((a_2)-1)*(lddt))
#define datest(a_0, a_1, a_2) (dAmgpu[a_0] + ((a_2)-1)*(ldda) + (a_1)-1)
char uplo_[2] = {uplo, 0};
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex c_neg_half = MAGMA_C_NEG_HALF;
magmaFloatComplex c_one = MAGMA_C_ONE ;
magmaFloatComplex c_zero = MAGMA_C_ZERO;
float d_one = MAGMA_D_ONE;
magma_int_t pm, pn, indi, indj, pk;
magma_int_t pm_old=0, pn_old=0, indi_old=0, indj_old=0, flipV=-1;
magma_int_t iblock, idev, di;
int i;
int lwkopt;
int lquery;
assert (nstream>=3);
assert (nstream>=(ngpu+1));
*info = 0;
int upper = lapackf77_lsame(uplo_, "U");
lquery = lwork == -1;
if (! upper && ! lapackf77_lsame(uplo_, "L")) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
} else if (lwork < 1 && ! lquery) {
*info = -9;
}
/* Determine the block size. */
lwkopt = n * nb;
if (*info == 0) {
MAGMA_C_SET2REAL( work[0], lwkopt );
}
if (*info != 0)
return *info;
else if (lquery)
return *info;
/* Quick return if possible */
if (n == 0) {
work[0] = c_one;
return *info;
}
magma_int_t mklth = min(threads,16);
magma_setlapack_numthreads(mklth);
magma_int_t gnode[MagmaMaxGPUs][MagmaMaxGPUs+2];
magma_int_t nbcmplx=0;
magma_buildconnection_mgpu(gnode, &nbcmplx, ngpu);
#ifdef ENABLE_DEBUG
printf(" Initializing communication pattern.... GPU-ncmplx %d\n\n" , nbcmplx);
#endif
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_t cstream;
magmablasGetKernelStream(&cstream);
magmaFloatComplex *dspace[MagmaMaxGPUs];
magmaFloatComplex *dwork[MagmaMaxGPUs], *dworkbis[MagmaMaxGPUs];
magmaFloatComplex *dvall[MagmaMaxGPUs], *dv[MagmaMaxGPUs], *dw[MagmaMaxGPUs];
magmaFloatComplex *workngpu[MagmaMaxGPUs+1];
magma_event_t redevents[MagmaMaxGPUs][MagmaMaxGPUs*MagmaMaxGPUs+10];
magma_int_t nbevents = MagmaMaxGPUs*MagmaMaxGPUs;
magma_int_t lddv = ldda;
magma_int_t lddw = lddv;
magma_int_t dwrk2siz = ldda*nb*(ngpu+1);
magma_int_t worksiz = n*nb;
magma_int_t devworksiz = 2*nb*lddv + nb*lddw + nb*ldda + dwrk2siz; // 2*dv(dv0+dv1) + dw + dwork +dworkbis
// local allocation and stream creation
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_cmalloc( &dspace[dev], devworksiz );
magma_cmalloc_pinned ( &workngpu[dev], worksiz);
dvall[dev] = dspace[dev];
dw[dev] = dvall[dev] + 2*nb*lddv;
dwork[dev] = dw[dev] + nb*lddw;
dworkbis[dev] = dwork[dev] + nb*ldda;
magmablasSetKernelStream( streams[ dev ][ 0 ] );
for( magma_int_t i = 0; i < nbevents; ++i ) {
cudaEventCreateWithFlags(&redevents[dev][i],cudaEventDisableTiming);
}
}
magma_cmalloc_pinned ( &workngpu[ngpu], worksiz);
magmaFloatComplex *worktest = NULL; //(magmaFloatComplex *) malloc(n*nb*sizeof(magmaFloatComplex)); // not used
// ======================
magmaFloatComplex *hT = work + lwork - nb*nb;
lwork -= nb*nb;
memset( hT, 0, nb*nb*sizeof(magmaFloatComplex));
if (upper) {
printf("CHETRD_HE2HB is not yet implemented for upper matrix storage. Exit.\n");
exit(1);
}else {
/* Reduce the lower triangle of A */
for (i = 1; i <= n-nb; i += nb)
{
indi = i+nb;
indj = i;
pm = n - i - nb + 1;
//pn = min(i+nb-1, n-nb) -i + 1;
pn = nb;
/* Get the current panel (no need for the 1st iteration) */
if (i > 1 ){
// cpanel_to_q copy the upper oof diagonal part of
// the matrix to work to be restored later. acctually
// the zero's and one's putted are not used this is only
// because we don't have a function that copy only the
// upper part of A to be restored after copying the
// lookahead panel that has been computted from GPU to CPU.
cpanel_to_q(MagmaUpper, pn-1, a_ref(i, i+1), lda, work);
// find the device who own the panel then send it to the CPU.
// below a -1 was added and then a -1 was done on di because of the fortran indexing
iblock = ((i-1) / distblk) / ngpu; // local block id
di = iblock*distblk + (i-1)%distblk; // local index in parent matrix
idev = ((i-1) / distblk) % ngpu; // device with this block
//printf("Receiving panel ofsize %d %d from idev %d A(%d,%d) \n",(pm+pn), pn,idev,i-1,di);
magma_setdevice( idev );
//magma_device_sync();
magma_cgetmatrix_async( (pm+pn), pn,
datest(idev, i, di+1), ldda,
a_ref ( i, i), lda, streams[ idev ][ nstream-1 ] );
/*
magma_device_sync();
cudaMemcpy2DAsync(a_ref(i,i), lda*sizeof(magmaFloatComplex),
datest(idev,i,di+1), ldda*sizeof(magmaFloatComplex),
(pm+pn)*sizeof(magmaFloatComplex), pn,
cudaMemcpyDeviceToHost, streams[ idev ][ nstream-1 ]);
*/
//magma_setdevice( 0 );
//printf("updating cher2k on A(%d,%d) of size %d %d \n",indi_old+pn_old-1,indi_old+pn_old-1,pm_old-pn_old,pn_old);
// compute CHER2K_MGPU
magmablas_cher2k_mgpu2(
MagmaLower, MagmaNoTrans, pm_old-pn_old, pn_old,
c_neg_one, dv, pm_old, pn_old,
dw, pm_old, pn_old,
d_one, dAmgpu, ldda, indi_old+pn_old-1,
ngpu, distblk, streams, 2 );
//magma_setdevice( 0 );
magma_setdevice( idev );
magma_queue_sync( streams[idev][ nstream-1 ] );
//magma_setdevice( 0 );
cq_to_panel(MagmaUpper, pn-1, a_ref(i, i+1), lda, work);
}
/* ==========================================================
QR factorization on a panel starting nb off of the diagonal.
Prepare the V and T matrices.
========================================================== */
lapackf77_cgeqrf(&pm, &pn, a_ref(indi, indj), &lda,
tau_ref(i), work, &lwork, info);
/* Form the matrix T */
pk=min(pm,pn);
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&pm, &pk, a_ref(indi, indj), &lda,
tau_ref(i), hT, &nb);
/* Prepare V - put 0s in the upper triangular part of the panel
(and 1s on the diagonal), temporaly storing the original in work */
cpanel_to_q(MagmaUpper, pk, a_ref(indi, indj), lda, work);
/* Send V and T from the CPU to the GPU */
// To be able to overlap the GET with the CHER2K
// it should be done on last stream.
// TO Avoid a BUG that is overwriting the old_V
// used atthis moment by cher2k with the new_V
// send it now, we decide to have a flipflop
// vector of Vs. if step%2=0 use V[0] else use V[nb*n]
flipV = ((i-1)/nb)%2;
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
dv[dev] = dvall[dev] + flipV*nb*lddv;
}
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
// send V
magma_csetmatrix_async( pm, pk,
a_ref(indi, indj), lda,
dv[dev], pm, streams[dev][nstream-1] );
// Send the triangular factor T to the GPU
magma_csetmatrix_async( pk, pk,
hT, nb,
dttest(dev, 1, i), lddt, streams[dev][nstream-1] );
}
/* ==========================================================
Compute W:
1. X = A (V T)
2. W = X - 0.5* V * (T' * (V' * X))
========================================================== */
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
// dwork = V T
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ nstream-1 ] );
magma_queue_sync( streams[dev][nstream-1] );
magma_cgemm(MagmaNoTrans, MagmaNoTrans, pm, pk, pk,
c_one, dv[dev], pm,
dttest(dev, 1, i), lddt,
c_zero, dwork[dev], pm);
}
// ===============================================
// SYNC TO BE SURE THAT BOTH V AND T WERE
// RECEIVED AND VT IS COMPUTED and SYR2K is done
// ===============================================
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
for( magma_int_t s = 0; s < nstream; ++s )
magma_queue_sync( streams[dev][s] );
}
// compute CHEMM_MGPU
// The broadcast of the result done inside this function
// should be done in stream [0] because i am assuming this
// for the GEMMs below otherwise I have to SYNC over the
// Broadcasting stream.
if(ngpu==1){
magmablasSetKernelStream( streams[ 0 ][ 0 ] );
magma_chemm(MagmaLeft, uplo, pm, pk,
c_one, dAmgpu[0]+(indi-1)*ldda+(indi-1), ldda,
dwork[0], pm,
c_zero, dw[0], pm);
}else{
magmablas_chemm_mgpu_com(
MagmaLeft, uplo, pm, pk,
c_one, dAmgpu, ldda, indi-1,
dwork, pm,
c_zero, dw, pm, dworkbis, dwrk2siz, worktest, pm, workngpu, worksiz,
ngpu, distblk, streams, nstream-1, redevents, nbevents, gnode, nbcmplx);
}
/* dwork = V*T already ==> dwork' = T'*V'
* compute T'*V'*X ==> dwork'*W ==>
* dwork + pm*nb = ((T' * V') * X) = dwork' * X = dwork' * W */
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
// Here we have to wait until the broadcast of CHEMM has been done.
// Note that the broadcast should be done on stream[0] so in a way
// we can continue here on the same stream and avoid a sync
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
// magma_queue_sync( streams[dev][0] );
magma_cgemm(MagmaConjTrans, MagmaNoTrans, pk, pk, pm,
c_one, dwork[dev], pm,
dw[dev], pm,
c_zero, dworkbis[dev], nb);
/* W = X - 0.5 * V * T'*V'*X
* = X - 0.5 * V * (dwork + pm*nb) = W - 0.5 * V * (dwork + pm*nb) */
magma_cgemm(MagmaNoTrans, MagmaNoTrans, pm, pk, pk,
c_neg_half, dv[dev], pm,
dworkbis[dev], nb,
c_one, dw[dev], pm);
}
/* restore the panel it is put here to overlap with the previous GEMM*/
cq_to_panel(MagmaUpper, pk, a_ref(indi, indj), lda, work);
// ===============================================
// SYNC TO BE SURE THAT BOTH V AND W ARE DONE
// ===============================================
// Synchronise to be sure that W has been computed
// because next CHER2K use streaming and may happen
// that lunch a gemm on stream 2 while stream 0
// which compute those 2 GEMM above has not been
// computed and also used for the same reason in
// the panel update below and also for the last HER2K
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_queue_sync( streams[dev][0] );
}
/* ==========================================================
Update the unreduced submatrix A(i+ib:n,i+ib:n), using
an update of the form: A := A - V*W' - W*V'
========================================================== */
if (i + nb <= n-nb){
/* There would be next iteration;
do lookahead - update the next panel */
// below a -1 was added and then a -1 was done on di because of the fortran indexing
iblock = ((indi-1) / distblk) / ngpu; // local block id
di = iblock*distblk + (indi-1)%distblk; // local index in parent matrix
idev = ((indi-1) / distblk) % ngpu; // device with this block
magma_setdevice( idev );
magmablasSetKernelStream( streams[ idev ][ nstream-1 ] );
//magma_queue_sync( streams[idev][0] ); removed because the sync has been done in the loop above
magma_cgemm(MagmaNoTrans, MagmaConjTrans, pm, pn, pn, c_neg_one,
dv[idev], pm,
dw[idev] , pm, c_one,
datest(idev, indi, di+1), ldda);
magma_cgemm(MagmaNoTrans, MagmaConjTrans, pm, pn, pn, c_neg_one,
dw[idev] , pm,
dv[idev], pm, c_one,
datest(idev, indi, di+1), ldda);
//printf("updating next panel distblk %d idev %d on A(%d,%d) of size %d %d %d \n",distblk,idev,indi-1,di,pm,pn,pn);
}
else {
/* no look-ahead as this is last iteration */
// below a -1 was added and then a -1 was done on di because of the fortran indexing
iblock = ((indi-1) / distblk) / ngpu; // local block id
di = iblock*distblk + (indi-1)%distblk; // local index in parent matrix
idev = ((indi-1) / distblk) % ngpu; // device with this block
magma_setdevice( idev );
magmablasSetKernelStream( streams[ idev ][ 0 ] );
//printf("LAST CHER2K idev %d on A(%d,%d) of size %d \n",idev, indi-1,di,pk);
magma_cher2k(MagmaLower, MagmaNoTrans, pk, pk, c_neg_one,
dv[idev], pm,
dw[idev] , pm, d_one,
datest(idev, indi, di+1), ldda);
/* Send the last block to the CPU */
cpanel_to_q(MagmaUpper, pk-1, a_ref(n-pk+1, n-pk+2), lda, work);
magma_cgetmatrix( pk, pk,
datest(idev, indi, di+1), ldda,
a_ref(n-pk+1, n-pk+1), lda );
cq_to_panel(MagmaUpper, pk-1, a_ref(n-pk+1, n-pk+2), lda, work);
}
indi_old = indi;
indj_old = indj;
pm_old = pm;
pn_old = pn;
} // end loop for(i)
}// end of LOWER
//magma_setdevice( 0 );
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_free( dspace[dev]);
magma_free_pinned(workngpu[dev]);
for( magma_int_t e = 0; e < nbevents; ++e ) {
cudaEventDestroy(redevents[dev][e]);
}
}
magma_free_pinned(workngpu[ngpu]);
free(worktest);
magma_setdevice( cdev );
magmablasSetKernelStream( cstream );
MAGMA_C_SET2REAL( work[0], lwkopt );
magma_setlapack_numthreads(1);
return *info;
} /* chetrd_he2hb_ */
extern "C" magma_err_t
magma_cgeqrf(magma_int_t m, magma_int_t n,
magmaFloatComplex *A, magma_int_t lda, magmaFloatComplex *tau,
magmaFloatComplex *work, magma_int_t lwork,
magma_int_t *info,
magma_queue_t* queue )
{
/* -- clMAGMA (version 1.1.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date January 2014
Purpose
=======
CGEQRF computes a QR factorization of a COMPLEX M-by-N matrix A:
A = Q * R. This version does not require work space on the GPU
passed as input. GPU memory is allocated in the routine.
If the current stream is NULL, this version replaces it with user defined
stream to overlap computation with communication.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. N >= 0.
A (input/output) COMPLEX array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, the elements on and above the diagonal of the array
contain the min(M,N)-by-N upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of min(m,n) elementary reflectors (see Further
Details).
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,M).
TAU (output) COMPLEX array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
WORK (workspace/output) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
Higher performance is achieved if WORK is in pinned memory, e.g.
allocated using magma_malloc_pinned.
LWORK (input) INTEGER
The dimension of the array WORK. LWORK >= max( N*NB, 2*NB*NB ),
where NB can be obtained through magma_get_cgeqrf_nb(M).
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.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
Further Details
===============
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
===================================================================== */
#define A(i,j) ( A + (i) + (j)*lda )
#define dA(i,j) dA, dA_offset + (i) + (j)*ldda
magmaFloatComplex_ptr dA, dwork, dT;
size_t dA_offset, dwork_offset, dT_offset;
magmaFloatComplex c_one = MAGMA_C_ONE;
magma_int_t i, k, lddwork, old_i, old_ib;
magma_int_t ib, ldda;
*info = 0;
magma_int_t nb = magma_get_cgeqrf_nb(min(m, n));
// need 2*nb*nb to store T and upper triangle of V simultaneously
magma_int_t lwkopt = max(n*nb, 2*nb*nb);
work[0] = MAGMA_C_MAKE( (float)lwkopt, 0 );
int lquery = (lwork == -1);
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,m)) {
*info = -4;
} else if (lwork < max(1, lwkopt) && ! lquery) {
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
k = min(m,n);
if (k == 0) {
work[0] = c_one;
return *info;
}
// largest N for larfb is n-nb (trailing matrix lacks 1st panel)
lddwork = ((n+31)/32)*32 - nb;
ldda = ((m+31)/32)*32;
magma_int_t num_gpus = magma_num_gpus();
if( num_gpus > 1 ) {
/* call multiple-GPU interface */
printf("multiple-GPU verison not implemented\n");
return MAGMA_ERR_NOT_IMPLEMENTED;
//return magma_cgeqrf4(num_gpus, m, n, A, lda, tau, work, lwork, info);
}
// allocate space for dA, dwork, and dT
if (MAGMA_SUCCESS != magma_cmalloc( &dA, (n*ldda + nb*lddwork + nb*nb) )) {
/* Switch to the "out-of-core" (out of GPU-memory) version */
printf("non-GPU-resident version not implemented\n");
return MAGMA_ERR_NOT_IMPLEMENTED;
//return magma_cgeqrf_ooc(m, n, A, lda, tau, work, lwork, info);
}
dA_offset = 0;
dwork = dA;
dwork_offset = n*ldda;
dT = dA;
dT_offset = n*ldda + nb*lddwork;
if ( (nb > 1) && (nb < k) ) {
/* Use blocked code initially.
Asynchronously send the matrix to the GPU except the first panel. */
magma_csetmatrix_async( m, n-nb,
A(0,nb), 0, lda,
dA(0,nb), ldda, queue[0], NULL );
old_i = 0;
old_ib = nb;
for (i = 0; i < k-nb; i += nb) {
ib = min(k-i, nb);
if (i>0) {
/* download i-th panel */
magma_queue_sync( queue[1] );
magma_cgetmatrix_async( m-i, ib,
dA(i,i), ldda,
A(i,i), 0, lda, queue[0], NULL );
/* Apply H' to A(i:m,i+2*ib:n) from the left */
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
m-old_i, n-old_i-2*old_ib, old_ib,
dA(old_i, old_i), ldda, dT, dT_offset, nb,
dA(old_i, old_i+2*old_ib), ldda, dwork, dwork_offset, lddwork, queue[1]);
magma_cgetmatrix_async( i, ib,
dA(0,i), ldda,
A(0,i), 0, lda, queue[1], NULL );
magma_queue_sync( queue[0] );
}
magma_int_t rows = m-i;
lapackf77_cgeqrf(&rows, &ib, A(i,i), &lda, tau+i, work, &lwork, info);
/* Form the triangular factor of the block reflector
H = H(i) H(i+1) . . . H(i+ib-1) */
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&rows, &ib, A(i,i), &lda, tau+i, work, &ib);
cpanel_to_q(MagmaUpper, ib, A(i,i), lda, work+ib*ib);
/* download the i-th V matrix */
magma_csetmatrix_async( rows, ib, A(i,i), 0, lda, dA(i,i), ldda, queue[0], NULL );
/* download the T matrix */
magma_queue_sync( queue[1] );
magma_csetmatrix_async( ib, ib, work, 0, ib, dT, dT_offset, nb, queue[0], NULL );
magma_queue_sync( queue[0] );
if (i + ib < n) {
if (i+ib < k-nb) {
/* Apply H' to A(i:m,i+ib:i+2*ib) from the left (look-ahead) */
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, ib, ib,
dA(i, i ), ldda, dT, dT_offset, nb,
dA(i, i+ib), ldda, dwork, dwork_offset, lddwork, queue[1]);
cq_to_panel(MagmaUpper, ib, A(i,i), lda, work+ib*ib);
}
else {
/* After last panel, update whole trailing matrix. */
/* Apply H' to A(i:m,i+ib:n) from the left */
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, n-i-ib, ib,
dA(i, i ), ldda, dT, dT_offset, nb,
dA(i, i+ib), ldda, dwork, dwork_offset, lddwork, queue[1]);
cq_to_panel(MagmaUpper, ib, A(i,i), lda, work+ib*ib);
}
old_i = i;
old_ib = ib;
}
}
} else {
i = 0;
}
/* Use unblocked code to factor the last or only block. */
if (i < k) {
ib = n-i;
if (i != 0) {
magma_cgetmatrix( m, ib, dA(0,i), ldda, A(0,i), 0, lda, queue[1] );
}
magma_int_t rows = m-i;
lapackf77_cgeqrf(&rows, &ib, A(i,i), &lda, tau+i, work, &lwork, info);
}
magma_queue_sync(queue[0]);
magma_queue_sync(queue[1]);
magma_free( dA );
return *info;
} /* magma_cgeqrf */
magma_err_t
magma_cgeqrf2_gpu( magma_int_t m, magma_int_t n,
magmaFloatComplex_ptr dA, size_t dA_offset, magma_int_t ldda,
magmaFloatComplex *tau, magma_err_t *info,
magma_queue_t* queue)
{
/* -- clMAGMA (version 1.1.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date January 2014
Purpose
=======
CGEQRF computes a QR factorization of a complex M-by-N matrix A:
A = Q * R.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. N >= 0.
dA (input/output) COMPLEX array on the GPU, dimension (LDDA,N)
On entry, the M-by-N matrix dA.
On exit, the elements on and above the diagonal of the array
contain the min(M,N)-by-N upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of min(m,n) elementary reflectors (see Further
Details).
LDDA (input) INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
To benefit from coalescent memory accesses LDDA must be
dividable by 16.
TAU (output) COMPLEX array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
if INFO = -9, internal GPU memory allocation failed.
Further Details
===============
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
===================================================================== */
#define dA(a_1,a_2) dA, (dA_offset + (a_1) + (a_2)*(ldda))
#define work_ref(a_1) work, (a_1)
#define work_href(a_1) ( work + (a_1))
#define hwork ( work + (nb)*(m))
#define hhwork work, ((nb)*(m))
magmaFloatComplex_ptr dwork;
magmaFloatComplex *work;
magma_int_t i, k, ldwork, lddwork, old_i, old_ib, rows;
magma_int_t nbmin, nx, ib, nb;
magma_int_t lhwork, lwork;
*info = 0;
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (ldda < max(1,m)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
k = min(m,n);
if (k == 0)
return MAGMA_SUCCESS;
nb = magma_get_cgeqrf_nb(m);
lwork = (m+n) * nb;
lhwork = lwork - (m)*nb;
if ( MAGMA_SUCCESS != magma_cmalloc( &dwork, n*nb )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/*
if ( MAGMA_SUCCESS != magma_cmalloc_cpu( &work, lwork ) ) {
*info = MAGMA_ERR_HOST_ALLOC;
magma_free( dwork );
return *info;
}
*/
cl_mem buffer = clCreateBuffer(gContext, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,
sizeof(magmaFloatComplex)*lwork, NULL, NULL);
work = (magmaFloatComplex*)clEnqueueMapBuffer(queue[0], buffer, CL_TRUE,
CL_MAP_READ | CL_MAP_WRITE,
0, lwork*sizeof(magmaFloatComplex),
0, NULL, NULL, NULL);
nbmin = 2;
nx = nb;
ldwork = m;
lddwork= n;
if (nb >= nbmin && nb < k && nx < k) {
/* Use blocked code initially */
old_i = 0; old_ib = nb;
for (i = 0; i < k-nx; i += nb) {
ib = min(k-i, nb);
rows = m -i;
magma_queue_sync( queue[1] );
chk(magma_cgetmatrix_async(rows, ib, dA(i, i), ldda, work_ref(i), ldwork, queue[0], NULL));
if (i>0){
/* Apply H' to A(i:m,i+2*ib:n) from the left */
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
m-old_i, n-old_i-2*old_ib, old_ib,
dA(old_i, old_i ), ldda, dwork,0, lddwork,
dA(old_i, old_i+2*old_ib), ldda, dwork,old_ib, lddwork, queue[1]);
chk(magma_csetmatrix_async( old_ib, old_ib, work_ref(old_i), ldwork,
dA(old_i, old_i), ldda, queue[1], NULL));
}
magma_queue_sync(queue[0]);
lapackf77_cgeqrf(&rows, &ib, work_href(i), &ldwork, tau+i, hwork, &lhwork, info);
/* Form the triangular factor of the block reflector
H = H(i) H(i+1) . . . H(i+ib-1) */
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&rows, &ib,
work_href(i), &ldwork, tau+i, hwork, &ib);
cpanel_to_q( MagmaUpper, ib, work_href(i), ldwork, hwork+ib*ib );
/* download the i-th V matrix */
chk(magma_csetmatrix_async(rows, ib, work_ref(i), ldwork, dA(i,i), ldda, queue[0], NULL));
/* download the T matrix */
magma_queue_sync( queue[1] );
chk(magma_csetmatrix_async( ib, ib, hhwork, ib, dwork, 0, lddwork, queue[0], NULL));
magma_queue_sync( queue[0] );
if (i + ib < n)
{
if (i+nb < k-nx) {
/* Apply H' to A(i:m,i+ib:i+2*ib) from the left */
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, ib, ib,
dA(i, i ), ldda, dwork,0, lddwork,
dA(i, i+ib), ldda, dwork,ib, lddwork, queue[1]);
cq_to_panel( MagmaUpper, ib, work_href(i), ldwork, hwork+ib*ib );
}
else {
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, n-i-ib, ib,
dA(i, i ), ldda, dwork,0, lddwork,
dA(i, i+ib), ldda, dwork,ib, lddwork, queue[1]);
cq_to_panel( MagmaUpper, ib, work_href(i), ldwork, hwork+ib*ib );
chk(magma_csetmatrix_async(ib, ib, work_ref(i), ldwork, dA(i,i), ldda, queue[1], NULL));
}
old_i = i;
old_ib = ib;
}
}
} else {
i = 0;
}
magma_free(dwork);
/* Use unblocked code to factor the last or only block. */
if (i < k) {
ib = n-i;
rows = m-i;
magma_cgetmatrix_async(rows, ib, dA(i, i), ldda, work, 0, rows, queue[1], NULL);
magma_queue_sync(queue[1]);
lhwork = lwork - rows*ib;
lapackf77_cgeqrf(&rows, &ib, work, &rows, tau+i, work+ib*rows, &lhwork, info);
magma_csetmatrix_async(rows, ib, work, 0, rows, dA(i, i), ldda, queue[1], NULL);
}
magma_queue_sync(queue[0]);
magma_queue_sync(queue[1]);
// magma_free_cpu(work);
clEnqueueUnmapMemObject(queue[0], buffer, work, 0, NULL, NULL);
clReleaseMemObject(buffer);
return *info;
} /* magma_cgeqrf2_gpu */
/**
Purpose
-------
CGEQRF computes a QR factorization of a COMPLEX M-by-N matrix A:
A = Q * R. This version does not require work space on the GPU
passed as input. GPU memory is allocated in the routine.
If the current stream is NULL, this version replaces it with user defined
stream to overlap computation with communication.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
A COMPLEX array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, the elements on and above the diagonal of the array
contain the min(M,N)-by-N upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of min(m,n) elementary reflectors (see Further
Details).
\n
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[out]
tau COMPLEX array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
@param[out]
work (workspace) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
\n
Higher performance is achieved if WORK is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lwork INTEGER
The dimension of the array WORK. LWORK >= max( N*NB, 2*NB*NB ),
where NB can be obtained through magma_get_cgeqrf_nb(M).
\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.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
Further Details
---------------
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
@ingroup magma_cgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_cgeqrf(magma_int_t m, magma_int_t n,
magmaFloatComplex *A, magma_int_t lda, magmaFloatComplex *tau,
magmaFloatComplex *work, magma_int_t lwork,
magma_int_t *info )
{
#define A(i,j) ( A + (i) + (j)*lda )
#define dA(i,j) (dA + (i) + (j)*ldda)
magmaFloatComplex *dA, *dwork, *dT;
magmaFloatComplex c_one = MAGMA_C_ONE;
magma_int_t i, k, lddwork, old_i, old_ib;
magma_int_t ib, ldda;
/* Function Body */
*info = 0;
magma_int_t nb = magma_get_cgeqrf_nb(min(m, n));
// need 2*nb*nb to store T and upper triangle of V simultaneously
magma_int_t lwkopt = max(n*nb, 2*nb*nb);
work[0] = MAGMA_C_MAKE( (float)lwkopt, 0 );
int lquery = (lwork == -1);
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,m)) {
*info = -4;
} else if (lwork < max(1, lwkopt) && ! lquery) {
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
k = min(m,n);
if (k == 0) {
work[0] = c_one;
return *info;
}
// largest N for larfb is n-nb (trailing matrix lacks 1st panel)
lddwork = ((n+31)/32)*32 - nb;
ldda = ((m+31)/32)*32;
magma_int_t num_gpus = magma_num_gpus();
if ( num_gpus > 1 ) {
/* call multiple-GPU interface */
return magma_cgeqrf4(num_gpus, m, n, A, lda, tau, work, lwork, info);
}
// allocate space for dA, dwork, and dT
if (MAGMA_SUCCESS != magma_cmalloc( &dA, n*ldda + nb*lddwork + nb*nb )) {
/* Switch to the "out-of-core" (out of GPU-memory) version */
return magma_cgeqrf_ooc(m, n, A, lda, tau, work, lwork, info);
}
/* Define user stream if current stream is NULL */
magma_queue_t stream[2], current_stream;
magmablasGetKernelStream(¤t_stream);
magma_queue_create( &stream[0] );
if (current_stream == NULL) {
magma_queue_create( &stream[1] );
magmablasSetKernelStream(stream[1]);
}
else {
stream[1] = current_stream;
}
dwork = dA + n*ldda;
dT = dA + n*ldda + nb*lddwork;
if ( (nb > 1) && (nb < k) ) {
/* Use blocked code initially.
Asynchronously send the matrix to the GPU except the first panel. */
magma_csetmatrix_async( m, n-nb,
A(0,nb), lda,
dA(0,nb), ldda, stream[0] );
old_i = 0;
old_ib = nb;
for (i = 0; i < k-nb; i += nb) {
ib = min(k-i, nb);
if (i > 0) {
/* download i-th panel */
magma_queue_sync( stream[1] );
magma_cgetmatrix_async( m-i, ib,
dA(i,i), ldda,
A(i,i), lda, stream[0] );
/* Apply H' to A(i:m,i+2*ib:n) from the left */
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
m-old_i, n-old_i-2*old_ib, old_ib,
dA(old_i, old_i), ldda, dT, nb,
dA(old_i, old_i+2*old_ib), ldda, dwork, lddwork);
magma_cgetmatrix_async( i, ib,
dA(0,i), ldda,
A(0,i), lda, stream[1] );
magma_queue_sync( stream[0] );
}
magma_int_t rows = m-i;
lapackf77_cgeqrf(&rows, &ib, A(i,i), &lda, tau+i, work, &lwork, info);
/* Form the triangular factor of the block reflector
H = H(i) H(i+1) . . . H(i+ib-1) */
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&rows, &ib, A(i,i), &lda, tau+i, work, &ib);
cpanel_to_q(MagmaUpper, ib, A(i,i), lda, work+ib*ib);
/* download the i-th V matrix */
magma_csetmatrix_async( rows, ib, A(i,i), lda, dA(i,i), ldda, stream[0] );
/* download the T matrix */
magma_queue_sync( stream[1] );
magma_csetmatrix_async( ib, ib, work, ib, dT, nb, stream[0] );
magma_queue_sync( stream[0] );
if (i + ib < n) {
if (i+ib < k-nb) {
/* Apply H' to A(i:m,i+ib:i+2*ib) from the left (look-ahead) */
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, ib, ib,
dA(i, i ), ldda, dT, nb,
dA(i, i+ib), ldda, dwork, lddwork);
cq_to_panel(MagmaUpper, ib, A(i,i), lda, work+ib*ib);
}
else {
/* After last panel, update whole trailing matrix. */
/* Apply H' to A(i:m,i+ib:n) from the left */
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, n-i-ib, ib,
dA(i, i ), ldda, dT, nb,
dA(i, i+ib), ldda, dwork, lddwork);
cq_to_panel(MagmaUpper, ib, A(i,i), lda, work+ib*ib);
}
old_i = i;
old_ib = ib;
}
}
} else {
i = 0;
}
/* Use unblocked code to factor the last or only block. */
if (i < k) {
ib = n-i;
if (i != 0) {
magma_cgetmatrix_async( m, ib, dA(0,i), ldda, A(0,i), lda, stream[1] );
magma_queue_sync( stream[1] );
}
magma_int_t rows = m-i;
lapackf77_cgeqrf(&rows, &ib, A(i,i), &lda, tau+i, work, &lwork, info);
}
magma_queue_destroy( stream[0] );
if (current_stream == NULL) {
magma_queue_destroy( stream[1] );
magmablasSetKernelStream(NULL);
}
magma_free( dA );
return *info;
} /* magma_cgeqrf */
extern "C" magma_int_t
magma_cgeqrf2_mgpu( magma_int_t num_gpus, magma_int_t m, magma_int_t n,
magmaFloatComplex **dlA, magma_int_t ldda,
magmaFloatComplex *tau,
magma_int_t *info )
{
/* -- MAGMA (version 1.4.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
December 2013
Purpose
=======
CGEQRF2_MGPU computes a QR factorization of a complex M-by-N matrix A:
A = Q * R. This is a GPU interface of the routine.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. N >= 0.
dA (input/output) COMPLEX array on the GPU, dimension (LDDA,N)
On entry, the M-by-N matrix dA.
On exit, the elements on and above the diagonal of the array
contain the min(M,N)-by-N upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of min(m,n) elementary reflectors (see Further
Details).
LDDA (input) INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
To benefit from coalescent memory accesses LDDA must be
dividable by 16.
TAU (output) COMPLEX array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
Further Details
===============
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
===================================================================== */
#define dlA(dev, i, j) (dlA[dev] + (i) + (j)*(ldda))
#define hpanel(i) (hpanel + (i))
// set to NULL to make cleanup easy: free(NULL) does nothing.
magmaFloatComplex *dwork[MagmaMaxGPUs]={NULL}, *dpanel[MagmaMaxGPUs]={NULL};
magmaFloatComplex *hwork=NULL, *hpanel=NULL;
magma_queue_t stream[MagmaMaxGPUs][2]={{NULL}};
magma_event_t panel_event[MagmaMaxGPUs]={NULL};
magma_int_t i, j, min_mn, dev, ldhpanel, lddwork, rows;
magma_int_t ib, nb;
magma_int_t lhwork, lwork;
magma_int_t panel_dev, i_local, i_nb_local, n_local[MagmaMaxGPUs], la_dev, dpanel_offset;
magma_queue_t cqueue;
magmablasGetKernelStream( &cqueue );
magma_device_t cdevice;
magma_getdevice( &cdevice );
*info = 0;
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (ldda < max(1,m)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
min_mn = min(m,n);
if (min_mn == 0)
return *info;
nb = magma_get_cgeqrf_nb( m );
/* dwork is (n*nb) --- for T (nb*nb) and clarfb work ((n-nb)*nb) ---
* + dpanel (ldda*nb), on each GPU.
* I think clarfb work could be smaller, max(n_local[:]).
* Oddly, T and clarfb work get stacked on top of each other, both with lddwork=n.
* on GPU that owns panel, set dpanel = dlA(dev,i,i_local).
* on other GPUs, set dpanel = dwork[dev] + dpanel_offset. */
lddwork = n;
dpanel_offset = lddwork*nb;
for( dev=0; dev < num_gpus; dev++ ) {
magma_setdevice( dev );
if ( MAGMA_SUCCESS != magma_cmalloc( &(dwork[dev]), (lddwork + ldda)*nb )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
goto CLEANUP;
}
}
/* hwork is MAX( workspace for cgeqrf (n*nb), two copies of T (2*nb*nb) )
* + hpanel (m*nb).
* for last block, need 2*n*nb total. */
ldhpanel = m;
lhwork = max( n*nb, 2*nb*nb );
lwork = max( lhwork + ldhpanel*nb, 2*n*nb );
if ( MAGMA_SUCCESS != magma_cmalloc_pinned( &hwork, lwork )) {
*info = MAGMA_ERR_HOST_ALLOC;
goto CLEANUP;
}
hpanel = hwork + lhwork;
/* Set the number of local n for each GPU */
for( dev=0; dev < num_gpus; dev++ ) {
n_local[dev] = ((n/nb)/num_gpus)*nb;
if (dev < (n/nb) % num_gpus)
n_local[dev] += nb;
else if (dev == (n/nb) % num_gpus)
n_local[dev] += n % nb;
}
for( dev=0; dev < num_gpus; dev++ ) {
magma_setdevice( dev );
magma_queue_create( &stream[dev][0] );
magma_queue_create( &stream[dev][1] );
magma_event_create( &panel_event[dev] );
}
if ( nb < min_mn ) {
/* Use blocked code initially */
// Note: as written, ib cannot be < nb.
for( i = 0; i < min_mn-nb; i += nb ) {
/* Set the GPU number that holds the current panel */
panel_dev = (i/nb) % num_gpus;
/* Set the local index where the current panel is (j==i) */
i_local = i/(nb*num_gpus)*nb;
ib = min(min_mn-i, nb);
rows = m-i;
/* Send current panel to the CPU, after panel_event indicates it has been updated */
magma_setdevice( panel_dev );
magma_queue_wait_event( stream[panel_dev][1], panel_event[panel_dev] );
magma_cgetmatrix_async( rows, ib,
dlA(panel_dev, i, i_local), ldda,
hpanel(i), ldhpanel, stream[panel_dev][1] );
magma_queue_sync( stream[panel_dev][1] );
// Factor panel
lapackf77_cgeqrf( &rows, &ib, hpanel(i), &ldhpanel, tau+i,
hwork, &lhwork, info );
if ( *info != 0 ) {
fprintf( stderr, "error %d\n", (int) *info );
}
// Form the triangular factor of the block reflector
// H = H(i) H(i+1) . . . H(i+ib-1)
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&rows, &ib,
hpanel(i), &ldhpanel, tau+i, hwork, &ib );
cpanel_to_q( MagmaUpper, ib, hpanel(i), ldhpanel, hwork + ib*ib );
// Send the current panel back to the GPUs
for( dev=0; dev < num_gpus; dev++ ) {
magma_setdevice( dev );
if (dev == panel_dev)
dpanel[dev] = dlA(dev, i, i_local);
else
dpanel[dev] = dwork[dev] + dpanel_offset;
magma_csetmatrix_async( rows, ib,
hpanel(i), ldhpanel,
dpanel[dev], ldda, stream[dev][0] );
}
for( dev=0; dev < num_gpus; dev++ ) {
magma_setdevice( dev );
magma_queue_sync( stream[dev][0] );
}
// TODO: if cpanel_to_q copied whole block, wouldn't need to restore
// -- just send the copy to the GPUs.
// TODO: also, could zero out the lower triangle and use Azzam's larfb w/ gemm.
/* Restore the panel */
cq_to_panel( MagmaUpper, ib, hpanel(i), ldhpanel, hwork + ib*ib );
if (i + ib < n) {
/* Send the T matrix to the GPU. */
for( dev=0; dev < num_gpus; dev++ ) {
magma_setdevice( dev );
magma_csetmatrix_async( ib, ib,
hwork, ib,
dwork[dev], lddwork, stream[dev][0] );
}
la_dev = (panel_dev+1) % num_gpus;
for( dev=0; dev < num_gpus; dev++ ) {
magma_setdevice( dev );
magmablasSetKernelStream( stream[dev][0] );
if (dev == la_dev && i+nb < min_mn-nb) {
// If not last panel,
// for look-ahead panel, apply H' to A(i:m,i+ib:i+2*ib)
i_nb_local = (i+nb)/(nb*num_gpus)*nb;
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, ib, ib,
dpanel[dev], ldda, // V
dwork[dev], lddwork, // T
dlA(dev, i, i_nb_local), ldda, // C
dwork[dev]+ib, lddwork ); // work
magma_event_record( panel_event[dev], stream[dev][0] );
// for trailing matrix, apply H' to A(i:m,i+2*ib:n)
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, n_local[dev]-(i_nb_local+ib), ib,
dpanel[dev], ldda, // V
dwork[dev], lddwork, // T
dlA(dev, i, i_nb_local+ib), ldda, // C
dwork[dev]+ib, lddwork ); // work
}
else {
// for trailing matrix, apply H' to A(i:m,i+ib:n)
i_nb_local = i_local;
if (dev <= panel_dev) {
i_nb_local += ib;
}
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, n_local[dev]-i_nb_local, ib,
dpanel[dev], ldda, // V
dwork[dev], lddwork, // T
dlA(dev, i, i_nb_local), ldda, // C
dwork[dev]+ib, lddwork ); // work
}
}
// Restore top of panel (after larfb is done)
magma_setdevice( panel_dev );
magma_csetmatrix_async( ib, ib,
hpanel(i), ldhpanel,
dlA(panel_dev, i, i_local), ldda, stream[panel_dev][0] );
}
}
}
else {
i = 0;
}
/* Use unblocked code to factor the last or only block row. */
if (i < min_mn) {
rows = m-i;
for( j=i; j < n; j += nb ) {
panel_dev = (j/nb) % num_gpus;
i_local = j/(nb*num_gpus)*nb;
ib = min( n-j, nb );
magma_setdevice( panel_dev );
magma_cgetmatrix( rows, ib,
dlA(panel_dev, i, i_local), ldda,
hwork + (j-i)*rows, rows );
}
// needs lwork >= 2*n*nb:
// needs (m-i)*(n-i) for last block row, bounded by nb*n.
// needs (n-i)*nb for cgeqrf work, bounded by n*nb.
ib = n-i; // total columns in block row
lhwork = lwork - ib*rows;
lapackf77_cgeqrf( &rows, &ib, hwork, &rows, tau+i, hwork + ib*rows, &lhwork, info );
if ( *info != 0 ) {
fprintf( stderr, "error %d\n", (int) *info );
}
for( j=i; j < n; j += nb ) {
panel_dev = (j/nb) % num_gpus;
i_local = j/(nb*num_gpus)*nb;
ib = min( n-j, nb );
magma_setdevice( panel_dev );
magma_csetmatrix( rows, ib,
hwork + (j-i)*rows, rows,
dlA(panel_dev, i, i_local), ldda );
}
}
CLEANUP:
// free(NULL) does nothing.
// check that queues and events are non-zero before destroying them, though.
for( dev=0; dev < num_gpus; dev++ ) {
magma_setdevice( dev );
if ( stream[dev][0] ) { magma_queue_destroy( stream[dev][0] ); }
if ( stream[dev][1] ) { magma_queue_destroy( stream[dev][1] ); }
if ( panel_event[dev] ) { magma_event_destroy( panel_event[dev] ); }
magma_free( dwork[dev] );
}
magma_free_pinned( hwork );
magma_setdevice( cdevice );
magmablasSetKernelStream( cqueue );
return *info;
} /* magma_cgeqrf2_mgpu */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cungqr
*/
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;
magmaFloatComplex *hA, *hR, *tau, *h_work;
magmaFloatComplex *dA, *dT;
magma_int_t m, n, k;
magma_int_t n2, lda, ldda, lwork, min_mn, nb, info;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
printf("Running version %d; available are (specified through --version num):\n",
(int) opts.version);
printf("1 - uses precomputed clarft matrices (default)\n");
printf("2 - recomputes the clarft matrices on the fly\n\n");
printf(" m n k CPU GFlop/s (sec) GPU GFlop/s (sec) ||R|| / ||A||\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];
k = opts.ksize[itest];
if ( m < n || n < k ) {
printf( "%5d %5d %5d skipping because m < n or n < k\n", (int) m, (int) n, (int) k );
continue;
}
lda = m;
ldda = ((m + 31)/32)*32;
n2 = lda*n;
min_mn = min(m, n);
nb = magma_get_cgeqrf_nb( m );
lwork = (m + 2*n+nb)*nb;
gflops = FLOPS_CUNGQR( m, n, k ) / 1e9;
TESTING_MALLOC_PIN( h_work, magmaFloatComplex, lwork );
TESTING_MALLOC_PIN( hR, magmaFloatComplex, lda*n );
TESTING_MALLOC_CPU( hA, magmaFloatComplex, lda*n );
TESTING_MALLOC_CPU( tau, magmaFloatComplex, min_mn );
TESTING_MALLOC_DEV( dA, magmaFloatComplex, ldda*n );
TESTING_MALLOC_DEV( dT, magmaFloatComplex, ( 2*min_mn + ((n + 31)/32)*32 )*nb );
lapackf77_clarnv( &ione, ISEED, &n2, hA );
lapackf77_clacpy( MagmaUpperLowerStr, &m, &n, hA, &lda, hR, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
// first, get QR factors
magma_csetmatrix( m, n, hA, lda, dA, ldda );
magma_cgeqrf_gpu( m, n, dA, ldda, tau, dT, &info );
if (info != 0)
printf("magma_cgeqrf_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_cgetmatrix( m, n, dA, ldda, hR, lda );
gpu_time = magma_wtime();
if (opts.version == 1)
magma_cungqr( m, n, k, hR, lda, tau, dT, nb, &info );
else
magma_cungqr2(m, n, k, hR, lda, tau, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cungqr_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
error = lapackf77_clange("f", &m, &n, hA, &lda, work );
lapackf77_cgeqrf( &m, &n, hA, &lda, tau, h_work, &lwork, &info );
if (info != 0)
printf("lapackf77_cgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
cpu_time = magma_wtime();
lapackf77_cungqr( &m, &n, &k, hA, &lda, tau, h_work, &lwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cungqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// compute relative error |R|/|A| := |Q_magma - Q_lapack|/|A|
blasf77_caxpy( &n2, &c_neg_one, hA, &ione, hR, &ione );
error = lapackf77_clange("f", &m, &n, hR, &lda, work) / error;
printf("%5d %5d %5d %7.1f (%7.2f) %7.1f (%7.2f) %8.2e %s\n",
(int) m, (int) n, (int) k,
cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed"));
status += ! (error < tol);
}
else {
printf("%5d %5d %5d --- ( --- ) %7.1f (%7.2f) --- \n",
(int) m, (int) n, (int) k,
gpu_perf, gpu_time );
}
TESTING_FREE_PIN( h_work );
TESTING_FREE_PIN( hR );
TESTING_FREE_CPU( hA );
TESTING_FREE_CPU( tau );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dT );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgegqr
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float e1, e2, e3, e4, e5, *work;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex *h_A, *h_R, *tau, *dtau, *h_work, *h_rwork, tmp[1];
magmaFloatComplex *d_A, *dwork;
magma_int_t M, N, n2, lda, ldda, lwork, info, min_mn;
magma_int_t ione = 1, ldwork;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
// versions 1...4 are valid
if (opts.version < 1 || opts.version > 4) {
printf("Unknown version %d; exiting\n", opts.version );
return -1;
}
float tol, eps = lapackf77_slamch("E");
tol = 10* opts.tolerance * eps;
printf(" M N CPU GFlop/s (ms) GPU GFlop/s (ms) ||I-Q'Q||_F / M ||I-Q'Q||_I / M ||A-Q R||_I\n");
printf(" MAGMA / LAPACK MAGMA / LAPACK\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];
if (N > 128) {
printf("%5d %5d skipping because cgegqr requires N <= 128\n",
(int) M, (int) N);
continue;
}
if (M < N) {
printf("%5d %5d skipping because cgegqr requires M >= N\n",
(int) M, (int) N);
continue;
}
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = FLOPS_CGEQRF( M, N ) / 1e9 + FLOPS_CUNGQR( M, N, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_cgeqrf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_C_REAL( tmp[0] );
lwork = max(lwork, 3*N*N);
ldwork = N*N;
if (opts.version == 2) {
ldwork = 3*N*N + min_mn;
}
TESTING_MALLOC_PIN( tau, magmaFloatComplex, min_mn );
TESTING_MALLOC_PIN( h_work, magmaFloatComplex, lwork );
TESTING_MALLOC_PIN(h_rwork, magmaFloatComplex, lwork );
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( h_R, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( work, float, M );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( dtau, magmaFloatComplex, min_mn );
TESTING_MALLOC_DEV( dwork, magmaFloatComplex, ldwork );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
magma_csetmatrix( M, N, h_R, lda, d_A, ldda );
// warmup
magma_cgegqr_gpu( 1, M, N, d_A, ldda, dwork, h_work, &info );
magma_csetmatrix( M, N, h_R, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_sync_wtime( 0 );
magma_cgegqr_gpu( opts.version, M, N, d_A, ldda, dwork, h_rwork, &info );
gpu_time = magma_sync_wtime( 0 ) - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cgegqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_cgetmatrix( M, N, d_A, ldda, h_R, M );
// Regenerate R
// blasf77_cgemm("t", "n", &N, &N, &M, &c_one, h_R, &M, h_A, &M, &c_zero, h_rwork, &N);
// magma_cprint(N, N, h_work, N);
blasf77_ctrmm("r", "u", "n", "n", &M, &N, &c_one, h_rwork, &N, h_R, &M);
blasf77_caxpy( &n2, &c_neg_one, h_A, &ione, h_R, &ione );
e5 = lapackf77_clange("i", &M, &N, h_R, &M, work) /
lapackf77_clange("i", &M, &N, h_A, &lda, work);
magma_cgetmatrix( M, N, d_A, ldda, h_R, M );
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
/* Orthogonalize on the CPU */
lapackf77_cgeqrf(&M, &N, h_A, &lda, tau, h_work, &lwork, &info);
lapackf77_cungqr(&M, &N, &N, h_A, &lda, tau, h_work, &lwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cungqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
blasf77_cgemm("t", "n", &N, &N, &M, &c_one, h_R, &M, h_R, &M, &c_zero, h_work, &N);
for(int ii = 0; ii < N*N; ii += N+1 ) {
h_work[ii] = MAGMA_C_SUB(h_work[ii], c_one);
}
e1 = lapackf77_clange("f", &N, &N, h_work, &N, work) / N;
e3 = lapackf77_clange("i", &N, &N, h_work, &N, work) / N;
blasf77_cgemm("t", "n", &N, &N, &M, &c_one, h_A, &M, h_A, &M, &c_zero, h_work, &N);
for(int ii = 0; ii < N*N; ii += N+1 ) {
h_work[ii] = MAGMA_C_SUB(h_work[ii], c_one);
}
e2 = lapackf77_clange("f", &N, &N, h_work, &N, work) / N;
e4 = lapackf77_clange("i", &N, &N, h_work, &N, work) / N;
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e / %8.2e %8.2e / %8.2e %8.2e %s\n",
(int) M, (int) N, cpu_perf, 1000.*cpu_time, gpu_perf, 1000.*gpu_time,
e1, e2, e3, e4, e5,
(e1 < tol ? "ok" : "failed"));
status += ! (e1 < tol);
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) M, (int) N, gpu_perf, 1000.*gpu_time );
}
TESTING_FREE_PIN( tau );
TESTING_FREE_PIN( h_work );
TESTING_FREE_PIN( h_rwork );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_R );
TESTING_FREE_CPU( work );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( dtau );
TESTING_FREE_DEV( dwork );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgegqr
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float e1, e2, work[1];
magmaFloatComplex *h_A, *h_R, *tau, *dtau, *h_work, tmp[1];
magmaFloatComplex *d_A, *dwork, *ddA, *d_T;
magma_int_t M, N, n2, lda, ldda, lwork, info, min_mn;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
printf(" M N CPU GFlop/s (ms) GPU GFlop/s (ms) ||I - Q'Q||_F \n");
printf("=======================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[i];
N = opts.nsize[i];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = FLOPS_CGEQRF( M, N ) / 1e9 + FLOPS_CUNGQR( M, N, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_cgeqrf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_C_REAL( tmp[0] );
lwork = max(lwork, 3*N*N);
TESTING_MALLOC_PIN( tau, magmaFloatComplex, min_mn );
TESTING_MALLOC_PIN( h_work, magmaFloatComplex, lwork );
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( h_R, magmaFloatComplex, n2 );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( dtau, magmaFloatComplex, min_mn );
TESTING_MALLOC_DEV( dwork, magmaFloatComplex, N*N );
TESTING_MALLOC_DEV( ddA, magmaFloatComplex, N*N );
TESTING_MALLOC_DEV( d_T, magmaFloatComplex, N*N );
cudaMemset( ddA, 0, N*N*sizeof(magmaFloatComplex) );
cudaMemset( d_T, 0, N*N*sizeof(magmaFloatComplex) );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
magma_csetmatrix( M, N, h_R, lda, d_A, ldda );
// warmup
magma_cgegqr_gpu( M, N, d_A, ldda, dwork, h_work, &info );
magma_csetmatrix( M, N, h_R, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_sync_wtime( 0 );
if (opts.version == 2) {
int min_mn = min(M, N);
int nb = N;
cuFloatComplex *dtau = dwork;
magma_cgeqr2x3_gpu(&M, &N, d_A, &ldda, dtau, d_T, ddA,
(float *)(dwork+min_mn), &info);
magma_cgetmatrix( min_mn, 1, dtau, min_mn, tau, min_mn);
magma_cungqr_gpu( M, N, N, d_A, ldda, tau, d_T, nb, &info );
}
else
magma_cgegqr_gpu( M, N, d_A, ldda, dwork, h_work, &info );
gpu_time = magma_sync_wtime( 0 ) - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cgegqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
/* Orthogonalize on the CPU */
lapackf77_cgeqrf(&M, &N, h_A, &lda, tau, h_work, &lwork, &info);
lapackf77_cungqr(&M, &N, &N, h_A, &lda, tau, h_work, &lwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cungqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_cgetmatrix( M, N, d_A, ldda, h_R, M );
magmaFloatComplex one = MAGMA_C_ONE, zero = MAGMA_C_ZERO;
blasf77_cgemm("t", "n", &N, &N, &M, &one, h_R, &M, h_R, &M, &zero, h_work, &N);
for(int ii=0; ii<N*N; ii+=(N+1)) h_work[ii] = MAGMA_C_SUB(h_work[ii], one);
e1 = lapackf77_clange("f", &N, &N, h_work, &N, work);
blasf77_cgemm("t", "n", &N, &N, &M, &one, h_A, &M, h_A, &M, &zero, h_work, &N);
for(int ii=0; ii<N*N; ii+=(N+1)) h_work[ii] = MAGMA_C_SUB(h_work[ii], one);
e2 = lapackf77_clange("f", &N, &N, h_work, &N, work);
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e\n",
(int) M, (int) N, cpu_perf, 1000.*cpu_time, gpu_perf, 1000.*gpu_time, e1, e2 );
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) M, (int) N, gpu_perf, 1000.*gpu_time );
}
TESTING_FREE_PIN( tau );
TESTING_FREE_PIN( h_work );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( dtau );
TESTING_FREE_DEV( dwork );
TESTING_FREE_DEV( ddA );
TESTING_FREE_DEV( d_T );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/***************************************************************************//**
Purpose
-------
CGEQRF computes a QR factorization of a complex M-by-N matrix A:
A = Q * R. This is a GPU interface of the routine.
Arguments
---------
@param[in]
ngpu INTEGER
Number of GPUs to use. ngpu > 0.
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
dlA COMPLEX array of pointers on the GPU, dimension (ngpu).
On entry, the M-by-N matrix A distributed over GPUs
(d_lA[d] points to the local matrix on d-th GPU).
It uses 1D block column cyclic format with the block size of nb,
and each local matrix is stored by column.
On exit, the elements on and above the diagonal of the array
contain the min(M,N)-by-N upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of min(m,n) elementary reflectors (see Further
Details).
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
@param[out]
tau COMPLEX array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
Further Details
---------------
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
@ingroup magma_geqrf
*******************************************************************************/
extern "C" magma_int_t
magma_cgeqrf2_mgpu(
magma_int_t ngpu,
magma_int_t m, magma_int_t n,
magmaFloatComplex_ptr dlA[], magma_int_t ldda,
magmaFloatComplex *tau,
magma_int_t *info )
{
#define dlA(dev, i, j) (dlA[dev] + (i) + (j)*(ldda))
#define hpanel(i) (hpanel + (i))
// set to NULL to make cleanup easy: free(NULL) does nothing.
magmaFloatComplex *dwork[MagmaMaxGPUs]={NULL}, *dpanel[MagmaMaxGPUs]={NULL};
magmaFloatComplex *hwork=NULL, *hpanel=NULL;
magma_queue_t queues[MagmaMaxGPUs][2]={{NULL}};
magma_event_t panel_event[MagmaMaxGPUs]={NULL};
magma_int_t i, j, min_mn, dev, ldhpanel, lddwork, rows;
magma_int_t ib, nb;
magma_int_t lhwork, lwork;
magma_int_t panel_dev, i_local, i_nb_local, n_local[MagmaMaxGPUs], la_dev, dpanel_offset;
*info = 0;
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (ldda < max(1,m)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
min_mn = min(m,n);
if (min_mn == 0)
return *info;
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
nb = magma_get_cgeqrf_nb( m, n );
/* dwork is (n*nb) --- for T (nb*nb) and clarfb work ((n-nb)*nb) ---
* + dpanel (ldda*nb), on each GPU.
* I think clarfb work could be smaller, max(n_local[:]).
* Oddly, T and clarfb work get stacked on top of each other, both with lddwork=n.
* on GPU that owns panel, set dpanel = dlA(dev,i,i_local).
* on other GPUs, set dpanel = dwork[dev] + dpanel_offset. */
lddwork = n;
dpanel_offset = lddwork*nb;
for( dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
if ( MAGMA_SUCCESS != magma_cmalloc( &(dwork[dev]), (lddwork + ldda)*nb )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
goto CLEANUP;
}
}
/* hwork is MAX( workspace for cgeqrf (n*nb), two copies of T (2*nb*nb) )
* + hpanel (m*nb).
* for last block, need 2*n*nb total. */
ldhpanel = m;
lhwork = max( n*nb, 2*nb*nb );
lwork = max( lhwork + ldhpanel*nb, 2*n*nb );
if ( MAGMA_SUCCESS != magma_cmalloc_pinned( &hwork, lwork )) {
*info = MAGMA_ERR_HOST_ALLOC;
goto CLEANUP;
}
hpanel = hwork + lhwork;
/* Set the number of local n for each GPU */
for( dev=0; dev < ngpu; dev++ ) {
n_local[dev] = ((n/nb)/ngpu)*nb;
if (dev < (n/nb) % ngpu)
n_local[dev] += nb;
else if (dev == (n/nb) % ngpu)
n_local[dev] += n % nb;
}
for( dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
magma_queue_create( dev, &queues[dev][0] );
magma_queue_create( dev, &queues[dev][1] );
magma_event_create( &panel_event[dev] );
}
if ( nb < min_mn ) {
/* Use blocked code initially */
// Note: as written, ib cannot be < nb.
for( i = 0; i < min_mn-nb; i += nb ) {
/* Set the GPU number that holds the current panel */
panel_dev = (i/nb) % ngpu;
/* Set the local index where the current panel is (j == i) */
i_local = i/(nb*ngpu)*nb;
ib = min(min_mn-i, nb);
rows = m-i;
/* Send current panel to the CPU, after panel_event indicates it has been updated */
magma_setdevice( panel_dev );
magma_queue_wait_event( queues[panel_dev][1], panel_event[panel_dev] );
magma_cgetmatrix_async( rows, ib,
dlA(panel_dev, i, i_local), ldda,
hpanel(i), ldhpanel,
queues[panel_dev][1] );
magma_queue_sync( queues[panel_dev][1] );
// Factor panel
lapackf77_cgeqrf( &rows, &ib, hpanel(i), &ldhpanel, tau+i,
hwork, &lhwork, info );
if ( *info != 0 ) {
fprintf( stderr, "error %lld\n", (long long) *info );
}
// Form the triangular factor of the block reflector
// H = H(i) H(i+1) . . . H(i+ib-1)
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&rows, &ib,
hpanel(i), &ldhpanel, tau+i, hwork, &ib );
magma_cpanel_to_q( MagmaUpper, ib, hpanel(i), ldhpanel, hwork + ib*ib );
// Send the current panel back to the GPUs
for( dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
if (dev == panel_dev)
dpanel[dev] = dlA(dev, i, i_local);
else
dpanel[dev] = dwork[dev] + dpanel_offset;
magma_csetmatrix_async( rows, ib,
hpanel(i), ldhpanel,
dpanel[dev], ldda,
queues[dev][0] );
}
for( dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
magma_queue_sync( queues[dev][0] );
}
// TODO: if magma_cpanel_to_q copied whole block, wouldn't need to restore
// -- just send the copy to the GPUs.
// TODO: also, could zero out the lower triangle and use Azzam's larfb w/ gemm.
/* Restore the panel */
magma_cq_to_panel( MagmaUpper, ib, hpanel(i), ldhpanel, hwork + ib*ib );
if (i + ib < n) {
/* Send the T matrix to the GPU. */
for( dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
magma_csetmatrix_async( ib, ib,
hwork, ib,
dwork[dev], lddwork,
queues[dev][0] );
}
la_dev = (panel_dev+1) % ngpu;
for( dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
if (dev == la_dev && i+nb < min_mn-nb) {
// If not last panel,
// for look-ahead panel, apply H' to A(i:m,i+ib:i+2*ib)
i_nb_local = (i+nb)/(nb*ngpu)*nb;
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, ib, ib,
dpanel[dev], ldda, // V
dwork[dev], lddwork, // T
dlA(dev, i, i_nb_local), ldda, // C
dwork[dev]+ib, lddwork, // work
queues[dev][0] );
magma_event_record( panel_event[dev], queues[dev][0] );
// for trailing matrix, apply H' to A(i:m,i+2*ib:n)
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, n_local[dev]-(i_nb_local+ib), ib,
dpanel[dev], ldda, // V
dwork[dev], lddwork, // T
dlA(dev, i, i_nb_local+ib), ldda, // C
dwork[dev]+ib, lddwork, // work
queues[dev][0] );
}
else {
// for trailing matrix, apply H' to A(i:m,i+ib:n)
i_nb_local = i_local;
if (dev <= panel_dev) {
i_nb_local += ib;
}
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, n_local[dev]-i_nb_local, ib,
dpanel[dev], ldda, // V
dwork[dev], lddwork, // T
dlA(dev, i, i_nb_local), ldda, // C
dwork[dev]+ib, lddwork, // work
queues[dev][0] );
}
}
// Restore top of panel (after larfb is done)
magma_setdevice( panel_dev );
magma_csetmatrix_async( ib, ib,
hpanel(i), ldhpanel,
dlA(panel_dev, i, i_local), ldda,
queues[panel_dev][0] );
}
}
}
else {
i = 0;
}
/* Use unblocked code to factor the last or only block row. */
if (i < min_mn) {
rows = m-i;
for( j=i; j < n; j += nb ) {
panel_dev = (j/nb) % ngpu;
i_local = j/(nb*ngpu)*nb;
ib = min( n-j, nb );
magma_setdevice( panel_dev );
magma_cgetmatrix( rows, ib,
dlA(panel_dev, i, i_local), ldda,
hwork + (j-i)*rows, rows,
queues[panel_dev][0] );
}
// needs lwork >= 2*n*nb:
// needs (m-i)*(n-i) for last block row, bounded by nb*n.
// needs (n-i)*nb for cgeqrf work, bounded by n*nb.
ib = n-i; // total columns in block row
lhwork = lwork - ib*rows;
lapackf77_cgeqrf( &rows, &ib, hwork, &rows, tau+i, hwork + ib*rows, &lhwork, info );
if ( *info != 0 ) {
fprintf( stderr, "error %lld\n", (long long) *info );
}
for( j=i; j < n; j += nb ) {
panel_dev = (j/nb) % ngpu;
i_local = j/(nb*ngpu)*nb;
ib = min( n-j, nb );
magma_setdevice( panel_dev );
magma_csetmatrix( rows, ib,
hwork + (j-i)*rows, rows,
dlA(panel_dev, i, i_local), ldda,
queues[panel_dev][0] );
}
}
CLEANUP:
// free(NULL) does nothing.
for( dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
magma_queue_destroy( queues[dev][0] );
magma_queue_destroy( queues[dev][1] );
magma_event_destroy( panel_event[dev] );
magma_free( dwork[dev] );
}
magma_free_pinned( hwork );
magma_setdevice( orig_dev );
return *info;
} /* magma_cgeqrf2_mgpu */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgeqrf_mgpu
*/
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
float error, work[1];
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *h_A, *h_R, *tau, *h_work, tmp[1];
magmaFloatComplex *d_lA[ MagmaMaxGPUs ];
magma_int_t M, N, n2, lda, ldda, n_local, ngpu;
magma_int_t info, min_mn, nb, lhwork;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1}, ISEED2[4];
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= (opts.check == 2); // check (-c2) implies lapack (-l)
magma_int_t status = 0;
float tol, eps = lapackf77_slamch("E");
tol = opts.tolerance * eps;
printf("ngpu %d\n", (int) opts.ngpu );
if ( opts.check == 1 ) {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R-Q'A||_1 / (M*||A||_1) ||I-Q'Q||_1 / M\n");
printf("================================================================================================\n");
} else {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F /(M*||A||_F)\n");
printf("==========================================================================\n");
}
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[i];
N = opts.nsize[i];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
nb = magma_get_cgeqrf_nb( M );
gflops = FLOPS_CGEQRF( M, N ) / 1e9;
// ngpu must be at least the number of blocks
ngpu = min( opts.ngpu, int((N+nb-1)/nb) );
if ( ngpu < opts.ngpu ) {
printf( " * too many GPUs for the matrix size, using %d GPUs\n", (int) ngpu );
}
// query for workspace size
lhwork = -1;
lapackf77_cgeqrf( &M, &N, NULL, &M, NULL, tmp, &lhwork, &info );
lhwork = (magma_int_t) MAGMA_C_REAL( tmp[0] );
// Allocate host memory for the matrix
TESTING_MALLOC_CPU( tau, magmaFloatComplex, min_mn );
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( h_work, magmaFloatComplex, lhwork );
TESTING_MALLOC_PIN( h_R, magmaFloatComplex, n2 );
// Allocate device memory
for( int dev = 0; dev < ngpu; dev++ ) {
n_local = ((N/nb)/ngpu)*nb;
if (dev < (N/nb) % ngpu)
n_local += nb;
else if (dev == (N/nb) % ngpu)
n_local += N % nb;
magma_setdevice( dev );
TESTING_MALLOC_DEV( d_lA[dev], magmaFloatComplex, ldda*n_local );
}
/* Initialize the matrix */
for ( int j=0; j<4; j++ ) ISEED2[j] = ISEED[j]; // saving seeds
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
magmaFloatComplex *tau2;
TESTING_MALLOC_CPU( tau2, magmaFloatComplex, min_mn );
cpu_time = magma_wtime();
lapackf77_cgeqrf( &M, &N, h_A, &M, tau2, h_work, &lhwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapack_cgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
TESTING_FREE_CPU( tau2 );
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_csetmatrix_1D_col_bcyclic( M, N, h_R, lda, d_lA, ldda, ngpu, nb );
gpu_time = magma_wtime();
magma_cgeqrf2_mgpu( ngpu, M, N, d_lA, ldda, tau, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cgeqrf2 returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_cgetmatrix_1D_col_bcyclic( M, N, d_lA, ldda, h_R, lda, ngpu, nb );
magma_queue_sync( NULL );
if ( opts.check == 1 ) {
/* =====================================================================
Check the result
=================================================================== */
magma_int_t lwork = n2+N;
magmaFloatComplex *h_W1, *h_W2, *h_W3;
float *h_RW, results[2];
TESTING_MALLOC_CPU( h_W1, magmaFloatComplex, n2 ); // Q
TESTING_MALLOC_CPU( h_W2, magmaFloatComplex, n2 ); // R
TESTING_MALLOC_CPU( h_W3, magmaFloatComplex, lwork ); // WORK
TESTING_MALLOC_CPU( h_RW, float, M ); // RWORK
lapackf77_clarnv( &ione, ISEED2, &n2, h_A );
lapackf77_cqrt02( &M, &N, &min_mn, h_A, h_R, h_W1, h_W2, &lda, tau, h_W3, &lwork,
h_RW, results );
results[0] *= eps;
results[1] *= eps;
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time, results[0],results[1] );
printf("%s\n", (results[0] < tol ? "" : " failed"));
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) %8.2e %8.2e",
(int) M, (int) N, gpu_perf, gpu_time, results[0],results[1] );
printf("%s\n", (results[0] < tol ? "" : " failed"));
}
status |= ! (results[0] < tol);
TESTING_FREE_CPU( h_W1 );
TESTING_FREE_CPU( h_W2 );
TESTING_FREE_CPU( h_W3 );
TESTING_FREE_CPU( h_RW );
}
else if ( opts.check == 2 ) {
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
error = lapackf77_clange("f", &M, &N, h_A, &lda, work );
blasf77_caxpy( &n2, &c_neg_one, h_A, &ione, h_R, &ione );
error = lapackf77_clange("f", &M, &N, h_R, &lda, work ) / (min_mn*error);
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time, error );
printf("%s\n", (error < tol ? "" : " failed"));
status |= ! (error < tol);
}
else {
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) ---\n",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time );
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) M, (int) N, gpu_perf, gpu_time);
}
}
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_work );
TESTING_FREE_PIN( h_R );
for( int dev=0; dev < ngpu; dev++ ){
magma_setdevice( dev );
TESTING_FREE_DEV( d_lA[dev] );
}
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cungqr_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;
magmaFloatComplex *hA, *hR, *tau, *h_work;
magmaFloatComplex *dA, *dT;
magma_int_t m, n, k;
magma_int_t n2, lda, ldda, lwork, min_mn, nb, info;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
printf(" m n k CPU GFlop/s (sec) GPU GFlop/s (sec) ||R|| / ||A||\n");
printf("=========================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
m = opts.msize[i];
n = opts.nsize[i];
k = opts.ksize[i];
if ( m < n || n < k ) {
printf( "skipping m %d, n %d, k %d because m < n or n < k\n", (int) m, (int) n, (int) k );
continue;
}
lda = m;
ldda = ((m + 31)/32)*32;
n2 = lda*n;
min_mn = min(m, n);
nb = magma_get_cgeqrf_nb( m );
lwork = (m + 2*n+nb)*nb;
gflops = FLOPS_CUNGQR( m, n, k ) / 1e9;
TESTING_HOSTALLOC( hA, magmaFloatComplex, lda*n );
TESTING_HOSTALLOC( h_work, magmaFloatComplex, lwork );
TESTING_MALLOC( hR, magmaFloatComplex, lda*n );
TESTING_MALLOC( tau, magmaFloatComplex, min_mn );
TESTING_DEVALLOC( dA, magmaFloatComplex, ldda*n );
TESTING_DEVALLOC( dT, magmaFloatComplex, ( 2*min_mn + ((n + 31)/32)*32 )*nb );
lapackf77_clarnv( &ione, ISEED, &n2, hA );
lapackf77_clacpy( MagmaUpperLowerStr, &m, &n, hA, &lda, hR, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_csetmatrix( m, n, hA, lda, dA, ldda );
magma_cgeqrf_gpu( m, n, dA, ldda, tau, dT, &info );
if (info != 0)
printf("magma_cgeqrf_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
gpu_time = magma_wtime();
magma_cungqr_gpu( m, n, k, dA, ldda, tau, dT, nb, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cungqr_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// Get dA back to the CPU to compare with the CPU result.
magma_cgetmatrix( m, n, dA, ldda, hR, lda );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
error = lapackf77_clange("f", &m, &n, hA, &lda, work );
lapackf77_cgeqrf( &m, &n, hA, &lda, tau, h_work, &lwork, &info );
if (info != 0)
printf("lapackf77_cgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
cpu_time = magma_wtime();
lapackf77_cungqr( &m, &n, &k, hA, &lda, tau, h_work, &lwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cungqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// compute relative error |R|/|A| := |Q_magma - Q_lapack|/|A|
blasf77_caxpy( &n2, &c_neg_one, hA, &ione, hR, &ione );
error = lapackf77_clange("f", &m, &n, hR, &lda, work) / error;
printf("%5d %5d %5d %7.1f (%7.2f) %7.1f (%7.2f) %8.2e\n",
(int) m, (int) n, (int) k,
cpu_perf, cpu_time, gpu_perf, gpu_time, error );
}
else {
printf("%5d %5d %5d --- ( --- ) %7.1f (%7.2f) --- \n",
(int) m, (int) n, (int) k,
gpu_perf, gpu_time );
}
TESTING_HOSTFREE( hA );
TESTING_HOSTFREE( h_work );
TESTING_FREE( hR );
TESTING_FREE( tau );
TESTING_DEVFREE( dA );
TESTING_DEVFREE( dT );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgeqrf
*/
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;
magmaFloatComplex *h_A, *h_T, *h_R, *tau, *h_work, tmp[1];
magmaFloatComplex *d_A, *d_T, *ddA, *dtau;
magmaFloatComplex *d_A2, *d_T2, *ddA2, *dtau2;
float *dwork, *dwork2;
magma_int_t M, N, lda, ldda, lwork, n2, info, min_mn;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
#define BLOCK_SIZE 64
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = 10. * opts.tolerance * lapackf77_slamch("E");
magma_queue_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
printf("version %d\n", (int) opts.version );
printf(" M N CPU GFlop/s (ms) GPU GFlop/s (ms) ||R||_F/||A||_F ||R_T||\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];
if (N > 128) {
printf("%5d %5d skipping because cgeqr2x requires N <= 128\n",
(int) M, (int) N);
continue;
}
if (M < N) {
printf("%5d %5d skipping because cgeqr2x requires M >= N\n",
(int) M, (int) N);
continue;
}
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = (FLOPS_CGEQRF( M, N ) + FLOPS_CGEQRT( M, N )) / 1e9;
/* Allocate memory for the matrix */
TESTING_MALLOC_CPU( tau, magmaFloatComplex, min_mn );
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( h_T, magmaFloatComplex, N*N );
TESTING_MALLOC_PIN( h_R, magmaFloatComplex, n2 );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( d_T, magmaFloatComplex, N*N );
TESTING_MALLOC_DEV( ddA, magmaFloatComplex, N*N );
TESTING_MALLOC_DEV( dtau, magmaFloatComplex, min_mn );
TESTING_MALLOC_DEV( d_A2, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( d_T2, magmaFloatComplex, N*N );
TESTING_MALLOC_DEV( ddA2, magmaFloatComplex, N*N );
TESTING_MALLOC_DEV( dtau2, magmaFloatComplex, min_mn );
TESTING_MALLOC_DEV( dwork, float, max(5*min_mn, (BLOCK_SIZE*2+2)*min_mn) );
TESTING_MALLOC_DEV( dwork2, float, max(5*min_mn, (BLOCK_SIZE*2+2)*min_mn) );
// todo replace with magma_claset
cudaMemset(ddA, 0, N*N*sizeof(magmaFloatComplex));
cudaMemset(d_T, 0, N*N*sizeof(magmaFloatComplex));
cudaMemset(ddA2, 0, N*N*sizeof(magmaFloatComplex));
cudaMemset(d_T2, 0, N*N*sizeof(magmaFloatComplex));
lwork = -1;
lapackf77_cgeqrf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_C_REAL( tmp[0] );
lwork = max(lwork, N*N);
TESTING_MALLOC_CPU( h_work, magmaFloatComplex, lwork );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
magma_csetmatrix( M, N, h_R, lda, d_A, ldda );
magma_csetmatrix( M, N, h_R, lda, d_A2, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_sync_wtime(0);
if (opts.version == 1)
magma_cgeqr2x_gpu(M, N, d_A, ldda, dtau, d_T, ddA, dwork, &info);
else if (opts.version == 2)
magma_cgeqr2x2_gpu(M, N, d_A, ldda, dtau, d_T, ddA, dwork, &info);
else if (opts.version == 3)
magma_cgeqr2x3_gpu(M, N, d_A, ldda, dtau, d_T, ddA, dwork, &info);
else {
printf( "call magma_cgeqr2x4_gpu\n" );
/*
Going through NULL stream is faster
Going through any stream is slower
Doing two streams in parallel is slower than doing them sequentially
Queuing happens on the NULL stream - user defined buffers are smaller?
*/
magma_cgeqr2x4_gpu(M, N, d_A, ldda, dtau, d_T, ddA, dwork, &info, NULL);
//magma_cgeqr2x4_gpu(M, N, d_A, ldda, dtau, d_T, ddA, dwork, &info, stream[1]);
//magma_cgeqr2x4_gpu(M, N, d_A2, ldda, dtau2, d_T2, ddA2, dwork2, &info, stream[0]);
//magma_cgeqr2x4_gpu(M, N, d_A2, ldda, dtau2, d_T2, ddA2, dwork2, &info, NULL);
//gflops *= 2;
}
gpu_time = magma_sync_wtime(0) - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0) {
printf("magma_cgeqr2x_gpu version %d returned error %d: %s.\n",
(int) opts.version, (int) info, magma_strerror( info ));
}
else {
if ( opts.check ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_cgeqrf(&M, &N, h_A, &lda, tau, h_work, &lwork, &info);
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&M, &N, h_A, &lda, tau, h_work, &N);
//magma_cgeqr2(&M, &N, h_A, &lda, tau, h_work, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_cgetmatrix( M, N, d_A, ldda, h_R, M );
magma_cgetmatrix( N, N, ddA, N, h_T, N );
// Restore the upper triangular part of A before the check
for(int col=0; col < N; col++){
for(int row=0; row <= col; row++)
h_R[row + col*M] = h_T[row + col*N];
}
error = lapackf77_clange("M", &M, &N, h_A, &lda, work);
blasf77_caxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_clange("M", &M, &N, h_R, &lda, work) / (N * error);
// Check if T is the same
magma_cgetmatrix( N, N, d_T, N, h_T, N );
float terr = 0.;
for(int col=0; col < N; col++)
for(int row=0; row <= col; row++)
terr += ( MAGMA_C_ABS(h_work[row + col*N] - h_T[row + col*N])*
MAGMA_C_ABS(h_work[row + col*N] - h_T[row + col*N]) );
terr = magma_ssqrt(terr);
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %s\n",
(int) M, (int) N, cpu_perf, 1000.*cpu_time, gpu_perf, 1000.*gpu_time,
error, terr, (error < tol ? "ok" : "failed") );
status += ! (error < tol);
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) M, (int) N, gpu_perf, 1000.*gpu_time);
}
}
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_T );
TESTING_FREE_CPU( h_work );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_T );
TESTING_FREE_DEV( ddA );
TESTING_FREE_DEV( dtau );
TESTING_FREE_DEV( dwork );
TESTING_FREE_DEV( d_A2 );
TESTING_FREE_DEV( d_T2 );
TESTING_FREE_DEV( ddA2 );
TESTING_FREE_DEV( dtau2 );
TESTING_FREE_DEV( dwork2 );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
TESTING_FINALIZE();
return status;
}
/**
Purpose
-------
CGEQRF3 computes a QR factorization of a complex M-by-N matrix A:
A = Q * R.
This version stores the triangular dT matrices used in
the block QR factorization so that they can be applied directly (i.e.,
without being recomputed) later. As a result, the application
of Q is much faster. Also, the upper triangular matrices for V have 0s
in them. The corresponding parts of the upper triangular R are
stored separately in dT.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
dA COMPLEX array on the GPU, dimension (LDDA,N)
On entry, the M-by-N matrix A.
On exit, the elements on and above the diagonal of the array
contain the min(M,N)-by-N upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of min(m,n) elementary reflectors (see Further
Details).
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
@param[out]
tau COMPLEX array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
@param[out]
dT (workspace) COMPLEX array on the GPU,
dimension (2*MIN(M, N) + ceil(N/32)*32 )*NB,
where NB can be obtained through magma_get_cgeqrf_nb( M, N ).
It starts with a MIN(M,N)*NB block that stores the triangular T
matrices, followed by a MIN(M,N)*NB block that stores
the diagonal blocks of the R matrix.
The rest of the array is used as workspace.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
Further Details
---------------
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v^H
where tau is a complex scalar, and v is a complex vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
@ingroup magma_cgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_cgeqrf3_gpu(
magma_int_t m, magma_int_t n,
magmaFloatComplex_ptr dA, magma_int_t ldda,
magmaFloatComplex *tau,
magmaFloatComplex_ptr dT,
magma_int_t *info )
{
#ifdef HAVE_clBLAS
#define dA(i_, j_) dA, (dA_offset + (i_) + (j_)*(ldda))
#define dT(i_) dT, (dT_offset + (i_)*nb)
#define dR(i_) dT, (dT_offset + ( minmn + (i_))*nb)
#define dwork(i_) dT, (dT_offset + (2*minmn + (i_))*nb)
#else
#define dA(i_, j_) (dA + (i_) + (j_)*(ldda))
#define dT(i_) (dT + (i_)*nb)
#define dR(i_) (dT + ( minmn + (i_))*nb)
#define dwork(i_) (dT + (2*minmn + (i_))*nb)
#endif
magmaFloatComplex *work, *hwork, *R;
magma_int_t cols, i, ib, ldwork, lddwork, lhwork, lwork, minmn, nb, old_i, old_ib, rows;
// check arguments
*info = 0;
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (ldda < max(1,m)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
minmn = min( m, n );
if (minmn == 0)
return *info;
// TODO: use min(m,n), but that affects dT
nb = magma_get_cgeqrf_nb( m, n );
// dT contains 3 blocks:
// dT is minmn*nb
// dR is minmn*nb
// dwork is n*nb
lddwork = n;
// work is m*nb for panel
// hwork is n*nb, and at least nb*nb for T in larft
// R is nb*nb
ldwork = m;
lhwork = max( n*nb, nb*nb );
lwork = ldwork*nb + lhwork + nb*nb;
// last block needs rows*cols for matrix and prefers cols*nb for work
// worst case is n > m*nb, m a small multiple of nb:
// needs n*nb + n > (m+n)*nb
// prefers 2*n*nb, about twice above (m+n)*nb.
i = ((minmn-1)/nb)*nb;
lwork = max( lwork, (m-i)*(n-i) + (n-i)*nb );
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &work, lwork )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
hwork = work + ldwork*nb;
R = work + ldwork*nb + lhwork;
memset( R, 0, nb*nb*sizeof(magmaFloatComplex) );
magma_queue_t queues[2];
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[0] );
magma_queue_create( cdev, &queues[1] );
if ( nb > 1 && nb < minmn ) {
// need nb*nb for T in larft
assert( lhwork >= nb*nb );
// Use blocked code initially
old_i = 0; old_ib = nb;
for (i = 0; i < minmn-nb; i += nb) {
ib = min( minmn-i, nb );
rows = m - i;
// get i-th panel from device
magma_cgetmatrix_async( rows, ib,
dA(i,i), ldda,
work, ldwork, queues[1] );
if (i > 0) {
// Apply H^H to A(i:m,i+2*ib:n) from the left
cols = n - old_i - 2*old_ib;
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
m-old_i, cols, old_ib,
dA(old_i, old_i ), ldda, dT(old_i), nb,
dA(old_i, old_i+2*old_ib), ldda, dwork(0), lddwork, queues[0] );
// Fix the diagonal block
magma_csetmatrix_async( old_ib, old_ib,
R, old_ib,
dR(old_i), old_ib, queues[0] );
}
magma_queue_sync( queues[1] ); // wait to get work(i)
lapackf77_cgeqrf( &rows, &ib, work, &ldwork, &tau[i], hwork, &lhwork, info );
// Form the triangular factor of the block reflector in hwork
// H = H(i) H(i+1) . . . H(i+ib-1)
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&rows, &ib,
work, &ldwork, &tau[i], hwork, &ib );
// wait for previous trailing matrix update (above) to finish with R
magma_queue_sync( queues[0] );
// copy the upper triangle of panel to R and invert it, and
// set the upper triangle of panel (V) to identity
csplit_diag_block( ib, work, ldwork, R );
// send i-th V matrix to device
magma_csetmatrix( rows, ib,
work, ldwork,
dA(i,i), ldda, queues[1] );
if (i + ib < n) {
// send T matrix to device
magma_csetmatrix( ib, ib,
hwork, ib,
dT(i), nb, queues[1] );
if (i+nb < minmn-nb) {
// Apply H^H to A(i:m,i+ib:i+2*ib) from the left
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, ib, ib,
dA(i, i ), ldda, dT(i), nb,
dA(i, i+ib), ldda, dwork(0), lddwork, queues[1] );
// wait for larfb to finish with dwork before larfb in next iteration starts
magma_queue_sync( queues[1] );
}
else {
// Apply H^H to A(i:m,i+ib:n) from the left
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, n-i-ib, ib,
dA(i, i ), ldda, dT(i), nb,
dA(i, i+ib), ldda, dwork(0), lddwork, queues[1] );
// Fix the diagonal block
magma_csetmatrix( ib, ib,
R, ib,
dR(i), ib, queues[1] );
}
old_i = i;
old_ib = ib;
}
}
} else {
i = 0;
}
// Use unblocked code to factor the last or only block.
if (i < minmn) {
rows = m-i;
cols = n-i;
magma_cgetmatrix( rows, cols, dA(i, i), ldda, work, rows, queues[1] );
// see comments for lwork above
lhwork = lwork - rows*cols;
lapackf77_cgeqrf( &rows, &cols, work, &rows, &tau[i], &work[rows*cols], &lhwork, info );
magma_csetmatrix( rows, cols, work, rows, dA(i, i), ldda, queues[1] );
}
magma_queue_destroy( queues[0] );
magma_queue_destroy( queues[1] );
magma_free_pinned( work );
return *info;
} // magma_cgeqrf_gpu
/**
Purpose
-------
CGEQRF4 computes a QR factorization of a COMPLEX M-by-N matrix A:
A = Q * R using multiple GPUs. This version does not require work space on the GPU
passed as input. GPU memory is allocated in the routine.
Arguments
---------
@param[in]
num_gpus INTEGER
The number of GPUs to be used for the factorization.
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
A COMPLEX array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, the elements on and above the diagonal of the array
contain the min(M,N)-by-N upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of min(m,n) elementary reflectors (see Further
Details).
\n
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[out]
tau COMPLEX array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
@param[out]
work (workspace) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
\n
Higher performance is achieved if WORK is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lwork INTEGER
The dimension of the array WORK. LWORK >= N*NB,
where NB can be obtained through magma_get_cgeqrf_nb(M).
\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.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
Further Details
---------------
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
@ingroup magma_cgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_cgeqrf4(magma_int_t num_gpus, magma_int_t m, magma_int_t n,
magmaFloatComplex *A, magma_int_t lda, magmaFloatComplex *tau,
magmaFloatComplex *work, magma_int_t lwork,
magma_int_t *info )
{
magmaFloatComplex *da[MagmaMaxGPUs];
magmaFloatComplex c_one = MAGMA_C_ONE;
int i, k, ldda;
*info = 0;
int nb = magma_get_cgeqrf_nb(min(m, n));
int lwkopt = n * nb;
work[0] = MAGMA_C_MAKE( (float)lwkopt, 0 );
int lquery = (lwork == -1);
if (num_gpus < 0 || num_gpus > MagmaMaxGPUs) {
*info = -1;
} else if (m < 0) {
*info = -2;
} else if (n < 0) {
*info = -3;
} else if (lda < max(1,m)) {
*info = -5;
} else if (lwork < max(1,n) && ! lquery) {
*info = -8;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
k = min(m,n);
if (k == 0) {
work[0] = c_one;
return *info;
}
ldda = ((m+31)/32)*32;
magma_int_t n_local[MagmaMaxGPUs];
for (i=0; i < num_gpus; i++) {
n_local[i] = ((n/nb)/num_gpus)*nb;
if (i < (n/nb)%num_gpus)
n_local[i] += nb;
else if (i == (n/nb)%num_gpus)
n_local[i] += n%nb;
magma_setdevice(i);
// TODO on failure, free previously allocated memory
if (MAGMA_SUCCESS != magma_cmalloc( &da[i], ldda*n_local[i] )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
}
if (m > nb && n > nb) {
/* Copy the matrix to the GPUs in 1D block cyclic distribution */
magma_csetmatrix_1D_col_bcyclic(m, n, A, lda, da, ldda, num_gpus, nb);
/* Factor using the GPU interface */
magma_cgeqrf2_mgpu( num_gpus, m, n, da, ldda, tau, info);
/* Copy the matrix back from the GPUs to the CPU */
magma_cgetmatrix_1D_col_bcyclic(m, n, da, ldda, A, lda, num_gpus, nb);
}
else {
lapackf77_cgeqrf(&m, &n, A, &lda, tau, work, &lwork, info);
}
/* Free the allocated GPU memory */
for (i=0; i < num_gpus; i++) {
magma_setdevice(i);
magma_free( da[i] );
}
return *info;
} /* magma_cgeqrf4 */
extern "C" magma_int_t
magma_cgeqrf4(magma_int_t num_gpus, magma_int_t m, magma_int_t n,
magmaFloatComplex *a, magma_int_t lda, magmaFloatComplex *tau,
magmaFloatComplex *work, magma_int_t lwork,
magma_int_t *info )
{
/* -- MAGMA (version 1.4.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
December 2013
Purpose
=======
CGEQRF4 computes a QR factorization of a COMPLEX M-by-N matrix A:
A = Q * R using multiple GPUs. This version does not require work space on the GPU
passed as input. GPU memory is allocated in the routine.
Arguments
=========
NUM_GPUS
(input) INTEGER
The number of GPUs to be used for the factorization.
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. N >= 0.
A (input/output) COMPLEX array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, the elements on and above the diagonal of the array
contain the min(M,N)-by-N upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of min(m,n) elementary reflectors (see Further
Details).
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,M).
TAU (output) COMPLEX array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
WORK (workspace/output) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
Higher performance is achieved if WORK is in pinned memory, e.g.
allocated using magma_malloc_pinned.
LWORK (input) INTEGER
The dimension of the array WORK. LWORK >= N*NB,
where NB can be obtained through magma_get_cgeqrf_nb(M).
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.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
Further Details
===============
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
===================================================================== */
magmaFloatComplex *da[MagmaMaxGPUs];
magmaFloatComplex c_one = MAGMA_C_ONE;
int i, k, ldda;
*info = 0;
int nb = magma_get_cgeqrf_nb(min(m, n));
int lwkopt = n * nb;
work[0] = MAGMA_C_MAKE( (float)lwkopt, 0 );
int lquery = (lwork == -1);
if (num_gpus <0 || num_gpus > 4) {
*info = -1;
} else if (m < 0) {
*info = -2;
} else if (n < 0) {
*info = -3;
} else if (lda < max(1,m)) {
*info = -5;
} else if (lwork < max(1,n) && ! lquery) {
*info = -8;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
k = min(m,n);
if (k == 0) {
work[0] = c_one;
return *info;
}
ldda = ((m+31)/32)*32;
magma_int_t n_local[MagmaMaxGPUs];
for(i=0; i<num_gpus; i++){
n_local[i] = ((n/nb)/num_gpus)*nb;
if (i < (n/nb)%num_gpus)
n_local[i] += nb;
else if (i == (n/nb)%num_gpus)
n_local[i] += n%nb;
magma_setdevice(i);
// TODO on failure, free previously allocated memory
if (MAGMA_SUCCESS != magma_cmalloc( &da[i], ldda*n_local[i] )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
}
if (m > nb && n > nb) {
/* Copy the matrix to the GPUs in 1D block cyclic distribution */
magma_csetmatrix_1D_col_bcyclic(m, n, a, lda, da, ldda, num_gpus, nb);
/* Factor using the GPU interface */
magma_cgeqrf2_mgpu( num_gpus, m, n, da, ldda, tau, info);
/* Copy the matrix back from the GPUs to the CPU */
magma_cgetmatrix_1D_col_bcyclic(m, n, da, ldda, a, lda, num_gpus, nb);
}
else {
lapackf77_cgeqrf(&m, &n, a, &lda, tau, work, &lwork, info);
}
/* Free the allocated GPU memory */
for(i=0; i<num_gpus; i++){
magma_setdevice(i);
magma_free( da[i] );
}
return *info;
} /* magma_cgeqrf4 */
/**
Purpose
-------
CGEQRF3 computes a QR factorization of a complex M-by-N matrix A:
A = Q * R.
This version stores the triangular dT matrices used in
the block QR factorization so that they can be applied directly (i.e.,
without being recomputed) later. As a result, the application
of Q is much faster. Also, the upper triangular matrices for V have 0s
in them and the corresponding parts of the upper triangular R are
stored separately in dT.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
dA COMPLEX array on the GPU, dimension (LDDA,N)
On entry, the M-by-N matrix A.
On exit, the elements on and above the diagonal of the array
contain the min(M,N)-by-N upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of min(m,n) elementary reflectors (see Further
Details).
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
@param[out]
tau COMPLEX array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
@param[out]
dT (workspace) COMPLEX array on the GPU,
dimension (2*MIN(M, N) + (N+31)/32*32 )*NB,
where NB can be obtained through magma_get_cgeqrf_nb(M).
It starts with MIN(M,N)*NB block that store the triangular T
matrices, followed by the MIN(M,N)*NB block of the diagonal
matrices for the R matrix. The rest of the array is used as workspace.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
Further Details
---------------
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
@ingroup magma_cgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_cgeqrf3_gpu( magma_int_t m, magma_int_t n,
magmaFloatComplex *dA, magma_int_t ldda,
magmaFloatComplex *tau, magmaFloatComplex *dT,
magma_int_t *info )
{
#define dA(a_1,a_2) (dA + (a_2)*(ldda) + (a_1))
#define dT(a_1) (dT + (a_1)*nb)
#define d_ref(a_1) (dT + ( minmn+(a_1))*nb)
#define dd_ref(a_1) (dT + (2*minmn+(a_1))*nb)
#define work(a_1) (work + (a_1))
#define hwork (work + (nb)*(m))
magma_int_t i, k, minmn, old_i, old_ib, rows, cols;
magma_int_t ib, nb;
magma_int_t ldwork, lddwork, lwork, lhwork;
magmaFloatComplex *work, *ut;
/* check arguments */
*info = 0;
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (ldda < max(1,m)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
k = minmn = min(m,n);
if (k == 0)
return *info;
nb = magma_get_cgeqrf_nb(m);
lwork = (m + n + nb)*nb;
lhwork = lwork - m*nb;
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &work, lwork )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
ut = hwork+nb*(n);
memset( ut, 0, nb*nb*sizeof(magmaFloatComplex));
magma_queue_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
ldwork = m;
lddwork= n;
if ( (nb > 1) && (nb < k) ) {
/* Use blocked code initially */
old_i = 0; old_ib = nb;
for (i = 0; i < k-nb; i += nb) {
ib = min(k-i, nb);
rows = m -i;
magma_cgetmatrix_async( rows, ib,
dA(i,i), ldda,
work(i), ldwork, stream[1] );
if (i > 0) {
/* Apply H' to A(i:m,i+2*ib:n) from the left */
cols = n-old_i-2*old_ib;
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
m-old_i, cols, old_ib,
dA(old_i, old_i ), ldda, dT(old_i), nb,
dA(old_i, old_i+2*old_ib), ldda, dd_ref(0), lddwork);
/* store the diagonal */
magma_csetmatrix_async( old_ib, old_ib,
ut, old_ib,
d_ref(old_i), old_ib, stream[0] );
}
magma_queue_sync( stream[1] );
lapackf77_cgeqrf(&rows, &ib, work(i), &ldwork, tau+i, hwork, &lhwork, info);
/* Form the triangular factor of the block reflector
H = H(i) H(i+1) . . . H(i+ib-1) */
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&rows, &ib,
work(i), &ldwork, tau+i, hwork, &ib);
/* Put 0s in the upper triangular part of a panel (and 1s on the
diagonal); copy the upper triangular in ut. */
magma_queue_sync( stream[0] );
csplit_diag_block3(ib, work(i), ldwork, ut);
magma_csetmatrix( rows, ib, work(i), ldwork, dA(i,i), ldda );
if (i + ib < n) {
/* Send the triangular factor T to the GPU */
magma_csetmatrix( ib, ib, hwork, ib, dT(i), nb );
if (i+nb < k-nb) {
/* Apply H' to A(i:m,i+ib:i+2*ib) from the left */
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, ib, ib,
dA(i, i ), ldda, dT(i), nb,
dA(i, i+ib), ldda, dd_ref(0), lddwork);
}
else {
cols = n-i-ib;
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, cols, ib,
dA(i, i ), ldda, dT(i), nb,
dA(i, i+ib), ldda, dd_ref(0), lddwork);
/* Fix the diagonal block */
magma_csetmatrix( ib, ib, ut, ib, d_ref(i), ib );
}
old_i = i;
old_ib = ib;
}
}
} else {
i = 0;
}
/* Use unblocked code to factor the last or only block. */
if (i < k) {
ib = n-i;
rows = m-i;
magma_cgetmatrix( rows, ib, dA(i, i), ldda, work, rows );
lhwork = lwork - rows*ib;
lapackf77_cgeqrf(&rows, &ib, work, &rows, tau+i, work+ib*rows, &lhwork, info);
magma_csetmatrix( rows, ib, work, rows, dA(i, i), ldda );
}
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free_pinned( work );
return *info;
} /* magma_cgeqrf_gpu */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgels
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float gpu_error, cpu_error, error, Anorm, work[1];
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *h_A, *h_A2, *h_B, *h_X, *h_R, *tau, *h_work, tmp[1];
magmaFloatComplex *d_A, *d_B;
magma_int_t M, N, size, nrhs, lda, ldb, ldda, lddb, min_mn, max_mn, nb, info;
magma_int_t lworkgpu, lhwork, lhwork2;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_int_t status = 0;
float tol = opts.tolerance * lapackf77_slamch("E");
nrhs = opts.nrhs;
printf(" ||b-Ax|| / (N||A||) ||dx-x||/(N||A||)\n");
printf(" M N NRHS CPU GFlop/s (sec) GPU GFlop/s (sec) CPU GPU \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];
if ( M < N ) {
printf( "%5d %5d %5d skipping because M < N is not yet supported.\n", (int) M, (int) N, (int) nrhs );
continue;
}
min_mn = min(M, N);
max_mn = max(M, N);
lda = M;
ldb = max_mn;
size = lda*N;
ldda = ((M+31)/32)*32;
lddb = ((max_mn+31)/32)*32;
nb = magma_get_cgeqrf_nb(M);
gflops = (FLOPS_CGEQRF( M, N ) + FLOPS_CGEQRS( M, N, nrhs )) / 1e9;
lworkgpu = (M - N + nb)*(nrhs + nb) + nrhs*nb;
// query for workspace size
lhwork = -1;
lapackf77_cgeqrf(&M, &N, NULL, &M, NULL, tmp, &lhwork, &info);
lhwork2 = (magma_int_t) MAGMA_C_REAL( tmp[0] );
lhwork = -1;
lapackf77_cunmqr( MagmaLeftStr, MagmaConjTransStr,
&M, &nrhs, &min_mn, NULL, &lda, NULL,
NULL, &ldb, tmp, &lhwork, &info);
lhwork = (magma_int_t) MAGMA_C_REAL( tmp[0] );
lhwork = max( max( lhwork, lhwork2 ), lworkgpu );
TESTING_MALLOC_CPU( tau, magmaFloatComplex, min_mn );
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, lda*N );
TESTING_MALLOC_CPU( h_A2, magmaFloatComplex, lda*N );
TESTING_MALLOC_CPU( h_B, magmaFloatComplex, ldb*nrhs );
TESTING_MALLOC_CPU( h_X, magmaFloatComplex, ldb*nrhs );
TESTING_MALLOC_CPU( h_R, magmaFloatComplex, ldb*nrhs );
TESTING_MALLOC_CPU( h_work, magmaFloatComplex, lhwork );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( d_B, magmaFloatComplex, lddb*nrhs );
/* Initialize the matrices */
lapackf77_clarnv( &ione, ISEED, &size, h_A );
lapackf77_clacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_A2, &lda );
// make random RHS
size = M*nrhs;
lapackf77_clarnv( &ione, ISEED, &size, h_B );
lapackf77_clacpy( MagmaUpperLowerStr, &M, &nrhs, h_B, &ldb, h_R, &ldb );
// make consistent RHS
//size = N*nrhs;
//lapackf77_clarnv( &ione, ISEED, &size, h_X );
//blasf77_cgemm( MagmaNoTransStr, MagmaNoTransStr, &M, &nrhs, &N,
// &c_one, h_A, &lda,
// h_X, &ldb,
// &c_zero, h_B, &ldb );
//lapackf77_clacpy( MagmaUpperLowerStr, &M, &nrhs, h_B, &ldb, h_R, &ldb );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_csetmatrix( M, N, h_A, lda, d_A, ldda );
magma_csetmatrix( M, nrhs, h_B, ldb, d_B, lddb );
gpu_time = magma_wtime();
magma_cgels3_gpu( MagmaNoTrans, M, N, nrhs, d_A, ldda,
d_B, lddb, h_work, lworkgpu, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cgels3_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// Get the solution in h_X
magma_cgetmatrix( N, nrhs, d_B, lddb, h_X, ldb );
// compute the residual
blasf77_cgemm( MagmaNoTransStr, MagmaNoTransStr, &M, &nrhs, &N,
&c_neg_one, h_A, &lda,
h_X, &ldb,
&c_one, h_R, &ldb);
Anorm = lapackf77_clange("f", &M, &N, h_A, &lda, work);
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
lapackf77_clacpy( MagmaUpperLowerStr, &M, &nrhs, h_B, &ldb, h_X, &ldb );
cpu_time = magma_wtime();
lapackf77_cgels( MagmaNoTransStr, &M, &N, &nrhs,
h_A, &lda, h_X, &ldb, h_work, &lhwork, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cgels returned error %d: %s.\n",
(int) info, magma_strerror( info ));
blasf77_cgemm( MagmaNoTransStr, MagmaNoTransStr, &M, &nrhs, &N,
&c_neg_one, h_A2, &lda,
h_X, &ldb,
&c_one, h_B, &ldb);
cpu_error = lapackf77_clange("f", &M, &nrhs, h_B, &ldb, work) / (min_mn*Anorm);
gpu_error = lapackf77_clange("f", &M, &nrhs, h_R, &ldb, work) / (min_mn*Anorm);
// error relative to LAPACK
size = M*nrhs;
blasf77_caxpy( &size, &c_neg_one, h_B, &ione, h_R, &ione );
error = lapackf77_clange("f", &M, &nrhs, h_R, &ldb, work) / (min_mn*Anorm);
printf("%5d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %8.2e",
(int) M, (int) N, (int) nrhs,
cpu_perf, cpu_time, gpu_perf, gpu_time, cpu_error, gpu_error, error );
if ( M == N ) {
printf( " %s\n", (gpu_error < tol && error < tol ? "ok" : "failed"));
status += ! (gpu_error < tol && error < tol);
}
else {
printf( " %s\n", (error < tol ? "ok" : "failed"));
status += ! (error < tol);
}
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_A2 );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_X );
TESTING_FREE_CPU( h_R );
TESTING_FREE_CPU( h_work );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/**
Purpose
-------
CGEQRF computes a QR factorization of a COMPLEX M-by-N matrix A:
A = Q * R. This version does not require work space on the GPU
passed as input. GPU memory is allocated in the routine.
This uses 2 queues to overlap communication and computation.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
A COMPLEX array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, the elements on and above the diagonal of the array
contain the min(M,N)-by-N upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of min(m,n) elementary reflectors (see Further
Details).
\n
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[out]
tau COMPLEX array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
@param[out]
work (workspace) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
\n
Higher performance is achieved if WORK is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lwork INTEGER
The dimension of the array WORK. LWORK >= max( N*NB, 2*NB*NB ),
where NB can be obtained through magma_get_cgeqrf_nb( M, 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.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
Further Details
---------------
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
@ingroup magma_cgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_cgeqrf(
magma_int_t m, magma_int_t n,
magmaFloatComplex *A, magma_int_t lda,
magmaFloatComplex *tau,
magmaFloatComplex *work, magma_int_t lwork,
magma_int_t *info )
{
#define A(i_,j_) (A + (i_) + (j_)*lda)
#ifdef HAVE_clBLAS
#define dA(i_,j_) dA, ((i_) + (j_)*ldda + dA_offset)
#define dT(i_,j_) dT, ((i_) + (j_)*nb + dT_offset)
#define dwork(i_) dwork, ((i_) + dwork_offset)
#else
#define dA(i_,j_) (dA + (i_) + (j_)*ldda)
#define dT(i_,j_) (dT + (i_) + (j_)*nb)
#define dwork(i_) (dwork + (i_))
#endif
/* Constants */
const magmaFloatComplex c_one = MAGMA_C_ONE;
/* Local variables */
magmaFloatComplex_ptr dA, dT, dwork;
magma_int_t i, ib, min_mn, ldda, lddwork, old_i, old_ib;
/* Function Body */
*info = 0;
magma_int_t nb = magma_get_cgeqrf_nb( m, n );
// need 2*nb*nb to store T and upper triangle of V simultaneously
magma_int_t lwkopt = max( n*nb, 2*nb*nb );
work[0] = magma_cmake_lwork( lwkopt );
bool lquery = (lwork == -1);
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,m)) {
*info = -4;
} else if (lwork < max(1, lwkopt) && ! lquery) {
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
min_mn = min( m, n );
if (min_mn == 0) {
work[0] = c_one;
return *info;
}
// largest N for larfb is n-nb (trailing matrix lacks 1st panel)
lddwork = magma_roundup( n, 32 ) - nb;
ldda = magma_roundup( m, 32 );
magma_int_t ngpu = magma_num_gpus();
if ( ngpu > 1 ) {
/* call multiple-GPU interface */
return magma_cgeqrf_m( ngpu, m, n, A, lda, tau, work, lwork, info );
}
// allocate space for dA, dwork, and dT
if (MAGMA_SUCCESS != magma_cmalloc( &dA, n*ldda + nb*lddwork + nb*nb )) {
/* alloc failed so call non-GPU-resident version */
return magma_cgeqrf_ooc( m, n, A, lda, tau, work, lwork, info );
}
dwork = dA + n*ldda;
dT = dA + n*ldda + nb*lddwork;
magma_queue_t queues[2];
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[0] );
magma_queue_create( cdev, &queues[1] );
if ( (nb > 1) && (nb < min_mn) ) {
/* Use blocked code initially.
Asynchronously send the matrix to the GPU except the first panel. */
magma_csetmatrix_async( m, n-nb,
A(0,nb), lda,
dA(0,nb), ldda, queues[0] );
old_i = 0;
old_ib = nb;
for (i = 0; i < min_mn-nb; i += nb) {
ib = min( min_mn-i, nb );
if (i > 0) {
/* get i-th panel from device */
magma_queue_sync( queues[1] );
magma_cgetmatrix_async( m-i, ib,
dA(i,i), ldda,
A(i,i), lda, queues[0] );
/* Apply H' to A(i:m,i+2*ib:n) from the left */
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
m-old_i, n-old_i-2*old_ib, old_ib,
dA(old_i, old_i), ldda, dT(0,0), nb,
dA(old_i, old_i+2*old_ib), ldda, dwork(0), lddwork, queues[1] );
magma_cgetmatrix_async( i, ib,
dA(0,i), ldda,
A(0,i), lda, queues[1] );
magma_queue_sync( queues[0] );
}
magma_int_t rows = m-i;
lapackf77_cgeqrf( &rows, &ib, A(i,i), &lda, tau+i, work, &lwork, info );
/* Form the triangular factor of the block reflector
H = H(i) H(i+1) . . . H(i+ib-1) */
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&rows, &ib, A(i,i), &lda, tau+i, work, &ib );
magma_cpanel_to_q( MagmaUpper, ib, A(i,i), lda, work+ib*ib );
/* put i-th V matrix onto device */
magma_csetmatrix_async( rows, ib, A(i,i), lda, dA(i,i), ldda, queues[0] );
/* put T matrix onto device */
magma_queue_sync( queues[1] );
magma_csetmatrix_async( ib, ib, work, ib, dT(0,0), nb, queues[0] );
magma_queue_sync( queues[0] );
if (i + ib < n) {
if (i+ib < min_mn-nb) {
/* Apply H' to A(i:m,i+ib:i+2*ib) from the left (look-ahead) */
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, ib, ib,
dA(i, i ), ldda, dT(0,0), nb,
dA(i, i+ib), ldda, dwork(0), lddwork, queues[1] );
magma_cq_to_panel( MagmaUpper, ib, A(i,i), lda, work+ib*ib );
}
else {
/* After last panel, update whole trailing matrix. */
/* Apply H' to A(i:m,i+ib:n) from the left */
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, n-i-ib, ib,
dA(i, i ), ldda, dT(0,0), nb,
dA(i, i+ib), ldda, dwork(0), lddwork, queues[1] );
magma_cq_to_panel( MagmaUpper, ib, A(i,i), lda, work+ib*ib );
}
old_i = i;
old_ib = ib;
}
}
} else {
i = 0;
}
/* Use unblocked code to factor the last or only block. */
if (i < min_mn) {
ib = n-i;
if (i != 0) {
magma_cgetmatrix( m, ib, dA(0,i), ldda, A(0,i), lda, queues[1] );
}
magma_int_t rows = m-i;
lapackf77_cgeqrf( &rows, &ib, A(i,i), &lda, tau+i, work, &lwork, info );
}
magma_queue_destroy( queues[0] );
magma_queue_destroy( queues[1] );
magma_free( dA );
return *info;
} /* magma_cgeqrf */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgeqrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
float error, work[1];
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *h_A, *h_R, *tau, *h_work, tmp[1];
magmaFloatComplex_ptr d_A, dT;
magma_int_t M, N, n2, lda, ldda, lwork, info, min_mn, nb, size;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1}, ISEED2[4];
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_int_t status = 0;
float tol;
opts.lapack |= (opts.version == 2 && opts.check == 2); // check (-c2) implies lapack (-l)
if ( opts.version != 2 && opts.check == 1 ) {
printf( "NOTE: version %d requires -c2 check due to the special structure of the\n"
"MAGMA cgeqrf results; using -c2.\n\n", (int) opts.version );
opts.check = 2;
}
printf( "version %d\n", (int) opts.version );
if ( opts.version == 2 ) {
if ( opts.check == 1 ) {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R-Q'A||_1 / (M*||A||_1*eps) ||I-Q'Q||_1 / (M*eps)\n");
printf("=========================================================================================================\n");
} else {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / ||A||_F\n");
printf("=======================================================================\n");
}
tol = 1.0;
} else {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||Ax-b||_F/(N*||A||_F*||x||_F)\n");
printf("====================================================================================\n");
tol = opts.tolerance * lapackf77_slamch("E");
}
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = FLOPS_CGEQRF( M, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_cgeqrf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_C_REAL( tmp[0] );
TESTING_MALLOC_CPU( tau, magmaFloatComplex, min_mn );
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( h_work, magmaFloatComplex, lwork );
TESTING_MALLOC_PIN( h_R, magmaFloatComplex, n2 );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N );
/* Initialize the matrix */
for ( int j=0; j<4; j++ )
ISEED2[j] = ISEED[j]; // save seeds
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
magma_csetmatrix( M, N, h_R, lda, d_A, 0, ldda, opts.queue );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
if ( opts.version == 2 ) {
magma_cgeqrf2_gpu( M, N, d_A, 0, ldda, tau, opts.queues2, &info );
}
else {
nb = magma_get_cgeqrf_nb( M );
size = (2*min(M, N) + (N+31)/32*32 )*nb;
TESTING_MALLOC_DEV( dT, magmaFloatComplex, size );
if ( opts.version == 1 ) {
magma_cgeqrf_gpu( M, N, d_A, 0, ldda, tau, dT, 0, opts.queue, &info );
}
#ifdef HAVE_CUBLAS
else if ( opts.version == 3 ) {
magma_cgeqrf3_gpu( M, N, d_A, 0, ldda, tau, dT, opts.queue, &info );
}
#endif
else {
printf( "Unknown version %d\n", opts.version );
exit(1);
}
}
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
magmaFloatComplex *tau2;
TESTING_MALLOC_CPU( tau2, magmaFloatComplex, min_mn );
cpu_time = magma_wtime();
lapackf77_cgeqrf(&M, &N, h_A, &lda, tau2, h_work, &lwork, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
TESTING_FREE_CPU( tau2 );
}
if ( opts.check == 1 && M >= N ) {
/* =====================================================================
Check the result -- only version 1, cqrt02 requires M >= N
=================================================================== */
magma_int_t lwork = n2+N;
magmaFloatComplex *h_W1, *h_W2, *h_W3;
float *h_RW, results[2];
magma_cgetmatrix( M, N, d_A, 0, ldda, h_R, M, opts.queue );
TESTING_MALLOC_CPU( h_W1, magmaFloatComplex, n2 ); // Q
TESTING_MALLOC_CPU( h_W2, magmaFloatComplex, n2 ); // R
TESTING_MALLOC_CPU( h_W3, magmaFloatComplex, lwork ); // WORK
TESTING_MALLOC_CPU( h_RW, float, M ); // RWORK
lapackf77_clarnv( &ione, ISEED2, &n2, h_A );
lapackf77_cqrt02( &M, &N, &min_mn, h_A, h_R, h_W1, h_W2, &lda, tau, h_W3, &lwork,
h_RW, results );
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time, results[0], results[1] );
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) %8.2e %8.2e",
(int) M, (int) N, gpu_perf, gpu_time, results[0], results[1] );
}
// todo also check results[1] < tol?
printf(" %s\n", (results[0] < tol ? "ok" : "failed"));
status += ! (results[0] < tol);
TESTING_FREE_CPU( h_W1 );
TESTING_FREE_CPU( h_W2 );
TESTING_FREE_CPU( h_W3 );
TESTING_FREE_CPU( h_RW );
}
else if ( opts.check == 2 && opts.version == 2 ) {
/* =====================================================================
Check the result compared to LAPACK -- only version 2
=================================================================== */
magma_cgetmatrix( M, N, d_A, 0, ldda, h_R, M, opts.queue );
error = lapackf77_clange("f", &M, &N, h_A, &lda, work);
blasf77_caxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_clange("f", &M, &N, h_R, &lda, work) / error;
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time, error );
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) %8.2e",
(int) M, (int) N, gpu_perf, gpu_time, error );
}
printf(" %s\n", (error < tol ? "ok" : "failed"));
status += ! (error < tol);
}
else if ( opts.check == 2 && M >= N ) {
/* =====================================================================
Check the result by solving linear system -- only versions 1 & 3, M >= N
=================================================================== */
magma_int_t lwork;
magmaFloatComplex *x, *b, *hwork;
magmaFloatComplex_ptr d_B;
const magmaFloatComplex c_zero = MAGMA_C_ZERO;
const magmaFloatComplex c_one = MAGMA_C_ONE;
const magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
const magma_int_t ione = 1;
// initialize RHS, b = A*random
TESTING_MALLOC_CPU( x, magmaFloatComplex, N );
TESTING_MALLOC_CPU( b, magmaFloatComplex, M );
lapackf77_clarnv( &ione, ISEED, &N, x );
blasf77_cgemv( "Notrans", &M, &N, &c_one, h_A, &lda, x, &ione, &c_zero, b, &ione );
// copy to GPU
TESTING_MALLOC_DEV( d_B, magmaFloatComplex, M );
magma_csetvector( M, b, 1, d_B, 0, 1, opts.queue );
if ( opts.version == 1 ) {
// allocate hwork
magma_cgeqrs_gpu( M, N, 1,
d_A, 0, ldda, tau, dT, 0,
d_B, 0, M, tmp, -1, opts.queue, &info );
lwork = (magma_int_t)MAGMA_C_REAL( tmp[0] );
TESTING_MALLOC_CPU( hwork, magmaFloatComplex, lwork );
// solve linear system
magma_cgeqrs_gpu( M, N, 1,
d_A, 0, ldda, tau, dT, 0,
d_B, 0, M, hwork, lwork, opts.queue, &info );
if (info != 0)
printf("magma_cgeqrs returned error %d: %s.\n",
(int) info, magma_strerror( info ));
TESTING_FREE_CPU( hwork );
}
#ifdef HAVE_CUBLAS
else if ( opts.version == 3 ) {
// allocate hwork
magma_cgeqrs3_gpu( M, N, 1,
d_A, 0, ldda, tau, dT, 0,
d_B, 0, M, tmp, -1, opts.queue, &info );
lwork = (magma_int_t)MAGMA_C_REAL( tmp[0] );
TESTING_MALLOC_CPU( hwork, magmaFloatComplex, lwork );
// solve linear system
magma_cgeqrs3_gpu( M, N, 1,
d_A, 0, ldda, tau, dT, 0,
d_B, 0, M, hwork, lwork, opts.queue, &info );
if (info != 0)
printf("magma_cgeqrs3 returned error %d: %s.\n",
(int) info, magma_strerror( info ));
TESTING_FREE_CPU( hwork );
}
#endif
else {
printf( "Unknown version %d\n", opts.version );
exit(1);
}
magma_cgetvector( N, d_B, 0, 1, x, 1, opts.queue );
// compute r = Ax - b, saved in b
lapackf77_clarnv( &ione, ISEED2, &n2, h_A );
blasf77_cgemv( "Notrans", &M, &N, &c_one, h_A, &lda, x, &ione, &c_neg_one, b, &ione );
// compute residual |Ax - b| / (n*|A|*|x|)
float norm_x, norm_A, norm_r, work[1];
norm_A = lapackf77_clange( "F", &M, &N, h_A, &lda, work );
norm_r = lapackf77_clange( "F", &M, &ione, b, &M, work );
norm_x = lapackf77_clange( "F", &N, &ione, x, &N, work );
TESTING_FREE_CPU( x );
TESTING_FREE_CPU( b );
TESTING_FREE_DEV( d_B );
error = norm_r / (N * norm_A * norm_x);
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time, error );
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) %8.2e",
(int) M, (int) N, gpu_perf, gpu_time, error );
}
printf(" %s\n", (error < tol ? "ok" : "failed"));
status += ! (error < tol);
}
else {
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) ---",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time );
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) ---",
(int) M, (int) N, gpu_perf, gpu_time);
}
printf("%s\n", (opts.check != 0 ? " (error check only for M >= N)" : ""));
}
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_work );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
if ( opts.version != 2 )
TESTING_FREE_DEV( dT );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
int main( int argc, char** argv)
{
real_Double_t gflops, gpu_perf, cpu_perf, gpu_time, cpu_time, error;
float matnorm, work[1];
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *h_A, *h_R, *tau, *h_work, tmp[1];
magmaFloatComplex_ptr d_lA[MagmaMaxGPUs];
/* Matrix size */
magma_int_t M = 0, N = 0, n2, n_local[4], lda, ldda, lhwork;
magma_int_t size[10] = {1000,2000,3000,4000,5000,6000,7000,8000,9000,10000};
magma_int_t i, k, nk, info, min_mn;
int max_num_gpus = 2, num_gpus = 2;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
if (argc != 1){
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0)
N = atoi(argv[++i]);
else if (strcmp("-M", argv[i])==0)
M = atoi(argv[++i]);
else if (strcmp("-NGPU", argv[i])==0)
num_gpus = atoi(argv[++i]);
}
if ( M == 0 ) {
M = N;
}
if ( N == 0 ) {
N = M;
}
if (M>0 && N>0)
printf(" testing_cgeqrf_gpu -M %d -N %d -NGPU %d\n\n", (int) M, (int) N, (int) num_gpus);
else
{
printf("\nUsage: \n");
printf(" testing_cgeqrf_gpu -M %d -N %d -NGPU %d\n\n",
1024, 1024, 1);
exit(1);
}
}
else {
printf("\nUsage: \n");
printf(" testing_cgeqrf_gpu -M %d -N %d -NGPU %d\n\n", 1024, 1024, 1);
M = N = size[9];
}
ldda = ((M+31)/32)*32;
n2 = M * N;
min_mn = min(M, N);
magma_int_t nb = magma_get_cgeqrf_nb(M);
if (num_gpus > max_num_gpus){
printf("More GPUs requested than available. Have to change it.\n");
num_gpus = max_num_gpus;
}
printf("Number of GPUs to be used = %d\n", (int) num_gpus);
/* Initialize */
magma_queue_t queues[MagmaMaxGPUs * 2];
magma_device_t devices[ MagmaMaxGPUs ];
magma_int_t num = 0;
magma_int_t err;
magma_init();
err = magma_getdevices( devices, MagmaMaxGPUs, &num );
if ( err != 0 || num < 1 ) {
fprintf( stderr, "magma_getdevices failed: %d\n", (int) err );
exit(-1);
}
for(i=0;i<num_gpus;i++){
err = magma_queue_create( devices[i], &queues[2*i] );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", (int) err );
exit(-1);
}
err = magma_queue_create( devices[i], &queues[2*i+1] );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", (int) err );
exit(-1);
}
}
/* Allocate host memory for the matrix */
TESTING_MALLOC_CPU( tau, magmaFloatComplex, min_mn );
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( h_R, magmaFloatComplex, n2 );
for(i=0; i<num_gpus; i++){
n_local[i] = ((N/nb)/num_gpus)*nb;
if (i < (N/nb)%num_gpus)
n_local[i] += nb;
else if (i == (N/nb)%num_gpus)
n_local[i] += N%nb;
TESTING_MALLOC_DEV( d_lA[i], magmaFloatComplex, ldda*n_local[i] );
printf("device %2d n_local = %4d\n", (int) i, (int) n_local[i]);
}
lhwork = -1;
lapackf77_cgeqrf(&M, &N, h_A, &M, tau, tmp, &lhwork, &info);
lhwork = (magma_int_t)MAGMA_C_REAL( tmp[0] );
TESTING_MALLOC_CPU( h_work, magmaFloatComplex, lhwork );
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / ||A||_F\n");
printf("======================================================================\n");
for(i=0; i<10; i++){
if (argc == 1){
M = N = size[i];
}
min_mn= min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = FLOPS( (float)M, (float)N ) * 1e-9;
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_cgeqrf(&M, &N, h_A, &M, tau, h_work, &lhwork, &info);
cpu_time = magma_wtime() - cpu_time;
if (info < 0)
printf("Argument %d of lapack_cgeqrf had an illegal value.\n", (int) -info);
cpu_perf = gflops / cpu_time;
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
int j;
magma_queue_t *trans_queues = (magma_queue_t*)malloc(num_gpus*sizeof(magma_queue_t));
for(j=0;j<num_gpus;j++){
trans_queues[j] = queues[2*j];
}
// warm-up
magma_csetmatrix_1D_col_bcyclic(M, N, h_R, lda, d_lA, ldda, num_gpus, nb, trans_queues);
magma_cgeqrf2_mgpu( num_gpus, M, N, d_lA, ldda, tau, queues, &info);
magma_csetmatrix_1D_col_bcyclic(M, N, h_R, lda, d_lA, ldda, num_gpus, nb, trans_queues);
gpu_time = magma_wtime();
magma_cgeqrf2_mgpu( num_gpus, M, N, d_lA, ldda, tau, queues, &info);
gpu_time = magma_wtime() - gpu_time;
if (info < 0)
printf("Argument %d of magma_cgeqrf2 had an illegal value.\n", (int) -info);
gpu_perf = gflops / gpu_time;
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_cgetmatrix_1D_col_bcyclic(M, N, d_lA, ldda, h_R, lda, num_gpus, nb, trans_queues);
matnorm = lapackf77_clange("f", &M, &N, h_A, &M, work);
blasf77_caxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
printf("%5d %5d %6.2f (%6.2f) %6.2f (%6.2f) %e\n",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time,
lapackf77_clange("f", &M, &N, h_R, &M, work) / matnorm);
if (argc != 1)
break;
}
/* Memory clean up */
TESTING_FREE_PIN( tau );
TESTING_FREE_PIN( h_A );
TESTING_FREE_PIN( h_work );
TESTING_FREE_PIN( h_R );
for(i=0; i<num_gpus; i++){
TESTING_FREE_DEV( d_lA[i] );
magma_queue_destroy(queues[2*i]);
magma_queue_destroy(queues[2*i+1]);
}
/* Shutdown */
magma_finalize();
}
extern "C" magma_int_t
magma_cgeqrf_gpu( magma_int_t m, magma_int_t n,
magmaFloatComplex *dA, magma_int_t ldda,
magmaFloatComplex *tau, magmaFloatComplex *dT,
magma_int_t *info )
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CGEQRF computes a QR factorization of a complex M-by-N matrix A:
A = Q * R.
This version stores the triangular dT matrices used in
the block QR factorization so that they can be applied directly (i.e.,
without being recomputed) later. As a result, the application
of Q is much faster. Also, the upper triangular matrices for V have 0s
in them. The corresponding parts of the upper triangular R are inverted
and stored separately in dT.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. N >= 0.
dA (input/output) COMPLEX array on the GPU, dimension (LDDA,N)
On entry, the M-by-N matrix A.
On exit, the elements on and above the diagonal of the array
contain the min(M,N)-by-N upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of min(m,n) elementary reflectors (see Further
Details).
LDDA (input) INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
To benefit from coalescent memory accesses LDDA must be
dividable by 16.
TAU (output) COMPLEX array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
dT (workspace/output) COMPLEX array on the GPU,
dimension (2*MIN(M, N) + (N+31)/32*32 )*NB,
where NB can be obtained through magma_get_cgeqrf_nb(M).
It starts with MIN(M,N)*NB block that store the triangular T
matrices, followed by the MIN(M,N)*NB block of the diagonal
inverses for the R matrix. The rest of the array is used as workspace.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
Further Details
===============
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
===================================================================== */
#define a_ref(a_1,a_2) (dA+(a_2)*(ldda) + (a_1))
#define t_ref(a_1) (dT+(a_1)*nb)
#define d_ref(a_1) (dT+(minmn+(a_1))*nb)
#define dd_ref(a_1) (dT+(2*minmn+(a_1))*nb)
#define work_ref(a_1) ( work + (a_1))
#define hwork ( work + (nb)*(m))
magma_int_t i, k, minmn, old_i, old_ib, rows, cols;
magma_int_t ib, nb;
magma_int_t ldwork, lddwork, lwork, lhwork;
magmaFloatComplex *work, *ut;
/* check arguments */
*info = 0;
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (ldda < max(1,m)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
k = minmn = min(m,n);
if (k == 0)
return *info;
nb = magma_get_cgeqrf_nb(m);
lwork = (m + n + nb)*nb;
lhwork = lwork - m*nb;
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &work, lwork )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
ut = hwork+nb*(n);
memset( ut, 0, nb*nb*sizeof(magmaFloatComplex));
magma_queue_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
ldwork = m;
lddwork= n;
if ( (nb > 1) && (nb < k) ) {
/* Use blocked code initially */
old_i = 0; old_ib = nb;
for (i = 0; i < k-nb; i += nb) {
ib = min(k-i, nb);
rows = m -i;
magma_cgetmatrix_async( rows, ib,
a_ref(i,i), ldda,
work_ref(i), ldwork, stream[1] );
if (i>0){
/* Apply H' to A(i:m,i+2*ib:n) from the left */
cols = n-old_i-2*old_ib;
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
m-old_i, cols, old_ib,
a_ref(old_i, old_i ), ldda, t_ref(old_i), nb,
a_ref(old_i, old_i+2*old_ib), ldda, dd_ref(0), lddwork);
/* store the diagonal */
magma_csetmatrix_async( old_ib, old_ib,
ut, old_ib,
d_ref(old_i), old_ib, stream[0] );
}
magma_queue_sync( stream[1] );
lapackf77_cgeqrf(&rows, &ib, work_ref(i), &ldwork, tau+i, hwork, &lhwork, info);
/* Form the triangular factor of the block reflector
H = H(i) H(i+1) . . . H(i+ib-1) */
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&rows, &ib,
work_ref(i), &ldwork, tau+i, hwork, &ib);
/* Put 0s in the upper triangular part of a panel (and 1s on the
diagonal); copy the upper triangular in ut and invert it. */
magma_queue_sync( stream[0] );
csplit_diag_block(ib, work_ref(i), ldwork, ut);
magma_csetmatrix( rows, ib, work_ref(i), ldwork, a_ref(i,i), ldda );
if (i + ib < n) {
/* Send the triangular factor T to the GPU */
magma_csetmatrix( ib, ib, hwork, ib, t_ref(i), nb );
if (i+nb < k-nb){
/* Apply H' to A(i:m,i+ib:i+2*ib) from the left */
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, ib, ib,
a_ref(i, i ), ldda, t_ref(i), nb,
a_ref(i, i+ib), ldda, dd_ref(0), lddwork);
}
else {
cols = n-i-ib;
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, cols, ib,
a_ref(i, i ), ldda, t_ref(i), nb,
a_ref(i, i+ib), ldda, dd_ref(0), lddwork);
/* Fix the diagonal block */
magma_csetmatrix( ib, ib, ut, ib, d_ref(i), ib );
}
old_i = i;
old_ib = ib;
}
}
} else {
i = 0;
}
/* Use unblocked code to factor the last or only block. */
if (i < k) {
ib = n-i;
rows = m-i;
magma_cgetmatrix( rows, ib, a_ref(i, i), ldda, work, rows );
lhwork = lwork - rows*ib;
lapackf77_cgeqrf(&rows, &ib, work, &rows, tau+i, work+ib*rows, &lhwork, info);
magma_csetmatrix( rows, ib, work, rows, a_ref(i, i), ldda );
}
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free_pinned( work );
return *info;
/* End of MAGMA_CGEQRF */
} /* magma_cgeqrf */
/**
Purpose
-------
CGEGQR orthogonalizes the N vectors given by a complex M-by-N matrix A:
A = Q * R.
On exit, if successful, the orthogonal vectors Q overwrite A
and R is given in work (on the CPU memory).
The routine is designed for tall-and-skinny matrices: M >> N, N <= 128.
This version uses normal equations and SVD in an iterative process that
makes the computation numerically accurate.
Arguments
---------
@param[in]
ikind INTEGER
Several versions are implemented indiceted by the ikind value:
1: This version uses normal equations and SVD in an iterative process
that makes the computation numerically accurate.
2: This version uses a standard LAPACK-based orthogonalization through
MAGMA's QR panel factorization (magma_cgeqr2x3_gpu) and magma_cungqr
3: MGS
4. Cholesky QR
@param[in]
m INTEGER
The number of rows of the matrix A. m >= n >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. 128 >= n >= 0.
@param[in,out]
dA COMPLEX array on the GPU, dimension (ldda,n)
On entry, the m-by-n matrix A.
On exit, the m-by-n matrix Q with orthogonal columns.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,m).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
@param
dwork (GPU workspace) COMPLEX array, dimension:
n^2 for ikind = 1
3 n^2 + min(m, n) for ikind = 2
0 (not used) for ikind = 3
n^2 for ikind = 4
@param[out]
work (CPU workspace) COMPLEX array, dimension 3 n^2.
On exit, work(1:n^2) holds the rectangular matrix R.
Preferably, for higher performance, work should be in pinned memory.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
@ingroup magma_cgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_cgegqr_gpu( magma_int_t ikind, magma_int_t m, magma_int_t n,
magmaFloatComplex *dA, magma_int_t ldda,
magmaFloatComplex *dwork, magmaFloatComplex *work,
magma_int_t *info )
{
#define work(i_,j_) (work + (i_) + (j_)*n)
#define dA(i_,j_) (dA + (i_) + (j_)*ldda)
magma_int_t i = 0, j, k, n2 = n*n;
magma_int_t ione = 1;
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
float cn = 200., mins, maxs;
/* check arguments */
*info = 0;
if (ikind < 1 || ikind > 4) {
*info = -1;
} else if (m < 0 || m < n) {
*info = -2;
} else if (n < 0 || n > 128) {
*info = -3;
} else if (ldda < max(1,m)) {
*info = -5;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (ikind == 1) {
// === Iterative, based on SVD ============================================================
magmaFloatComplex *U, *VT, *vt, *R, *G, *hwork, *tau;
float *S;
R = work; // Size n * n
G = R + n*n; // Size n * n
VT = G + n*n; // Size n * n
magma_cmalloc_cpu( &hwork, 32 + 2*n*n + 2*n);
if ( hwork == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_int_t lwork=n*n+32; // First part f hwork; used as workspace in svd
U = hwork + n*n + 32; // Size n*n
S = (float *)(U+n*n); // Size n
tau = U + n*n + n; // Size n
#if defined(PRECISION_c) || defined(PRECISION_z)
float *rwork;
magma_smalloc_cpu( &rwork, 5*n);
if ( rwork == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
#endif
do {
i++;
magma_cgemm(MagmaConjTrans, MagmaNoTrans, n, n, m, c_one, dA, ldda, dA, ldda, c_zero, dwork, n );
magma_cgetmatrix(n, n, dwork, n, G, n);
#if defined(PRECISION_s) || defined(PRECISION_d)
lapackf77_cgesvd("n", "a", &n, &n, G, &n, S, U, &n, VT, &n,
hwork, &lwork, info);
#else
lapackf77_cgesvd("n", "a", &n, &n, G, &n, S, U, &n, VT, &n,
hwork, &lwork, rwork, info);
#endif
mins = 100.f, maxs = 0.f;
for (k=0; k < n; k++) {
S[k] = magma_ssqrt( S[k] );
if (S[k] < mins) mins = S[k];
if (S[k] > maxs) maxs = S[k];
}
for (k=0; k < n; k++) {
vt = VT + k*n;
for (j=0; j < n; j++)
vt[j] *= S[j];
}
lapackf77_cgeqrf(&n, &n, VT, &n, tau, hwork, &lwork, info);
if (i == 1)
blasf77_ccopy(&n2, VT, &ione, R, &ione);
else
blasf77_ctrmm("l", "u", "n", "n", &n, &n, &c_one, VT, &n, R, &n);
magma_csetmatrix(n, n, VT, n, dwork, n);
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit, m, n, c_one, dwork, n, dA, ldda);
if (mins > 0.00001f)
cn = maxs/mins;
//fprintf(stderr, "Iteration %d, cond num = %f \n", i, cn);
} while (cn > 10.f);
magma_free_cpu( hwork );
#if defined(PRECISION_c) || defined(PRECISION_z)
magma_free_cpu( rwork );
#endif
// ================== end of ikind == 1 ===================================================
}
else if (ikind == 2) {
// ================== LAPACK based ===================================================
magma_int_t min_mn = min(m, n);
magma_int_t nb = n;
magmaFloatComplex *dtau = dwork + 2*n*n, *d_T = dwork, *ddA = dwork + n*n;
magmaFloatComplex *tau = work+n*n;
magmablas_claset( MagmaFull, n, n, c_zero, c_zero, d_T, n );
magma_cgeqr2x3_gpu(&m, &n, dA, &ldda, dtau, d_T, ddA,
(float *)(dwork+min_mn+2*n*n), info);
magma_cgetmatrix( min_mn, 1, dtau, min_mn, tau, min_mn);
magma_cgetmatrix( n, n, ddA, n, work, n);
magma_cungqr_gpu( m, n, n, dA, ldda, tau, d_T, nb, info );
// ================== end of ikind == 2 ===================================================
}
else if (ikind == 3) {
// ================== MGS ===================================================
for(magma_int_t j = 0; j<n; j++){
for(magma_int_t i = 0; i<j; i++){
*work(i, j) = magma_cdotc(m, dA(0,i), 1, dA(0,j), 1);
magma_caxpy(m, -(*work(i,j)), dA(0,i), 1, dA(0,j), 1);
}
for(magma_int_t i = j; i<n; i++)
*work(i, j) = MAGMA_C_ZERO;
//*work(j,j) = MAGMA_C_MAKE( magma_scnrm2(m, dA(0,j), 1), 0. );
*work(j,j) = magma_cdotc(m, dA(0,j), 1, dA(0,j), 1);
*work(j,j) = MAGMA_C_MAKE( sqrt(MAGMA_C_REAL( *work(j,j) )), 0.);
magma_cscal(m, 1./ *work(j,j), dA(0,j), 1);
}
// ================== end of ikind == 3 ===================================================
}
else if (ikind == 4) {
// ================== Cholesky QR ===================================================
magma_cgemm(MagmaConjTrans, MagmaNoTrans, n, n, m, c_one, dA, ldda, dA, ldda, c_zero, dwork, n );
magma_cgetmatrix(n, n, dwork, n, work, n);
lapackf77_cpotrf("u", &n, work, &n, info);
magma_csetmatrix(n, n, work, n, dwork, n);
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit, m, n, c_one, dwork, n, dA, ldda);
// ================== end of ikind == 4 ===================================================
}
return *info;
} /* magma_cgegqr_gpu */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgeqrf
*/
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;
magmaFloatComplex *h_A, *h_R, *tau, *dtau, *h_work, tmp[1];
magmaFloatComplex *d_A;
float *dwork;
magma_int_t M, N, n2, lda, ldda, lwork, info, min_mn;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
printf(" M N CPU GFlop/s (ms) GPU GFlop/s (ms) ||R||_F / ||A||_F\n");
printf("=======================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[i];
N = opts.nsize[i];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = FLOPS_CGEQRF( M, N ) / 1e9;
lwork = -1;
lapackf77_cgeqrf(&M, &N, h_A, &M, tau, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_C_REAL( tmp[0] );
TESTING_MALLOC( tau, magmaFloatComplex, min_mn );
TESTING_MALLOC( h_A, magmaFloatComplex, n2 );
TESTING_HOSTALLOC( h_R, magmaFloatComplex, n2 );
TESTING_DEVALLOC( d_A, magmaFloatComplex, ldda*N );
TESTING_DEVALLOC( dtau, magmaFloatComplex, min_mn );
TESTING_DEVALLOC(dwork, float, min_mn );
TESTING_MALLOC( h_work, magmaFloatComplex, lwork );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
magma_csetmatrix( M, N, h_R, lda, d_A, ldda );
// warmup
magma_cgeqr2_gpu( M, N, d_A, ldda, dtau, dwork, &info );
magma_csetmatrix( M, N, h_R, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_sync_wtime( 0 );
magma_cgeqr2_gpu( M, N, d_A, ldda, dtau, dwork, &info );
gpu_time = magma_sync_wtime( 0 ) - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_cgeqrf(&M, &N, h_A, &lda, tau, h_work, &lwork, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_cgetmatrix( M, N, d_A, ldda, h_R, M );
error = lapackf77_clange("f", &M, &N, h_A, &lda, work);
blasf77_caxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_clange("f", &M, &N, h_R, &lda, work) / error;
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e\n",
(int) M, (int) N, cpu_perf, 1000.*cpu_time, gpu_perf, 1000.*gpu_time, error );
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) M, (int) N, gpu_perf, 1000.*gpu_time );
}
TESTING_FREE( tau );
TESTING_FREE( h_A );
TESTING_FREE( h_work );
TESTING_HOSTFREE( h_R );
TESTING_DEVFREE( d_A );
TESTING_DEVFREE( dtau );
TESTING_DEVFREE( dwork );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/**
Purpose
-------
CHETRD_HE2HB reduces a complex Hermitian matrix A to real symmetric
band-diagonal form T by an orthogonal similarity transformation:
Q**H * A * Q = T.
This version stores the triangular matrices T used in the accumulated
Householder transformations (I - V T V').
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored;
- = MagmaLower: Lower triangle of A is stored.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
A COMPLEX array, dimension (LDA,N)
On entry, the Hermitian matrix A. If UPLO = MagmaUpper, the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit, if UPLO = MagmaUpper, the Upper band-diagonal of A is
overwritten by the corresponding elements of the
band-diagonal matrix T, and the elements above the band
diagonal, with the array TAU, represent the orthogonal
matrix Q as a product of elementary reflectors; if UPLO
= MagmaLower, the the Lower band-diagonal of A is overwritten by
the corresponding elements of the band-diagonal
matrix T, and the elements below the band-diagonal, with
the array TAU, represent the orthogonal matrix Q as a product
of elementary reflectors. See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
tau COMPLEX array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details).
@param[out]
work (workspace) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK. LWORK >= 1.
For optimum performance LWORK >= N*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]
dT COMPLEX array on the GPU, dimension N*NB,
where NB is the optimal blocksize.
On exit dT holds the upper triangular matrices T from the
accumulated Householder transformations (I - V T V') used
in the factorization. The nb x nb matrices T are ordered
consecutively in memory one after another.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
Further Details
---------------
If UPLO = MagmaUpper, the matrix Q is represented as a product of elementary
reflectors
Q = H(n-1) . . . H(2) H(1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(i+1:n) = 0 and v(i) = 1; v(1:i-1) is stored on exit in
A(1:i-1,i+1), and tau in TAU(i).
If UPLO = MagmaLower, the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(n-1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i) = 0 and v(i+1) = 1; v(i+2:n) is stored on exit in A(i+2:n,i),
and tau in TAU(i).
The contents of A on exit are illustrated by the following examples
with n = 5:
if UPLO = MagmaUpper: if UPLO = MagmaLower:
( d e v2 v3 v4 ) ( d )
( d e v3 v4 ) ( e d )
( d e v4 ) ( v1 e d )
( d e ) ( v1 v2 e d )
( d ) ( v1 v2 v3 e d )
where d and e denote diagonal and off-diagonal elements of T, and vi
denotes an element of the vector defining H(i).
@ingroup magma_cheev_2stage
********************************************************************/
extern "C" magma_int_t
magma_chetrd_he2hb( magma_uplo_t uplo, magma_int_t n, magma_int_t nb,
magmaFloatComplex *A, magma_int_t lda,
magmaFloatComplex *tau,
magmaFloatComplex *work, magma_int_t lwork,
magmaFloatComplex *dT,
magma_int_t *info)
{
#define A(a_1,a_2) ( A + ((a_2)-1)*( lda) + (a_1)-1)
#define dA(a_1,a_2) (dA + ((a_2)-1)*(ldda) + (a_1)-1)
#define tau_ref(a_1) (tau + (a_1)-1)
#define dT(a_1) (dT + ((a_1)-1)*(lddt))
int ldda = ((n+31)/32)*32;
int lddt = nb;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex c_neg_half = MAGMA_C_NEG_HALF;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_zero = MAGMA_C_ZERO;
float d_one = MAGMA_D_ONE;
magma_int_t pm, pn, indi, indj, pk;
magma_int_t pm_old=0, pn_old=0, indi_old=0, indj_old=0;
int i;
int lwkopt;
int lquery;
*info = 0;
int upper = (uplo == MagmaUpper);
lquery = (lwork == -1);
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
} else if (lwork < 1 && ! lquery) {
*info = -9;
}
/* Determine the block size. */
lwkopt = n * nb;
if (*info == 0) {
work[0] = MAGMA_C_MAKE( lwkopt, 0 );
}
if (*info != 0)
return *info;
else if (lquery)
return *info;
/* Quick return if possible */
if (n == 0) {
work[0] = c_one;
return *info;
}
magmaFloatComplex *dA;
if (MAGMA_SUCCESS != magma_cmalloc( &dA, (n + 2*nb)*ldda )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_int_t threads = magma_get_lapack_numthreads();
magma_int_t mklth = min(threads,16);
magma_set_lapack_numthreads(mklth);
/* Use the first panel of dA as work space */
magmaFloatComplex *dwork = dA + n*ldda;
magmaFloatComplex *dW = dwork + nb*ldda;
#ifdef TRACING
char buf[80];
#endif
magma_queue_t stream[3];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
stream[2] = 0; // default stream
trace_init( 1, 1, 3, stream );
magmaFloatComplex *hT = work + lwork - nb*nb;
lwork -= nb*nb;
memset( hT, 0, nb*nb*sizeof(magmaFloatComplex));
magmablasSetKernelStream( stream[0] );
magma_event_t Pupdate_event;
cudaEventCreateWithFlags(&Pupdate_event,cudaEventDisableTiming);
//cudaEventCreate(&Pupdate_event);
if (upper) {
printf("CHETRD_HE2HB is not yet implemented for upper matrix storage. Exit.\n");
exit(1);
} else {
/* Copy the matrix to the GPU */
if (1 <= n-nb) {
trace_gpu_start( 0, 0, "set", "set A" );
magma_csetmatrix_async( (n-nb), (n-nb),
A(nb+1, nb+1), lda,
dA(nb+1, nb+1), ldda, stream[0] );
trace_gpu_end( 0, 0 );
}
/* Reduce the lower triangle of A */
for (i = 1; i <= n-nb; i += nb) {
indi = i+nb;
indj = i;
pm = n - i - nb + 1;
//pn = min(i+nb-1, n-nb) -i + 1;
pn = nb;
/* Get the current panel (no need for the 1st iteration) */
if (i > 1 ) {
// cpanel_to_q copy the upper oof diagonal part of
// the matrix to work to be restored later. acctually
// the zero's and one's putted are not used this is only
// because we don't have a function that copy only the
// upper part of A to be restored after copying the
// lookahead panel that has been computted from GPU to CPU.
cpanel_to_q(MagmaUpper, pn-1, A(i, i+1), lda, work);
trace_gpu_start( 0, 1, "get", "get panel" );
//magma_queue_sync( stream[0] );
cudaStreamWaitEvent(stream[1], Pupdate_event, 0);
magma_cgetmatrix_async( (pm+pn), pn,
dA( i, i), ldda,
A ( i, i), lda, stream[1] );
trace_gpu_end( 0, 1 );
trace_gpu_start( 0, 2, "her2k", "her2k" );
magma_cher2k(MagmaLower, MagmaNoTrans, pm_old-pn_old, pn_old, c_neg_one,
dA(indi_old+pn_old, indj_old), ldda,
dW + pn_old, pm_old, d_one,
dA(indi_old+pn_old, indi_old+pn_old), ldda);
trace_gpu_end( 0, 2 );
trace_cpu_start( 0, "sync", "sync on 1" );
magma_queue_sync( stream[1] );
trace_cpu_end( 0 );
cq_to_panel(MagmaUpper, pn-1, A(i, i+1), lda, work);
}
/* ==========================================================
QR factorization on a panel starting nb off of the diagonal.
Prepare the V and T matrices.
========================================================== */
#ifdef TRACING
snprintf( buf, sizeof(buf), "panel %d", i );
#endif
trace_cpu_start( 0, "geqrf", buf );
lapackf77_cgeqrf(&pm, &pn, A(indi, indj), &lda,
tau_ref(i), work, &lwork, info);
/* Form the matrix T */
pk=min(pm,pn);
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&pm, &pk, A(indi, indj), &lda,
tau_ref(i), hT, &nb);
/* Prepare V - put 0s in the upper triangular part of the panel
(and 1s on the diagonal), temporaly storing the original in work */
cpanel_to_q(MagmaUpper, pk, A(indi, indj), lda, work);
trace_cpu_end( 0 );
/* Send V from the CPU to the GPU */
trace_gpu_start( 0, 0, "set", "set V and T" );
magma_csetmatrix_async( pm, pk,
A(indi, indj), lda,
dA(indi, indj), ldda, stream[0] );
/* Send the triangular factor T to the GPU */
magma_csetmatrix_async( pk, pk,
hT, nb,
dT(i), lddt, stream[0] );
trace_gpu_end( 0, 0 );
/* ==========================================================
Compute W:
1. X = A (V T)
2. W = X - 0.5* V * (T' * (V' * X))
========================================================== */
/* dwork = V T */
trace_cpu_start( 0, "sync", "sync on 0" );
// this sync is done here to be sure that the copy has been finished
// because below we made a restore cq_to_panel and this restore need
// to ensure that the copy has been finished. we did it here to allow
// overlapp of restore with next gemm and symm.
magma_queue_sync( stream[0] );
trace_cpu_end( 0 );
trace_gpu_start( 0, 2, "gemm", "work = V*T" );
magma_cgemm(MagmaNoTrans, MagmaNoTrans, pm, pk, pk,
c_one, dA(indi, indj), ldda,
dT(i), lddt,
c_zero, dwork, pm);
trace_gpu_end( 0, 2 );
/* dW = X = A*V*T. dW = A*dwork */
trace_gpu_start( 0, 2, "hemm", "X = A*work" );
magma_chemm(MagmaLeft, uplo, pm, pk,
c_one, dA(indi, indi), ldda,
dwork, pm,
c_zero, dW, pm);
trace_gpu_end( 0, 2 );
/* restore the panel */
cq_to_panel(MagmaUpper, pk, A(indi, indj), lda, work);
/* dwork = V*T already ==> dwork' = T'*V'
* compute T'*V'*X ==> dwork'*W ==>
* dwork + pm*nb = ((T' * V') * X) = dwork' * X = dwork' * W */
trace_gpu_start( 0, 2, "gemm", "work = T'*V'*X" );
magma_cgemm(MagmaConjTrans, MagmaNoTrans, pk, pk, pm,
c_one, dwork, pm,
dW, pm,
c_zero, dwork + pm*nb, nb);
trace_gpu_end( 0, 2 );
/* W = X - 0.5 * V * T'*V'*X
* = X - 0.5 * V * (dwork + pm*nb) = W - 0.5 * V * (dwork + pm*nb) */
trace_gpu_start( 0, 2, "gemm", "W = X - 0.5*V*(T'*V'*X)" );
magma_cgemm(MagmaNoTrans, MagmaNoTrans, pm, pk, pk,
c_neg_half, dA(indi, indj), ldda,
dwork + pm*nb, nb,
c_one, dW, pm);
trace_gpu_end( 0, 2 );
/* ==========================================================
Update the unreduced submatrix A(i+ib:n,i+ib:n), using
an update of the form: A := A - V*W' - W*V'
========================================================== */
if (i + nb <= n-nb) {
/* There would be next iteration;
do lookahead - update the next panel */
trace_gpu_start( 0, 2, "gemm", "gemm 4 next panel left" );
magma_cgemm(MagmaNoTrans, MagmaConjTrans, pm, pn, pn, c_neg_one,
dA(indi, indj), ldda,
dW, pm, c_one,
dA(indi, indi), ldda);
trace_gpu_end( 0, 2 );
trace_gpu_start( 0, 2, "gemm", "gemm 5 next panel right" );
magma_cgemm(MagmaNoTrans, MagmaConjTrans, pm, pn, pn, c_neg_one,
dW, pm,
dA(indi, indj), ldda, c_one,
dA(indi, indi), ldda);
trace_gpu_end( 0, 2 );
cudaEventRecord(Pupdate_event, stream[0]);
}
else {
/* no look-ahead as this is last iteration */
trace_gpu_start( 0, 2, "her2k", "her2k last iteration" );
magma_cher2k(MagmaLower, MagmaNoTrans, pk, pk, c_neg_one,
dA(indi, indj), ldda,
dW, pm, d_one,
dA(indi, indi), ldda);
trace_gpu_end( 0, 2 );
}
indi_old = indi;
indj_old = indj;
pm_old = pm;
pn_old = pn;
} // end loop for (i)
/* Send the last block to the CPU */
pk = min(pm,pn);
if (1 <= n-nb) {
cpanel_to_q(MagmaUpper, pk-1, A(n-pk+1, n-pk+2), lda, work);
trace_gpu_start( 0, 2, "get", "get last block" );
magma_cgetmatrix( pk, pk,
dA(n-pk+1, n-pk+1), ldda,
A(n-pk+1, n-pk+1), lda );
trace_gpu_end( 0, 2 );
cq_to_panel(MagmaUpper, pk-1, A(n-pk+1, n-pk+2), lda, work);
}
}// end of LOWER
trace_finalize( "chetrd_he2hb.svg", "trace.css" );
cudaEventDestroy(Pupdate_event);
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free( dA );
work[0] = MAGMA_C_MAKE( lwkopt, 0 );
magmablasSetKernelStream( 0 );
magma_set_lapack_numthreads(threads);
return *info;
} /* magma_chetrd_he2hb */
