/**
@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
=======
SSYTRF_nopiv_gpu computes the LDLt factorization of a real symmetric
matrix A.
The factorization has the form
A = U^H * D * U , if UPLO = 'U', or
A = L * D * L^H, if UPLO = 'L',
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 CHARACTER*1
- = 'U': Upper triangle of A is stored;
- = 'L': Lower triangle of A is stored.
@param[in]
N INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
dA REAL array on the GPU, dimension (LDA,N)
On entry, the symmetric 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.
\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]
LDA INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
INFO INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
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_ssytrf_comp
******************************************************************* */
extern "C" magma_int_t
magma_ssytrf_nopiv_gpu(
magma_uplo_t uplo, magma_int_t n,
magmaFloat_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))
/* Local variables */
float zone = MAGMA_S_ONE;
float mzone = MAGMA_S_NEG_ONE;
int 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 MAGMA_ERR_ILLEGAL_VALUE;
}
/* Quick return */
if ( n == 0 )
return MAGMA_SUCCESS;
nb = magma_get_ssytrf_nopiv_nb(n);
ib = min(32, nb); // inner-block for diagonal factorization
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
magma_queue_t stream[2];
magma_event_t event;
magma_queue_create(&stream[0]);
magma_queue_create(&stream[1]);
magma_event_create( &event );
trace_init( 1, 1, 2, stream );
// CPU workspace
float *A;
if (MAGMA_SUCCESS != magma_smalloc_pinned( &A, nb*nb )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
// GPU workspace
magmaFloat_ptr dW;
if (MAGMA_SUCCESS != magma_smalloc( &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_queue_wait_event( stream[1], event );
magma_event_sync(event);
magma_sgetmatrix_async(jb, jb, dA(j, j), ldda, A(j,j), nb, stream[1]);
trace_gpu_end( 0, 0 );
// factorize the diagonal block
magma_queue_sync(stream[1]);
trace_cpu_start( 0, "potrf", "potrf" );
ssytrf_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_ssetmatrix_async(jb, jb, A(j, j), nb, dA(j, j), ldda, stream[0]);
trace_gpu_end( 0, 0 );
if ( (j+jb) < n) {
// compute the off-diagonal blocks of current block column
magmablasSetKernelStream( stream[0] );
trace_gpu_start( 0, 0, "trsm", "trsm" );
magma_strsm(MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaUnit,
jb, (n-j-jb),
zone, dA(j, j), ldda,
dA(j, j+jb), ldda);
magma_scopymatrix( jb, n-j-jb, dA( j, j+jb ), ldda, dWt( 0, j+jb ), nb );
// update the trailing submatrix with D
magmablas_slascl_diag(MagmaUpper, jb, n-j-jb,
dA(j, j), ldda,
dA(j, j+jb), ldda,
&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_sgemm(MagmaConjTrans, MagmaNoTrans, kb, n-k, jb,
mzone, dWt(0, k), nb,
dA(j, k), ldda,
zone, dA(k, k), ldda);
if (k==j+jb)
magma_event_record( event, stream[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_queue_wait_event( stream[0], event );
magma_event_sync(event);
magma_sgetmatrix_async(jb, jb, dA(j, j), ldda, A(j,j), nb, stream[1]);
trace_gpu_end( 0, 0 );
// factorize the diagonal block
magma_queue_sync(stream[1]);
trace_cpu_start( 0, "potrf", "potrf" );
ssytrf_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_ssetmatrix_async(jb, jb, A(j, j), nb, dA(j, j), ldda, stream[0]);
trace_gpu_end( 0, 0 );
if ( (j+jb) < n) {
// compute the off-diagonal blocks of current block column
magmablasSetKernelStream( stream[0] );
trace_gpu_start( 0, 0, "trsm", "trsm" );
magma_strsm(MagmaRight, MagmaLower, MagmaConjTrans, MagmaUnit,
(n-j-jb), jb,
zone, dA(j, j), ldda,
dA(j+jb, j), ldda);
magma_scopymatrix( n-j-jb,jb, dA( j+jb, j ), ldda, dW( j+jb, 0 ), ldda );
// update the trailing submatrix with D
magmablas_slascl_diag(MagmaLower, n-j-jb, jb,
dA(j, j), ldda,
dA(j+jb, j), ldda,
&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_sgemm(MagmaNoTrans, MagmaConjTrans, n-k, kb, jb,
mzone, dA(k, j), ldda,
dW(k, 0), ldda,
zone, dA(k, k), ldda);
if (k==j+jb)
magma_event_record( event, stream[0] );
}
trace_gpu_end( 0, 0 );
}
}
}
trace_finalize( "ssytrf.svg","trace.css" );
magma_queue_destroy(stream[0]);
magma_queue_destroy(stream[1]);
magma_event_destroy( event );
magma_free( dW );
magma_free_pinned( A );
magmablasSetKernelStream( orig_stream );
return MAGMA_SUCCESS;
} /* magma_ssytrf_nopiv */
/**
Purpose
-------
CGEQP3 computes a QR factorization with column pivoting of a
matrix A: A*P = Q*R using Level 3 BLAS.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
dA COMPLEX array on the GPU, dimension (LDDA,N)
On entry, the M-by-N matrix A.
On exit, the upper triangle of the array contains the
min(M,N)-by-N upper trapezoidal matrix R; the elements below
the diagonal, together with the array TAU, represent the
unitary matrix Q as a product of min(M,N) elementary
reflectors.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,M).
@param[in,out]
jpvt INTEGER array, dimension (N)
On entry, if JPVT(J).ne.0, the J-th column of A is permuted
to the front of A*P (a leading column); if JPVT(J)=0,
the J-th column of A is a free column.
On exit, if JPVT(J)=K, then the J-th column of A*P was the
the K-th column of A.
@param[out]
tau COMPLEX array, dimension (min(M,N))
The scalar factors of the elementary reflectors.
@param[out]
dwork (workspace) COMPLEX array on the GPU, dimension (MAX(1,LWORK))
On exit, if INFO=0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK.
For [sd]geqp3, LWORK >= (N+1)*NB + 2*N;
for [cz]geqp3, LWORK >= (N+1)*NB,
where NB is the optimal blocksize.
\n
Note: unlike the CPU interface of this routine, the GPU interface
does not support a workspace query.
@param
rwork (workspace, for [cz]geqp3 only) REAL array, dimension (2*N)
@param[out]
info INTEGER
- = 0: successful exit.
- < 0: if INFO = -i, the i-th argument had an illegal value.
Further Details
---------------
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector
with v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in
A(i+1:m,i), and tau in TAU(i).
@ingroup magma_cgeqp3_comp
********************************************************************/
extern "C" magma_int_t
magma_cgeqp3_gpu(
magma_int_t m, magma_int_t n,
magmaFloatComplex_ptr dA, magma_int_t ldda,
magma_int_t *jpvt, magmaFloatComplex *tau,
magmaFloatComplex_ptr dwork, magma_int_t lwork,
#ifdef COMPLEX
float *rwork,
#endif
magma_int_t *info )
{
#define dA(i_, j_) (dA + (i_) + (j_)*ldda)
const magmaFloatComplex c_zero = MAGMA_C_ZERO;
const magma_int_t ione = 1;
//magma_int_t na;
magma_int_t n_j;
magma_int_t j, jb, nb, sm, sn, fjb, nfxd, minmn;
magma_int_t topbmn, lwkopt;
*info = 0;
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (ldda < max(1,m)) {
*info = -4;
}
nb = magma_get_cgeqp3_nb( m, n );
minmn = min(m,n);
if (*info == 0) {
if (minmn == 0) {
lwkopt = 1;
} else {
lwkopt = (n + 1)*nb;
#ifdef REAL
lwkopt += 2*n;
#endif
}
if (lwork < lwkopt) {
*info = -8;
}
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (minmn == 0)
return *info;
#ifdef REAL
float *rwork = dwork + (n + 1)*nb;
#endif
magmaFloatComplex_ptr df;
if (MAGMA_SUCCESS != magma_cmalloc( &df, (n+1)*nb )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magmaFloat_ptr dlsticcs;
if (MAGMA_SUCCESS != magma_smalloc( &dlsticcs, 1+256*(n+255)/256 )) {
magma_free( df );
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_queue_t queue;
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queue );
magmablas_claset( MagmaFull, n+1, nb, c_zero, c_zero, df, n+1, queue );
nfxd = 0;
/* Move initial columns up front.
* Note jpvt uses 1-based indices for historical compatibility. */
for (j = 0; j < n; ++j) {
if (jpvt[j] != 0) {
if (j != nfxd) {
blasf77_cswap(&m, dA(0, j), &ione, dA(0, nfxd), &ione); // TODO: ERROR, matrix not on CPU!
jpvt[j] = jpvt[nfxd];
jpvt[nfxd] = j + 1;
}
else {
jpvt[j] = j + 1;
}
++nfxd;
}
else {
jpvt[j] = j + 1;
}
}
/*
// TODO:
Factorize fixed columns
=======================
Compute the QR factorization of fixed columns and update
remaining columns.
if (nfxd > 0) {
na = min(m,nfxd);
lapackf77_cgeqrf(&m, &na, dA, &ldda, tau, dwork, &lwork, info);
if (na < n) {
n_j = n - na;
lapackf77_cunmqr( MagmaLeftStr, MagmaConjTransStr, &m, &n_j, &na,
dA, &ldda, tau, dA(0, na), &ldda,
dwork, &lwork, info );
}
}*/
/* Factorize free columns */
if (nfxd < minmn) {
sm = m - nfxd;
sn = n - nfxd;
//sminmn = minmn - nfxd;
/* Initialize partial column norms. */
magmablas_scnrm2_cols( sm, sn, dA(nfxd,nfxd), ldda, &rwork[nfxd], queue );
magma_scopymatrix( sn, 1, &rwork[nfxd], sn, &rwork[n+nfxd], sn, queue );
j = nfxd;
//if (nb < sminmn)
{
/* Use blocked code initially. */
/* Compute factorization: while loop. */
topbmn = minmn; // - nb;
while(j < topbmn) {
jb = min(nb, topbmn - j);
/* Factorize JB columns among columns J:N. */
n_j = n - j;
//magma_claqps_gpu // this is a cpp-file
magma_claqps2_gpu // this is a cuda-file
( m, n_j, j, jb, &fjb,
dA(0, j), ldda,
&jpvt[j], &tau[j], &rwork[j], &rwork[n + j],
dwork,
&df[jb], n_j,
dlsticcs, queue );
j += fjb; /* fjb is actual number of columns factored */
}
}
/*
// Use unblocked code to factor the last or only block.
if (j < minmn) {
n_j = n - j;
if (j > nfxd) {
magma_cgetmatrix( m-j, n_j,
dA(j,j), ldda,
A(j,j), lda, queue );
}
lapackf77_claqp2(&m, &n_j, &j, dA(0, j), &ldda, &jpvt[j],
&tau[j], &rwork[j], &rwork[n+j], dwork );
}*/
}
magma_queue_destroy( queue );
magma_free( df );
magma_free( dlsticcs );
return *info;
} /* magma_cgeqp3_gpu */
/**
Purpose
-------
SSYEVDX computes selected eigenvalues and, optionally, eigenvectors
of a real symmetric 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 REAL 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. 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) REAL array, dimension (LDWA, N)
@param[in]
ldwa INTEGER
The leading dimension of the array wA. LDWA >= max(1,N).
@param[out]
work (workspace) REAL array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The length of the array WORK.
If N <= 1, LWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LWORK >= 2*N + N*NB.
If JOBZ = MagmaVec and N > 1, LWORK >= max( 2*N + N*NB, 1 + 6*N + 2*N**2 ).
NB can be obtained through magma_get_ssytrd_nb(N).
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal sizes of the WORK and IWORK
arrays, returns these values as the first entries of the WORK
and IWORK arrays, and no error message related to LWORK 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 and
IWORK arrays, returns these values as the first entries of
the WORK and IWORK arrays, and no error message related to
LWORK 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_ssyev_driver
********************************************************************/
extern "C" magma_int_t
magma_ssyevdx_gpu(magma_vec_t jobz, magma_range_t range, magma_uplo_t uplo,
magma_int_t n,
float *dA, magma_int_t ldda,
float vl, float vu, magma_int_t il, magma_int_t iu,
magma_int_t *m, float *w,
float *wA, magma_int_t ldwa,
float *work, magma_int_t lwork,
magma_int_t *iwork, magma_int_t liwork,
magma_int_t *info)
{
magma_int_t ione = 1;
float d__1;
float eps;
magma_int_t inde;
float anrm;
float rmin, rmax;
float sigma;
magma_int_t iinfo, lwmin;
magma_int_t lower;
magma_int_t wantz;
magma_int_t indwk2, llwrk2;
magma_int_t iscale;
float safmin;
float bignum;
magma_int_t indtau;
magma_int_t indwrk, liwmin;
magma_int_t llwork;
float smlnum;
magma_int_t lquery;
magma_int_t alleig, valeig, indeig;
float *dwork;
magma_int_t lddc = ldda;
wantz = (jobz == MagmaVec);
lower = (uplo == MagmaLower);
alleig = (range == MagmaRangeAll);
valeig = (range == MagmaRangeV);
indeig = (range == MagmaRangeI);
lquery = (lwork == -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_ssytrd_nb( n );
if ( n <= 1 ) {
lwmin = 1;
liwmin = 1;
}
else if ( wantz ) {
lwmin = max( 2*n + n*nb, 1 + 6*n + 2*n*n );
liwmin = 3 + 5*n;
}
else {
lwmin = 2*n + n*nb;
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] = lwmin * one_eps;
iwork[0] = liwmin;
if ((lwork < lwmin) && !lquery) {
*info = -16;
} else if ((liwork < liwmin) && ! lquery) {
*info = -18;
}
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
const char* jobz_ = lapack_vec_const( jobz );
const char* uplo_ = lapack_uplo_const( uplo );
float *A;
magma_smalloc_cpu( &A, n*n );
magma_sgetmatrix(n, n, dA, ldda, A, n);
lapackf77_ssyevd(jobz_, uplo_,
&n, A, &n,
w, work, &lwork,
iwork, &liwork, info);
magma_ssetmatrix( n, n, A, n, dA, ldda);
magma_free_cpu(A);
return *info;
}
magma_queue_t stream;
magma_queue_create( &stream );
// n*lddc for ssytrd2_gpu
// n for slansy
magma_int_t ldwork = n*lddc;
if ( wantz ) {
// need 3n^2/2 for sstedx
ldwork = max( ldwork, 3*n*(n/2 + 1));
}
if (MAGMA_SUCCESS != magma_smalloc( &dwork, ldwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/* 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_slansy(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_slascl(uplo, 0, 0, 1., sigma, n, n, dA, ldda, info);
}
/* Call SSYTRD to reduce symmetric matrix to tridiagonal form. */
// ssytrd work: e (n) + tau (n) + llwork (n*nb) ==> 2n + n*nb
// sstedx work: e (n) + tau (n) + z (n*n) + llwrk2 (1 + 4*n + n^2) ==> 1 + 6n + 2n^2
inde = 0;
indtau = inde + n;
indwrk = indtau + n;
indwk2 = indwrk + n*n;
llwork = lwork - indwrk;
llwrk2 = lwork - indwk2;
magma_timer_t time=0;
timer_start( time );
#ifdef FAST_SYMV
magma_ssytrd2_gpu(uplo, n, dA, ldda, w, &work[inde],
&work[indtau], wA, ldwa, &work[indwrk], llwork,
dwork, n*lddc, &iinfo);
#else
magma_ssytrd_gpu(uplo, n, dA, ldda, w, &work[inde],
&work[indtau], wA, ldwa, &work[indwrk], llwork,
&iinfo);
#endif
timer_stop( time );
timer_printf( "time ssytrd = %6.2f\n", time );
/* For eigenvalues only, call SSTERF. For eigenvectors, first call
SSTEDC to generate the eigenvector matrix, WORK(INDWRK), of the
tridiagonal matrix, then call SORMTR to multiply it to the Householder
transformations represented as Householder vectors in A. */
if (! wantz) {
lapackf77_ssterf(&n, w, &work[inde], info);
magma_smove_eig(range, n, w, &il, &iu, vl, vu, m);
}
else {
timer_start( time );
magma_sstedx(range, n, vl, vu, il, iu, w, &work[inde],
&work[indwrk], n, &work[indwk2],
llwrk2, iwork, liwork, dwork, info);
timer_stop( time );
timer_printf( "time sstedx = %6.2f\n", time );
timer_start( time );
magma_smove_eig(range, n, w, &il, &iu, vl, vu, m);
magma_ssetmatrix( n, *m, &work[indwrk + n* (il-1) ], n, dwork, lddc );
magma_sormtr_gpu(MagmaLeft, uplo, MagmaNoTrans, n, *m, dA, ldda, &work[indtau],
dwork, lddc, wA, ldwa, &iinfo);
magma_scopymatrix( n, *m, dwork, lddc, dA, ldda );
timer_stop( time );
timer_printf( "time sormtr + copy = %6.2f\n", time );
}
/* If matrix was scaled, then rescale eigenvalues appropriately. */
if (iscale == 1) {
d__1 = 1. / sigma;
blasf77_sscal(&n, &d__1, w, &ione);
}
work[0] = lwmin * one_eps; // round up
iwork[0] = liwmin;
magma_queue_destroy( stream );
magma_free( dwork );
return *info;
} /* magma_ssyevd_gpu */
/**
@deprecated
Purpose
-------
SLAQPS computes a step of QR factorization with column pivoting
of a real 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 REAL 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 REAL 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 REAL array, dimension (NB), on the GPU
Auxiliary vector.
@param[in,out]
dF REAL 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_sgeqp3_aux
********************************************************************/
extern "C" magma_int_t
magma_slaqps_gpu(
magma_int_t m, magma_int_t n, magma_int_t offset,
magma_int_t nb, magma_int_t *kb,
magmaFloat_ptr dA, magma_int_t ldda,
magma_int_t *jpvt, float *tau,
float *vn1, float *vn2,
magmaFloat_ptr dauxv,
magmaFloat_ptr dF, magma_int_t lddf)
{
#define dA(i, j) (dA + (i) + (j)*(ldda))
#define dF(i, j) (dF + (i) + (j)*(lddf))
float c_zero = MAGMA_S_MAKE( 0.,0.);
float c_one = MAGMA_S_MAKE( 1.,0.);
float c_neg_one = MAGMA_S_MAKE(-1.,0.);
magma_int_t ione = 1;
magma_int_t i__1, i__2;
float z__1;
magma_int_t k, rk;
magmaFloat_ptr dAks;
float tauk = MAGMA_S_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_smalloc( &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_sswap( m, dA(0, pvt), ione, dA(0, k), ione, queue );
magmablas_sswap( 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_sgemv( 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_sgemv( 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_slarfg_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_ssetvector( 1, &c_one, 1, dA(rk, k), 1, queue );
else magma_scopymatrix( 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_sgetvector( 1, &tau[k], 1, &tauk, 1, queue );
if (k < n-1) {
i__1 = m - rk;
i__2 = n - k - 1;
/* Multiply on GPU */
magma_sgemv( 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_S_NEGATE( tauk );
#ifdef RIGHT_UPDATE
i__1 = m - offset - nb;
i__2 = k;
magma_sgemv( 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_sgemv( 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_sgemv( 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_sgemv( 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_sgemm( 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_sgemm( 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_snrm2_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_scopymatrix( 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_sgemm( 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_snrm2_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_slaqps */
/**
Purpose
=======
SSYTRF_nopiv computes the LDLt factorization of a real symmetric
matrix A. This version does not require work space on the GPU passed
as input. GPU memory is allocated in the routine.
The factorization has the form
A = U^H * D * U, if UPLO = MagmaUpper, or
A = L * D * L^H, 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]
A REAL array, dimension (LDA,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.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
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_ssysv_comp
******************************************************************* */
extern "C" magma_int_t
magma_ssytrf_nopiv(
magma_uplo_t uplo, magma_int_t n,
float *A, magma_int_t lda,
magma_int_t *info)
{
#define A(i, j) ( A +(j)*lda + (i))
#define dA(i, j) (dA +(j)*ldda + (i))
#define dW(i, j) (dW +(j)*ldda + (i))
#define dWt(i, j) (dW +(j)*nb + (i))
/* Constants */
const float c_one = MAGMA_S_ONE;
const float c_neg_one = MAGMA_S_NEG_ONE;
/* Local variables */
bool upper = (uplo == MagmaUpper);
magma_int_t j, k, jb, ldda, nb, ib, iinfo;
magmaFloat_ptr dA;
magmaFloat_ptr dW;
*info = 0;
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return MAGMA_ERR_ILLEGAL_VALUE;
}
/* Quick return */
if ( n == 0 )
return MAGMA_SUCCESS;
ldda = magma_roundup( n, 32 );
nb = magma_get_ssytrf_nopiv_nb(n);
ib = min(32, nb); // inner-block for diagonal factorization
if ((MAGMA_SUCCESS != magma_smalloc(&dA, n *ldda)) ||
(MAGMA_SUCCESS != magma_smalloc(&dW, nb*ldda))) {
/* alloc failed so call the non-GPU-resident version */
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_device_t cdev;
magma_queue_t queues[2];
magma_event_t event;
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 );
/* Use hybrid blocked code. */
if (upper) {
//=========================================================
// Compute the LDLt factorization A = U'*D*U without pivoting.
// copy matrix to GPU
for (j=0; j < n; j += nb) {
jb = min(nb, (n-j));
trace_gpu_start( 0, 0, "set", "set" );
magma_ssetmatrix_async(j+jb, jb, A(0, j), lda, dA(0, j), ldda, queues[0]);
trace_gpu_end( 0, 0 );
}
// 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" );
if ( j != 0) {
//magma_event_sync(event);
magma_sgetmatrix_async(jb, jb, dA(j, j), ldda, A(j,j), lda, queues[1]);
}
trace_gpu_end( 0, 0 );
// factorize the diagonal block
magma_queue_sync(queues[1]);
trace_cpu_start( 0, "potrf", "potrf" );
magma_ssytrf_nopiv_cpu( MagmaUpper, jb, ib, A(j, j), lda, 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_ssetmatrix_async(jb, jb, A(j, j), lda, dA(j, j), ldda, queues[0]);
trace_gpu_end( 0, 0 );
// copy j-th column of U back to CPU
trace_gpu_start( 0, 1, "get", "get" );
magma_sgetmatrix_async(j, jb, dA(0, j), ldda, A(0, j), lda, queues[1]);
trace_gpu_end( 0, 1 );
if ( (j+jb) < n) {
// compute the off-diagonal blocks of current block column
trace_gpu_start( 0, 0, "trsm", "trsm" );
magma_strsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaUnit,
jb, (n-j-jb),
c_one, dA(j, j), ldda,
dA(j, j+jb), ldda,
queues[0] );
magma_scopymatrix( jb, n-j-jb, dA( j, j+jb ), ldda, dWt( 0, j+jb ), nb, queues[0] );
// update the trailing submatrix with D
magmablas_slascl_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_sgemm( MagmaConjTrans, 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] );
magma_queue_sync( queues[0] );
}
}
trace_gpu_end( 0, 0 );
}
}
} else {
//=========================================================
// Compute the LDLt factorization A = L*D*L' without pivoting.
// copy the matrix to GPU
for (j=0; j < n; j += nb) {
jb = min(nb, (n-j));
trace_gpu_start( 0, 0, "set", "set" );
magma_ssetmatrix_async((n-j), jb, A(j, j), lda, dA(j, j), ldda, queues[0]);
trace_gpu_end( 0, 0 );
}
// 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" );
if (j != 0) {
//magma_event_sync(event);
magma_sgetmatrix_async(jb, jb, dA(j, j), ldda, A(j,j), lda, queues[1]);
}
trace_gpu_end( 0, 0 );
// factorize the diagonal block
magma_queue_sync(queues[1]);
trace_cpu_start( 0, "potrf", "potrf" );
magma_ssytrf_nopiv_cpu( MagmaLower, jb, ib, A(j, j), lda, 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_ssetmatrix_async(jb, jb, A(j, j), lda, dA(j, j), ldda, queues[0]);
trace_gpu_end( 0, 0 );
// copy j-th row of L back to CPU
trace_gpu_start( 0, 1, "get", "get" );
magma_sgetmatrix_async(jb, j, dA(j, 0), ldda, A(j, 0), lda, queues[1]);
trace_gpu_end( 0, 1 );
if ( (j+jb) < n) {
// compute the off-diagonal blocks of current block column
trace_gpu_start( 0, 0, "trsm", "trsm" );
magma_strsm( MagmaRight, MagmaLower, MagmaConjTrans, MagmaUnit,
(n-j-jb), jb,
c_one, dA(j, j), ldda,
dA(j+jb, j), ldda,
queues[0] );
magma_scopymatrix( n-j-jb,jb, dA( j+jb, j ), ldda, dW( j+jb, 0 ), ldda, queues[0] );
// update the trailing submatrix with D
magmablas_slascl_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_sgemm( MagmaNoTrans, MagmaConjTrans, 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] );
magma_queue_sync(queues[0]);
}
}
trace_gpu_end( 0, 0 );
}
}
}
trace_finalize( "ssytrf.svg","trace.css" );
magma_queue_destroy(queues[0]);
magma_queue_destroy(queues[1]);
magma_event_destroy( event );
magma_free(dW);
magma_free(dA);
return MAGMA_SUCCESS;
} /* magma_ssytrf_nopiv */
extern "C" magma_int_t
magma_zlaqps_gpu(magma_int_t m, magma_int_t n, magma_int_t offset,
magma_int_t nb, magma_int_t *kb,
magmaDoubleComplex *A, magma_int_t lda,
magma_int_t *jpvt, magmaDoubleComplex *tau,
double *vn1, double *vn2,
magmaDoubleComplex *auxv,
magmaDoubleComplex *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
=======
ZLAQPS 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 ))
magmaDoubleComplex c_zero = MAGMA_Z_MAKE( 0.,0.);
magmaDoubleComplex c_one = MAGMA_Z_MAKE( 1.,0.);
magmaDoubleComplex c_neg_one = MAGMA_Z_MAKE(-1.,0.);
magma_int_t ione = 1;
magma_int_t i__1, i__2;
//double d__1;
magmaDoubleComplex z__1;
//magma_int_t j;
magma_int_t k, rk;
//magmaDoubleComplex Akk;
magmaDoubleComplex *Aks;
magmaDoubleComplex tauk;
magma_int_t pvt;
//double temp, temp2;
double tol3z;
magma_int_t itemp;
double lsticc, *lsticcs;
magma_int_t lastrk;
magma_dmalloc( &lsticcs, 1+256*(n+255)/256 );
lastrk = min( m, n + offset );
tol3z = magma_dsqrt( lapackf77_dlamch("Epsilon"));
lsticc = 0;
k = 0;
magma_zmalloc( &Aks, nb );
while( k < nb && lsticc == 0 ) {
rk = offset + k;
/* Determine ith pivot column and swap if necessary */
// Fortran: pvt, k, idamax are all 1-based; subtract 1 from k.
// C: pvt, k, idamax are all 0-based; don't subtract 1.
pvt = k - 1 + magma_idamax( n-k, &vn1[k], ione );
if (pvt != k) {
/*if (pvt >= nb) {
// 1. Start copy from GPU
magma_zgetmatrix_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_zswap( &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_zswap(&m, A(0, pvt), &ione, A(0, k), &ione);
// 3. Restore the GPU
magma_zsetmatrix_async( m - offset - nb, 1,
A (offset + nb, pvt), lda,
dA(offset + nb, pvt), ldda, stream);
}*/
magmablas_zswap( m, A(0, pvt), ione, A(0, k), ione );
//blasf77_zswap( &i__1, F(pvt,0), &ldf, F(k,0), &ldf );
magmablas_zswap( 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_Z_CNJG( *F(k,j) );
}
#endif*/
//#define RIGHT_UPDATE
#ifdef RIGHT_UPDATE
i__1 = m - offset - nb;
i__2 = k;
magma_zgemv( 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_zgemv( MagmaNoTransStr, &i__1, &i__2,
&c_neg_one, A(rk, 0), &lda,
F(k, 0), &ldf,
&c_one, A(rk, k), &ione );*/
magma_zgemv( 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_Z_CNJG( *F(k,j) );
}
#endif*/
}
/* Generate elementary reflector H(k). */
magma_zlarfg_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_zgetvector( 1, &Aks[k], 1, &Akk, 1 );
/* needed to avoid the race condition */
if (k == 0) magma_zsetvector( 1, &c_one, 1, A(rk, k), 1 );
else magma_zcopymatrix( 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_zgetvector( 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_zsetmatrix( i__1, 1, A(rk, k), lda, dA(rk,k), ldda );
/* Multiply on GPU */
// was CALL ZGEMV( '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_zgetvector( 1, &tau[k], 1, &tauk, 1 );
magma_zgemv( MagmaConjTrans, m-rk, n-k-1,
tauk, A( rk, k+1 ), lda,
A( rk, k ), 1,
c_zero, F( k+1, k ), 1 );
//magma_zscal( 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_zgemv( 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_zgetmatrix_async( i__2-i__3, 1,
// dF(k + 1 +i__3, k), i__2,
// F (k + 1 +i__3, k), i__2, stream );
//blasf77_zgemv( 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_zgemv( 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_zsetvector( 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_zgetvector( 1, &tau[k], 1, &tauk, 1 );
z__1 = MAGMA_Z_NEGATE( tauk );
#ifdef RIGHT_UPDATE
i__1 = m - offset - nb;
i__2 = k;
magma_zgemv( 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_zgemv( 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_zgemv( MagmaConjTransStr, &i__1, &i__2,
// &z__1, A(rk, 0), &lda,
// A(rk, k), &ione,
// &c_zero, auxv, &ione );
magma_zgemv( MagmaConjTrans, i__1, i__2,
z__1, A(rk, 0), lda,
A(rk, k), ione,
c_zero, auxv, ione );
//i__1 = k;
//blasf77_zgemv( MagmaNoTransStr, &n, &i__1,
// &c_one, F(0,0), &ldf,
// auxv, &ione,
// &c_one, F(0,k), &ione );
/*magma_zgemv( 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_zgemv( 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_zgemm( 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_zgemm( 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_zgemm( 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_dznrm2_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_Z_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] = (double) lsticc;
lsticc = j;
} else {
vn1[j] *= magma_dsqrt(temp);
}
}
}
}*/
//*A(rk, k) = Akk;
//magma_zsetvector( 1, &Akk, 1, A(rk, k), 1 );
//magma_zswap( 1, &Aks[k], 1, A(rk, k), 1 );
++k;
}
magma_zcopymatrix( 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_zsetmatrix( i__2, *kb,
F (*kb, 0), ldf,
dF(*kb, 0), i__2 );*/
magma_zgemm( 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_dznrm2_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_dznrm2(i__1, A(rk + 1, lsticc), ione);
else {
// Where is the data, CPU or GPU ?
double r1, r2;
r1 = cblas_dznrm2(nb-k, A(rk + 1, lsticc), ione);
r2 = magma_dznrm2(m-offset-nb, dA(offset + nb + 1, lsticc), ione);
vn1[lsticc] = magma_dsqrt(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(DLAMCH('S'))
vn2[lsticc] = vn1[lsticc];
lsticc = itemp;*/
}
magma_free(Aks);
magma_free(lsticcs);
return MAGMA_SUCCESS;
} /* magma_zlaqps */
/**
Purpose
-------
SSYEVD_GPU computes all eigenvalues and, optionally, eigenvectors of
a real symmetric matrix A. 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]
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 REAL 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. 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, A contains the
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[out]
w REAL array, dimension (N)
If INFO = 0, the eigenvalues in ascending order.
@param
wA (workspace) REAL array, dimension (LDWA, N)
@param[in]
ldwa INTEGER
The leading dimension of the array wA. LDWA >= max(1,N).
@param[out]
work (workspace) REAL array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The length of the array WORK.
If N <= 1, LWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LWORK >= 2*N + N*NB.
If JOBZ = MagmaVec and N > 1, LWORK >= max( 2*N + N*NB, 1 + 6*N + 2*N**2 ).
NB can be obtained through magma_get_ssytrd_nb(N).
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal sizes of the WORK and IWORK
arrays, returns these values as the first entries of the WORK
and IWORK arrays, and no error message related to LWORK 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 and
IWORK arrays, returns these values as the first entries of
the WORK and IWORK arrays, and no error message related to
LWORK 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_ssyev_driver
********************************************************************/
extern "C" magma_int_t
magma_ssyevd_gpu(
magma_vec_t jobz, magma_uplo_t uplo,
magma_int_t n,
magmaFloat_ptr dA, magma_int_t ldda,
float *w,
float *wA, magma_int_t ldwa,
float *work, magma_int_t lwork,
#ifdef COMPLEX
float *rwork, magma_int_t lrwork,
#endif
magma_int_t *iwork, magma_int_t liwork,
magma_int_t *info)
{
magma_int_t ione = 1;
float d__1;
float eps;
magma_int_t inde;
float anrm;
float rmin, rmax;
float sigma;
magma_int_t iinfo, lwmin;
magma_int_t lower;
magma_int_t wantz;
magma_int_t indwk2, llwrk2;
magma_int_t iscale;
float safmin;
float bignum;
magma_int_t indtau;
magma_int_t indwrk, liwmin;
magma_int_t llwork;
float smlnum;
magma_int_t lquery;
magmaFloat_ptr dwork;
magma_int_t lddc = ldda;
wantz = (jobz == MagmaVec);
lower = (uplo == MagmaLower);
lquery = (lwork == -1 || liwork == -1);
*info = 0;
if (! (wantz || (jobz == MagmaNoVec))) {
*info = -1;
} else if (! (lower || (uplo == MagmaUpper))) {
*info = -2;
} else if (n < 0) {
*info = -3;
} else if (ldda < max(1,n)) {
*info = -5;
}
magma_int_t nb = magma_get_ssytrd_nb( n );
if ( n <= 1 ) {
lwmin = 1;
liwmin = 1;
}
else if ( wantz ) {
lwmin = max( 2*n + n*nb, 1 + 6*n + 2*n*n );
liwmin = 3 + 5*n;
}
else {
lwmin = 2*n + n*nb;
liwmin = 1;
}
work[0] = magma_smake_lwork( lwmin );
iwork[0] = liwmin;
if ((lwork < lwmin) && !lquery) {
*info = -10;
} else if ((liwork < liwmin) && ! lquery) {
*info = -12;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
magma_queue_t queue;
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queue );
/* If matrix is very small, then just call LAPACK on CPU, no need for GPU */
if (n <= 128) {
magma_int_t lda = n;
float *A;
magma_smalloc_cpu( &A, lda*n );
magma_sgetmatrix( n, n, dA, ldda, A, lda, queue );
lapackf77_ssyevd( lapack_vec_const(jobz), lapack_uplo_const(uplo),
&n, A, &lda,
w, work, &lwork,
iwork, &liwork, info );
magma_ssetmatrix( n, n, A, lda, dA, ldda, queue );
magma_free_cpu( A );
magma_queue_destroy( queue );
return *info;
}
// ssytrd2_gpu requires ldda*ceildiv(n,64) + 2*ldda*nb
// sormtr_gpu requires lddc*n
// slansy requires n
magma_int_t ldwork = max( ldda*magma_ceildiv(n,64) + 2*ldda*nb, lddc*n );
ldwork = max( ldwork, n );
if ( wantz ) {
// sstedx requires 3n^2/2
ldwork = max( ldwork, 3*n*(n/2 + 1) );
}
if (MAGMA_SUCCESS != magma_smalloc( &dwork, ldwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/* 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_slansy( MagmaMaxNorm, uplo, n, dA, ldda, dwork, ldwork, queue );
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_slascl( uplo, 0, 0, 1., sigma, n, n, dA, ldda, queue, info );
}
/* Call SSYTRD to reduce symmetric matrix to tridiagonal form. */
// ssytrd work: e (n) + tau (n) + llwork (n*nb) ==> 2n + n*nb
// sstedx work: e (n) + tau (n) + z (n*n) + llwrk2 (1 + 4*n + n^2) ==> 1 + 6n + 2n^2
inde = 0;
indtau = inde + n;
indwrk = indtau + n;
indwk2 = indwrk + n*n;
llwork = lwork - indwrk;
llwrk2 = lwork - indwk2;
magma_timer_t time=0;
timer_start( time );
#ifdef FAST_SYMV
magma_ssytrd2_gpu( uplo, n, dA, ldda, w, &work[inde],
&work[indtau], wA, ldwa, &work[indwrk], llwork,
dwork, ldwork, &iinfo );
#else
magma_ssytrd_gpu( uplo, n, dA, ldda, w, &work[inde],
&work[indtau], wA, ldwa, &work[indwrk], llwork,
&iinfo );
#endif
timer_stop( time );
#ifdef FAST_SYMV
timer_printf( "time ssytrd2 = %6.2f\n", time );
#else
timer_printf( "time ssytrd = %6.2f\n", time );
#endif
/* For eigenvalues only, call SSTERF. For eigenvectors, first call
SSTEDC to generate the eigenvector matrix, WORK(INDWRK), of the
tridiagonal matrix, then call SORMTR to multiply it to the Householder
transformations represented as Householder vectors in A. */
if (! wantz) {
lapackf77_ssterf( &n, w, &work[inde], info );
}
else {
timer_start( time );
magma_sstedx( MagmaRangeAll, n, 0., 0., 0, 0, w, &work[inde],
&work[indwrk], n, &work[indwk2],
llwrk2, iwork, liwork, dwork, info );
timer_stop( time );
timer_printf( "time sstedx = %6.2f\n", time );
timer_start( time );
magma_ssetmatrix( n, n, &work[indwrk], n, dwork, lddc, queue );
magma_sormtr_gpu( MagmaLeft, uplo, MagmaNoTrans, n, n, dA, ldda, &work[indtau],
dwork, lddc, wA, ldwa, &iinfo );
magma_scopymatrix( n, n, dwork, lddc, dA, ldda, queue );
timer_stop( time );
timer_printf( "time sormtr + copy = %6.2f\n", time );
}
/* If matrix was scaled, then rescale eigenvalues appropriately. */
if (iscale == 1) {
d__1 = 1. / sigma;
blasf77_sscal( &n, &d__1, w, &ione );
}
work[0] = magma_smake_lwork( lwmin );
iwork[0] = liwmin;
magma_queue_destroy( queue );
magma_free( dwork );
return *info;
} /* magma_ssyevd_gpu */
extern "C" magma_int_t
magma_ssytrf_nopiv(magma_uplo_t uplo, magma_int_t n,
float *A, magma_int_t lda,
magma_int_t *info)
{
/* -- MAGMA (version 1.6.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2011
Purpose
=======
SSYTRF_nopiv computes the LDLt factorization of a real symmetric
matrix A. This version does not require work space on the GPU passed
as input. GPU memory is allocated in the routine.
The factorization has the form
A = U\*\*H * D * U, if UPLO = 'U', or
A = L * D * L\*\*H, if UPLO = 'L',
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
=========
UPLO (input) CHARACTER*1
= 'U': Upper triangle of A is stored;
= 'L': Lower triangle of A is stored.
N (input) INTEGER
The order of the matrix A. N >= 0.
A (input/output) REAL array, dimension (LDA,N)
On entry, the symmetric matrix A. If UPLO = 'U', the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = 'L', the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit, if INFO = 0, the factor U or L from the Cholesky
factorization A = U\*\*H*U or A = L*L\*\*H.
Higher performance is achieved if A is in pinned memory, e.g.
allocated using cudaMallocHost.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
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.
===================================================================== */
/* Local variables */
float zone = MAGMA_S_ONE;
float mzone = MAGMA_S_NEG_ONE;
int upper = (uplo == MagmaUpper);
magma_int_t j, k, jb, ldda, nb, ib, iinfo;
magmaFloat_ptr dA;
magmaFloat_ptr dW;
*info = 0;
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return MAGMA_ERR_ILLEGAL_VALUE;
}
/* Quick return */
if ( n == 0 )
return MAGMA_SUCCESS;
ldda = ((n+31)/32)*32;
nb = magma_get_ssytrf_nopiv_nb(n);
ib = min(32, nb); // inner-block for diagonal factorization
if ((MAGMA_SUCCESS != magma_smalloc(&dA, n *ldda)) ||
(MAGMA_SUCCESS != magma_smalloc(&dW, nb*ldda))) {
/* alloc failed so call the non-GPU-resident version */
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_queue_t stream[2];
magma_event_t event;
magma_queue_create(&stream[0]);
magma_queue_create(&stream[1]);
magma_event_create( &event );
trace_init( 1, 1, 2, (CUstream_st**)stream );
//if (nb <= 1 || nb >= n)
//{
// lapackf77_spotrf(uplo_, &n, a, &lda, info);
//} else
{
/* Use hybrid blocked code. */
if (upper) {
//=========================================================
// Compute the LDLt factorization A = U'*D*U without pivoting.
// copy matrix to GPU
for (j=0; j<n; j+=nb) {
jb = min(nb, (n-j));
trace_gpu_start( 0, 0, "set", "set" );
magma_ssetmatrix_async(j+jb, jb, A(0, j), lda, dA(0, j), ldda, stream[0]);
trace_gpu_end( 0, 0 );
}
// 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_sgetmatrix_async(jb, jb, dA(j, j), ldda, A(j,j), lda, stream[0]);
trace_gpu_end( 0, 0 );
// copy j-th column of U back to CPU
magma_queue_wait_event( stream[1], event );
trace_gpu_start( 0, 1, "get", "get" );
magma_sgetmatrix_async(j, jb, dA(0, j), ldda, A(0, j), lda, stream[1]);
trace_gpu_end( 0, 1 );
// factorize the diagonal block
magma_queue_sync(stream[0]);
trace_cpu_start( 0, "potrf", "potrf" );
ssytrf_nopiv_cpu(MagmaUpper, jb, ib, A(j, j), lda, 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_ssetmatrix_async(jb, jb, A(j, j), lda, dA(j, j), ldda, stream[0]);
trace_gpu_end( 0, 0 );
if ( (j+jb) < n) {
// compute the off-diagonal blocks of current block column
magmablasSetKernelStream( stream[0] );
trace_gpu_start( 0, 0, "trsm", "trsm" );
magma_strsm(MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaUnit,
jb, (n-j-jb),
zone, dA(j, j), ldda,
dA(j, j+jb), ldda);
magma_scopymatrix( jb, n-j-jb, dA( j, j+jb ), ldda, dWt( 0, j+jb ), nb );
// update the trailing submatrix with D
magmablas_slascl_diag(MagmaUpper, jb, n-j-jb,
dA(j, j), ldda,
dA(j, j+jb), ldda,
&iinfo);
magma_event_record( event, stream[0] );
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_sgemm(MagmaConjTrans, MagmaNoTrans, kb, n-k, jb,
mzone, dWt(0, k), nb,
dA(j, k), ldda,
zone, dA(k, k), ldda);
}
trace_gpu_end( 0, 0 );
}
}
} else {
//=========================================================
// Compute the LDLt factorization A = L*D*L' without pivoting.
// copy the matrix to GPU
for (j=0; j<n; j+=nb) {
jb = min(nb, (n-j));
trace_gpu_start( 0, 0, "set", "set" );
magma_ssetmatrix_async((n-j), jb, A(j, j), lda, dA(j, j), ldda, stream[0]);
trace_gpu_end( 0, 0 );
}
// 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_sgetmatrix_async(jb, jb, dA(j, j), ldda, A(j,j), lda, stream[0]);
trace_gpu_end( 0, 0 );
// copy j-th row of L back to CPU
magma_queue_wait_event( stream[1], event );
trace_gpu_start( 0, 1, "get", "get" );
magma_sgetmatrix_async(jb, j, dA(j, 0), ldda, A(j, 0), lda, stream[1]);
trace_gpu_end( 0, 1 );
// factorize the diagonal block
magma_queue_sync(stream[0]);
trace_cpu_start( 0, "potrf", "potrf" );
ssytrf_nopiv_cpu(MagmaLower, jb, ib, A(j, j), lda, 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_ssetmatrix_async(jb, jb, A(j, j), lda, dA(j, j), ldda, stream[0]);
trace_gpu_end( 0, 0 );
if ( (j+jb) < n) {
// compute the off-diagonal blocks of current block column
magmablasSetKernelStream( stream[0] );
trace_gpu_start( 0, 0, "trsm", "trsm" );
magma_strsm(MagmaRight, MagmaLower, MagmaConjTrans, MagmaUnit,
(n-j-jb), jb,
zone, dA(j, j), ldda,
dA(j+jb, j), ldda);
magma_scopymatrix( n-j-jb,jb, dA( j+jb, j ), ldda, dW( j+jb, 0 ), ldda );
// update the trailing submatrix with D
magmablas_slascl_diag(MagmaLower, n-j-jb, jb,
dA(j, j), ldda,
dA(j+jb, j), ldda,
&iinfo);
magma_event_record( event, stream[0] );
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_sgemm(MagmaNoTrans, MagmaConjTrans, n-k, kb, jb,
mzone, dA(k, j), ldda,
dW(k, 0), ldda,
zone, dA(k, k), ldda);
}
trace_gpu_end( 0, 0 );
}
}
}
}
trace_finalize( "ssytrf.svg","trace.css" );
magma_queue_destroy(stream[0]);
magma_queue_destroy(stream[1]);
magma_event_destroy( event );
magma_free(dW);
magma_free(dA);
return MAGMA_SUCCESS;
} /* magma_ssytrf_nopiv */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing sswap, sswapblk, spermute, slaswp, slaswpx
*/
int main( int argc, char** argv)
{
TESTING_INIT();
float *h_A1, *h_A2;
float *d_A1, *d_A2;
float *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_opts opts;
parse_opts( argc, argv, &opts );
magma_queue_t queue = 0;
printf(" cublasSswap sswap sswapblk slaswp spermute slaswp2 slaswpx scopymatrix 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 slaswp (GByte/s)\n");
printf("==================================================================================================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
// each test is assigned one bit in the check bitmask, bit=1 is failure.
// shift keeps track of which bit is for current test
int shift = 1;
int check = 0;
N = opts.nsize[i];
lda = N;
ldda = ((N+31)/32)*32;
nb = (opts.nb > 0 ? opts.nb : magma_get_sgetrf_nb( N ));
// for each swap, does 2N loads and 2N stores
gbytes = sizeof(float) * 4.*N*nb / 1e9;
TESTING_MALLOC_PIN( h_A1, float, lda*N );
TESTING_MALLOC_PIN( h_A2, float, lda*N );
TESTING_MALLOC_PIN( h_R1, float, lda*N );
TESTING_MALLOC_PIN( h_R2, float, 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, float, ldda*N );
TESTING_MALLOC_DEV( d_A2, float, ldda*N );
for( j=0; j < nb; j++ ) {
ipiv[j] = (magma_int_t) ((rand()*1.*N) / (RAND_MAX * 1.)) + 1;
}
/* =====================================================================
* cublasSswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_ssetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
cublasSswap( 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_sswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_sgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_sgetmatrix( 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_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_ssetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
cublasSswap( 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_sswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_sgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_sgetmatrix( 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;
/* =====================================================================
* sswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_ssetmatrix( 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_sswap( 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_sswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_sgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_sgetmatrix( 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_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_ssetmatrix( 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_sswap( 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_sswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_sgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_sgetmatrix( 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;
/* =====================================================================
* sswapblk, blocked version (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_ssetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_sswapblk( 'R', 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_sswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_sgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_sgetmatrix( 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_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_ssetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_sswapblk( 'C', 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_sswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_sgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_sgetmatrix( 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;
/* =====================================================================
* spermute_long (1 matrix)
*/
/* Row Major */
memcpy( ipiv2, ipiv, nb*sizeof(magma_int_t) ); // spermute updates ipiv2
init_matrix( N, N, h_A1, lda, 0 );
magma_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_spermute_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_sswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_sgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style slaswp (1 matrix)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_slaswp( 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_sswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_sgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style slaswp (1 matrix) - d_ipiv on GPU
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_ssetmatrix( 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_slaswp2( N, d_A1, ldda, 1, nb, d_ipiv );
time = magma_sync_wtime( queue ) - time;
row_perf7 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_sswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_sgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style slaswpx (extended for row- and col-major) (1 matrix)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_slaswpx( 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_sswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_sgetmatrix( 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_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_slaswpx( N, d_A1, 1, ldda, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
col_perf5 = gbytes / time;
time = magma_wtime();
lapackf77_slaswp( &N, h_A1, &lda, &ione, &nb, ipiv, &ione);
time = magma_wtime() - time;
cpu_perf = gbytes / time;
magma_sgetmatrix( 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_scopymatrix( 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_scopymatrix( 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" : "* failures") );
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 );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
extern "C" magma_int_t
magma_ssyevd_gpu(char jobz, char uplo,
magma_int_t n,
float *da, magma_int_t ldda,
float *w,
float *wa, magma_int_t ldwa,
float *work, magma_int_t lwork,
magma_int_t *iwork, magma_int_t liwork,
magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
SSYEVD_GPU computes all eigenvalues and, optionally, eigenvectors of
a real symmetric matrix A. 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
=========
JOBZ (input) CHARACTER*1
= 'N': Compute eigenvalues only;
= 'V': Compute eigenvalues and eigenvectors.
UPLO (input) CHARACTER*1
= 'U': Upper triangle of A is stored;
= 'L': Lower triangle of A is stored.
N (input) INTEGER
The order of the matrix A. N >= 0.
DA (device input/output) REAL array on the GPU,
dimension (LDDA, N).
On entry, the symmetric matrix A. If UPLO = 'U', the
leading N-by-N upper triangular part of A contains the
upper triangular part of the matrix A. If UPLO = 'L',
the leading N-by-N lower triangular part of A contains
the lower triangular part of the matrix A.
On exit, if JOBZ = 'V', then if INFO = 0, A contains the
orthonormal eigenvectors of the matrix A.
If JOBZ = 'N', then on exit the lower triangle (if UPLO='L')
or the upper triangle (if UPLO='U') of A, including the
diagonal, is destroyed.
LDDA (input) INTEGER
The leading dimension of the array DA. LDDA >= max(1,N).
W (output) DOUBLE PRECISION array, dimension (N)
If INFO = 0, the eigenvalues in ascending order.
WA (workspace) DOUBLE PRECISION array, dimension (LDWA, N)
LDWA (input) INTEGER
The leading dimension of the array WA. LDWA >= max(1,N).
WORK (workspace/output) REAL array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
LWORK (input) INTEGER
The length of the array WORK.
If N <= 1, LWORK >= 1.
If JOBZ = 'N' and N > 1, LWORK >= 2*N + N*NB.
If JOBZ = 'V' and N > 1, LWORK >= max( 2*N + N*NB, 1 + 6*N + 2*N**2 ).
NB can be obtained through magma_get_ssytrd_nb(N).
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal sizes of the WORK and IWORK
arrays, returns these values as the first entries of the WORK
and IWORK arrays, and no error message related to LWORK or
LIWORK is issued by XERBLA.
IWORK (workspace/output) INTEGER array, dimension (MAX(1,LIWORK))
On exit, if INFO = 0, IWORK[0] returns the optimal LIWORK.
LIWORK (input) INTEGER
The dimension of the array IWORK.
If N <= 1, LIWORK >= 1.
If JOBZ = 'N' and N > 1, LIWORK >= 1.
If JOBZ = 'V' and N > 1, LIWORK >= 3 + 5*N.
If LIWORK = -1, then a workspace query is assumed; the
routine only calculates the optimal sizes of the WORK and
IWORK arrays, returns these values as the first entries of
the WORK and IWORK arrays, and no error message related to
LWORK or LIWORK is issued by XERBLA.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
> 0: if INFO = i and JOBZ = 'N', 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 = 'V', 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.
===================================================================== */
char uplo_[2] = {uplo, 0};
char jobz_[2] = {jobz, 0};
magma_int_t ione = 1;
float d__1;
float eps;
magma_int_t inde;
float anrm;
float rmin, rmax;
float sigma;
magma_int_t iinfo, lwmin;
magma_int_t lower;
magma_int_t wantz;
magma_int_t indwk2, llwrk2;
magma_int_t iscale;
float safmin;
float bignum;
magma_int_t indtau;
magma_int_t indwrk, liwmin;
magma_int_t llwork;
float smlnum;
magma_int_t lquery;
float *dwork;
magma_int_t lddc = ldda;
wantz = lapackf77_lsame(jobz_, MagmaVecStr);
lower = lapackf77_lsame(uplo_, MagmaLowerStr);
lquery = lwork == -1 || liwork == -1;
*info = 0;
if (! (wantz || lapackf77_lsame(jobz_, MagmaNoVecStr))) {
*info = -1;
} else if (! (lower || lapackf77_lsame(uplo_, MagmaUpperStr))) {
*info = -2;
} else if (n < 0) {
*info = -3;
} else if (ldda < max(1,n)) {
*info = -5;
}
magma_int_t nb = magma_get_ssytrd_nb( n );
if ( n <= 1 ) {
lwmin = 1;
liwmin = 1;
}
else if ( wantz ) {
lwmin = max( 2*n + n*nb, 1 + 6*n + 2*n*n );
liwmin = 3 + 5*n;
}
else {
lwmin = 2*n + n*nb;
liwmin = 1;
}
// multiply by 1+eps to ensure length gets rounded up,
// if it cannot be exactly represented in floating point.
work[0] = lwmin * (1. + lapackf77_slamch("Epsilon"));
iwork[0] = liwmin;
if ((lwork < lwmin) && !lquery) {
*info = -10;
} else if ((liwork < liwmin) && ! lquery) {
*info = -12;
}
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
char jobz_[2] = {jobz, 0}, uplo_[2] = {uplo, 0};
float *a = (float *) malloc( n * n * sizeof(float) );
magma_sgetmatrix(n, n, da, ldda, a, n);
lapackf77_ssyevd(jobz_, uplo_,
&n, a, &n,
w, work, &lwork,
iwork, &liwork, info);
magma_ssetmatrix( n, n, a, n, da, ldda);
free(a);
return *info;
}
magma_queue_t stream;
magma_queue_create( &stream );
// n*lddc for ssytrd2_gpu
// n for slansy
magma_int_t ldwork = n*lddc;
if ( wantz ) {
// need 3n^2/2 for sstedx
ldwork = max( ldwork, 3*n*(n/2 + 1));
}
if (MAGMA_SUCCESS != magma_smalloc( &dwork, ldwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/* 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_slansy('M', 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_slascl(uplo, 0, 0, 1., sigma, n, n, da, ldda, info);
}
/* Call SSYTRD to reduce symmetric matrix to tridiagonal form. */
// ssytrd work: e (n) + tau (n) + llwork (n*nb) ==> 2n + n*nb
// sstedx work: e (n) + tau (n) + z (n*n) + llwrk2 (1 + 4*n + n^2) ==> 1 + 6n + 2n^2
inde = 0;
indtau = inde + n;
indwrk = indtau + n;
indwk2 = indwrk + n*n;
llwork = lwork - indwrk;
llwrk2 = lwork - indwk2;
//
#ifdef ENABLE_TIMER
magma_timestr_t start, end;
start = get_current_time();
#endif
#ifdef FAST_SYMV
magma_ssytrd2_gpu(uplo, n, da, ldda, w, &work[inde],
&work[indtau], wa, ldwa, &work[indwrk], llwork,
dwork, n*lddc, &iinfo);
#else
magma_ssytrd_gpu(uplo, n, da, ldda, w, &work[inde],
&work[indtau], wa, ldwa, &work[indwrk], llwork,
&iinfo);
#endif
#ifdef ENABLE_TIMER
end = get_current_time();
#ifdef FAST_SYMV
printf("time ssytrd2 = %6.2f\n", GetTimerValue(start,end)/1000.);
#else
printf("time ssytrd = %6.2f\n", GetTimerValue(start,end)/1000.);
#endif
#endif
/* For eigenvalues only, call SSTERF. For eigenvectors, first call
SSTEDC to generate the eigenvector matrix, WORK(INDWRK), of the
tridiagonal matrix, then call SORMTR to multiply it to the Householder
transformations represented as Householder vectors in A. */
if (! wantz) {
lapackf77_ssterf(&n, w, &work[inde], info);
} else {
#ifdef ENABLE_TIMER
start = get_current_time();
#endif
magma_sstedx('A', n, 0., 0., 0, 0, w, &work[inde],
&work[indwrk], n, &work[indwk2],
llwrk2, iwork, liwork, dwork, info);
#ifdef ENABLE_TIMER
end = get_current_time();
printf("time sstedx = %6.2f\n", GetTimerValue(start,end)/1000.);
#endif
magma_ssetmatrix( n, n, &work[indwrk], n, dwork, lddc );
#ifdef ENABLE_TIMER
start = get_current_time();
#endif
magma_sormtr_gpu(MagmaLeft, uplo, MagmaNoTrans, n, n, da, ldda, &work[indtau],
dwork, lddc, wa, ldwa, &iinfo);
magma_scopymatrix( n, n, dwork, lddc, da, ldda );
#ifdef ENABLE_TIMER
end = get_current_time();
printf("time sormtr + copy = %6.2f\n", GetTimerValue(start,end)/1000.);
#endif
}
/* If matrix was scaled, then rescale eigenvalues appropriately. */
if (iscale == 1) {
d__1 = 1. / sigma;
blasf77_sscal(&n, &d__1, w, &ione);
}
work[0] = lwmin * (1. + lapackf77_slamch("Epsilon")); // round up
iwork[0] = liwmin;
magma_queue_destroy( stream );
magma_free( dwork );
return *info;
} /* magma_ssyevd_gpu */
/**
Purpose
-------
Solves the least squares problem
min || A*X - C ||
using the QR factorization A = Q*R computed by SGEQRF_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 REAL 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
SGEQRF_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 REAL array, dimension (N)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by MAGMA_SGEQRF_GPU.
@param[in,out]
dB REAL 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]
dT REAL array that is the output (the 6th argument)
of magma_sgeqrf_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).
@param[in]
lddb INTEGER
The leading dimension of the array dB. LDDB >= M.
@param[out]
hwork (workspace) REAL array, dimension (LWORK)
On exit, if INFO = 0, WORK(1) 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_sgeqrf_nb( M ).
\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_sgels_comp
********************************************************************/
extern "C" magma_int_t
magma_sgeqrs_gpu(magma_int_t m, magma_int_t n, magma_int_t nrhs,
float *dA, magma_int_t ldda,
float *tau, float *dT,
float *dB, magma_int_t lddb,
float *hwork, magma_int_t lwork,
magma_int_t *info)
{
#define dA(a_1,a_2) (dA + (a_2)*(ldda) + (a_1))
#define dT(a_1) (dT + (lddwork+(a_1))*nb)
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
float *dwork;
magma_int_t i, k, lddwork, rows, ib;
magma_int_t ione = 1;
magma_int_t nb = magma_get_sgeqrf_nb(m);
magma_int_t lwkopt = (m - n + nb)*(nrhs + nb) + nrhs*nb;
int lquery = (lwork == -1);
hwork[0] = MAGMA_S_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 = -9;
else if (lwork < lwkopt && ! lquery)
*info = -11;
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_sormqr_gpu( MagmaLeft, MagmaTrans,
m, nrhs, n,
dA(0,0), ldda, tau,
dB, lddb, hwork, lwork, dT, nb, info );
if ( *info != 0 ) {
return *info;
}
/* Solve R*X = B(1:n,:) */
lddwork= k;
if (nb < k)
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 = (k-1)/nb * nb;
ib = n-i;
rows = m-i;
// TODO: this assumes that, on exit from magma_sormqr_gpu, hwork contains
// the last block of A and B (i.e., C in sormqr). This should be fixed.
// Seems this data should already be on the GPU, so could switch to
// magma_strsm and drop the ssetmatrix.
if ( nrhs == 1 ) {
blasf77_strsv( MagmaUpperStr, MagmaNoTransStr, MagmaNonUnitStr,
&ib, hwork, &rows,
hwork+rows*ib, &ione);
} else {
blasf77_strsm( MagmaLeftStr, MagmaUpperStr, MagmaNoTransStr, MagmaNonUnitStr,
&ib, &nrhs,
&c_one, hwork, &rows,
hwork+rows*ib, &rows);
}
// update the solution vector
magma_ssetmatrix( ib, nrhs, hwork+rows*ib, rows, dwork+i, lddwork );
// update c
if (nrhs == 1)
magma_sgemv( MagmaNoTrans, i, ib,
c_neg_one, dA(0, i), ldda,
dwork + i, 1,
c_one, dB, 1);
else
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
i, nrhs, ib,
c_neg_one, dA(0, i), ldda,
dwork + i, lddwork,
c_one, dB, lddb);
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_sgemv( MagmaNoTrans, ib, ib,
c_one, dT(i), ib,
dB+i, 1,
c_zero, dwork+i, 1);
magma_sgemv( MagmaNoTrans, i, ib,
c_neg_one, dA(0, i), ldda,
dwork + i, 1,
c_one, dB, 1);
}
else {
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
ib, nrhs, ib,
c_one, dT(i), ib,
dB+i, lddb,
c_zero, dwork+i, lddwork);
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
i, nrhs, ib,
c_neg_one, dA(0, i), ldda,
dwork + i, lddwork,
c_one, dB, lddb);
}
}
}
}
magma_scopymatrix( (n), nrhs,
dwork, lddwork,
dB, lddb );
return *info;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing sswap, sswapblk, slaswp, slaswpx
*/
int main( int argc, char** argv)
{
TESTING_INIT();
float *h_A1, *h_A2;
float *h_R1, *h_R2;
magmaFloat_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 sswap sswap sswapblk slaswp slaswp2 slaswpx scopymatrix 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 slaswp (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_sgetrf_nb( N ));
nb = min( N, nb );
// each swap does 2N loads and 2N stores, for nb swaps
gbytes = sizeof(float) * 4.*N*nb / 1e9;
TESTING_MALLOC_PIN( h_A1, float, lda*N );
TESTING_MALLOC_PIN( h_A2, float, lda*N );
TESTING_MALLOC_PIN( h_R1, float, lda*N );
TESTING_MALLOC_PIN( h_R2, float, 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, float, ldda*N );
TESTING_MALLOC_DEV( d_A2, float, 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 sswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_ssetmatrix( 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
cublasSswap( opts.handle, N, d_A1+ldda*j, 1, d_A2+ldda*(ipiv[j]-1), 1 );
#else
magma_sswap( 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_sswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_sgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_sgetmatrix( 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_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_ssetmatrix( 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
cublasSswap( opts.handle, N, d_A1+j, ldda, d_A2+ipiv[j]-1, ldda );
#else
magma_sswap( 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_sswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_sgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_sgetmatrix( 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;
/* =====================================================================
* sswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_ssetmatrix( 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_sswap( 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_sswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_sgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_sgetmatrix( 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_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_ssetmatrix( 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_sswap( 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_sswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_sgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_sgetmatrix( 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;
/* =====================================================================
* sswapblk, 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_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_ssetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_sswapblk( 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_sswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_sgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_sgetmatrix( 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_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_ssetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_sswapblk( 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_sswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_sgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_sgetmatrix( 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 slaswp (1 matrix)
*/
#ifdef HAVE_CUBLAS
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_slaswp( 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_sswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_sgetmatrix( 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 slaswp (1 matrix) - d_ipiv on GPU
*/
#ifdef HAVE_CUBLAS
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_ssetmatrix( 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_slaswp2( 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_sswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_sgetmatrix( 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 slaswpx (extended for row- and col-major) (1 matrix)
*/
#ifdef HAVE_CUBLAS
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_slaswpx( 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_sswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_sgetmatrix( 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_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_slaswpx( 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_slaswp( &N, h_A1, &lda, &ione, &nb, ipiv, &ione);
time = magma_wtime() - time;
cpu_perf = gbytes / time;
#ifdef HAVE_CUBLAS
magma_sgetmatrix( 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_scopymatrix( 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_scopymatrix( 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;
}
extern "C" magma_err_t
magma_sgeqrs_gpu(magma_int_t m, magma_int_t n, magma_int_t nrhs,
magmaFloat_ptr dA, size_t dA_offset, magma_int_t ldda,
float *tau, magmaFloat_ptr dT, size_t dT_offset,
magmaFloat_ptr dB, size_t dB_offset, magma_int_t lddb,
float *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 SGEQRF_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) REAL 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
SGEQRF_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) REAL array, dimension (N)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by MAGMA_SGEQRF_GPU.
DB (input/output) REAL 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) REAL array that is the output (the 6th argument)
of magma_sgeqrf_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) REAL 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_sgeqrf_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)
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
magmaFloat_ptr dwork;
magma_int_t i, k, lddwork, rows, ib;
magma_int_t ione = 1;
magma_int_t nb = magma_get_sgeqrf_nb(m);
magma_int_t lwkopt = (m-n+nb)*(nrhs+2*nb);
long int lquery = (lwork == -1);
hwork[0] = MAGMA_S_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_sormqr_gpu( MagmaLeft, MagmaTrans,
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_smalloc( &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_strsv( MagmaUpperStr, MagmaNoTransStr, MagmaNonUnitStr,
&ib, hwork, &rows,
hwork+rows*ib, &ione);
} else {
blasf77_strsm( MagmaLeftStr, MagmaUpperStr, MagmaNoTransStr, MagmaNonUnitStr,
&ib, &nrhs,
&c_one, hwork, &rows,
hwork+rows*ib, &rows);
}
// update the solution vector
magma_ssetmatrix( ib, nrhs, hwork+rows*ib, 0, rows, dwork, dwork_offset+i, lddwork, queue );
// update c
if (nrhs == 1)
magma_sgemv( 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_sgemm( 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_sgemv( MagmaNoTrans, ib, ib,
c_one, d_ref(i), ib,
dB, dB_offset+i, 1,
c_zero, dwork, dwork_offset+i, 1, queue );
magma_sgemv( 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_sgemm( 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_sgemm( 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_scopymatrix( (n), nrhs,
dwork, dwork_offset, lddwork,
dB, dB_offset, lddb, queue );
if (nb >= k)
magma_free(dwork);
magma_queue_sync( queue );
return *info;
}
extern "C" magma_int_t
magma_sgeqp3_gpu( magma_int_t m, magma_int_t n,
float *A, magma_int_t lda,
magma_int_t *jpvt, float *tau,
float *work, magma_int_t lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
float *rwork,
#endif
magma_int_t *info )
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
SGEQP3 computes a QR factorization with column pivoting of a
matrix A: A*P = Q*R using Level 3 BLAS.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. N >= 0.
A (input/output) REAL array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, the upper triangle of the array contains the
min(M,N)-by-N upper trapezoidal matrix R; the elements below
the diagonal, together with the array TAU, represent the
unitary matrix Q as a product of min(M,N) elementary
reflectors.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,M).
JPVT (input/output) INTEGER array, dimension (N)
On entry, if JPVT(J).ne.0, the J-th column of A is permuted
to the front of A*P (a leading column); if JPVT(J)=0,
the J-th column of A is a free column.
On exit, if JPVT(J)=K, then the J-th column of A*P was the
the K-th column of A.
TAU (output) REAL array, dimension (min(M,N))
The scalar factors of the elementary reflectors.
WORK (workspace/output) REAL array, dimension (MAX(1,LWORK))
On exit, if INFO=0, WORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The dimension of the array WORK.
For [sd]geqp3, LWORK >= (N+1)*NB + 2*N;
for [cz]geqp3, LWORK >= (N+1)*NB,
where NB is the optimal blocksize.
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
For [cz]geqp3 only:
RWORK (workspace) DOUBLE PRECISION array, dimension (2*N)
INFO (output) INTEGER
= 0: successful exit.
< 0: if INFO = -i, the i-th argument had an illegal value.
Further Details
===============
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector
with v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in
A(i+1:m,i), and tau in TAU(i).
===================================================================== */
#define A(i, j) (A + (i) + (j)*(lda ))
magma_int_t ione = 1;
//magma_int_t na;
magma_int_t n_j;
magma_int_t j, jb, nb, sm, sn, fjb, nfxd, minmn;
magma_int_t topbmn, sminmn, lwkopt, lquery;
*info = 0;
lquery = (lwork == -1);
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,m)) {
*info = -4;
}
nb = magma_get_sgeqp3_nb(min(m, n));
if (*info == 0) {
minmn = min(m,n);
if (minmn == 0) {
lwkopt = 1;
} else {
lwkopt = (n + 1)*nb;
#if defined(PRECISION_d) || defined(PRECISION_s)
lwkopt += 2*n;
#endif
}
//work[0] = MAGMA_S_MAKE( lwkopt, 0. );
if (lwork < lwkopt && ! lquery) {
*info = -8;
}
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
} else if (lquery) {
return *info;
}
if (minmn == 0)
return *info;
#if defined(PRECISION_d) || defined(PRECISION_s)
float *rwork = work + (n + 1)*nb;
#endif
float *df;
if (MAGMA_SUCCESS != magma_smalloc( &df, (n+1)*nb )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
cudaMemset( df, 0, (n+1)*nb*sizeof(float) );
nfxd = 0;
/* Move initial columns up front.
* Note jpvt uses 1-based indices for historical compatibility. */
for (j = 0; j < n; ++j) {
if (jpvt[j] != 0) {
if (j != nfxd) {
blasf77_sswap(&m, A(0, j), &ione, A(0, nfxd), &ione);
jpvt[j] = jpvt[nfxd];
jpvt[nfxd] = j + 1;
}
else {
jpvt[j] = j + 1;
}
++nfxd;
}
else {
jpvt[j] = j + 1;
}
}
/* Factorize fixed columns
=======================
Compute the QR factorization of fixed columns and update
remaining columns.
if (nfxd > 0) {
na = min(m,nfxd);
lapackf77_sgeqrf(&m, &na, A, &lda, tau, work, &lwork, info);
if (na < n) {
n_j = n - na;
lapackf77_sormqr( MagmaLeftStr, MagmaTransStr, &m, &n_j, &na,
A, &lda, tau, A(0, na), &lda,
work, &lwork, info );
}
}*/
/* Factorize free columns */
if (nfxd < minmn) {
sm = m - nfxd;
sn = n - nfxd;
sminmn = minmn - nfxd;
/*if (nb < sminmn) {
j = nfxd;
// Set the original matrix to the GPU
magma_ssetmatrix_async( m, sn,
A (0,j), lda,
dA(0,j), ldda, stream[0] );
}*/
/* Initialize partial column norms. */
magmablas_snrm2_cols(sm, sn, A(nfxd,nfxd), lda, &rwork[nfxd]);
#if defined(PRECISION_d) || defined(PRECISION_z)
magma_dcopymatrix( sn, 1, &rwork[nfxd], sn, &rwork[n+nfxd], sn);
#else
magma_scopymatrix( sn, 1, &rwork[nfxd], sn, &rwork[n+nfxd], sn);
#endif
/*for (j = nfxd; j < n; ++j) {
rwork[j] = cblas_snrm2(sm, A(nfxd, j), ione);
rwork[n + j] = rwork[j];
}*/
j = nfxd;
//if (nb < sminmn)
{
/* Use blocked code initially. */
//magma_queue_sync( stream[0] );
/* Compute factorization: while loop. */
topbmn = minmn;// - nb;
while(j < topbmn) {
jb = min(nb, topbmn - j);
/* Factorize JB columns among columns J:N. */
n_j = n - j;
/*if (j>nfxd) {
// Get panel to the CPU
magma_sgetmatrix( m-j, jb,
dA(j,j), ldda,
A (j,j), lda );
// Get the rows
magma_sgetmatrix( jb, n_j - jb,
dA(j,j + jb), ldda,
A (j,j + jb), lda );
}*/
//magma_slaqps_gpu // this is a cpp-file
magma_slaqps2_gpu // this is a cuda-file
( m, n_j, j, jb, &fjb,
A (0, j), lda,
&jpvt[j], &tau[j], &rwork[j], &rwork[n + j],
work,
&df[jb], n_j );
j += fjb; /* fjb is actual number of columns factored */
}
}
/* Use unblocked code to factor the last or only block.
if (j < minmn) {
n_j = n - j;
if (j > nfxd) {
magma_sgetmatrix( m-j, n_j,
dA(j,j), ldda,
A (j,j), lda );
}
lapackf77_slaqp2(&m, &n_j, &j, A(0, j), &lda, &jpvt[j],
&tau[j], &rwork[j], &rwork[n+j], work );
}*/
}
//work[0] = MAGMA_S_MAKE( lwkopt, 0. );
magma_free(df);
return *info;
} /* sgeqp3 */
