// --------------------
extern "C" void
magma_ccopyvector(
magma_int_t n,
magmaFloatComplex_const_ptr dx_src, size_t dx_offset, magma_int_t incx,
magmaFloatComplex_ptr dy_dst, size_t dy_offset, magma_int_t incy,
magma_queue_t queue )
{
if (n <= 0)
return;
if (incx == 1 && incy == 1) {
cl_int err = clEnqueueReadBuffer(
queue, dx_src, CL_TRUE,
dx_offset*sizeof(magmaFloatComplex), n*sizeof(magmaFloatComplex),
dy_dst, dy_offset*sizeof(magmaFloatComplex), NULL, g_event);
check_error( err );
}
else {
magma_int_t ldda = incx;
magma_int_t lddb = incy;
magma_ccopymatrix( 1, n,
dx_src, dx_offset, ldda,
dy_dst, dy_offset, lddb,
queue);
}
}
/**
@deprecated
Purpose
-------
CLAQPS computes a step of QR factorization with column pivoting
of a complex M-by-N matrix A by using Blas-3. It tries to factorize
NB columns from A starting from the row OFFSET+1, and updates all
of the matrix with Blas-3 xGEMM.
In some cases, due to catastrophic cancellations, it cannot
factorize NB columns. Hence, the actual number of factorized
columns is returned in KB.
Block A(1:OFFSET,1:N) is accordingly pivoted, but not factorized.
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]
offset INTEGER
The number of rows of A that have been factorized in
previous steps.
@param[in]
nb INTEGER
The number of columns to factorize.
@param[out]
kb INTEGER
The number of columns actually factorized.
@param[in,out]
A COMPLEX array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, block A(OFFSET+1:M,1:KB) is the triangular
factor obtained and block A(1:OFFSET,1:N) has been
accordingly pivoted, but no factorized.
The rest of the matrix, block A(OFFSET+1:M,KB+1:N) has
been updated.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[in,out]
jpvt INTEGER array, dimension (N)
JPVT(I) = K <==> Column K of the full matrix A has been
permuted into position I in AP.
@param[out]
tau COMPLEX array, dimension (KB)
The scalar factors of the elementary reflectors.
@param[in,out]
vn1 REAL array, dimension (N)
The vector with the partial column norms.
@param[in,out]
vn2 REAL array, dimension (N)
The vector with the exact column norms.
@param[in,out]
auxv COMPLEX array, dimension (NB)
Auxiliar vector.
@param[in,out]
F COMPLEX array, dimension (LDF,NB)
Matrix F' = L*Y'*A.
@param[in]
ldf INTEGER
The leading dimension of the array F. LDF >= max(1,N).
@ingroup magma_cgeqp3_aux
********************************************************************/
extern "C" magma_int_t
magma_claqps_gpu(magma_int_t m, magma_int_t n, magma_int_t offset,
magma_int_t nb, magma_int_t *kb,
magmaFloatComplex *A, magma_int_t lda,
magma_int_t *jpvt, magmaFloatComplex *tau,
float *vn1, float *vn2,
magmaFloatComplex *auxv,
magmaFloatComplex *F, magma_int_t ldf)
{
#define A(i, j) (A + (i) + (j)*(lda ))
#define F(i, j) (F + (i) + (j)*(ldf ))
magmaFloatComplex c_zero = MAGMA_C_MAKE( 0.,0.);
magmaFloatComplex c_one = MAGMA_C_MAKE( 1.,0.);
magmaFloatComplex c_neg_one = MAGMA_C_MAKE(-1.,0.);
magma_int_t ione = 1;
magma_int_t i__1, i__2;
//float d__1;
magmaFloatComplex z__1;
//magma_int_t j;
magma_int_t k, rk;
//magmaFloatComplex Akk;
magmaFloatComplex *Aks;
magmaFloatComplex tauk = MAGMA_C_ZERO;
magma_int_t pvt;
//float temp, temp2;
float tol3z;
magma_int_t itemp;
float lsticc, *lsticcs;
magma_int_t lastrk;
magma_smalloc( &lsticcs, 1+256*(n+255)/256 );
lastrk = min( m, n + offset );
tol3z = magma_ssqrt( lapackf77_slamch("Epsilon"));
lsticc = 0;
k = 0;
magma_cmalloc( &Aks, nb );
while( k < nb && lsticc == 0 ) {
rk = offset + k;
/* Determine ith pivot column and swap if necessary */
// subtract 1 from Fortran/CUBLAS isamax; pvt, k are 0-based.
pvt = k + magma_isamax( n-k, &vn1[k], ione ) - 1;
if (pvt != k) {
/*if (pvt >= nb) {
// 1. Start copy from GPU
magma_cgetmatrix_async( m - offset - nb, 1,
dA(offset + nb, pvt), ldda,
A (offset + nb, pvt), lda, stream );
}*/
/* F gets swapped so F must be sent at the end to GPU */
i__1 = k;
/*if (pvt < nb) {
// no need of transfer if pivot is within the panel
blasf77_cswap( &m, A(0, pvt), &ione, A(0, k), &ione );
}
else {
// 1. Finish copy from GPU
magma_queue_sync( stream );
// 2. Swap as usual on CPU
blasf77_cswap(&m, A(0, pvt), &ione, A(0, k), &ione);
// 3. Restore the GPU
magma_csetmatrix_async( m - offset - nb, 1,
A (offset + nb, pvt), lda,
dA(offset + nb, pvt), ldda, stream);
}*/
magmablas_cswap( m, A(0, pvt), ione, A(0, k), ione );
//blasf77_cswap( &i__1, F(pvt,0), &ldf, F(k,0), &ldf );
magmablas_cswap( i__1, F(pvt, 0), ldf, F(k, 0), ldf);
itemp = jpvt[pvt];
jpvt[pvt] = jpvt[k];
jpvt[k] = itemp;
//vn1[pvt] = vn1[k];
//vn2[pvt] = vn2[k];
#if defined(PRECISION_d) || defined(PRECISION_z)
//magma_dswap( 1, &vn1[pvt], 1, &vn1[k], 1 );
//magma_dswap( 1, &vn2[pvt], 1, &vn2[k], 1 );
magma_dswap( 2, &vn1[pvt], n+offset, &vn1[k], n+offset );
#else
//magma_sswap( 1, &vn1[pvt], 1, &vn1[k], 1 );
//magma_sswap( 1, &vn2[pvt], 1, &vn2[k], 1 );
magma_sswap(2, &vn1[pvt], n+offset, &vn1[k], n+offset);
#endif
}
/* Apply previous Householder reflectors to column K:
A(RK:M,K) := A(RK:M,K) - A(RK:M,1:K-1)*F(K,1:K-1)'.
Optimization: multiply with beta=0; wait for vector and subtract */
if (k > 0) {
/*#if (defined(PRECISION_c) || defined(PRECISION_z))
for (j = 0; j < k; ++j) {
*F(k,j) = MAGMA_C_CNJG( *F(k,j) );
}
#endif*/
//#define RIGHT_UPDATE
#ifdef RIGHT_UPDATE
i__1 = m - offset - nb;
i__2 = k;
magma_cgemv( MagmaNoTrans, i__1, i__2,
c_neg_one, A(offset+nb, 0), lda,
F(k, 0), ldf,
c_one, A(offset+nb, k), ione );
#else
i__1 = m - rk;
i__2 = k;
/*blasf77_cgemv( MagmaNoTransStr, &i__1, &i__2,
&c_neg_one, A(rk, 0), &lda,
F(k, 0), &ldf,
&c_one, A(rk, k), &ione );*/
magma_cgemv( MagmaNoTrans, i__1, i__2,
c_neg_one, A(rk, 0), lda,
F(k, 0), ldf,
c_one, A(rk, k), ione );
#endif
/*#if (defined(PRECISION_c) || defined(PRECISION_z))
for (j = 0; j < k; ++j) {
*F(k,j) = MAGMA_C_CNJG( *F(k,j) );
}
#endif*/
}
/* Generate elementary reflector H(k). */
magma_clarfg_gpu(m-rk, A(rk, k), A(rk + 1, k), &tau[k], &vn1[k], &Aks[k]);
//Akk = *A(rk, k);
//*A(rk, k) = c_one;
//magma_cgetvector( 1, &Aks[k], 1, &Akk, 1 );
/* needed to avoid the race condition */
if (k == 0) magma_csetvector( 1, &c_one, 1, A(rk, k), 1 );
else magma_ccopymatrix( 1, 1, A(offset, 0), 1, A(rk, k), 1 );
/* Compute Kth column of F:
Compute F(K+1:N,K) := tau(K)*A(RK:M,K+1:N)'*A(RK:M,K) on the GPU */
if (k < n-1 || k > 0) magma_cgetvector( 1, &tau[k], 1, &tauk, 1 );
if (k < n-1) {
i__1 = m - rk;
i__2 = n - k - 1;
/* Send the vector to the GPU */
//magma_csetmatrix( i__1, 1, A(rk, k), lda, dA(rk,k), ldda );
/* Multiply on GPU */
// was CALL CGEMV( 'Conjugate transpose', M-RK+1, N-K,
// TAU( K ), A( RK, K+1 ), LDA,
// A( RK, K ), 1,
// CZERO, F( K+1, K ), 1 )
//magma_cgetvector( 1, &tau[k], 1, &tauk, 1 );
magma_cgemv( MagmaConjTrans, m-rk, n-k-1,
tauk, A( rk, k+1 ), lda,
A( rk, k ), 1,
c_zero, F( k+1, k ), 1 );
//magma_cscal( m-rk, tau[k], F( k+1, k), 1 );
//magma_int_t i__3 = nb-k-1;
//magma_int_t i__4 = i__2 - i__3;
//magma_int_t i__5 = nb-k;
//magma_cgemv( MagmaConjTrans, i__1 - i__5, i__2 - i__3,
// tau[k], dA(rk +i__5, k+1+i__3), ldda,
// dA(rk +i__5, k ), ione,
// c_zero, dF(k+1+i__3, k ), ione );
//magma_cgetmatrix_async( i__2-i__3, 1,
// dF(k + 1 +i__3, k), i__2,
// F (k + 1 +i__3, k), i__2, stream );
//blasf77_cgemv( MagmaConjTransStr, &i__1, &i__3,
// &tau[k], A(rk, k+1), &lda,
// A(rk, k ), &ione,
// &c_zero, F(k+1, k ), &ione );
//magma_queue_sync( stream );
//blasf77_cgemv( MagmaConjTransStr, &i__5, &i__4,
// &tau[k], A(rk, k+1+i__3), &lda,
// A(rk, k ), &ione,
// &c_one, F(k+1+i__3, k ), &ione );
}
/* Padding F(1:K,K) with zeros.
for (j = 0; j <= k; ++j) {
magma_csetvector( 1, &c_zero, 1, F(j, k), 1 );
}*/
/* Incremental updating of F:
F(1:N,K) := F(1:N,K) - tau(K)*F(1:N,1:K-1)*A(RK:M,1:K-1)'*A(RK:M,K).
F(1:N,K) := tau(K)*A(RK:M,K+1:N)'*A(RK:M,K) - tau(K)*F(1:N,1:K-1)*A(RK:M,1:K-1)'*A(RK:M,K)
:= tau(K)(A(RK:M,K+1:N)' - F(1:N,1:K-1)*A(RK:M,1:K-1)') A(RK:M,K)
so, F is (updated A)*V */
//if (k > 0 && k < n-1) {
if (k > 0) {
//magma_cgetvector( 1, &tau[k], 1, &tauk, 1 );
z__1 = MAGMA_C_NEGATE( tauk );
#ifdef RIGHT_UPDATE
i__1 = m - offset - nb;
i__2 = k;
magma_cgemv( MagmaConjTrans, i__1, i__2,
z__1, A(offset+nb, 0), lda,
A(offset+nb, k), ione,
c_zero, auxv, ione );
i__1 = k;
magma_cgemv( MagmaNoTrans, n-k-1, i__1,
c_one, F(k+1,0), ldf,
auxv, ione,
c_one, F(k+1,k), ione );
#else
i__1 = m - rk;
i__2 = k;
//blasf77_cgemv( MagmaConjTransStr, &i__1, &i__2,
// &z__1, A(rk, 0), &lda,
// A(rk, k), &ione,
// &c_zero, auxv, &ione );
magma_cgemv( MagmaConjTrans, i__1, i__2,
z__1, A(rk, 0), lda,
A(rk, k), ione,
c_zero, auxv, ione );
//i__1 = k;
//blasf77_cgemv( MagmaNoTransStr, &n, &i__1,
// &c_one, F(0,0), &ldf,
// auxv, &ione,
// &c_one, F(0,k), &ione );
/*magma_cgemv( MagmaNoTrans, n, i__1,
c_one, F(0,0), ldf,
auxv, ione,
c_one, F(0,k), ione );*/
/* I think we only need stricly lower-triangular part :) */
magma_cgemv( MagmaNoTrans, n-k-1, i__2,
c_one, F(k+1,0), ldf,
auxv, ione,
c_one, F(k+1,k), ione );
#endif
}
/* Optimization: On the last iteration start sending F back to the GPU */
/* Update the current row of A:
A(RK,K+1:N) := A(RK,K+1:N) - A(RK,1:K)*F(K+1:N,1:K)'. */
if (k < n-1) {
i__1 = n - k - 1;
i__2 = k + 1;
//blasf77_cgemm( MagmaNoTransStr, MagmaConjTransStr, &ione, &i__1, &i__2,
// &c_neg_one, A(rk, 0 ), &lda,
// F(k+1,0 ), &ldf,
// &c_one, A(rk, k+1), &lda );
#ifdef RIGHT_UPDATE
/* right-looking update of rows, */
magma_cgemm( MagmaNoTrans, MagmaConjTrans, nb-k, i__1, ione,
c_neg_one, A(rk, k ), lda,
F(k+1, k ), ldf,
c_one, A(rk, k+1), lda );
#else
/* left-looking update of rows, *
* since F=A'v with original A, so no right-looking */
magma_cgemm( MagmaNoTrans, MagmaConjTrans, ione, i__1, i__2,
c_neg_one, A(rk, 0 ), lda,
F(k+1,0 ), ldf,
c_one, A(rk, k+1), lda );
#endif
}
/* Update partial column norms. */
if (rk < min(m, n+offset)-1 ) {
magmablas_scnrm2_row_check_adjust(n-k-1, tol3z, &vn1[k+1], &vn2[k+1], A(rk,k+1), lda, lsticcs);
magma_device_sync();
#if defined(PRECISION_d) || defined(PRECISION_z)
magma_sgetvector( 1, &lsticcs[0], 1, &lsticc, 1 );
#else
magma_sgetvector( 1, &lsticcs[0], 1, &lsticc, 1 );
#endif
}
/*if (rk < lastrk) {
for (j = k + 1; j < n; ++j) {
if (vn1[j] != 0.) {
// NOTE: The following 4 lines follow from the analysis in
// Lapack Working Note 176.
temp = MAGMA_C_ABS( *A(rk,j) ) / vn1[j];
temp = max( 0., ((1. + temp) * (1. - temp)) );
d__1 = vn1[j] / vn2[j];
temp2 = temp * (d__1 * d__1);
if (temp2 <= tol3z) {
vn2[j] = (float) lsticc;
lsticc = j;
} else {
vn1[j] *= magma_ssqrt(temp);
}
}
}
}*/
//*A(rk, k) = Akk;
//magma_csetvector( 1, &Akk, 1, A(rk, k), 1 );
//magma_cswap( 1, &Aks[k], 1, A(rk, k), 1 );
++k;
}
magma_ccopymatrix( 1, k, Aks, 1, A(offset, 0), lda+1 );
// leave k as the last column done
--k;
*kb = k + 1;
rk = offset + *kb - 1;
/* Apply the block reflector to the rest of the matrix:
A(OFFSET+KB+1:M,KB+1:N) := A(OFFSET+KB+1:M,KB+1:N) - A(OFFSET+KB+1:M,1:KB)*F(KB+1:N,1:KB)' */
if (*kb < min(n, m - offset)) {
i__1 = m - rk - 1;
i__2 = n - *kb;
/* Send F to the GPU
magma_csetmatrix( i__2, *kb,
F (*kb, 0), ldf,
dF(*kb, 0), i__2 );*/
magma_cgemm( MagmaNoTrans, MagmaConjTrans, i__1, i__2, *kb,
c_neg_one, A(rk+1, 0 ), lda,
F(*kb, 0 ), ldf,
c_one, A(rk+1, *kb), lda );
}
/* Recomputation of difficult columns. */
if ( lsticc > 0 ) {
// printf( " -- recompute dnorms --\n" );
magmablas_scnrm2_check(m-rk-1, n-*kb, A(rk+1,*kb), lda,
&vn1[*kb], lsticcs);
magma_scopymatrix( n-*kb, 1, &vn1[*kb], *kb, &vn2[*kb], *kb);
/*while( lsticc > 0 ) {
itemp = (magma_int_t)(vn2[lsticc] >= 0. ? floor(vn2[lsticc] + .5) : -floor(.5 - vn2[lsticc]));
i__1 = m - rk - 1;
if (lsticc <= nb)
vn1[lsticc] = magma_cblas_scnrm2( i__1, A(rk+1,lsticc), ione );
else {
// Where is the data, CPU or GPU ?
float r1, r2;
r1 = magma_cblas_scnrm2( nb-k, A(rk+1,lsticc), ione );
r2 = magma_scnrm2(m-offset-nb, dA(offset + nb + 1, lsticc), ione);
vn1[lsticc] = magma_ssqrt(r1*r1+r2*r2);
}
// NOTE: The computation of VN1( LSTICC ) relies on the fact that
// SNRM2 does not fail on vectors with norm below the value of SQRT(SLAMCH('S'))
vn2[lsticc] = vn1[lsticc];
lsticc = itemp;*/
}
magma_free(Aks);
magma_free(lsticcs);
return MAGMA_SUCCESS;
} /* magma_claqps */
/**
Purpose
-------
CSYTRF_nopiv_gpu computes the LDLt factorization of a complex symmetric
matrix A.
The factorization has the form
A = U^T * D * U, if UPLO = MagmaUpper, or
A = L * D * L^T, if UPLO = MagmaLower,
where U is an upper triangular matrix, L is lower triangular, and
D is a diagonal matrix.
This is the block version of the algorithm, calling Level 3 BLAS.
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored;
- = MagmaLower: Lower triangle of A is stored.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
dA COMPLEX array on the GPU, dimension (LDDA,N)
On entry, the symmetric matrix A. If UPLO = MagmaUpper, the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
\n
On exit, if INFO = 0, the factor U or L from the Cholesky
factorization A = U^H D U or A = L D L^H.
\n
Higher performance is achieved if A is in pinned memory, e.g.
allocated using cudaMallocHost.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,N).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
if INFO = -6, the GPU memory allocation failed
- > 0: if INFO = i, the leading minor of order i is not
positive definite, and the factorization could not be
completed.
@ingroup magma_csysv_comp
******************************************************************* */
extern "C" magma_int_t
magma_csytrf_nopiv_gpu(
magma_uplo_t uplo, magma_int_t n,
magmaFloatComplex_ptr dA, magma_int_t ldda,
magma_int_t *info)
{
#define A(i, j) (A)
#define dA(i, j) (dA +(j)*ldda + (i))
#define dW(i, j) (dW +(j)*ldda + (i))
#define dWt(i, j) (dW +(j)*nb + (i))
/* Constants */
const magmaFloatComplex c_one = MAGMA_C_ONE;
const magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
/* Local variables */
bool upper = (uplo == MagmaUpper);
magma_int_t j, k, jb, nb, ib, iinfo;
*info = 0;
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (ldda < max(1,n)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return */
if ( n == 0 )
return *info;
nb = magma_get_chetrf_nopiv_nb(n);
ib = min(32, nb); // inner-block for diagonal factorization
magma_queue_t queues[2];
magma_event_t event;
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[0] );
magma_queue_create( cdev, &queues[1] );
magma_event_create( &event );
trace_init( 1, 1, 2, queues );
// CPU workspace
magmaFloatComplex *A;
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &A, nb*nb )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
// GPU workspace
magmaFloatComplex_ptr dW;
if (MAGMA_SUCCESS != magma_cmalloc( &dW, (1+nb)*ldda )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/* Use hybrid blocked code. */
if (upper) {
//=========================================================
// Compute the LDLt factorization A = U'*D*U without pivoting.
// main loop
for (j=0; j < n; j += nb) {
jb = min(nb, (n-j));
// copy A(j,j) back to CPU
trace_gpu_start( 0, 0, "get", "get" );
magma_event_sync( event );
magma_cgetmatrix_async( jb, jb, dA(j, j), ldda, A(j,j), nb, queues[1] );
trace_gpu_end( 0, 0 );
// factorize the diagonal block
magma_queue_sync( queues[1] );
trace_cpu_start( 0, "potrf", "potrf" );
magma_csytrf_nopiv_cpu( MagmaUpper, jb, ib, A(j, j), nb, info );
trace_cpu_end( 0 );
if (*info != 0) {
*info = *info + j;
break;
}
// copy A(j,j) back to GPU
trace_gpu_start( 0, 0, "set", "set" );
magma_csetmatrix_async( jb, jb, A(j, j), nb, dA(j, j), ldda, queues[0] );
trace_gpu_end( 0, 0 );
if ( (j+jb) < n) {
// compute the off-diagonal blocks of current block column
trace_gpu_start( 0, 0, "trsm", "trsm" );
magma_ctrsm( MagmaLeft, MagmaUpper, MagmaTrans, MagmaUnit,
jb, (n-j-jb),
c_one, dA(j, j), ldda,
dA(j, j+jb), ldda, queues[0] );
magma_ccopymatrix( jb, n-j-jb, dA( j, j+jb ), ldda, dWt( 0, j+jb ), nb, queues[0] );
// update the trailing submatrix with D
magmablas_clascl_diag( MagmaUpper, jb, n-j-jb,
dA(j, j), ldda,
dA(j, j+jb), ldda,
queues[0], &iinfo );
trace_gpu_end( 0, 0 );
// update the trailing submatrix with U and W
trace_gpu_start( 0, 0, "gemm", "gemm" );
for (k=j+jb; k < n; k += nb) {
magma_int_t kb = min(nb,n-k);
magma_cgemm( MagmaTrans, MagmaNoTrans, kb, n-k, jb,
c_neg_one, dWt(0, k), nb,
dA(j, k), ldda,
c_one, dA(k, k), ldda, queues[0] );
if (k == j+jb)
magma_event_record( event, queues[0] );
}
trace_gpu_end( 0, 0 );
}
}
} else {
//=========================================================
// Compute the LDLt factorization A = L*D*L' without pivoting.
// main loop
for (j=0; j < n; j += nb) {
jb = min(nb, (n-j));
// copy A(j,j) back to CPU
trace_gpu_start( 0, 0, "get", "get" );
magma_event_sync( event );
magma_cgetmatrix_async( jb, jb, dA(j, j), ldda, A(j,j), nb, queues[1] );
trace_gpu_end( 0, 0 );
// factorize the diagonal block
magma_queue_sync( queues[1] );
trace_cpu_start( 0, "potrf", "potrf" );
magma_csytrf_nopiv_cpu( MagmaLower, jb, ib, A(j, j), nb, info );
trace_cpu_end( 0 );
if (*info != 0) {
*info = *info + j;
break;
}
// copy A(j,j) back to GPU
trace_gpu_start( 0, 0, "set", "set" );
magma_csetmatrix_async( jb, jb, A(j, j), nb, dA(j, j), ldda, queues[0] );
trace_gpu_end( 0, 0 );
if ( (j+jb) < n) {
// compute the off-diagonal blocks of current block column
trace_gpu_start( 0, 0, "trsm", "trsm" );
magma_ctrsm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
(n-j-jb), jb,
c_one, dA(j, j), ldda,
dA(j+jb, j), ldda, queues[0] );
magma_ccopymatrix( n-j-jb,jb, dA( j+jb, j ), ldda, dW( j+jb, 0 ), ldda, queues[0] );
// update the trailing submatrix with D
magmablas_clascl_diag(MagmaLower, n-j-jb, jb,
dA(j, j), ldda,
dA(j+jb, j), ldda,
queues[0], &iinfo);
trace_gpu_end( 0, 0 );
// update the trailing submatrix with L and W
trace_gpu_start( 0, 0, "gemm", "gemm" );
for (k=j+jb; k < n; k += nb) {
magma_int_t kb = min(nb,n-k);
magma_cgemm( MagmaNoTrans, MagmaTrans, n-k, kb, jb,
c_neg_one, dA(k, j), ldda,
dW(k, 0), ldda,
c_one, dA(k, k), ldda, queues[0] );
if (k == j+jb)
magma_event_record( event, queues[0] );
}
trace_gpu_end( 0, 0 );
}
}
}
trace_finalize( "chetrf.svg","trace.css" );
magma_queue_destroy( queues[0] );
magma_queue_destroy( queues[1] );
magma_event_destroy( event );
magma_free( dW );
magma_free_pinned( A );
return *info;
} /* magma_csytrf_nopiv */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cswap, cswapblk, claswp, claswpx
*/
int main( int argc, char** argv)
{
TESTING_INIT();
magmaFloatComplex *h_A1, *h_A2;
magmaFloatComplex *h_R1, *h_R2;
magmaFloatComplex_ptr d_A1, d_A2;
// row-major and column-major performance
real_Double_t row_perf0 = MAGMA_D_NAN, col_perf0 = MAGMA_D_NAN;
real_Double_t row_perf1 = MAGMA_D_NAN, col_perf1 = MAGMA_D_NAN;
real_Double_t row_perf2 = MAGMA_D_NAN, col_perf2 = MAGMA_D_NAN;
real_Double_t row_perf4 = MAGMA_D_NAN;
real_Double_t row_perf5 = MAGMA_D_NAN, col_perf5 = MAGMA_D_NAN;
real_Double_t row_perf6 = MAGMA_D_NAN, col_perf6 = MAGMA_D_NAN;
real_Double_t row_perf7 = MAGMA_D_NAN;
real_Double_t cpu_perf = MAGMA_D_NAN;
real_Double_t time, gbytes;
magma_int_t N, lda, ldda, nb, j;
magma_int_t ione = 1;
magma_int_t *ipiv, *ipiv2;
magmaInt_ptr d_ipiv;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_queue_t queue = 0;
printf(" %8s cswap cswap cswapblk claswp claswp2 claswpx ccopymatrix CPU (all in )\n", g_platform_str );
printf(" N nb row-maj/col-maj row-maj/col-maj row-maj/col-maj row-maj row-maj row-maj/col-maj row-blk/col-blk claswp (GByte/s)\n");
printf("=========================================================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
// For an N x N matrix, swap nb rows or nb columns using various methods.
// Each test is assigned one bit in the 'check' bitmask; bit=1 indicates failure.
// The variable 'shift' keeps track of which bit is for current test
int shift = 1;
int check = 0;
N = opts.nsize[itest];
lda = N;
ldda = ((N+31)/32)*32;
nb = (opts.nb > 0 ? opts.nb : magma_get_cgetrf_nb( N ));
nb = min( N, nb );
// each swap does 2N loads and 2N stores, for nb swaps
gbytes = sizeof(magmaFloatComplex) * 4.*N*nb / 1e9;
TESTING_MALLOC_PIN( h_A1, magmaFloatComplex, lda*N );
TESTING_MALLOC_PIN( h_A2, magmaFloatComplex, lda*N );
TESTING_MALLOC_PIN( h_R1, magmaFloatComplex, lda*N );
TESTING_MALLOC_PIN( h_R2, magmaFloatComplex, lda*N );
TESTING_MALLOC_CPU( ipiv, magma_int_t, nb );
TESTING_MALLOC_CPU( ipiv2, magma_int_t, nb );
TESTING_MALLOC_DEV( d_ipiv, magma_int_t, nb );
TESTING_MALLOC_DEV( d_A1, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( d_A2, magmaFloatComplex, ldda*N );
// getrf always makes ipiv[j] >= j+1, where ipiv is one based and j is zero based
// some implementations (e.g., MacOS dlaswp) assume this
for( j=0; j < nb; j++ ) {
ipiv[j] = (rand() % (N-j)) + j + 1;
assert( ipiv[j] >= j+1 );
assert( ipiv[j] <= N );
}
/* =====================================================================
* cublas / clBLAS / Xeon Phi cswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
#ifdef HAVE_CUBLAS
cublasCswap( opts.handle, N, d_A1+ldda*j, 1, d_A2+ldda*(ipiv[j]-1), 1 );
#else
magma_cswap( N, d_A1, ldda*j, 1, d_A2, ldda*(ipiv[j]-1), 1, opts.queue );
#endif
}
}
time = magma_sync_wtime( queue ) - time;
row_perf0 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
#ifdef HAVE_CUBLAS
cublasCswap( opts.handle, N, d_A1+j, ldda, d_A2+ipiv[j]-1, ldda );
#else
magma_cswap( N, d_A1, j, ldda, d_A2, ipiv[j]-1, ldda, opts.queue );
#endif
}
}
time = magma_sync_wtime( queue ) - time;
col_perf0 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* =====================================================================
* cswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
magmablas_cswap( N, d_A1+ldda*j, 1, d_A2+ldda*(ipiv[j]-1), 1);
}
}
time = magma_sync_wtime( queue ) - time;
row_perf1 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
magmablas_cswap( N, d_A1+j, ldda, d_A2+ipiv[j]-1, ldda );
}
}
time = magma_sync_wtime( queue ) - time;
col_perf1 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* =====================================================================
* cswapblk, blocked version (2 matrices)
*/
#ifdef HAVE_CUBLAS
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_cswapblk( MagmaRowMajor, N, d_A1, ldda, d_A2, ldda, 1, nb, ipiv, 1, 0);
time = magma_sync_wtime( queue ) - time;
row_perf2 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_cswapblk( MagmaColMajor, N, d_A1, ldda, d_A2, ldda, 1, nb, ipiv, 1, 0);
time = magma_sync_wtime( queue ) - time;
col_perf2 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
#endif
/* =====================================================================
* LAPACK-style claswp (1 matrix)
*/
#ifdef HAVE_CUBLAS
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_claswp( N, d_A1, ldda, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
row_perf4 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
#endif
/* =====================================================================
* LAPACK-style claswp (1 matrix) - d_ipiv on GPU
*/
#ifdef HAVE_CUBLAS
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magma_setvector( nb, sizeof(magma_int_t), ipiv, 1, d_ipiv, 1 );
magmablas_claswp2( N, d_A1, ldda, 1, nb, d_ipiv, 1 );
time = magma_sync_wtime( queue ) - time;
row_perf7 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
#endif
/* =====================================================================
* LAPACK-style claswpx (extended for row- and col-major) (1 matrix)
*/
#ifdef HAVE_CUBLAS
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_claswpx( N, d_A1, ldda, 1, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
row_perf5 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* Col Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_claswpx( N, d_A1, 1, ldda, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
col_perf5 = gbytes / time;
#endif
/* LAPACK swap on CPU for comparison */
time = magma_wtime();
lapackf77_claswp( &N, h_A1, &lda, &ione, &nb, ipiv, &ione);
time = magma_wtime() - time;
cpu_perf = gbytes / time;
#ifdef HAVE_CUBLAS
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
#endif
/* =====================================================================
* Copy matrix.
*/
time = magma_sync_wtime( queue );
magma_ccopymatrix( N, nb, d_A1, ldda, d_A2, ldda );
time = magma_sync_wtime( queue ) - time;
// copy reads 1 matrix and writes 1 matrix, so has half gbytes of swap
col_perf6 = 0.5 * gbytes / time;
time = magma_sync_wtime( queue );
magma_ccopymatrix( nb, N, d_A1, ldda, d_A2, ldda );
time = magma_sync_wtime( queue ) - time;
// copy reads 1 matrix and writes 1 matrix, so has half gbytes of swap
row_perf6 = 0.5 * gbytes / time;
printf("%5d %3d %6.2f%c/ %6.2f%c %6.2f%c/ %6.2f%c %6.2f%c/ %6.2f%c %6.2f%c %6.2f%c %6.2f%c/ %6.2f%c %6.2f / %6.2f %6.2f %10s\n",
(int) N, (int) nb,
row_perf0, ((check & 0x001) != 0 ? '*' : ' '),
col_perf0, ((check & 0x002) != 0 ? '*' : ' '),
row_perf1, ((check & 0x004) != 0 ? '*' : ' '),
col_perf1, ((check & 0x008) != 0 ? '*' : ' '),
row_perf2, ((check & 0x010) != 0 ? '*' : ' '),
col_perf2, ((check & 0x020) != 0 ? '*' : ' '),
row_perf4, ((check & 0x040) != 0 ? '*' : ' '),
row_perf7, ((check & 0x080) != 0 ? '*' : ' '),
row_perf5, ((check & 0x100) != 0 ? '*' : ' '),
col_perf5, ((check & 0x200) != 0 ? '*' : ' '),
row_perf6,
col_perf6,
cpu_perf,
(check == 0 ? "ok" : "* failed") );
status += ! (check == 0);
TESTING_FREE_PIN( h_A1 );
TESTING_FREE_PIN( h_A2 );
TESTING_FREE_PIN( h_R1 );
TESTING_FREE_PIN( h_R2 );
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( ipiv2 );
TESTING_FREE_DEV( d_ipiv );
TESTING_FREE_DEV( d_A1 );
TESTING_FREE_DEV( d_A2 );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/***************************************************************************//**
Purpose
-------
CGEQRS solves the least squares problem
min || A*X - C ||
using the QR factorization A = Q*R computed by CGEQRF_GPU.
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. M >= N >= 0.
@param[in]
nrhs INTEGER
The number of columns of the matrix C. NRHS >= 0.
@param[in]
dA COMPLEX array on the GPU, dimension (LDDA,N)
The i-th column must contain the vector which defines the
elementary reflector H(i), for i = 1,2,...,n, as returned by
CGEQRF_GPU in the first n columns of its array argument A.
@param[in]
ldda INTEGER
The leading dimension of the array A, LDDA >= M.
@param[in]
tau COMPLEX array, dimension (N)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by MAGMA_CGEQRF_GPU.
@param[in,out]
dB COMPLEX array on the GPU, dimension (LDDB,NRHS)
On entry, the M-by-NRHS matrix C.
On exit, the N-by-NRHS solution matrix X.
@param[in,out]
dT COMPLEX array that is the output (the 6th argument)
of magma_cgeqrf_gpu of size
2*MIN(M, N)*NB + ceil(N/32)*32 )* MAX(NB, NRHS).
The array starts with a block of size MIN(M,N)*NB that stores
the triangular T matrices used in the QR factorization,
followed by MIN(M,N)*NB block storing the diagonal block
inverses for the R matrix, followed by work space of size
(ceil(N/32)*32)* MAX(NB, NRHS).
@param[in]
lddb INTEGER
The leading dimension of the array dB. LDDB >= M.
@param[out]
hwork (workspace) COMPLEX array, dimension (LWORK)
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK,
LWORK >= (M - N + NB)*(NRHS + NB) + NRHS*NB,
where NB is the blocksize given by 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 HWORK array, returns
this value as the first entry of the WORK array.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_geqrs
*******************************************************************************/
extern "C" magma_int_t
magma_cgeqrs_gpu(
magma_int_t m, magma_int_t n, magma_int_t nrhs,
magmaFloatComplex_const_ptr dA, magma_int_t ldda,
magmaFloatComplex const *tau,
magmaFloatComplex_ptr dT,
magmaFloatComplex_ptr dB, magma_int_t lddb,
magmaFloatComplex *hwork, magma_int_t lwork,
magma_int_t *info)
{
#define dA(i_,j_) (dA + (i_) + (j_)*ldda)
#define dT(i_) (dT + (lddwork + (i_))*nb)
/* Constants */
const magmaFloatComplex c_zero = MAGMA_C_ZERO;
const magmaFloatComplex c_one = MAGMA_C_ONE;
const magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
const magma_int_t ione = 1;
/* Local variables */
magmaFloatComplex_ptr dwork;
magma_int_t i, min_mn, lddwork, rows, ib;
magma_int_t nb = magma_get_cgeqrf_nb( m, n );
magma_int_t lwkopt = (m - n + nb)*(nrhs + nb) + nrhs*nb;
bool lquery = (lwork == -1);
hwork[0] = magma_cmake_lwork( lwkopt );
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0 || m < n)
*info = -2;
else if (nrhs < 0)
*info = -3;
else if (ldda < max(1,m))
*info = -5;
else if (lddb < max(1,m))
*info = -9;
else if (lwork < lwkopt && ! lquery)
*info = -11;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
min_mn = min(m,n);
if (min_mn == 0) {
hwork[0] = c_one;
return *info;
}
magma_queue_t queue;
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queue );
/* B := Q^H * B */
magma_cunmqr_gpu( MagmaLeft, Magma_ConjTrans,
m, nrhs, n,
dA(0,0), ldda, tau,
dB, lddb, hwork, lwork, dT, nb, info );
if ( *info != 0 ) {
magma_queue_destroy( queue );
return *info;
}
/* Solve R*X = B(1:n,:) */
lddwork= min_mn;
if (nb < min_mn)
dwork = dT+2*lddwork*nb;
else
dwork = dT;
// To do: Why did we have this line originally; seems to be a bug (Stan)?
// dwork = dT;
i = (min_mn - 1)/nb * nb;
ib = n-i;
rows = m-i;
// TODO: this assumes that, on exit from magma_cunmqr_gpu, hwork contains
// the last block of A and B (i.e., C in cunmqr). This should be fixed.
// Seems this data should already be on the GPU, so could switch to
// magma_ctrsm and drop the csetmatrix.
if ( nrhs == 1 ) {
blasf77_ctrsv( MagmaUpperStr, MagmaNoTransStr, MagmaNonUnitStr,
&ib, hwork, &rows,
hwork+rows*ib, &ione);
} else {
blasf77_ctrsm( MagmaLeftStr, MagmaUpperStr, MagmaNoTransStr, MagmaNonUnitStr,
&ib, &nrhs,
&c_one, hwork, &rows,
hwork+rows*ib, &rows);
}
// update the solution vector
magma_csetmatrix( ib, nrhs,
hwork+rows*ib, rows,
dwork+i, lddwork, queue );
// update c
if (nrhs == 1) {
magma_cgemv( MagmaNoTrans, i, ib,
c_neg_one, dA(0, i), ldda,
dwork + i, 1,
c_one, dB, 1, queue );
}
else {
magma_cgemm( MagmaNoTrans, MagmaNoTrans, i, nrhs, ib,
c_neg_one, dA(0, i), ldda,
dwork + i, lddwork,
c_one, dB, lddb, queue );
}
magma_int_t start = i-nb;
if (nb < min_mn) {
for (i = start; i >= 0; i -= nb) {
ib = min(min_mn - i, nb);
rows = m - i;
if (i + ib < n) {
if (nrhs == 1) {
magma_cgemv( MagmaNoTrans, ib, ib,
c_one, dT(i), ib,
dB+i, 1,
c_zero, dwork+i, 1, queue );
magma_cgemv( MagmaNoTrans, i, ib,
c_neg_one, dA(0, i), ldda,
dwork + i, 1,
c_one, dB, 1, queue );
}
else {
magma_cgemm( MagmaNoTrans, MagmaNoTrans, ib, nrhs, ib,
c_one, dT(i), ib,
dB+i, lddb,
c_zero, dwork+i, lddwork, queue );
magma_cgemm( MagmaNoTrans, MagmaNoTrans, i, nrhs, ib,
c_neg_one, dA(0, i), ldda,
dwork + i, lddwork,
c_one, dB, lddb, queue );
}
}
}
}
magma_ccopymatrix( n, nrhs,
dwork, lddwork,
dB, lddb, queue );
magma_queue_destroy( queue );
return *info;
}
/**
Purpose
-------
CHEEVDX_GPU computes selected eigenvalues and, optionally, eigenvectors
of a complex Hermitian matrix A. Eigenvalues and eigenvectors can
be selected by specifying either a range of values or a range of
indices for the desired eigenvalues.
If eigenvectors are desired, it uses a divide and conquer algorithm.
The divide and conquer algorithm makes very mild assumptions about
floating point arithmetic. It will work on machines with a guard
digit in add/subtract, or on those binary machines without guard
digits which subtract like the Cray X-MP, Cray Y-MP, Cray C-90, or
Cray-2. It could conceivably fail on hexadecimal or decimal machines
without guard digits, but we know of none.
Arguments
---------
@param[in]
jobz magma_vec_t
- = MagmaNoVec: Compute eigenvalues only;
- = MagmaVec: Compute eigenvalues and eigenvectors.
@param[in]
range magma_range_t
- = MagmaRangeAll: all eigenvalues will be found.
- = MagmaRangeV: all eigenvalues in the half-open interval (VL,VU]
will be found.
- = MagmaRangeI: the IL-th through IU-th eigenvalues will be found.
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored;
- = MagmaLower: Lower triangle of A is stored.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
dA COMPLEX array on the GPU,
dimension (LDDA, N).
On entry, the Hermitian matrix A. If UPLO = MagmaUpper, the
leading N-by-N upper triangular part of A contains the
upper triangular part of the matrix A. If UPLO = MagmaLower,
the leading N-by-N lower triangular part of A contains
the lower triangular part of the matrix A.
On exit, if JOBZ = MagmaVec, then if INFO = 0, the first m columns
of A contains the required
orthonormal eigenvectors of the matrix A.
If JOBZ = MagmaNoVec, then on exit the lower triangle (if UPLO=MagmaLower)
or the upper triangle (if UPLO=MagmaUpper) of A, including the
diagonal, is destroyed.
@param[in]
ldda INTEGER
The leading dimension of the array DA. LDDA >= max(1,N).
@param[in]
vl REAL
@param[in]
vu REAL
If RANGE=MagmaRangeV, the lower and upper bounds of the interval to
be searched for eigenvalues. VL < VU.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeI.
@param[in]
il INTEGER
@param[in]
iu INTEGER
If RANGE=MagmaRangeI, the indices (in ascending order) of the
smallest and largest eigenvalues to be returned.
1 <= IL <= IU <= N, if N > 0; IL = 1 and IU = 0 if N = 0.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeV.
@param[out]
m INTEGER
The total number of eigenvalues found. 0 <= M <= N.
If RANGE = MagmaRangeAll, M = N, and if RANGE = MagmaRangeI, M = IU-IL+1.
@param[out]
w REAL array, dimension (N)
If INFO = 0, the required m eigenvalues in ascending order.
@param
wA (workspace) COMPLEX array, dimension (LDWA, N)
@param[in]
ldwa INTEGER
The leading dimension of the array wA. LDWA >= max(1,N).
@param[out]
work (workspace) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The length of the array WORK.
If N <= 1, LWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LWORK >= N + N*NB.
If JOBZ = MagmaVec and N > 1, LWORK >= max( N + N*NB, 2*N + N**2 ).
NB can be obtained through magma_get_chetrd_nb(N).
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal sizes of the WORK, RWORK and
IWORK arrays, returns these values as the first entries of
the WORK, RWORK and IWORK arrays, and no error message
related to LWORK or LRWORK or LIWORK is issued by XERBLA.
@param[out]
rwork (workspace) REAL array, dimension (LRWORK)
On exit, if INFO = 0, RWORK[0] returns the optimal LRWORK.
@param[in]
lrwork INTEGER
The dimension of the array RWORK.
If N <= 1, LRWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LRWORK >= N.
If JOBZ = MagmaVec and N > 1, LRWORK >= 1 + 5*N + 2*N**2.
\n
If LRWORK = -1, then a workspace query is assumed; the
routine only calculates the optimal sizes of the WORK, RWORK
and IWORK arrays, returns these values as the first entries
of the WORK, RWORK and IWORK arrays, and no error message
related to LWORK or LRWORK or LIWORK is issued by XERBLA.
@param[out]
iwork (workspace) INTEGER array, dimension (MAX(1,LIWORK))
On exit, if INFO = 0, IWORK[0] returns the optimal LIWORK.
@param[in]
liwork INTEGER
The dimension of the array IWORK.
If N <= 1, LIWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LIWORK >= 1.
If JOBZ = MagmaVec and N > 1, LIWORK >= 3 + 5*N.
\n
If LIWORK = -1, then a workspace query is assumed; the
routine only calculates the optimal sizes of the WORK, RWORK
and IWORK arrays, returns these values as the first entries
of the WORK, RWORK and IWORK arrays, and no error message
related to LWORK or LRWORK or LIWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: if INFO = i and JOBZ = MagmaNoVec, then the algorithm failed
to converge; i off-diagonal elements of an intermediate
tridiagonal form did not converge to zero;
if INFO = i and JOBZ = MagmaVec, then the algorithm failed
to compute an eigenvalue while working on the submatrix
lying in rows and columns INFO/(N+1) through
mod(INFO,N+1).
Further Details
---------------
Based on contributions by
Jeff Rutter, Computer Science Division, University of California
at Berkeley, USA
Modified description of INFO. Sven, 16 Feb 05.
@ingroup magma_cheev_driver
********************************************************************/
extern "C" magma_int_t
magma_cheevdx_gpu(magma_vec_t jobz, magma_range_t range, magma_uplo_t uplo,
magma_int_t n,
magmaFloatComplex *dA, magma_int_t ldda,
float vl, float vu, magma_int_t il, magma_int_t iu,
magma_int_t *m, float *w,
magmaFloatComplex *wA, magma_int_t ldwa,
magmaFloatComplex *work, magma_int_t lwork,
float *rwork, magma_int_t lrwork,
magma_int_t *iwork, magma_int_t liwork,
magma_int_t *info)
{
const char* uplo_ = lapack_uplo_const( uplo );
const char* jobz_ = lapack_vec_const( jobz );
magma_int_t ione = 1;
float d__1;
float eps;
magma_int_t inde;
float anrm;
magma_int_t imax;
float rmin, rmax;
float sigma;
magma_int_t iinfo, lwmin;
magma_int_t lower;
magma_int_t llrwk;
magma_int_t wantz;
magma_int_t indwk2, llwrk2;
magma_int_t iscale;
float safmin;
float bignum;
magma_int_t indtau;
magma_int_t indrwk, indwrk, liwmin;
magma_int_t lrwmin, llwork;
float smlnum;
magma_int_t lquery;
magma_int_t alleig, valeig, indeig;
float *dwork;
magmaFloatComplex *dC;
magma_int_t lddc = ldda;
wantz = (jobz == MagmaVec);
lower = (uplo == MagmaLower);
alleig = (range == MagmaRangeAll);
valeig = (range == MagmaRangeV);
indeig = (range == MagmaRangeI);
lquery = (lwork == -1 || lrwork == -1 || liwork == -1);
*info = 0;
if (! (wantz || (jobz == MagmaNoVec))) {
*info = -1;
} else if (! (alleig || valeig || indeig)) {
*info = -2;
} else if (! (lower || (uplo == MagmaUpper))) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (ldda < max(1,n)) {
*info = -6;
} else if (ldwa < max(1,n)) {
*info = -14;
} else {
if (valeig) {
if (n > 0 && vu <= vl) {
*info = -8;
}
} else if (indeig) {
if (il < 1 || il > max(1,n)) {
*info = -9;
} else if (iu < min(n,il) || iu > n) {
*info = -10;
}
}
}
magma_int_t nb = magma_get_chetrd_nb( n );
if ( n <= 1 ) {
lwmin = 1;
lrwmin = 1;
liwmin = 1;
}
else if ( wantz ) {
lwmin = max( n + n*nb, 2*n + n*n );
lrwmin = 1 + 5*n + 2*n*n;
liwmin = 3 + 5*n;
}
else {
lwmin = n + n*nb;
lrwmin = n;
liwmin = 1;
}
// multiply by 1+eps (in Double!) to ensure length gets rounded up,
// if it cannot be exactly represented in floating point.
real_Double_t one_eps = 1. + lapackf77_slamch("Epsilon");
work[0] = MAGMA_C_MAKE( lwmin * one_eps, 0.);
rwork[0] = lrwmin * one_eps;
iwork[0] = liwmin;
if ((lwork < lwmin) && !lquery) {
*info = -16;
} else if ((lrwork < lrwmin) && ! lquery) {
*info = -18;
} else if ((liwork < liwmin) && ! lquery) {
*info = -20;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info));
return *info;
}
else if (lquery) {
return *info;
}
/* Check if matrix is very small then just call LAPACK on CPU, no need for GPU */
if (n <= 128) {
#ifdef ENABLE_DEBUG
printf("--------------------------------------------------------------\n");
printf(" warning matrix too small N=%d NB=%d, calling lapack on CPU \n", (int) n, (int) nb);
printf("--------------------------------------------------------------\n");
#endif
magmaFloatComplex *A;
magma_cmalloc_cpu( &A, n*n );
magma_cgetmatrix(n, n, dA, ldda, A, n);
lapackf77_cheevd(jobz_, uplo_,
&n, A, &n,
w, work, &lwork,
rwork, &lrwork,
iwork, &liwork, info);
magma_csetmatrix( n, n, A, n, dA, ldda);
magma_free_cpu(A);
*m=n;
return *info;
}
magma_queue_t stream;
magma_queue_create( &stream );
// dC and dwork are never used together, so use one buffer for both;
// unfortunately they're different types (complex and float).
// (this works better in dsyevd_gpu where they're both float).
// n*lddc for chetrd2_gpu, *2 for complex
// n for clanhe
magma_int_t ldwork = n*lddc*2;
if ( wantz ) {
// need 3n^2/2 for cstedx
ldwork = max( ldwork, 3*n*(n/2 + 1) );
}
if (MAGMA_SUCCESS != magma_smalloc( &dwork, ldwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dC = (magmaFloatComplex*) dwork;
/* Get machine constants. */
safmin = lapackf77_slamch("Safe minimum");
eps = lapackf77_slamch("Precision");
smlnum = safmin / eps;
bignum = 1. / smlnum;
rmin = magma_ssqrt(smlnum);
rmax = magma_ssqrt(bignum);
/* Scale matrix to allowable range, if necessary. */
anrm = magmablas_clanhe(MagmaMaxNorm, uplo, n, dA, ldda, dwork);
iscale = 0;
sigma = 1;
if (anrm > 0. && anrm < rmin) {
iscale = 1;
sigma = rmin / anrm;
} else if (anrm > rmax) {
iscale = 1;
sigma = rmax / anrm;
}
if (iscale == 1) {
magmablas_clascl(uplo, 0, 0, 1., sigma, n, n, dA, ldda, info);
}
/* Call CHETRD to reduce Hermitian matrix to tridiagonal form. */
// chetrd rwork: e (n)
// cstedx rwork: e (n) + llrwk (1 + 4*N + 2*N**2) ==> 1 + 5n + 2n^2
inde = 0;
indrwk = inde + n;
llrwk = lrwork - indrwk;
// chetrd work: tau (n) + llwork (n*nb) ==> n + n*nb
// cstedx work: tau (n) + z (n^2)
// cunmtr work: tau (n) + z (n^2) + llwrk2 (n or n*nb) ==> 2n + n^2, or n + n*nb + n^2
indtau = 0;
indwrk = indtau + n;
indwk2 = indwrk + n*n;
llwork = lwork - indwrk;
llwrk2 = lwork - indwk2;
magma_timer_t time=0;
timer_start( time );
#ifdef FAST_HEMV
magma_chetrd2_gpu(uplo, n, dA, ldda, w, &rwork[inde],
&work[indtau], wA, ldwa, &work[indwrk], llwork,
dC, n*lddc, &iinfo);
#else
magma_chetrd_gpu (uplo, n, dA, ldda, w, &rwork[inde],
&work[indtau], wA, ldwa, &work[indwrk], llwork, &iinfo);
#endif
timer_stop( time );
timer_printf( "time chetrd_gpu = %6.2f\n", time );
/* For eigenvalues only, call SSTERF. For eigenvectors, first call
CSTEDC to generate the eigenvector matrix, WORK(INDWRK), of the
tridiagonal matrix, then call CUNMTR to multiply it to the Householder
transformations represented as Householder vectors in A. */
if (! wantz) {
lapackf77_ssterf(&n, w, &rwork[inde], info);
magma_smove_eig(range, n, w, &il, &iu, vl, vu, m);
}
else {
timer_start( time );
magma_cstedx(range, n, vl, vu, il, iu, w, &rwork[inde],
&work[indwrk], n, &rwork[indrwk],
llrwk, iwork, liwork, dwork, info);
timer_stop( time );
timer_printf( "time cstedx = %6.2f\n", time );
timer_start( time );
magma_smove_eig(range, n, w, &il, &iu, vl, vu, m);
magma_csetmatrix( n, *m, &work[indwrk + n * (il-1) ], n, dC, lddc );
magma_cunmtr_gpu(MagmaLeft, uplo, MagmaNoTrans, n, *m, dA, ldda, &work[indtau],
dC, lddc, wA, ldwa, &iinfo);
magma_ccopymatrix( n, *m, dC, lddc, dA, ldda );
timer_stop( time );
timer_printf( "time cunmtr_gpu + copy = %6.2f\n", time );
}
/* If matrix was scaled, then rescale eigenvalues appropriately. */
if (iscale == 1) {
if (*info == 0) {
imax = n;
} else {
imax = *info - 1;
}
d__1 = 1. / sigma;
blasf77_sscal(&imax, &d__1, w, &ione);
}
work[0] = MAGMA_C_MAKE( lwmin * one_eps, 0.); // round up
rwork[0] = lrwmin * one_eps;
iwork[0] = liwmin;
magma_queue_destroy( stream );
magma_free( dwork );
return *info;
} /* magma_cheevdx_gpu */
extern "C" magma_int_t
magma_cgesv_rbt(
magma_bool_t ref, magma_int_t n, magma_int_t nrhs,
magmaFloatComplex *A, magma_int_t lda,
magmaFloatComplex *B, magma_int_t ldb,
magma_int_t *info)
{
/* Function Body */
*info = 0;
if ( ! (ref == MagmaTrue) &&
! (ref == MagmaFalse) ) {
*info = -1;
}
else if (n < 0) {
*info = -2;
} else if (nrhs < 0) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
} else if (ldb < max(1,n)) {
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (nrhs == 0 || n == 0)
return *info;
magma_int_t nn = n + ((4-(n % 4))%4);
magmaFloatComplex *dA, *hu, *hv, *db, *dAo, *dBo, *dwork;
magma_int_t n2;
magma_int_t iter;
n2 = nn*nn;
if (MAGMA_SUCCESS != magma_cmalloc( &dA, n2 )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
if (MAGMA_SUCCESS != magma_cmalloc( &db, nn*nrhs )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
if (ref == MagmaTrue) {
if (MAGMA_SUCCESS != magma_cmalloc( &dAo, n2 )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
if (MAGMA_SUCCESS != magma_cmalloc( &dwork, nn*nrhs )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
if (MAGMA_SUCCESS != magma_cmalloc( &dBo, nn*nrhs )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
}
if (MAGMA_SUCCESS != magma_cmalloc_cpu( &hu, 2*nn )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
if (MAGMA_SUCCESS != magma_cmalloc_cpu( &hv, 2*nn )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magmablas_claset(MagmaFull, nn, nn, MAGMA_C_ZERO, MAGMA_C_ONE, dA, nn);
/* Send matrix on the GPU*/
magma_csetmatrix(n, n, A, lda, dA, nn);
/* Send b on the GPU*/
magma_csetmatrix(n, nrhs, B, ldb, db, nn);
*info = magma_cgerbt_gpu(MagmaTrue, nn, nrhs, dA, nn, db, nn, hu, hv, info);
if (*info != MAGMA_SUCCESS) {
return *info;
}
if (ref == MagmaTrue) {
magma_ccopymatrix(nn, nn, dA, nn, dAo, nn);
magma_ccopymatrix(nn, nrhs, db, nn, dBo, nn);
}
/* Solve the system U^TAV.y = U^T.b on the GPU*/
magma_cgesv_nopiv_gpu( nn, nrhs, dA, nn, db, nn, info);
/* Iterative refinement */
if (ref == MagmaTrue) {
magma_cgerfs_nopiv_gpu(MagmaNoTrans, nn, nrhs, dAo, nn, dBo, nn, db, nn, dwork, dA, &iter, info);
}
//printf("iter = %d\n", iter);
/* The solution of A.x = b is Vy computed on the GPU */
magmaFloatComplex *dv;
if (MAGMA_SUCCESS != magma_cmalloc( &dv, 2*nn )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_csetvector(2*nn, hv, 1, dv, 1);
for(int i = 0; i < nrhs; i++) {
magmablas_cprbt_mv(nn, dv, db+(i*nn));
}
magma_cgetmatrix(n, nrhs, db, nn, B, ldb);
magma_free_cpu( hu);
magma_free_cpu( hv);
magma_free( dA );
magma_free( dv );
magma_free( db );
if (ref == MagmaTrue) {
magma_free( dAo );
magma_free( dBo );
magma_free( dwork );
}
return *info;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cswap, cswapblk, cpermute, claswp, claswpx
*/
int main( int argc, char** argv)
{
TESTING_INIT();
magmaFloatComplex *h_A1, *h_A2;
magmaFloatComplex *d_A1, *d_A2;
magmaFloatComplex *h_R1, *h_R2;
// row-major and column-major performance
real_Double_t row_perf0, col_perf0;
real_Double_t row_perf1, col_perf1;
real_Double_t row_perf2, col_perf2;
real_Double_t row_perf3;
real_Double_t row_perf4;
real_Double_t row_perf5, col_perf5;
real_Double_t row_perf6, col_perf6;
real_Double_t row_perf7;
real_Double_t cpu_perf;
real_Double_t time, gbytes;
magma_int_t N, lda, ldda, nb, j;
magma_int_t ione = 1;
magma_int_t *ipiv, *ipiv2;
magma_int_t *d_ipiv;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_queue_t queue = 0;
printf(" cublasCswap cswap cswapblk claswp cpermute claswp2 claswpx ccopymatrix CPU (all in )\n");
printf(" N nb row-maj/col-maj row-maj/col-maj row-maj/col-maj row-maj row-maj row-maj row-maj/col-maj row-blk/col-blk claswp (GByte/s)\n");
printf("==================================================================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
// For an N x N matrix, swap nb rows or nb columns using various methods.
// Each test is assigned one bit in the 'check' bitmask; bit=1 indicates failure.
// The variable 'shift' keeps track of which bit is for current test
int shift = 1;
int check = 0;
N = opts.nsize[itest];
lda = N;
ldda = ((N+31)/32)*32;
nb = (opts.nb > 0 ? opts.nb : magma_get_cgetrf_nb( N ));
nb = min( N, nb );
// each swap does 2N loads and 2N stores, for nb swaps
gbytes = sizeof(magmaFloatComplex) * 4.*N*nb / 1e9;
TESTING_MALLOC_PIN( h_A1, magmaFloatComplex, lda*N );
TESTING_MALLOC_PIN( h_A2, magmaFloatComplex, lda*N );
TESTING_MALLOC_PIN( h_R1, magmaFloatComplex, lda*N );
TESTING_MALLOC_PIN( h_R2, magmaFloatComplex, lda*N );
TESTING_MALLOC_CPU( ipiv, magma_int_t, nb );
TESTING_MALLOC_CPU( ipiv2, magma_int_t, nb );
TESTING_MALLOC_DEV( d_ipiv, magma_int_t, nb );
TESTING_MALLOC_DEV( d_A1, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( d_A2, magmaFloatComplex, ldda*N );
for( j=0; j < nb; j++ ) {
ipiv[j] = (magma_int_t) ((rand()*1.*N) / (RAND_MAX * 1.)) + 1;
}
/* =====================================================================
* cublasCswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
cublasCswap( N, d_A1+ldda*j, 1, d_A2+ldda*(ipiv[j]-1), 1);
}
}
time = magma_sync_wtime( queue ) - time;
row_perf0 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
cublasCswap( N, d_A1+j, ldda, d_A2+ipiv[j]-1, ldda);
}
}
time = magma_sync_wtime( queue ) - time;
col_perf0 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* =====================================================================
* cswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
magmablas_cswap( N, d_A1+ldda*j, 1, d_A2+ldda*(ipiv[j]-1), 1);
}
}
time = magma_sync_wtime( queue ) - time;
row_perf1 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
magmablas_cswap( N, d_A1+j, ldda, d_A2+ipiv[j]-1, ldda );
}
}
time = magma_sync_wtime( queue ) - time;
col_perf1 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* =====================================================================
* cswapblk, blocked version (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_cswapblk( MagmaRowMajor, N, d_A1, ldda, d_A2, ldda, 1, nb, ipiv, 1, 0);
time = magma_sync_wtime( queue ) - time;
row_perf2 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_cswapblk( MagmaColMajor, N, d_A1, ldda, d_A2, ldda, 1, nb, ipiv, 1, 0);
time = magma_sync_wtime( queue ) - time;
col_perf2 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* =====================================================================
* cpermute_long (1 matrix)
*/
/* Row Major */
memcpy( ipiv2, ipiv, nb*sizeof(magma_int_t) ); // cpermute updates ipiv2
init_matrix( N, N, h_A1, lda, 0 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_cpermute_long2( N, d_A1, ldda, ipiv2, nb, 0 );
time = magma_sync_wtime( queue ) - time;
row_perf3 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style claswp (1 matrix)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_claswp( N, d_A1, ldda, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
row_perf4 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style claswp (1 matrix) - d_ipiv on GPU
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magma_setvector( nb, sizeof(magma_int_t), ipiv, 1, d_ipiv, 1 );
magmablas_claswp2( N, d_A1, ldda, 1, nb, d_ipiv, 1 );
time = magma_sync_wtime( queue ) - time;
row_perf7 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style claswpx (extended for row- and col-major) (1 matrix)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_claswpx( N, d_A1, ldda, 1, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
row_perf5 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* Col Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_claswpx( N, d_A1, 1, ldda, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
col_perf5 = gbytes / time;
time = magma_wtime();
lapackf77_claswp( &N, h_A1, &lda, &ione, &nb, ipiv, &ione);
time = magma_wtime() - time;
cpu_perf = gbytes / time;
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* Copy matrix.
*/
time = magma_sync_wtime( queue );
magma_ccopymatrix( N, nb, d_A1, ldda, d_A2, ldda );
time = magma_sync_wtime( queue ) - time;
// copy reads 1 matrix and writes 1 matrix, so has half gbytes of swap
col_perf6 = 0.5 * gbytes / time;
time = magma_sync_wtime( queue );
magma_ccopymatrix( nb, N, d_A1, ldda, d_A2, ldda );
time = magma_sync_wtime( queue ) - time;
// copy reads 1 matrix and writes 1 matrix, so has half gbytes of swap
row_perf6 = 0.5 * gbytes / time;
printf("%5d %3d %6.2f%c/ %6.2f%c %6.2f%c/ %6.2f%c %6.2f%c/ %6.2f%c %6.2f%c %6.2f%c %6.2f%c %6.2f%c/ %6.2f%c %6.2f / %6.2f %6.2f %10s\n",
(int) N, (int) nb,
row_perf0, ((check & 0x001) != 0 ? '*' : ' '),
col_perf0, ((check & 0x002) != 0 ? '*' : ' '),
row_perf1, ((check & 0x004) != 0 ? '*' : ' '),
col_perf1, ((check & 0x008) != 0 ? '*' : ' '),
row_perf2, ((check & 0x010) != 0 ? '*' : ' '),
col_perf2, ((check & 0x020) != 0 ? '*' : ' '),
row_perf3, ((check & 0x040) != 0 ? '*' : ' '),
row_perf4, ((check & 0x080) != 0 ? '*' : ' '),
row_perf7, ((check & 0x100) != 0 ? '*' : ' '),
row_perf5, ((check & 0x200) != 0 ? '*' : ' '),
col_perf5, ((check & 0x400) != 0 ? '*' : ' '),
row_perf6,
col_perf6,
cpu_perf,
(check == 0 ? "ok" : "* failed") );
status += ! (check == 0);
TESTING_FREE_PIN( h_A1 );
TESTING_FREE_PIN( h_A2 );
TESTING_FREE_PIN( h_R1 );
TESTING_FREE_PIN( h_R2 );
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( ipiv2 );
TESTING_FREE_DEV( d_ipiv );
TESTING_FREE_DEV( d_A1 );
TESTING_FREE_DEV( d_A2 );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/**
@deprecated
Purpose
-------
CLAQPS computes a step of QR factorization with column pivoting
of a complex M-by-N matrix A by using Blas-3. It tries to factorize
NB columns from A starting from the row OFFSET+1, and updates all
of the matrix with Blas-3 xGEMM.
In some cases, due to catastrophic cancellations, it cannot
factorize NB columns. Hence, the actual number of factorized
columns is returned in KB.
Block A(1:OFFSET,1:N) is accordingly pivoted, but not factorized.
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]
offset INTEGER
The number of rows of A that have been factorized in
previous steps.
@param[in]
nb INTEGER
The number of columns to factorize.
@param[out]
kb INTEGER
The number of columns actually factorized.
@param[in,out]
dA COMPLEX array, dimension (LDDA,N), on the GPU.
On entry, the M-by-N matrix A.
On exit, block A(OFFSET+1:M,1:KB) is the triangular
factor obtained and block A(1:OFFSET,1:N) has been
accordingly pivoted, but no factorized.
The rest of the matrix, block A(OFFSET+1:M,KB+1:N) has
been updated.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,M).
@param[in,out]
jpvt INTEGER array, dimension (N)
JPVT(I) = K <==> Column K of the full matrix A has been
permuted into position I in AP.
@param[out]
tau COMPLEX array, dimension (KB)
The scalar factors of the elementary reflectors.
@param[in,out]
vn1 REAL array, dimension (N)
The vector with the partial column norms.
@param[in,out]
vn2 REAL array, dimension (N)
The vector with the exact column norms.
@param[in,out]
dauxv COMPLEX array, dimension (NB), on the GPU
Auxiliary vector.
@param[in,out]
dF COMPLEX array, dimension (LDDF,NB), on the GPU
Matrix F' = L*Y'*A.
@param[in]
lddf INTEGER
The leading dimension of the array F. LDDF >= max(1,N).
@ingroup magma_cgeqp3_aux
********************************************************************/
extern "C" magma_int_t
magma_claqps_gpu(
magma_int_t m, magma_int_t n, magma_int_t offset,
magma_int_t nb, magma_int_t *kb,
magmaFloatComplex_ptr dA, magma_int_t ldda,
magma_int_t *jpvt, magmaFloatComplex *tau,
float *vn1, float *vn2,
magmaFloatComplex_ptr dauxv,
magmaFloatComplex_ptr dF, magma_int_t lddf)
{
#define dA(i, j) (dA + (i) + (j)*(ldda))
#define dF(i, j) (dF + (i) + (j)*(lddf))
magmaFloatComplex c_zero = MAGMA_C_MAKE( 0.,0.);
magmaFloatComplex c_one = MAGMA_C_MAKE( 1.,0.);
magmaFloatComplex c_neg_one = MAGMA_C_MAKE(-1.,0.);
magma_int_t ione = 1;
magma_int_t i__1, i__2;
magmaFloatComplex z__1;
magma_int_t k, rk;
magmaFloatComplex_ptr dAks;
magmaFloatComplex tauk = MAGMA_C_ZERO;
magma_int_t pvt;
float tol3z;
magma_int_t itemp;
float lsticc;
magmaFloat_ptr dlsticcs;
magma_smalloc( &dlsticcs, 1+256*(n+255)/256 );
tol3z = magma_ssqrt( lapackf77_slamch("Epsilon"));
lsticc = 0;
k = 0;
magma_cmalloc( &dAks, nb );
magma_queue_t queue;
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queue );
while( k < nb && lsticc == 0 ) {
rk = offset + k;
/* Determine ith pivot column and swap if necessary */
// subtract 1 from Fortran/CUBLAS isamax; pvt, k are 0-based.
pvt = k + magma_isamax( n-k, &vn1[k], ione, queue ) - 1;
if (pvt != k) {
/* F gets swapped so F must be sent at the end to GPU */
i__1 = k;
magmablas_cswap( m, dA(0, pvt), ione, dA(0, k), ione, queue );
magmablas_cswap( i__1, dF(pvt, 0), lddf, dF(k, 0), lddf, queue );
itemp = jpvt[pvt];
jpvt[pvt] = jpvt[k];
jpvt[k] = itemp;
magma_sswap( 2, &vn1[pvt], n+offset, &vn1[k], n+offset, queue );
}
/* Apply previous Householder reflectors to column K:
A(RK:M,K) := A(RK:M,K) - A(RK:M,1:K-1)*F(K,1:K-1)'.
Optimization: multiply with beta=0; wait for vector and subtract */
if (k > 0) {
//#define RIGHT_UPDATE
#ifdef RIGHT_UPDATE
i__1 = m - offset - nb;
i__2 = k;
magma_cgemv( MagmaNoTrans, i__1, i__2,
c_neg_one, A(offset+nb, 0), lda,
F(k, 0), ldf,
c_one, A(offset+nb, k), ione, queue );
#else
i__1 = m - rk;
i__2 = k;
magma_cgemv( MagmaNoTrans, i__1, i__2,
c_neg_one, dA(rk, 0), ldda,
dF(k, 0), lddf,
c_one, dA(rk, k), ione, queue );
#endif
}
/* Generate elementary reflector H(k). */
magma_clarfg_gpu( m-rk, dA(rk, k), dA(rk + 1, k), &tau[k], &vn1[k], &dAks[k], queue );
/* needed to avoid the race condition */
if (k == 0) magma_csetvector( 1, &c_one, 1, dA(rk, k), 1, queue );
else magma_ccopymatrix( 1, 1, dA(offset, 0), 1, dA(rk, k), 1, queue );
/* Compute Kth column of F:
Compute F(K+1:N,K) := tau(K)*A(RK:M,K+1:N)'*A(RK:M,K) on the GPU */
if (k < n-1 || k > 0) magma_cgetvector( 1, &tau[k], 1, &tauk, 1, queue );
if (k < n-1) {
i__1 = m - rk;
i__2 = n - k - 1;
/* Multiply on GPU */
magma_cgemv( MagmaConjTrans, m-rk, n-k-1,
tauk, dA( rk, k+1 ), ldda,
dA( rk, k ), 1,
c_zero, dF( k+1, k ), 1, queue );
}
/* Incremental updating of F:
F(1:N,K) := F(1:N,K) - tau(K)*F(1:N,1:K-1)*A(RK:M,1:K-1)'*A(RK:M,K).
F(1:N,K) := tau(K)*A(RK:M,K+1:N)'*A(RK:M,K) - tau(K)*F(1:N,1:K-1)*A(RK:M,1:K-1)'*A(RK:M,K)
:= tau(K)(A(RK:M,K+1:N)' - F(1:N,1:K-1)*A(RK:M,1:K-1)') A(RK:M,K)
so, F is (updated A)*V */
if (k > 0) {
z__1 = MAGMA_C_NEGATE( tauk );
#ifdef RIGHT_UPDATE
i__1 = m - offset - nb;
i__2 = k;
magma_cgemv( MagmaConjTrans, i__1, i__2,
z__1, dA(offset+nb, 0), lda,
dA(offset+nb, k), ione,
c_zero, dauxv, ione, queue );
i__1 = k;
magma_cgemv( MagmaNoTrans, n-k-1, i__1,
c_one, F(k+1,0), ldf,
dauxv, ione,
c_one, F(k+1,k), ione, queue );
#else
i__1 = m - rk;
i__2 = k;
magma_cgemv( MagmaConjTrans, i__1, i__2,
z__1, dA(rk, 0), ldda,
dA(rk, k), ione,
c_zero, dauxv, ione, queue );
/* I think we only need stricly lower-triangular part :) */
magma_cgemv( MagmaNoTrans, n-k-1, i__2,
c_one, dF(k+1,0), lddf,
dauxv, ione,
c_one, dF(k+1,k), ione, queue );
#endif
}
/* Optimization: On the last iteration start sending F back to the GPU */
/* Update the current row of A:
A(RK,K+1:N) := A(RK,K+1:N) - A(RK,1:K)*F(K+1:N,1:K)'. */
if (k < n-1) {
i__1 = n - k - 1;
i__2 = k + 1;
#ifdef RIGHT_UPDATE
/* right-looking update of rows, */
magma_cgemm( MagmaNoTrans, MagmaConjTrans, nb-k, i__1, ione,
c_neg_one, dA(rk, k ), ldda,
dF(k+1, k ), lddf,
c_one, dA(rk, k+1), ldda, queue );
#else
/* left-looking update of rows, *
* since F=A'v with original A, so no right-looking */
magma_cgemm( MagmaNoTrans, MagmaConjTrans, ione, i__1, i__2,
c_neg_one, dA(rk, 0 ), ldda,
dF(k+1,0 ), lddf,
c_one, dA(rk, k+1), ldda, queue );
#endif
}
/* Update partial column norms. */
if (rk < min(m, n+offset)-1 ) {
magmablas_scnrm2_row_check_adjust( n-k-1, tol3z, &vn1[k+1], &vn2[k+1],
dA(rk,k+1), ldda, dlsticcs, queue );
//magma_device_sync();
magma_sgetvector( 1, &dlsticcs[0], 1, &lsticc, 1, queue );
}
++k;
}
magma_ccopymatrix( 1, k, dAks, 1, dA(offset, 0), ldda+1, queue );
// leave k as the last column done
--k;
*kb = k + 1;
rk = offset + *kb - 1;
/* Apply the block reflector to the rest of the matrix:
A(OFFSET+KB+1:M,KB+1:N) := A(OFFSET+KB+1:M,KB+1:N) - A(OFFSET+KB+1:M,1:KB)*F(KB+1:N,1:KB)' */
if (*kb < min(n, m - offset)) {
i__1 = m - rk - 1;
i__2 = n - *kb;
magma_cgemm( MagmaNoTrans, MagmaConjTrans, i__1, i__2, *kb,
c_neg_one, dA(rk+1, 0 ), ldda,
dF(*kb, 0 ), lddf,
c_one, dA(rk+1, *kb), ldda, queue );
}
/* Recomputation of difficult columns. */
if ( lsticc > 0 ) {
// printf( " -- recompute dnorms --\n" );
magmablas_scnrm2_check( m-rk-1, n-*kb, dA(rk+1,*kb), ldda,
&vn1[*kb], dlsticcs, queue );
magma_scopymatrix( n-*kb, 1, &vn1[*kb], *kb, &vn2[*kb], *kb, queue );
}
magma_free( dAks );
magma_free( dlsticcs );
magma_queue_destroy( queue );
return MAGMA_SUCCESS;
} /* magma_claqps */
extern "C" magma_int_t
magma_claqps_gpu(magma_int_t m, magma_int_t n, magma_int_t offset,
magma_int_t nb, magma_int_t *kb,
magmaFloatComplex *A, magma_int_t lda,
magma_int_t *jpvt, magmaFloatComplex *tau,
float *vn1, float *vn2,
magmaFloatComplex *auxv,
magmaFloatComplex *F, magma_int_t ldf)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CLAQPS computes a step of QR factorization with column pivoting
of a complex M-by-N matrix A by using Blas-3. It tries to factorize
NB columns from A starting from the row OFFSET+1, and updates all
of the matrix with Blas-3 xGEMM.
In some cases, due to catastrophic cancellations, it cannot
factorize NB columns. Hence, the actual number of factorized
columns is returned in KB.
Block A(1:OFFSET,1:N) is accordingly pivoted, but not factorized.
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
OFFSET (input) INTEGER
The number of rows of A that have been factorized in
previous steps.
NB (input) INTEGER
The number of columns to factorize.
KB (output) INTEGER
The number of columns actually factorized.
A (input/output) COMPLEX*16 array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, block A(OFFSET+1:M,1:KB) is the triangular
factor obtained and block A(1:OFFSET,1:N) has been
accordingly pivoted, but no factorized.
The rest of the matrix, block A(OFFSET+1:M,KB+1:N) has
been updated.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,M).
JPVT (input/output) INTEGER array, dimension (N)
JPVT(I) = K <==> Column K of the full matrix A has been
permuted into position I in AP.
TAU (output) COMPLEX*16 array, dimension (KB)
The scalar factors of the elementary reflectors.
VN1 (input/output) DOUBLE PRECISION array, dimension (N)
The vector with the partial column norms.
VN2 (input/output) DOUBLE PRECISION array, dimension (N)
The vector with the exact column norms.
AUXV (input/output) COMPLEX*16 array, dimension (NB)
Auxiliar vector.
F (input/output) COMPLEX*16 array, dimension (LDF,NB)
Matrix F' = L*Y'*A.
LDF (input) INTEGER
The leading dimension of the array F. LDF >= max(1,N).
===================================================================== */
#define A(i, j) (A + (i) + (j)*(lda ))
#define F(i, j) (F + (i) + (j)*(ldf ))
magmaFloatComplex c_zero = MAGMA_C_MAKE( 0.,0.);
magmaFloatComplex c_one = MAGMA_C_MAKE( 1.,0.);
magmaFloatComplex c_neg_one = MAGMA_C_MAKE(-1.,0.);
magma_int_t ione = 1;
magma_int_t i__1, i__2;
//float d__1;
magmaFloatComplex z__1;
//magma_int_t j;
magma_int_t k, rk;
//magmaFloatComplex Akk;
magmaFloatComplex *Aks;
magmaFloatComplex tauk;
magma_int_t pvt;
//float temp, temp2;
float tol3z;
magma_int_t itemp;
float lsticc, *lsticcs;
magma_int_t lastrk;
magma_smalloc( &lsticcs, 1+256*(n+255)/256 );
lastrk = min( m, n + offset );
tol3z = magma_ssqrt( lapackf77_slamch("Epsilon"));
lsticc = 0;
k = 0;
magma_cmalloc( &Aks, nb );
while( k < nb && lsticc == 0 ) {
rk = offset + k;
/* Determine ith pivot column and swap if necessary */
// Fortran: pvt, k, isamax are all 1-based; subtract 1 from k.
// C: pvt, k, isamax are all 0-based; don't subtract 1.
pvt = k - 1 + magma_isamax( n-k, &vn1[k], ione );
if (pvt != k) {
/*if (pvt >= nb) {
// 1. Start copy from GPU
magma_cgetmatrix_async( m - offset - nb, 1,
dA(offset + nb, pvt), ldda,
A (offset + nb, pvt), lda, stream );
}*/
/* F gets swapped so F must be sent at the end to GPU */
i__1 = k;
/*if (pvt < nb){
// no need of transfer if pivot is within the panel
blasf77_cswap( &m, A(0, pvt), &ione, A(0, k), &ione );
}
else {
// 1. Finish copy from GPU
magma_queue_sync( stream );
// 2. Swap as usual on CPU
blasf77_cswap(&m, A(0, pvt), &ione, A(0, k), &ione);
// 3. Restore the GPU
magma_csetmatrix_async( m - offset - nb, 1,
A (offset + nb, pvt), lda,
dA(offset + nb, pvt), ldda, stream);
}*/
magmablas_cswap( m, A(0, pvt), ione, A(0, k), ione );
//blasf77_cswap( &i__1, F(pvt,0), &ldf, F(k,0), &ldf );
magmablas_cswap( i__1, F(pvt, 0), ldf, F(k, 0), ldf);
itemp = jpvt[pvt];
jpvt[pvt] = jpvt[k];
jpvt[k] = itemp;
//vn1[pvt] = vn1[k];
//vn2[pvt] = vn2[k];
#if defined(PRECISION_d) || defined(PRECISION_z)
//magma_dswap( 1, &vn1[pvt], 1, &vn1[k], 1 );
//magma_dswap( 1, &vn2[pvt], 1, &vn2[k], 1 );
magma_dswap( 2, &vn1[pvt], n+offset, &vn1[k], n+offset );
#else
//magma_sswap( 1, &vn1[pvt], 1, &vn1[k], 1 );
//magma_sswap( 1, &vn2[pvt], 1, &vn2[k], 1 );
magma_sswap(2, &vn1[pvt], n+offset, &vn1[k], n+offset);
#endif
}
/* Apply previous Householder reflectors to column K:
A(RK:M,K) := A(RK:M,K) - A(RK:M,1:K-1)*F(K,1:K-1)'.
Optimization: multiply with beta=0; wait for vector and subtract */
if (k > 0) {
/*#if (defined(PRECISION_c) || defined(PRECISION_z))
for (j = 0; j < k; ++j){
*F(k,j) = MAGMA_C_CNJG( *F(k,j) );
}
#endif*/
//#define RIGHT_UPDATE
#ifdef RIGHT_UPDATE
i__1 = m - offset - nb;
i__2 = k;
magma_cgemv( MagmaNoTrans, i__1, i__2,
c_neg_one, A(offset+nb, 0), lda,
F(k, 0), ldf,
c_one, A(offset+nb, k), ione );
#else
i__1 = m - rk;
i__2 = k;
/*blasf77_cgemv( MagmaNoTransStr, &i__1, &i__2,
&c_neg_one, A(rk, 0), &lda,
F(k, 0), &ldf,
&c_one, A(rk, k), &ione );*/
magma_cgemv( MagmaNoTrans, i__1, i__2,
c_neg_one, A(rk, 0), lda,
F(k, 0), ldf,
c_one, A(rk, k), ione );
#endif
/*#if (defined(PRECISION_c) || defined(PRECISION_z))
for (j = 0; j < k; ++j) {
*F(k,j) = MAGMA_C_CNJG( *F(k,j) );
}
#endif*/
}
/* Generate elementary reflector H(k). */
magma_clarfg_gpu(m-rk, A(rk, k), A(rk + 1, k), &tau[k], &vn1[k], &Aks[k]);
//Akk = *A(rk, k);
//*A(rk, k) = c_one;
//magma_cgetvector( 1, &Aks[k], 1, &Akk, 1 );
/* needed to avoid the race condition */
if (k == 0) magma_csetvector( 1, &c_one, 1, A(rk, k), 1 );
else magma_ccopymatrix( 1, 1, A(offset, 0), 1, A(rk, k), 1 );
/* Compute Kth column of F:
Compute F(K+1:N,K) := tau(K)*A(RK:M,K+1:N)'*A(RK:M,K) on the GPU */
if (k < n-1 || k > 0) magma_cgetvector( 1, &tau[k], 1, &tauk, 1 );
if (k < n-1) {
i__1 = m - rk;
i__2 = n - k - 1;
/* Send the vector to the GPU */
//magma_csetmatrix( i__1, 1, A(rk, k), lda, dA(rk,k), ldda );
/* Multiply on GPU */
// was CALL CGEMV( 'Conjugate transpose', M-RK+1, N-K,
// TAU( K ), A( RK, K+1 ), LDA,
// A( RK, K ), 1,
// CZERO, F( K+1, K ), 1 )
//magma_cgetvector( 1, &tau[k], 1, &tauk, 1 );
magma_cgemv( MagmaConjTrans, m-rk, n-k-1,
tauk, A( rk, k+1 ), lda,
A( rk, k ), 1,
c_zero, F( k+1, k ), 1 );
//magma_cscal( m-rk, tau[k], F( k+1, k), 1 );
//magma_int_t i__3 = nb-k-1;
//magma_int_t i__4 = i__2 - i__3;
//magma_int_t i__5 = nb-k;
//magma_cgemv( MagmaConjTrans, i__1 - i__5, i__2 - i__3,
// tau[k], dA(rk +i__5, k+1+i__3), ldda,
// dA(rk +i__5, k ), ione,
// c_zero, dF(k+1+i__3, k ), ione );
//magma_cgetmatrix_async( i__2-i__3, 1,
// dF(k + 1 +i__3, k), i__2,
// F (k + 1 +i__3, k), i__2, stream );
//blasf77_cgemv( MagmaConjTransStr, &i__1, &i__3,
// &tau[k], A(rk, k+1), &lda,
// A(rk, k ), &ione,
// &c_zero, F(k+1, k ), &ione );
//magma_queue_sync( stream );
//blasf77_cgemv( MagmaConjTransStr, &i__5, &i__4,
// &tau[k], A(rk, k+1+i__3), &lda,
// A(rk, k ), &ione,
// &c_one, F(k+1+i__3, k ), &ione );
}
/* Padding F(1:K,K) with zeros.
for (j = 0; j <= k; ++j) {
magma_csetvector( 1, &c_zero, 1, F(j, k), 1 );
}*/
/* Incremental updating of F:
F(1:N,K) := F(1:N,K) - tau(K)*F(1:N,1:K-1)*A(RK:M,1:K-1)'*A(RK:M,K).
F(1:N,K) := tau(K)*A(RK:M,K+1:N)'*A(RK:M,K) - tau(K)*F(1:N,1:K-1)*A(RK:M,1:K-1)'*A(RK:M,K)
:= tau(K)(A(RK:M,K+1:N)' - F(1:N,1:K-1)*A(RK:M,1:K-1)') A(RK:M,K)
so, F is (updated A)*V */
//if (k > 0 && k<n-1) {
if (k > 0) {
//magma_cgetvector( 1, &tau[k], 1, &tauk, 1 );
z__1 = MAGMA_C_NEGATE( tauk );
#ifdef RIGHT_UPDATE
i__1 = m - offset - nb;
i__2 = k;
magma_cgemv( MagmaConjTrans, i__1, i__2,
z__1, A(offset+nb, 0), lda,
A(offset+nb, k), ione,
c_zero, auxv, ione );
i__1 = k;
magma_cgemv( MagmaNoTrans, n-k-1, i__1,
c_one, F(k+1,0), ldf,
auxv, ione,
c_one, F(k+1,k), ione );
#else
i__1 = m - rk;
i__2 = k;
//blasf77_cgemv( MagmaConjTransStr, &i__1, &i__2,
// &z__1, A(rk, 0), &lda,
// A(rk, k), &ione,
// &c_zero, auxv, &ione );
magma_cgemv( MagmaConjTrans, i__1, i__2,
z__1, A(rk, 0), lda,
A(rk, k), ione,
c_zero, auxv, ione );
//i__1 = k;
//blasf77_cgemv( MagmaNoTransStr, &n, &i__1,
// &c_one, F(0,0), &ldf,
// auxv, &ione,
// &c_one, F(0,k), &ione );
/*magma_cgemv( MagmaNoTrans, n, i__1,
c_one, F(0,0), ldf,
auxv, ione,
c_one, F(0,k), ione );*/
/* I think we only need stricly lower-triangular part :) */
magma_cgemv( MagmaNoTrans, n-k-1, i__2,
c_one, F(k+1,0), ldf,
auxv, ione,
c_one, F(k+1,k), ione );
#endif
}
/* Optimization: On the last iteration start sending F back to the GPU */
/* Update the current row of A:
A(RK,K+1:N) := A(RK,K+1:N) - A(RK,1:K)*F(K+1:N,1:K)'. */
if (k < n-1) {
i__1 = n - k - 1;
i__2 = k + 1;
//blasf77_cgemm( MagmaNoTransStr, MagmaConjTransStr, &ione, &i__1, &i__2,
// &c_neg_one, A(rk, 0 ), &lda,
// F(k+1,0 ), &ldf,
// &c_one, A(rk, k+1), &lda );
#ifdef RIGHT_UPDATE
/* right-looking update of rows, */
magma_cgemm( MagmaNoTrans, MagmaConjTrans, nb-k, i__1, ione,
c_neg_one, A(rk, k ), lda,
F(k+1, k ), ldf,
c_one, A(rk, k+1), lda );
#else
/* left-looking update of rows, *
* since F=A'v with original A, so no right-looking */
magma_cgemm( MagmaNoTrans, MagmaConjTrans, ione, i__1, i__2,
c_neg_one, A(rk, 0 ), lda,
F(k+1,0 ), ldf,
c_one, A(rk, k+1), lda );
#endif
}
/* Update partial column norms. */
if (rk < min(m, n+offset)-1 ){
magmablas_scnrm2_row_check_adjust(n-k-1, tol3z, &vn1[k+1], &vn2[k+1], A(rk,k+1), lda, lsticcs);
magma_device_sync();
#if defined(PRECISION_d) || defined(PRECISION_z)
magma_dgetvector( 1, &lsticcs[0], 1, &lsticc, 1 );
#else
magma_sgetvector( 1, &lsticcs[0], 1, &lsticc, 1 );
#endif
}
/*if (rk < lastrk) {
for (j = k + 1; j < n; ++j) {
if (vn1[j] != 0.) {
// NOTE: The following 4 lines follow from the analysis in
// Lapack Working Note 176.
temp = MAGMA_C_ABS( *A(rk,j) ) / vn1[j];
temp = max( 0., ((1. + temp) * (1. - temp)) );
d__1 = vn1[j] / vn2[j];
temp2 = temp * (d__1 * d__1);
if (temp2 <= tol3z) {
vn2[j] = (float) lsticc;
lsticc = j;
} else {
vn1[j] *= magma_ssqrt(temp);
}
}
}
}*/
//*A(rk, k) = Akk;
//magma_csetvector( 1, &Akk, 1, A(rk, k), 1 );
//magma_cswap( 1, &Aks[k], 1, A(rk, k), 1 );
++k;
}
magma_ccopymatrix( 1, k, Aks, 1, A(offset, 0), lda+1 );
// leave k as the last column done
--k;
*kb = k + 1;
rk = offset + *kb - 1;
/* Apply the block reflector to the rest of the matrix:
A(OFFSET+KB+1:M,KB+1:N) := A(OFFSET+KB+1:M,KB+1:N) - A(OFFSET+KB+1:M,1:KB)*F(KB+1:N,1:KB)' */
if (*kb < min(n, m - offset)) {
i__1 = m - rk - 1;
i__2 = n - *kb;
/* Send F to the GPU
magma_csetmatrix( i__2, *kb,
F (*kb, 0), ldf,
dF(*kb, 0), i__2 );*/
magma_cgemm( MagmaNoTrans, MagmaConjTrans, i__1, i__2, *kb,
c_neg_one, A(rk+1, 0 ), lda,
F(*kb, 0 ), ldf,
c_one, A(rk+1, *kb), lda );
}
/* Recomputation of difficult columns. */
if( lsticc > 0 ) {
printf( " -- recompute dnorms --\n" );
magmablas_scnrm2_check(m-rk-1, n-*kb, A(rk+1,*kb), lda,
&vn1[*kb], lsticcs);
#if defined(PRECISION_d) || defined(PRECISION_z)
magma_dcopymatrix( n-*kb, 1, &vn1[*kb], *kb, &vn2[*kb], *kb);
#else
magma_scopymatrix( n-*kb, 1, &vn1[*kb], *kb, &vn2[*kb], *kb);
#endif
/*while( lsticc > 0 ) {
itemp = (magma_int_t)(vn2[lsticc] >= 0. ? floor(vn2[lsticc] + .5) : -floor(.5 - vn2[lsticc]));
i__1 = m - rk - 1;
if (lsticc <= nb)
vn1[lsticc] = cblas_scnrm2(i__1, A(rk + 1, lsticc), ione);
else {
// Where is the data, CPU or GPU ?
float r1, r2;
r1 = cblas_scnrm2(nb-k, A(rk + 1, lsticc), ione);
r2 = magma_scnrm2(m-offset-nb, dA(offset + nb + 1, lsticc), ione);
vn1[lsticc] = magma_ssqrt(r1*r1+r2*r2);
}
// NOTE: The computation of VN1( LSTICC ) relies on the fact that
// SNRM2 does not fail on vectors with norm below the value of SQRT(SLAMCH('S'))
vn2[lsticc] = vn1[lsticc];
lsticc = itemp;*/
}
magma_free(Aks);
magma_free(lsticcs);
return MAGMA_SUCCESS;
} /* magma_claqps */
extern "C" magma_err_t
magma_cgeqrs_gpu(magma_int_t m, magma_int_t n, magma_int_t nrhs,
magmaFloatComplex_ptr dA, size_t dA_offset, magma_int_t ldda,
magmaFloatComplex *tau, magmaFloatComplex_ptr dT, size_t dT_offset,
magmaFloatComplex_ptr dB, size_t dB_offset, magma_int_t lddb,
magmaFloatComplex *hwork, magma_int_t lwork,
magma_int_t *info, magma_queue_t queue)
{
/* -- clMagma (version 0.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date January 2014
Purpose
=======
Solves the least squares problem
min || A*X - C ||
using the QR factorization A = Q*R computed by CGEQRF_GPU.
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. M >= N >= 0.
NRHS (input) INTEGER
The number of columns of the matrix C. NRHS >= 0.
A (input) COMPLEX array on the GPU, dimension (LDDA,N)
The i-th column must contain the vector which defines the
elementary reflector H(i), for i = 1,2,...,n, as returned by
CGEQRF_GPU in the first n columns of its array argument A.
LDDA (input) INTEGER
The leading dimension of the array A, LDDA >= M.
TAU (input) COMPLEX array, dimension (N)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by MAGMA_CGEQRF_GPU.
DB (input/output) COMPLEX array on the GPU, dimension (LDDB,NRHS)
On entry, the M-by-NRHS matrix C.
On exit, the N-by-NRHS solution matrix X.
DT (input) COMPLEX array that is the output (the 6th argument)
of magma_cgeqrf_gpu of size
2*MIN(M, N)*NB + ((N+31)/32*32 )* MAX(NB, NRHS).
The array starts with a block of size MIN(M,N)*NB that stores
the triangular T matrices used in the QR factorization,
followed by MIN(M,N)*NB block storing the diagonal block
inverses for the R matrix, followed by work space of size
((N+31)/32*32 )* MAX(NB, NRHS).
LDDB (input) INTEGER
The leading dimension of the array DB. LDDB >= M.
HWORK (workspace/output) COMPLEX array, dimension (LWORK)
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The dimension of the array WORK, LWORK >= max(1,NRHS).
For optimum performance LWORK >= (M-N+NB)*(NRHS + 2*NB), where
NB is the blocksize given by magma_get_cgeqrf_nb( M ).
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the HWORK array, returns
this value as the first entry of the WORK array.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
===================================================================== */
#define a_ref(a_1,a_2) dA, (dA_offset + (a_1) + (a_2)*(ldda))
#define d_ref(a_1) dT, (dT_offset + (lddwork+(a_1))*nb)
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex_ptr dwork;
magma_int_t i, k, lddwork, rows, ib;
magma_int_t ione = 1;
magma_int_t nb = magma_get_cgeqrf_nb(m);
magma_int_t lwkopt = (m-n+nb)*(nrhs+2*nb);
long int lquery = (lwork == -1);
hwork[0] = MAGMA_C_MAKE( (float)lwkopt, 0. );
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0 || m < n)
*info = -2;
else if (nrhs < 0)
*info = -3;
else if (ldda < max(1,m))
*info = -5;
else if (lddb < max(1,m))
*info = -8;
else if (lwork < lwkopt && ! lquery)
*info = -10;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
k = min(m,n);
if (k == 0) {
hwork[0] = c_one;
return *info;
}
/* B := Q' * B */
magma_cunmqr_gpu( MagmaLeft, MagmaConjTrans,
m, nrhs, n,
a_ref(0,0), ldda, tau,
dB, dB_offset, lddb, hwork, lwork, dT, dT_offset, nb, info, queue );
if ( *info != 0 ) {
return *info;
}
/* Solve R*X = B(1:n,:) */
lddwork= k;
int ldtwork;
size_t dwork_offset = 0;
if (nb < k)
{
dwork = dT;
dwork_offset = dT_offset+2*lddwork*nb;
}
else
{
ldtwork = ( 2*k + ((n+31)/32)*32 )*nb;
magma_cmalloc( &dwork, ldtwork );
}
// To do: Why did we have this line originally; seems to be a bug (Stan)?
//dwork = dT;
i = (k-1)/nb * nb;
ib = n-i;
rows = m-i;
if ( nrhs == 1 ) {
blasf77_ctrsv( MagmaUpperStr, MagmaNoTransStr, MagmaNonUnitStr,
&ib, hwork, &rows,
hwork+rows*ib, &ione);
} else {
blasf77_ctrsm( MagmaLeftStr, MagmaUpperStr, MagmaNoTransStr, MagmaNonUnitStr,
&ib, &nrhs,
&c_one, hwork, &rows,
hwork+rows*ib, &rows);
}
// update the solution vector
magma_csetmatrix( ib, nrhs, hwork+rows*ib, 0, rows, dwork, dwork_offset+i, lddwork, queue );
// update c
if (nrhs == 1)
magma_cgemv( MagmaNoTrans, i, ib,
c_neg_one, a_ref(0, i), ldda,
dwork, dwork_offset+i, 1,
c_one, dB, dB_offset, 1, queue );
else
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
i, nrhs, ib,
c_neg_one, a_ref(0, i), ldda,
dwork, dwork_offset + i, lddwork,
c_one, dB, dB_offset, lddb, queue );
int start = i-nb;
if (nb < k) {
for (i = start; i >=0; i -= nb) {
ib = min(k-i, nb);
rows = m -i;
if (i + ib < n) {
if (nrhs == 1) {
magma_cgemv( MagmaNoTrans, ib, ib,
c_one, d_ref(i), ib,
dB, dB_offset+i, 1,
c_zero, dwork, dwork_offset+i, 1, queue );
magma_cgemv( MagmaNoTrans, i, ib,
c_neg_one, a_ref(0, i), ldda,
dwork, dwork_offset+i, 1,
c_one, dB, dB_offset, 1, queue );
} else {
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
ib, nrhs, ib,
c_one, d_ref(i), ib,
dB, dB_offset+i, lddb,
c_zero, dwork, dwork_offset+i, lddwork, queue );
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
i, nrhs, ib,
c_neg_one, a_ref(0, i), ldda,
dwork, dwork_offset+i, lddwork,
c_one, dB, dB_offset, lddb, queue );
}
}
}
}
magma_ccopymatrix( (n), nrhs,
dwork, dwork_offset, lddwork,
dB, dB_offset, lddb, queue );
if (nb >= k)
magma_free(dwork);
magma_queue_sync( queue );
return *info;
}
/**
Purpose
=======
CLAHEF computes a partial factorization of a complex Hermitian
matrix A using the Bunch-Kaufman diagonal pivoting method. The
partial factorization has the form:
A = ( I U12 ) ( A11 0 ) ( I 0 ) if UPLO = 'U', or:
( 0 U22 ) ( 0 D ) ( U12' U22' )
A = ( L11 0 ) ( D 0 ) ( L11' L21' ) if UPLO = 'L'
( L21 I ) ( 0 A22 ) ( 0 I )
where the order of D is at most NB. The actual order is returned in
the argument KB, and is either NB or NB-1, or N if N <= NB.
Note that U' denotes the conjugate transpose of U.
CLAHEF is an auxiliary routine called by CHETRF. It uses blocked code
(calling Level 3 BLAS) to update the submatrix A11 (if UPLO = 'U') or
A22 (if UPLO = 'L').
Arguments
---------
@param[in]
UPLO CHARACTER
Specifies whether the upper or lower triangular part of the
Hermitian matrix A is stored:
- = 'U': Upper triangular
- = 'L': Lower triangular
@param[in]
N INTEGER
The order of the matrix A. N >= 0.
@param[in]
NB INTEGER
The maximum number of columns of the matrix A that should be
factored. NB should be at least 2 to allow for 2-by-2 pivot
blocks.
@param[out]
KB INTEGER
The number of columns of A that were actually factored.
KB is either NB-1 or NB, or N if N <= NB.
@param[in,out]
A COMPLEX array, dimension (LDA,N)
On entry, the Hermitian matrix A. If UPLO = 'U', the leading
n-by-n upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = 'L', the
leading n-by-n lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit, A contains details of the partial factorization.
@param[in]
LDA INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
ipiv INTEGER array, dimension (N)
Details of the interchanges and the block structure of D.
If UPLO = 'U', only the last KB elements of ipiv are set;
if UPLO = 'L', only the first KB elements are set.
\n
If ipiv(k) > 0, then rows and columns k and ipiv(k) were
interchanged and D(k,k) is a 1-by-1 diagonal block.
If UPLO = 'U' and ipiv(k) = ipiv(k-1) < 0, then rows and
columns k-1 and -ipiv(k) were interchanged and D(k-1:k,k-1:k)
is a 2-by-2 diagonal block. If UPLO = 'L' and ipiv(k) =
ipiv(k+1) < 0, then rows and columns k+1 and -ipiv(k) were
interchanged and D(k:k+1,k:k+1) is a 2-by-2 diagonal block.
@param[out]
W (workspace) COMPLEX array, dimension (LDW,NB)
@param[in]
LDW INTEGER
The leading dimension of the array W. LDW >= max(1,N).
@param[out]
INFO INTEGER
- = 0: successful exit
- > 0: if INFO = k, D(k,k) is exactly zero. The factorization
has been completed, but the block diagonal matrix D is
exactly singular.
@ingroup magma_chetrf_comp
********************************************************************/
extern "C" magma_int_t
magma_clahef_gpu(
magma_uplo_t uplo, magma_int_t n, magma_int_t nb, magma_int_t *kb,
magmaFloatComplex *hA, magma_int_t lda,
magmaFloatComplex_ptr dA, size_t dA_offset, magma_int_t ldda,
magma_int_t *ipiv,
magmaFloatComplex_ptr dW, size_t dW_offset, magma_int_t lddw,
magma_queue_t queue,
magma_int_t *info)
{
/* .. Parameters .. */
float d_one = 1.0;
float d_zero = 0.0;
float d_eight = 8.0;
float d_seven = 7.0;
#if defined(PRECISION_c)
float f_zero = 0.0;
#endif
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_mone = -MAGMA_C_ONE;
magma_int_t upper = (uplo == MagmaUpper);
magma_int_t ione = 1;
/* .. Local Scalars .. */
magma_int_t imax = 0, jmax = 0, kk, kkW, kp, kstep, iinfo;
float abs_akk, alpha, colmax, R1, rowmax;
magmaFloatComplex Zimax, Z;
#define dA(i, j) dA, dA_offset + (j)*ldda + (i)
#define dW(i, j) dW, dW_offset + (j)*lddw + (i)
#define A(i, j) (hA + (j)*lda + (i))
/* .. Executable Statements .. */
*info = 0;
/* Initialize alpha for use in choosing pivot block size. */
alpha = ( d_one+sqrt( d_seven ) ) / d_eight;
magma_event_t event = NULL;
if( upper ) {
/* Factorize the trailing columns of A using the upper triangle
of A and working backwards, and compute the matrix W = U12*D
for use in updating A11 (note that conjg(W) is actually stored)
K is the main loop index, decreasing from N in steps of 1 or 2
KW is the column of W which corresponds to column K of A */
int k, kw = 0;
for (k = n-1; k+1 > max(n-nb+1, nb); k -= kstep) {
kw = nb - (n-k);
/* Copy column K of A to column KW of W and update it */
magma_ccopy( k+1, dA( 0, k ), 1, dW( 0, kw ), 1, queue );
// set imaginary part of diagonal to be zero
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dW, 2*(k+ kw*lddw+dW_offset)+1, 1, queue, &event);
magma_queue_sync( queue );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dW, 2*(k+ kw*lddw+dW_offset)+1, 1, queue, &event);
magma_queue_sync( queue );
#endif
if (k+1 < n) {
magma_cgemv( MagmaNoTrans, k+1, n-(k+1), c_mone, dA( 0, k+1 ), ldda,
dW( k, kw+1 ), lddw, c_one, dW( 0, kw ), ione, queue );
// set imaginary part of diagonal to be zero
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dW, 2*(k+ kw*lddw+dW_offset)+1, 1, queue, &event );
magma_queue_sync( queue );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dW, 2*(k+ kw*lddw+dW_offset)+1, 1, queue, &event );
magma_queue_sync( queue );
#endif
}
kstep = 1;
/* Determine rows and columns to be interchanged and whether
a 1-by-1 or 2-by-2 pivot block will be used */
magma_cgetvector_async( 1, dW( k, kw ), 1, &Z, 1, queue, &event );
magma_queue_sync( queue );
abs_akk = fabs( MAGMA_C_REAL( Z ) );
/* imax is the row-index of the largest off-diagonal element in
column K, and colmax is its absolute value */
if( k > 0 ) {
// magma is one-base
imax = magma_icamax( k, dW( 0, kw ), 1, queue ) - 1;
magma_cgetvector( 1, dW( imax, kw ), 1, &Z, 1, queue );
colmax = MAGMA_C_ABS1( Z );
} else {
colmax = d_zero;
}
if( max( abs_akk, colmax ) == 0.0 ) {
/* Column K is zero: set INFO and continue */
if ( *info == 0 ) *info = k;
kp = k;
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dA, 2*(k+ k*ldda+dA_offset)+1, 1, queue, &event );
magma_queue_sync( queue );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dA, 2*(k+ k*ldda+dA_offset)+1, 1, queue, &event );
magma_queue_sync( queue );
#endif
} else {
if( abs_akk >= alpha*colmax ) {
/* no interchange, use 1-by-1 pivot block */
kp = k;
} else {
/* Copy column imax to column KW-1 of W and update it */
magma_ccopy( imax+1, dA( 0, imax ), 1, dW( 0, kw-1 ), 1, queue );
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dW, 2*(imax+ (kw-1)*lddw+dW_offset)+1, 1, queue, &event );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dW, 2*(imax+ (kw-1)*lddw+dW_offset)+1, 1, queue, &event );
#endif
#if defined(PRECISION_z) || defined(PRECISION_c)
magmablas_clacpy_cnjg( k-imax, dA(imax,imax+1), ldda, dW(imax+1,kw-1), 1, queue );
#else
magma_ccopy( k-imax, dA(imax,imax+1), ldda, dW(imax+1,kw-1), 1, queue );
#endif
if( k+1 < n ) {
magma_cgemv( MagmaNoTrans, k+1, n-(k+1), c_mone,
dA( 0, k+1 ), ldda, dW( imax, kw+1 ), lddw,
c_one, dW( 0, kw-1 ), ione, queue );
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dW, 2*(imax+ (kw-1)*lddw+dW_offset)+1, 1, queue, &event );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dW, 2*(imax+ (kw-1)*lddw+dW_offset)+1, 1, queue, &event );
#endif
}
magma_cgetvector_async( 1, dW( imax, kw-1 ), 1, &Zimax, 1, queue, &event );
magma_queue_sync( queue );
/* jmax is the column-index of the largest off-diagonal
element in row imax, and rowmax is its absolute value */
jmax = imax + magma_icamax( k-imax, dW( imax+1, kw-1 ), 1, queue );
magma_cgetvector( 1, dW( jmax, kw-1 ), 1, &Z, 1, queue );
rowmax = MAGMA_C_ABS1( Z );
if ( imax > 0 ) {
// magma is one-base
jmax = magma_icamax( imax, dW( 0, kw-1 ), 1, queue ) - 1;
magma_cgetvector( 1, dW( jmax, kw-1 ), 1, &Z, 1, queue );
rowmax = max( rowmax, MAGMA_C_ABS1( Z ) );
}
if( abs_akk >= alpha*colmax*( colmax / rowmax ) ) {
/* no interchange, use 1-by-1 pivot block */
kp = k;
} else if ( fabs( MAGMA_C_REAL( Zimax ) ) >= alpha*rowmax ) {
/* interchange rows and columns K and imax, use 1-by-1
pivot block */
kp = imax;
/* copy column KW-1 of W to column KW */
magma_ccopy( k+1, dW( 0, kw-1 ), 1, dW( 0, kw ), 1, queue );
} else {
/* interchange rows and columns K-1 and imax, use 2-by-2
pivot block */
kp = imax;
kstep = 2;
}
}
kk = k - kstep + 1;
kkW = nb - (n - kk);
/* Updated column kp is already stored in column kkW of W */
if( kp != kk ) {
/* Interchange rows kk and kp in last kk columns of A and W */
// note: row-swap A(:,kk)
magmablas_cswap( n-kk, dA( kk, kk ), ldda, dA( kp, kk ), ldda, queue );
magmablas_cswap( n-kk, dW( kk, kkW), lddw, dW( kp, kkW), lddw, queue );
/* Copy non-updated column kk to column kp */
#if defined(PRECISION_z) || defined(PRECISION_c)
magmablas_clacpy_cnjg( kk-kp-1, dA( kp+1, kk ), 1, dA( kp, kp+1 ), ldda, queue );
#else
magma_ccopy( kk-kp-1, dA( kp+1, kk ), 1, dA( kp, kp+1 ), ldda, queue );
#endif
// now A(kp,kk) should be A(kk,kk), and copy to A(kp,kp)
magma_ccopy( kp+1, dA( 0, kk ), 1, dA( 0, kp ), 1, queue );
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dA, 2*(kp+ kp*ldda+dA_offset)+1, 1, queue, &event );
magma_queue_sync( queue );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dA, 2*(kp+ kp*ldda+dA_offset)+1, 1, queue, &event );
#endif
}
if( kstep == 1 ) {
/* 1-by-1 pivot block D(k): column KW of W now holds
W(k) = U(k)*D(k)
where U(k) is the k-th column of U
Store U(k) in column k of A */
magma_ccopy( k+1, dW( 0, kw ), 1, dA( 0, k ), 1, queue );
if ( k > 0 ) {
magma_cgetvector_async( 1, dA( k, k ), 1, &Z, 1, queue, &event );
magma_queue_sync( queue );
R1 = d_one / MAGMA_C_REAL( Z );
magma_csscal( k, R1, dA( 0, k ), 1, queue );
/* Conjugate W(k) */
#if defined(PRECISION_z) || defined(PRECISION_c)
magmablas_clacpy_cnjg( k, dW( 0, kw ), 1, dW( 0, kw ), 1, queue );
#endif
}
} else {
/* 2-by-2 pivot block D(k): columns KW and KW-1 of W now hold
( W(k-1) W(k) ) = ( U(k-1) U(k) )*D(k)
where U(k) and U(k-1) are the k-th and (k-1)-th columns of U */
if( k > 1 ) {
/* Store U(k) and U(k-1) in columns k and k-1 of A */
magmablas_clascl_2x2( MagmaUpper,
k-1, dW(0, kw-1), lddw, dA(0,k-1), ldda, &iinfo, queue );
}
/* Copy D(k) to A */
magma_ccopymatrix( 2, 2, dW( k-1, kw-1 ), lddw, dA( k-1, k-1 ), ldda, queue );
/* Conjugate W(k) and W(k-1) */
#if defined(PRECISION_z) || defined(PRECISION_c)
magmablas_clacpy_cnjg( k, dW( 0, kw ), 1, dW( 0, kw ), 1, queue );
magmablas_clacpy_cnjg( k-1, dW( 0, kw-1 ), 1, dW( 0, kw-1 ), 1, queue );
#endif
}
}
/* Store details of the interchanges in ipiv */
if( kstep == 1 ) {
ipiv[ k ] = 1+kp;
} else {
ipiv[ k ] = -(1+kp);
ipiv[ k-1 ] = -(1+kp);
}
}
/* Update the upper triangle of A11 (= A(1:k,1:k)) as
A11 := A11 - U12*D*U12' = A11 - U12*W'
computing blocks of NB columns at a time (note that conjg(W) is
actually stored) */
kw = nb - (n-k);
for (int j = ( k / nb )*nb; j >= 0; j -= nb ) {
int jb = min( nb, k-j+1 );
#ifdef SYMMETRIC_UPDATE
/* Update the upper triangle of the diagonal block */
for (int jj = j; jj < j + jb; jj++) {
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dA, 2*(jj+ jj*ldda+dA_offset)+1, 1, queue, &event );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dA, 2*(jj+ jj*ldda+dA_offset)+1, 1, queue, &event );
#endif
magma_cgemv( MagmaNoTrans, jj-j+1, n-(k+1), c_mone,
dA( j, k+1 ), ldda, dW( jj, kw+1 ), lddw, c_one,
dA( j, jj ), 1, queue );
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dA, 2*(jj+ jj*ldda+dA_offset)+1, 1, queue, &event );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dA, 2*(jj+ jj*ldda+dA_offset)+1, 1, queue, &event );
#endif
}
/* Update the rectangular superdiagonal block */
magma_cgemm( MagmaNoTrans, MagmaTrans, j, jb, n-(k+1),
c_mone, dA( 0, k+1 ), ldda, dW( j, kw+1 ), lddw,
c_one, dA( 0, j ), ldda, queue );
#else
#if defined(PRECISION_z)
magmablas_dlaset(MagmaUpperLower, 1, jb,
0, 0, dA, 2*(j+ j*ldda+dA_offset)+1, 2*(1+ldda), queue );
#elif defined(PRECISION_c)
magmablas_slaset(MagmaUpperLower, 1, jb,
0, 0, dA, 2*(j+ j*ldda+dA_offset)+1, 2*(1+ldda), queue );
#endif
magma_cgemm( MagmaNoTrans, MagmaTrans, j+jb, jb, n-(k+1),
c_mone, dA( 0, k+1 ), ldda,
dW( j, kw+1 ), lddw,
c_one, dA( 0, j ), ldda, queue );
#if defined(PRECISION_z)
magmablas_dlaset(MagmaUpperLower, 1, jb,
0, 0, dA, 2*(j+ j*ldda+dA_offset)+1, 2*(1+ldda), queue );
#elif defined(PRECISION_c)
magmablas_slaset(MagmaUpperLower, 1, jb,
0, 0, dA, 2*(j+ j*ldda+dA_offset)+1, 2*(1+ldda), queue );
#endif
#endif
}
/* Put U12 in standard form by partially undoing the interchanges in columns k+1:n */
for (int j = k+1; j < n;)
{
int jj = j;
int jp = ipiv[ j ];
if( jp < 0 ) {
jp = -jp;
j = j + 1;
}
j = j + 1;
jp = jp - 1;
if( jp != jj && j < n )
magmablas_cswap( n-j, dA( jp, j ), ldda, dA( jj, j ), ldda, queue );
}
// copying the panel back to CPU
magma_cgetmatrix_async( n, n-(k+1), dA(0,k+1), ldda, A(0,k+1), lda, queue, &event );
magma_queue_sync( queue );
/* Set KB to the number of columns factorized */
*kb = n - (k+1);
} else {
/* Factorize the leading columns of A using the lower triangle
of A and working forwards, and compute the matrix W = L21*D
for use in updating A22 (note that conjg(W) is actually stored)
K is the main loop index, increasing from 1 in steps of 1 or 2 */
int k;
for (k = 0; k < min(nb-1,n); k += kstep) {
/* Copy column K of A to column K of W and update it */
/* -------------------------------------------------------------- */
magma_ccopy( n-k, dA( k, k ), 1, dW( k, k ), 1, queue );
// set imaginary part of diagonal to be zero
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dW, 2*(k*lddw+k+dW_offset)+1, 1, queue, &event);
magma_queue_sync( queue );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dW, 2*(k*lddw+k+dW_offset)+1, 1, queue, &event);
magma_queue_sync( queue );
#endif
/* -------------------------------------------------------------- */
magma_cgemv( MagmaNoTrans, n-k, k, c_mone, dA( k, 0 ), ldda,
dW( k, 0 ), lddw, c_one, dW( k, k ), ione, queue );
// re-set imaginary part of diagonal to be zero
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dW, 2*(k*lddw+k+dW_offset)+1, 1, queue, &event );
magma_queue_sync( queue );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dW, 2*(k*lddw+k+dW_offset)+1, 1, queue, &event );
magma_queue_sync( queue );
#endif
kstep = 1;
/* Determine rows and columns to be interchanged and whether
a 1-by-1 or 2-by-2 pivot block will be used */
magma_cgetvector_async( 1, dW( k, k ), 1, &Z, 1, queue, &event );
magma_queue_sync( queue );
abs_akk = fabs( MAGMA_C_REAL( Z ) );
/* imax is the row-index of the largest off-diagonal element in
column K, and colmax is its absolute value */
if( k < n-1 ) {
// magmablas is one-base
imax = k + magma_icamax( n-k-1, dW(k+1,k), 1, queue );
magma_cgetvector( 1, dW( imax,k ), 1, &Z, 1, queue );
colmax = MAGMA_C_ABS1( Z );
} else {
colmax = d_zero;
}
if ( max( abs_akk, colmax ) == 0.0 ) {
/* Column K is zero: set INFO and continue */
if( *info == 0 ) *info = k;
kp = k;
// make sure the imaginary part of diagonal is zero
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dA, 2*(k*ldda+k+dA_offset)+1, 1, queue, &event );
magma_queue_sync( queue );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dA, 2*(k*ldda+k+dA_offset)+1, 1, queue, &event );
magma_queue_sync( queue );
#endif
} else {
if ( abs_akk >= alpha*colmax ) {
/* no interchange, use 1-by-1 pivot block */
kp = k;
} else {
/* Copy column imax to column K+1 of W and update it */
#if defined(PRECISION_z) || defined(PRECISION_c)
magmablas_clacpy_cnjg( imax-k, dA(imax,k), ldda, dW(k,k+1), 1, queue );
#else
magma_ccopy( imax-k, dA( imax, k ), ldda, dW( k, k+1 ), 1, queue );
#endif
magma_ccopy( n-imax, dA( imax, imax ), 1, dW( imax, k+1 ), 1, queue );
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dW, 2*((k+1)*lddw+imax+dW_offset)+1, 1, queue, &event);
magma_queue_sync( queue );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dW, 2*((k+1)*lddw+imax+dW_offset)+1, 1, queue, &event);
magma_queue_sync( queue );
#endif
magma_cgemv( MagmaNoTrans, n-k, k, c_mone, dA( k, 0 ), ldda,
dW( imax, 0 ), lddw, c_one, dW( k, k+1 ), ione, queue );
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dW, 2*((k+1)*lddw+imax+dW_offset)+1, 1, queue, &event);
magma_queue_sync( queue );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dW, 2*((k+1)*lddw+imax+dW_offset)+1, 1, queue, &event);
magma_queue_sync( queue );
#endif
magma_cgetvector_async( 1, dW(imax,k+1), 1, &Zimax, 1, queue, &event);
magma_queue_sync( queue );
/* jmax is the column-index of the largest off-diagonal
element in row imax, and rowmax is its absolute value */
// magmablas is one-base
jmax = k-1 + magma_icamax( imax-k, dW(k, k+1), 1, queue );
magma_cgetvector( 1, dW(jmax,k+1), 1, &Z, 1, queue );
rowmax = MAGMA_C_ABS1( Z );
if( imax < n-1 ) {
// magmablas is one-base
jmax = imax + magma_icamax( (n-1)-imax, dW(imax+1,k+1), 1, queue);
magma_cgetvector( 1, dW(jmax,k+1), 1, &Z, 1, queue );
rowmax = max( rowmax, MAGMA_C_ABS1( Z ) );
}
if( abs_akk >= alpha*colmax*( colmax / rowmax ) ) {
/* no interchange, use 1-by-1 pivot block */
kp = k;
} else if( fabs( MAGMA_C_REAL( Zimax ) ) >= alpha*rowmax ) {
/* interchange rows and columns K and imax, use 1-by-1
pivot block */
kp = imax;
/* copy column K+1 of W to column K */
magma_ccopy( n-k, dW( k, k+1 ), 1, dW( k, k ), 1, queue );
} else {
/* interchange rows and columns K+1 and imax, use 2-by-2
pivot block */
kp = imax;
kstep = 2;
}
}
kk = k + kstep - 1;
/* Updated column kp is already stored in column kk of W */
if( kp != kk ) {
/* Copy non-updated column kk to column kp */
/* ------------------------------------------------------------------ */
#if defined(PRECISION_z) || defined(PRECISION_c)
magmablas_clacpy_cnjg( kp-kk, dA( kk, kk ), 1, dA( kp, kk ), ldda, queue );
#else
magma_ccopy( kp-kk, dA( kk, kk ), 1, dA( kp, kk ), ldda, queue );
#endif
if ( kp < n ) {
magma_ccopy( n-kp, dA( kp, kk), 1, dA( kp, kp ), 1, queue );
}
/* ------------------------------------------------------------------ */
/* Interchange rows kk and kp in first kk columns of A and W */
magmablas_cswap( kk+1, dA( kk, 0 ), ldda, dA( kp, 0 ), ldda, queue );
magmablas_cswap( kk+1, dW( kk, 0 ), lddw, dW( kp, 0 ), lddw, queue );
}
if ( kstep == 1 ) {
/* 1-by-1 pivot block D(k): column k of W now holds
W(k) = L(k)*D(k)
where L(k) is the k-th column of L
Store L(k) in column k of A */
magma_ccopy( n-k, dW( k, k ), 1, dA( k, k ), 1, queue );
if ( k < n-1 ) {
magma_cgetvector_async( 1, dA(k,k), 1, &Z, 1, queue, &event );
magma_queue_sync( queue );
R1 = d_one / MAGMA_C_REAL( Z );
magma_csscal((n-1)-k, R1, dA( k+1,k ), 1, queue);
/* Conjugate W(k) */
#if defined(PRECISION_z) || defined(PRECISION_c)
magmablas_clacpy_cnjg( (n-1)-k, dW( k+1,k ), 1, dW( k+1,k ), 1, queue );
#endif
}
} else {
/* 2-by-2 pivot block D(k): columns k and k+1 of W now hold
( W(k) W(k+1) ) = ( L(k) L(k+1) )*D(k)
where L(k) and L(k+1) are the k-th and (k+1)-th columns
of L */
magmablas_clascl_2x2( MagmaLower,
n-(k+2), dW(k,k), lddw, dA(k+2,k), ldda, &iinfo,
queue );
/* Copy D(k) to A */
magma_ccopymatrix( 2, 2, dW( k, k ), lddw, dA( k, k ), ldda, queue );
/* Conjugate W(k) and W(k+1) */
#if defined(PRECISION_z) || defined(PRECISION_c)
magmablas_clacpy_cnjg( (n-1)-k, dW( k+1,k ), 1, dW( k+1,k ), 1, queue );
magmablas_clacpy_cnjg( (n-1)-k-1, dW( k+2,k+1), 1, dW( k+2,k+1 ), 1, queue );
#endif
}
}
/* Store details of the interchanges in ipiv */
if ( kstep == 1 ) {
ipiv[k] = kp+1;
} else {
ipiv[k] = -kp-1;
ipiv[k+1] = -kp-1;
}
}
/* Update the lower triangle of A22 (= A(k:n,k:n)) as
A22 := A22 - L21*D*L21' = A22 - L21*W'
computing blocks of NB columns at a time (note that conjg(W) is
actually stored) */
for( int j = k; j < n; j += nb ) {
int jb = min( nb, n-j );
/* Update the lower triangle of the diagonal block */
#ifdef SYMMETRIC_UPDATE
for (int jj = j; jj < j + jb; jj++) {
int jnb = j + jb - jj;
/* -------------------------------------------------------- */
magma_cgemv( MagmaNoTrans, jnb, k, c_mone, dA( jj, 0 ), ldda,
dW( jj, 0 ), lddw, c_one, dA( jj, jj ), ione, queue );
/* -------------------------------------------------------- */
}
/* Update the rectangular subdiagonal block */
if( j+jb < n ) {
int nk = n - (j+jb);
/* -------------------------------------------- */
magma_cgemm( MagmaNoTrans, MagmaTrans, nk, jb, k,
c_mone, dA( j+jb, 0 ), ldda,
dW( j, 0 ), lddw,
c_one, dA( j+jb, j ), ldda, queue );
/* ------------------------------------------- */
}
#else
#if defined(PRECISION_z)
magmablas_dlaset(MagmaUpperLower, 1, jb,
0, 0, dA, 2*(j*ldda+j+dA_offset)+1, 2*(1+ldda), queue );
#elif defined(PRECISION_c)
magmablas_slaset(MagmaUpperLower, 1, jb,
0, 0, dA, 2*(j*ldda+j+dA_offset)+1, 2*(1+ldda), queue );
#endif
magma_cgemm( MagmaNoTrans, MagmaTrans, n-j, jb, k,
c_mone, dA( j, 0 ), ldda,
dW( j, 0 ), lddw,
c_one, dA( j, j ), ldda, queue );
#if defined(PRECISION_z)
magmablas_dlaset(MagmaUpperLower, 1, jb,
0, 0, dA, 2*(j*ldda+j+dA_offset)+1, 2*(1+ldda), queue );
#elif defined(PRECISION_c)
magmablas_slaset(MagmaUpperLower, 1, jb,
0, 0, dA, 2*(j*ldda+j+dA_offset)+1, 2*(1+ldda), queue );
#endif
#endif
}
/* Put L21 in standard form by partially undoing the interchanges
in columns 1:k-1 */
for (int j = k; j > 0;) {
int jj = j;
int jp = ipiv[j-1];
if( jp < 0 ) {
jp = -jp;
j--;
}
j--;
if ( jp != jj && j >= 1 ) {
magmablas_cswap( j, dA( jp-1,0 ), ldda, dA( jj-1,0 ), ldda, queue );
}
}
// copying the panel back to CPU
magma_cgetmatrix_async( n, k, dA(0,0), ldda, A(0,0), lda, queue, &event );
magma_queue_sync( queue );
/* Set KB to the number of columns factorized */
*kb = k;
}
return *info;
/* End of CLAHEF */
}
