extern "C" magma_int_t
magma_dgeqp3( magma_int_t m, magma_int_t n,
double *A, magma_int_t lda,
magma_int_t *jpvt, double *tau,
double *work, magma_int_t lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
double *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
=======
DGEQP3 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) DOUBLE_PRECISION 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) DOUBLE_PRECISION array, dimension (min(M,N))
The scalar factors of the elementary reflectors.
WORK (workspace/output) DOUBLE_PRECISION 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 ))
#define dA(i, j) (dwork + (i) + (j)*(ldda))
double *dwork, *df;
magma_int_t ione = 1;
magma_int_t n_j, ldda, ldwork;
magma_int_t j, jb, na, 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_dgeqp3_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_D_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)
double *rwork = work + (n + 1)*nb;
#endif
ldda = ((m+31)/32)*32;
ldwork = n*ldda + (n+1)*nb;
if (MAGMA_SUCCESS != magma_dmalloc( &dwork, ldwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
df = dwork + n*ldda;
// dwork used for dA
magma_queue_t stream;
magma_queue_create( &stream );
/* Move initial columns up front.
* Note jpvt uses 1-based indices for historical compatibility. */
nfxd = 0;
for (j = 0; j < n; ++j) {
if (jpvt[j] != 0) {
if (j != nfxd) {
blasf77_dswap(&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_dgeqrf(&m, &na, A, &lda, tau, work, &lwork, info);
if (na < n) {
n_j = n - na;
lapackf77_dormqr( 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_dsetmatrix_async( m, sn,
A (0,j), lda,
dA(0,j), ldda, stream );
}
/* Initialize partial column norms. */
for (j = nfxd; j < n; ++j) {
rwork[j] = cblas_dnrm2(sm, A(nfxd, j), ione);
rwork[n + j] = rwork[j];
}
j = nfxd;
if (nb < sminmn) {
/* Use blocked code initially. */
magma_queue_sync( stream );
/* 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_dgetmatrix( m-j, jb,
dA(j,j), ldda,
A (j,j), lda );
// Get the rows
magma_dgetmatrix( jb, n_j - jb,
dA(j,j + jb), ldda,
A (j,j + jb), lda );
}
magma_dlaqps( m, n_j, j, jb, &fjb,
A (0, j), lda,
dA(0, j), ldda,
&jpvt[j], &tau[j], &rwork[j], &rwork[n + j],
work,
&work[jb], n_j,
&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_dgetmatrix( m-j, n_j,
dA(j,j), ldda,
A (j,j), lda );
}
lapackf77_dlaqp2(&m, &n_j, &j, A(0, j), &lda, &jpvt[j],
&tau[j], &rwork[j], &rwork[n+j], work );
}
}
work[0] = MAGMA_D_MAKE( lwkopt, 0. );
magma_free( dwork );
magma_queue_destroy( stream );
return *info;
} /* dgeqp3 */
/**
Purpose
-------
DGEQP3 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]
A DOUBLE_PRECISION 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.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= 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 DOUBLE_PRECISION array, dimension (min(M,N))
The scalar factors of the elementary reflectors.
@param[out]
work (workspace) DOUBLE_PRECISION array, dimension (MAX(1,LWORK))
On exit, if INFO=0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK.
For [sd]geqp3, LWORK >= (N+1)*NB + 2*N;
for [cz]geqp3, LWORK >= (N+1)*NB,
where NB is the optimal blocksize.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param
rwork (workspace, for [cz]geqp3 only) DOUBLE PRECISION array, dimension (2*N)
@param[out]
info INTEGER
- = 0: successful exit.
- < 0: if INFO = -i, the i-th argument had an illegal value.
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).
@ingroup magma_dgeqp3_comp
********************************************************************/
extern "C" magma_int_t
magma_dgeqp3(
magma_int_t m, magma_int_t n,
double *A, magma_int_t lda,
magma_int_t *jpvt, double *tau,
double *work, magma_int_t lwork,
#ifdef COMPLEX
double *rwork,
#endif
magma_int_t *info )
{
#define A(i, j) (A + (i) + (j)*(lda ))
#define dA(i, j) (dwork + (i) + (j)*(ldda))
double *dwork, *df;
magma_int_t ione = 1;
magma_int_t n_j, ldda, ldwork;
magma_int_t j, jb, na, nb, sm, sn, fjb, nfxd, minmn;
magma_int_t topbmn, sminmn, lwkopt=0, 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_dgeqp3_nb(min(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
}
work[0] = MAGMA_D_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;
#ifdef REAL
double *rwork = work + (n + 1)*nb;
#endif
ldda = ((m+31)/32)*32;
ldwork = n*ldda + (n+1)*nb;
if (MAGMA_SUCCESS != magma_dmalloc( &dwork, ldwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
df = dwork + n*ldda;
// dwork used for dA
magma_queue_t stream;
magma_queue_create( &stream );
/* Move initial columns up front.
* Note jpvt uses 1-based indices for historical compatibility. */
nfxd = 0;
for (j = 0; j < n; ++j) {
if (jpvt[j] != 0) {
if (j != nfxd) {
blasf77_dswap(&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_dgeqrf(&m, &na, A, &lda, tau, work, &lwork, info);
if (na < n) {
n_j = n - na;
lapackf77_dormqr( MagmaLeftStr, MagmaConjTransStr, &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_dsetmatrix_async( m, sn,
A (0,j), lda,
dA(0,j), ldda, stream );
}
/* Initialize partial column norms. */
for (j = nfxd; j < n; ++j) {
rwork[j] = magma_cblas_dnrm2( sm, A(nfxd,j), ione );
rwork[n + j] = rwork[j];
}
j = nfxd;
if (nb < sminmn) {
/* Use blocked code initially. */
magma_queue_sync( stream );
/* 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_dgetmatrix( m-j, jb,
dA(j,j), ldda,
A (j,j), lda );
// Get the rows
magma_dgetmatrix( jb, n_j - jb,
dA(j,j + jb), ldda,
A (j,j + jb), lda );
}
magma_dlaqps( m, n_j, j, jb, &fjb,
A (0, j), lda,
dA(0, j), ldda,
&jpvt[j], &tau[j], &rwork[j], &rwork[n + j],
work,
&work[jb], n_j,
&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_dgetmatrix( m-j, n_j,
dA(j,j), ldda,
A (j,j), lda );
}
lapackf77_dlaqp2(&m, &n_j, &j, A(0, j), &lda, &jpvt[j],
&tau[j], &rwork[j], &rwork[n+j], work );
}
}
work[0] = MAGMA_D_MAKE( lwkopt, 0. );
magma_free( dwork );
magma_queue_destroy( stream );
return *info;
} /* magma_dgeqp3 */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dswap, dswapblk, dpermute, dlaswp, dlaswpx
*/
int main( int argc, char** argv)
{
TESTING_INIT();
double *h_A1, *h_A2;
double *d_A1, *d_A2;
double *h_R1, *h_R2;
// row-major and column-major performance
real_Double_t row_perf0, col_perf0;
real_Double_t row_perf1, col_perf1;
real_Double_t row_perf2, col_perf2;
real_Double_t row_perf3;
real_Double_t row_perf4;
real_Double_t row_perf5, col_perf5;
real_Double_t row_perf6, col_perf6;
real_Double_t row_perf7;
real_Double_t cpu_perf;
real_Double_t time, gbytes;
magma_int_t N, lda, ldda, nb, j;
magma_int_t ione = 1;
magma_int_t *ipiv, *ipiv2;
magma_int_t *d_ipiv;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_queue_t queue = 0;
printf(" cublasDswap dswap dswapblk dlaswp dpermute dlaswp2 dlaswpx dcopymatrix 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 dlaswp (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_dgetrf_nb( N ));
nb = min( N, nb );
// each swap does 2N loads and 2N stores, for nb swaps
gbytes = sizeof(double) * 4.*N*nb / 1e9;
TESTING_MALLOC_PIN( h_A1, double, lda*N );
TESTING_MALLOC_PIN( h_A2, double, lda*N );
TESTING_MALLOC_PIN( h_R1, double, lda*N );
TESTING_MALLOC_PIN( h_R2, double, 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, double, ldda*N );
TESTING_MALLOC_DEV( d_A2, double, ldda*N );
for( j=0; j < nb; j++ ) {
ipiv[j] = (magma_int_t) ((rand()*1.*N) / (RAND_MAX * 1.)) + 1;
}
/* =====================================================================
* cublasDswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_dsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
cublasDswap( 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_dswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_dgetmatrix( 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_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_dsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
cublasDswap( 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_dswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_dgetmatrix( 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;
/* =====================================================================
* dswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_dsetmatrix( 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_dswap( 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_dswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_dgetmatrix( 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_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_dsetmatrix( 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_dswap( 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_dswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_dgetmatrix( 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;
/* =====================================================================
* dswapblk, blocked version (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_dsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_dswapblk( 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_dswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_dgetmatrix( 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_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_dsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_dswapblk( 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_dswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_dgetmatrix( 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;
/* =====================================================================
* dpermute_long (1 matrix)
*/
/* Row Major */
memcpy( ipiv2, ipiv, nb*sizeof(magma_int_t) ); // dpermute updates ipiv2
init_matrix( N, N, h_A1, lda, 0 );
magma_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_dpermute_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_dswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style dlaswp (1 matrix)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_dlaswp( 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_dswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style dlaswp (1 matrix) - d_ipiv on GPU
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_dsetmatrix( 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_dlaswp2( 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_dswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style dlaswpx (extended for row- and col-major) (1 matrix)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_dlaswpx( 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_dswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_dgetmatrix( 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_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_dlaswpx( N, d_A1, 1, ldda, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
col_perf5 = gbytes / time;
time = magma_wtime();
lapackf77_dlaswp( &N, h_A1, &lda, &ione, &nb, ipiv, &ione);
time = magma_wtime() - time;
cpu_perf = gbytes / time;
magma_dgetmatrix( 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_dcopymatrix( 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_dcopymatrix( nb, N, d_A1, ldda, d_A2, ldda );
time = magma_sync_wtime( queue ) - time;
// copy reads 1 matrix and writes 1 matrix, so has half gbytes of swap
row_perf6 = 0.5 * gbytes / time;
printf("%5d %3d %6.2f%c/ %6.2f%c %6.2f%c/ %6.2f%c %6.2f%c/ %6.2f%c %6.2f%c %6.2f%c %6.2f%c %6.2f%c/ %6.2f%c %6.2f / %6.2f %6.2f %10s\n",
(int) N, (int) nb,
row_perf0, ((check & 0x001) != 0 ? '*' : ' '),
col_perf0, ((check & 0x002) != 0 ? '*' : ' '),
row_perf1, ((check & 0x004) != 0 ? '*' : ' '),
col_perf1, ((check & 0x008) != 0 ? '*' : ' '),
row_perf2, ((check & 0x010) != 0 ? '*' : ' '),
col_perf2, ((check & 0x020) != 0 ? '*' : ' '),
row_perf3, ((check & 0x040) != 0 ? '*' : ' '),
row_perf4, ((check & 0x080) != 0 ? '*' : ' '),
row_perf7, ((check & 0x100) != 0 ? '*' : ' '),
row_perf5, ((check & 0x200) != 0 ? '*' : ' '),
col_perf5, ((check & 0x400) != 0 ? '*' : ' '),
row_perf6,
col_perf6,
cpu_perf,
(check == 0 ? "ok" : "* failed") );
status += ! (check == 0);
TESTING_FREE_PIN( h_A1 );
TESTING_FREE_PIN( h_A2 );
TESTING_FREE_PIN( h_R1 );
TESTING_FREE_PIN( h_R2 );
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( ipiv2 );
TESTING_FREE_DEV( d_ipiv );
TESTING_FREE_DEV( d_A1 );
TESTING_FREE_DEV( d_A2 );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/**
Purpose
-------
DSTEDX computes some eigenvalues and, optionally, eigenvectors of a
symmetric tridiagonal matrix using the divide and conquer method.
This code 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. See DLAEX3 for details.
Arguments
---------
@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]
n INTEGER
The dimension of the symmetric tridiagonal matrix. N >= 0.
@param[in]
vl DOUBLE PRECISION
@param[in]
vu DOUBLE PRECISION
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[in,out]
d DOUBLE PRECISION array, dimension (N)
On entry, the diagonal elements of the tridiagonal matrix.
On exit, if INFO = 0, the eigenvalues in ascending order.
@param[in,out]
e DOUBLE PRECISION array, dimension (N-1)
On entry, the subdiagonal elements of the tridiagonal matrix.
On exit, E has been destroyed.
@param[in,out]
Z DOUBLE PRECISION array, dimension (LDZ,N)
On exit, if INFO = 0, Z contains the orthonormal eigenvectors
of the symmetric tridiagonal matrix.
@param[in]
ldz INTEGER
The leading dimension of the array Z. LDZ >= max(1,N).
@param[out]
work (workspace) DOUBLE PRECISION array, dimension (LWORK)
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK.
If N > 1 then LWORK >= ( 1 + 4*N + N**2 ).
Note that if N is less than or
equal to the minimum divide size, usually 25, then LWORK need
only be max(1,2*(N-1)).
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
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.
LIWORK >= ( 3 + 5*N ).
Note that if N is less than or
equal to the minimum divide size, usually 25, then LIWORK
need only be 1.
\n
If LIWORK = -1, then a workspace query is assumed; the
routine only calculates the optimal size of the IWORK array,
returns this value as the first entry of the IWORK array, and
no error message related to LIWORK is issued by XERBLA.
@param
dwork (workspace) DOUBLE PRECISION array, dimension (3*N*N/2+3*N)
@param[out]
info INTEGER
- = 0: successful exit.
- < 0: if INFO = -i, the i-th argument had an illegal value.
- > 0: 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 by Francoise Tisseur, University of Tennessee.
@ingroup magma_dsyev_comp
********************************************************************/
extern "C" magma_int_t
magma_dstedx(
magma_range_t range, magma_int_t n, double vl, double vu,
magma_int_t il, magma_int_t iu, double *d, double *e,
double *Z, magma_int_t ldz,
double *work, magma_int_t lwork,
magma_int_t *iwork, magma_int_t liwork,
magmaDouble_ptr dwork,
magma_int_t *info)
{
#define Z(i_,j_) (Z + (i_) + (j_)*ldz)
double d_zero = 0.;
double d_one = 1.;
magma_int_t izero = 0;
magma_int_t ione = 1;
magma_int_t alleig, indeig, valeig, lquery;
magma_int_t i, j, k, m;
magma_int_t liwmin, lwmin;
magma_int_t start, end, smlsiz;
double eps, orgnrm, p, tiny;
// Test the input parameters.
alleig = (range == MagmaRangeAll);
valeig = (range == MagmaRangeV);
indeig = (range == MagmaRangeI);
lquery = (lwork == -1 || liwork == -1);
*info = 0;
if (! (alleig || valeig || indeig)) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (ldz < max(1,n)) {
*info = -10;
} else {
if (valeig) {
if (n > 0 && vu <= vl) {
*info = -4;
}
} else if (indeig) {
if (il < 1 || il > max(1,n)) {
*info = -5;
} else if (iu < min(n,il) || iu > n) {
*info = -6;
}
}
}
if (*info == 0) {
// Compute the workspace requirements
smlsiz = magma_get_smlsize_divideconquer();
if ( n <= 1 ) {
lwmin = 1;
liwmin = 1;
} else {
lwmin = 1 + 4*n + n*n;
liwmin = 3 + 5*n;
}
work[0] = magma_dmake_lwork( lwmin );
iwork[0] = liwmin;
if (lwork < lwmin && ! lquery) {
*info = -12;
} else if (liwork < liwmin && ! lquery) {
*info = -14;
}
}
if (*info != 0) {
magma_xerbla( __func__, -(*info));
return *info;
} else if (lquery) {
return *info;
}
// Quick return if possible
if (n == 0)
return *info;
if (n == 1) {
*Z = 1.;
return *info;
}
/* determine the number of threads *///not needed here to be checked Azzam
//magma_int_t threads = magma_get_parallel_numthreads();
//magma_int_t mklth = magma_get_lapack_numthreads();
//magma_set_lapack_numthreads(mklth);
#ifdef ENABLE_DEBUG
//printf(" D&C is using %d threads\n", threads);
#endif
// If N is smaller than the minimum divide size (SMLSIZ+1), then
// solve the problem with another solver.
if (n < smlsiz) {
lapackf77_dsteqr("I", &n, d, e, Z, &ldz, work, info);
} else {
lapackf77_dlaset("F", &n, &n, &d_zero, &d_one, Z, &ldz);
//Scale.
orgnrm = lapackf77_dlanst("M", &n, d, e);
if (orgnrm == 0) {
work[0] = magma_dmake_lwork( lwmin );
iwork[0] = liwmin;
return *info;
}
eps = lapackf77_dlamch( "Epsilon" );
if (alleig) {
start = 0;
while ( start < n ) {
// Let FINISH be the position of the next subdiagonal entry
// such that E( END ) <= TINY or FINISH = N if no such
// subdiagonal exists. The matrix identified by the elements
// between START and END constitutes an independent
// sub-problem.
for (end = start+1; end < n; ++end) {
tiny = eps * sqrt( MAGMA_D_ABS(d[end-1]*d[end]));
if (MAGMA_D_ABS(e[end-1]) <= tiny)
break;
}
// (Sub) Problem determined. Compute its size and solve it.
m = end - start;
if (m == 1) {
start = end;
continue;
}
if (m > smlsiz) {
// Scale
orgnrm = lapackf77_dlanst("M", &m, &d[start], &e[start]);
lapackf77_dlascl("G", &izero, &izero, &orgnrm, &d_one, &m, &ione, &d[start], &m, info);
magma_int_t mm = m-1;
lapackf77_dlascl("G", &izero, &izero, &orgnrm, &d_one, &mm, &ione, &e[start], &mm, info);
magma_dlaex0( m, &d[start], &e[start], Z(start, start), ldz, work, iwork, dwork, MagmaRangeAll, vl, vu, il, iu, info);
if ( *info != 0) {
return *info;
}
// Scale Back
lapackf77_dlascl("G", &izero, &izero, &d_one, &orgnrm, &m, &ione, &d[start], &m, info);
} else {
lapackf77_dsteqr( "I", &m, &d[start], &e[start], Z(start, start), &ldz, work, info);
if (*info != 0) {
*info = (start+1) *(n+1) + end;
}
}
start = end;
}
// If the problem split any number of times, then the eigenvalues
// will not be properly ordered. Here we permute the eigenvalues
// (and the associated eigenvectors) into ascending order.
if (m < n) {
// Use Selection Sort to minimize swaps of eigenvectors
for (i = 1; i < n; ++i) {
k = i-1;
p = d[i-1];
for (j = i; j < n; ++j) {
if (d[j] < p) {
k = j;
p = d[j];
}
}
if (k != i-1) {
d[k] = d[i-1];
d[i-1] = p;
blasf77_dswap(&n, Z(0,i-1), &ione, Z(0,k), &ione);
}
}
}
} else {
// Scale
lapackf77_dlascl("G", &izero, &izero, &orgnrm, &d_one, &n, &ione, d, &n, info);
magma_int_t nm = n-1;
lapackf77_dlascl("G", &izero, &izero, &orgnrm, &d_one, &nm, &ione, e, &nm, info);
magma_dlaex0( n, d, e, Z, ldz, work, iwork, dwork, range, vl, vu, il, iu, info);
if ( *info != 0) {
return *info;
}
// Scale Back
lapackf77_dlascl("G", &izero, &izero, &d_one, &orgnrm, &n, &ione, d, &n, info);
}
}
work[0] = magma_dmake_lwork( lwmin );
iwork[0] = liwmin;
return *info;
} /* magma_dstedx */