extern "C" magma_int_t
magma_cunmqr2_gpu(const char side, const char trans,
magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloatComplex *da, magma_int_t ldda,
magmaFloatComplex *tau,
magmaFloatComplex *dc, magma_int_t lddc,
magmaFloatComplex *wa, magma_int_t ldwa,
magma_int_t *info)
{
/* -- MAGMA (version 1.4.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
December 2013
Purpose
=======
CUNMQR overwrites the general complex M-by-N matrix C with
SIDE = 'L' SIDE = 'R'
TRANS = 'N': Q * C C * Q
TRANS = 'T': Q**H * C C * Q**H
where Q is a complex orthogonal matrix defined as the product of k
elementary reflectors
Q = H(1) H(2) . . . H(k)
as returned by CGEQRF. Q is of order M if SIDE = 'L' and of order N
if SIDE = 'R'.
Arguments
=========
SIDE (input) CHARACTER*1
= 'L': apply Q or Q**H from the Left;
= 'R': apply Q or Q**H from the Right.
TRANS (input) CHARACTER*1
= 'N': No transpose, apply Q;
= 'T': Transpose, apply Q**H.
M (input) INTEGER
The number of rows of the matrix C. M >= 0.
N (input) INTEGER
The number of columns of the matrix C. N >= 0.
K (input) INTEGER
The number of elementary reflectors whose product defines
the matrix Q.
If SIDE = 'L', M >= K >= 0;
if SIDE = 'R', N >= K >= 0.
DA (input) COMPLEX array, dimension (LDA,K)
The i-th column must contain the vector which defines the
elementary reflector H(i), for i = 1,2,...,k, as returned by
CGEQRF in the first k columns of its array argument A.
The diagonal and the upper part
are destroyed, the reflectors are not modified.
LDDA (input) INTEGER
The leading dimension of the array DA.
LDDA >= max(1,M) if SIDE = 'L'; LDDA >= max(1,N) if SIDE = 'R'.
TAU (input) COMPLEX array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by CGEQRF.
DC (device input/output) COMPLEX array, dimension (LDDC,N)
On entry, the M-by-N matrix C.
On exit, C is overwritten by (Q*C) or (Q**H * C) or (C * Q**H) or (C*Q).
LDDC (input) INTEGER
The leading dimension of the array C. LDDC >= max(1,M).
WA (input/workspace) COMPLEX array, dimension
(LDWA,M) if SIDE = 'L'
(LDWA,N) if SIDE = 'R'
The vectors which define the elementary reflectors, as
returned by CHETRD_GPU.
LDWA (input) INTEGER
The leading dimension of the array A.
LDWA >= max(1,M) if SIDE = 'L'; LDWA >= max(1,N) if SIDE = 'R'.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
===================================================================== */
char side_[2] = {side, 0};
char trans_[2] = {trans, 0};
/* Allocate work space on the GPU */
magmaFloatComplex *dwork;
magma_int_t wa_offset, dc_offset, i__4, lddwork;
magma_int_t i;
magmaFloatComplex t[2*4160] /* was [65][64] */;
magma_int_t i1, i2, step, ib, ic, jc, nb, mi, ni, nq, nw;
int left, notran;
wa_offset = 1 + ldwa;
wa -= wa_offset;
--tau;
dc_offset = 1 + lddc;
dc -= dc_offset;
*info = 0;
left = lapackf77_lsame(side_, "L");
notran = lapackf77_lsame(trans_, "N");
/* NQ is the order of Q and NW is the minimum dimension of WORK */
if (left) {
nq = m;
nw = n;
magma_cmalloc( &dwork, (n + 64)*64 );
} else {
nq = n;
nw = m;
magma_cmalloc( &dwork, (m + 64)*64 );
}
if (! left && ! lapackf77_lsame(side_, "R")) {
*info = -1;
} else if (! notran && ! lapackf77_lsame(trans_, "T")) {
*info = -2;
} else if (m < 0) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (k < 0 || k > nq) {
*info = -5;
} else if (ldda < max(1,nq)) {
*info = -7;
} else if (lddc < max(1,m)) {
*info = -10;
} else if (ldwa < max(1,nq)) {
*info = -12;
}
// size of the block
nb = 64;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0 || k == 0) {
return *info;
}
/* Use hybrid CPU-GPU code */
if ( ( left && (! notran) ) || ( (! left) && notran ) ) {
i1 = 1;
i2 = k;
step = nb;
} else {
i1 = ((k - 1)/nb)*nb + 1;
i2 = 1;
step = -nb;
}
// silence "uninitialized" warnings
mi = 0;
ni = 0;
if (left) {
ni = n;
jc = 1;
} else {
mi = m;
ic = 1;
}
magmablas_csetdiag1subdiag0('L', k, nb, da, ldda);
// for i=i1 to i2 by step
for (i = i1; (step < 0 ? i >= i2 : i <= i2); i += step) {
ib = min(nb, k - i + 1);
/* Form the triangular factor of the block reflector
H = H(i) H(i+1) . . . H(i+ib-1) */
i__4 = nq - i + 1;
lapackf77_clarft("F", "C", &i__4, &ib, &wa[i + i*ldwa], &ldwa,
&tau[i], t, &ib);
if (left) {
/* H or H' is applied to C(i:m,1:n) */
mi = m - i + 1;
ic = i;
}
else {
/* H or H' is applied to C(1:m,i:n) */
ni = n - i + 1;
jc = i;
}
if (left)
lddwork = ni;
else
lddwork = mi;
/* Apply H or H'; First copy T to the GPU */
magma_csetmatrix( ib, ib, t, ib, dwork, ib );
magma_clarfb_gpu( side, trans, MagmaForward, MagmaColumnwise,
mi, ni, ib,
da + (i - 1) + (i - 1)*ldda , ldda, dwork, ib,
&dc[ic + jc*lddc], lddc,
dwork + ib*ib, lddwork);
}
magma_free( dwork );
return *info;
} /* magma_cunmqr */
extern "C" magma_int_t
magma_cunmqr(magma_side_t side, magma_trans_t trans,
magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloatComplex *a, magma_int_t lda,
magmaFloatComplex *tau,
magmaFloatComplex *c, magma_int_t ldc,
magmaFloatComplex *work, magma_int_t lwork,
magma_int_t *info, magma_queue_t queue)
{
/* -- MAGMA (version 1.0.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
September 2012
Purpose
=======
CUNMQR overwrites the general complex M-by-N matrix C with
SIDE = 'L' SIDE = 'R'
TRANS = 'N': Q * C C * Q
TRANS = 'T': Q**H * C C * Q**H
where Q is a complex orthogonal matrix defined as the product of k
elementary reflectors
Q = H(1) H(2) . . . H(k)
as returned by CGEQRF. Q is of order M if SIDE = 'L' and of order N
if SIDE = 'R'.
Arguments
=========
SIDE (input) CHARACTER*1
= 'L': apply Q or Q**H from the Left;
= 'R': apply Q or Q**H from the Right.
TRANS (input) CHARACTER*1
= 'N': No transpose, apply Q;
= 'T': Transpose, apply Q**H.
M (input) INTEGER
The number of rows of the matrix C. M >= 0.
N (input) INTEGER
The number of columns of the matrix C. N >= 0.
K (input) INTEGER
The number of elementary reflectors whose product defines
the matrix Q.
If SIDE = 'L', M >= K >= 0;
if SIDE = 'R', N >= K >= 0.
A (input) COMPLEX array, dimension (LDA,K)
The i-th column must contain the vector which defines the
elementary reflector H(i), for i = 1,2,...,k, as returned by
CGEQRF in the first k columns of its array argument A.
A is modified by the routine but restored on exit.
LDA (input) INTEGER
The leading dimension of the array A.
If SIDE = 'L', LDA >= max(1,M);
if SIDE = 'R', LDA >= max(1,N).
TAU (input) COMPLEX array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by CGEQRF.
C (input/output) COMPLEX array, dimension (LDC,N)
On entry, the M-by-N matrix C.
On exit, C is overwritten by Q*C or Q**H * C or C * Q**H or C*Q.
LDC (input) INTEGER
The leading dimension of the array C. LDC >= max(1,M).
WORK (workspace/output) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(0) returns the optimal LWORK.
LWORK (input) INTEGER
The dimension of the array WORK.
If SIDE = 'L', LWORK >= max(1,N);
if SIDE = 'R', LWORK >= max(1,M).
For optimum performance LWORK >= N*NB if SIDE = 'L', and
LWORK >= M*NB if SIDE = 'R', 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.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
===================================================================== */
magmaFloatComplex c_one = MAGMA_C_ONE;
magma_side_t side_ = side;
magma_trans_t trans_ = trans;
/* Allocate work space on the GPU */
magmaFloatComplex_ptr dwork, dc;
magma_malloc( &dc, (m)*(n)*sizeof(magmaFloatComplex) );
magma_malloc( &dwork, (m + n + 64)*64*sizeof(magmaFloatComplex) );
/* Copy matrix C from the CPU to the GPU */
magma_csetmatrix( m, n, c, 0, ldc, dc, 0, m, queue );
//dc -= (1 + m);
size_t dc_offset = -(1+m);
magma_int_t a_offset, c_offset, i__4, lddwork;
magma_int_t i__;
magmaFloatComplex t[2*4160] /* was [65][64] */;
magma_int_t i1, i2, i3, ib, ic, jc, nb, mi, ni, nq, nw;
int left, notran, lquery;
magma_int_t iinfo, lwkopt;
a_offset = 1 + lda;
a -= a_offset;
--tau;
c_offset = 1 + ldc;
c -= c_offset;
*info = 0;
left = lapackf77_lsame(lapack_const(side_), "L");
notran = lapackf77_lsame(lapack_const(trans_), "N");
lquery = (lwork == -1);
/* NQ is the order of Q and NW is the minimum dimension of WORK */
if (left) {
nq = m;
nw = n;
} else {
nq = n;
nw = m;
}
if (! left && ! lapackf77_lsame(lapack_const(side_), "R")) {
*info = -1;
} else if (! notran && ! lapackf77_lsame(lapack_const(trans_), "T")) {
*info = -2;
} else if (m < 0) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (k < 0 || k > nq) {
*info = -5;
} else if (lda < max(1,nq)) {
*info = -7;
} else if (ldc < max(1,m)) {
*info = -10;
} else if (lwork < max(1,nw) && ! lquery) {
*info = -12;
}
if (*info == 0)
{
/* Determine the block size. NB may be at most NBMAX, where NBMAX
is used to define the local array T. */
nb = 64;
lwkopt = max(1,nw) * nb;
MAGMA_C_SET2REAL( work[0], lwkopt );
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0 || k == 0) {
work[0] = c_one;
return *info;
}
if (nb >= k)
{
/* Use CPU code */
lapackf77_cunmqr(lapack_const(side_), lapack_const(trans_), &m, &n, &k, &a[a_offset], &lda, &tau[1],
&c[c_offset], &ldc, work, &lwork, &iinfo);
}
else
{
/* Use hybrid CPU-GPU code */
if ( ( left && (! notran) ) || ( (! left) && notran ) ) {
i1 = 1;
i2 = k;
i3 = nb;
} else {
i1 = (k - 1) / nb * nb + 1;
i2 = 1;
i3 = -nb;
}
if (left) {
ni = n;
jc = 1;
} else {
mi = m;
ic = 1;
}
for (i__ = i1; i3 < 0 ? i__ >= i2 : i__ <= i2; i__ += i3)
{
ib = min(nb, k - i__ + 1);
/* Form the triangular factor of the block reflector
H = H(i) H(i+1) . . . H(i+ib-1) */
i__4 = nq - i__ + 1;
lapackf77_clarft("F", "C", &i__4, &ib, &a[i__ + i__ * lda], &lda,
&tau[i__], t, &ib);
/* 1) Put 0s in the upper triangular part of A;
2) copy the panel from A to the GPU, and
3) restore A */
cpanel_to_q(MagmaUpper, ib, &a[i__ + i__ * lda], lda, t+ib*ib);
magma_csetmatrix( i__4, ib, &a[i__ + i__ * lda], 0, lda, dwork, 0, i__4, queue );
cq_to_panel(MagmaUpper, ib, &a[i__ + i__ * lda], lda, t+ib*ib);
if (left)
{
/* H or H' is applied to C(i:m,1:n) */
mi = m - i__ + 1;
ic = i__;
}
else
{
/* H or H' is applied to C(1:m,i:n) */
ni = n - i__ + 1;
jc = i__;
}
if (left)
lddwork = ni;
else
lddwork = mi;
/* Apply H or H'; First copy T to the GPU */
magma_csetmatrix( ib, ib, t, 0, ib, dwork, i__4*ib, ib, queue );
magma_clarfb_gpu( side, trans, MagmaForward, MagmaColumnwise,
mi, ni, ib,
dwork, 0, i__4, dwork, i__4*ib, ib,
dc, dc_offset+(ic + jc * m), m,
dwork, (i__4*ib + ib*ib), lddwork, queue);
}
magma_cgetmatrix( m, n, dc, dc_offset+(1+m), m, &c[c_offset], 0, ldc, queue );
}
MAGMA_C_SET2REAL( work[0], lwkopt );
//dc += (1 + m);
magma_free( dc );
magma_free( dwork );
return *info;
} /* magma_cunmqr */
extern "C" magma_int_t
magma_cgeqrf_ooc(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.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CGEQRF_OOC 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 is an out-of-core (ooc) version that is similar to magma_cgeqrf but
the difference is that this version can use a GPU even if the matrix
does not fit into the GPU memory at once.
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 >= 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).
===================================================================== */
#define a_ref(a_1,a_2) ( a+(a_2)*(lda) + (a_1))
#define da_ref(a_1,a_2) (da+(a_2)*ldda + (a_1))
magmaFloatComplex *da, *dwork;
magmaFloatComplex c_one = MAGMA_C_ONE;
int k, lddwork, 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 (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,m)) {
*info = -4;
} else if (lwork < max(1,n) && ! lquery) {
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
/* Check how much memory do we have */
size_t freeMem, totalMem;
cudaMemGetInfo( &freeMem, &totalMem );
freeMem /= sizeof(magmaFloatComplex);
magma_int_t IB, NB = (magma_int_t)(0.8*freeMem/m);
NB = (NB / nb) * nb;
if (NB >= n)
return magma_cgeqrf(m, n, a, lda, tau, work, lwork, info);
k = min(m,n);
if (k == 0) {
work[0] = c_one;
return *info;
}
lddwork = ((NB+31)/32)*32+nb;
ldda = ((m+31)/32)*32;
if (MAGMA_SUCCESS != magma_cmalloc( &da, (NB + nb)*ldda + nb*lddwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_queue_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
// magmablasSetKernelStream(stream[1]);
magmaFloatComplex *ptr = da + ldda * NB;
dwork = da + ldda*(NB + nb);
/* start the main loop over the blocks that fit in the GPU memory */
for(int i=0; i<n; i+=NB) {
IB = min(n-i, NB);
//printf("Processing %5d columns -- %5d to %5d ... \n", IB, i, i+IB);
/* 1. Copy the next part of the matrix to the GPU */
magma_csetmatrix_async( (m), IB,
a_ref(0,i), lda,
da_ref(0,0), ldda, stream[0] );
magma_queue_sync( stream[0] );
/* 2. Update it with the previous transformations */
for(int j=0; j<min(i,k); j+=nb) {
magma_int_t ib = min(k-j, nb);
/* Get a panel in ptr. */
// 1. Form the triangular factor of the block reflector
// 2. Send it to the GPU.
// 3. Put 0s in the upper triangular part of V.
// 4. Send V to the GPU in ptr.
// 5. Update the matrix.
// 6. Restore the upper part of V.
magma_int_t rows = m-j;
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&rows, &ib, a_ref(j,j), &lda, tau+j, work, &ib);
magma_csetmatrix_async( ib, ib,
work, ib,
dwork, lddwork, stream[1] );
cpanel_to_q(MagmaUpper, ib, a_ref(j,j), lda, work+ib*ib);
magma_csetmatrix_async( rows, ib,
a_ref(j,j), lda,
ptr, rows, stream[1] );
magma_queue_sync( stream[1] );
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, IB, ib,
ptr, rows, dwork, lddwork,
da_ref(j, 0), ldda, dwork+ib, lddwork);
cq_to_panel(MagmaUpper, ib, a_ref(j,j), lda, work+ib*ib);
}
/* 3. Do a QR on the current part */
if (i<k)
magma_cgeqrf2_gpu(m-i, IB, da_ref(i,0), ldda, tau+i, info);
/* 4. Copy the current part back to the CPU */
magma_cgetmatrix_async( (m), IB,
da_ref(0,0), ldda,
a_ref(0,i), lda, stream[0] );
}
magma_queue_sync( stream[0] );
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free( da );
return *info;
} /* magma_cgeqrf_ooc */
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 */
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 */
/**
Purpose
-------
CUNGQR generates an M-by-N COMPLEX matrix Q with orthonormal columns,
which is defined as the first N columns of a product of K elementary
reflectors of order M
Q = H(1) H(2) . . . H(k)
as returned by CGEQRF.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix Q. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix Q. M >= N >= 0.
@param[in]
k INTEGER
The number of elementary reflectors whose product defines the
matrix Q. N >= K >= 0.
@param[in,out]
A COMPLEX array A, dimension (LDDA,N).
On entry, the i-th column must contain the vector
which defines the elementary reflector H(i), for
i = 1,2,...,k, as returned by CGEQRF_GPU in the
first k columns of its array argument A.
On exit, the M-by-N matrix Q.
@param[in]
lda INTEGER
The first dimension of the array A. LDA >= max(1,M).
@param[in]
tau COMPLEX array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by CGEQRF_GPU.
@param[in]
T COMPLEX array, dimension (NB, min(M,N)).
T contains the T matrices used in blocking the elementary
reflectors H(i), e.g., this can be the 6th argument of
magma_cgeqrf_gpu (except stored on the CPU, not the GPU).
@param[in]
nb INTEGER
This is the block size used in CGEQRF_GPU, and correspondingly
the size of the T matrices, used in the factorization, and
stored in T.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_cgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_cungqr_m(
magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloatComplex *A, magma_int_t lda,
magmaFloatComplex *tau,
magmaFloatComplex *T, magma_int_t nb,
magma_int_t *info)
{
#define A(i,j) ( A + (i) + (j)*lda )
#define dA(d,i,j) (dA[d] + (i) + (j)*ldda)
#define dT(d,i,j) (dT[d] + (i) + (j)*nb)
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magma_int_t m_kk, n_kk, k_kk, mi;
magma_int_t lwork, ldwork;
magma_int_t d, i, ib, j, jb, ki, kk;
magmaFloatComplex *work=NULL;
*info = 0;
if (m < 0) {
*info = -1;
} else if ((n < 0) || (n > m)) {
*info = -2;
} else if ((k < 0) || (k > n)) {
*info = -3;
} else if (lda < max(1,m)) {
*info = -5;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (n <= 0) {
return *info;
}
magma_int_t di, dn;
magma_int_t dpanel;
magma_int_t ngpu = magma_num_gpus();
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
// Allocate memory on GPUs for A and workspaces
magma_int_t ldda = magma_roundup( m, 32 );
magma_int_t lddwork = magma_roundup( n, 32 );
magma_int_t min_lblocks = (n / nb) / ngpu; // min. blocks per gpu
magma_int_t last_dev = (n / nb) % ngpu; // device with last block
magma_int_t nlocal[ MagmaMaxGPUs ] = { 0 };
magmaFloatComplex *dA[ MagmaMaxGPUs ] = { NULL };
magmaFloatComplex *dT[ MagmaMaxGPUs ] = { NULL };
magmaFloatComplex *dV[ MagmaMaxGPUs ] = { NULL };
magmaFloatComplex *dW[ MagmaMaxGPUs ] = { NULL };
magma_queue_t queues[ MagmaMaxGPUs ] = { NULL };
for( d = 0; d < ngpu; ++d ) {
// example with n = 75, nb = 10, ngpu = 3
// min_lblocks = 2
// last_dev = 1
// gpu 0: 2 blocks, cols: 0- 9, 30-39, 60-69
// gpu 1: 1+ blocks, cols: 10-19, 40-49, 70-74 (partial)
// gpu 2: 1 block, cols: 20-29, 50-59
magma_setdevice( d );
nlocal[d] = min_lblocks*nb;
if ( d < last_dev ) {
nlocal[d] += nb;
}
else if ( d == last_dev ) {
nlocal[d] += (n % nb);
}
ldwork = nlocal[d]*ldda // dA
+ nb*m // dT
+ nb*ldda // dV
+ nb*lddwork; // dW
if ( MAGMA_SUCCESS != magma_cmalloc( &dA[d], ldwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
goto cleanup;
}
dT[d] = dA[d] + nlocal[d]*ldda;
dV[d] = dT[d] + nb*m;
dW[d] = dV[d] + nb*ldda;
magma_queue_create( d, &queues[d] );
}
trace_init( 1, ngpu, 1, queues );
// first kk columns are handled by blocked method.
// ki is start of 2nd-to-last block
if ((nb > 1) && (nb < k)) {
ki = (k - nb - 1) / nb * nb;
kk = min(k, ki + nb);
} else {
ki = 0;
kk = 0;
}
// Allocate CPU work space
// n*nb for larfb work
// m*nb for V
// nb*nb for T
lwork = (n + m + nb) * nb;
magma_cmalloc_cpu( &work, lwork );
if (work == NULL) {
*info = MAGMA_ERR_HOST_ALLOC;
goto cleanup;
}
magmaFloatComplex *work_T, *work_V;
work_T = work + n*nb;
work_V = work + n*nb + nb*nb;
// Use unblocked code for the last or only block.
if (kk < n) {
trace_cpu_start( 0, "ungqr", "ungqr last block" );
m_kk = m - kk;
n_kk = n - kk;
k_kk = k - kk;
// cungqr requires less workspace (n*nb), but is slow if k < cungqr's block size.
// replacing it with the 4 routines below is much faster (e.g., 60x).
//magma_int_t iinfo;
//lapackf77_cungqr( &m_kk, &n_kk, &k_kk,
// A(kk, kk), &lda,
// &tau[kk], work, &lwork, &iinfo );
lapackf77_clacpy( MagmaFullStr, &m_kk, &k_kk, A(kk,kk), &lda, work_V, &m_kk);
lapackf77_claset( MagmaFullStr, &m_kk, &n_kk, &c_zero, &c_one, A(kk, kk), &lda );
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&m_kk, &k_kk,
work_V, &m_kk, &tau[kk], work_T, &k_kk);
lapackf77_clarfb( MagmaLeftStr, MagmaNoTransStr, MagmaForwardStr, MagmaColumnwiseStr,
&m_kk, &n_kk, &k_kk,
work_V, &m_kk, work_T, &k_kk, A(kk, kk), &lda, work, &n_kk );
if (kk > 0) {
for( j=kk; j < n; j += nb ) {
jb = min( n-j, nb );
d = (j / nb) % ngpu;
di = ((j / nb) / ngpu) * nb;
magma_setdevice( d );
magma_csetmatrix( m_kk, jb,
A(kk, j), lda,
dA(d, kk, di), ldda, queues[d] );
// Set A(1:kk,kk+1:n) to zero.
magmablas_claset( MagmaFull, kk, jb, c_zero, c_zero, dA(d, 0, di), ldda, queues[d] );
}
}
trace_cpu_end( 0 );
}
if (kk > 0) {
// Use blocked code
// send T to all GPUs
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
trace_gpu_start( d, 0, "set", "set T" );
magma_csetmatrix_async( nb, min(m,n), T, nb, dT[d], nb, queues[d] );
trace_gpu_end( d, 0 );
}
// queue: set Aii (V) --> laset --> laset --> larfb --> [next]
// CPU has no computation
for( i = ki; i >= 0; i -= nb ) {
ib = min(nb, k - i);
mi = m - i;
dpanel = (i / nb) % ngpu;
di = ((i / nb) / ngpu) * nb;
// Send current panel to dV on the GPUs
lapackf77_claset( "Upper", &ib, &ib, &c_zero, &c_one, A(i, i), &lda );
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
trace_gpu_start( d, 0, "set", "set V" );
magma_csetmatrix_async( mi, ib,
A(i, i), lda,
dV[d], ldda, queues[d] );
trace_gpu_end( d, 0 );
}
// set panel to identity
magma_setdevice( dpanel );
trace_gpu_start( dpanel, 0, "laset", "laset" );
magmablas_claset( MagmaFull, i, ib, c_zero, c_zero, dA(dpanel, 0, di), ldda, queues[dpanel] );
magmablas_claset( MagmaFull, mi, ib, c_zero, c_one, dA(dpanel, i, di), ldda, queues[dpanel] );
trace_gpu_end( dpanel, 0 );
if (i < n) {
// Apply H to A(i:m,i:n) from the left
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magma_indices_1D_bcyclic( nb, ngpu, d, i, n, &di, &dn );
trace_gpu_start( d, 0, "larfb", "larfb" );
magma_clarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
mi, dn-di, ib,
dV[d], ldda, dT(d,0,i), nb,
dA(d, i, di), ldda, dW[d], lddwork, queues[d] );
trace_gpu_end( d, 0 );
}
}
}
// copy result back to CPU
trace_cpu_start( 0, "get", "get A" );
magma_cgetmatrix_1D_col_bcyclic( m, n, dA, ldda, A, lda, ngpu, nb, queues );
trace_cpu_end( 0 );
}
#ifdef TRACING
char name[80];
snprintf( name, sizeof(name), "cungqr-n%d-ngpu%d.svg", m, ngpu );
trace_finalize( name, "trace.css" );
#endif
cleanup:
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magma_free( dA[d] );
magma_queue_destroy( queues[d] );
}
magma_free_cpu( work );
magma_setdevice( orig_dev );
return *info;
} /* magma_cungqr */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing clarfb_gpu
*/
int main( int argc, char** argv )
{
TESTING_INIT();
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t M, N, K, size, ldc, ldv, ldt, ldw, nv;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
float error, work[1];
magma_int_t status = 0;
// test all combinations of input parameters
magma_side_t side [] = { MagmaLeft, MagmaRight };
magma_trans_t trans [] = { MagmaConjTrans, MagmaNoTrans };
magma_direct_t direct[] = { MagmaForward, MagmaBackward };
magma_storev_t storev[] = { MagmaColumnwise, MagmaRowwise };
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
printf(" M N K storev side direct trans ||R||_F / ||HC||_F\n");
printf("========================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
M = opts.msize[itest];
N = opts.nsize[itest];
K = opts.ksize[itest];
if ( M < K || N < K || K <= 0 ) {
printf( "%5d %5d %5d skipping because clarfb requires M >= K, N >= K, K >= 0\n",
(int) M, (int) N, (int) K );
continue;
}
for( int istor = 0; istor < 2; ++istor ) {
for( int iside = 0; iside < 2; ++iside ) {
for( int idir = 0; idir < 2; ++idir ) {
for( int itran = 0; itran < 2; ++itran ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
ldc = ((M+31)/32)*32;
ldt = ((K+31)/32)*32;
ldw = (side[iside] == MagmaLeft ? N : M);
// (ldv, nv) get swapped later if rowwise
ldv = (side[iside] == MagmaLeft ? M : N);
nv = K;
// Allocate memory for matrices
magmaFloatComplex *C, *R, *V, *T, *W;
TESTING_MALLOC_CPU( C, magmaFloatComplex, ldc*N );
TESTING_MALLOC_CPU( R, magmaFloatComplex, ldc*N );
TESTING_MALLOC_CPU( V, magmaFloatComplex, ldv*K );
TESTING_MALLOC_CPU( T, magmaFloatComplex, ldt*K );
TESTING_MALLOC_CPU( W, magmaFloatComplex, ldw*K );
magmaFloatComplex_ptr dC, dV, dT, dW;
TESTING_MALLOC_DEV( dC, magmaFloatComplex, ldc*N );
TESTING_MALLOC_DEV( dV, magmaFloatComplex, ldv*K );
TESTING_MALLOC_DEV( dT, magmaFloatComplex, ldt*K );
TESTING_MALLOC_DEV( dW, magmaFloatComplex, ldw*K );
// C is M x N.
size = ldc*N;
lapackf77_clarnv( &ione, ISEED, &size, C );
//printf( "C=" ); magma_cprint( M, N, C, ldc );
// V is ldv x nv. See larfb docs for description.
// if column-wise and left, M x K
// if column-wise and right, N x K
// if row-wise and left, K x M
// if row-wise and right, K x N
size = ldv*nv;
lapackf77_clarnv( &ione, ISEED, &size, V );
if ( storev[istor] == MagmaColumnwise ) {
if ( direct[idir] == MagmaForward ) {
lapackf77_claset( MagmaUpperStr, &K, &K, &c_zero, &c_one, V, &ldv );
}
else {
lapackf77_claset( MagmaLowerStr, &K, &K, &c_zero, &c_one, &V[(ldv-K)], &ldv );
}
}
else {
// rowwise, swap V's dimensions
std::swap( ldv, nv );
if ( direct[idir] == MagmaForward ) {
lapackf77_claset( MagmaLowerStr, &K, &K, &c_zero, &c_one, V, &ldv );
}
else {
lapackf77_claset( MagmaUpperStr, &K, &K, &c_zero, &c_one, &V[(nv-K)*ldv], &ldv );
}
}
//printf( "# ldv %d, nv %d\n", ldv, nv );
//printf( "V=" ); magma_cprint( ldv, nv, V, ldv );
// T is K x K, upper triangular for forward, and lower triangular for backward
magma_int_t k1 = K-1;
size = ldt*K;
lapackf77_clarnv( &ione, ISEED, &size, T );
if ( direct[idir] == MagmaForward ) {
lapackf77_claset( MagmaLowerStr, &k1, &k1, &c_zero, &c_zero, &T[1], &ldt );
}
else {
lapackf77_claset( MagmaUpperStr, &k1, &k1, &c_zero, &c_zero, &T[1*ldt], &ldt );
}
//printf( "T=" ); magma_cprint( K, K, T, ldt );
magma_csetmatrix( M, N, C, ldc, dC, ldc );
magma_csetmatrix( ldv, nv, V, ldv, dV, ldv );
magma_csetmatrix( K, K, T, ldt, dT, ldt );
lapackf77_clarfb( lapack_side_const( side[iside] ), lapack_trans_const( trans[itran] ),
lapack_direct_const( direct[idir] ), lapack_storev_const( storev[istor] ),
&M, &N, &K,
V, &ldv, T, &ldt, C, &ldc, W, &ldw );
//printf( "HC=" ); magma_cprint( M, N, C, ldc );
magma_clarfb_gpu( side[iside], trans[itran], direct[idir], storev[istor],
M, N, K,
dV, ldv, dT, ldt, dC, ldc, dW, ldw );
magma_cgetmatrix( M, N, dC, ldc, R, ldc );
//printf( "dHC=" ); magma_cprint( M, N, R, ldc );
// compute relative error |HC_magma - HC_lapack| / |HC_lapack|
error = lapackf77_clange( "Fro", &M, &N, C, &ldc, work );
size = ldc*N;
blasf77_caxpy( &size, &c_neg_one, C, &ione, R, &ione );
error = lapackf77_clange( "Fro", &M, &N, R, &ldc, work ) / error;
printf( "%5d %5d %5d %c %c %c %c %8.2e %s\n",
(int) M, (int) N, (int) K,
lapacke_storev_const(storev[istor]), lapacke_side_const(side[iside]),
lapacke_direct_const(direct[idir]), lapacke_trans_const(trans[itran]),
error, (error < tol ? "ok" : "failed") );
status += ! (error < tol);
TESTING_FREE_CPU( C );
TESTING_FREE_CPU( R );
TESTING_FREE_CPU( V );
TESTING_FREE_CPU( T );
TESTING_FREE_CPU( W );
TESTING_FREE_DEV( dC );
TESTING_FREE_DEV( dV );
TESTING_FREE_DEV( dT );
TESTING_FREE_DEV( dW );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}}}}
printf( "\n" );
}
TESTING_FINALIZE();
return status;
}
/**
Purpose
-------
CUNMLQ overwrites the general complex M-by-N matrix C with
@verbatim
SIDE = MagmaLeft SIDE = MagmaRight
TRANS = MagmaNoTrans: Q * C C * Q
TRANS = Magma_ConjTrans: Q**H * C C * Q**H
@endverbatim
where Q is a complexunitary matrix defined as the product of k
elementary reflectors
Q = H(k)**H . . . H(2)**H H(1)**H
as returned by CGELQF. Q is of order M if SIDE = MagmaLeft and of order N
if SIDE = MagmaRight.
Arguments
---------
@param[in]
side magma_side_t
- = MagmaLeft: apply Q or Q**H from the Left;
- = MagmaRight: apply Q or Q**H from the Right.
@param[in]
trans magma_trans_t
- = MagmaNoTrans: No transpose, apply Q;
- = Magma_ConjTrans: Conjugate transpose, apply Q**H.
@param[in]
m INTEGER
The number of rows of the matrix C. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix C. N >= 0.
@param[in]
k INTEGER
The number of elementary reflectors whose product defines
the matrix Q.
If SIDE = MagmaLeft, M >= K >= 0;
if SIDE = MagmaRight, N >= K >= 0.
@param[in]
A COMPLEX array, dimension
(LDA,M) if SIDE = MagmaLeft,
(LDA,N) if SIDE = MagmaRight.
The i-th row must contain the vector which defines the
elementary reflector H(i), for i = 1,2,...,k, as returned by
CGELQF in the first k rows of its array argument A.
A is modified by the routine but restored on exit.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,K).
@param[in]
tau COMPLEX array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by CGELQF.
@param[in,out]
C COMPLEX array, dimension (LDC,N)
On entry, the M-by-N matrix C.
On exit, C is overwritten by Q*C or Q**H*C or C*Q**H or C*Q.
@param[in]
ldc INTEGER
The leading dimension of the array C. LDC >= max(1,M).
@param[out]
work (workspace) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK.
If SIDE = MagmaLeft, LWORK >= max(1,N);
if SIDE = MagmaRight, LWORK >= max(1,M).
For optimum performance
if SIDE = MagmaLeft, LWORK >= N*NB;
if SIDE = MagmaRight, LWORK >= M*NB,
where NB is the optimal blocksize.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_cgelqf_comp
********************************************************************/
extern "C" magma_int_t
magma_cunmlq(
magma_side_t side, magma_trans_t trans,
magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloatComplex *A, magma_int_t lda,
magmaFloatComplex *tau,
magmaFloatComplex *C, magma_int_t ldc,
magmaFloatComplex *work, magma_int_t lwork,
magma_int_t *info)
{
#define A(i_,j_) ( A + (i_) + (j_)*lda)
#define dC(i_,j_) (dC + (i_) + (j_)*lddc)
#define dV(i_,j_) (dV + (i_) + (j_)*ib)
#define dT(i_,j_) (dT + (i_) + (j_)*ib)
#define dwork(i_) (dwork + (i_))
magmaFloatComplex *T, *T2;
magma_int_t i, i1, i2, ib, ic, jc, nb, mi, ni, nq, nq_i, nw, step;
magma_int_t iinfo, ldwork, lwkopt;
magma_int_t left, notran, lquery;
magma_trans_t transt;
*info = 0;
left = (side == MagmaLeft);
notran = (trans == MagmaNoTrans);
lquery = (lwork == -1);
/* NQ is the order of Q and NW is the minimum dimension of WORK */
if (left) {
nq = m;
nw = n;
} else {
nq = n;
nw = m;
}
/* Test the input arguments */
if (! left && side != MagmaRight) {
*info = -1;
} else if (! notran && trans != Magma_ConjTrans) {
*info = -2;
} else if (m < 0) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (k < 0 || k > nq) {
*info = -5;
} else if (lda < max(1,k)) {
*info = -7;
} else if (ldc < max(1,m)) {
*info = -10;
} else if (lwork < max(1,nw) && ! lquery) {
*info = -12;
}
if (*info == 0) {
nb = magma_get_cgelqf_nb( min( m, n ));
lwkopt = max(1,nw)*nb;
work[0] = MAGMA_C_MAKE( lwkopt, 0 );
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0 || k == 0) {
work[0] = MAGMA_C_ONE;
return *info;
}
ldwork = nw;
if (nb >= k) {
/* Use CPU code */
lapackf77_cunmlq( lapack_side_const(side), lapack_trans_const(trans),
&m, &n, &k, A, &lda, tau, C, &ldc, work, &lwork, &iinfo);
}
else {
/* Use hybrid CPU-GPU code */
/* Allocate work space on the GPU.
* nw*nb for dwork (m or n) by nb
* nq*nb for dV (n or m) by nb
* nb*nb for dT
* lddc*n for dC.
*/
magma_int_t lddc = ((m+31)/32)*32;
magmaFloatComplex_ptr dwork, dV, dT, dC;
magma_cmalloc( &dwork, (nw + nq + nb)*nb + lddc*n );
if ( dwork == NULL ) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dV = dwork + nw*nb;
dT = dV + nq*nb;
dC = dT + nb*nb;
/* work space on CPU.
* nb*nb for T
* nb*nb for T2, used to save and restore diagonal block of panel */
magma_cmalloc_cpu( &T, 2*nb*nb );
if ( T == NULL ) {
magma_free( dwork );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
T2 = T + nb*nb;
/* Copy matrix C from the CPU to the GPU */
magma_csetmatrix( m, n, C, ldc, dC(0,0), lddc );
if ( (left && notran) || (! left && ! notran) ) {
i1 = 0;
i2 = k;
step = nb;
} else {
i1 = ((k - 1) / nb)*nb;
i2 = 0;
step = -nb;
}
// silence "uninitialized" warnings
mi = 0;
ni = 0;
if (left) {
ni = n;
jc = 0;
} else {
mi = m;
ic = 0;
}
if (notran) {
transt = Magma_ConjTrans;
} else {
transt = MagmaNoTrans;
}
for (i = i1; (step < 0 ? i >= i2 : i < i2); i += step) {
ib = min(nb, k - i);
/* Form the triangular factor of the block reflector
H = H(i) H(i + 1) . . . H(i + ib-1) */
nq_i = nq - i;
lapackf77_clarft("Forward", "Rowwise", &nq_i, &ib,
A(i,i), &lda, &tau[i], T, &ib);
/* 1) set upper triangle of panel in A to identity,
2) copy the panel from A to the GPU, and
3) restore A */
cpanel_to_q( MagmaLower, ib, A(i,i), lda, T2 );
magma_csetmatrix( ib, nq_i, A(i,i), lda, dV(0,0), ib );
cq_to_panel( MagmaLower, ib, A(i,i), lda, T2 );
if (left) {
/* H or H**H is applied to C(i:m,1:n) */
mi = m - i;
ic = i;
}
else {
/* H or H**H is applied to C(1:m,i:n) */
ni = n - i;
jc = i;
}
/* Apply H or H**H; First copy T to the GPU */
magma_csetmatrix( ib, ib, T, ib, dT(0,0), ib );
magma_clarfb_gpu( side, transt, MagmaForward, MagmaRowwise,
mi, ni, ib,
dV(0,0), ib,
dT(0,0), ib,
dC(ic,jc), lddc,
dwork(0), ldwork );
}
magma_cgetmatrix( m, n, dC(0,0), lddc, C, ldc );
magma_free( dwork );
magma_free_cpu( T );
}
work[0] = MAGMA_C_MAKE( lwkopt, 0 );
return *info;
} /* magma_cunmlq */
magma_int_t magma_cungqr_2stage_gpu(magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloatComplex *da, magma_int_t ldda,
magmaFloatComplex *tau, magmaFloatComplex *dT,
magma_int_t nb, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CUNGQR generates an M-by-N COMPLEX matrix Q with orthonormal columns,
which is defined as the first N columns of a product of K elementary
reflectors of order M
Q = H(1) H(2) . . . H(k)
as returned by CGEQRF_GPU.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix Q. M >= 0.
N (input) INTEGER
The number of columns of the matrix Q. M >= N >= 0.
K (input) INTEGER
The number of elementary reflectors whose product defines the
matrix Q. N >= K >= 0.
DA (input/output) COMPLEX array A on the GPU device,
dimension (LDDA,N). On entry, the i-th column must contain
the vector which defines the elementary reflector H(i), for
i = 1,2,...,k, as returned by CGEQRF_GPU in the first k
columns of its array argument A.
On exit, the M-by-N matrix Q.
LDDA (input) INTEGER
The first dimension of the array A. LDDA >= max(1,M).
TAU (input) COMPLEX array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by CGEQRF_GPU.
DT (input) COMPLEX work space array on the GPU device,
dimension (MIN(M, N) )*NB.
This must be the 6th argument of magma_cgeqrf_gpu
[ note that if N here is bigger than N in magma_cgeqrf_gpu,
the workspace requirement DT in magma_cgeqrf_gpu must be
as specified in this routine ].
NB (input) INTEGER
This is the block size used in CGEQRF_GPU, and correspondingly
the size of the T matrices, used in the factorization, and
stored in DT.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument has an illegal value
===================================================================== */
#define da_ref(a_1,a_2) (da+(a_2)*(ldda) + (a_1))
#define t_ref(a_1) (dT+(a_1)*nb)
magma_int_t i__1, i__2, i__3;
//magma_int_t lwork;
magma_int_t i, ib, ki, kk; //, iinfo;
//magma_int_t lddwork = min(m, n);
//magmaFloatComplex *work, *panel;
magmaFloatComplex *dwork;
//magma_queue_t stream[2];
magma_int_t ldt=nb; // need to be an input parameter
*info = 0;
if (m < 0) {
*info = -1;
} else if ((n < 0) || (n > m)) {
*info = -2;
} else if ((k < 0) || (k > n)) {
*info = -3;
} else if (ldda < max(1,m)) {
*info = -5;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (n <= 0)
return *info;
if(MAGMA_SUCCESS != magma_cmalloc( &dwork, n*nb )) {
printf ("!!!! cungqr_2stage magma_alloc failed for: dwork\n" );
exit(-1);
}
if ( (nb > 1) && (nb < k) ) {
/* Use blocked code after the last block.
The first kk columns are handled by the block method.
ki is start of 2nd-to-last block. */
ki = (k - nb - 1) / nb * nb;
kk = min(k, ki + nb);
/* Set A(1:kk,kk+1:n) to zero. */
magmablas_claset(MagmaUpperLower, kk, n-kk, da_ref(0,kk), ldda);
/* A(kk+1:m, kk+1:n) = I */
magmablas_claset_identity(m-kk, n-kk, da_ref(kk,kk), ldda);
}
else {
ki = 0;
kk = 0;
}
/* Allocate work space on CPU in pinned memory */
//lwork = (n+m) * nb;
//if (kk < n)
// lwork = max(lwork, n * nb + (m-kk)*(n-kk));
//if (MAGMA_SUCCESS != magma_cmalloc_pinned( &work, (lwork) )) {
// *info = MAGMA_ERR_HOST_ALLOC;
// return *info;
//}
//panel = work + n * nb;
//magma_queue_create( &stream[0] );
//magma_queue_create( &stream[1] );
/* Use unblocked code for the last or only block. */
if (kk < n) {
i__1 = m - kk;
i__2 = n - kk;
i__3 = k - kk;
//cublasGetMatrix(i__1, i__2, sizeof(magmaFloatComplex),
// da_ref(kk, kk), ldda, panel, i__1);
//lapackf77_cungqr(&i__1, &i__2, &i__3, panel, &i__1, &tau[kk],
// work, &lwork, &iinfo);
//
//cublasSetMatrix(i__1, i__2, sizeof(magmaFloatComplex),
// panel, i__1, da_ref(kk, kk), ldda);
magma_clarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
i__1, i__2, i__3,
da_ref(kk, kk-nb), ldda, t_ref(kk-nb), ldt,
da_ref(kk, kk), ldda, dwork, i__2);
//magmablas_claset(MagmaUpperLower, kk-nb, nb, da_ref(0,kk-nb), ldda);
//magmablas_claset_identity(m-(kk-nb), nb, da_ref(kk-nb,kk-nb), ldda);
}
if (kk > 0) {
/* Use blocked code */
for (i = ki; i >= nb; i-=nb) {
ib = min(nb, k - i);
/* Send current panel to the CPU for update */
i__2 = m - i;
//cudaMemcpy2DAsync(panel, i__2 * sizeof(magmaFloatComplex),
// da_ref(i,i), ldda * sizeof(magmaFloatComplex),
// sizeof(magmaFloatComplex)*i__2, ib,
// cudaMemcpyDeviceToHost,stream[0]);
if (i + ib < n) {
/* Apply H to A(i:m,i+ib:n) from the left */
i__3 = n - i;
magmablas_claset(MagmaUpperLower, i, ib, da_ref(0,i), ldda);
magmablas_claset_identity(m-i, ib, da_ref(i,i), ldda);
magma_clarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
i__2, i__3, ib,
da_ref(i, i-nb), ldda, t_ref(i-nb), ldt,
da_ref(i, i), ldda, dwork, i__3);
}
/* Apply H to rows i:m of current block on the CPU */
//magma_queue_sync( stream[0] );
//lapackf77_cungqr(&i__2, &ib, &ib, panel, &i__2, &tau[i],
// work, &lwork, &iinfo);
//cudaMemcpy2DAsync(da_ref(i,i), ldda * sizeof(magmaFloatComplex),
// panel, i__2 * sizeof(magmaFloatComplex),
// sizeof(magmaFloatComplex)*i__2, ib,
// cudaMemcpyHostToDevice,stream[1]);
/* Set rows 1:i-1 of current block to zero */
i__2 = i + ib;
//magmablas_claset(MagmaUpperLower, i-ib, ib, da_ref(0,i-ib), ldda);
//magmablas_claset_identity(m-(i-ib), ib, da_ref(i-ib,i-ib), ldda);
}
}
magmablas_claset_identity(m, nb, da_ref(0,0), ldda);
magma_free( dwork );
//magma_free_pinned( work );
//magma_queue_destroy( stream[0] );
//magma_queue_destroy( stream[1] );
return *info;
} /* magma_cungqr_gpu */
/**
Purpose
-------
CGEQLF computes a QL factorization of a COMPLEX M-by-N matrix A:
A = Q * L.
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, if m >= n, the lower triangle of the subarray
A(m-n+1:m,1:n) contains the N-by-N lower triangular matrix L;
if m <= n, the elements on and below the (n-m)-th
superdiagonal contain the M-by-N lower trapezoidal matrix L;
the remaining elements, with the array TAU, represent the
orthogonal matrix Q as a product of 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(1,N,2*NB^2).
For optimum performance LWORK >= max(N*NB, 2*NB^2) where NB can be obtained
through magma_get_cgeqlf_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 by XERBLA.
@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(k) . . . H(2) H(1), 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(m-k+i+1:m) = 0 and v(m-k+i) = 1; v(1:m-k+i-1) is stored on exit in
A(1:m-k+i-1,n-k+i), and tau in TAU(i).
@ingroup magma_cgeqlf_comp
********************************************************************/
extern "C" magma_int_t
magma_cgeqlf(
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)
#define dwork(i_) (dwork + (i_))
magmaFloatComplex_ptr dA, dwork;
magmaFloatComplex c_one = MAGMA_C_ONE;
magma_int_t i, k, lddwork, old_i, old_ib, nb;
magma_int_t rows, cols;
magma_int_t ib, ki, kk, mu, nu, iinfo, ldda;
int lquery;
nb = magma_get_cgeqlf_nb(m);
*info = 0;
lquery = (lwork == -1);
// silence "uninitialized" warnings
old_ib = nb;
old_i = 0;
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,m)) {
*info = -4;
}
k = min(m,n);
if (*info == 0) {
if (k == 0)
work[0] = c_one;
else {
work[0] = MAGMA_C_MAKE( max(n*nb, 2*nb*nb), 0 );
}
if (lwork < max(max(1,n), 2*nb*nb) && ! lquery)
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
/* Quick return if possible */
if (k == 0)
return *info;
lddwork = ((n+31)/32)*32;
ldda = ((m+31)/32)*32;
if (MAGMA_SUCCESS != magma_cmalloc( &dA, (n)*ldda + nb*lddwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dwork = dA + ldda*n;
magma_queue_t queues[2];
magma_queue_create( &queues[0] );
magma_queue_create( &queues[1] );
if ( (nb > 1) && (nb < k) ) {
/* Use blocked code initially.
The last kk columns are handled by the block method.
First, copy the matrix on the GPU except the last kk columns */
magma_csetmatrix_async( m, n-nb,
A(0, 0), lda,
dA(0, 0), ldda, queues[0] );
ki = ((k - nb - 1) / nb) * nb;
kk = min(k, ki + nb);
for (i = k - kk + ki; i >= k -kk; i -= nb) {
ib = min(k-i,nb);
if (i < k - kk + ki) {
/* 1. Copy asynchronously the current panel to the CPU.
2. Copy asynchronously the submatrix below the panel
to the CPU) */
rows = m - k + i + ib;
magma_cgetmatrix_async( rows, ib,
dA(0, n-k+i), ldda,
A(0, n-k+i), lda, queues[1] );
magma_cgetmatrix_async( m-rows, ib,
dA(rows, n-k+i), ldda,
A(rows, n-k+i), lda, queues[0] );
/* Apply H' to A(1:m-k+i+ib-1,1:n-k+i-1) from the left in
two steps - implementing the lookahead techniques.
This is the main update from the lookahead techniques. */
rows = m - k + old_i + old_ib;
cols = n - k + old_i - old_ib;
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaBackward, MagmaColumnwise,
rows, cols, old_ib,
dA(0, cols+old_ib), ldda, dwork(0), lddwork,
dA(0, 0 ), ldda, dwork(old_ib), lddwork);
}
magma_queue_sync( queues[1] );
/* Compute the QL factorization of the current block
A(1:m-k+i+ib-1,n-k+i:n-k+i+ib-1) */
rows = m - k + i + ib;
cols = n - k + i;
lapackf77_cgeqlf( &rows, &ib, A(0,cols), &lda, tau+i, work, &lwork, &iinfo );
if (cols > 0) {
/* Form the triangular factor of the block reflector
H = H(i+ib-1) . . . H(i+1) H(i) */
lapackf77_clarft( MagmaBackwardStr, MagmaColumnwiseStr,
&rows, &ib,
A(0, cols), &lda, tau + i, work, &ib);
cpanel_to_q( MagmaLower, ib, A(rows-ib,cols), lda, work+ib*ib);
magma_csetmatrix( rows, ib,
A(0,cols), lda,
dA(0,cols), ldda );
cq_to_panel( MagmaLower, ib, A(rows-ib,cols), lda, work+ib*ib);
// Send the triangular part on the GPU
magma_csetmatrix( ib, ib, work, ib, dwork(0), lddwork );
/* Apply H' to A(1:m-k+i+ib-1,1:n-k+i-1) from the left in
two steps - implementing the lookahead techniques.
This is the update of first ib columns. */
if (i-ib >= k -kk)
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaBackward, MagmaColumnwise,
rows, ib, ib,
dA(0, cols), ldda, dwork(0), lddwork,
dA(0,cols-ib), ldda, dwork(ib), lddwork);
else {
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaBackward, MagmaColumnwise,
rows, cols, ib,
dA(0, cols), ldda, dwork(0), lddwork,
dA(0, 0 ), ldda, dwork(ib), lddwork);
}
old_i = i;
old_ib = ib;
}
}
mu = m - k + i + nb;
nu = n - k + i + nb;
magma_cgetmatrix( m, nu, dA(0,0), ldda, A(0,0), lda );
} else {
mu = m;
nu = n;
}
/* Use unblocked code to factor the last or only block */
if (mu > 0 && nu > 0)
lapackf77_cgeqlf(&mu, &nu, A(0,0), &lda, tau, work, &lwork, &iinfo);
magma_queue_destroy( queues[0] );
magma_queue_destroy( queues[1] );
magma_free( dA );
return *info;
} /* magma_cgeqlf */
extern "C" magma_int_t
magma_cunmql(const char side, const char trans,
magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloatComplex *a, magma_int_t lda,
magmaFloatComplex *tau,
magmaFloatComplex *c, magma_int_t ldc,
magmaFloatComplex *work, magma_int_t lwork,
magma_int_t *info)
{
/* -- MAGMA (version 1.4.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
December 2013
Purpose
=======
CUNMQL overwrites the general complex M-by-N matrix C with
SIDE = 'L' SIDE = 'R'
TRANS = 'N': Q * C C * Q
TRANS = 'C': Q**H * C C * Q**H
where Q is a complex unitary matrix defined as the product of k
elementary reflectors
Q = H(k) . . . H(2) H(1)
as returned by CGEQLF. Q is of order M if SIDE = 'L' and of order N
if SIDE = 'R'.
Arguments
=========
SIDE (input) CHARACTER*1
= 'L': apply Q or Q**H from the Left;
= 'R': apply Q or Q**H from the Right.
TRANS (input) CHARACTER*1
= 'N': No transpose, apply Q;
= 'C': Transpose, apply Q**H.
M (input) INTEGER
The number of rows of the matrix C. M >= 0.
N (input) INTEGER
The number of columns of the matrix C. N >= 0.
K (input) INTEGER
The number of elementary reflectors whose product defines
the matrix Q.
If SIDE = 'L', M >= K >= 0;
if SIDE = 'R', N >= K >= 0.
A (input) COMPLEX array, dimension (LDA,K)
The i-th column must contain the vector which defines the
elementary reflector H(i), for i = 1,2,...,k, as returned by
CGEQLF in the last k columns of its array argument A.
A is modified by the routine but restored on exit.
LDA (input) INTEGER
The leading dimension of the array A.
If SIDE = 'L', LDA >= max(1,M);
if SIDE = 'R', LDA >= max(1,N).
TAU (input) COMPLEX array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by CGEQLF.
C (input/output) COMPLEX array, dimension (LDC,N)
On entry, the M-by-N matrix C.
On exit, C is overwritten by Q*C or Q**H*C or C*Q**H or C*Q.
LDC (input) INTEGER
The leading dimension of the array C. LDC >= max(1,M).
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.
If SIDE = 'L', LWORK >= max(1,N);
if SIDE = 'R', LWORK >= max(1,M).
For optimum performance LWORK >= N*NB if SIDE = 'L', and
LWORK >= M*NB if SIDE = 'R', 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.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
===================================================================== */
char side_[2] = {side, 0};
char trans_[2] = {trans, 0};
magma_int_t i__4, i__;
magmaFloatComplex *T;
magma_int_t i1, i2, i3, ib, nb, mi, ni, nq, nw;
magma_int_t iinfo, ldwork, lwkopt=0;
int lquery, left, notran;
*info = 0;
left = lapackf77_lsame(side_, "L");
notran = lapackf77_lsame(trans_, "N");
lquery = (lwork == -1);
/* NQ is the order of Q and NW is the minimum dimension of WORK */
if (left) {
nq = m;
nw = max(1,n);
} else {
nq = n;
nw = max(1,m);
}
if (! left && ! lapackf77_lsame(side_, "R")) {
*info = -1;
} else if (! notran && ! lapackf77_lsame(trans_, "C")) {
*info = -2;
} else if (m < 0) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (k < 0 || k > nq) {
*info = -5;
} else if (lda < max(1,nq)) {
*info = -7;
} else if (ldc < max(1,m)) {
*info = -10;
}
if (*info == 0) {
if (m == 0 || n == 0) {
lwkopt = 1;
} else {
/* Determine the block size. NB may be at most NBMAX, where
NBMAX is used to define the local array T. */
nb = 64;
lwkopt = nw * nb;
}
work[0] = MAGMA_C_MAKE( lwkopt, 0 );
if (lwork < nw && ! lquery) {
*info = -12;
}
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0) {
return *info;
}
/* Allocate work space on the GPU */
magmaFloatComplex *dwork, *dc;
magma_cmalloc( &dc, (m)*(n) );
magma_cmalloc( &dwork, 2*(m + 64)*64 );
/* Copy matrix C from the CPU to the GPU */
magma_csetmatrix( m, n, c, ldc, dc, m );
/* work space on CPU */
if ( MAGMA_SUCCESS != magma_cmalloc_pinned( &T, 2*nb*nb ) ) {
magma_free( dc );
magma_free( dwork );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
ldwork = nw;
if ( nb >= k ) {
/* Use CPU code */
lapackf77_cunmql(side_, trans_, &m, &n, &k, a, &lda, tau,
c, &ldc, work, &lwork, &iinfo);
}
else {
/* Use hybrid CPU-GPU code */
if ((left && notran) || (! left && ! notran)) {
i1 = 1;
i2 = k;
i3 = nb;
} else {
i1 = (k - 1) / nb * nb + 1;
i2 = 1;
i3 = -nb;
}
// silence "uninitialized" warnings
mi = 0;
ni = 0;
if (left) {
ni = n;
} else {
mi = m;
}
for (i__ = i1; (i3 < 0 ? i__ >= i2 : i__ <= i2); i__ += i3) {
ib = min(nb, k - i__ + 1);
/* Form the triangular factor of the block reflector
H = H(i+ib-1) . . . H(i+1) H(i) */
i__4 = nq - k + i__ + ib - 1;
lapackf77_clarft("Backward", "Columnwise", &i__4, &ib,
&a[(i__-1) * lda], &lda, &tau[i__-1], T, &ib);
/* 1) Put 0s in the lower triangular part of A;
2) copy the panel from A to the GPU, and
3) restore A */
cpanel_to_q('L', ib, &a[i__-1 + (i__-1) * lda], lda, T+ib*ib);
magma_csetmatrix( i__4, ib, &a[(i__-1) * lda], lda, dwork, i__4 );
cq_to_panel('L', ib, &a[i__-1 + (i__-1) * lda], lda, T+ib*ib);
if (left) {
/* H or H' is applied to C(1:m-k+i+ib-1,1:n) */
mi = m - k + i__ + ib - 1;
}
else {
/* H or H' is applied to C(1:m,1:n-k+i+ib-1) */
ni = n - k + i__ + ib - 1;
}
/* Apply H or H'; First copy T to the GPU */
magma_csetmatrix( ib, ib, T, ib, dwork+i__4*ib, ib );
magma_clarfb_gpu(side, trans, MagmaBackward, MagmaColumnwise,
mi, ni, ib,
dwork, i__4, dwork+i__4*ib, ib,
dc, m,
dwork+i__4*ib + ib*ib, ldwork);
}
magma_cgetmatrix( m, n, dc, m, c, ldc );
}
work[0] = MAGMA_C_MAKE( lwkopt, 0 );
magma_free( dc );
magma_free( dwork );
magma_free_pinned( T);
return *info;
} /* magma_cunmql */
extern "C" magma_int_t
magma_cunmqr(const char side, const char trans,
magma_int_t m, magma_int_t n, magma_int_t k,
cuFloatComplex *A, magma_int_t lda,
cuFloatComplex *tau,
cuFloatComplex *C, magma_int_t ldc,
cuFloatComplex *work, magma_int_t lwork,
magma_int_t *info)
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
Purpose
=======
CUNMQR overwrites the general complex M-by-N matrix C with
SIDE = 'L' SIDE = 'R'
TRANS = 'N': Q * C C * Q
TRANS = 'T': Q**H * C C * Q**H
where Q is a complex orthogonal matrix defined as the product of k
elementary reflectors
Q = H(1) H(2) . . . H(k)
as returned by CGEQRF. Q is of order M if SIDE = 'L' and of order N
if SIDE = 'R'.
Arguments
=========
SIDE (input) CHARACTER*1
= 'L': apply Q or Q**H from the Left;
= 'R': apply Q or Q**H from the Right.
TRANS (input) CHARACTER*1
= 'N': No transpose, apply Q;
= 'T': Transpose, apply Q**H.
M (input) INTEGER
The number of rows of the matrix C. M >= 0.
N (input) INTEGER
The number of columns of the matrix C. N >= 0.
K (input) INTEGER
The number of elementary reflectors whose product defines
the matrix Q.
If SIDE = 'L', M >= K >= 0;
if SIDE = 'R', N >= K >= 0.
A (input) COMPLEX array, dimension (LDA,K)
The i-th column must contain the vector which defines the
elementary reflector H(i), for i = 1,2,...,k, as returned by
CGEQRF in the first k columns of its array argument A.
A is modified by the routine but restored on exit.
LDA (input) INTEGER
The leading dimension of the array A.
If SIDE = 'L', LDA >= max(1,M);
if SIDE = 'R', LDA >= max(1,N).
TAU (input) COMPLEX array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by CGEQRF.
C (input/output) COMPLEX array, dimension (LDC,N)
On entry, the M-by-N matrix C.
On exit, C is overwritten by Q*C or Q**H * C or C * Q**H or C*Q.
LDC (input) INTEGER
The leading dimension of the array C. LDC >= max(1,M).
WORK (workspace/output) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(0) returns the optimal LWORK.
LWORK (input) INTEGER
The dimension of the array WORK.
If SIDE = 'L', LWORK >= max(1,N);
if SIDE = 'R', LWORK >= max(1,M).
For optimum performance
LWORK >= N*NB if SIDE = 'L', and
LWORK >= M*NB if SIDE = 'R',
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.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
===================================================================== */
#define A(a_1,a_2) ( A + (a_1) + (a_2)*lda)
#define dC(a_1,a_2) (dC + (a_1) + (a_2)*lddc)
magma_int_t nb = magma_get_cgeqrf_nb( min( m, n ));
cuFloatComplex c_one = MAGMA_C_ONE;
char side_[2] = {side, 0};
char trans_[2] = {trans, 0};
magma_int_t nq_i, lddwork;
magma_int_t i;
cuFloatComplex T[ 2*nb*nb ];
magma_int_t i1, i2, step, ib, ic, jc, mi, ni, nq, nw;
int left, notran, lquery;
magma_int_t iinfo, lwkopt;
*info = 0;
left = lapackf77_lsame(side_, "L");
notran = lapackf77_lsame(trans_, "N");
lquery = (lwork == -1);
/* NQ is the order of Q and NW is the minimum dimension of WORK */
if (left) {
nq = m;
nw = n;
} else {
nq = n;
nw = m;
}
lwkopt = max(1,nw) * nb;
work[0] = MAGMA_C_MAKE( lwkopt, 0 );
if (! left && ! lapackf77_lsame(side_, "R")) {
*info = -1;
} else if (! notran && ! lapackf77_lsame(trans_, MagmaConjTransStr)) {
*info = -2;
} else if (m < 0) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (k < 0 || k > nq) {
*info = -5;
} else if (lda < max(1,nq)) {
*info = -7;
} else if (ldc < max(1,m)) {
*info = -10;
} else if (lwork < max(1,nw) && ! lquery) {
*info = -12;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0 || k == 0) {
work[0] = c_one;
return *info;
}
/* Allocate work space on the GPU */
magma_int_t lddc = m;
cuFloatComplex *dwork, *dC;
magma_cmalloc( &dC, lddc*n );
magma_cmalloc( &dwork, (m + n + nb)*nb );
/* Copy matrix C from the CPU to the GPU */
magma_csetmatrix( m, n, C, ldc, dC, lddc );
if (nb >= k) {
/* Use CPU code */
lapackf77_cunmqr(side_, trans_, &m, &n, &k, A, &lda, &tau[1],
C, &ldc, work, &lwork, &iinfo);
}
else {
/* Use hybrid CPU-GPU code */
if ( (left && (! notran)) || ((! left) && notran) ) {
i1 = 0;
i2 = k;
step = nb;
} else {
i1 = ((k - 1) / nb) * nb;
i2 = 0;
step = -nb;
}
if (left) {
ni = n;
jc = 0;
} else {
mi = m;
ic = 0;
}
for( i=i1; (step<0 ? i>=i2 : i<i2); i += step ) {
ib = min(nb, k - i);
/* Form the triangular factor of the block reflector
H = H(i) H(i+1) . . . H(i+ib-1) */
nq_i = nq - i;
lapackf77_clarft("F", "C", &nq_i, &ib, A(i,i), &lda,
&tau[i], T, &ib);
/* 1) Put 0s in the upper triangular part of A;
2) copy the panel from A to the GPU, and
3) restore A */
cpanel_to_q('U', ib, A(i,i), lda, T+ib*ib);
magma_csetmatrix( nq_i, ib, A(i,i), lda, dwork, nq_i );
cq_to_panel('U', ib, A(i,i), lda, T+ib*ib);
if (left) {
/* H or H' is applied to C(i:m,1:n) */
mi = m - i;
ic = i;
}
else {
/* H or H' is applied to C(1:m,i:n) */
ni = n - i;
jc = i;
}
if (left)
lddwork = ni;
else
lddwork = mi;
/* Apply H or H'; First copy T to the GPU */
magma_csetmatrix( ib, ib, T, ib, dwork+nq_i*ib, ib );
magma_clarfb_gpu( side, trans, MagmaForward, MagmaColumnwise,
mi, ni, ib,
dwork, nq_i, dwork+nq_i*ib, ib,
dC(ic,jc), lddc,
dwork+nq_i*ib + ib*ib, lddwork);
}
magma_cgetmatrix( m, n, dC, lddc, C, ldc );
}
work[0] = MAGMA_C_MAKE( lwkopt, 0 );
magma_free( dC );
magma_free( dwork );
return *info;
} /* magma_cunmqr */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing clarfb_gpu
*/
int main( int argc, char** argv )
{
TESTING_CUDA_INIT();
cuFloatComplex c_zero = MAGMA_C_ZERO;
cuFloatComplex c_one = MAGMA_C_ONE;
cuFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t ione = 1;
printf( "\nUsage: %s -M m -N n -K k\n\n", argv[0] );
magma_int_t m = 500;
magma_int_t n = 300;
magma_int_t k = 32;
for( int i = 1; i < argc; i++ ) {
if (strcmp("-M", argv[i]) == 0 && i+1 < argc) {
m = atoi( argv[++i] );
}
else if (strcmp("-N", argv[i]) == 0 && i+1 < argc) {
n = atoi( argv[++i] );
}
else if (strcmp("-K", argv[i]) == 0 && i+1 < argc) {
k = atoi( argv[++i] );
}
else {
printf( "invalid argument: %s\n", argv[i] );
exit(1);
}
}
if ( k <= 0 || k > m || k > n ) {
printf( "requires 0 < k <= min(m,n)\n" );
exit(1);
}
magma_int_t ldc = m;
magma_int_t ldv = max(m,n);
magma_int_t ldt = k;
magma_int_t ldw = max(m,n);
magma_int_t nv;
ldc = ((ldc+31)/32)*32;
ldv = ((ldv+31)/32)*32;
ldt = ((ldt+31)/32)*32;
ldw = ((ldw+31)/32)*32;
// Allocate memory for matrices
cuFloatComplex *C, *R, *V, *T, *W;
TESTING_MALLOC( C, cuFloatComplex, ldc*n );
TESTING_MALLOC( R, cuFloatComplex, ldc*n );
TESTING_MALLOC( V, cuFloatComplex, ldv*k );
TESTING_MALLOC( T, cuFloatComplex, ldt*k );
TESTING_MALLOC( W, cuFloatComplex, ldw*k );
cuFloatComplex *dC, *dV, *dT, *dW;
TESTING_DEVALLOC( dC, cuFloatComplex, ldc*n );
TESTING_DEVALLOC( dV, cuFloatComplex, ldv*k );
TESTING_DEVALLOC( dT, cuFloatComplex, ldt*k );
TESTING_DEVALLOC( dW, cuFloatComplex, ldw*k );
magma_int_t size;
magma_int_t iseed[4] = { 1, 2, 3, 4 };
float error, work[1];
// test all combinations of input parameters
const char* side[] = { MagmaLeftStr, MagmaRightStr };
const char* trans[] = { MagmaConjTransStr, MagmaNoTransStr };
const char* direct[] = { MagmaForwardStr, MagmaBackwardStr };
const char* storev[] = { MagmaColumnwiseStr, MagmaRowwiseStr };
printf(" M N K storev side direct trans ||R||_F / ||HC||_F\n");
printf("==================================================================================\n");
for( int istor = 0; istor < 2; ++istor ) {
for( int iside = 0; iside < 2; ++iside ) {
for( int idir = 0; idir < 2; ++idir ) {
for( int itran = 0; itran < 2; ++itran ) {
//printf( "# ----------\n" );
//printf( "# %-10s %-10s %-10s %-10s\n", storev[istor], side[iside], direct[idir], trans[itran] );
// C is full
size = ldc*n;
lapackf77_clarnv( &ione, iseed, &size, C );
//printf( "C=" ); magma_cprint( m, n, C, ldc );
// V is ldv x nv. See larfb docs for description.
ldv = (*side[iside] == 'L' ? m : n);
nv = k;
size = ldv*nv;
lapackf77_clarnv( &ione, iseed, &size, V );
if ( *storev[istor] == MagmaColumnwise ) {
if ( *direct[idir] == MagmaForward ) {
lapackf77_claset( MagmaUpperStr, &k, &k, &c_zero, &c_one, V, &ldv );
}
else {
lapackf77_claset( MagmaLowerStr, &k, &k, &c_zero, &c_one, &V[(ldv-k)], &ldv );
}
}
else {
// rowwise, swap V's dimensions
std::swap( ldv, nv );
if ( *direct[idir] == MagmaForward ) {
lapackf77_claset( MagmaLowerStr, &k, &k, &c_zero, &c_one, V, &ldv );
}
else {
lapackf77_claset( MagmaUpperStr, &k, &k, &c_zero, &c_one, &V[(nv-k)*ldv], &ldv );
}
}
//printf( "# ldv %d, nv %d\n", ldv, nv );
//printf( "V=" ); magma_cprint( ldv, nv, V, ldv );
// T is upper triangular for forward, and lower triangular for backward
magma_int_t k1 = k-1;
size = ldt*k;
lapackf77_clarnv( &ione, iseed, &size, T );
if ( *direct[idir] == MagmaForward ) {
lapackf77_claset( MagmaLowerStr, &k1, &k1, &c_zero, &c_zero, &T[1], &ldt );
}
else {
lapackf77_claset( MagmaUpperStr, &k1, &k1, &c_zero, &c_zero, &T[1*ldt], &ldt );
}
//printf( "T=" ); magma_cprint( k, k, T, ldt );
magma_csetmatrix( m, n, C, ldc, dC, ldc );
magma_csetmatrix( ldv, nv, V, ldv, dV, ldv );
magma_csetmatrix( k, k, T, ldt, dT, ldt );
lapackf77_clarfb( side[iside], trans[itran], direct[idir], storev[istor],
&m, &n, &k,
V, &ldv, T, &ldt, C, &ldc, W, &ldw );
//printf( "HC=" ); magma_cprint( m, n, C, ldc );
magma_clarfb_gpu( *side[iside], *trans[itran], *direct[idir], *storev[istor],
m, n, k,
dV, ldv, dT, ldt, dC, ldc, dW, ldw );
magma_cgetmatrix( m, n, dC, ldc, R, ldc );
//printf( "dHC=" ); magma_cprint( m, n, R, ldc );
// compute relative error |HC_magma - HC_lapack| / |HC_lapack|
error = lapackf77_clange( "Fro", &m, &n, C, &ldc, work );
size = ldc*n;
blasf77_caxpy( &size, &c_neg_one, C, &ione, R, &ione );
error = lapackf77_clange( "Fro", &m, &n, R, &ldc, work ) / error;
printf( "%5d %5d %5d %-10s %-10s %-10s %-10s %8.2e\n",
(int) m, (int) n, (int) k,
storev[istor], side[iside], direct[idir], trans[itran], error );
}}}}
// Memory clean up
TESTING_FREE( C );
TESTING_FREE( R );
TESTING_FREE( V );
TESTING_FREE( T );
TESTING_FREE( W );
TESTING_DEVFREE( dC );
TESTING_DEVFREE( dV );
TESTING_DEVFREE( dT );
TESTING_DEVFREE( dW );
// Shutdown
TESTING_CUDA_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 */
extern "C" magma_int_t
magma_cungqr_m(
magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloatComplex *A, magma_int_t lda,
magmaFloatComplex *tau,
magmaFloatComplex *T, magma_int_t nb,
magma_int_t *info)
{
/* -- MAGMA (version 1.4.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
December 2013
Purpose
=======
CUNGQR generates an M-by-N COMPLEX matrix Q with orthonormal columns,
which is defined as the first N columns of a product of K elementary
reflectors of order M
Q = H(1) H(2) . . . H(k)
as returned by CGEQRF.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix Q. M >= 0.
N (input) INTEGER
The number of columns of the matrix Q. M >= N >= 0.
K (input) INTEGER
The number of elementary reflectors whose product defines the
matrix Q. N >= K >= 0.
A (input/output) COMPLEX array A, dimension (LDDA,N).
On entry, the i-th column must contain the vector
which defines the elementary reflector H(i), for
i = 1,2,...,k, as returned by CGEQRF_GPU in the
first k columns of its array argument A.
On exit, the M-by-N matrix Q.
LDA (input) INTEGER
The first dimension of the array A. LDA >= max(1,M).
TAU (input) COMPLEX array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by CGEQRF_GPU.
T (input) COMPLEX array, dimension (NB, min(M,N)).
T contains the T matrices used in blocking the elementary
reflectors H(i), e.g., this can be the 6th argument of
magma_cgeqrf_gpu (except stored on the CPU, not the GPU).
NB (input) INTEGER
This is the block size used in CGEQRF_GPU, and correspondingly
the size of the T matrices, used in the factorization, and
stored in T.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument has an illegal value
===================================================================== */
#define A(i,j) ( A + (i) + (j)*lda )
#define dA(d,i,j) (dA[d] + (i) + (j)*ldda)
#define dT(d,i,j) (dT[d] + (i) + (j)*nb)
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magma_int_t m_kk, n_kk, k_kk, mi;
magma_int_t lwork, ldwork;
magma_int_t i, ib, ki, kk, iinfo;
magmaFloatComplex *work;
*info = 0;
if (m < 0) {
*info = -1;
} else if ((n < 0) || (n > m)) {
*info = -2;
} else if ((k < 0) || (k > n)) {
*info = -3;
} else if (lda < max(1,m)) {
*info = -5;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (n <= 0) {
return *info;
}
magma_int_t di, dn;
int dpanel;
int ngpu = magma_num_gpus();
int doriginal;
magma_getdevice( &doriginal );
// Allocate memory on GPUs for A and workspaces
magma_int_t ldda = ((m + 31) / 32) * 32;
magma_int_t lddwork = ((n + 31) / 32) * 32;
magma_int_t min_lblocks = (n / nb) / ngpu; // min. blocks per gpu
magma_int_t last_dev = (n / nb) % ngpu; // device with last block
magma_int_t nlocal[ MagmaMaxGPUs ] = { 0 };
magmaFloatComplex *dA[ MagmaMaxGPUs ] = { NULL };
magmaFloatComplex *dT[ MagmaMaxGPUs ] = { NULL };
magmaFloatComplex *dV[ MagmaMaxGPUs ] = { NULL };
magmaFloatComplex *dW[ MagmaMaxGPUs ] = { NULL };
magma_queue_t stream[ MagmaMaxGPUs ] = { NULL };
for( int d = 0; d < ngpu; ++d ) {
// example with n = 75, nb = 10, ngpu = 3
// min_lblocks = 2
// last_dev = 1
// gpu 0: 2 blocks, cols: 0- 9, 30-39, 60-69
// gpu 1: 1+ blocks, cols: 10-19, 40-49, 70-74 (partial)
// gpu 2: 1 block , cols: 20-29, 50-59
magma_setdevice( d );
nlocal[d] = min_lblocks*nb;
if ( d < last_dev ) {
nlocal[d] += nb;
}
else if ( d == last_dev ) {
nlocal[d] += (n % nb);
}
ldwork = nlocal[d]*ldda // dA
+ nb*m // dT
+ nb*ldda // dV
+ nb*lddwork; // dW
if ( MAGMA_SUCCESS != magma_cmalloc( &dA[d], ldwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
goto CLEANUP;
}
dT[d] = dA[d] + nlocal[d]*ldda;
dV[d] = dT[d] + nb*m;
dW[d] = dV[d] + nb*ldda;
magma_queue_create( &stream[d] );
}
trace_init( 1, ngpu, 1, stream );
// first kk columns are handled by blocked method.
// ki is start of 2nd-to-last block
if ((nb > 1) && (nb < k)) {
ki = (k - nb - 1) / nb * nb;
kk = min(k, ki + nb);
} else {
ki = 0;
kk = 0;
}
// Allocate CPU work space
// n*nb for cungqr workspace
lwork = n * nb;
magma_cmalloc_cpu( &work, lwork );
if (work == NULL) {
*info = MAGMA_ERR_HOST_ALLOC;
goto CLEANUP;
}
// Use unblocked code for the last or only block.
if (kk < n) {
trace_cpu_start( 0, "ungqr", "ungqr last block" );
m_kk = m - kk;
n_kk = n - kk;
k_kk = k - kk;
dpanel = (kk / nb) % ngpu;
di = ((kk / nb) / ngpu) * nb;
magma_setdevice( dpanel );
lapackf77_cungqr( &m_kk, &n_kk, &k_kk,
A(kk, kk), &lda,
&tau[kk], work, &lwork, &iinfo );
magma_csetmatrix( m_kk, n_kk,
A(kk, kk), lda,
dA(dpanel, kk, di), ldda );
// Set A(1:kk,kk+1:n) to zero.
magmablas_claset( MagmaUpperLower, kk, n - kk, dA(dpanel, 0, di), ldda );
trace_cpu_end( 0 );
}
if (kk > 0) {
// Use blocked code
// send T to all GPUs
for( int d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
trace_gpu_start( d, 0, "set", "set T" );
magma_csetmatrix_async( nb, min(m,n), T, nb, dT[d], nb, stream[d] );
trace_gpu_end( d, 0 );
}
// stream: set Aii (V) --> laset --> laset --> larfb --> [next]
// CPU has no computation
for( i = ki; i >= 0; i -= nb ) {
ib = min(nb, k - i);
mi = m - i;
dpanel = (i / nb) % ngpu;
di = ((i / nb) / ngpu) * nb;
// Send current panel to the GPUs
lapackf77_claset( "Upper", &ib, &ib, &c_zero, &c_one, A(i, i), &lda );
for( int d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
trace_gpu_start( d, 0, "set", "set V" );
magma_csetmatrix_async( mi, ib,
A(i, i), lda,
dV[d], ldda, stream[d] );
trace_gpu_end( d, 0 );
}
// set panel to identity
magma_setdevice( dpanel );
magmablasSetKernelStream( stream[dpanel] );
trace_gpu_start( dpanel, 0, "laset", "laset" );
magmablas_claset( MagmaUpperLower, i, ib, dA(dpanel, 0, di), ldda );
magmablas_claset_identity( mi, ib, dA(dpanel, i, di), ldda );
trace_gpu_end( dpanel, 0 );
if (i < n) {
// Apply H to A(i:m,i:n) from the left
for( int d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magmablasSetKernelStream( stream[d] );
magma_indices_1D_bcyclic( nb, ngpu, d, i, n, &di, &dn );
trace_gpu_start( d, 0, "larfb", "larfb" );
magma_clarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
mi, dn-di, ib,
dV[d], ldda, dT(d,0,i), nb,
dA(d, i, di), ldda, dW[d], lddwork );
trace_gpu_end( d, 0 );
}
}
}
}
// copy result back to CPU
trace_cpu_start( 0, "get", "get A" );
magma_cgetmatrix_1D_col_bcyclic( m, n, dA, ldda, A, lda, ngpu, nb );
trace_cpu_end( 0 );
#ifdef TRACING
char name[80];
snprintf( name, sizeof(name), "cungqr-n%d-ngpu%d.svg", m, ngpu );
trace_finalize( name, "trace.css" );
#endif
CLEANUP:
for( int d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magmablasSetKernelStream( NULL );
magma_free( dA[d] );
dA[d] = NULL;
if ( stream[d] != NULL ) {
magma_queue_destroy( stream[d] );
}
}
magma_free_cpu( work );
magma_setdevice( doriginal );
return *info;
} /* magma_cungqr */
/**
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 */
/**
Purpose
-------
CUNMQL overwrites the general complex M-by-N matrix C with
@verbatim
SIDE = MagmaLeft SIDE = MagmaRight
TRANS = MagmaNoTrans: Q * C C * Q
TRANS = MagmaConjTrans: Q**H * C C * Q**H
@endverbatim
where Q is a complex unitary matrix defined as the product of k
elementary reflectors
Q = H(k) . . . H(2) H(1)
as returned by CGEQLF. Q is of order M if SIDE = MagmaLeft and of order N
if SIDE = MagmaRight.
Arguments
---------
@param[in]
side magma_side_t
- = MagmaLeft: apply Q or Q**H from the Left;
- = MagmaRight: apply Q or Q**H from the Right.
@param[in]
trans magma_trans_t
- = MagmaNoTrans: No transpose, apply Q;
- = MagmaConjTrans: Conjugate transpose, apply Q**H.
@param[in]
m INTEGER
The number of rows of the matrix C. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix C. N >= 0.
@param[in]
k INTEGER
The number of elementary reflectors whose product defines
the matrix Q.
If SIDE = MagmaLeft, M >= K >= 0;
if SIDE = MagmaRight, N >= K >= 0.
@param[in]
dA COMPLEX array, dimension (LDA,K)
The i-th column must contain the vector which defines the
elementary reflector H(i), for i = 1,2,...,k, as returned by
CGEQLF in the last k columns of its array argument A.
The diagonal and the lower part
are destroyed, the reflectors are not modified.
@param[in]
ldda INTEGER
The leading dimension of the array DA.
LDDA >= max(1,M) if SIDE = MagmaLeft;
LDDA >= max(1,N) if SIDE = MagmaRight.
@param[in]
tau COMPLEX array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by CGEQLF.
@param[in,out]
dC COMPLEX array, dimension (LDDC,N)
On entry, the M-by-N matrix C.
On exit, C is overwritten by Q*C or Q**H*C or C*Q**H or C*Q.
@param[in]
lddc INTEGER
The leading dimension of the array C. LDDC >= max(1,M).
@param[in]
wA (workspace) COMPLEX array, dimension
(LDWA,M) if SIDE = MagmaLeft
(LDWA,N) if SIDE = MagmaRight
The vectors which define the elementary reflectors, as
returned by CHETRD_GPU.
@param[in]
ldwa INTEGER
The leading dimension of the array wA.
LDWA >= max(1,M) if SIDE = MagmaLeft; LDWA >= max(1,N) if SIDE = MagmaRight.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_cgeqlf_comp
********************************************************************/
extern "C" magma_int_t
magma_cunmql2_gpu(magma_side_t side, magma_trans_t trans,
magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloatComplex *dA, magma_int_t ldda,
magmaFloatComplex *tau,
magmaFloatComplex *dC, magma_int_t lddc,
magmaFloatComplex *wA, magma_int_t ldwa,
magma_int_t *info)
{
#define dA(i_,j_) (dA + (i_) + (j_)*ldda)
#define dC(i_,j_) (dC + (i_) + (j_)*lddc)
#define wA(i_,j_) (wA + (i_) + (j_)*ldwa)
/* Allocate work space on the GPU */
magmaFloatComplex *dwork;
magma_cmalloc( &dwork, 2*(m + 64)*64 );
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magma_int_t i, i__4;
magmaFloatComplex T[2*4160] /* was [65][64] */;
magma_int_t i1, i2, step, ib, nb, mi, ni, nq, nw;
magma_int_t ldwork;
int left, notran;
wA -= 1 + ldwa;
dC -= 1 + lddc;
--tau;
*info = 0;
left = (side == MagmaLeft);
notran = (trans == MagmaNoTrans);
/* NQ is the order of Q and NW is the minimum dimension of WORK */
if (left) {
nq = m;
nw = max(1,n);
} else {
nq = n;
nw = max(1,m);
}
if (! left && side != MagmaRight) {
*info = -1;
} else if (! notran && trans != MagmaConjTrans) {
*info = -2;
} else if (m < 0) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (k < 0 || k > nq) {
*info = -5;
} else if (ldda < max(1,nq)) {
*info = -7;
} else if (lddc < max(1,m)) {
*info = -10;
} else if (ldwa < max(1,nq)) {
*info = -12;
}
// size of the block
nb = 64;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0) {
return *info;
}
ldwork = nw;
/* Use hybrid CPU-GPU code */
if ((left && notran) || (! left && ! notran)) {
i1 = 1;
i2 = k;
step = nb;
} else {
i1 = (k - 1) / nb * nb + 1;
i2 = 1;
step = -nb;
}
// silence "uninitialized" warnings
mi = 0;
ni = 0;
if (left) {
ni = n;
} else {
mi = m;
}
// set nb-1 sub-diagonals to 0, and diagonal to 1.
// This way we can copy V directly to the GPU,
// already with the lower triangle parts already set to identity.
magmablas_claset_band( MagmaLower, k, k, nb, c_zero, c_one, dA, ldda );
for (i = i1; (step < 0 ? i >= i2 : i <= i2); i += step) {
ib = min(nb, k - i + 1);
/* Form the triangular factor of the block reflector
H = H(i+ib-1) . . . H(i+1) H(i) */
i__4 = nq - k + i + ib - 1;
lapackf77_clarft("Backward", "Columnwise", &i__4, &ib,
wA(1,i), &ldwa, &tau[i], T, &ib);
if (left) {
/* H or H' is applied to C(1:m-k+i+ib-1,1:n) */
mi = m - k + i + ib - 1;
}
else {
/* H or H' is applied to C(1:m,1:n-k+i+ib-1) */
ni = n - k + i + ib - 1;
}
/* Apply H or H'; First copy T to the GPU */
magma_csetmatrix( ib, ib, T, ib, dwork+i__4*ib, ib );
magma_clarfb_gpu(side, trans, MagmaBackward, MagmaColumnwise,
mi, ni, ib,
dA(0,i-1), ldda, dwork+i__4*ib, ib, // dA using 0-based indices here
dC(1,1), lddc,
dwork+i__4*ib + ib*ib, ldwork);
}
magma_free( dwork );
return *info;
} /* magma_cunmql */
/***************************************************************************//**
Purpose
-------
CUNGQR generates an M-by-N COMPLEX matrix Q with orthonormal columns,
which is defined as the first N columns of a product of K elementary
reflectors of order M
Q = H(1) H(2) . . . H(k)
as returned by CGEQRF.
This version recomputes the T matrices on the CPU and sends them to the GPU.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix Q. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix Q. M >= N >= 0.
@param[in]
k INTEGER
The number of elementary reflectors whose product defines the
matrix Q. N >= K >= 0.
@param[in,out]
A COMPLEX array A, dimension (LDDA,N).
On entry, the i-th column must contain the vector
which defines the elementary reflector H(i), for
i = 1,2,...,k, as returned by CGEQRF_GPU in the
first k columns of its array argument A.
On exit, the M-by-N matrix Q.
@param[in]
lda INTEGER
The first dimension of the array A. LDA >= max(1,M).
@param[in]
tau COMPLEX array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by CGEQRF_GPU.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument has an illegal value
@ingroup magma_ungqr
*******************************************************************************/
extern "C" magma_int_t
magma_cungqr2(
magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloatComplex *A, magma_int_t lda,
magmaFloatComplex *tau,
magma_int_t *info)
{
#define A(i,j) ( A + (i) + (j)*lda )
#define dA(i,j) (dA + (i) + (j)*ldda)
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magma_int_t nb = magma_get_cgeqrf_nb( m, n );
magma_int_t m_kk, n_kk, k_kk, mi;
magma_int_t lwork, ldda;
magma_int_t i, ib, ki, kk; //, iinfo;
magma_int_t lddwork;
magmaFloatComplex *dA, *dV, *dW, *dT, *T;
magmaFloatComplex *work;
*info = 0;
if (m < 0) {
*info = -1;
} else if ((n < 0) || (n > m)) {
*info = -2;
} else if ((k < 0) || (k > n)) {
*info = -3;
} else if (lda < max(1,m)) {
*info = -5;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (n <= 0) {
return *info;
}
// first kk columns are handled by blocked method.
// ki is start of 2nd-to-last block
if ((nb > 1) && (nb < k)) {
ki = (k - nb - 1) / nb * nb;
kk = min(k, ki + nb);
} else {
ki = 0;
kk = 0;
}
// Allocate GPU work space
// ldda*n for matrix dA
// ldda*nb for dV
// lddwork*nb for dW larfb workspace
ldda = magma_roundup( m, 32 );
lddwork = magma_roundup( n, 32 );
if (MAGMA_SUCCESS != magma_cmalloc( &dA, ldda*n + ldda*nb + lddwork*nb + nb*nb)) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dV = dA + ldda*n;
dW = dA + ldda*n + ldda*nb;
dT = dA + ldda*n + ldda*nb + lddwork*nb;
// Allocate CPU work space
lwork = (n+m+nb) * nb;
magma_cmalloc_cpu( &work, lwork );
T = work;
if (work == NULL) {
magma_free( dA );
magma_free_cpu( work );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magmaFloatComplex *V = work + (n+nb)*nb;
magma_queue_t queue;
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queue );
// Use unblocked code for the last or only block.
if (kk < n) {
m_kk = m - kk;
n_kk = n - kk;
k_kk = k - kk;
/*
lapackf77_cungqr( &m_kk, &n_kk, &k_kk,
A(kk, kk), &lda,
&tau[kk], work, &lwork, &iinfo );
*/
lapackf77_clacpy( MagmaFullStr, &m_kk, &k_kk, A(kk,kk), &lda, V, &m_kk);
lapackf77_claset( MagmaFullStr, &m_kk, &n_kk, &c_zero, &c_one, A(kk, kk), &lda );
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&m_kk, &k_kk,
V, &m_kk, &tau[kk], work, &k_kk);
lapackf77_clarfb( MagmaLeftStr, MagmaNoTransStr, MagmaForwardStr, MagmaColumnwiseStr,
&m_kk, &n_kk, &k_kk,
V, &m_kk, work, &k_kk, A(kk, kk), &lda, work+k_kk*k_kk, &n_kk );
if (kk > 0) {
magma_csetmatrix( m_kk, n_kk,
A(kk, kk), lda,
dA(kk, kk), ldda, queue );
// Set A(1:kk,kk+1:n) to zero.
magmablas_claset( MagmaFull, kk, n - kk, c_zero, c_zero, dA(0, kk), ldda, queue );
}
}
if (kk > 0) {
// Use blocked code
// queue: set Aii (V) --> laset --> laset --> larfb --> [next]
// CPU has no computation
for (i = ki; i >= 0; i -= nb) {
ib = min(nb, k - i);
// Send current panel to the GPU
mi = m - i;
lapackf77_claset( "Upper", &ib, &ib, &c_zero, &c_one, A(i, i), &lda );
magma_csetmatrix_async( mi, ib,
A(i, i), lda,
dV, ldda, queue );
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&mi, &ib,
A(i,i), &lda, &tau[i], T, &nb);
magma_csetmatrix_async( ib, ib,
T, nb,
dT, nb, queue );
// set panel to identity
magmablas_claset( MagmaFull, i, ib, c_zero, c_zero, dA(0, i), ldda, queue );
magmablas_claset( MagmaFull, mi, ib, c_zero, c_one, dA(i, i), ldda, queue );
magma_queue_sync( queue );
if (i < n) {
// Apply H to A(i:m,i:n) from the left
magma_clarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
mi, n-i, ib,
dV, ldda, dT, nb,
dA(i, i), ldda, dW, lddwork, queue );
}
}
// copy result back to CPU
magma_cgetmatrix( m, n,
dA(0, 0), ldda, A(0, 0), lda, queue );
}
magma_queue_destroy( queue );
magma_free( dA );
magma_free_cpu( work );
return *info;
} /* magma_cungqr */
/**
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_cunmqr_gpu_2stages(const char side, const char trans,
magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloatComplex *da, magma_int_t ldda,
magmaFloatComplex *dc, magma_int_t lddc,
magmaFloatComplex *dT, magma_int_t nb,
magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CUNMQR_GPU overwrites the general complex M-by-N matrix C with
SIDE = 'L' SIDE = 'R'
TRANS = 'N': Q * C C * Q
TRANS = 'T': Q**H * C C * Q**H
where Q is a complex orthogonal matrix defined as the product of k
elementary reflectors
Q = H(1) H(2) . . . H(k)
as returned by CGEQRF. Q is of order M if SIDE = 'L' and of order N
if SIDE = 'R'.
Arguments
=========
SIDE (input) CHARACTER*1
= 'L': apply Q or Q**H from the Left;
= 'R': apply Q or Q**H from the Right.
TRANS (input) CHARACTER*1
= 'N': No transpose, apply Q;
= 'T': Transpose, apply Q**H.
M (input) INTEGER
The number of rows of the matrix C. M >= 0.
N (input) INTEGER
The number of columns of the matrix C. N >= 0.
K (input) INTEGER
The number of elementary reflectors whose product defines
the matrix Q.
If SIDE = 'L', M >= K >= 0;
if SIDE = 'R', N >= K >= 0.
DA (input) COMPLEX array on the GPU, dimension (LDDA,K)
The i-th column must contain the vector which defines the
elementary reflector H(i), for i = 1,2,...,k, as returned by
CGEQRF in the first k columns of its array argument DA.
DA is modified by the routine but restored on exit.
LDDA (input) INTEGER
The leading dimension of the array DA.
If SIDE = 'L', LDDA >= max(1,M);
if SIDE = 'R', LDDA >= max(1,N).
DC (input/output) COMPLEX array on the GPU, dimension (LDDC,N)
On entry, the M-by-N matrix C.
On exit, C is overwritten by Q*C or Q**H * C or C * Q**H or C*Q.
LDDC (input) INTEGER
The leading dimension of the array DC. LDDC >= max(1,M).
DT (input) COMPLEX array on the GPU that is the output
(the 9th argument) of magma_cgeqrf_gpu.
NB (input) INTEGER
This is the blocking size that was used in pre-computing DT, e.g.,
the blocking size used in magma_cgeqrf_gpu.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
===================================================================== */
char side_[2] = {side, 0};
char trans_[2] = {trans, 0};
magmaFloatComplex *dwork;
magma_int_t i1, i2, i3, ib, ic, jc, mi, ni, nq, nw, ret;
int left, notran;
//magma_int_t lwkopt;
*info = 0;
left = lapackf77_lsame(side_, "L");
notran = lapackf77_lsame(trans_, "N");
/* NQ is the order of Q and NW is the minimum dimension of WORK */
if (left) {
nq = m;
nw = n;
} else {
nq = n;
nw = m;
}
if ( (!left) && (!lapackf77_lsame(side_, "R")) ) {
*info = -1;
} else if ( (!notran) && (!lapackf77_lsame(trans_, MagmaConjTransStr)) ) {
*info = -2;
} else if (m < 0) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (k < 0 || k > nq) {
*info = -5;
} else if (ldda < max(1,nq)) {
*info = -7;
} else if (lddc < max(1,m)) {
*info = -10;
}
if(MAGMA_SUCCESS != magma_cmalloc( &dwork, n*nb )) {
printf ("!!!! cungqr_2stage magma_alloc failed for: dwork\n" );
exit(-1);
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0 || k == 0) {
return *info;
}
if ( (left && (! notran)) || ( (!left) && notran ) ) {
i1 = 0;
i2 = k;
i3 = nb;
} else {
i1 = (k - 1) / nb * nb;
i2 = 0;
i3 = -nb;
}
// silence "uninitialized" warnings
mi = 0;
ni = 0;
if (left) {
ni = n;
jc = 0;
} else {
mi = m;
ic = 0;
}
for (magma_int_t i=i1; (i3<0 ? i>=i2 : i<i2); i+=i3)
{
ib = min(nb, k - i);
if (left){
mi = m - i;
ic = i;
}
else {
ni = n - i;
jc = i;
}
ret = magma_clarfb_gpu( MagmaLeft, trans, MagmaForward, MagmaColumnwise,
mi, ni, ib, da+i+i*ldda, ldda, dT+i*nb, nb,
dc+ic+jc*lddc, lddc, dwork, nw);
if ( ret != MAGMA_SUCCESS ){
magma_free(dwork);
return ret;
}
}
return MAGMA_SUCCESS;
} /* End of MAGMA_CUNMQR_GPU_2stages */
extern "C" magma_int_t
magma_cungqr(
magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloatComplex *a, magma_int_t lda,
magmaFloatComplex *tau, magmaFloatComplex_ptr dT, size_t dT_offset,
magma_int_t nb,
magma_queue_t queue,
magma_int_t *info )
{
/* -- clMAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date November 2014
Purpose
=======
CUNGQR generates an M-by-N COMPLEX matrix Q with orthonormal columns,
which is defined as the first N columns of a product of K elementary
reflectors of order M
Q = H(1) H(2) . . . H(k)
as returned by CGEQRF.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix Q. M >= 0.
N (input) INTEGER
The number of columns of the matrix Q. M >= N >= 0.
K (input) INTEGER
The number of elementary reflectors whose product defines the
matrix Q. N >= K >= 0.
A (input/output) COMPLEX array A, dimension (LDDA,N).
On entry, the i-th column must contain the vector
which defines the elementary reflector H(i), for
i = 1,2,...,k, as returned by CGEQRF_GPU in the
first k columns of its array argument A.
On exit, the M-by-N matrix Q.
LDA (input) INTEGER
The first dimension of the array A. LDA >= max(1,M).
TAU (input) COMPLEX array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by CGEQRF_GPU.
DT (input) COMPLEX array on the GPU device.
DT contains the T matrices used in blocking the elementary
reflectors H(i), e.g., this can be the 6th argument of
magma_cgeqrf_gpu.
NB (input) INTEGER
This is the block size used in CGEQRF_GPU, and correspondingly
the size of the T matrices, used in the factorization, and
stored in DT.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument has an illegal value
===================================================================== */
#define a_ref(i,j) ( a + (j)*lda + (i))
#define da_ref(i,j) da, (da_offset + (j)*ldda + (i))
#define t_ref(a_1) dT, (dT_offset + (a_1)*nb)
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magma_int_t i__1, i__2, i__3;
magma_int_t lwork, ldda;
magma_int_t i, ib, ki, kk, iinfo;
magma_int_t lddwork = min(m, n);
magmaFloatComplex *work;
magmaFloatComplex_ptr da, dwork;
size_t da_offset, dwork_offset;
magma_event_t event = NULL;
*info = 0;
if (m < 0) {
*info = -1;
} else if ((n < 0) || (n > m)) {
*info = -2;
} else if ((k < 0) || (k > n)) {
*info = -3;
} else if (lda < max(1,m)) {
*info = -5;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (n <= 0)
return *info;
/* Allocate GPU work space */
ldda = ((m+31)/32)*32;
lddwork = ((lddwork+31)/32)*32;
if (MAGMA_SUCCESS != magma_cmalloc( &da, ((n)*ldda + nb*lddwork ) )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
da_offset = 0;
dwork = da;
dwork_offset = da_offset + (n)*ldda;
/* Allocate CPU work space */
lwork = n * nb;
magma_cmalloc_cpu( &work, lwork );
if( work == NULL ) {
magma_free( da );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
if ( (nb > 1) && (nb < k) )
{
/* Use blocked code after the last block.
The first kk columns are handled by the block method. */
ki = (k - nb - 1) / nb * nb;
kk = min(k, ki + nb);
/* Set A(1:kk,kk+1:n) to zero. */
magmablas_claset(MagmaFull, kk, n-kk, c_zero, c_zero, da_ref(0,kk), ldda, queue);
}
else
kk = 0;
/* Use unblocked code for the last or only block. */
if (kk < n)
{
i__1 = m - kk;
i__2 = n - kk;
i__3 = k - kk;
lapackf77_cungqr(&i__1, &i__2, &i__3,
a_ref(kk, kk), &lda,
&tau[kk], work, &lwork, &iinfo);
magma_csetmatrix(i__1, i__2, a_ref(kk, kk), lda, da_ref(kk, kk), ldda, queue);
}
if (kk > 0)
{
/* Use blocked code */
for (i = ki; i >= 0; i-=nb)
{
ib = min(nb, k - i);
/* Send the current panel to the GPU */
i__2 = m - i;
cpanel_to_q(MagmaUpper, ib, a_ref(i,i), lda, work);
magma_csetmatrix(i__2, ib, a_ref(i, i), lda, da_ref(i, i), ldda, queue);
if (i + ib < n)
{
/* Apply H to A(i:m,i+ib:n) from the left */
i__3 = n - i - ib;
magma_clarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
i__2, i__3, ib,
da_ref(i, i ), ldda, t_ref(i), nb,
da_ref(i, i+ib), ldda, dwork, dwork_offset, lddwork, queue);
}
/* Apply H to rows i:m of current block on the CPU */
lapackf77_cungqr(&i__2, &ib, &ib,
a_ref(i, i), &lda,
&tau[i], work, &lwork, &iinfo);
magma_csetmatrix_async( i__2, ib,
a_ref(i,i), lda,
da_ref(i,i), ldda, queue, &event );
/* Set rows 1:i-1 of current block to zero */
i__2 = i + ib;
magmablas_claset(MagmaFull, i, i__2 - i, c_zero, c_zero, da_ref(0,i), ldda, queue);
}
}
magma_cgetmatrix(m, n, da_ref(0, 0), ldda, a_ref(0, 0), lda, queue);
//cudaStreamDestroy(stream);
magma_free( da );
magma_free_cpu(work);
return *info;
} /* magma_cungqr */
extern "C" void magma_cbulge_applyQ(
magma_int_t WANTZ, magma_side_t SIDE, magma_int_t NE, magma_int_t N, magma_int_t NB,
magma_int_t Vblksiz, magmaFloatComplex *E, magma_int_t LDE,
magmaFloatComplex *V, magmaFloatComplex *TAU, magmaFloatComplex *T,
magma_int_t *INFO, magmaFloatComplex *dV, magmaFloatComplex *dT,
magmaFloatComplex *dE, magma_int_t copytype )
{
//%===========================
//% local variables
//%===========================
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magma_int_t LDT, LDV, blklen, firstcolj;
magma_int_t bg, nbGblk, rownbm, k, m, n;
magma_int_t st, ed, fst, vlen, vnb, colj, len;
magma_int_t blkid, vpos, taupos, tpos;
//magmaFloatComplex *WORK;
magma_int_t LWORK;
magma_int_t cur_blksiz, avai_blksiz, ncolinvolvd;
magma_int_t nbgr, colst, coled, versionL, versionR;
magma_int_t blkcnt=-1;
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
*INFO=0;
versionL = 113;
versionR = 92;
LDT = Vblksiz;
LDV = NB+Vblksiz-1;
blklen = LDV*Vblksiz;
nbGblk = plasma_ceildiv((N-1), Vblksiz);
//magma_cmalloc_cpu( &WORK, LWORK );
/* find the size of the matrix T V*/
findVTsiz(N, NB, Vblksiz, &blkcnt, &LDV);
/* Copy E & V & T to the GPU in dE and dV and dT
* depending on copytype:
* 1: mean copy only V
* 2: mean copy V and T
* 3: mean copy V, T and E
* */
if (copytype > 0) magma_csetmatrix( LDV, blkcnt*Vblksiz, V, LDV, dV, LDV );
if (copytype > 1) magma_csetmatrix( LDT, blkcnt*Vblksiz, T, LDT, dT, LDT );
if (copytype > 2) magma_csetmatrix( N, NE, E, N, dE, N );
magmaFloatComplex *dwork;
magma_int_t ldwork;
ldwork = NE;
LWORK = 2*N*max(Vblksiz, 64);
if (MAGMA_SUCCESS != magma_cmalloc( &dwork, LWORK )) {
printf ("!!!! magma_cbulge_applyQ magma_alloc failed for: dwork\n" );
exit(-1);
}
/* SIDE LEFT meaning apply E = Q*E = (q_1*q_2*.....*q_n) * E ==> so traverse Vs in reverse order (forward) from q_n to q_1
* Also E is splitten by row meaning each apply consist in a block of row (horizontal block) */
/* SIDE RIGHT meaning apply E = E*Q = E * (q_1*q_2*.....*q_n) ==> so tarverse Vs in normal order (forward) from q_1 to q_n
* Also E is splitten by col meaning each apply consist in a block of col (vertical block) */
/* WANTZ = 1 meaning E is IDENTITY so form Q using optimized update.
* So we use the reverse order from small q to large one,
* so from q_n to q_1 so Left update to Identity.
* Use versionL 113 because in 114 we need to update the whole matrix and not in icreasing order.
* WANTZ = 2 meaning E is a full matrix and need to be updated from Left or Right so use normal update
* */
if (WANTZ == 1) {
versionL=113;
SIDE = MagmaLeft;
//set the matrix to Identity here to avoid copying it from the CPU
magmablas_claset( MagmaFull, N, N, c_zero, c_one, dE, N );
}
printf(" APPLY Q_v115 GPU with N %d NB %d Vblksiz %d SIDE %c versionL %d versionR %d WANTZ %d \n",
(int) N, (int) NB, (int) Vblksiz, SIDE, (int) versionL, (int) versionR, (int) WANTZ);
#if defined(USESTREAM)
magma_int_t N2=N/2;
magma_int_t N1=N-N2;
printf("using stream\n");
magma_queue_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
#endif
if (SIDE == MagmaLeft) {
if (versionL == 113) {
for (bg = nbGblk; bg > 0; bg--) {
firstcolj = (bg-1)*Vblksiz + 1;
if (bg == nbGblk)
rownbm = plasma_ceildiv((N-(firstcolj)), NB); // last blk has size=1 used for complex to handle A(N,N-1)
else
rownbm = plasma_ceildiv((N-(firstcolj+1)), NB);
for (m = rownbm; m > 0; m--) {
vlen = 0;
vnb = 0;
colj = (bg-1)*Vblksiz; // for k=0; I compute the fst and then can remove it from the loop
fst = (rownbm -m)*NB+colj +1;
for (k=0; k < Vblksiz; k++) {
colj = (bg-1)*Vblksiz + k;
st = (rownbm -m)*NB+colj +1;
ed = min(st+NB-1, N-1);
if (st > ed) break;
if ((st == ed) && (colj != N-2)) break;
vlen=ed-fst+1;
vnb=k+1;
}
colst = (bg-1)*Vblksiz;
findVTpos(N, NB, Vblksiz, colst, fst, &vpos, &taupos, &tpos, &blkid);
printf("voici bg %d m %d vlen %d vnb %d fcolj %d vpos %d taupos %d \n", (int) bg, (int) m, (int) vlen, (int) vnb, (int) colst+1, (int) vpos+1, (int) taupos+1);
if ((vlen > 0) && (vnb > 0)) {
if (WANTZ == 1) {
len = N-colst;
magma_clarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, vlen, len, vnb, dV(vpos), LDV, dT(tpos), LDT, dE(fst,colst), LDE, dwork, len);
} else {
magma_clarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, vlen, NE, vnb, dV(vpos), LDV, dT(tpos), LDT, dE(fst,0), LDE, dwork, NE);
}
}
}
}
} else if (versionL == 114) {
rownbm = plasma_ceildiv((N-1), NB);
for (m = rownbm; m > 0; m--) {
ncolinvolvd = min(N-1, m*NB);
avai_blksiz=min(Vblksiz, ncolinvolvd);
nbgr = plasma_ceildiv(ncolinvolvd, avai_blksiz);
for (n = nbgr; n > 0; n--) {
vlen = 0;
vnb = 0;
cur_blksiz = min(ncolinvolvd-(n-1)*avai_blksiz, avai_blksiz);
colst = (n-1)*avai_blksiz;
coled = colst + cur_blksiz -1;
fst = (rownbm -m)*NB+colst +1;
for (colj=colst; colj <= coled; colj++) {
st = (rownbm -m)*NB+colj +1;
ed = min(st+NB-1, N-1);
if (st > ed) break;
if ((st == ed) && (colj != N-2)) break;
vlen=ed-fst+1;
vnb=vnb+1;
}
findVTpos(N, NB, Vblksiz, colst, fst, &vpos, &taupos, &tpos, &blkid);
//printf("voici bg %d m %d vlen %d vnb %d fcolj %d vpos %d taupos %d \n", bg, m, vlen, vnb, colst+1, vpos+1, taupos+1);
if ((vlen > 0) && (vnb > 0)) {
#if defined(USESTREAM)
magmablasSetKernelStream(stream[0]);
magma_clarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, vlen, N1, vnb, dV(vpos), LDV, dT(tpos), LDT, dE(fst,0), LDE, dwork, N1);
magmablasSetKernelStream(stream[1]);
magma_clarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, vlen, N2, vnb, dV(vpos), LDV, dT(tpos), LDT, dE(fst,N1), LDE, &dwork[N1*Vblksiz], N2);
#else
magma_clarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, vlen, NE, vnb, dV(vpos), LDV, dT(tpos), LDT, dE(fst,0), LDE, dwork, NE);
#endif
}
}
}
}
} else if (SIDE == MagmaRight) {
if (versionR == 91) {
for (bg =1; bg <= nbGblk; bg++) {
firstcolj = (bg-1)*Vblksiz + 1;
rownbm = plasma_ceildiv((N-(firstcolj+1)), NB);
if (bg == nbGblk) rownbm = plasma_ceildiv((N-(firstcolj)), NB); // last blk has size=1 used for complex to handle A(N,N-1)
for (m = 1; m <= rownbm; m++) {
vlen = 0;
vnb = 0;
// for k=0; I compute the fst and then can remove it from the loop
colj = (bg-1)*Vblksiz;
fst = (rownbm -m)*NB+colj +1;
for (k=0; k < Vblksiz; k++) {
colj = (bg-1)*Vblksiz + k;
st = (rownbm -m)*NB+colj +1;
ed = min(st+NB-1, N-1);
if (st > ed) break;
if ((st == ed) && (colj != N-2)) break;
vlen=ed-fst+1;
vnb=k+1;
}
colj = (bg-1)*Vblksiz;
findVTpos(N, NB, Vblksiz, colj, fst, &vpos, &taupos, &tpos, &blkid);
//printf("voici bg %d m %d vlen %d vnb %d fcolj %d vpos %d taupos %d \n", bg, m, vlen, vnb, colj, vpos, taupos);
if ((vlen > 0) && (vnb > 0)) {
#if defined(USESTREAM)
magmablasSetKernelStream(stream[0]);
magma_clarfb_gpu( MagmaRight, MagmaNoTrans, MagmaForward, MagmaColumnwise, N1, vlen, vnb, dV(vpos), LDV, dT(tpos), LDT, dE(0, fst), LDE, dwork, N1);
magmablasSetKernelStream(stream[1]);
magma_clarfb_gpu( MagmaRight, MagmaNoTrans, MagmaForward, MagmaColumnwise, N2, vlen, vnb, dV(vpos), LDV, dT(tpos), LDT, dE(N1, fst), LDE, &dwork[N1*Vblksiz], N2);
#else
magma_clarfb_gpu( MagmaRight, MagmaNoTrans, MagmaForward, MagmaColumnwise, NE, vlen, vnb, dV(vpos), LDV, dT(tpos), LDT, dE(0, fst), LDE, dwork, NE);
#endif
}
}
}
} else if (versionR == 92) {
rownbm = plasma_ceildiv((N-1), NB);
for (m = 1; m <= rownbm; m++) {
ncolinvolvd = min(N-1, m*NB);
avai_blksiz=min(Vblksiz, ncolinvolvd);
nbgr = plasma_ceildiv(ncolinvolvd, avai_blksiz);
for (n = 1; n <= nbgr; n++) {
vlen = 0;
vnb = 0;
cur_blksiz = min(ncolinvolvd-(n-1)*avai_blksiz, avai_blksiz);
colst = (n-1)*avai_blksiz;
coled = colst + cur_blksiz -1;
fst = (rownbm -m)*NB+colst +1;
for (colj=colst; colj <= coled; colj++) {
st = (rownbm -m)*NB+colj +1;
ed = min(st+NB-1, N-1);
if (st > ed) break;
if ((st == ed) && (colj != N-2)) break;
vlen=ed-fst+1;
vnb=vnb+1;
}
findVTpos(N, NB, Vblksiz, colst, fst, &vpos, &taupos, &tpos, &blkid);
if ((vlen > 0) && (vnb > 0)) {
#if defined(USESTREAM)
magmablasSetKernelStream(stream[0]);
magma_clarfb_gpu( MagmaRight, MagmaNoTrans, MagmaForward, MagmaColumnwise, N1, vlen, vnb, dV(vpos), LDV, dT(tpos), LDT, dE(0, fst), LDE, dwork, N1);
magmablasSetKernelStream(stream[1]);
magma_clarfb_gpu( MagmaRight, MagmaNoTrans, MagmaForward, MagmaColumnwise, N2, vlen, vnb, dV(vpos), LDV, dT(tpos), LDT, dE(N1, fst), LDE, &dwork[N1*Vblksiz], N2);
#else
magma_clarfb_gpu( MagmaRight, MagmaNoTrans, MagmaForward, MagmaColumnwise, NE, vlen, vnb, dV(vpos), LDV, dT(tpos), LDT, dE(0, fst), LDE, dwork, NE);
#endif
}
}
}
}
} else {
printf("ERROR SIDE %d\n", SIDE);
}
#if defined(USESTREAM)
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
#endif
magmablasSetKernelStream( orig_stream );
}
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
-------
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 */
extern "C" magma_int_t
magma_cungqr(magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloatComplex *A, magma_int_t lda,
magmaFloatComplex *tau,
magmaFloatComplex *dT, magma_int_t nb,
magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CUNGQR generates an M-by-N COMPLEX matrix Q with orthonormal columns,
which is defined as the first N columns of a product of K elementary
reflectors of order M
Q = H(1) H(2) . . . H(k)
as returned by CGEQRF.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix Q. M >= 0.
N (input) INTEGER
The number of columns of the matrix Q. M >= N >= 0.
K (input) INTEGER
The number of elementary reflectors whose product defines the
matrix Q. N >= K >= 0.
A (input/output) COMPLEX array A, dimension (LDDA,N).
On entry, the i-th column must contain the vector
which defines the elementary reflector H(i), for
i = 1,2,...,k, as returned by CGEQRF_GPU in the
first k columns of its array argument A.
On exit, the M-by-N matrix Q.
LDA (input) INTEGER
The first dimension of the array A. LDA >= max(1,M).
TAU (input) COMPLEX array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by CGEQRF_GPU.
DT (input) COMPLEX array on the GPU device.
DT contains the T matrices used in blocking the elementary
reflectors H(i), e.g., this can be the 6th argument of
magma_cgeqrf_gpu.
NB (input) INTEGER
This is the block size used in CGEQRF_GPU, and correspondingly
the size of the T matrices, used in the factorization, and
stored in DT.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument has an illegal value
===================================================================== */
#define A(i,j) ( A + (i) + (j)*lda )
#define dA(i,j) (dA + (i) + (j)*ldda)
#define dT(j) (dT + (j)*nb)
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magma_int_t m_kk, n_kk, k_kk, mi;
magma_int_t lwork, ldda;
magma_int_t i, ib, ki, kk; //, iinfo;
magma_int_t lddwork;
magmaFloatComplex *dA, *dV, *dW;
magmaFloatComplex *work;
*info = 0;
if (m < 0) {
*info = -1;
} else if ((n < 0) || (n > m)) {
*info = -2;
} else if ((k < 0) || (k > n)) {
*info = -3;
} else if (lda < max(1,m)) {
*info = -5;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (n <= 0) {
return *info;
}
// first kk columns are handled by blocked method.
// ki is start of 2nd-to-last block
if ((nb > 1) && (nb < k)) {
ki = (k - nb - 1) / nb * nb;
kk = min(k, ki + nb);
} else {
ki = 0;
kk = 0;
}
// Allocate GPU work space
// ldda*n for matrix dA
// ldda*nb for dV
// lddwork*nb for dW larfb workspace
ldda = ((m + 31) / 32) * 32;
lddwork = ((n + 31) / 32) * 32;
if (MAGMA_SUCCESS != magma_cmalloc( &dA, ldda*n + ldda*nb + lddwork*nb )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dV = dA + ldda*n;
dW = dA + ldda*n + ldda*nb;
// Allocate CPU work space
lwork = (n+m+nb) * nb;
magma_cmalloc_cpu( &work, lwork );
if (work == NULL) {
magma_free( dA );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magmaFloatComplex *V = work + (n+nb)*nb;
magma_queue_t stream;
magma_queue_create( &stream );
// Use unblocked code for the last or only block.
if (kk < n) {
m_kk = m - kk;
n_kk = n - kk;
k_kk = k - kk;
/*
// Replacing this with the following 4 routines works but cungqr is slow for
// k smaller than the cungqr's blocking size (new version can be up to 60x faster)
lapackf77_cungqr( &m_kk, &n_kk, &k_kk,
A(kk, kk), &lda,
&tau[kk], work, &lwork, &iinfo );
*/
lapackf77_clacpy( MagmaUpperLowerStr, &m_kk, &k_kk, A(kk,kk), &lda, V, &m_kk);
lapackf77_claset( MagmaUpperLowerStr, &m_kk, &n_kk, &c_zero, &c_one, A(kk, kk), &lda );
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&m_kk, &k_kk,
V, &m_kk, &tau[kk], work, &k_kk);
lapackf77_clarfb( MagmaLeftStr, MagmaNoTransStr, MagmaForwardStr, MagmaColumnwiseStr,
&m_kk, &n_kk, &k_kk,
V, &m_kk, work, &k_kk, A(kk, kk), &lda, work+k_kk*k_kk, &n_kk );
if (kk > 0) {
magma_csetmatrix( m_kk, n_kk,
A(kk, kk), lda,
dA(kk, kk), ldda );
// Set A(1:kk,kk+1:n) to zero.
magmablas_claset( MagmaUpperLower, kk, n - kk, dA(0, kk), ldda );
}
}
if (kk > 0) {
// Use blocked code
// stream: set Aii (V) --> laset --> laset --> larfb --> [next]
// CPU has no computation
magmablasSetKernelStream( stream );
for (i = ki; i >= 0; i -= nb) {
ib = min(nb, k - i);
// Send current panel to the GPU
mi = m - i;
lapackf77_claset( "Upper", &ib, &ib, &c_zero, &c_one, A(i, i), &lda );
magma_csetmatrix_async( mi, ib,
A(i, i), lda,
dV, ldda, stream );
// set panel to identity
magmablas_claset( MagmaUpperLower, i, ib, dA(0, i), ldda );
magmablas_claset_identity( mi, ib, dA(i, i), ldda );
if (i < n) {
// Apply H to A(i:m,i:n) from the left
magma_clarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
mi, n-i, ib,
dV, ldda, dT(i), nb,
dA(i, i), ldda, dW, lddwork );
}
}
// copy result back to CPU
magma_cgetmatrix( m, n,
dA(0, 0), ldda, A(0, 0), lda);
}
magmablasSetKernelStream( NULL );
magma_queue_destroy( stream );
magma_free( dA );
magma_free_cpu( work );
return *info;
} /* magma_cungqr */
/**
Purpose
-------
CUNMQR_GPU overwrites the general complex M-by-N matrix C with
@verbatim
SIDE = MagmaLeft SIDE = MagmaRight
TRANS = MagmaNoTrans: Q * C C * Q
TRANS = Magma_ConjTrans: Q**H * C C * Q**H
@endverbatim
where Q is a complex unitary matrix defined as the product of k
elementary reflectors
Q = H(1) H(2) . . . H(k)
as returned by CGEQRF. Q is of order M if SIDE = MagmaLeft and of order N
if SIDE = MagmaRight.
Arguments
---------
@param[in]
side magma_side_t
- = MagmaLeft: apply Q or Q**H from the Left;
- = MagmaRight: apply Q or Q**H from the Right.
@param[in]
trans magma_trans_t
- = MagmaNoTrans: No transpose, apply Q;
- = Magma_ConjTrans: Conjugate transpose, apply Q**H.
@param[in]
m INTEGER
The number of rows of the matrix C. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix C. N >= 0.
@param[in]
k INTEGER
The number of elementary reflectors whose product defines
the matrix Q.
If SIDE = MagmaLeft, M >= K >= 0;
if SIDE = MagmaRight, N >= K >= 0.
@param[in]
dA COMPLEX array on the GPU, dimension (LDDA,K)
The i-th column must contain the vector which defines the
elementary reflector H(i), for i = 1,2,...,k, as returned by
CGEQRF in the first k columns of its array argument DA.
DA is modified by the routine but restored on exit.
@param[in]
ldda INTEGER
The leading dimension of the array DA.
If SIDE = MagmaLeft, LDDA >= max(1,M);
if SIDE = MagmaRight, LDDA >= max(1,N).
@param[in,out]
dC COMPLEX array on the GPU, dimension (LDDC,N)
On entry, the M-by-N matrix C.
On exit, C is overwritten by Q*C or Q**H * C or C * Q**H or C*Q.
@param[in]
lddc INTEGER
The leading dimension of the array DC. LDDC >= max(1,M).
@param[in]
dT COMPLEX array on the GPU that is the output
(the 9th argument) of magma_cgeqrf_gpu.
@param[in]
nb INTEGER
This is the blocking size that was used in pre-computing DT, e.g.,
the blocking size used in magma_cgeqrf_gpu.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_cheev_2stage
********************************************************************/
extern "C" magma_int_t
magma_cunmqr_gpu_2stages(magma_side_t side, magma_trans_t trans,
magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloatComplex *dA, magma_int_t ldda,
magmaFloatComplex *dC, magma_int_t lddc,
magmaFloatComplex *dT, magma_int_t nb,
magma_int_t *info)
{
#define dA(i_,j_) (dA + (i_) + (j_)*ldda)
#define dC(i_,j_) (dC + (i_) + (j_)*lddc)
magmaFloatComplex *dwork;
magma_int_t i, i1, i2, step, ib, ic, jc, mi, ni, nq, nw;
int left, notran;
//magma_int_t lwkopt;
*info = 0;
left = (side == MagmaLeft);
notran = (trans == MagmaNoTrans);
/* NQ is the order of Q and NW is the minimum dimension of WORK */
if (left) {
nq = m;
nw = n;
} else {
nq = n;
nw = m;
}
if ( ! left && side != MagmaRight ) {
*info = -1;
} else if ( ! notran && trans != Magma_ConjTrans ) {
*info = -2;
} else if (m < 0) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (k < 0 || k > nq) {
*info = -5;
} else if (ldda < max(1,nq)) {
*info = -7;
} else if (lddc < max(1,m)) {
*info = -10;
}
// TODO alloc after xerbla & quick return, else memory leak
if (MAGMA_SUCCESS != magma_cmalloc( &dwork, n*nb )) {
printf ("!!!! cungqr_2stage magma_alloc failed for: dwork\n" );
exit(-1);
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0 || k == 0) {
return *info;
}
if ( (left && (! notran)) || ( (! left) && notran ) ) {
i1 = 0;
i2 = k;
step = nb;
} else {
i1 = (k - 1) / nb * nb;
i2 = 0;
step = -nb;
}
// silence "uninitialized" warnings
mi = 0;
ni = 0;
if (left) {
ni = n;
jc = 0;
} else {
mi = m;
ic = 0;
}
for (i=i1; (step < 0 ? i >= i2 : i < i2); i += step) {
ib = min(nb, k - i);
if (left) {
mi = m - i;
ic = i;
}
else {
ni = n - i;
jc = i;
}
magma_clarfb_gpu( MagmaLeft, trans, MagmaForward, MagmaColumnwise,
mi, ni, ib, dA(i,i), ldda, dT+i*nb, nb,
dC(ic,jc), lddc, dwork, nw );
}
magma_free( dwork );
return *info;
} /* magma_cunmqr_gpu_2stages */
/**
Purpose
-------
CUNGQR generates an M-by-N COMPLEX matrix Q with orthonormal columns,
which is defined as the first N columns of a product of K elementary
reflectors of order M
Q = H(1) H(2) . . . H(k)
as returned by CGEQRF_GPU.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix Q. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix Q. M >= N >= 0.
@param[in]
k INTEGER
The number of elementary reflectors whose product defines the
matrix Q. N >= K >= 0.
@param[in,out]
dA COMPLEX array A on the GPU device,
dimension (LDDA,N). On entry, the i-th column must contain
the vector which defines the elementary reflector H(i), for
i = 1,2,...,k, as returned by CGEQRF_GPU in the first k
columns of its array argument A.
On exit, the M-by-N matrix Q.
@param[in]
ldda INTEGER
The first dimension of the array A. LDDA >= max(1,M).
@param[in]
tau COMPLEX array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by CGEQRF_GPU.
@param[in]
dT COMPLEX work space array on the GPU device,
dimension (MIN(M, N) )*NB.
This must be the 6th argument of magma_cgeqrf_gpu
[ note that if N here is bigger than N in magma_cgeqrf_gpu,
the workspace requirement DT in magma_cgeqrf_gpu must be
as specified in this routine ].
@param[in]
nb INTEGER
This is the block size used in CGEQRF_GPU, and correspondingly
the size of the T matrices, used in the factorization, and
stored in DT.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument has an illegal value
@ingroup magma_cheev_2stage
********************************************************************/
extern "C" magma_int_t
magma_cungqr_2stage_gpu(
magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloatComplex_ptr dA, magma_int_t ldda,
magmaFloatComplex *tau,
magmaFloatComplex_ptr dT, magma_int_t nb,
magma_int_t *info)
{
#define dA(a_1,a_2) (dA + (a_2)*(ldda) + (a_1))
#define dT(a_1) (dT + (a_1)*nb)
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magma_int_t i__1, i__2, i__3;
//magma_int_t lwork;
magma_int_t i, ib, ki, kk; //, iinfo;
//magma_int_t lddwork = min(m, n);
//magmaFloatComplex *work, *panel;
magmaFloatComplex_ptr dwork;
//magma_queue_t stream[2];
magma_int_t ldt=nb; // need to be an input parameter
*info = 0;
if (m < 0) {
*info = -1;
} else if ((n < 0) || (n > m)) {
*info = -2;
} else if ((k < 0) || (k > n)) {
*info = -3;
} else if (ldda < max(1,m)) {
*info = -5;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (n <= 0)
return *info;
if (MAGMA_SUCCESS != magma_cmalloc( &dwork, n*nb )) {
printf ("!!!! cungqr_2stage magma_alloc failed for: dwork\n" );
exit(-1);
}
if ( (nb > 1) && (nb < k) ) {
/* Use blocked code after the last block.
The first kk columns are handled by the block method.
ki is start of 2nd-to-last block. */
ki = (k - nb - 1) / nb * nb;
kk = min(k, ki + nb);
/* Set A(1:kk,kk+1:n) to zero. */
/* and A(kk+1:m, kk+1:n) = I */
magmablas_claset( MagmaFull, kk, n-kk, c_zero, c_zero, dA(0, kk), ldda );
magmablas_claset( MagmaFull, m-kk, n-kk, c_zero, c_one, dA(kk,kk), ldda );
}
else {
ki = 0;
kk = 0;
}
/* Allocate work space on CPU in pinned memory */
//lwork = (n+m) * nb;
//if (kk < n)
// lwork = max(lwork, n * nb + (m-kk)*(n-kk));
//if (MAGMA_SUCCESS != magma_cmalloc_pinned( &work, (lwork) )) {
// *info = MAGMA_ERR_HOST_ALLOC;
// return *info;
//}
//panel = work + n * nb;
//magma_queue_create( &stream[0] );
//magma_queue_create( &stream[1] );
/* Use unblocked code for the last or only block. */
if (kk < n) {
i__1 = m - kk;
i__2 = n - kk;
i__3 = k - kk;
//magma_cgetmatrix(i__1, i__2, dA(kk, kk), ldda, panel, i__1);
//lapackf77_cungqr(&i__1, &i__2, &i__3, panel, &i__1, &tau[kk],
// work, &lwork, &iinfo);
//
//magma_csetmatrix(i__1, i__2, panel, i__1, dA(kk, kk), ldda);
magma_clarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
i__1, i__2, i__3,
dA(kk, kk-nb), ldda, dT(kk-nb), ldt,
dA(kk, kk), ldda, dwork, i__2);
//magmablas_claset(MagmaFull, kk-nb, nb, c_zero, c_zero, dA(0,kk-nb), ldda);
//magmablas_claset(MagmaFull, m-(kk-nb), nb, c_zero, c_one, dA(kk-nb,kk-nb), ldda);
}
if (kk > 0) {
/* Use blocked code */
for (i = ki; i >= nb; i -= nb) {
ib = min(nb, k - i);
/* Send current panel to the CPU for update */
i__2 = m - i;
//magma_cgetmatrix_async( i__2, ib, dA(i,i), ldda, panel, i__2, stream[0] ); // verify
if (i + ib < n) {
/* Apply H to A(i:m,i+ib:n) from the left */
i__3 = n - i;
magmablas_claset( MagmaFull, i, ib, c_zero, c_zero, dA(0,i), ldda );
magmablas_claset( MagmaFull, m-i, ib, c_zero, c_one, dA(i,i), ldda );
magma_clarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
i__2, i__3, ib,
dA(i, i-nb), ldda, dT(i-nb), ldt,
dA(i, i), ldda, dwork, i__3);
}
/* Apply H to rows i:m of current block on the CPU */
//magma_queue_sync( stream[0] );
//lapackf77_cungqr(&i__2, &ib, &ib, panel, &i__2, &tau[i],
// work, &lwork, &iinfo);
//magma_csetmatrix_async( i__2, ib, panel, i__2, dA(i,i), ldda, stream[1] ); // verify
/* Set rows 1:i-1 of current block to zero */
i__2 = i + ib;
//magmablas_claset(MagmaFull, i-ib, ib, c_zero, c_zero, dA(0,i-ib), ldda);
//magmablas_claset(MagmaFull, m-(i-ib), ib, c_zero, c_one, dA(i-ib,i-ib), ldda);
}
}
magmablas_claset( MagmaFull, m, nb, c_zero, c_one, dA(0,0), ldda );
magma_free( dwork );
//magma_free_pinned( work );
//magma_queue_destroy( stream[0] );
//magma_queue_destroy( stream[1] );
return *info;
} /* magma_cungqr_gpu */
/**
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
-------
CGEQRF_OOC 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 is an out-of-core (ooc) version that is similar to magma_cgeqrf but
the difference is that this version can use a GPU even if the matrix
does not fit into the GPU memory at once.
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 >= 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_cgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_cgeqrf_ooc(
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)
/* Constants */
const magmaFloatComplex c_one = MAGMA_C_ONE;
/* Local variables */
magmaFloatComplex_ptr dA, dwork;
magma_int_t i, ib, IB, j, min_mn, lddwork, ldda, rows;
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);
*info = 0;
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,m)) {
*info = -4;
} else if (lwork < max(1,n) && ! lquery) {
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Check how much memory do we have */
size_t freeMem, totalMem;
cudaMemGetInfo( &freeMem, &totalMem );
freeMem /= sizeof(magmaFloatComplex);
magma_int_t NB = magma_int_t(0.8*freeMem/m);
NB = (NB / nb) * nb;
if (NB >= n)
return magma_cgeqrf(m, n, A, lda, tau, work, lwork, info);
min_mn = min(m,n);
if (min_mn == 0) {
work[0] = c_one;
return *info;
}
lddwork = magma_roundup( NB, 32 ) + nb;
ldda = magma_roundup( m, 32 );
if (MAGMA_SUCCESS != magma_cmalloc( &dA, (NB + nb)*ldda + nb*lddwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_queue_t queues[2];
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[0] );
magma_queue_create( cdev, &queues[1] );
magmaFloatComplex_ptr ptr = dA + ldda*NB;
dwork = dA + ldda*(NB + nb);
/* start the main loop over the blocks that fit in the GPU memory */
for (i=0; i < n; i += NB) {
IB = min( n-i, NB );
//printf("Processing %5d columns -- %5d to %5d ... \n", IB, i, i+IB);
/* 1. Copy the next part of the matrix to the GPU */
magma_csetmatrix_async( m, IB,
A(0,i), lda,
dA(0,0), ldda, queues[0] );
magma_queue_sync( queues[0] );
/* 2. Update it with the previous transformations */
for (j=0; j < min(i,min_mn); j += nb) {
ib = min( min_mn-j, nb );
/* Get a panel in ptr. */
// 1. Form the triangular factor of the block reflector
// 2. Send it to the GPU.
// 3. Put 0s in the upper triangular part of V.
// 4. Send V to the GPU in ptr.
// 5. Update the matrix.
// 6. Restore the upper part of V.
rows = m-j;
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&rows, &ib, A(j,j), &lda, tau+j, work, &ib);
magma_csetmatrix_async( ib, ib,
work, ib,
dwork, lddwork, queues[1] );
magma_cpanel_to_q( MagmaUpper, ib, A(j,j), lda, work+ib*ib );
magma_csetmatrix_async( rows, ib,
A(j,j), lda,
ptr, rows, queues[1] );
magma_queue_sync( queues[1] );
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, IB, ib,
ptr, rows, dwork, lddwork,
dA(j, 0), ldda, dwork+ib, lddwork, queues[1] );
magma_cq_to_panel( MagmaUpper, ib, A(j,j), lda, work+ib*ib );
}
/* 3. Do a QR on the current part */
if (i < min_mn)
magma_cgeqrf2_gpu( m-i, IB, dA(i,0), ldda, tau+i, info );
/* 4. Copy the current part back to the CPU */
magma_cgetmatrix_async( m, IB,
dA(0,0), ldda,
A(0,i), lda, queues[0] );
}
magma_queue_sync( queues[0] );
magma_queue_destroy( queues[0] );
magma_queue_destroy( queues[1] );
magma_free( dA );
return *info;
} /* magma_cgeqrf_ooc */
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 */
