extern "C" magma_int_t
magma_ccg(
magma_c_matrix A, magma_c_matrix b, magma_c_matrix *x,
magma_c_solver_par *solver_par,
magma_queue_t queue )
{
magma_int_t info = 0;
// set queue for old dense routines
magma_queue_t orig_queue=NULL;
magmablasGetKernelStream( &orig_queue );
// prepare solver feedback
solver_par->solver = Magma_CG;
solver_par->numiter = 0;
solver_par->info = MAGMA_SUCCESS;
// local variables
magmaFloatComplex c_zero = MAGMA_C_ZERO, c_one = MAGMA_C_ONE;
magma_int_t dofs = A.num_rows * b.num_cols;
// GPU workspace
magma_c_matrix r={Magma_CSR}, p={Magma_CSR}, q={Magma_CSR};
CHECK( magma_cvinit( &r, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_cvinit( &p, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_cvinit( &q, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
// solver variables
magmaFloatComplex alpha, beta;
float nom, nom0, r0, betanom, betanomsq, den;
// solver setup
CHECK( magma_cresidualvec( A, b, *x, &r, &nom0, queue));
magma_ccopy( dofs, r.dval, 1, p.dval, 1 ); // p = r
betanom = nom0;
nom = nom0 * nom0; // nom = r' * r
CHECK( magma_c_spmv( c_one, A, p, c_zero, q, queue )); // q = A p
den = MAGMA_C_REAL( magma_cdotc(dofs, p.dval, 1, q.dval, 1) );// den = p dot q
solver_par->init_res = nom0;
if ( (r0 = nom * solver_par->epsilon) < ATOLERANCE )
r0 = ATOLERANCE;
if ( nom < r0 ) {
solver_par->final_res = solver_par->init_res;
solver_par->iter_res = solver_par->init_res;
goto cleanup;
}
// check positive definite
if (den <= 0.0) {
printf("Operator A is not postive definite. (Ar,r) = %f\n", den);
magmablasSetKernelStream( orig_queue );
info = MAGMA_NONSPD;
goto cleanup;
}
//Chronometry
real_Double_t tempo1, tempo2;
tempo1 = magma_sync_wtime( queue );
if ( solver_par->verbose > 0 ) {
solver_par->res_vec[0] = (real_Double_t)nom0;
solver_par->timing[0] = 0.0;
}
solver_par->numiter = 0;
// start iteration
do
{
solver_par->numiter++;
alpha = MAGMA_C_MAKE(nom/den, 0.);
magma_caxpy(dofs, alpha, p.dval, 1, x->dval, 1); // x = x + alpha p
magma_caxpy(dofs, -alpha, q.dval, 1, r.dval, 1); // r = r - alpha q
betanom = magma_scnrm2(dofs, r.dval, 1); // betanom = || r ||
betanomsq = betanom * betanom; // betanoms = r' * r
if ( solver_par->verbose > 0 ) {
tempo2 = magma_sync_wtime( queue );
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
if ( betanom < r0 ) {
break;
}
beta = MAGMA_C_MAKE(betanomsq/nom, 0.); // beta = betanoms/nom
magma_cscal(dofs, beta, p.dval, 1); // p = beta*p
magma_caxpy(dofs, c_one, r.dval, 1, p.dval, 1); // p = p + r
CHECK( magma_c_spmv( c_one, A, p, c_zero, q, queue )); // q = A p
den = MAGMA_C_REAL(magma_cdotc(dofs, p.dval, 1, q.dval, 1));
// den = p dot q
nom = betanomsq;
}
while ( solver_par->numiter+1 <= solver_par->maxiter );
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t) tempo2-tempo1;
float residual;
CHECK( magma_cresidualvec( A, b, *x, &r, &residual, queue));
solver_par->final_res = residual;
if ( solver_par->numiter < solver_par->maxiter ) {
solver_par->info = MAGMA_SUCCESS;
} else if ( solver_par->init_res > solver_par->final_res ) {
if ( solver_par->verbose > 0 ) {
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
info = MAGMA_SLOW_CONVERGENCE;
if( solver_par->iter_res < solver_par->epsilon*solver_par->init_res ){
info = MAGMA_SUCCESS;
}
}
else {
if ( solver_par->verbose > 0 ) {
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
info = MAGMA_DIVERGENCE;
}
cleanup:
magma_cmfree(&r, queue );
magma_cmfree(&p, queue );
magma_cmfree(&q, queue );
magmablasSetKernelStream( orig_queue );
solver_par->info = info;
return info;
} /* magma_ccg */
extern "C" magma_int_t
magma_cbicgstab_merge(
magma_c_sparse_matrix A, magma_c_vector b,
magma_c_vector *x, magma_c_solver_par *solver_par,
magma_queue_t queue )
{
// set queue for old dense routines
magma_queue_t orig_queue;
magmablasGetKernelStream( &orig_queue );
// prepare solver feedback
solver_par->solver = Magma_BICGSTABMERGE;
solver_par->numiter = 0;
solver_par->info = MAGMA_SUCCESS;
// some useful variables
magmaFloatComplex c_zero = MAGMA_C_ZERO, c_one = MAGMA_C_ONE;
magma_int_t dofs = A.num_rows;
// GPU stream
magma_queue_t stream[2];
magma_event_t event[1];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
magma_event_create( &event[0] );
// workspace
magma_c_vector q, r,rr,p,v,s,t;
magmaFloatComplex *d1, *d2, *skp;
d1 = NULL;
d2 = NULL;
skp = NULL;
magma_int_t stat_dev = 0, stat_cpu = 0;
stat_dev += magma_cmalloc( &d1, dofs*(2) );
stat_dev += magma_cmalloc( &d2, dofs*(2) );
// array for the parameters
stat_dev += magma_cmalloc( &skp, 8 );
if( stat_dev != 0 ){
magma_free( d1 );
magma_free( d2 );
magma_free( skp );
printf("error: memory allocation.\n");
return MAGMA_ERR_DEVICE_ALLOC;
}
// skp = [alpha|beta|omega|rho_old|rho|nom|tmp1|tmp2]
magma_c_vinit( &q, Magma_DEV, dofs*6, c_zero, queue );
// q = rr|r|p|v|s|t
rr.memory_location = Magma_DEV; rr.dval = NULL; rr.num_rows = rr.nnz = dofs; rr.num_cols = 1;
r.memory_location = Magma_DEV; r.dval = NULL; r.num_rows = r.nnz = dofs; r.num_cols = 1;
p.memory_location = Magma_DEV; p.dval = NULL; p.num_rows = p.nnz = dofs; p.num_cols = 1;
v.memory_location = Magma_DEV; v.dval = NULL; v.num_rows = v.nnz = dofs; v.num_cols = 1;
s.memory_location = Magma_DEV; s.dval = NULL; s.num_rows = s.nnz = dofs; s.num_cols = 1;
t.memory_location = Magma_DEV; t.dval = NULL; t.num_rows = t.nnz = dofs; t.num_cols = 1;
rr.dval = q(0);
r.dval = q(1);
p.dval = q(2);
v.dval = q(3);
s.dval = q(4);
t.dval = q(5);
// solver variables
magmaFloatComplex alpha, beta, omega, rho_old, rho_new, *skp_h;
float nom, nom0, betanom, r0, den;
// solver setup
magma_cscal( dofs, c_zero, x->dval, 1) ; // x = 0
magma_ccopy( dofs, b.dval, 1, q(0), 1 ); // rr = b
magma_ccopy( dofs, b.dval, 1, q(1), 1 ); // r = b
rho_new = magma_cdotc( dofs, r.dval, 1, r.dval, 1 ); // rho=<rr,r>
nom = MAGMA_C_REAL(magma_cdotc( dofs, r.dval, 1, r.dval, 1 ));
nom0 = betanom = sqrt(nom); // nom = || r ||
rho_old = omega = alpha = MAGMA_C_MAKE( 1.0, 0. );
beta = rho_new;
solver_par->init_res = nom0;
// array on host for the parameters
stat_cpu = magma_cmalloc_cpu( &skp_h, 8 );
if( stat_cpu != 0 ){
magma_free( d1 );
magma_free( d2 );
magma_free( skp );
magma_free_cpu( skp_h );
printf("error: memory allocation.\n");
return MAGMA_ERR_HOST_ALLOC;
}
skp_h[0]=alpha;
skp_h[1]=beta;
skp_h[2]=omega;
skp_h[3]=rho_old;
skp_h[4]=rho_new;
skp_h[5]=MAGMA_C_MAKE(nom, 0.0);
magma_csetvector( 8, skp_h, 1, skp, 1 );
magma_c_spmv( c_one, A, r, c_zero, v, queue ); // z = A r
den = MAGMA_C_REAL( magma_cdotc(dofs, v.dval, 1, r.dval, 1) );// den = z dot r
if ( (r0 = nom * solver_par->epsilon) < ATOLERANCE )
r0 = ATOLERANCE;
if ( nom < r0 ) {
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
}
// check positive definite
if (den <= 0.0) {
printf("Operator A is not postive definite. (Ar,r) = %f\n", den);
magmablasSetKernelStream( orig_queue );
return MAGMA_NONSPD;
solver_par->info = MAGMA_NONSPD;;
}
//Chronometry
real_Double_t tempo1, tempo2;
tempo1 = magma_sync_wtime( queue );
if ( solver_par->verbose > 0 ) {
solver_par->res_vec[0] = nom0;
solver_par->timing[0] = 0.0;
}
// start iteration
for( solver_par->numiter= 1; solver_par->numiter<solver_par->maxiter;
solver_par->numiter++ ) {
magmablasSetKernelStream(stream[0]);
// computes p=r+beta*(p-omega*v)
magma_cbicgmerge1( dofs, skp, v.dval, r.dval, p.dval, queue );
magma_c_spmv( c_one, A, p, c_zero, v, queue ); // v = Ap
magma_cmdotc( dofs, 1, q.dval, v.dval, d1, d2, skp, queue );
magma_cbicgmerge4( 1, skp, queue );
magma_cbicgmerge2( dofs, skp, r.dval, v.dval, s.dval, queue ); // s=r-alpha*v
magma_c_spmv( c_one, A, s, c_zero, t, queue ); // t=As
magma_cmdotc( dofs, 2, q.dval+4*dofs, t.dval, d1, d2, skp+6, queue );
magma_cbicgmerge4( 2, skp, queue );
magma_cbicgmerge_xrbeta( dofs, d1, d2, q.dval, r.dval, p.dval,
s.dval, t.dval, x->dval, skp, queue );
// check stopping criterion (asynchronous copy)
magma_cgetvector_async( 1 , skp+5, 1,
skp_h+5, 1, stream[1] );
betanom = sqrt(MAGMA_C_REAL(skp_h[5]));
if ( solver_par->verbose > 0 ) {
tempo2 = magma_sync_wtime( queue );
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
if ( betanom < r0 ) {
break;
}
}
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t) tempo2-tempo1;
float residual;
magma_cresidual( A, b, *x, &residual, queue );
solver_par->iter_res = betanom;
solver_par->final_res = residual;
if ( solver_par->numiter < solver_par->maxiter) {
solver_par->info = MAGMA_SUCCESS;
} else if ( solver_par->init_res > solver_par->final_res ) {
if ( solver_par->verbose > 0 ) {
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
solver_par->info = MAGMA_SLOW_CONVERGENCE;
}
else {
if ( solver_par->verbose > 0 ) {
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
solver_par->info = MAGMA_DIVERGENCE;
}
magma_c_vfree(&q, queue ); // frees all vectors
magma_free(d1);
magma_free(d2);
magma_free( skp );
magma_free_cpu( skp_h );
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
} /* cbicgstab_merge */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing claset
Code is very similar to testing_clacpy.cpp
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gbytes, gpu_perf, gpu_time, cpu_perf, cpu_time;
float error, work[1];
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *h_A, *h_R;
magmaFloatComplex *d_A;
magmaFloatComplex offdiag = MAGMA_C_MAKE( 1.2000, 6.7000 );
magmaFloatComplex diag = MAGMA_C_MAKE( 3.1415, 2.7183 );
magma_int_t M, N, size, lda, ldb, ldda;
magma_int_t ione = 1;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_uplo_t uplo[] = { MagmaLower, MagmaUpper, MagmaFull };
printf("uplo M N CPU GByte/s (ms) GPU GByte/s (ms) check\n");
printf("==================================================================\n");
for( int iuplo = 0; iuplo < 3; ++iuplo ) {
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
//M += 2; // space for insets
//N += 2;
lda = M;
ldb = lda;
ldda = ((M+31)/32)*32;
size = lda*N;
if ( uplo[iuplo] == MagmaLower || uplo[iuplo] == MagmaUpper ) {
// save triangle (with diagonal)
// TODO wrong for trapezoid
gbytes = sizeof(magmaFloatComplex) * 0.5*N*(N+1) / 1e9;
}
else {
// save entire matrix
gbytes = sizeof(magmaFloatComplex) * 1.*M*N / 1e9;
}
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, size );
TESTING_MALLOC_CPU( h_R, magmaFloatComplex, size );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N );
/* Initialize the matrix */
for( int j = 0; j < N; ++j ) {
for( int i = 0; i < M; ++i ) {
h_A[i + j*lda] = MAGMA_C_MAKE( i + j/10000., j );
}
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_csetmatrix( M, N, h_A, lda, d_A, ldda );
gpu_time = magma_sync_wtime( 0 );
//magmablas_claset( uplo[iuplo], M-2, N-2, offdiag, diag, d_A+1+ldda, ldda ); // inset by 1 row & col
magmablas_claset( uplo[iuplo], M, N, offdiag, diag, d_A, ldda );
gpu_time = magma_sync_wtime( 0 ) - gpu_time;
gpu_perf = gbytes / gpu_time;
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
//magma_int_t M2 = M-2; // inset by 1 row & col
//magma_int_t N2 = N-2;
//lapackf77_claset( lapack_uplo_const( uplo[iuplo] ), &M2, &N2, &offdiag, &diag, h_A+1+lda, &lda );
lapackf77_claset( lapack_uplo_const( uplo[iuplo] ), &M, &N, &offdiag, &diag, h_A, &lda );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gbytes / cpu_time;
/* =====================================================================
Check the result
=================================================================== */
magma_cgetmatrix( M, N, d_A, ldda, h_R, lda );
blasf77_caxpy(&size, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_clange("f", &M, &N, h_R, &lda, work);
printf("%4c %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %s\n",
lapacke_uplo_const( uplo[iuplo] ), (int) M, (int) N,
cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000.,
(error == 0. ? "ok" : "failed") );
status += ! (error == 0.);
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_R );
TESTING_FREE_DEV( d_A );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
printf( "\n" );
}
TESTING_FINALIZE();
return status;
}
magma_int_t
magma_cm_5stencil(
magma_int_t n,
magma_c_matrix *A,
magma_queue_t queue )
{
magma_int_t info = 0;
magma_int_t i,j,k;
magma_c_matrix hA={Magma_CSR};
// generate matrix of desired structure and size (2d 5-point stencil)
magma_int_t nn = n*n;
magma_int_t offdiags = 2;
magma_index_t *diag_offset=NULL;
magmaFloatComplex *diag_vals=NULL;
CHECK( magma_cmalloc_cpu( &diag_vals, offdiags+1 ));
CHECK( magma_index_malloc_cpu( &diag_offset, offdiags+1 ));
diag_offset[0] = 0;
diag_offset[1] = 1;
diag_offset[2] = n;
#define COMPLEX
#ifdef COMPLEX
// complex case
diag_vals[0] = MAGMA_C_MAKE( 4.0, 4.0 );
diag_vals[1] = MAGMA_C_MAKE( -1.0, -1.0 );
diag_vals[2] = MAGMA_C_MAKE( -1.0, -1.0 );
#else
// real case
diag_vals[0] = MAGMA_C_MAKE( 4.0, 0.0 );
diag_vals[1] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[2] = MAGMA_C_MAKE( -1.0, 0.0 );
#endif
CHECK( magma_cmgenerator( nn, offdiags, diag_offset, diag_vals, &hA, queue ));
// now set some entries to zero (boundary...)
for( i=0; i<n; i++ ) {
for( j=0; j<n; j++ ) {
magma_index_t row = i*n+j;
for( k=hA.row[row]; k<hA.row[row+1]; k++) {
if ((hA.col[k] == row-1 ) && (row+1)%n == 1 )
hA.val[k] = MAGMA_C_MAKE( 0.0, 0.0 );
if ((hA.col[k] == row+1 ) && (row)%n ==n-1 )
hA.val[k] = MAGMA_C_MAKE( 0.0, 0.0 );
}
}
}
CHECK( magma_cmconvert( hA, A, Magma_CSR, Magma_CSR, queue ));
magma_cmcsrcompressor( A, queue );
A->true_nnz = A->nnz;
cleanup:
magma_free_cpu( diag_vals );
magma_free_cpu( diag_offset );
magma_cmfree( &hA, queue );
return info;
}
/* ////////////////////////////////////////////////////////////////////////////
-- testing sparse matrix vector product
*/
int main( int argc, char** argv )
{
TESTING_INIT();
magma_queue_t queue;
magma_queue_create( /*devices[ opts->device ],*/ &queue );
magma_c_sparse_matrix hA, hA_SELLP, hA_ELL, dA, dA_SELLP, dA_ELL;
hA_SELLP.blocksize = 8;
hA_SELLP.alignment = 8;
real_Double_t start, end, res;
magma_int_t *pntre;
magmaFloatComplex c_one = MAGMA_C_MAKE(1.0, 0.0);
magmaFloatComplex c_zero = MAGMA_C_MAKE(0.0, 0.0);
magma_int_t i, j;
for( i = 1; i < argc; ++i ) {
if ( strcmp("--blocksize", argv[i]) == 0 ) {
hA_SELLP.blocksize = atoi( argv[++i] );
} else if ( strcmp("--alignment", argv[i]) == 0 ) {
hA_SELLP.alignment = atoi( argv[++i] );
} else
break;
}
printf( "\n# usage: ./run_cspmv"
" [ --blocksize %d --alignment %d (for SELLP) ]"
" matrices \n\n", (int) hA_SELLP.blocksize, (int) hA_SELLP.alignment );
while( i < argc ) {
if ( strcmp("LAPLACE2D", argv[i]) == 0 && i+1 < argc ) { // Laplace test
i++;
magma_int_t laplace_size = atoi( argv[i] );
magma_cm_5stencil( laplace_size, &hA, queue );
} else { // file-matrix test
magma_c_csr_mtx( &hA, argv[i], queue );
}
printf( "\n# matrix info: %d-by-%d with %d nonzeros\n\n",
(int) hA.num_rows,(int) hA.num_cols,(int) hA.nnz );
real_Double_t FLOPS = 2.0*hA.nnz/1e9;
magma_c_vector hx, hy, dx, dy, hrefvec, hcheck;
// init CPU vectors
magma_c_vinit( &hx, Magma_CPU, hA.num_rows, c_zero, queue );
magma_c_vinit( &hy, Magma_CPU, hA.num_rows, c_zero, queue );
// init DEV vectors
magma_c_vinit( &dx, Magma_DEV, hA.num_rows, c_one, queue );
magma_c_vinit( &dy, Magma_DEV, hA.num_rows, c_zero, queue );
#ifdef MAGMA_WITH_MKL
// calling MKL with CSR
pntre = (magma_int_t*)malloc( (hA.num_rows+1)*sizeof(magma_int_t) );
pntre[0] = 0;
for (j=0; j<hA.num_rows; j++ ) {
pntre[j] = hA.row[j+1];
}
MKL_INT num_rows = hA.num_rows;
MKL_INT num_cols = hA.num_cols;
MKL_INT nnz = hA.nnz;
MKL_INT *col;
TESTING_MALLOC_CPU( col, MKL_INT, nnz );
for( magma_int_t t=0; t < hA.nnz; ++t ) {
col[ t ] = hA.col[ t ];
}
MKL_INT *row;
TESTING_MALLOC_CPU( row, MKL_INT, num_rows );
for( magma_int_t t=0; t < hA.num_rows; ++t ) {
row[ t ] = hA.col[ t ];
}
start = magma_wtime();
for (j=0; j<10; j++ ) {
mkl_ccsrmv( "N", &num_rows, &num_cols,
MKL_ADDR(&c_one), "GFNC", MKL_ADDR(hA.val),
col, row, pntre,
MKL_ADDR(hx.val),
MKL_ADDR(&c_zero), MKL_ADDR(hy.val) );
}
end = magma_wtime();
printf( "\n > MKL : %.2e seconds %.2e GFLOP/s (CSR).\n",
(end-start)/10, FLOPS*10/(end-start) );
TESTING_FREE_CPU( row );
TESTING_FREE_CPU( col );
free(pntre);
#endif // MAGMA_WITH_MKL
// copy matrix to GPU
magma_c_mtransfer( hA, &dA, Magma_CPU, Magma_DEV, queue );
// SpMV on GPU (CSR) -- this is the reference!
start = magma_sync_wtime( queue );
for (j=0; j<10; j++)
magma_c_spmv( c_one, dA, dx, c_zero, dy, queue );
end = magma_sync_wtime( queue );
printf( " > MAGMA: %.2e seconds %.2e GFLOP/s (standard CSR).\n",
(end-start)/10, FLOPS*10/(end-start) );
magma_c_mfree(&dA, queue );
magma_c_vtransfer( dy, &hrefvec , Magma_DEV, Magma_CPU, queue );
// convert to ELL and copy to GPU
magma_c_mconvert( hA, &hA_ELL, Magma_CSR, Magma_ELL, queue );
magma_c_mtransfer( hA_ELL, &dA_ELL, Magma_CPU, Magma_DEV, queue );
magma_c_mfree(&hA_ELL, queue );
magma_c_vfree( &dy, queue );
magma_c_vinit( &dy, Magma_DEV, hA.num_rows, c_zero, queue );
// SpMV on GPU (ELL)
start = magma_sync_wtime( queue );
for (j=0; j<10; j++)
magma_c_spmv( c_one, dA_ELL, dx, c_zero, dy, queue );
end = magma_sync_wtime( queue );
printf( " > MAGMA: %.2e seconds %.2e GFLOP/s (standard ELL).\n",
(end-start)/10, FLOPS*10/(end-start) );
magma_c_mfree(&dA_ELL, queue );
magma_c_vtransfer( dy, &hcheck , Magma_DEV, Magma_CPU, queue );
res = 0.0;
for(magma_int_t k=0; k<hA.num_rows; k++ )
res=res + MAGMA_C_REAL(hcheck.val[k]) - MAGMA_C_REAL(hrefvec.val[k]);
if ( res < .000001 )
printf("# tester spmv ELL: ok\n");
else
printf("# tester spmv ELL: failed\n");
magma_c_vfree( &hcheck, queue );
// convert to SELLP and copy to GPU
magma_c_mconvert( hA, &hA_SELLP, Magma_CSR, Magma_SELLP, queue );
magma_c_mtransfer( hA_SELLP, &dA_SELLP, Magma_CPU, Magma_DEV, queue );
magma_c_mfree(&hA_SELLP, queue );
magma_c_vfree( &dy, queue );
magma_c_vinit( &dy, Magma_DEV, hA.num_rows, c_zero, queue );
// SpMV on GPU (SELLP)
start = magma_sync_wtime( queue );
for (j=0; j<10; j++)
magma_c_spmv( c_one, dA_SELLP, dx, c_zero, dy, queue );
end = magma_sync_wtime( queue );
printf( " > MAGMA: %.2e seconds %.2e GFLOP/s (SELLP).\n",
(end-start)/10, FLOPS*10/(end-start) );
magma_c_vtransfer( dy, &hcheck , Magma_DEV, Magma_CPU, queue );
res = 0.0;
for(magma_int_t k=0; k<hA.num_rows; k++ )
res=res + MAGMA_C_REAL(hcheck.val[k]) - MAGMA_C_REAL(hrefvec.val[k]);
printf("# |x-y|_F = %8.2e\n", res);
if ( res < .000001 )
printf("# tester spmv SELL-P: ok\n");
else
printf("# tester spmv SELL-P: failed\n");
magma_c_vfree( &hcheck, queue );
magma_c_mfree(&dA_SELLP, queue );
// SpMV on GPU (CUSPARSE - CSR)
// CUSPARSE context //
cusparseHandle_t cusparseHandle = 0;
cusparseStatus_t cusparseStatus;
cusparseStatus = cusparseCreate(&cusparseHandle);
cusparseSetStream( cusparseHandle, queue );
cusparseMatDescr_t descr = 0;
cusparseStatus = cusparseCreateMatDescr(&descr);
cusparseSetMatType(descr,CUSPARSE_MATRIX_TYPE_GENERAL);
cusparseSetMatIndexBase(descr,CUSPARSE_INDEX_BASE_ZERO);
magmaFloatComplex alpha = c_one;
magmaFloatComplex beta = c_zero;
magma_c_vfree( &dy, queue );
magma_c_vinit( &dy, Magma_DEV, hA.num_rows, c_zero, queue );
// copy matrix to GPU
magma_c_mtransfer( hA, &dA, Magma_CPU, Magma_DEV, queue );
start = magma_sync_wtime( queue );
for (j=0; j<10; j++)
cusparseStatus =
cusparseCcsrmv(cusparseHandle,CUSPARSE_OPERATION_NON_TRANSPOSE,
hA.num_rows, hA.num_cols, hA.nnz, &alpha, descr,
dA.dval, dA.drow, dA.dcol, dx.dval, &beta, dy.dval);
end = magma_sync_wtime( queue );
if (cusparseStatus != 0) printf("error in cuSPARSE CSR\n");
printf( " > CUSPARSE: %.2e seconds %.2e GFLOP/s (CSR).\n",
(end-start)/10, FLOPS*10/(end-start) );
cusparseMatDescr_t descrA;
cusparseStatus = cusparseCreateMatDescr(&descrA);
if (cusparseStatus != 0) printf("error\n");
cusparseHybMat_t hybA;
cusparseStatus = cusparseCreateHybMat( &hybA );
if (cusparseStatus != 0) printf("error\n");
magma_c_vtransfer( dy, &hcheck , Magma_DEV, Magma_CPU, queue );
res = 0.0;
for(magma_int_t k=0; k<hA.num_rows; k++ )
res=res + MAGMA_C_REAL(hcheck.val[k]) - MAGMA_C_REAL(hrefvec.val[k]);
printf("# |x-y|_F = %8.2e\n", res);
if ( res < .000001 )
printf("# tester spmv cuSPARSE CSR: ok\n");
else
printf("# tester spmv cuSPARSE CSR: failed\n");
magma_c_vfree( &hcheck, queue );
magma_c_vfree( &dy, queue );
magma_c_vinit( &dy, Magma_DEV, hA.num_rows, c_zero, queue );
cusparseCcsr2hyb(cusparseHandle, hA.num_rows, hA.num_cols,
descrA, dA.dval, dA.drow, dA.dcol,
hybA, 0, CUSPARSE_HYB_PARTITION_AUTO);
start = magma_sync_wtime( queue );
for (j=0; j<10; j++)
cusparseStatus =
cusparseChybmv( cusparseHandle, CUSPARSE_OPERATION_NON_TRANSPOSE,
&alpha, descrA, hybA,
dx.dval, &beta, dy.dval);
end = magma_sync_wtime( queue );
if (cusparseStatus != 0) printf("error in cuSPARSE HYB\n");
printf( " > CUSPARSE: %.2e seconds %.2e GFLOP/s (HYB).\n",
(end-start)/10, FLOPS*10/(end-start) );
magma_c_vtransfer( dy, &hcheck , Magma_DEV, Magma_CPU, queue );
res = 0.0;
for(magma_int_t k=0; k<hA.num_rows; k++ )
res=res + MAGMA_C_REAL(hcheck.val[k]) - MAGMA_C_REAL(hrefvec.val[k]);
printf("# |x-y|_F = %8.2e\n", res);
if ( res < .000001 )
printf("# tester spmv cuSPARSE HYB: ok\n");
else
printf("# tester spmv cuSPARSE HYB: failed\n");
magma_c_vfree( &hcheck, queue );
cusparseDestroyMatDescr( descrA );
cusparseDestroyHybMat( hybA );
cusparseDestroy( cusparseHandle );
magma_c_mfree(&dA, queue );
printf("\n\n");
// free CPU memory
magma_c_mfree(&hA, queue );
magma_c_vfree(&hx, queue );
magma_c_vfree(&hy, queue );
magma_c_vfree(&hrefvec, queue );
// free GPU memory
magma_c_vfree(&dx, queue );
magma_c_vfree(&dy, queue );
i++;
}
magma_queue_destroy( queue );
TESTING_FINALIZE();
return 0;
}
int main(int argc, char **argv)
{
TESTING_INIT();
const magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
const magma_int_t ione = 1;
real_Double_t gflops, magma_perf, magma_time, cpu_perf, cpu_time;
float magma_error, work[1];
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t N, lda, ldda, sizeA, sizeX, sizeY, blocks, ldwork;
magma_int_t incx = 1;
magma_int_t incy = 1;
magma_int_t nb = 64;
magmaFloatComplex alpha = MAGMA_C_MAKE( 1.5, -2.3 );
magmaFloatComplex beta = MAGMA_C_MAKE( -0.6, 0.8 );
magmaFloatComplex *A, *X, *Y, *Ymagma;
magmaFloatComplex_ptr dA, dX, dY, dwork;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
printf("uplo = %s\n", lapack_uplo_const(opts.uplo) );
printf(" N MAGMA Gflop/s (ms) CPU Gflop/s (ms) MAGMA error\n");
printf("=========================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
ldda = ((N + 31)/32)*32;
sizeA = N*lda;
sizeX = N*incx;
sizeY = N*incy;
gflops = FLOPS_CSYMV( N ) / 1e9;
TESTING_MALLOC_CPU( A, magmaFloatComplex, sizeA );
TESTING_MALLOC_CPU( X, magmaFloatComplex, sizeX );
TESTING_MALLOC_CPU( Y, magmaFloatComplex, sizeY );
TESTING_MALLOC_CPU( Ymagma, magmaFloatComplex, sizeY );
TESTING_MALLOC_DEV( dA, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( dX, magmaFloatComplex, sizeX );
TESTING_MALLOC_DEV( dY, magmaFloatComplex, sizeY );
blocks = (N + nb - 1) / nb;
ldwork = ldda*blocks;
TESTING_MALLOC_DEV( dwork, magmaFloatComplex, ldwork );
magmablas_claset( MagmaFull, ldwork, 1, MAGMA_C_NAN, MAGMA_C_NAN, dwork, ldwork );
magmablas_claset( MagmaFull, ldda, N, MAGMA_C_NAN, MAGMA_C_NAN, dA, ldda );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &sizeA, A );
magma_cmake_hermitian( N, A, lda );
// should not use data from the opposite triangle -- fill with NAN to check
magma_int_t N1 = N-1;
if ( opts.uplo == MagmaUpper ) {
lapackf77_claset( "Lower", &N1, &N1, &MAGMA_C_NAN, &MAGMA_C_NAN, &A[1], &lda );
}
else {
lapackf77_claset( "Upper", &N1, &N1, &MAGMA_C_NAN, &MAGMA_C_NAN, &A[lda], &lda );
}
lapackf77_clarnv( &ione, ISEED, &sizeX, X );
lapackf77_clarnv( &ione, ISEED, &sizeY, Y );
/* Note: CUBLAS does not implement csymv */
/* =====================================================================
Performs operation using MAGMABLAS
=================================================================== */
magma_csetmatrix( N, N, A, lda, dA, ldda );
magma_csetvector( N, X, incx, dX, incx );
magma_csetvector( N, Y, incy, dY, incy );
magma_time = magma_sync_wtime( 0 );
if ( opts.version == 1 ) {
magmablas_csymv_work( opts.uplo, N, alpha, dA, ldda, dX, incx, beta, dY, incy, dwork, ldwork, opts.queue );
}
else {
// non-work interface (has added overhead)
magmablas_csymv( opts.uplo, N, alpha, dA, ldda, dX, incx, beta, dY, incy );
}
magma_time = magma_sync_wtime( 0 ) - magma_time;
magma_perf = gflops / magma_time;
magma_cgetvector( N, dY, incy, Ymagma, incy );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
cpu_time = magma_wtime();
lapackf77_csymv( lapack_uplo_const(opts.uplo), &N, &alpha, A, &lda, X, &incx, &beta, Y, &incy );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/* =====================================================================
Check the result
=================================================================== */
blasf77_caxpy( &N, &c_neg_one, Y, &incy, Ymagma, &incy );
magma_error = lapackf77_clange( "M", &N, &ione, Ymagma, &N, work ) / N;
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) N,
magma_perf, 1000.*magma_time,
cpu_perf, 1000.*cpu_time,
magma_error, (magma_error < tol ? "ok" : "failed"));
status += ! (magma_error < tol);
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( X );
TESTING_FREE_CPU( Y );
TESTING_FREE_CPU( Ymagma );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dX );
TESTING_FREE_DEV( dY );
TESTING_FREE_DEV( dwork );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
extern "C" magma_int_t
magma_cgeqrf4(magma_int_t num_gpus, magma_int_t m, magma_int_t n,
magmaFloatComplex *a, magma_int_t lda, magmaFloatComplex *tau,
magmaFloatComplex *work, magma_int_t lwork,
magma_int_t *info )
{
/* -- MAGMA (version 1.4.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
December 2013
Purpose
=======
CGEQRF4 computes a QR factorization of a COMPLEX M-by-N matrix A:
A = Q * R using multiple GPUs. This version does not require work space on the GPU
passed as input. GPU memory is allocated in the routine.
Arguments
=========
NUM_GPUS
(input) INTEGER
The number of GPUs to be used for the factorization.
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. N >= 0.
A (input/output) COMPLEX array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, the elements on and above the diagonal of the array
contain the min(M,N)-by-N upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of min(m,n) elementary reflectors (see Further
Details).
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,M).
TAU (output) COMPLEX array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
WORK (workspace/output) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
Higher performance is achieved if WORK is in pinned memory, e.g.
allocated using magma_malloc_pinned.
LWORK (input) INTEGER
The dimension of the array WORK. LWORK >= N*NB,
where NB can be obtained through magma_get_cgeqrf_nb(M).
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
Further Details
===============
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
===================================================================== */
magmaFloatComplex *da[MagmaMaxGPUs];
magmaFloatComplex c_one = MAGMA_C_ONE;
int i, k, ldda;
*info = 0;
int nb = magma_get_cgeqrf_nb(min(m, n));
int lwkopt = n * nb;
work[0] = MAGMA_C_MAKE( (float)lwkopt, 0 );
int lquery = (lwork == -1);
if (num_gpus <0 || num_gpus > 4) {
*info = -1;
} else if (m < 0) {
*info = -2;
} else if (n < 0) {
*info = -3;
} else if (lda < max(1,m)) {
*info = -5;
} else if (lwork < max(1,n) && ! lquery) {
*info = -8;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
k = min(m,n);
if (k == 0) {
work[0] = c_one;
return *info;
}
ldda = ((m+31)/32)*32;
magma_int_t n_local[MagmaMaxGPUs];
for(i=0; i<num_gpus; i++){
n_local[i] = ((n/nb)/num_gpus)*nb;
if (i < (n/nb)%num_gpus)
n_local[i] += nb;
else if (i == (n/nb)%num_gpus)
n_local[i] += n%nb;
magma_setdevice(i);
// TODO on failure, free previously allocated memory
if (MAGMA_SUCCESS != magma_cmalloc( &da[i], ldda*n_local[i] )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
}
if (m > nb && n > nb) {
/* Copy the matrix to the GPUs in 1D block cyclic distribution */
magma_csetmatrix_1D_col_bcyclic(m, n, a, lda, da, ldda, num_gpus, nb);
/* Factor using the GPU interface */
magma_cgeqrf2_mgpu( num_gpus, m, n, da, ldda, tau, info);
/* Copy the matrix back from the GPUs to the CPU */
magma_cgetmatrix_1D_col_bcyclic(m, n, da, ldda, a, lda, num_gpus, nb);
}
else {
lapackf77_cgeqrf(&m, &n, a, &lda, tau, work, &lwork, info);
}
/* Free the allocated GPU memory */
for(i=0; i<num_gpus; i++){
magma_setdevice(i);
magma_free( da[i] );
}
return *info;
} /* magma_cgeqrf4 */
int main( int argc, char** argv)
{
#define hA(i,j) (hA + (i) + (j)*lda)
TESTING_CUDA_INIT();
cuFloatComplex c_zero = MAGMA_C_ZERO;
cuFloatComplex c_one = MAGMA_C_ONE;
cuFloatComplex *hA, *hR, *dA;
//real_Double_t gpu_time, gpu_perf;
//int ione = 1;
//int ISEED[4] = {0, 0, 0, 1};
int nsize[] = { 32, 64, 96, 256, 100, 200, 512 };
int ntest = sizeof(nsize) / sizeof(int);
int n = nsize[ntest-1];
int lda = ((n + 31)/32)*32;
int ntile, nb;
TESTING_MALLOC ( hA, cuFloatComplex, lda*n );
TESTING_MALLOC ( hR, cuFloatComplex, lda*n );
TESTING_DEVALLOC ( dA, cuFloatComplex, lda*n );
for( int t = 0; t < ntest; ++t ) {
n = nsize[t];
lda = ((n + 31)/32)*32;
// initialize matrices; entries are (i.j) for A
float nf = 100.;
for( int j = 0; j < n; ++j ) {
// upper
for( int i = 0; i < j; ++i ) {
*hA(i,j) = MAGMA_C_MAKE( (i + j/nf)/nf, 0. );
}
// lower
for( int i = j; i < n; ++i ) {
*hA(i,j) = MAGMA_C_MAKE( i + j/nf, 0. );
}
}
printf( "A%d = ", n );
magma_cprint( n, n, hA, lda );
magma_csetmatrix( n, n, hA, lda, dA, lda );
magmablas_csymmetrize( MagmaLower, n, dA, lda );
magma_cgetmatrix( n, n, dA, lda, hR, lda );
printf( "L%d = ", n );
magma_cprint( n, n, hR, lda );
magma_csetmatrix( n, n, hA, lda, dA, lda );
magmablas_csymmetrize( MagmaUpper, n, dA, lda );
magma_cgetmatrix( n, n, dA, lda, hR, lda );
printf( "U%d = ", n );
magma_cprint( n, n, hR, lda );
// -----
//lapackf77_claset( "u", &n, &n, &c_zero, &c_one, hA, &lda );
nb = 64;
ntile = n / nb;
magma_csetmatrix( n, n, hA, lda, dA, lda );
magmablas_csymmetrize_tiles( MagmaLower, nb, dA, lda, ntile, nb, nb );
magma_cgetmatrix( n, n, dA, lda, hR, lda );
printf( "L%d_%d = ", n, nb );
magma_cprint( n, n, hR, lda );
nb = 32;
ntile = n / nb;
magma_csetmatrix( n, n, hA, lda, dA, lda );
magmablas_csymmetrize_tiles( MagmaLower, nb, dA, lda, ntile, nb, nb );
magma_cgetmatrix( n, n, dA, lda, hR, lda );
printf( "L%d_%d = ", n, nb );
magma_cprint( n, n, hR, lda );
ntile = (n - nb < 0 ? 0 : (n - nb) / (2*nb) + 1);
magma_csetmatrix( n, n, hA, lda, dA, lda );
magmablas_csymmetrize_tiles( MagmaLower, nb, dA, lda, ntile, 2*nb, nb );
magma_cgetmatrix( n, n, dA, lda, hR, lda );
printf( "L%d_%d_2m = ", n, nb );
magma_cprint( n, n, hR, lda );
nb = 25;
ntile = n / nb;
magma_csetmatrix( n, n, hA, lda, dA, lda );
magmablas_csymmetrize_tiles( MagmaLower, nb, dA, lda, ntile, nb, nb );
magma_cgetmatrix( n, n, dA, lda, hR, lda );
printf( "L%d_%d = ", n, nb );
magma_cprint( n, n, hR, lda );
nb = 25;
ntile = (n - nb < 0 ? 0 : (n - nb) / (3*nb) + 1);
magma_csetmatrix( n, n, hA, lda, dA, lda );
magmablas_csymmetrize_tiles( MagmaLower, nb, dA, lda, ntile, nb, 3*nb );
magma_cgetmatrix( n, n, dA, lda, hR, lda );
printf( "L%d_%d_3n = ", n, nb );
magma_cprint( n, n, hR, lda );
nb = 100;
ntile = n / nb;
magma_csetmatrix( n, n, hA, lda, dA, lda );
magmablas_csymmetrize_tiles( MagmaLower, nb, dA, lda, ntile, nb, nb );
magmablas_csymmetrize( MagmaLower, n%nb, &dA[ ntile*nb*(1+lda) ], lda ); // last partial block
magma_cgetmatrix( n, n, dA, lda, hR, lda );
printf( "L%d_%d = ", n, nb );
magma_cprint( n, n, hR, lda );
// -----
nb = 64;
ntile = n / nb;
magma_csetmatrix( n, n, hA, lda, dA, lda );
magmablas_csymmetrize_tiles( MagmaUpper, nb, dA, lda, ntile, nb, nb );
magma_cgetmatrix( n, n, dA, lda, hR, lda );
printf( "U%d_%d = ", n, nb );
magma_cprint( n, n, hR, lda );
}
TESTING_FREE( hA );
TESTING_FREE( hR );
TESTING_DEVFREE( dA );
/* Shutdown */
TESTING_CUDA_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgeadd_batched
Code is very similar to testing_clacpy_batched.cpp
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float error, work[1];
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *h_A, *h_B;
magmaFloatComplex *d_A, *d_B;
magmaFloatComplex **hAarray, **hBarray, **dAarray, **dBarray;
magmaFloatComplex alpha = MAGMA_C_MAKE( 3.1415, 2.718 );
magma_int_t M, N, mb, nb, size, lda, ldda, mstride, nstride, ntile;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
mb = (opts.nb == 0 ? 32 : opts.nb);
nb = (opts.nb == 0 ? 64 : opts.nb);
mstride = 2*mb;
nstride = 3*nb;
printf("mb=%d, nb=%d, mstride=%d, nstride=%d\n", (int) mb, (int) nb, (int) mstride, (int) nstride );
printf(" M N ntile CPU GFlop/s (ms) GPU GFlop/s (ms) error \n");
printf("====================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
lda = M;
ldda = ((M+31)/32)*32;
size = lda*N;
if ( N < nb || M < nb ) {
ntile = 0;
} else {
ntile = min( (M - nb)/mstride + 1,
(N - nb)/nstride + 1 );
}
gflops = 2.*mb*nb*ntile / 1e9;
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, lda *N );
TESTING_MALLOC_CPU( h_B, magmaFloatComplex, lda *N );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( d_B, magmaFloatComplex, ldda*N );
TESTING_MALLOC_CPU( hAarray, magmaFloatComplex*, ntile );
TESTING_MALLOC_CPU( hBarray, magmaFloatComplex*, ntile );
TESTING_MALLOC_DEV( dAarray, magmaFloatComplex*, ntile );
TESTING_MALLOC_DEV( dBarray, magmaFloatComplex*, ntile );
lapackf77_clarnv( &ione, ISEED, &size, h_A );
lapackf77_clarnv( &ione, ISEED, &size, h_B );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_csetmatrix( M, N, h_A, lda, d_A, ldda );
magma_csetmatrix( M, N, h_B, lda, d_B, ldda );
// setup pointers
for( int tile = 0; tile < ntile; ++tile ) {
int offset = tile*mstride + tile*nstride*ldda;
hAarray[tile] = &d_A[offset];
hBarray[tile] = &d_B[offset];
}
magma_setvector( ntile, sizeof(magmaFloatComplex*), hAarray, 1, dAarray, 1 );
magma_setvector( ntile, sizeof(magmaFloatComplex*), hBarray, 1, dBarray, 1 );
gpu_time = magma_sync_wtime( 0 );
magmablas_cgeadd_batched( mb, nb, alpha, dAarray, ldda, dBarray, ldda, ntile );
gpu_time = magma_sync_wtime( 0 ) - gpu_time;
gpu_perf = gflops / gpu_time;
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
for( int tile = 0; tile < ntile; ++tile ) {
int offset = tile*mstride + tile*nstride*lda;
for( int j = 0; j < nb; ++j ) {
blasf77_caxpy( &mb, &alpha,
&h_A[offset + j*lda], &ione,
&h_B[offset + j*lda], &ione );
}
}
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/* =====================================================================
Check the result
=================================================================== */
magma_cgetmatrix( M, N, d_B, ldda, h_A, lda );
error = lapackf77_clange( "F", &M, &N, h_B, &lda, work );
blasf77_caxpy(&size, &c_neg_one, h_A, &ione, h_B, &ione);
error = lapackf77_clange("f", &M, &N, h_B, &lda, work) / error;
printf("%5d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) M, (int) N, (int) ntile,
cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000.,
error, (error < tol ? "ok" : "failed"));
status += ! (error < tol);
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
TESTING_FREE_CPU( hAarray );
TESTING_FREE_CPU( hBarray );
TESTING_FREE_DEV( dAarray );
TESTING_FREE_DEV( dBarray );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
extern "C" magma_int_t
magma_cpgmres(
magma_c_sparse_matrix A, magma_c_vector b, magma_c_vector *x,
magma_c_solver_par *solver_par,
magma_c_preconditioner *precond_par,
magma_queue_t queue )
{
magma_int_t stat = 0;
// set queue for old dense routines
magma_queue_t orig_queue;
magmablasGetKernelStream( &orig_queue );
magma_int_t stat_cpu = 0, stat_dev = 0;
// prepare solver feedback
solver_par->solver = Magma_PGMRES;
solver_par->numiter = 0;
solver_par->info = MAGMA_SUCCESS;
// local variables
magmaFloatComplex c_zero = MAGMA_C_ZERO, c_one = MAGMA_C_ONE,
c_mone = MAGMA_C_NEG_ONE;
magma_int_t dofs = A.num_rows;
magma_int_t i, j, k, m = 0;
magma_int_t restart = min( dofs-1, solver_par->restart );
magma_int_t ldh = restart+1;
float nom, rNorm, RNorm, nom0, betanom, r0 = 0.;
// CPU workspace
//magma_setdevice(0);
magmaFloatComplex *H, *HH, *y, *h1;
stat_cpu += magma_cmalloc_pinned( &H, (ldh+1)*ldh );
stat_cpu += magma_cmalloc_pinned( &y, ldh );
stat_cpu += magma_cmalloc_pinned( &HH, ldh*ldh );
stat_cpu += magma_cmalloc_pinned( &h1, ldh );
if( stat_cpu != 0){
magma_free_pinned( H );
magma_free_pinned( y );
magma_free_pinned( HH );
magma_free_pinned( h1 );
magmablasSetKernelStream( orig_queue );
return MAGMA_ERR_HOST_ALLOC;
}
// GPU workspace
magma_c_vector r, q, q_t, z, z_t, t;
magma_c_vinit( &t, Magma_DEV, dofs, c_zero, queue );
magma_c_vinit( &r, Magma_DEV, dofs, c_zero, queue );
magma_c_vinit( &q, Magma_DEV, dofs*(ldh+1), c_zero, queue );
magma_c_vinit( &z, Magma_DEV, dofs*(ldh+1), c_zero, queue );
magma_c_vinit( &z_t, Magma_DEV, dofs, c_zero, queue );
q_t.memory_location = Magma_DEV;
q_t.dval = NULL;
q_t.num_rows = q_t.nnz = dofs; q_t.num_cols = 1;
magmaFloatComplex *dy = NULL, *dH = NULL;
stat_dev += magma_cmalloc( &dy, ldh );
stat_dev += magma_cmalloc( &dH, (ldh+1)*ldh );
if( stat_dev != 0){
magma_free_pinned( H );
magma_free_pinned( y );
magma_free_pinned( HH );
magma_free_pinned( h1 );
magma_free( dH );
magma_free( dy );
magmablasSetKernelStream( orig_queue );
return MAGMA_ERR_DEVICE_ALLOC;
}
magma_cscal( dofs, c_zero, x->dval, 1 ); // x = 0
magma_ccopy( dofs, b.dval, 1, r.dval, 1 ); // r = b
nom0 = betanom = magma_scnrm2( dofs, r.dval, 1 ); // nom0= || r||
nom = nom0 * nom0;
solver_par->init_res = nom0;
H(1,0) = MAGMA_C_MAKE( nom0, 0. );
magma_csetvector(1, &H(1,0), 1, &dH(1,0), 1);
if ( (r0 = nom0 * solver_par->epsilon ) < ATOLERANCE ){
r0 = solver_par->epsilon;
}
if ( nom < r0 ) {
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
}
//Chronometry
real_Double_t tempo1, tempo2;
tempo1 = magma_sync_wtime( queue );
if ( solver_par->verbose > 0 ) {
solver_par->res_vec[0] = nom0;
solver_par->timing[0] = 0.0;
}
// start iteration
for( solver_par->numiter= 1; solver_par->numiter<solver_par->maxiter;
solver_par->numiter++ ) {
for(k=1; k<=restart; k++) {
magma_ccopy(dofs, r.dval, 1, q(k-1), 1); // q[0] = 1.0/||r||
magma_cscal(dofs, 1./H(k,k-1), q(k-1), 1); // (to be fused)
q_t.dval = q(k-1);
// preconditioner
// z[k] = M^(-1) q(k)
magma_c_applyprecond_left( A, q_t, &t, precond_par, queue );
magma_c_applyprecond_right( A, t, &z_t, precond_par, queue );
magma_ccopy(dofs, z_t.dval, 1, z(k-1), 1);
// r = A q[k]
magma_c_spmv( c_one, A, z_t, c_zero, r, queue );
// if (solver_par->ortho == Magma_MGS ) {
// modified Gram-Schmidt
for (i=1; i<=k; i++) {
H(i,k) =magma_cdotc(dofs, q(i-1), 1, r.dval, 1);
// H(i,k) = q[i] . r
magma_caxpy(dofs,-H(i,k), q(i-1), 1, r.dval, 1);
// r = r - H(i,k) q[i]
}
H(k+1,k) = MAGMA_C_MAKE( magma_scnrm2(dofs, r.dval, 1), 0. ); // H(k+1,k) = ||r||
/*} else if (solver_par->ortho == Magma_FUSED_CGS ) {
// fusing cgemv with scnrm2 in classical Gram-Schmidt
magmablasSetKernelStream(stream[0]);
magma_ccopy(dofs, r.dval, 1, q(k), 1);
// dH(1:k+1,k) = q[0:k] . r
magmablas_cgemv(MagmaTrans, dofs, k+1, c_one, q(0),
dofs, r.dval, 1, c_zero, &dH(1,k), 1);
// r = r - q[0:k-1] dH(1:k,k)
magmablas_cgemv(MagmaNoTrans, dofs, k, c_mone, q(0),
dofs, &dH(1,k), 1, c_one, r.dval, 1);
// 1) dH(k+1,k) = sqrt( dH(k+1,k) - dH(1:k,k) )
magma_ccopyscale( dofs, k, r.dval, q(k), &dH(1,k) );
// 2) q[k] = q[k] / dH(k+1,k)
magma_event_record( event[0], stream[0] );
magma_queue_wait_event( stream[1], event[0] );
magma_cgetvector_async(k+1, &dH(1,k), 1, &H(1,k), 1, stream[1]);
// asynch copy dH(1:(k+1),k) to H(1:(k+1),k)
} else {
// classical Gram-Schmidt (default)
// > explicitly calling magmabls
magmablasSetKernelStream(stream[0]);
magmablas_cgemv(MagmaTrans, dofs, k, c_one, q(0),
dofs, r.dval, 1, c_zero, &dH(1,k), 1, queue );
// dH(1:k,k) = q[0:k-1] . r
#ifndef SCNRM2SCALE
// start copying dH(1:k,k) to H(1:k,k)
magma_event_record( event[0], stream[0] );
magma_queue_wait_event( stream[1], event[0] );
magma_cgetvector_async(k, &dH(1,k), 1, &H(1,k),
1, stream[1]);
#endif
// r = r - q[0:k-1] dH(1:k,k)
magmablas_cgemv(MagmaNoTrans, dofs, k, c_mone, q(0),
dofs, &dH(1,k), 1, c_one, r.dval, 1);
#ifdef SCNRM2SCALE
magma_ccopy(dofs, r.dval, 1, q(k), 1);
// q[k] = r / H(k,k-1)
magma_scnrm2scale(dofs, q(k), dofs, &dH(k+1,k) );
// dH(k+1,k) = sqrt(r . r) and r = r / dH(k+1,k)
magma_event_record( event[0], stream[0] );
// start sending dH(1:k,k) to H(1:k,k)
magma_queue_wait_event( stream[1], event[0] );
// can we keep H(k+1,k) on GPU and combine?
magma_cgetvector_async(k+1, &dH(1,k), 1, &H(1,k), 1, stream[1]);
#else
H(k+1,k) = MAGMA_C_MAKE( magma_scnrm2(dofs, r.dval, 1), 0. );
// H(k+1,k) = sqrt(r . r)
if ( k<solver_par->restart ) {
magmablasSetKernelStream(stream[0]);
magma_ccopy(dofs, r.dval, 1, q(k), 1);
// q[k] = 1.0/H[k][k-1] r
magma_cscal(dofs, 1./H(k+1,k), q(k), 1);
// (to be fused)
}
#endif
}*/
/* Minimization of || b-Ax || in H_k */
for (i=1; i<=k; i++) {
HH(k,i) = magma_cblas_cdotc( i+1, &H(1,k), 1, &H(1,i), 1 );
}
h1[k] = H(1,k)*H(1,0);
if (k != 1) {
for (i=1; i<k; i++) {
HH(k,i) = HH(k,i)/HH(i,i);//
for (m=i+1; m<=k; m++) {
HH(k,m) -= HH(k,i) * HH(m,i) * HH(i,i);
}
h1[k] -= h1[i] * HH(k,i);
}
}
y[k] = h1[k]/HH(k,k);
if (k != 1)
for (i=k-1; i>=1; i--) {
y[i] = h1[i]/HH(i,i);
for (j=i+1; j<=k; j++)
y[i] -= y[j] * HH(j,i);
}
m = k;
rNorm = fabs(MAGMA_C_REAL(H(k+1,k)));
}/* Minimization done */
// compute solution approximation
magma_csetmatrix(m, 1, y+1, m, dy, m );
magma_cgemv(MagmaNoTrans, dofs, m, c_one, z(0), dofs, dy, 1,
c_one, x->dval, 1);
// compute residual
magma_c_spmv( c_mone, A, *x, c_zero, r, queue ); // r = - A * x
magma_caxpy(dofs, c_one, b.dval, 1, r.dval, 1); // r = r + b
H(1,0) = MAGMA_C_MAKE( magma_scnrm2(dofs, r.dval, 1), 0. );
// RNorm = H[1][0] = || r ||
RNorm = MAGMA_C_REAL( H(1,0) );
betanom = fabs(RNorm);
if ( solver_par->verbose > 0 ) {
tempo2 = magma_sync_wtime( queue );
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
if ( betanom < r0 ) {
break;
}
}
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t) tempo2-tempo1;
float residual;
magma_cresidual( A, b, *x, &residual, queue );
solver_par->iter_res = betanom;
solver_par->final_res = residual;
if ( solver_par->numiter < solver_par->maxiter) {
solver_par->info = MAGMA_SUCCESS;
} else if ( solver_par->init_res > solver_par->final_res ) {
if ( solver_par->verbose > 0 ) {
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
solver_par->info = MAGMA_SLOW_CONVERGENCE;
}
else {
if ( solver_par->verbose > 0 ) {
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
solver_par->info = MAGMA_DIVERGENCE;
}
// free pinned memory
magma_free_pinned( H );
magma_free_pinned( y );
magma_free_pinned( HH );
magma_free_pinned( h1 );
// free GPU memory
magma_free(dy);
if (dH != NULL ) magma_free(dH);
magma_c_vfree(&t, queue );
magma_c_vfree(&r, queue );
magma_c_vfree(&q, queue );
magma_c_vfree(&z, queue );
magma_c_vfree(&z_t, queue );
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
} /* magma_cpgmres */
extern "C" magma_int_t
magma_clsqr(
magma_c_matrix A, magma_c_matrix b, magma_c_matrix *x,
magma_c_solver_par *solver_par,
magma_c_preconditioner *precond_par,
magma_queue_t queue )
{
magma_int_t info = MAGMA_NOTCONVERGED;
// prepare solver feedback
solver_par->solver = Magma_LSQR;
solver_par->numiter = 0;
solver_par->spmv_count = 0;
magma_int_t m = A.num_rows * b.num_cols;
magma_int_t n = A.num_cols * b.num_cols;
// local variables
magmaFloatComplex c_zero = MAGMA_C_ZERO, c_one = MAGMA_C_ONE;
// solver variables
float s, nom0, r0, res=0, nomb, phibar, beta, alpha, c, rho, rhot, phi, thet, normr, normar, norma, sumnormd2, normd;
// need to transpose the matrix
magma_c_matrix AT={Magma_CSR}, Ah1={Magma_CSR}, Ah2={Magma_CSR};
// GPU workspace
magma_c_matrix r={Magma_CSR},
v={Magma_CSR}, z={Magma_CSR}, zt={Magma_CSR},
d={Magma_CSR}, vt={Magma_CSR}, q={Magma_CSR},
w={Magma_CSR}, u={Magma_CSR};
CHECK( magma_cvinit( &r, Magma_DEV, A.num_cols, b.num_cols, c_zero, queue ));
CHECK( magma_cvinit( &v, Magma_DEV, A.num_cols, b.num_cols, c_zero, queue ));
CHECK( magma_cvinit( &z, Magma_DEV, A.num_cols, b.num_cols, c_zero, queue ));
CHECK( magma_cvinit( &d, Magma_DEV, A.num_cols, b.num_cols, c_zero, queue ));
CHECK( magma_cvinit( &vt,Magma_DEV, A.num_cols, b.num_cols, c_zero, queue ));
CHECK( magma_cvinit( &q, Magma_DEV, A.num_cols, b.num_cols, c_zero, queue ));
CHECK( magma_cvinit( &w, Magma_DEV, A.num_cols, b.num_cols, c_zero, queue ));
CHECK( magma_cvinit( &u, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_cvinit( &zt,Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
// transpose the matrix
magma_cmtransfer( A, &Ah1, Magma_DEV, Magma_CPU, queue );
magma_cmconvert( Ah1, &Ah2, A.storage_type, Magma_CSR, queue );
magma_cmfree(&Ah1, queue );
magma_cmtransposeconjugate( Ah2, &Ah1, queue );
magma_cmfree(&Ah2, queue );
Ah2.blocksize = A.blocksize;
Ah2.alignment = A.alignment;
magma_cmconvert( Ah1, &Ah2, Magma_CSR, A.storage_type, queue );
magma_cmfree(&Ah1, queue );
magma_cmtransfer( Ah2, &AT, Magma_CPU, Magma_DEV, queue );
magma_cmfree(&Ah2, queue );
// solver setup
CHECK( magma_cresidualvec( A, b, *x, &r, &nom0, queue));
solver_par->init_res = nom0;
nomb = magma_scnrm2( m, b.dval, 1, queue );
if ( nomb == 0.0 ){
nomb=1.0;
}
if ( (r0 = nomb * solver_par->rtol) < ATOLERANCE ){
r0 = ATOLERANCE;
}
solver_par->final_res = solver_par->init_res;
solver_par->iter_res = solver_par->init_res;
if ( solver_par->verbose > 0 ) {
solver_par->res_vec[0] = (real_Double_t)nom0;
solver_par->timing[0] = 0.0;
}
if ( nom0 < r0 ) {
info = MAGMA_SUCCESS;
goto cleanup;
}
magma_ccopy( m, b.dval, 1, u.dval, 1, queue );
beta = magma_scnrm2( m, u.dval, 1, queue );
magma_cscal( m, MAGMA_C_MAKE(1./beta, 0.0 ), u.dval, 1, queue );
normr = beta;
c = 1.0;
s = 0.0;
phibar = beta;
CHECK( magma_c_spmv( c_one, AT, u, c_zero, v, queue ));
if( precond_par->solver == Magma_NONE ){
;
} else {
CHECK( magma_c_applyprecond_right( MagmaTrans, A, v, &zt, precond_par, queue ));
CHECK( magma_c_applyprecond_left( MagmaTrans, A, zt, &v, precond_par, queue ));
}
alpha = magma_scnrm2( n, v.dval, 1, queue );
magma_cscal( n, MAGMA_C_MAKE(1./alpha, 0.0 ), v.dval, 1, queue );
normar = alpha * beta;
norma = 0;
sumnormd2 = 0;
//Chronometry
real_Double_t tempo1, tempo2;
tempo1 = magma_sync_wtime( queue );
solver_par->numiter = 0;
// start iteration
do
{
solver_par->numiter++;
if( precond_par->solver == Magma_NONE || A.num_rows != A.num_cols ) {
magma_ccopy( n, v.dval, 1 , z.dval, 1, queue );
} else {
CHECK( magma_c_applyprecond_left( MagmaNoTrans, A, v, &zt, precond_par, queue ));
CHECK( magma_c_applyprecond_right( MagmaNoTrans, A, zt, &z, precond_par, queue ));
}
//CHECK( magma_c_spmv( c_one, A, z, MAGMA_C_MAKE(-alpha,0.0), u, queue ));
CHECK( magma_c_spmv( c_one, A, z, c_zero, zt, queue ));
magma_cscal( m, MAGMA_C_MAKE(-alpha, 0.0 ), u.dval, 1, queue );
magma_caxpy( m, c_one, zt.dval, 1, u.dval, 1, queue );
solver_par->spmv_count++;
beta = magma_scnrm2( m, u.dval, 1, queue );
magma_cscal( m, MAGMA_C_MAKE(1./beta, 0.0 ), u.dval, 1, queue );
// norma = norm([norma alpha beta]);
norma = sqrt(norma*norma + alpha*alpha + beta*beta );
//lsvec( solver_par->numiter-1 ) = normar / norma;
thet = -s * alpha;
rhot = c * alpha;
rho = sqrt( rhot * rhot + beta * beta );
c = rhot / rho;
s = - beta / rho;
phi = c * phibar;
phibar = s * phibar;
// d = (z - thet * d) / rho;
magma_cscal( n, MAGMA_C_MAKE(-thet, 0.0 ), d.dval, 1, queue );
magma_caxpy( n, c_one, z.dval, 1, d.dval, 1, queue );
magma_cscal( n, MAGMA_C_MAKE(1./rho, 0.0 ), d.dval, 1, queue );
normd = magma_scnrm2( n, d.dval, 1, queue );
sumnormd2 = sumnormd2 + normd*normd;
// convergence check
res = normr;
if ( solver_par->verbose > 0 ) {
tempo2 = magma_sync_wtime( queue );
if ( (solver_par->numiter)%solver_par->verbose == c_zero ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) res;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
// check for convergence in A*x=b
if ( res/nomb <= solver_par->rtol || res <= solver_par->atol ){
info = MAGMA_SUCCESS;
break;
}
// check for convergence in min{|b-A*x|}
if ( A.num_rows != A.num_cols &&
( normar/(norma*normr) <= solver_par->rtol || normar <= solver_par->atol ) ){
printf("%% warning: quit from minimization convergence check.\n");
info = MAGMA_SUCCESS;
break;
}
magma_caxpy( n, MAGMA_C_MAKE( phi, 0.0 ), d.dval, 1, x->dval, 1, queue );
normr = fabs(s) * normr;
CHECK( magma_c_spmv( c_one, AT, u, c_zero, vt, queue ));
solver_par->spmv_count++;
if( precond_par->solver == Magma_NONE ){
;
} else {
CHECK( magma_c_applyprecond_right( MagmaTrans, A, vt, &zt, precond_par, queue ));
CHECK( magma_c_applyprecond_left( MagmaTrans, A, zt, &vt, precond_par, queue ));
}
magma_cscal( n, MAGMA_C_MAKE(-beta, 0.0 ), v.dval, 1, queue );
magma_caxpy( n, c_one, vt.dval, 1, v.dval, 1, queue );
alpha = magma_scnrm2( n, v.dval, 1, queue );
magma_cscal( n, MAGMA_C_MAKE(1./alpha, 0.0 ), v.dval, 1, queue );
normar = alpha * fabs(s*phi);
}
while ( solver_par->numiter+1 <= solver_par->maxiter );
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t) tempo2-tempo1;
float residual;
CHECK( magma_cresidualvec( A, b, *x, &r, &residual, queue));
solver_par->iter_res = res;
solver_par->final_res = residual;
if ( solver_par->numiter < solver_par->maxiter && info == MAGMA_SUCCESS ) {
info = MAGMA_SUCCESS;
} else if ( solver_par->init_res > solver_par->final_res ) {
if ( solver_par->verbose > 0 ) {
if ( (solver_par->numiter)%solver_par->verbose == c_zero ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) res;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
info = MAGMA_SLOW_CONVERGENCE;
if( solver_par->iter_res < solver_par->rtol*solver_par->init_res ||
solver_par->iter_res < solver_par->atol ) {
info = MAGMA_SUCCESS;
}
}
else {
if ( solver_par->verbose > 0 ) {
if ( (solver_par->numiter)%solver_par->verbose == c_zero ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) res;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
info = MAGMA_DIVERGENCE;
}
cleanup:
magma_cmfree(&r, queue );
magma_cmfree(&v, queue );
magma_cmfree(&z, queue );
magma_cmfree(&zt, queue );
magma_cmfree(&d, queue );
magma_cmfree(&vt, queue );
magma_cmfree(&q, queue );
magma_cmfree(&u, queue );
magma_cmfree(&w, queue );
magma_cmfree(&AT, queue );
magma_cmfree(&Ah1, queue );
magma_cmfree(&Ah2, queue );
solver_par->info = info;
return info;
} /* magma_cqmr */
extern "C" magma_int_t
magma_cheevx(char jobz, char range, char uplo, magma_int_t n,
magmaFloatComplex *a, magma_int_t lda, float vl, float vu,
magma_int_t il, magma_int_t iu, float abstol, magma_int_t *m,
float *w, magmaFloatComplex *z, magma_int_t ldz, magmaFloatComplex *work, magma_int_t lwork,
float *rwork, magma_int_t *iwork, magma_int_t *ifail, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CHEEVX computes selected eigenvalues and, optionally, eigenvectors
of a complex Hermitian matrix A. Eigenvalues and eigenvectors can
be selected by specifying either a range of values or a range of
indices for the desired eigenvalues.
Arguments
=========
JOBZ (input) CHARACTER*1
= 'N': Compute eigenvalues only;
= 'V': Compute eigenvalues and eigenvectors.
RANGE (input) CHARACTER*1
= 'A': all eigenvalues will be found.
= 'V': all eigenvalues in the half-open interval (VL,VU]
will be found.
= 'I': the IL-th through IU-th eigenvalues will be found.
UPLO (input) CHARACTER*1
= 'U': Upper triangle of A is stored;
= 'L': Lower triangle of A is stored.
N (input) INTEGER
The order of the matrix A. N >= 0.
A (input/output) COMPLEX array, dimension (LDA, N)
On entry, the Hermitian matrix A. If UPLO = 'U', the
leading N-by-N upper triangular part of A contains the
upper triangular part of the matrix A. If UPLO = 'L',
the leading N-by-N lower triangular part of A contains
the lower triangular part of the matrix A.
On exit, the lower triangle (if UPLO='L') or the upper
triangle (if UPLO='U') of A, including the diagonal, is
destroyed.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
VL (input) DOUBLE PRECISION
VU (input) DOUBLE PRECISION
If RANGE='V', the lower and upper bounds of the interval to
be searched for eigenvalues. VL < VU.
Not referenced if RANGE = 'A' or 'I'.
IL (input) INTEGER
IU (input) INTEGER
If RANGE='I', 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 = 'A' or 'V'.
ABSTOL (input) DOUBLE PRECISION
The absolute error tolerance for the eigenvalues.
An approximate eigenvalue is accepted as converged
when it is determined to lie in an interval [a,b]
of width less than or equal to
ABSTOL + EPS * max( |a|,|b| ) ,
where EPS is the machine precision. If ABSTOL is less than
or equal to zero, then EPS*|T| will be used in its place,
where |T| is the 1-norm of the tridiagonal matrix obtained
by reducing A to tridiagonal form.
Eigenvalues will be computed most accurately when ABSTOL is
set to twice the underflow threshold 2*SLAMCH('S'), not zero.
If this routine returns with INFO>0, indicating that some
eigenvectors did not converge, try setting ABSTOL to
2*SLAMCH('S').
See "Computing Small Singular Values of Bidiagonal Matrices
with Guaranteed High Relative Accuracy," by Demmel and
Kahan, LAPACK Working Note #3.
M (output) INTEGER
The total number of eigenvalues found. 0 <= M <= N.
If RANGE = 'A', M = N, and if RANGE = 'I', M = IU-IL+1.
W (output) DOUBLE PRECISION array, dimension (N)
On normal exit, the first M elements contain the selected
eigenvalues in ascending order.
Z (output) COMPLEX array, dimension (LDZ, max(1,M))
If JOBZ = 'V', then if INFO = 0, the first M columns of Z
contain the orthonormal eigenvectors of the matrix A
corresponding to the selected eigenvalues, with the i-th
column of Z holding the eigenvector associated with W(i).
If an eigenvector fails to converge, then that column of Z
contains the latest approximation to the eigenvector, and the
index of the eigenvector is returned in IFAIL.
If JOBZ = 'N', then Z is not referenced.
Note: the user must ensure that at least max(1,M) columns are
supplied in the array Z; if RANGE = 'V', the exact value of M
is not known in advance and an upper bound must be used.
LDZ (input) INTEGER
The leading dimension of the array Z. LDZ >= 1, and if
JOBZ = 'V', LDZ >= max(1,N).
WORK (workspace/output) COMPLEX array, dimension (LWORK)
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The length of the array WORK. LWORK >= max(1,2*N-1).
For optimal efficiency, LWORK >= (NB+1)*N,
where NB is the max of the blocksize for CHETRD and for
CUNMTR as returned by ILAENV.
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.
RWORK (workspace) DOUBLE PRECISION array, dimension (7*N)
IWORK (workspace) INTEGER array, dimension (5*N)
IFAIL (output) INTEGER array, dimension (N)
If JOBZ = 'V', then if INFO = 0, the first M elements of
IFAIL are zero. If INFO > 0, then IFAIL contains the
indices of the eigenvectors that failed to converge.
If JOBZ = 'N', then IFAIL is not referenced.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
> 0: if INFO = i, then i eigenvectors failed to converge.
Their indices are stored in array IFAIL.
===================================================================== */
char uplo_[2] = {uplo, 0};
char jobz_[2] = {jobz, 0};
char range_[2] = {range, 0};
magma_int_t izero = 0;
magma_int_t ione = 1;
char order[1];
magma_int_t indd, inde;
magma_int_t imax;
magma_int_t lopt, itmp1, indee;
magma_int_t lower, wantz;
magma_int_t i, j, jj, i__1;
magma_int_t alleig, valeig, indeig;
magma_int_t iscale, indibl;
magma_int_t indiwk, indisp, indtau;
magma_int_t indrwk, indwrk;
magma_int_t llwork, nsplit;
magma_int_t lquery;
magma_int_t iinfo;
float safmin;
float bignum;
float smlnum;
float eps, tmp1;
float anrm;
float sigma, d__1;
float rmin, rmax;
/* Function Body */
lower = lapackf77_lsame(uplo_, MagmaLowerStr);
wantz = lapackf77_lsame(jobz_, MagmaVecStr);
alleig = lapackf77_lsame(range_, "A");
valeig = lapackf77_lsame(range_, "V");
indeig = lapackf77_lsame(range_, "I");
lquery = lwork == -1;
*info = 0;
if (! (wantz || lapackf77_lsame(jobz_, MagmaNoVecStr))) {
*info = -1;
} else if (! (alleig || valeig || indeig)) {
*info = -2;
} else if (! (lower || lapackf77_lsame(uplo_, MagmaUpperStr))) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (lda < max(1,n)) {
*info = -6;
} else if (ldz < 1 || (wantz && ldz < n)) {
*info = -15;
} else {
if (valeig) {
if (n > 0 && vu <= vl) {
*info = -8;
}
} else if (indeig) {
if (il < 1 || il > max(1,n)) {
*info = -9;
} else if (iu < min(n,il) || iu > n) {
*info = -10;
}
}
}
magma_int_t nb = magma_get_chetrd_nb(n);
lopt = n * (nb + 1);
MAGMA_C_SET2REAL(work[0],(float)lopt);
if (lwork < lopt && ! lquery) {
*info = -17;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info));
return *info;
} else if (lquery) {
return *info;
}
*m = 0;
/* 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
lapackf77_cheevx(jobz_, range_, uplo_,
&n, a, &lda, &vl, &vu, &il, &iu, &abstol, m,
w, z, &ldz, work, &lwork,
rwork, iwork, ifail, info);
return *info;
}
--w;
--work;
--rwork;
--iwork;
--ifail;
/* 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 = lapackf77_clanhe("M", uplo_, &n, a, &lda, &rwork[1]);
iscale = 0;
if (anrm > 0. && anrm < rmin) {
iscale = 1;
sigma = rmin / anrm;
} else if (anrm > rmax) {
iscale = 1;
sigma = rmax / anrm;
}
if (iscale == 1) {
d__1 = 1.;
lapackf77_clascl(uplo_, &izero, &izero, &d__1, &sigma, &n, &n, a,
&lda, info);
if (abstol > 0.) {
abstol *= sigma;
}
if (valeig) {
vl *= sigma;
vu *= sigma;
}
}
/* Call CHETRD to reduce Hermitian matrix to tridiagonal form. */
indd = 1;
inde = indd + n;
indrwk = inde + n;
indtau = 1;
indwrk = indtau + n;
llwork = lwork - indwrk + 1;
magma_chetrd(uplo, n, a, lda, &rwork[indd], &rwork[inde], &work[indtau], &work[indwrk], llwork, &iinfo);
lopt = n + (magma_int_t)MAGMA_C_REAL(work[indwrk]);
/* If all eigenvalues are desired and ABSTOL is less than or equal to
zero, then call SSTERF or CUNGTR and CSTEQR. If this fails for
some eigenvalue, then try SSTEBZ. */
if ((alleig || (indeig && il == 1 && iu == n)) && abstol <= 0.) {
blasf77_scopy(&n, &rwork[indd], &ione, &w[1], &ione);
indee = indrwk + 2*n;
if (! wantz) {
i__1 = n - 1;
blasf77_scopy(&i__1, &rwork[inde], &ione, &rwork[indee], &ione);
lapackf77_ssterf(&n, &w[1], &rwork[indee], info);
}
else {
lapackf77_clacpy("A", &n, &n, a, &lda, z, &ldz);
lapackf77_cungtr(uplo_, &n, z, &ldz, &work[indtau], &work[indwrk], &llwork, &iinfo);
i__1 = n - 1;
blasf77_scopy(&i__1, &rwork[inde], &ione, &rwork[indee], &ione);
lapackf77_csteqr(jobz_, &n, &w[1], &rwork[indee], z, &ldz, &rwork[indrwk], info);
if (*info == 0) {
for (i = 1; i <= n; ++i) {
ifail[i] = 0;
}
}
}
if (*info == 0) {
*m = n;
}
}
/* Otherwise, call SSTEBZ and, if eigenvectors are desired, CSTEIN. */
if (*m == 0) {
*info = 0;
if (wantz) {
*(unsigned char *)order = 'B';
} else {
*(unsigned char *)order = 'E';
}
indibl = 1;
indisp = indibl + n;
indiwk = indisp + n;
lapackf77_sstebz(range_, order, &n, &vl, &vu, &il, &iu, &abstol, &rwork[indd], &rwork[inde], m,
&nsplit, &w[1], &iwork[indibl], &iwork[indisp], &rwork[indrwk], &iwork[indiwk], info);
if (wantz) {
lapackf77_cstein(&n, &rwork[indd], &rwork[inde], m, &w[1], &iwork[indibl], &iwork[indisp],
z, &ldz, &rwork[indrwk], &iwork[indiwk], &ifail[1], info);
/* Apply unitary matrix used in reduction to tridiagonal
form to eigenvectors returned by CSTEIN. */
magma_cunmtr(MagmaLeft, uplo, MagmaNoTrans, n, *m, a, lda, &work[indtau],
z, ldz, &work[indwrk], llwork, &iinfo);
}
}
/* If matrix was scaled, then rescale eigenvalues appropriately. */
if (iscale == 1) {
if (*info == 0) {
imax = *m;
} else {
imax = *info - 1;
}
d__1 = 1. / sigma;
blasf77_sscal(&imax, &d__1, &w[1], &ione);
}
/* If eigenvalues are not in order, then sort them, along with
eigenvectors. */
if (wantz) {
for (j = 1; j <= *m-1; ++j) {
i = 0;
tmp1 = w[j];
for (jj = j + 1; jj <= *m; ++jj) {
if (w[jj] < tmp1) {
i = jj;
tmp1 = w[jj];
}
}
if (i != 0) {
itmp1 = iwork[indibl + i - 1];
w[i] = w[j];
iwork[indibl + i - 1] = iwork[indibl + j - 1];
w[j] = tmp1;
iwork[indibl + j - 1] = itmp1;
blasf77_cswap(&n, z + (i-1)*ldz, &ione, z + (j-1)*ldz, &ione);
if (*info != 0) {
itmp1 = ifail[i];
ifail[i] = ifail[j];
ifail[j] = itmp1;
}
}
}
}
/* Set WORK(1) to optimal complex workspace size. */
work[1] = MAGMA_C_MAKE((float) lopt, 0.);
return *info;
} /* magma_cheevx */
magma_int_t magma_clatrsd(
magma_uplo_t uplo, magma_trans_t trans, magma_diag_t diag, magma_bool_t normin,
magma_int_t n, const magmaFloatComplex *A, magma_int_t lda,
magmaFloatComplex lambda,
magmaFloatComplex *x,
float *scale, float *cnorm,
magma_int_t *info)
{
#define A(i,j) (A + (i) + (j)*lda)
/* constants */
const magma_int_t ione = 1;
const float d_half = 0.5;
const magmaFloatComplex c_zero = MAGMA_C_ZERO;
const magmaFloatComplex c_one = MAGMA_C_ONE;
/* System generated locals */
magma_int_t len;
magmaFloatComplex ztmp;
/* Local variables */
magma_int_t i, j;
float xj, rec, tjj;
magma_int_t jinc;
float xbnd;
magma_int_t imax;
float tmax;
magmaFloatComplex tjjs;
float xmax, grow;
float tscal;
magmaFloatComplex uscal;
magma_int_t jlast;
magmaFloatComplex csumj;
float bignum;
magma_int_t jfirst;
float smlnum;
/* Function Body */
*info = 0;
magma_int_t upper = (uplo == MagmaUpper);
magma_int_t notran = (trans == MagmaNoTrans);
magma_int_t nounit = (diag == MagmaNonUnit);
/* Test the input parameters. */
if ( ! upper && uplo != MagmaLower ) {
*info = -1;
}
else if (! notran &&
trans != MagmaTrans &&
trans != MagmaConjTrans) {
*info = -2;
}
else if ( ! nounit && diag != MagmaUnit ) {
*info = -3;
}
else if ( ! (normin == MagmaTrue) &&
! (normin == MagmaFalse) ) {
*info = -4;
}
else if ( n < 0 ) {
*info = -5;
}
else if ( lda < max(1,n) ) {
*info = -7;
}
if ( *info != 0 ) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if ( n == 0 ) {
return *info;
}
/* Determine machine dependent parameters to control overflow. */
smlnum = lapackf77_slamch( "Safe minimum" );
bignum = 1. / smlnum;
lapackf77_slabad( &smlnum, &bignum );
smlnum /= lapackf77_slamch( "Precision" );
bignum = 1. / smlnum;
*scale = 1.;
if ( normin == MagmaFalse ) {
/* Compute the 1-norm of each column, not including the diagonal. */
if ( upper ) {
/* A is upper triangular. */
cnorm[0] = 0.;
for( j = 1; j < n; ++j ) {
cnorm[j] = magma_cblas_scasum( j, A(0,j), ione );
}
}
else {
/* A is lower triangular. */
for( j = 0; j < n-1; ++j ) {
cnorm[j] = magma_cblas_scasum( n-(j+1), A(j+1,j), ione );
}
cnorm[n-1] = 0.;
}
}
/* Scale the column norms by TSCAL if the maximum element in CNORM is */
/* greater than BIGNUM/2. */
imax = blasf77_isamax( &n, &cnorm[0], &ione ) - 1;
tmax = cnorm[imax];
if ( tmax <= bignum * 0.5 ) {
tscal = 1.;
}
else {
tscal = 0.5 / (smlnum * tmax);
blasf77_sscal( &n, &tscal, &cnorm[0], &ione );
}
/* ================================================================= */
/* Compute a bound on the computed solution vector to see if the */
/* Level 2 BLAS routine CTRSV can be used. */
xmax = 0.;
for( j = 0; j < n; ++j ) {
xmax = max( xmax, 0.5*MAGMA_C_ABS1( x[j] ));
}
xbnd = xmax;
if ( notran ) {
/* ---------------------------------------- */
/* Compute the growth in A * x = b. */
if ( upper ) {
jfirst = n-1;
jlast = 0;
jinc = -1;
}
else {
jfirst = 0;
jlast = n;
jinc = 1;
}
if ( tscal != 1. ) {
grow = 0.;
goto L60;
}
/* A is non-unit triangular. */
/* Compute GROW = 1/G(j) and XBND = 1/M(j). */
/* Initially, G(0) = max{x(i), i=1,...,n}. */
grow = 0.5 / max( xbnd, smlnum );
xbnd = grow;
for( j = jfirst; (jinc < 0 ? j >= jlast : j < jlast); j += jinc ) {
/* Exit the loop if the growth factor is too small. */
if ( grow <= smlnum ) {
goto L60;
}
if ( nounit ) {
tjjs = *A(j,j) - lambda;
}
else {
tjjs = c_one - lambda;
}
tjj = MAGMA_C_ABS1( tjjs );
if ( tjj >= smlnum ) {
/* M(j) = G(j-1) / abs(A(j,j)) */
xbnd = min( xbnd, min(1.,tjj)*grow );
}
else {
/* M(j) could overflow, set XBND to 0. */
xbnd = 0.;
}
if ( tjj + cnorm[j] >= smlnum ) {
/* G(j) = G(j-1)*( 1 + CNORM(j) / abs(A(j,j)) ) */
grow *= (tjj / (tjj + cnorm[j]));
}
else {
/* G(j) could overflow, set GROW to 0. */
grow = 0.;
}
}
grow = xbnd;
L60:
;
}
else {
/* ---------------------------------------- */
/* Compute the growth in A**T * x = b or A**H * x = b. */
if ( upper ) {
jfirst = 0;
jlast = n;
jinc = 1;
}
else {
jfirst = n-1;
jlast = 0;
jinc = -1;
}
if ( tscal != 1. ) {
grow = 0.;
goto L90;
}
/* A is non-unit triangular. */
/* Compute GROW = 1/G(j) and XBND = 1/M(j). */
/* Initially, M(0) = max{x(i), i=1,...,n}. */
grow = 0.5 / max( xbnd, smlnum );
xbnd = grow;
for( j = jfirst; (jinc < 0 ? j >= jlast : j < jlast); j += jinc ) {
/* Exit the loop if the growth factor is too small. */
if ( grow <= smlnum ) {
goto L90;
}
/* G(j) = max( G(j-1), M(j-1)*( 1 + CNORM(j) ) ) */
xj = 1. + cnorm[j];
grow = min( grow, xbnd / xj );
if ( nounit ) {
tjjs = *A(j,j) - lambda;
}
else {
tjjs = c_one - lambda;
}
tjj = MAGMA_C_ABS1( tjjs );
if ( tjj >= smlnum ) {
/* M(j) = M(j-1)*( 1 + CNORM(j) ) / abs(A(j,j)) */
if ( xj > tjj ) {
xbnd *= (tjj / xj);
}
}
else {
/* M(j) could overflow, set XBND to 0. */
xbnd = 0.;
}
}
grow = min( grow, xbnd );
L90:
;
}
/* ================================================================= */
/* Due to modified diagonal, we can't use regular BLAS ctrsv. */
/* Use a Level 1 BLAS solve, scaling intermediate results. */
if ( xmax > bignum * 0.5 ) {
/* Scale X so that its components are less than or equal to */
/* BIGNUM in absolute value. */
*scale = (bignum * 0.5) / xmax;
blasf77_csscal( &n, scale, &x[0], &ione );
xmax = bignum;
}
else {
xmax *= 2.;
}
if ( notran ) {
/* ---------------------------------------- */
/* Solve A * x = b */
for( j = jfirst; (jinc < 0 ? j >= jlast : j < jlast); j += jinc ) {
/* Compute x(j) = b(j) / A(j,j), scaling x if necessary. */
xj = MAGMA_C_ABS1( x[j] );
if ( nounit ) {
tjjs = (*A(j,j) - lambda ) * tscal;
}
else {
tjjs = (c_one - lambda) * tscal;
if ( tscal == 1. ) {
goto L110;
}
}
tjj = MAGMA_C_ABS1( tjjs );
if ( tjj > smlnum ) {
/* abs(A(j,j)) > SMLNUM: */
if ( tjj < 1. ) {
if ( xj > tjj * bignum ) {
/* Scale x by 1/b(j). */
rec = 1. / xj;
blasf77_csscal( &n, &rec, &x[0], &ione );
*scale *= rec;
xmax *= rec;
}
}
x[j] = x[j] / tjjs;
xj = MAGMA_C_ABS1( x[j] );
}
else if ( tjj > 0. ) {
/* 0 < abs(A(j,j)) <= SMLNUM: */
if ( xj > tjj * bignum ) {
/* Scale x by (1/abs(x(j)))*abs(A(j,j))*BIGNUM */
/* to avoid overflow when dividing by A(j,j). */
rec = (tjj * bignum) / xj;
if ( cnorm[j] > 1. ) {
/* Scale by 1/CNORM(j) to avoid overflow when */
/* multiplying x(j) times column j. */
rec /= cnorm[j];
}
blasf77_csscal( &n, &rec, &x[0], &ione );
*scale *= rec;
xmax *= rec;
}
x[j] = x[j] / tjjs;
xj = MAGMA_C_ABS1( x[j] );
}
else {
/* A(j,j) = 0: Set x(1:n) = 0, x(j) = 1, and */
/* scale = 0, and compute a solution to A*x = 0. */
for( i = 0; i < n; ++i ) {
x[i] = c_zero;
}
x[j] = c_one;
xj = 1.;
*scale = 0.;
xmax = 0.;
}
L110:
/* Scale x if necessary to avoid overflow when adding a */
/* multiple of column j of A. */
if ( xj > 1. ) {
rec = 1. / xj;
if ( cnorm[j] > (bignum - xmax) * rec ) {
/* Scale x by 1/(2*abs(x(j))). */
rec *= 0.5;
blasf77_csscal( &n, &rec, &x[0], &ione );
*scale *= rec;
}
}
else if ( xj * cnorm[j] > bignum - xmax ) {
/* Scale x by 1/2. */
blasf77_csscal( &n, &d_half, &x[0], &ione );
*scale *= 0.5;
}
if ( upper ) {
if ( j > 0 ) {
/* Compute the update */
/* x(1:j-1) := x(1:j-1) - x(j) * A(1:j-1,j) */
len = j;
ztmp = -tscal * x[j];
blasf77_caxpy( &len, &ztmp, A(0,j), &ione, &x[0], &ione );
i = blasf77_icamax( &len, &x[0], &ione ) - 1;
xmax = MAGMA_C_ABS1( x[i] );
}
}
else {
if ( j < n-1 ) {
/* Compute the update */
/* x(j+1:n) := x(j+1:n) - x(j) * A(j+1:n,j) */
len = n - (j+1);
ztmp = -tscal * x[j];
blasf77_caxpy( &len, &ztmp, A(j+1,j), &ione, &x[j + 1], &ione );
i = j + blasf77_icamax( &len, &x[j + 1], &ione );
xmax = MAGMA_C_ABS1( x[i] );
}
}
}
}
else if ( trans == MagmaTrans ) {
/* ---------------------------------------- */
/* Solve A**T * x = b */
for( j = jfirst; (jinc < 0 ? j >= jlast : j < jlast); j += jinc ) {
/* Compute x(j) = b(j) - sum A(k,j)*x(k). */
/* k<>j */
xj = MAGMA_C_ABS1( x[j] );
uscal = MAGMA_C_MAKE( tscal, 0. );
rec = 1. / max( xmax, 1. );
if ( cnorm[j] > (bignum - xj) * rec ) {
/* If x(j) could overflow, scale x by 1/(2*XMAX). */
rec *= 0.5;
if ( nounit ) {
tjjs = (*A(j,j) - lambda) * tscal;
}
else {
tjjs = (c_one - lambda) * tscal;
}
tjj = MAGMA_C_ABS1( tjjs );
if ( tjj > 1. ) {
/* Divide by A(j,j) when scaling x if A(j,j) > 1. */
rec = min( 1., rec * tjj );
uscal = uscal / tjjs;
}
if ( rec < 1. ) {
blasf77_csscal( &n, &rec, &x[0], &ione );
*scale *= rec;
xmax *= rec;
}
}
csumj = c_zero;
if ( uscal == c_one ) {
/* If the scaling needed for A in the dot product is 1, */
/* call CDOTU to perform the dot product. */
if ( upper ) {
csumj = magma_cblas_cdotu( j, A(0,j), ione, &x[0], ione );
}
else if ( j < n-1 ) {
csumj = magma_cblas_cdotu( n-(j+1), A(j+1,j), ione, &x[j+1], ione );
}
}
else {
/* Otherwise, use in-line code for the dot product. */
if ( upper ) {
for( i = 0; i < j; ++i ) {
csumj += (*A(i,j) * uscal) * x[i];
}
}
else if ( j < n-1 ) {
for( i = j+1; i < n; ++i ) {
csumj += (*A(i,j) * uscal) * x[i];
}
}
}
if ( uscal == MAGMA_C_MAKE( tscal, 0. )) {
/* Compute x(j) := ( x(j) - CSUMJ ) / A(j,j) if 1/A(j,j) */
/* was not used to scale the dotproduct. */
x[j] -= csumj;
xj = MAGMA_C_ABS1( x[j] );
if ( nounit ) {
tjjs = (*A(j,j) - lambda) * tscal;
}
else {
tjjs = (c_one - lambda) * tscal;
if ( tscal == 1. ) {
goto L160;
}
}
/* Compute x(j) = x(j) / A(j,j), scaling if necessary. */
tjj = MAGMA_C_ABS1( tjjs );
if ( tjj > smlnum ) {
/* abs(A(j,j)) > SMLNUM: */
if ( tjj < 1. ) {
if ( xj > tjj * bignum ) {
/* Scale X by 1/abs(x(j)). */
rec = 1. / xj;
blasf77_csscal( &n, &rec, &x[0], &ione );
*scale *= rec;
xmax *= rec;
}
}
x[j] = x[j] / tjjs;
}
else if ( tjj > 0. ) {
/* 0 < abs(A(j,j)) <= SMLNUM: */
if ( xj > tjj * bignum ) {
/* Scale x by (1/abs(x(j)))*abs(A(j,j))*BIGNUM. */
rec = (tjj * bignum) / xj;
blasf77_csscal( &n, &rec, &x[0], &ione );
*scale *= rec;
xmax *= rec;
}
x[j] = x[j] / tjjs;
}
else {
/* A(j,j) = 0: Set x(1:n) = 0, x(j) = 1, and */
/* scale = 0 and compute a solution to A**T *x = 0. */
for( i = 0; i < n; ++i ) {
x[i] = c_zero;
}
x[j] = c_one;
*scale = 0.;
xmax = 0.;
}
L160:
;
}
else {
/* Compute x(j) := x(j) / A(j,j) - CSUMJ if the dot */
/* product has already been divided by 1/A(j,j). */
x[j] = (x[j] / tjjs) - csumj;
}
xmax = max( xmax, MAGMA_C_ABS1( x[j] ));
}
}
else {
/* ---------------------------------------- */
/* Solve A**H * x = b */
for( j = jfirst; (jinc < 0 ? j >= jlast : j < jlast); j += jinc ) {
/* Compute x(j) = b(j) - sum A(k,j)*x(k). */
/* k<>j */
xj = MAGMA_C_ABS1( x[j] );
uscal = MAGMA_C_MAKE( tscal, 0. );
rec = 1. / max(xmax, 1.);
if ( cnorm[j] > (bignum - xj) * rec ) {
/* If x(j) could overflow, scale x by 1/(2*XMAX). */
rec *= 0.5;
if ( nounit ) {
tjjs = MAGMA_C_CONJ( *A(j,j) - lambda ) * tscal;
}
else {
tjjs = (c_one - lambda) * tscal;
}
tjj = MAGMA_C_ABS1( tjjs );
if ( tjj > 1. ) {
/* Divide by A(j,j) when scaling x if A(j,j) > 1. */
rec = min( 1., rec * tjj );
uscal = uscal / tjjs;
}
if ( rec < 1. ) {
blasf77_csscal( &n, &rec, &x[0], &ione );
*scale *= rec;
xmax *= rec;
}
}
csumj = c_zero;
if ( uscal == c_one ) {
/* If the scaling needed for A in the dot product is 1, */
/* call CDOTC to perform the dot product. */
if ( upper ) {
csumj = magma_cblas_cdotc( j, A(0,j), ione, &x[0], ione );
}
else if ( j < n-1 ) {
csumj = magma_cblas_cdotc( n-(j+1), A(j+1,j), ione, &x[j+1], ione );
}
}
else {
/* Otherwise, use in-line code for the dot product. */
if ( upper ) {
for( i = 0; i < j; ++i ) {
csumj += (MAGMA_C_CONJ( *A(i,j) ) * uscal) * x[i];
}
}
else if ( j < n-1 ) {
for( i = j + 1; i < n; ++i ) {
csumj += (MAGMA_C_CONJ( *A(i,j) ) * uscal) * x[i];
}
}
}
if ( uscal == tscal ) {
/* Compute x(j) := ( x(j) - CSUMJ ) / A(j,j) if 1/A(j,j) */
/* was not used to scale the dotproduct. */
x[j] -= csumj;
xj = MAGMA_C_ABS1( x[j] );
if ( nounit ) {
tjjs = MAGMA_C_CONJ( *A(j,j) - lambda ) * tscal;
}
else {
tjjs = (c_one - lambda) * tscal;
if ( tscal == 1. ) {
goto L210;
}
}
/* Compute x(j) = x(j) / A(j,j), scaling if necessary. */
tjj = MAGMA_C_ABS1( tjjs );
if ( tjj > smlnum ) {
/* abs(A(j,j)) > SMLNUM: */
if ( tjj < 1. ) {
if ( xj > tjj * bignum ) {
/* Scale X by 1/abs(x(j)). */
rec = 1. / xj;
blasf77_csscal( &n, &rec, &x[0], &ione );
*scale *= rec;
xmax *= rec;
}
}
x[j] = x[j] / tjjs;
}
else if ( tjj > 0. ) {
/* 0 < abs(A(j,j)) <= SMLNUM: */
if ( xj > tjj * bignum ) {
/* Scale x by (1/abs(x(j)))*abs(A(j,j))*BIGNUM. */
rec = (tjj * bignum) / xj;
blasf77_csscal( &n, &rec, &x[0], &ione );
*scale *= rec;
xmax *= rec;
}
x[j] = x[j] / tjjs;
}
else {
/* A(j,j) = 0: Set x(1:n) = 0, x(j) = 1, and */
/* scale = 0 and compute a solution to A**H *x = 0. */
for( i = 0; i < n; ++i ) {
x[i] = c_zero;
}
x[j] = c_one;
*scale = 0.;
xmax = 0.;
}
L210:
;
}
else {
/* Compute x(j) := x(j) / A(j,j) - CSUMJ if the dot */
/* product has already been divided by 1/A(j,j). */
x[j] = (x[j] / tjjs) - csumj;
}
xmax = max( xmax, MAGMA_C_ABS1( x[j] ));
}
}
*scale /= tscal;
/* Scale the column norms by 1/TSCAL for return. */
if ( tscal != 1. ) {
float d = 1. / tscal;
blasf77_sscal( &n, &d, &cnorm[0], &ione );
}
return *info;
} /* end clatrsd */
extern "C" magma_int_t
magma_cbicgstab(
magma_c_sparse_matrix A, magma_c_vector b, magma_c_vector *x,
magma_c_solver_par *solver_par,
magma_queue_t queue )
{
// set queue for old dense routines
magma_queue_t orig_queue;
magmablasGetKernelStream( &orig_queue );
// prepare solver feedback
solver_par->solver = Magma_BICGSTAB;
solver_par->numiter = 0;
solver_par->info = MAGMA_SUCCESS;
// some useful variables
magmaFloatComplex c_zero = MAGMA_C_ZERO, c_one = MAGMA_C_ONE,
c_mone = MAGMA_C_NEG_ONE;
magma_int_t dofs = A.num_rows;
// workspace
magma_c_vector r,rr,p,v,s,t;
magma_c_vinit( &r, Magma_DEV, dofs, c_zero, queue );
magma_c_vinit( &rr, Magma_DEV, dofs, c_zero, queue );
magma_c_vinit( &p, Magma_DEV, dofs, c_zero, queue );
magma_c_vinit( &v, Magma_DEV, dofs, c_zero, queue );
magma_c_vinit( &s, Magma_DEV, dofs, c_zero, queue );
magma_c_vinit( &t, Magma_DEV, dofs, c_zero, queue );
// solver variables
magmaFloatComplex alpha, beta, omega, rho_old, rho_new;
float nom, betanom, nom0, r0, den, res;
// solver setup
magma_cscal( dofs, c_zero, x->dval, 1) ; // x = 0
magma_ccopy( dofs, b.dval, 1, r.dval, 1 ); // r = b
magma_ccopy( dofs, b.dval, 1, rr.dval, 1 ); // rr = b
nom0 = betanom = magma_scnrm2( dofs, r.dval, 1 ); // nom = || r ||
nom = nom0*nom0;
rho_old = omega = alpha = MAGMA_C_MAKE( 1.0, 0. );
solver_par->init_res = nom0;
magma_c_spmv( c_one, A, r, c_zero, v, queue ); // z = A r
den = MAGMA_C_REAL( magma_cdotc(dofs, v.dval, 1, r.dval, 1) ); // den = z' * r
if ( (r0 = nom * solver_par->epsilon) < ATOLERANCE )
r0 = ATOLERANCE;
if ( nom < r0 ) {
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
}
// check positive definite
if (den <= 0.0) {
printf("Operator A is not postive definite. (Ar,r) = %f\n", den);
magmablasSetKernelStream( orig_queue );
return MAGMA_NONSPD;
solver_par->info = MAGMA_NONSPD;;
}
//Chronometry
real_Double_t tempo1, tempo2;
tempo1 = magma_sync_wtime( queue );
if ( solver_par->verbose > 0 ) {
solver_par->res_vec[0] = nom0;
solver_par->timing[0] = 0.0;
}
// start iteration
for( solver_par->numiter= 1; solver_par->numiter<solver_par->maxiter;
solver_par->numiter++ ) {
rho_new = magma_cdotc( dofs, rr.dval, 1, r.dval, 1 ); // rho=<rr,r>
beta = rho_new/rho_old * alpha/omega; // beta=rho/rho_old *alpha/omega
magma_cscal( dofs, beta, p.dval, 1 ); // p = beta*p
magma_caxpy( dofs, c_mone * omega * beta, v.dval, 1 , p.dval, 1 );
// p = p-omega*beta*v
magma_caxpy( dofs, c_one, r.dval, 1, p.dval, 1 ); // p = p+r
magma_c_spmv( c_one, A, p, c_zero, v, queue ); // v = Ap
alpha = rho_new / magma_cdotc( dofs, rr.dval, 1, v.dval, 1 );
magma_ccopy( dofs, r.dval, 1 , s.dval, 1 ); // s=r
magma_caxpy( dofs, c_mone * alpha, v.dval, 1 , s.dval, 1 ); // s=s-alpha*v
magma_c_spmv( c_one, A, s, c_zero, t, queue ); // t=As
omega = magma_cdotc( dofs, t.dval, 1, s.dval, 1 ) // omega = <s,t>/<t,t>
/ magma_cdotc( dofs, t.dval, 1, t.dval, 1 );
magma_caxpy( dofs, alpha, p.dval, 1 , x->dval, 1 ); // x=x+alpha*p
magma_caxpy( dofs, omega, s.dval, 1 , x->dval, 1 ); // x=x+omega*s
magma_ccopy( dofs, s.dval, 1 , r.dval, 1 ); // r=s
magma_caxpy( dofs, c_mone * omega, t.dval, 1 , r.dval, 1 ); // r=r-omega*t
res = betanom = magma_scnrm2( dofs, r.dval, 1 );
nom = betanom*betanom;
rho_old = rho_new; // rho_old=rho
if ( solver_par->verbose > 0 ) {
tempo2 = magma_sync_wtime( queue );
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) res;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
if ( res/nom0 < solver_par->epsilon ) {
break;
}
}
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t) tempo2-tempo1;
float residual;
magma_cresidual( A, b, *x, &residual, queue );
solver_par->iter_res = res;
solver_par->final_res = residual;
if ( solver_par->numiter < solver_par->maxiter) {
solver_par->info = MAGMA_SUCCESS;
} else if ( solver_par->init_res > solver_par->final_res ) {
if ( solver_par->verbose > 0 ) {
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
solver_par->info = MAGMA_SLOW_CONVERGENCE;
}
else {
if ( solver_par->verbose > 0 ) {
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
solver_par->info = MAGMA_DIVERGENCE;
}
magma_c_vfree(&r, queue );
magma_c_vfree(&rr, queue );
magma_c_vfree(&p, queue );
magma_c_vfree(&v, queue );
magma_c_vfree(&s, queue );
magma_c_vfree(&t, queue );
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
} /* magma_cbicgstab */
extern "C" magma_int_t
magma_chegvd_m(magma_int_t nrgpu, magma_int_t itype, char jobz, char uplo, magma_int_t n,
magmaFloatComplex *a, magma_int_t lda, magmaFloatComplex *b, magma_int_t ldb,
float *w, magmaFloatComplex *work, magma_int_t lwork,
float *rwork, magma_int_t lrwork,
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
=======
CHEGVD computes all the eigenvalues, and optionally, the eigenvectors
of a complex generalized Hermitian-definite eigenproblem, of the form
A*x=(lambda)*B*x, A*Bx=(lambda)*x, or B*A*x=(lambda)*x. Here A and
B are assumed to be Hermitian and B is also positive definite.
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
=========
ITYPE (input) INTEGER
Specifies the problem type to be solved:
= 1: A*x = (lambda)*B*x
= 2: A*B*x = (lambda)*x
= 3: B*A*x = (lambda)*x
JOBZ (input) CHARACTER*1
= 'N': Compute eigenvalues only;
= 'V': Compute eigenvalues and eigenvectors.
UPLO (input) CHARACTER*1
= 'U': Upper triangles of A and B are stored;
= 'L': Lower triangles of A and B are stored.
N (input) INTEGER
The order of the matrices A and B. N >= 0.
A (input/output) COMPLEX array, dimension (LDA, N)
On entry, the Hermitian matrix A. If UPLO = 'U', the
leading N-by-N upper triangular part of A contains the
upper triangular part of the matrix A. 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
matrix Z of eigenvectors. The eigenvectors are normalized
as follows:
if ITYPE = 1 or 2, Z**H*B*Z = I;
if ITYPE = 3, Z**H*inv(B)*Z = I.
If JOBZ = 'N', then on exit the upper triangle (if UPLO='U')
or the lower triangle (if UPLO='L') of A, including the
diagonal, is destroyed.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
B (input/output) COMPLEX array, dimension (LDB, N)
On entry, the Hermitian matrix B. If UPLO = 'U', the
leading N-by-N upper triangular part of B contains the
upper triangular part of the matrix B. If UPLO = 'L',
the leading N-by-N lower triangular part of B contains
the lower triangular part of the matrix B.
On exit, if INFO <= N, the part of B containing the matrix is
overwritten by the triangular factor U or L from the Cholesky
factorization B = U**H*U or B = L*L**H.
LDB (input) INTEGER
The leading dimension of the array B. LDB >= max(1,N).
W (output) DOUBLE PRECISION array, dimension (N)
If INFO = 0, the eigenvalues in ascending order.
WORK (workspace/output) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) 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 >= N + 1.
If JOBZ = 'V' and N > 1, LWORK >= 2*N*nb + N**2.
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal sizes of the WORK, RWORK and
IWORK arrays, returns these values as the first entries of
the WORK, RWORK and IWORK arrays, and no error message
related to LWORK or LRWORK or LIWORK is issued by XERBLA.
RWORK (workspace/output) DOUBLE PRECISION array, dimension (MAX(1,LRWORK))
On exit, if INFO = 0, RWORK(1) returns the optimal LRWORK.
LRWORK (input) INTEGER
The dimension of the array RWORK.
If N <= 1, LRWORK >= 1.
If JOBZ = 'N' and N > 1, LRWORK >= N.
If JOBZ = 'V' and N > 1, LRWORK >= 1 + 5*N + 2*N**2.
If LRWORK = -1, then a workspace query is assumed; the
routine only calculates the optimal sizes of the WORK, RWORK
and IWORK arrays, returns these values as the first entries
of the WORK, RWORK and IWORK arrays, and no error message
related to LWORK or LRWORK or LIWORK is issued by XERBLA.
IWORK (workspace/output) INTEGER array, dimension (MAX(1,LIWORK))
On exit, if INFO = 0, IWORK(1) 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, RWORK
and IWORK arrays, returns these values as the first entries
of the WORK, RWORK and IWORK arrays, and no error message
related to LWORK or LRWORK or LIWORK is issued by XERBLA.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
> 0: CPOTRF or CHEEVD returned an error code:
<= N: 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);
> N: if INFO = N + i, for 1 <= i <= N, then the leading
minor of order i of B is not positive definite.
The factorization of B could not be completed and
no eigenvalues or eigenvectors were computed.
Further Details
===============
Based on contributions by
Mark Fahey, Department of Mathematics, Univ. of Kentucky, USA
Modified so that no backsubstitution is performed if CHEEVD fails to
converge (NEIG in old code could be greater than N causing out of
bounds reference to A - reported by Ralf Meyer). Also corrected the
description of INFO and the test on ITYPE. Sven, 16 Feb 05.
===================================================================== */
char uplo_[2] = {uplo, 0};
char jobz_[2] = {jobz, 0};
magmaFloatComplex c_one = MAGMA_C_ONE;
magma_int_t lower;
char trans[1];
magma_int_t wantz;
magma_int_t lquery;
magma_int_t lwmin;
magma_int_t liwmin;
magma_int_t lrwmin;
magma_queue_t stream;
magma_queue_create( &stream );
wantz = lapackf77_lsame(jobz_, MagmaVecStr);
lower = lapackf77_lsame(uplo_, MagmaLowerStr);
lquery = lwork == -1 || lrwork == -1 || liwork == -1;
*info = 0;
if (itype < 1 || itype > 3) {
*info = -1;
} else if (! (wantz || lapackf77_lsame(jobz_, MagmaNoVecStr))) {
*info = -2;
} else if (! (lower || lapackf77_lsame(uplo_, MagmaUpperStr))) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (lda < max(1,n)) {
*info = -6;
} else if (ldb < max(1,n)) {
*info = -8;
}
magma_int_t nb = magma_get_chetrd_nb( n );
if ( n <= 1 ) {
lwmin = 1;
lrwmin = 1;
liwmin = 1;
}
else if ( wantz ) {
lwmin = 2*n + n*n;
lrwmin = 1 + 5*n + 2*n*n;
liwmin = 3 + 5*n;
}
else {
lwmin = n + n*nb;
lrwmin = n;
liwmin = 1;
}
work[0] = MAGMA_C_MAKE( lwmin * (1. + lapackf77_slamch("Epsilon")), 0.); // round up
rwork[0] = lrwmin * (1. + lapackf77_slamch("Epsilon"));
iwork[0] = liwmin;
if (lwork < lwmin && ! lquery) {
*info = -11;
} else if (lrwork < lrwmin && ! lquery) {
*info = -13;
} else if (liwork < liwmin && ! lquery) {
*info = -15;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (n == 0) {
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
lapackf77_chegvd(&itype, jobz_, uplo_,
&n, a, &lda, b, &ldb,
w, work, &lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, &lrwork,
#endif
iwork, &liwork, info);
return *info;
}
//
#ifdef ENABLE_TIMER
magma_timestr_t start, end;
start = get_current_time();
#endif
magma_cpotrf_m(nrgpu, uplo_[0], n, b, ldb, info);
if (*info != 0) {
*info = n + *info;
return *info;
}
#ifdef ENABLE_TIMER
end = get_current_time();
printf("time cpotrf = %6.2f\n", GetTimerValue(start,end)/1000.);
start = get_current_time();
#endif
/* Transform problem to standard eigenvalue problem and solve. */
magma_chegst_m(nrgpu, itype, uplo_[0], n, a, lda, b, ldb, info);
#ifdef ENABLE_TIMER
end = get_current_time();
printf("time chegst = %6.2f\n", GetTimerValue(start,end)/1000.);
start = get_current_time();
#endif
magma_cheevd_m(nrgpu, jobz_[0], uplo_[0], n, a, lda, w, work, lwork, rwork, lrwork, iwork, liwork, info);
#ifdef ENABLE_TIMER
end = get_current_time();
printf("time cheevd = %6.2f\n", GetTimerValue(start,end)/1000.);
#endif
if (wantz && *info == 0)
{
#ifdef ENABLE_TIMER
start = get_current_time();
#endif
/* Backtransform eigenvectors to the original problem. */
if (itype == 1 || itype == 2)
{
/* For A*x=(lambda)*B*x and A*B*x=(lambda)*x;
backtransform eigenvectors: x = inv(L)'*y or inv(U)*y */
if (lower) {
*(unsigned char *)trans = MagmaConjTrans;
} else {
*(unsigned char *)trans = MagmaNoTrans;
}
magma_ctrsm_m(nrgpu, MagmaLeft, uplo_[0], *trans, MagmaNonUnit,
n, n, c_one, b, ldb, a, lda);
}
else if (itype == 3)
{
/* For B*A*x=(lambda)*x;
backtransform eigenvectors: x = L*y or U'*y */
if (lower) {
*(unsigned char *)trans = MagmaNoTrans;
} else {
*(unsigned char *)trans = MagmaConjTrans;
}
//magma_ctrmm(MagmaLeft, uplo_[0], *trans, MagmaNonUnit,
// n, n, c_one, db, lddb, da, ldda);
}
#ifdef ENABLE_TIMER
end = get_current_time();
printf("time setmatrices trsm/mm + getmatrices = %6.2f\n", GetTimerValue(start,end)/1000.);
#endif
}
work[0] = MAGMA_C_MAKE( lwmin * (1. + lapackf77_slamch("Epsilon")), 0.); // round up
rwork[0] = lrwmin * (1. + lapackf77_slamch("Epsilon"));
iwork[0] = liwmin;
return *info;
} /* magma_chegvd_m */
/**
Purpose
-------
If VECT = MagmaQ, CUNMBR overwrites the general complex M-by-N matrix C with
SIDE = MagmaLeft SIDE = MagmaRight
TRANS = MagmaNoTrans: Q*C C*Q
TRANS = Magma_ConjTrans: Q**H*C C*Q**H
If VECT = MagmaP, CUNMBR overwrites the general complex M-by-N matrix C with
SIDE = MagmaLeft SIDE = MagmaRight
TRANS = MagmaNoTrans: P*C C*P
TRANS = Magma_ConjTrans: P**H*C C*P**H
Here Q and P**H are the unitary matrices determined by CGEBRD when
reducing A complex matrix A to bidiagonal form: A = Q*B * P**H. Q
and P**H are defined as products of elementary reflectors H(i) and
G(i) respectively.
Let nq = m if SIDE = MagmaLeft and nq = n if SIDE = MagmaRight. Thus nq is the
order of the unitary matrix Q or P**H that is applied.
If VECT = MagmaQ, A is assumed to have been an NQ-by-K matrix:
if nq >= k, Q = H(1) H(2) . . . H(k);
if nq < k, Q = H(1) H(2) . . . H(nq-1).
If VECT = MagmaP, A is assumed to have been A K-by-NQ matrix:
if k < nq, P = G(1) G(2) . . . G(k);
if k >= nq, P = G(1) G(2) . . . G(nq-1).
Arguments
---------
@param[in]
vect magma_vect_t
- = MagmaQ: apply Q or Q**H;
- = MagmaP: apply P or P**H.
@param[in]
side magma_side_t
- = MagmaLeft: apply Q, Q**H, P or P**H from the Left;
- = MagmaRight: apply Q, Q**H, P or P**H from the Right.
@param[in]
trans magma_trans_t
- = MagmaNoTrans: No transpose, apply Q or P;
- = Magma_ConjTrans: Conjugate transpose, apply Q**H or P**H.
@param[in]
m INTEGER
The number of rows of the matrix C. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix C. N >= 0.
@param[in]
k INTEGER
If VECT = MagmaQ, the number of columns in the original
matrix reduced by CGEBRD.
If VECT = MagmaP, the number of rows in the original
matrix reduced by CGEBRD.
K >= 0.
@param[in]
A COMPLEX array, dimension
(LDA,min(nq,K)) if VECT = MagmaQ
(LDA,nq) if VECT = MagmaP
The vectors which define the elementary reflectors H(i) and
G(i), whose products determine the matrices Q and P, as
returned by CGEBRD.
@param[in]
lda INTEGER
The leading dimension of the array A.
If VECT = MagmaQ, LDA >= max(1,nq);
if VECT = MagmaP, LDA >= max(1,min(nq,K)).
@param[in]
tau COMPLEX array, dimension (min(nq,K))
TAU(i) must contain the scalar factor of the elementary
reflector H(i) or G(i) which determines Q or P, as returned
by CGEBRD in the array argument TAUQ or TAUP.
@param[in,out]
C COMPLEX array, dimension (LDC,N)
On entry, the M-by-N matrix C.
On exit, C is overwritten by Q*C or Q**H*C or C*Q**H or C*Q
or P*C or P**H*C or C*P or C*P**H.
@param[in]
ldc INTEGER
The leading dimension of the array C. LDC >= max(1,M).
@param[out]
work (workspace) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK.
If SIDE = MagmaLeft, LWORK >= max(1,N);
if SIDE = MagmaRight, LWORK >= max(1,M);
if N = 0 or M = 0, LWORK >= 1.
For optimum performance
if SIDE = MagmaLeft, LWORK >= max(1,N*NB);
if SIDE = MagmaRight, LWORK >= max(1,M*NB),
where NB is the optimal blocksize. (NB = 0 if M = 0 or N = 0.)
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_cgesvd_comp
********************************************************************/
extern "C" magma_int_t
magma_cunmbr(
magma_vect_t vect, magma_side_t side, magma_trans_t trans,
magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloatComplex *A, magma_int_t lda,
magmaFloatComplex *tau,
magmaFloatComplex *C, magma_int_t ldc,
magmaFloatComplex *work, magma_int_t lwork,
magma_int_t *info)
{
#define A(i,j) (A + (i) + (j)*lda)
#define C(i,j) (C + (i) + (j)*ldc)
magma_int_t i1, i2, nb, mi, ni, nq, nq_1, nw, iinfo, lwkopt;
magma_int_t left, notran, applyq, lquery;
magma_trans_t transt;
*info = 0;
applyq = (vect == MagmaQ);
left = (side == MagmaLeft);
notran = (trans == MagmaNoTrans);
lquery = (lwork == -1);
/* NQ is the order of Q or P and NW is the minimum dimension of WORK */
if (left) {
nq = m;
nw = n;
}
else {
nq = n;
nw = m;
}
if (m == 0 || n == 0) {
nw = 0;
}
/* check arguments */
if (! applyq && vect != MagmaP) {
*info = -1;
}
else if (! left && side != MagmaRight) {
*info = -2;
}
else if (! notran && trans != Magma_ConjTrans) {
*info = -3;
}
else if (m < 0) {
*info = -4;
}
else if (n < 0) {
*info = -5;
}
else if (k < 0) {
*info = -6;
}
else if ( ( applyq && lda < max(1,nq) ) ||
( ! applyq && lda < max(1,min(nq,k)) ) ) {
*info = -8;
}
else if (ldc < max(1,m)) {
*info = -11;
}
else if (lwork < max(1,nw) && ! lquery) {
*info = -13;
}
if (*info == 0) {
if (nw > 0) {
// TODO have get_cunmqr_nb and get_cunmlq_nb routines? see original LAPACK cunmbr.
// TODO make them dependent on m, n, and k?
nb = magma_get_cgebrd_nb( min( m, n ));
lwkopt = max(1, nw*nb);
}
else {
lwkopt = 1;
}
work[0] = MAGMA_C_MAKE( lwkopt, 0 );
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0) {
return *info;
}
if (applyq) {
/* Apply Q */
if (nq >= k) {
/* Q was determined by a call to CGEBRD with nq >= k */
#if VERSION == 1
lapackf77_cunmqr( lapack_side_const(side), lapack_trans_const(trans),
&m, &n, &k, A, &lda, tau, C, &ldc, work, &lwork, &iinfo);
#else
magma_cunmqr( side, trans,
m, n, k, A, lda, tau, C, ldc, work, lwork, &iinfo);
#endif
}
else if (nq > 1) {
/* Q was determined by a call to CGEBRD with nq < k */
if (left) {
mi = m - 1;
ni = n;
i1 = 1;
i2 = 0;
}
else {
mi = m;
ni = n - 1;
i1 = 0;
i2 = 1;
}
nq_1 = nq - 1;
#if VERSION == 1
lapackf77_cunmqr( lapack_side_const(side), lapack_trans_const(trans),
&mi, &ni, &nq_1, A(1,0), &lda, tau, C(i1,i2), &ldc, work, &lwork, &iinfo);
#else
magma_cunmqr( side, trans,
mi, ni, nq-1, A(1,0), lda, tau, C(i1,i2), ldc, work, lwork, &iinfo);
#endif
}
}
else {
/* Apply P */
if (notran) {
transt = Magma_ConjTrans;
}
else {
transt = MagmaNoTrans;
}
if (nq > k) {
/* P was determined by a call to CGEBRD with nq > k */
#if VERSION == 1
lapackf77_cunmlq( lapack_side_const(side), lapack_trans_const(transt),
&m, &n, &k, A, &lda, tau, C, &ldc, work, &lwork, &iinfo);
#else
magma_cunmlq( side, transt,
m, n, k, A, lda, tau, C, ldc, work, lwork, &iinfo);
#endif
}
else if (nq > 1) {
/* P was determined by a call to CGEBRD with nq <= k */
if (left) {
mi = m - 1;
ni = n;
i1 = 1;
i2 = 0;
}
else {
mi = m;
ni = n - 1;
i1 = 0;
i2 = 1;
}
nq_1 = nq - 1;
#if VERSION == 1
lapackf77_cunmlq( lapack_side_const(side), lapack_trans_const(transt),
&mi, &ni, &nq_1, A(0,1), &lda, tau, C(i1,i2), &ldc, work, &lwork, &iinfo);
#else
magma_cunmlq( side, transt,
mi, ni, nq-1, A(0,1), lda, tau, C(i1,i2), ldc, work, lwork, &iinfo);
#endif
}
}
work[0] = MAGMA_C_MAKE( lwkopt, 0 );
return *info;
} /* magma_cunmbr */
/**
Purpose
-------
CHEGVD computes all the eigenvalues, and optionally, the eigenvectors
of a complex generalized Hermitian-definite eigenproblem, of the form
A*x=(lambda)*B*x, A*Bx=(lambda)*x, or B*A*x=(lambda)*x. Here A and
B are assumed to be Hermitian and B is also positive definite.
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]
nrgpu INTEGER
Number of GPUs to use.
@param[in]
itype INTEGER
Specifies the problem type to be solved:
= 1: A*x = (lambda)*B*x
= 2: A*B*x = (lambda)*x
= 3: B*A*x = (lambda)*x
@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 triangles of A and B are stored;
- = MagmaLower: Lower triangles of A and B are stored.
@param[in]
n INTEGER
The order of the matrices A and B. N >= 0.
@param[in,out]
A COMPLEX array, dimension (LDA, N)
On entry, the Hermitian matrix A. If UPLO = MagmaUpper, the
leading N-by-N upper triangular part of A contains the
upper triangular part of the matrix A. If UPLO = MagmaLower,
the leading N-by-N lower triangular part of A contains
the lower triangular part of the matrix A.
\n
On exit, if JOBZ = MagmaVec, then if INFO = 0, A contains the
matrix Z of eigenvectors. The eigenvectors are normalized
as follows:
if ITYPE = 1 or 2, Z**H*B*Z = I;
if ITYPE = 3, Z**H*inv(B)*Z = I.
If JOBZ = MagmaNoVec, then on exit the upper triangle (if UPLO=MagmaUpper)
or the lower triangle (if UPLO=MagmaLower) of A, including the
diagonal, is destroyed.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[in,out]
B COMPLEX array, dimension (LDB, N)
On entry, the Hermitian matrix B. If UPLO = MagmaUpper, the
leading N-by-N upper triangular part of B contains the
upper triangular part of the matrix B. If UPLO = MagmaLower,
the leading N-by-N lower triangular part of B contains
the lower triangular part of the matrix B.
\n
On exit, if INFO <= N, the part of B containing the matrix is
overwritten by the triangular factor U or L from the Cholesky
factorization B = U**H*U or B = L*L**H.
@param[in]
ldb INTEGER
The leading dimension of the array B. LDB >= 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 eigenvalues in ascending order.
@param[out]
work (workspace) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
@param[in]
lwork INTEGER
The length of the array WORK.
If N <= 1, LWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LWORK >= N + 1.
If JOBZ = MagmaVec and N > 1, LWORK >= 2*N*nb + N**2.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal sizes of the WORK, RWORK and
IWORK arrays, returns these values as the first entries of
the WORK, RWORK and IWORK arrays, and no error message
related to LWORK or LRWORK or LIWORK is issued by XERBLA.
@param[out]
rwork (workspace) REAL array, dimension (MAX(1,LRWORK))
On exit, if INFO = 0, RWORK(1) returns the optimal LRWORK.
@param[in]
lrwork INTEGER
The dimension of the array RWORK.
If N <= 1, LRWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LRWORK >= N.
If JOBZ = MagmaVec and N > 1, LRWORK >= 1 + 5*N + 2*N**2.
\n
If LRWORK = -1, then a workspace query is assumed; the
routine only calculates the optimal sizes of the WORK, RWORK
and IWORK arrays, returns these values as the first entries
of the WORK, RWORK and IWORK arrays, and no error message
related to LWORK or LRWORK or LIWORK is issued by XERBLA.
@param[out]
iwork (workspace) INTEGER array, dimension (MAX(1,LIWORK))
On exit, if INFO = 0, IWORK(1) returns the optimal LIWORK.
@param[in]
liwork INTEGER
The dimension of the array IWORK.
If N <= 1, LIWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LIWORK >= 1.
If JOBZ = MagmaVec and N > 1, LIWORK >= 3 + 5*N.
\n
If LIWORK = -1, then a workspace query is assumed; the
routine only calculates the optimal sizes of the WORK, RWORK
and IWORK arrays, returns these values as the first entries
of the WORK, RWORK and IWORK arrays, and no error message
related to LWORK or LRWORK or LIWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: CPOTRF or CHEEVD returned an error code:
<= N: 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);
> N: if INFO = N + i, for 1 <= i <= N, then the leading
minor of order i of B is not positive definite.
The factorization of B could not be completed and
no eigenvalues or eigenvectors were computed.
Further Details
---------------
Based on contributions by
Mark Fahey, Department of Mathematics, Univ. of Kentucky, USA
Modified so that no backsubstitution is performed if CHEEVD fails to
converge (NEIG in old code could be greater than N causing out of
bounds reference to A - reported by Ralf Meyer). Also corrected the
description of INFO and the test on ITYPE. Sven, 16 Feb 05.
@ingroup magma_chegv_driver
********************************************************************/
extern "C" magma_int_t
magma_chegvdx_m(magma_int_t nrgpu, magma_int_t itype, magma_vec_t jobz, magma_range_t range, magma_uplo_t uplo, magma_int_t n,
magmaFloatComplex *A, magma_int_t lda, magmaFloatComplex *B, magma_int_t ldb,
float vl, float vu, magma_int_t il, magma_int_t iu,
magma_int_t *m, float *w, magmaFloatComplex *work, magma_int_t lwork,
float *rwork, magma_int_t lrwork,
magma_int_t *iwork, magma_int_t liwork, magma_int_t *info)
{
const char* uplo_ = lapack_uplo_const( uplo );
const char* jobz_ = lapack_vec_const( jobz );
magmaFloatComplex c_one = MAGMA_C_ONE;
magma_int_t lower;
magma_trans_t trans;
magma_int_t wantz;
magma_int_t lquery;
magma_int_t alleig, valeig, indeig;
magma_int_t lwmin;
magma_int_t liwmin;
magma_int_t lrwmin;
wantz = (jobz == MagmaVec);
lower = (uplo == MagmaLower);
alleig = (range == MagmaRangeAll);
valeig = (range == MagmaRangeV);
indeig = (range == MagmaRangeI);
lquery = (lwork == -1 || lrwork == -1 || liwork == -1);
*info = 0;
if (itype < 1 || itype > 3) {
*info = -1;
} else if (! (alleig || valeig || indeig)) {
*info = -2;
} else if (! (wantz || (jobz == MagmaNoVec))) {
*info = -3;
} else if (! (lower || (uplo == MagmaUpper))) {
*info = -4;
} else if (n < 0) {
*info = -5;
} else if (lda < max(1,n)) {
*info = -7;
} else if (ldb < max(1,n)) {
*info = -9;
} else {
if (valeig) {
if (n > 0 && vu <= vl) {
*info = -11;
}
} else if (indeig) {
if (il < 1 || il > max(1,n)) {
*info = -12;
} else if (iu < min(n,il) || iu > n) {
*info = -13;
}
}
}
magma_int_t nb = magma_get_chetrd_nb( n );
if ( n <= 1 ) {
lwmin = 1;
lrwmin = 1;
liwmin = 1;
}
else if ( wantz ) {
lwmin = 2*n + n*n;
lrwmin = 1 + 5*n + 2*n*n;
liwmin = 3 + 5*n;
}
else {
lwmin = n + n*nb;
lrwmin = n;
liwmin = 1;
}
// multiply by 1+eps (in Double!) to ensure length gets rounded up,
// if it cannot be exactly represented in floating point.
real_Double_t one_eps = 1. + lapackf77_slamch("Epsilon");
work[0] = MAGMA_C_MAKE( lwmin * one_eps, 0.); // round up
rwork[0] = lrwmin * one_eps;
iwork[0] = liwmin;
if (lwork < lwmin && ! lquery) {
*info = -17;
} else if (lrwork < lrwmin && ! lquery) {
*info = -19;
} else if (liwork < liwmin && ! lquery) {
*info = -21;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info));
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (n == 0) {
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
lapackf77_chegvd(&itype, jobz_, uplo_,
&n, A, &lda, B, &ldb,
w, work, &lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, &lrwork,
#endif
iwork, &liwork, info);
*m = n;
return *info;
}
magma_timer_t time=0;
timer_start( time );
magma_cpotrf_m(nrgpu, uplo, n, B, ldb, info);
if (*info != 0) {
*info = n + *info;
return *info;
}
timer_stop( time );
timer_printf("time cpotrf = %6.2f\n", time );
timer_start( time );
/* Transform problem to standard eigenvalue problem and solve. */
magma_chegst_m(nrgpu, itype, uplo, n, A, lda, B, ldb, info);
timer_stop( time );
timer_printf( "time chegst = %6.2f\n", time );
timer_start( time );
magma_cheevdx_m(nrgpu, jobz, range, uplo, n, A, lda, vl, vu, il, iu, m, w, work, lwork, rwork, lrwork, iwork, liwork, info);
timer_stop( time );
timer_printf( "time cheevd = %6.2f\n", time );
if (wantz && *info == 0) {
timer_start( time );
/* Backtransform eigenvectors to the original problem. */
if (itype == 1 || itype == 2) {
/* For A*x=(lambda)*B*x and A*B*x=(lambda)*x;
backtransform eigenvectors: x = inv(L)'*y or inv(U)*y */
if (lower) {
trans = MagmaConjTrans;
} else {
trans = MagmaNoTrans;
}
magma_ctrsm_m(nrgpu, MagmaLeft, uplo, trans, MagmaNonUnit,
n, *m, c_one, B, ldb, A, lda);
}
else if (itype == 3) {
/* For B*A*x=(lambda)*x;
backtransform eigenvectors: x = L*y or U'*y */
if (lower) {
trans = MagmaNoTrans;
} else {
trans = MagmaConjTrans;
}
//magma_ctrmm(MagmaLeft, uplo, trans, MagmaNonUnit,
// n, n, c_one, db, lddb, da, ldda);
}
timer_stop( time );
timer_printf( "time setmatrices trsm/mm + getmatrices = %6.2f\n", time );
}
work[0] = MAGMA_C_MAKE( lwmin * one_eps, 0.); // round up
rwork[0] = lrwmin * one_eps;
iwork[0] = liwmin;
return *info;
} /* magma_chegvd_m */
int main(int argc, char **argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, cublas_perf, cublas_time, cpu_perf, cpu_time;
float magma_error, cublas_error, work[1];
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t M, N, Xm, Ym, lda, sizeA, sizeX, sizeY;
magma_int_t incx = 1;
magma_int_t incy = 1;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex alpha = MAGMA_C_MAKE( 1.5, -2.3 );
magmaFloatComplex beta = MAGMA_C_MAKE( -0.6, 0.8 );
magmaFloatComplex *A, *X, *Y, *Ycublas, *Ymagma;
magmaFloatComplex *dA, *dX, *dY;
magma_opts opts;
parse_opts( argc, argv, &opts );
printf(" M N MAGMA Gflop/s (ms) CUBLAS Gflop/s (ms) CPU Gflop/s (ms) MAGMA error CUBLAS error\n");
printf("===================================================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[i];
N = opts.nsize[i];
lda = ((M+31)/32)*32;
gflops = FLOPS_CGEMV( M, N ) / 1e9;
if ( opts.transA == MagmaNoTrans ) {
Xm = N;
Ym = M;
} else {
Xm = M;
Ym = N;
}
sizeA = lda*N;
sizeX = incx*Xm;
sizeY = incy*Ym;
TESTING_MALLOC_CPU( A, magmaFloatComplex, sizeA );
TESTING_MALLOC_CPU( X, magmaFloatComplex, sizeX );
TESTING_MALLOC_CPU( Y, magmaFloatComplex, sizeY );
TESTING_MALLOC_CPU( Ycublas, magmaFloatComplex, sizeY );
TESTING_MALLOC_CPU( Ymagma, magmaFloatComplex, sizeY );
TESTING_MALLOC_DEV( dA, magmaFloatComplex, sizeA );
TESTING_MALLOC_DEV( dX, magmaFloatComplex, sizeX );
TESTING_MALLOC_DEV( dY, magmaFloatComplex, sizeY );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &sizeA, A );
lapackf77_clarnv( &ione, ISEED, &sizeX, X );
lapackf77_clarnv( &ione, ISEED, &sizeY, Y );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_csetmatrix( M, N, A, lda, dA, lda );
magma_csetvector( Xm, X, incx, dX, incx );
magma_csetvector( Ym, Y, incy, dY, incy );
cublas_time = magma_sync_wtime( 0 );
cublasCgemv( opts.transA, M, N, alpha, dA, lda, dX, incx, beta, dY, incy );
cublas_time = magma_sync_wtime( 0 ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_cgetvector( Ym, dY, incy, Ycublas, incy );
/* =====================================================================
Performs operation using MAGMABLAS
=================================================================== */
magma_csetvector( Ym, Y, incy, dY, incy );
magma_time = magma_sync_wtime( 0 );
magmablas_cgemv( opts.transA, M, N, alpha, dA, lda, dX, incx, beta, dY, incy );
magma_time = magma_sync_wtime( 0 ) - magma_time;
magma_perf = gflops / magma_time;
magma_cgetvector( Ym, dY, incx, Ymagma, incx );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
cpu_time = magma_wtime();
blasf77_cgemv( &opts.transA, &M, &N,
&alpha, A, &lda,
X, &incx,
&beta, Y, &incy );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/* =====================================================================
Check the result
=================================================================== */
blasf77_caxpy( &Ym, &c_neg_one, Y, &incy, Ymagma, &incy );
magma_error = lapackf77_clange( "M", &Ym, &ione, Ymagma, &Ym, work ) / Ym;
blasf77_caxpy( &Ym, &c_neg_one, Y, &incy, Ycublas, &incy );
cublas_error = lapackf77_clange( "M", &Ym, &ione, Ycublas, &Ym, work ) / Ym;
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e\n",
(int) M, (int) N,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
magma_error, cublas_error );
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( X );
TESTING_FREE_CPU( Y );
TESTING_FREE_CPU( Ycublas );
TESTING_FREE_CPU( Ymagma );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dX );
TESTING_FREE_DEV( dY );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cher2k
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, cublas_perf, cublas_time, cpu_perf, cpu_time;
float cublas_error, Cnorm, work[1];
magma_int_t N, K;
magma_int_t Ak, An, Bk, Bn;
magma_int_t sizeA, sizeB, sizeC;
magma_int_t lda, ldb, ldc, ldda, lddb, lddc;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magmaFloatComplex *h_A, *h_B, *h_C, *h_Ccublas;
magmaFloatComplex *d_A, *d_B, *d_C;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex alpha = MAGMA_C_MAKE( 0.29, -0.86 );
float beta = MAGMA_D_MAKE( -0.48, 0.38 );
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
float tol = opts.tolerance * lapackf77_slamch("E");
printf("If running lapack (option --lapack), CUBLAS error is computed\n"
"relative to CPU BLAS result.\n\n");
printf("uplo = %s, transA = %s\n",
lapack_uplo_const(opts.uplo), lapack_trans_const(opts.transA) );
printf(" N K CUBLAS Gflop/s (ms) CPU Gflop/s (ms) CUBLAS error\n");
printf("==================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.msize[itest];
K = opts.ksize[itest];
gflops = FLOPS_CHER2K(K, N) / 1e9;
if ( opts.transA == MagmaNoTrans ) {
lda = An = N;
Ak = K;
ldb = Bn = N;
Bk = K;
} else {
lda = An = K;
Ak = N;
ldb = Bn = K;
Bk = N;
}
ldc = N;
ldda = ((lda+31)/32)*32;
lddb = ((ldb+31)/32)*32;
lddc = ((ldc+31)/32)*32;
sizeA = lda*Ak;
sizeB = ldb*Ak;
sizeC = ldc*N;
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, lda*Ak );
TESTING_MALLOC_CPU( h_B, magmaFloatComplex, ldb*Bk );
TESTING_MALLOC_CPU( h_C, magmaFloatComplex, ldc*N );
TESTING_MALLOC_CPU( h_Ccublas, magmaFloatComplex, ldc*N );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*Ak );
TESTING_MALLOC_DEV( d_B, magmaFloatComplex, lddb*Bk );
TESTING_MALLOC_DEV( d_C, magmaFloatComplex, lddc*N );
/* Initialize the matrices */
lapackf77_clarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_clarnv( &ione, ISEED, &sizeB, h_B );
lapackf77_clarnv( &ione, ISEED, &sizeC, h_C );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_csetmatrix( An, Ak, h_A, lda, d_A, ldda );
magma_csetmatrix( Bn, Bk, h_B, ldb, d_B, lddb );
magma_csetmatrix( N, N, h_C, ldc, d_C, lddc );
cublas_time = magma_sync_wtime( NULL );
cublasCher2k( handle, cublas_uplo_const(opts.uplo), cublas_trans_const(opts.transA), N, K,
&alpha, d_A, ldda,
d_B, lddb,
&beta, d_C, lddc );
cublas_time = magma_sync_wtime( NULL ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_cgetmatrix( N, N, d_C, lddc, h_Ccublas, ldc );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
blasf77_cher2k( lapack_uplo_const(opts.uplo), lapack_trans_const(opts.transA), &N, &K,
&alpha, h_A, &lda,
h_B, &ldb,
&beta, h_C, &ldc );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
}
/* =====================================================================
Check the result
=================================================================== */
if ( opts.lapack ) {
// compute relative error for both magma & cublas, relative to lapack,
// |C_magma - C_lapack| / |C_lapack|
Cnorm = lapackf77_clange( "M", &N, &N, h_C, &ldc, work );
blasf77_caxpy( &sizeC, &c_neg_one, h_C, &ione, h_Ccublas, &ione );
cublas_error = lapackf77_clange( "M", &N, &N, h_Ccublas, &ldc, work ) / Cnorm;
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) N, (int) K,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
cublas_error, (cublas_error < tol ? "ok" : "failed"));
status += ! (cublas_error < tol);
}
else {
printf("%5d %5d %7.2f (%7.2f) --- ( --- ) --- ---\n",
(int) N, (int) K,
cublas_perf, 1000.*cublas_time);
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_C );
TESTING_FREE_CPU( h_Ccublas );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
TESTING_FREE_DEV( d_C );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/**
Purpose
-------
CGEGQR orthogonalizes the N vectors given by a complex M-by-N matrix A:
A = Q * R.
On exit, if successful, the orthogonal vectors Q overwrite A
and R is given in work (on the CPU memory).
The routine is designed for tall-and-skinny matrices: M >> N, N <= 128.
This version uses normal equations and SVD in an iterative process that
makes the computation numerically accurate.
Arguments
---------
@param[in]
ikind INTEGER
Several versions are implemented indiceted by the ikind value:
1: This version uses normal equations and SVD in an iterative process
that makes the computation numerically accurate.
2: This version uses a standard LAPACK-based orthogonalization through
MAGMA's QR panel factorization (magma_cgeqr2x3_gpu) and magma_cungqr
3: MGS
4. Cholesky QR
@param[in]
m INTEGER
The number of rows of the matrix A. m >= n >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. 128 >= n >= 0.
@param[in,out]
dA COMPLEX array on the GPU, dimension (ldda,n)
On entry, the m-by-n matrix A.
On exit, the m-by-n matrix Q with orthogonal columns.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,m).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
@param
dwork (GPU workspace) COMPLEX array, dimension:
n^2 for ikind = 1
3 n^2 + min(m, n) for ikind = 2
0 (not used) for ikind = 3
n^2 for ikind = 4
@param[out]
work (CPU workspace) COMPLEX array, dimension 3 n^2.
On exit, work(1:n^2) holds the rectangular matrix R.
Preferably, for higher performance, work should be in pinned memory.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
@ingroup magma_cgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_cgegqr_gpu( magma_int_t ikind, magma_int_t m, magma_int_t n,
magmaFloatComplex *dA, magma_int_t ldda,
magmaFloatComplex *dwork, magmaFloatComplex *work,
magma_int_t *info )
{
#define work(i_,j_) (work + (i_) + (j_)*n)
#define dA(i_,j_) (dA + (i_) + (j_)*ldda)
magma_int_t i = 0, j, k, n2 = n*n;
magma_int_t ione = 1;
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
float cn = 200., mins, maxs;
/* check arguments */
*info = 0;
if (ikind < 1 || ikind > 4) {
*info = -1;
} else if (m < 0 || m < n) {
*info = -2;
} else if (n < 0 || n > 128) {
*info = -3;
} else if (ldda < max(1,m)) {
*info = -5;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (ikind == 1) {
// === Iterative, based on SVD ============================================================
magmaFloatComplex *U, *VT, *vt, *R, *G, *hwork, *tau;
float *S;
R = work; // Size n * n
G = R + n*n; // Size n * n
VT = G + n*n; // Size n * n
magma_cmalloc_cpu( &hwork, 32 + 2*n*n + 2*n);
if ( hwork == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_int_t lwork=n*n+32; // First part f hwork; used as workspace in svd
U = hwork + n*n + 32; // Size n*n
S = (float *)(U+n*n); // Size n
tau = U + n*n + n; // Size n
#if defined(PRECISION_c) || defined(PRECISION_z)
float *rwork;
magma_smalloc_cpu( &rwork, 5*n);
if ( rwork == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
#endif
do {
i++;
magma_cgemm(MagmaConjTrans, MagmaNoTrans, n, n, m, c_one, dA, ldda, dA, ldda, c_zero, dwork, n );
magma_cgetmatrix(n, n, dwork, n, G, n);
#if defined(PRECISION_s) || defined(PRECISION_d)
lapackf77_cgesvd("n", "a", &n, &n, G, &n, S, U, &n, VT, &n,
hwork, &lwork, info);
#else
lapackf77_cgesvd("n", "a", &n, &n, G, &n, S, U, &n, VT, &n,
hwork, &lwork, rwork, info);
#endif
mins = 100.f, maxs = 0.f;
for (k=0; k < n; k++) {
S[k] = magma_ssqrt( S[k] );
if (S[k] < mins) mins = S[k];
if (S[k] > maxs) maxs = S[k];
}
for (k=0; k < n; k++) {
vt = VT + k*n;
for (j=0; j < n; j++)
vt[j] *= S[j];
}
lapackf77_cgeqrf(&n, &n, VT, &n, tau, hwork, &lwork, info);
if (i == 1)
blasf77_ccopy(&n2, VT, &ione, R, &ione);
else
blasf77_ctrmm("l", "u", "n", "n", &n, &n, &c_one, VT, &n, R, &n);
magma_csetmatrix(n, n, VT, n, dwork, n);
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit, m, n, c_one, dwork, n, dA, ldda);
if (mins > 0.00001f)
cn = maxs/mins;
//fprintf(stderr, "Iteration %d, cond num = %f \n", i, cn);
} while (cn > 10.f);
magma_free_cpu( hwork );
#if defined(PRECISION_c) || defined(PRECISION_z)
magma_free_cpu( rwork );
#endif
// ================== end of ikind == 1 ===================================================
}
else if (ikind == 2) {
// ================== LAPACK based ===================================================
magma_int_t min_mn = min(m, n);
magma_int_t nb = n;
magmaFloatComplex *dtau = dwork + 2*n*n, *d_T = dwork, *ddA = dwork + n*n;
magmaFloatComplex *tau = work+n*n;
magmablas_claset( MagmaFull, n, n, c_zero, c_zero, d_T, n );
magma_cgeqr2x3_gpu(&m, &n, dA, &ldda, dtau, d_T, ddA,
(float *)(dwork+min_mn+2*n*n), info);
magma_cgetmatrix( min_mn, 1, dtau, min_mn, tau, min_mn);
magma_cgetmatrix( n, n, ddA, n, work, n);
magma_cungqr_gpu( m, n, n, dA, ldda, tau, d_T, nb, info );
// ================== end of ikind == 2 ===================================================
}
else if (ikind == 3) {
// ================== MGS ===================================================
for(magma_int_t j = 0; j<n; j++){
for(magma_int_t i = 0; i<j; i++){
*work(i, j) = magma_cdotc(m, dA(0,i), 1, dA(0,j), 1);
magma_caxpy(m, -(*work(i,j)), dA(0,i), 1, dA(0,j), 1);
}
for(magma_int_t i = j; i<n; i++)
*work(i, j) = MAGMA_C_ZERO;
//*work(j,j) = MAGMA_C_MAKE( magma_scnrm2(m, dA(0,j), 1), 0. );
*work(j,j) = magma_cdotc(m, dA(0,j), 1, dA(0,j), 1);
*work(j,j) = MAGMA_C_MAKE( sqrt(MAGMA_C_REAL( *work(j,j) )), 0.);
magma_cscal(m, 1./ *work(j,j), dA(0,j), 1);
}
// ================== end of ikind == 3 ===================================================
}
else if (ikind == 4) {
// ================== Cholesky QR ===================================================
magma_cgemm(MagmaConjTrans, MagmaNoTrans, n, n, m, c_one, dA, ldda, dA, ldda, c_zero, dwork, n );
magma_cgetmatrix(n, n, dwork, n, work, n);
lapackf77_cpotrf("u", &n, work, &n, info);
magma_csetmatrix(n, n, work, n, dwork, n);
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit, m, n, c_one, dwork, n, dA, ldda);
// ================== end of ikind == 4 ===================================================
}
return *info;
} /* magma_cgegqr_gpu */
magma_int_t
magma_cm_27stencil(
magma_int_t n,
magma_c_matrix *A,
magma_queue_t queue )
{
magma_int_t info = 0;
magma_int_t i,j,k;
magma_c_matrix hA={Magma_CSR};
// generate matrix of desired structure and size (3d 27-point stencil)
magma_int_t nn = n*n*n;
magma_int_t offdiags = 13;
magma_index_t *diag_offset=NULL;
magmaFloatComplex *diag_vals=NULL;
CHECK( magma_cmalloc_cpu( &diag_vals, offdiags+1 ));
CHECK( magma_index_malloc_cpu( &diag_offset, offdiags+1 ));
diag_offset[0] = 0;
diag_offset[1] = 1;
diag_offset[2] = n-1;
diag_offset[3] = n;
diag_offset[4] = n+1;
diag_offset[5] = n*n-n-1;
diag_offset[6] = n*n-n;
diag_offset[7] = n*n-n+1;
diag_offset[8] = n*n-1;
diag_offset[9] = n*n;
diag_offset[10] = n*n+1;
diag_offset[11] = n*n+n-1;
diag_offset[12] = n*n+n;
diag_offset[13] = n*n+n+1;
diag_vals[0] = MAGMA_C_MAKE( 26.0, 0.0 );
diag_vals[1] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[2] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[3] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[4] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[5] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[6] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[7] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[8] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[9] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[10] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[11] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[12] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[13] = MAGMA_C_MAKE( -1.0, 0.0 );
CHECK( magma_cmgenerator( nn, offdiags, diag_offset, diag_vals, &hA, queue ));
// now set some entries to zero (boundary...)
for( i=0; i < n*n; i++ ) {
for( j=0; j < n; j++ ) {
magma_index_t row = i*n+j;
for( k=hA.row[row]; k<hA.row[row+1]; k++) {
if ((hA.col[k] == row-1 ||
hA.col[k] == row-n-1 ||
hA.col[k] == row+n-1 ||
hA.col[k] == row-n*n+n-1 ||
hA.col[k] == row+n*n-n-1 ||
hA.col[k] == row-n*n-1 ||
hA.col[k] == row+n*n-1 ||
hA.col[k] == row-n*n-n-1 ||
hA.col[k] == row+n*n+n-1 ) && (row+1)%n == 1 )
hA.val[k] = MAGMA_C_MAKE( 0.0, 0.0 );
if ((hA.col[k] == row+1 ||
hA.col[k] == row-n+1 ||
hA.col[k] == row+n+1 ||
hA.col[k] == row-n*n+n+1 ||
hA.col[k] == row+n*n-n+1 ||
hA.col[k] == row-n*n+1 ||
hA.col[k] == row+n*n+1 ||
hA.col[k] == row-n*n-n+1 ||
hA.col[k] == row+n*n+n+1 ) && (row)%n ==n-1 )
hA.val[k] = MAGMA_C_MAKE( 0.0, 0.0 );
}
}
}
hA.true_nnz = hA.nnz;
CHECK( magma_cmconvert( hA, A, Magma_CSR, Magma_CSR, queue ));
cleanup:
magma_free_cpu( diag_vals );
magma_free_cpu( diag_offset );
magma_cmfree( &hA, queue );
return info;
}
int main(int argc, char **argv)
{
TESTING_INIT();
magma_setdevice(0);
magma_timestr_t start, end;
float flops, magma_perf, cuda_perf, error, work[1];
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t n_local[4];
FILE *fp ;
magma_int_t N, m, i, j, lda, LDA, M;
magma_int_t matsize;
magma_int_t vecsize;
magma_int_t istart = 64;
magma_int_t incx = 1;
char uplo = MagmaLower;
magmaFloatComplex alpha = MAGMA_C_MAKE(1., 0.); // MAGMA_C_MAKE( 1.5, -2.3 );
magmaFloatComplex beta = MAGMA_C_MAKE(0., 0.); // MAGMA_C_MAKE( -0.6, 0.8 );
magmaFloatComplex *A, *X, *Y[4], *Ycublas, *Ymagma;
magmaFloatComplex *dA, *dX[4], *dY[4], *d_lA[4], *dYcublas ;
magma_queue_t stream[4][10];
magmaFloatComplex *C_work;
magmaFloatComplex *dC_work[4];
int max_num_gpus;
magma_int_t num_gpus = 1, nb;
magma_int_t blocks, lwork;
magma_int_t offset = 0;
M = 0;
N = 0;
if (argc != 1){
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0)
{
N = atoi(argv[++i]);
istart = N;
}
else if (strcmp("-M", argv[i])==0)
M = atoi(argv[++i]);
else if (strcmp("-NGPU", argv[i])==0)
num_gpus = atoi(argv[++i]);
else if (strcmp("-offset", argv[i])==0)
offset = atoi(argv[++i]);
}
if ( M == 0 ) {
M = N;
}
if ( N == 0 ) {
N = M;
}
if (M>0 && N>0)
{ printf(" testing_chemv_mgpu -M %d -N %d -NGPU %d\n\n", (int) M, (int) N, (int) num_gpus);
printf(" in %c side \n", uplo);
}
else
{
printf("\nUsage: \n");
printf(" testing_chemv_mgpu -M %d -N %d -NGPU %d\n\n",
1024, 1024, 1);
exit(1);
}
}
else {
#if defined(PRECISION_z)
M = N = 8000;
#else
M = N = 12480;
#endif
num_gpus = 2;
offset = 0;
printf("\nUsage: \n");
printf(" testing_chemv_mgpu -M %d -N %d -NGPU %d\n\n", (int) M, (int) N, (int) num_gpus);
}
//////////////////////////////////////////////////////////////////////////
cudaGetDeviceCount(&max_num_gpus);
if (num_gpus > max_num_gpus){
printf("More GPUs requested than available. Have to change it.\n");
num_gpus = max_num_gpus;
}
printf("Number of GPUs to be used = %d\n", (int) num_gpus);
for(int i=0; i< num_gpus; i++)
{
magma_queue_create(&stream[i][0]);
}
LDA = ((N+31)/32)*32;
matsize = N*LDA;
vecsize = N*incx;
nb = 32;
//nb = 64;
printf("block size = %d\n", (int) nb);
TESTING_MALLOC_CPU( A, magmaFloatComplex, matsize );
TESTING_MALLOC_CPU( X, magmaFloatComplex, vecsize );
TESTING_MALLOC_CPU( Ycublas, magmaFloatComplex, vecsize );
TESTING_MALLOC_CPU( Ymagma, magmaFloatComplex, vecsize );
for(i=0; i<num_gpus; i++)
{
TESTING_MALLOC_CPU( Y[i], magmaFloatComplex, vecsize );
}
magma_setdevice(0);
TESTING_MALLOC_DEV( dA, magmaFloatComplex, matsize );
TESTING_MALLOC_DEV( dYcublas, magmaFloatComplex, vecsize );
for(i=0; i<num_gpus; i++)
{
n_local[i] = ((N/nb)/num_gpus)*nb;
if (i < (N/nb)%num_gpus)
n_local[i] += nb;
else if (i == (N/nb)%num_gpus)
n_local[i] += N%nb;
magma_setdevice(i);
TESTING_MALLOC_DEV( d_lA[i], magmaFloatComplex, LDA*n_local[i] );// potentially bugged
TESTING_MALLOC_DEV( dX[i], magmaFloatComplex, vecsize );
TESTING_MALLOC_DEV( dY[i], magmaFloatComplex, vecsize );
printf("device %2d n_local = %4d\n", (int) i, (int) n_local[i]);
}
magma_setdevice(0);
//////////////////////////////////////////////////////////////////////////
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &matsize, A );
magma_cmake_hermitian( N, A, LDA );
blocks = N / nb + (N % nb != 0);
lwork = LDA * (blocks + 1);
TESTING_MALLOC_CPU( C_work, magmaFloatComplex, lwork );
for(i=0; i<num_gpus; i++){
magma_setdevice(i);
TESTING_MALLOC_DEV( dC_work[i], magmaFloatComplex, lwork );
//fillZero(dC_work[i], lwork);
}
magma_setdevice(0);
//////////////////////////////////////////////////////////////////////////
fp = fopen ("results_chemv_mgpu.csv", "w") ;
if( fp == NULL ){ printf("Couldn't open output file\n"); exit(1);}
printf("CHEMV magmaFloatComplex precision\n\n");
printf( " n CUBLAS,Gflop/s MAGMABLAS,Gflop/s \"error\"\n"
"==============================================================\n");
fprintf(fp, " n CUBLAS,Gflop/s MAGMABLAS,Gflop/s \"error\"\n"
"==============================================================\n");
// for( offset = 0; offset< N; offset ++ )
for(int size = istart ; size <= N ; size += 128)
{
// printf("offset = %d ", offset);
m = size ;
// m = N;
// lda = ((m+31)/32)*32;//
lda = LDA;
flops = FLOPS( (float)m ) / 1e6;
printf( "N %5d ", (int) m );
fprintf( fp, "%5d, ", (int) m );
vecsize = m * incx;
lapackf77_clarnv( &ione, ISEED, &vecsize, X );
lapackf77_clarnv( &ione, ISEED, &vecsize, Y[0] );
/* =====================================================================
Performs operation using CUDA-BLAS
=================================================================== */
magma_setdevice(0);
magma_csetmatrix_1D_col_bcyclic(m, m, A, LDA, d_lA, lda, num_gpus, nb);
magma_setdevice(0);
magma_csetmatrix( m, m, A, LDA, dA, lda );
magma_csetvector( m, Y[0], incx, dYcublas, incx );
for(i=0; i<num_gpus; i++){
magma_setdevice(i);
magma_csetvector( m, X, incx, dX[i], incx );
magma_csetvector( m, Y[0], incx, dY[i], incx );
blocks = m / nb + (m % nb != 0);
magma_csetmatrix( lda, blocks, C_work, LDA, dC_work[i], lda );
}
magma_setdevice(0);
start = get_current_time();
cublasChemv( uplo, m-offset, alpha, dA + offset + offset * lda, lda, dX[0] + offset, incx, beta, dYcublas + offset, incx );
end = get_current_time();
magma_cgetvector( m, dYcublas, incx, Ycublas, incx );
cuda_perf = flops / GetTimerValue(start,end);
printf( "%11.2f", cuda_perf );
fprintf(fp, "%11.2f,", cuda_perf );
magma_setdevice(0);
start = get_current_time();
if(nb == 32)
{
magmablas_chemv2_mgpu_32_offset( uplo, m, alpha, d_lA, lda, dX, incx, beta, dY, incx,
dC_work, lwork, num_gpus, nb, offset);
}
else // nb = 64
{
magmablas_chemv2_mgpu_offset( uplo, m, alpha, d_lA, lda, dX, incx, beta, dY, incx,
dC_work, lwork, num_gpus, nb, offset);
}
for(i=1; i<num_gpus; i++)
{
magma_setdevice(i);
cudaDeviceSynchronize();
}
end = get_current_time();
magma_perf = flops / GetTimerValue(start,end);
printf( "%11.2f", magma_perf );
fprintf(fp, "%11.2f,", magma_perf );
for(i=0; i<num_gpus; i++)
{
magma_setdevice(i);
magma_cgetvector( m, dY[i], incx, Y[i], incx );
}
magma_setdevice(0);
#ifdef validate
for( j= offset;j<m;j++)
{
for(i=1; i<num_gpus; i++)
{
// printf("Y[%d][%d] = %15.14f\n", i, j, Y[i][j].x);
#if defined(PRECISION_z) || defined(PRECISION_c)
Y[0][j].x = Y[0][j].x + Y[i][j].x;
Y[0][j].y = Y[0][j].y + Y[i][j].y;
#else
Y[0][j] = Y[0][j] + Y[i][j];
#endif
}
}
/*
#if defined(PRECISION_z) || defined(PRECISION_c)
for( j=offset;j<m;j++)
{
if(Y[0][j].x != Ycublas[j].x)
{
printf("Y-multi[%d] = %f, %f\n", j, Y[0][j].x, Y[0][j].y );
printf("Ycublas[%d] = %f, %f\n", j, Ycublas[j].x, Ycublas[j].y);
}
}
#else
for( j=offset;j<m;j++)
{
if(Y[0][j] != Ycublas[j])
{
printf("Y-multi[%d] = %f\n", j, Y[0][j] );
printf("Ycublas[%d] = %f\n", j, Ycublas[j]);
}
}
#endif
*/
/* =====================================================================
Computing the Difference Cublas VS Magma
=================================================================== */
magma_int_t nw = m - offset ;
blasf77_caxpy( &nw, &c_neg_one, Y[0] + offset, &incx, Ycublas + offset, &incx);
error = lapackf77_clange( "M", &nw, &ione, Ycublas + offset, &nw, work );
#if 0
printf( "\t\t %8.6e", error / m );
fprintf( fp, "\t\t %8.6e", error / m );
/*
* Extra check with cblas vs magma
*/
cblas_ccopy( m, Y, incx, Ycublas, incx );
cblas_chemv( CblasColMajor, CblasLower, m,
CBLAS_SADDR(alpha), A, LDA, X, incx,
CBLAS_SADDR(beta), Ycublas, incx );
blasf77_caxpy( &m, &c_neg_one, Ymagma, &incx, Ycublas, &incx);
error = lapackf77_clange( "M", &m, &ione, Ycublas, &m, work );
#endif
printf( "\t\t %8.6e", error / m );
fprintf( fp, "\t\t %8.6e", error / m );
#endif
printf("\n");
fprintf(fp, "\n");
}
fclose( fp ) ;
/* Free Memory */
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( X );
TESTING_FREE_CPU( Ycublas );
TESTING_FREE_CPU( Ymagma );
TESTING_FREE_CPU( C_work );
magma_setdevice(0);
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dYcublas );
for(i=0; i<num_gpus; i++)
{
TESTING_FREE_CPU( Y[i] );
magma_setdevice(i);
TESTING_FREE_DEV( d_lA[i] );
TESTING_FREE_DEV( dX[i] );
TESTING_FREE_DEV( dY[i] );
TESTING_FREE_DEV( dC_work[i] );
}
magma_setdevice(0);
///////////////////////////////////////////////////////////
/* Free device */
TESTING_FINALIZE();
return 0;
}
magma_int_t
magma_cmgenerator(
magma_int_t n,
magma_int_t offdiags,
magma_index_t *diag_offset,
magmaFloatComplex *diag_vals,
magma_c_matrix *A,
magma_queue_t queue )
{
magma_int_t info = 0;
magma_c_matrix B={Magma_CSR};
B.val = NULL;
B.col = NULL;
B.row = NULL;
B.rowidx = NULL;
B.blockinfo = NULL;
B.diag = NULL;
B.dval = NULL;
B.dcol = NULL;
B.drow = NULL;
B.drowidx = NULL;
B.ddiag = NULL;
B.list = NULL;
B.dlist = NULL;
B.num_rows = n;
B.num_cols = n;
B.fill_mode = MagmaFull;
B.memory_location = Magma_CPU;
B.storage_type = Magma_ELLPACKT;
B.max_nnz_row = (2*offdiags+1);
CHECK( magma_cmalloc_cpu( &B.val, B.max_nnz_row*n ));
CHECK( magma_index_malloc_cpu( &B.col, B.max_nnz_row*n ));
for( int i=0; i<n; i++ ) { // stride over rows
// stride over the number of nonzeros in each row
// left of diagonal
for( int j=0; j<offdiags; j++ ) {
B.val[ i*B.max_nnz_row + j ] = diag_vals[ offdiags - j ];
B.col[ i*B.max_nnz_row + j ] = -1 * diag_offset[ offdiags-j ] + i;
}
// elements on the diagonal
B.val[ i*B.max_nnz_row + offdiags ] = diag_vals[ 0 ];
B.col[ i*B.max_nnz_row + offdiags ] = i;
// right of diagonal
for( int j=0; j<offdiags; j++ ) {
B.val[ i*B.max_nnz_row + j + offdiags +1 ] = diag_vals[ j+1 ];
B.col[ i*B.max_nnz_row + j + offdiags +1 ] = diag_offset[ j+1 ] + i;
}
}
// set invalid entries to zero
for( int i=0; i<n; i++ ) { // stride over rows
for( int j=0; j<B.max_nnz_row; j++ ) { // nonzeros in every row
if ( (B.col[i*B.max_nnz_row + j] < 0) ||
(B.col[i*B.max_nnz_row + j] >= n) ) {
B.val[ i*B.max_nnz_row + j ] = MAGMA_C_MAKE( 0.0, 0.0 );
}
}
}
B.nnz = 0;
for( int i=0; i<n; i++ ) { // stride over rows
for( int j=0; j<B.max_nnz_row; j++ ) { // nonzeros in every row
if ( MAGMA_C_REAL( B.val[i*B.max_nnz_row + j]) != 0.0 )
B.nnz++;
}
}
B.true_nnz = B.nnz;
// converting it to CSR will remove the invalit entries completely
CHECK( magma_cmconvert( B, A, Magma_ELLPACKT, Magma_CSR, queue ));
cleanup:
if( info != 0 ){
magma_cmfree( &B, queue );
}
return info;
}
extern "C" magma_int_t
magma_cheevdx_2stage_m(magma_int_t nrgpu, char jobz, char range, char uplo,
magma_int_t n,
magmaFloatComplex *a, magma_int_t lda,
float vl, float vu, magma_int_t il, magma_int_t iu,
magma_int_t *m, float *w,
magmaFloatComplex *work, magma_int_t lwork,
float *rwork, magma_int_t lrwork,
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
=======
CHEEVD_2STAGE computes all eigenvalues and, optionally, eigenvectors of a
complex Hermitian matrix A. It uses a two-stage algorithm for the tridiagonalization.
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.
RANGE (input) CHARACTER*1
= 'A': all eigenvalues will be found.
= 'V': all eigenvalues in the half-open interval (VL,VU]
will be found.
= 'I': the IL-th through IU-th eigenvalues will be found.
UPLO (input) CHARACTER*1
= 'U': Upper triangle of A is stored;
= 'L': Lower triangle of A is stored.
N (input) INTEGER
The order of the matrix A. N >= 0.
A (input/output) COMPLEX array, dimension (LDA, N)
On entry, the Hermitian matrix A. If UPLO = 'U', the
leading N-by-N upper triangular part of A contains the
upper triangular part of the matrix A. 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, the first m columns
of A contains the required
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.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
VL (input) DOUBLE PRECISION
VU (input) DOUBLE PRECISION
If RANGE='V', the lower and upper bounds of the interval to
be searched for eigenvalues. VL < VU.
Not referenced if RANGE = 'A' or 'I'.
IL (input) INTEGER
IU (input) INTEGER
If RANGE='I', 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 = 'A' or 'V'.
M (output) INTEGER
The total number of eigenvalues found. 0 <= M <= N.
If RANGE = 'A', M = N, and if RANGE = 'I', M = IU-IL+1.
W (output) DOUBLE PRECISION array, dimension (N)
If INFO = 0, the required m eigenvalues in ascending order.
WORK (workspace/output) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) 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 >= LQ2 + N * (NB + 1).
If JOBZ = 'V' and N > 1, LWORK >= LQ2 + 2*N + N**2.
where LQ2 is the size needed to store
the Q2 matrix and is returned by
MAGMA_BULGE_GET_LQ2.
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal sizes of the WORK, RWORK and
IWORK arrays, returns these values as the first entries of
the WORK, RWORK and IWORK arrays, and no error message
related to LWORK or LRWORK or LIWORK is issued by XERBLA.
RWORK (workspace/output) DOUBLE PRECISION array,
dimension (LRWORK)
On exit, if INFO = 0, RWORK(1) returns the optimal LRWORK.
LRWORK (input) INTEGER
The dimension of the array RWORK.
If N <= 1, LRWORK >= 1.
If JOBZ = 'N' and N > 1, LRWORK >= N.
If JOBZ = 'V' and N > 1, LRWORK >= 1 + 5*N + 2*N**2.
If LRWORK = -1, then a workspace query is assumed; the
routine only calculates the optimal sizes of the WORK, RWORK
and IWORK arrays, returns these values as the first entries
of the WORK, RWORK and IWORK arrays, and no error message
related to LWORK or LRWORK or LIWORK is issued by XERBLA.
IWORK (workspace/output) INTEGER array, dimension (MAX(1,LIWORK))
On exit, if INFO = 0, IWORK(1) 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, RWORK
and IWORK arrays, returns these values as the first entries
of the WORK, RWORK and IWORK arrays, and no error message
related to LWORK or LRWORK or LIWORK is issued by XERBLA.
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};
char range_[2] = {range, 0};
magmaFloatComplex c_one = MAGMA_C_ONE;
magma_int_t ione = 1;
magma_int_t izero = 0;
float d_one = 1.;
float d__1;
float eps;
float anrm;
magma_int_t imax;
float rmin, rmax;
float sigma;
//magma_int_t iinfo;
magma_int_t lwmin, lrwmin, liwmin;
magma_int_t lower;
magma_int_t wantz;
magma_int_t iscale;
float safmin;
float bignum;
float smlnum;
magma_int_t lquery;
magma_int_t alleig, valeig, indeig;
/* determine the number of threads */
magma_int_t threads = magma_get_numthreads();
magma_setlapack_numthreads(threads);
wantz = lapackf77_lsame(jobz_, MagmaVecStr);
lower = lapackf77_lsame(uplo_, MagmaLowerStr);
alleig = lapackf77_lsame( range_, "A" );
valeig = lapackf77_lsame( range_, "V" );
indeig = lapackf77_lsame( range_, "I" );
lquery = lwork == -1 || lrwork == -1 || liwork == -1;
*info = 0;
if (! (wantz || lapackf77_lsame(jobz_, MagmaNoVecStr))) {
*info = -1;
} else if (! (alleig || valeig || indeig)) {
*info = -2;
} else if (! (lower || lapackf77_lsame(uplo_, MagmaUpperStr))) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (lda < max(1,n)) {
*info = -6;
} 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_cbulge_nb(n, threads);
magma_int_t Vblksiz = magma_cbulge_get_Vblksiz(n, nb, threads);
magma_int_t ldt = Vblksiz;
magma_int_t ldv = nb + Vblksiz;
magma_int_t blkcnt = magma_bulge_get_blkcnt(n, nb, Vblksiz);
magma_int_t lq2 = magma_cbulge_get_lq2(n, threads);
if (wantz) {
lwmin = lq2 + 2 * n + n * n;
lrwmin = 1 + 5 * n + 2 * n * n;
liwmin = 5 * n + 3;
} else {
lwmin = lq2 + n * (nb + 1);
lrwmin = n;
liwmin = 1;
}
work[0] = MAGMA_C_MAKE( lwmin * (1. + lapackf77_slamch("Epsilon")), 0.); // round up
rwork[0] = lrwmin * (1. + lapackf77_slamch("Epsilon"));
iwork[0] = liwmin;
if ((lwork < lwmin) && !lquery) {
*info = -14;
} else if ((lrwork < lrwmin) && ! lquery) {
*info = -16;
} else if ((liwork < liwmin) && ! lquery) {
*info = -18;
}
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) {
w[0] = MAGMA_C_REAL(a[0]);
if (wantz) {
a[0] = MAGMA_C_ONE;
}
return *info;
}
#ifdef ENABLE_DEBUG
printf("using %d threads\n", threads);
#endif
/* Check if matrix is very small then just call LAPACK on CPU, no need for GPU */
magma_int_t ntiles = n/nb;
if( ( ntiles < 2 ) || ( 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
lapackf77_cheevd(jobz_, uplo_, &n,
a, &lda, w,
work, &lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, &lrwork,
#endif
iwork, &liwork,
info);
*m = n;
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 = lapackf77_clanhe("M", uplo_, &n, a, &lda, rwork);
iscale = 0;
if (anrm > 0. && anrm < rmin) {
iscale = 1;
sigma = rmin / anrm;
} else if (anrm > rmax) {
iscale = 1;
sigma = rmax / anrm;
}
if (iscale == 1) {
lapackf77_clascl(uplo_, &izero, &izero, &d_one, &sigma, &n, &n, a,
&lda, info);
}
magma_int_t indT2 = 0;
magma_int_t indTAU2 = indT2 + blkcnt*ldt*Vblksiz;
magma_int_t indV2 = indTAU2+ blkcnt*Vblksiz;
magma_int_t indtau1 = indV2 + blkcnt*ldv*Vblksiz;
magma_int_t indwrk = indtau1+ n;
magma_int_t indwk2 = indwrk + n * n;
magma_int_t llwork = lwork - indwrk;
magma_int_t llwrk2 = lwork - indwk2;
magma_int_t inde = 0;
magma_int_t indrwk = inde + n;
magma_int_t llrwk = lrwork - indrwk;
#ifdef ENABLE_TIMER
magma_timestr_t start, st1, st2, end;
start = get_current_time();
#endif
#ifdef HE2HB_SINGLEGPU
magmaFloatComplex *dT1;
if (MAGMA_SUCCESS != magma_cmalloc( &dT1, n*nb)) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
#ifdef ENABLE_TIMER
tband1 = get_current_time();
#endif
magma_chetrd_he2hb(uplo, n, nb, a, lda, &work[indtau1], &work[indwrk], llwork, dT1, threads, info);
#ifdef ENABLE_TIMER
tband2 = get_current_time();
printf(" 1 GPU seq code time chetrd_he2hb only = %7.4f\n" , GetTimerValue(tband1,tband2)/1000.);
#endif
magma_free(dT1);
#else
magma_int_t nstream = max(3,nrgpu+2);
magma_queue_t streams[MagmaMaxGPUs][20];
magmaFloatComplex *da[MagmaMaxGPUs],*dT1[MagmaMaxGPUs];
magma_int_t ldda = ((n+31)/32)*32;
magma_int_t ver = 0;
magma_int_t distblk = max(256, 4*nb);
#ifdef ENABLE_DEBUG
printf("voici ngpu %d distblk %d NB %d nstream %d version %d \n ",nrgpu,distblk,nb,nstream,ver);
#endif
#ifdef ENABLE_TIMER
magma_timestr_t tband1, tband2, t1, t2, ta1, ta2;
t1 = get_current_time();
#endif
for( magma_int_t dev = 0; dev < nrgpu; ++dev ) {
magma_int_t mlocal = ((n / distblk) / nrgpu + 1) * distblk;
magma_setdevice( dev );
magma_cmalloc(&da[dev], ldda*mlocal );
magma_cmalloc(&dT1[dev], (n*nb) );
for( int i = 0; i < nstream; ++i ) {
magma_queue_create( &streams[dev][i] );
}
}
#ifdef ENABLE_TIMER
t2 = get_current_time();
#endif
magma_csetmatrix_1D_col_bcyclic( n, n, a, lda, da, ldda, nrgpu, distblk);
magma_setdevice(0);
#ifdef ENABLE_TIMER
tband1 = get_current_time();
#endif
if(ver==30){
magma_chetrd_he2hb_mgpu_spec(uplo, n, nb, a, lda, &work[indtau1], &work[indwrk], llwork, da, ldda, dT1, nb, nrgpu, distblk, streams, nstream, threads, info);
}else{
magma_chetrd_he2hb_mgpu(uplo, n, nb, a, lda, &work[indtau1], &work[indwrk], llwork, da, ldda, dT1, nb, nrgpu, distblk, streams, nstream, threads, info);
}
#ifdef ENABLE_TIMER
tband2 = get_current_time();
printf(" time alloc %7.4f, ditribution %7.4f, chetrd_he2hb only = %7.4f\n" , GetTimerValue(t1,t2)/1000., GetTimerValue(t2,tband1)/1000., GetTimerValue(tband1,tband2)/1000.);
#endif
for( magma_int_t dev = 0; dev < nrgpu; ++dev ) {
magma_setdevice( dev );
magma_free( da[dev] );
magma_free( dT1[dev] );
for( int i = 0; i < nstream; ++i ) {
magma_queue_destroy( streams[dev][i] );
}
}
#endif
#ifdef ENABLE_TIMER
st1 = get_current_time();
printf(" time chetrd_he2hb_mgpu = %6.2f\n" , GetTimerValue(start,st1)/1000.);
#endif
/* copy the input matrix into WORK(INDWRK) with band storage */
/* PAY ATTENTION THAT work[indwrk] should be able to be of size lda2*n which it should be checked in any future modification of lwork.*/
magma_int_t lda2 = 2*nb; //nb+1+(nb-1);
magmaFloatComplex* A2 = &work[indwrk];
memset(A2 , 0, n*lda2*sizeof(magmaFloatComplex));
for (magma_int_t j = 0; j < n-nb; j++)
{
cblas_ccopy(nb+1, &a[j*(lda+1)], 1, &A2[j*lda2], 1);
memset(&a[j*(lda+1)], 0, (nb+1)*sizeof(magmaFloatComplex));
a[nb + j*(lda+1)] = c_one;
}
for (magma_int_t j = 0; j < nb; j++)
{
cblas_ccopy(nb-j, &a[(j+n-nb)*(lda+1)], 1, &A2[(j+n-nb)*lda2], 1);
memset(&a[(j+n-nb)*(lda+1)], 0, (nb-j)*sizeof(magmaFloatComplex));
}
#ifdef ENABLE_TIMER
st2 = get_current_time();
printf(" time chetrd_convert = %6.2f\n" , GetTimerValue(st1,st2)/1000.);
#endif
magma_chetrd_hb2st(threads, uplo, n, nb, Vblksiz, A2, lda2, w, &rwork[inde], &work[indV2], ldv, &work[indTAU2], wantz, &work[indT2], ldt);
#ifdef ENABLE_TIMER
end = get_current_time();
printf(" time chetrd_hb2st = %6.2f\n" , GetTimerValue(st2,end)/1000.);
printf(" time chetrd = %6.2f\n", GetTimerValue(start,end)/1000.);
#endif
/* For eigenvalues only, call SSTERF. For eigenvectors, first call
CSTEDC to generate the eigenvector matrix, WORK(INDWRK), of the
tridiagonal matrix, then call CUNMTR to multiply it to the Householder
transformations represented as Householder vectors in A. */
if (! wantz) {
#ifdef ENABLE_TIMER
start = get_current_time();
#endif
lapackf77_ssterf(&n, w, &rwork[inde], info);
magma_smove_eig(range, n, w, &il, &iu, vl, vu, m);
#ifdef ENABLE_TIMER
end = get_current_time();
printf(" time dstedc = %6.2f\n", GetTimerValue(start,end)/1000.);
#endif
} else {
#ifdef ENABLE_TIMER
start = get_current_time();
#endif
magma_cstedx_m(nrgpu, range, n, vl, vu, il, iu, w, &rwork[inde],
&work[indwrk], n, &rwork[indrwk],
llrwk, iwork, liwork, info);
#ifdef ENABLE_TIMER
end = get_current_time();
printf(" time cstedx_m = %6.2f\n", GetTimerValue(start,end)/1000.);
start = get_current_time();
#endif
magma_smove_eig(range, n, w, &il, &iu, vl, vu, m);
/*
magmaFloatComplex *dZ;
magma_int_t lddz = n;
if (MAGMA_SUCCESS != magma_cmalloc( &dZ, *m*lddz)) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_cbulge_back(threads, uplo, n, nb, *m, Vblksiz, &work[indwrk + n * (il-1)], n, dZ, lddz,
&work[indV2], ldv, &work[indTAU2], &work[indT2], ldt, info);
magma_cgetmatrix( n, *m, dZ, lddz, &work[indwrk], n);
magma_free(dZ);
*/
magma_cbulge_back_m(nrgpu, threads, uplo, n, nb, *m, Vblksiz, &work[indwrk + n * (il-1)], n,
&work[indV2], ldv, &work[indTAU2], &work[indT2], ldt, info);
#ifdef ENABLE_TIMER
st1 = get_current_time();
printf(" time cbulge_back_m = %6.2f\n" , GetTimerValue(start,st1)/1000.);
#endif
magma_cunmqr_m(nrgpu, MagmaLeft, MagmaNoTrans, n-nb, *m, n-nb, a+nb, lda, &work[indtau1],
&work[indwrk + n * (il-1) + nb], n, &work[indwk2], llwrk2, info);
lapackf77_clacpy("A", &n, m, &work[indwrk + n * (il-1)], &n, a, &lda);
#ifdef ENABLE_TIMER
end = get_current_time();
printf(" time cunmqr_m + copy = %6.2f\n", GetTimerValue(st1,end)/1000.);
printf(" time eigenvectors backtransf. = %6.2f\n" , GetTimerValue(start,end)/1000.);
#endif
}
/* If matrix was scaled, then rescale eigenvalues appropriately. */
if (iscale == 1) {
if (*info == 0) {
imax = n;
} else {
imax = *info - 1;
}
d__1 = 1. / sigma;
blasf77_sscal(&imax, &d__1, w, &ione);
}
work[0] = MAGMA_C_MAKE((float) lwmin, 0.);
rwork[0] = (float) lrwmin;
iwork[0] = liwmin;
return *info;
} /* magma_cheevdx_2stage_m */
extern "C" magma_int_t
magma_cgelqf( magma_int_t m, magma_int_t n,
magmaFloatComplex *a, magma_int_t lda, magmaFloatComplex *tau,
magmaFloatComplex *work, magma_int_t lwork, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CGELQF computes an LQ factorization of a COMPLEX M-by-N matrix A:
A = L * Q.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. N >= 0.
A (input/output) COMPLEX array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, the elements on and below the diagonal of the array
contain the m-by-min(m,n) lower trapezoidal matrix L (L is
lower triangular if m <= n); the elements above the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of elementary reflectors (see Further Details).
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,M).
TAU (output) COMPLEX array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
WORK (workspace/output) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
Higher performance is achieved if WORK is in pinned memory, e.g.
allocated using magma_malloc_pinned.
LWORK (input) INTEGER
The dimension of the array WORK. LWORK >= max(1,M).
For optimum performance LWORK >= M*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.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
if INFO = -10 internal GPU memory allocation failed.
Further Details
===============
The matrix Q is represented as a product of elementary reflectors
Q = H(k) . . . H(2) H(1), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:n) is stored on exit in A(i,i+1:n),
and tau in TAU(i).
===================================================================== */
#define a_ref(a_1,a_2) ( a+(a_2)*(lda) + (a_1))
magmaFloatComplex *dA, *dAT;
magmaFloatComplex c_one = MAGMA_C_ONE;
magma_int_t maxm, maxn, maxdim, nb;
magma_int_t iinfo, ldda;
int lquery;
/* Function Body */
*info = 0;
nb = magma_get_cgelqf_nb(m);
work[0] = MAGMA_C_MAKE( (float)(m*nb), 0 );
lquery = (lwork == -1);
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,m)) {
*info = -4;
} else if (lwork < max(1,m) && ! lquery) {
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (min(m, n) == 0) {
work[0] = c_one;
return *info;
}
maxm = ((m + 31)/32)*32;
maxn = ((n + 31)/32)*32;
maxdim = max(maxm, maxn);
if (maxdim*maxdim < 2*maxm*maxn)
{
ldda = maxdim;
if (MAGMA_SUCCESS != magma_cmalloc( &dA, maxdim*maxdim )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_csetmatrix( m, n, a, lda, dA, ldda );
dAT = dA;
magmablas_ctranspose_inplace( ldda, dAT, ldda );
}
else
{
ldda = maxn;
if (MAGMA_SUCCESS != magma_cmalloc( &dA, 2*maxn*maxm )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_csetmatrix( m, n, a, lda, dA, maxm );
dAT = dA + maxn * maxm;
magmablas_ctranspose2( dAT, ldda, dA, maxm, m, n );
}
magma_cgeqrf2_gpu(n, m, dAT, ldda, tau, &iinfo);
if (maxdim*maxdim < 2*maxm*maxn) {
magmablas_ctranspose_inplace( ldda, dAT, ldda );
magma_cgetmatrix( m, n, dA, ldda, a, lda );
} else {
magmablas_ctranspose2( dA, maxm, dAT, ldda, n, m );
magma_cgetmatrix( m, n, dA, maxm, a, lda );
}
magma_free( dA );
return *info;
} /* magma_cgelqf */
extern "C" magma_int_t
magma_ccg_merge(
magma_c_matrix A, magma_c_matrix b, magma_c_matrix *x,
magma_c_solver_par *solver_par,
magma_queue_t queue )
{
magma_int_t info = MAGMA_NOTCONVERGED;
// prepare solver feedback
solver_par->solver = Magma_CGMERGE;
solver_par->numiter = 0;
solver_par->spmv_count = 0;
// solver variables
magmaFloatComplex alpha, beta, gamma, rho, tmp1, *skp_h={0};
float nom, nom0, betanom, den, nomb;
// some useful variables
magmaFloatComplex c_zero = MAGMA_C_ZERO, c_one = MAGMA_C_ONE;
magma_int_t dofs = A.num_rows*b.num_cols;
magma_c_matrix r={Magma_CSR}, d={Magma_CSR}, z={Magma_CSR}, B={Magma_CSR}, C={Magma_CSR};
magmaFloatComplex *d1=NULL, *d2=NULL, *skp=NULL;
// GPU workspace
CHECK( magma_cvinit( &r, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_cvinit( &d, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_cvinit( &z, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_cmalloc( &d1, dofs*(1) ));
CHECK( magma_cmalloc( &d2, dofs*(1) ));
// array for the parameters
CHECK( magma_cmalloc( &skp, 6 ));
// skp = [alpha|beta|gamma|rho|tmp1|tmp2]
// solver setup
magma_cscal( dofs, c_zero, x->dval, 1, queue ); // x = 0
//CHECK( magma_cresidualvec( A, b, *x, &r, nom0, queue));
magma_ccopy( dofs, b.dval, 1, r.dval, 1, queue ); // r = b
magma_ccopy( dofs, r.dval, 1, d.dval, 1, queue ); // d = r
nom0 = betanom = magma_scnrm2( dofs, r.dval, 1, queue );
nom = nom0 * nom0; // nom = r' * r
CHECK( magma_c_spmv( c_one, A, d, c_zero, z, queue )); // z = A d
den = MAGMA_C_ABS( magma_cdotc( dofs, d.dval, 1, z.dval, 1, queue ) ); // den = d'* z
solver_par->init_res = nom0;
nomb = magma_scnrm2( dofs, b.dval, 1, queue );
if ( nomb == 0.0 ){
nomb=1.0;
}
// array on host for the parameters
CHECK( magma_cmalloc_cpu( &skp_h, 6 ));
alpha = rho = gamma = tmp1 = c_one;
beta = magma_cdotc( dofs, r.dval, 1, r.dval, 1, queue );
skp_h[0]=alpha;
skp_h[1]=beta;
skp_h[2]=gamma;
skp_h[3]=rho;
skp_h[4]=tmp1;
skp_h[5]=MAGMA_C_MAKE(nom, 0.0);
magma_csetvector( 6, skp_h, 1, skp, 1, queue );
if( nom0 < solver_par->atol ||
nom0/nomb < solver_par->rtol ){
info = MAGMA_SUCCESS;
goto cleanup;
}
solver_par->final_res = solver_par->init_res;
solver_par->iter_res = solver_par->init_res;
if ( solver_par->verbose > 0 ) {
solver_par->res_vec[0] = (real_Double_t) nom0;
solver_par->timing[0] = 0.0;
}
// check positive definite
if (den <= 0.0) {
info = MAGMA_NONSPD;
goto cleanup;
}
//Chronometry
real_Double_t tempo1, tempo2;
tempo1 = magma_sync_wtime( queue );
solver_par->numiter = 0;
solver_par->spmv_count = 0;
// start iteration
do
{
solver_par->numiter++;
// computes SpMV and dot product
CHECK( magma_ccgmerge_spmv1( A, d1, d2, d.dval, z.dval, skp, queue ));
solver_par->spmv_count++;
// updates x, r, computes scalars and updates d
CHECK( magma_ccgmerge_xrbeta( dofs, d1, d2, x->dval, r.dval, d.dval, z.dval, skp, queue ));
// check stopping criterion (asynchronous copy)
magma_cgetvector( 1 , skp+1, 1, skp_h+1, 1, queue );
betanom = sqrt(MAGMA_C_ABS(skp_h[1]));
if ( solver_par->verbose > 0 ) {
tempo2 = magma_sync_wtime( queue );
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
if ( betanom < solver_par->atol ||
betanom/nomb < solver_par->rtol ) {
break;
}
}
while ( solver_par->numiter+1 <= solver_par->maxiter );
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t) tempo2-tempo1;
float residual;
CHECK( magma_cresidualvec( A, b, *x, &r, &residual, queue));
solver_par->iter_res = betanom;
solver_par->final_res = residual;
if ( solver_par->numiter < solver_par->maxiter ) {
info = MAGMA_SUCCESS;
} else if ( solver_par->init_res > solver_par->final_res ) {
if ( solver_par->verbose > 0 ) {
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
info = MAGMA_SLOW_CONVERGENCE;
if( solver_par->iter_res < solver_par->atol ||
solver_par->iter_res/solver_par->init_res < solver_par->rtol ){
info = MAGMA_SUCCESS;
}
}
else {
if ( solver_par->verbose > 0 ) {
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
solver_par->info = MAGMA_DIVERGENCE;
}
cleanup:
magma_cmfree(&r, queue );
magma_cmfree(&z, queue );
magma_cmfree(&d, queue );
magma_cmfree(&B, queue );
magma_cmfree(&C, queue );
magma_free( d1 );
magma_free( d2 );
magma_free( skp );
magma_free_cpu( skp_h );
solver_par->info = info;
return info;
} /* magma_ccg_merge */
/**
Purpose
-------
CHEEVX computes selected eigenvalues and, optionally, eigenvectors
of a complex Hermitian matrix A. Eigenvalues and eigenvectors can
be selected by specifying either a range of values or a range of
indices for the desired eigenvalues.
Arguments
---------
@param[in]
jobz magma_vec_t
- = MagmaNoVec: Compute eigenvalues only;
- = MagmaVec: Compute eigenvalues and eigenvectors.
@param[in]
range magma_range_t
- = MagmaRangeAll: all eigenvalues will be found.
- = MagmaRangeV: all eigenvalues in the half-open interval (VL,VU]
will be found.
- = MagmaRangeI: the IL-th through IU-th eigenvalues will be found.
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored;
- = MagmaLower: Lower triangle of A is stored.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
dA COMPLEX array, dimension (LDDA, N)
On entry, the Hermitian matrix A. If UPLO = MagmaUpper, the
leading N-by-N upper triangular part of A contains the
upper triangular part of the matrix A. If UPLO = MagmaLower,
the leading N-by-N lower triangular part of A contains
the lower triangular part of the matrix A.
On exit, 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[in]
abstol REAL
The absolute error tolerance for the eigenvalues.
An approximate eigenvalue is accepted as converged
when it is determined to lie in an interval [a,b]
of width less than or equal to
ABSTOL + EPS * max( |a|,|b| ),
\n
where EPS is the machine precision. If ABSTOL is less than
or equal to zero, then EPS*|T| will be used in its place,
where |T| is the 1-norm of the tridiagonal matrix obtained
by reducing A to tridiagonal form.
\n
Eigenvalues will be computed most accurately when ABSTOL is
set to twice the underflow threshold 2*SLAMCH('S'), not zero.
If this routine returns with INFO > 0, indicating that some
eigenvectors did not converge, try setting ABSTOL to
2*SLAMCH('S').
\n
See "Computing Small Singular Values of Bidiagonal Matrices
with Guaranteed High Relative Accuracy," by Demmel and
Kahan, LAPACK Working Note #3.
@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)
On normal exit, the first M elements contain the selected
eigenvalues in ascending order.
@param[out]
dZ COMPLEX array, dimension (LDDZ, max(1,M))
If JOBZ = MagmaVec, then if INFO = 0, the first M columns of Z
contain the orthonormal eigenvectors of the matrix A
corresponding to the selected eigenvalues, with the i-th
column of Z holding the eigenvector associated with W(i).
If an eigenvector fails to converge, then that column of Z
contains the latest approximation to the eigenvector, and the
index of the eigenvector is returned in IFAIL.
If JOBZ = MagmaNoVec, then Z is not referenced.
Note: the user must ensure that at least max(1,M) columns are
supplied in the array Z; if RANGE = MagmaRangeV, the exact value of M
is not known in advance and an upper bound must be used.
********* (workspace) If FAST_HEMV is defined DZ should be (LDDZ, max(1,N)) in both cases.
@param[in]
lddz INTEGER
The leading dimension of the array DZ. LDDZ >= 1, and if
JOBZ = MagmaVec, LDDZ >= max(1,N).
@param
wA (workspace) COMPLEX array, dimension (LDWA, N)
@param[in]
ldwa INTEGER
The leading dimension of the array wA. LDWA >= max(1,N).
@param
wZ (workspace) COMPLEX array, dimension (LDWZ, max(1,M))
@param[in]
ldwz INTEGER
The leading dimension of the array wZ. LDWZ >= 1, and if
JOBZ = MagmaVec, LDWZ >= max(1,N).
@param[out]
work (workspace) COMPLEX array, dimension (LWORK)
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The length of the array WORK. LWORK >= (NB+1)*N,
where NB is the max of the blocksize for CHETRD.
\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) REAL array, dimension (7*N)
@param
iwork (workspace) INTEGER array, dimension (5*N)
@param[out]
ifail INTEGER array, dimension (N)
If JOBZ = MagmaVec, then if INFO = 0, the first M elements of
IFAIL are zero. If INFO > 0, then IFAIL contains the
indices of the eigenvectors that failed to converge.
If JOBZ = MagmaNoVec, then IFAIL is not referenced.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: if INFO = i, then i eigenvectors failed to converge.
Their indices are stored in array IFAIL.
@ingroup magma_cheev_driver
********************************************************************/
extern "C" magma_int_t
magma_cheevx_gpu(
magma_vec_t jobz, magma_range_t range, magma_uplo_t uplo, magma_int_t n,
magmaFloatComplex_ptr dA, magma_int_t ldda,
float vl, float vu,
magma_int_t il, magma_int_t iu, float abstol, magma_int_t *m,
float *w,
magmaFloatComplex_ptr dZ, magma_int_t lddz,
magmaFloatComplex *wA, magma_int_t ldwa,
magmaFloatComplex *wZ, magma_int_t ldwz,
magmaFloatComplex *work, magma_int_t lwork,
float *rwork, magma_int_t *iwork, magma_int_t *ifail,
magma_int_t *info)
{
const char* uplo_ = lapack_uplo_const( uplo );
const char* jobz_ = lapack_vec_const( jobz );
const char* range_ = lapack_range_const( range );
magma_int_t ione = 1;
const char* order_;
magma_int_t indd, inde;
magma_int_t imax;
magma_int_t lopt, itmp1, indee;
magma_int_t lower, wantz;
magma_int_t i, j, jj, i__1;
magma_int_t alleig, valeig, indeig;
magma_int_t iscale, indibl;
magma_int_t indiwk, indisp, indtau;
magma_int_t indrwk, indwrk;
magma_int_t llwork, nsplit;
magma_int_t lquery;
magma_int_t iinfo;
float safmin;
float bignum;
float smlnum;
float eps, tmp1;
float anrm;
float sigma, d__1;
float rmin, rmax;
magmaFloat_ptr dwork;
/* Function Body */
lower = (uplo == MagmaLower);
wantz = (jobz == MagmaVec);
alleig = (range == MagmaRangeAll);
valeig = (range == MagmaRangeV);
indeig = (range == MagmaRangeI);
lquery = (lwork == -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 (lddz < 1 || (wantz && lddz < n)) {
*info = -15;
} else if (ldwa < max(1,n)) {
*info = -17;
} else if (ldwz < 1 || (wantz && ldwz < n)) {
*info = -19;
} else {
if (valeig) {
if (n > 0 && vu <= vl) {
*info = -8;
}
} else if (indeig) {
if (il < 1 || il > max(1,n)) {
*info = -9;
} else if (iu < min(n,il) || iu > n) {
*info = -10;
}
}
}
magma_int_t nb = magma_get_chetrd_nb(n);
lopt = n * (nb + 1);
work[0] = MAGMA_C_MAKE( lopt, 0 );
if (lwork < lopt && ! lquery) {
*info = -21;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info));
return *info;
} else if (lquery) {
return *info;
}
*m = 0;
/* Check if matrix is very small then just call LAPACK on CPU, no need for GPU */
if (n <= 128) {
#ifdef ENABLE_DEBUG
printf("--------------------------------------------------------------\n");
printf(" warning matrix too small N=%d NB=%d, calling lapack on CPU \n", (int) n, (int) nb);
printf("--------------------------------------------------------------\n");
#endif
magmaFloatComplex *a;
magma_cmalloc_cpu( &a, n*n );
magma_cgetmatrix(n, n, dA, ldda, a, n);
lapackf77_cheevx(jobz_, range_, uplo_,
&n, a, &n, &vl, &vu, &il, &iu, &abstol, m,
w, wZ, &ldwz, work, &lwork,
rwork, iwork, ifail, info);
magma_csetmatrix( n, n, a, n, dA, ldda);
magma_csetmatrix( n, *m, wZ, ldwz, dZ, lddz);
magma_free_cpu(a);
return *info;
}
if (MAGMA_SUCCESS != magma_smalloc( &dwork, n )) {
fprintf (stderr, "!!!! device memory allocation error (magma_cheevx_gpu)\n");
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
--w;
--work;
--rwork;
--iwork;
--ifail;
/* Get machine constants. */
safmin = lapackf77_slamch("Safe minimum");
eps = lapackf77_slamch("Precision");
smlnum = safmin / eps;
bignum = 1. / smlnum;
rmin = magma_ssqrt(smlnum);
rmax = magma_ssqrt(bignum);
/* Scale matrix to allowable range, if necessary. */
anrm = magmablas_clanhe(MagmaMaxNorm, uplo, n, dA, ldda, dwork);
iscale = 0;
sigma = 1;
if (anrm > 0. && anrm < rmin) {
iscale = 1;
sigma = rmin / anrm;
} else if (anrm > rmax) {
iscale = 1;
sigma = rmax / anrm;
}
if (iscale == 1) {
d__1 = 1.;
magmablas_clascl(uplo, 0, 0, 1., sigma, n, n, dA,
ldda, info);
if (abstol > 0.) {
abstol *= sigma;
}
if (valeig) {
vl *= sigma;
vu *= sigma;
}
}
/* Call CHETRD to reduce Hermitian matrix to tridiagonal form. */
indd = 1;
inde = indd + n;
indrwk = inde + n;
indtau = 1;
indwrk = indtau + n;
llwork = lwork - indwrk + 1;
#ifdef FAST_HEMV
magma_chetrd2_gpu(uplo, n, dA, ldda, &rwork[indd], &rwork[inde],
&work[indtau], wA, ldwa, &work[indwrk], llwork, dZ, lddz*n, &iinfo);
#else
magma_chetrd_gpu (uplo, n, dA, ldda, &rwork[indd], &rwork[inde],
&work[indtau], wA, ldwa, &work[indwrk], llwork, &iinfo);
#endif
lopt = n + (magma_int_t)MAGMA_C_REAL(work[indwrk]);
/* If all eigenvalues are desired and ABSTOL is less than or equal to
zero, then call SSTERF or CUNGTR and CSTEQR. If this fails for
some eigenvalue, then try SSTEBZ. */
if ((alleig || (indeig && il == 1 && iu == n)) && abstol <= 0.) {
blasf77_scopy(&n, &rwork[indd], &ione, &w[1], &ione);
indee = indrwk + 2*n;
if (! wantz) {
i__1 = n - 1;
blasf77_scopy(&i__1, &rwork[inde], &ione, &rwork[indee], &ione);
lapackf77_ssterf(&n, &w[1], &rwork[indee], info);
}
else {
lapackf77_clacpy("A", &n, &n, wA, &ldwa, wZ, &ldwz);
lapackf77_cungtr(uplo_, &n, wZ, &ldwz, &work[indtau], &work[indwrk], &llwork, &iinfo);
i__1 = n - 1;
blasf77_scopy(&i__1, &rwork[inde], &ione, &rwork[indee], &ione);
lapackf77_csteqr(jobz_, &n, &w[1], &rwork[indee], wZ, &ldwz, &rwork[indrwk], info);
if (*info == 0) {
for (i = 1; i <= n; ++i) {
ifail[i] = 0;
}
magma_csetmatrix( n, n, wZ, ldwz, dZ, lddz );
}
}
if (*info == 0) {
*m = n;
}
}
/* Otherwise, call SSTEBZ and, if eigenvectors are desired, CSTEIN. */
if (*m == 0) {
*info = 0;
if (wantz) {
order_ = "B";
} else {
order_ = "E";
}
indibl = 1;
indisp = indibl + n;
indiwk = indisp + n;
lapackf77_sstebz(range_, order_, &n, &vl, &vu, &il, &iu, &abstol, &rwork[indd], &rwork[inde], m,
&nsplit, &w[1], &iwork[indibl], &iwork[indisp], &rwork[indrwk], &iwork[indiwk], info);
if (wantz) {
lapackf77_cstein(&n, &rwork[indd], &rwork[inde], m, &w[1], &iwork[indibl], &iwork[indisp],
wZ, &ldwz, &rwork[indrwk], &iwork[indiwk], &ifail[1], info);
magma_csetmatrix( n, *m, wZ, ldwz, dZ, lddz );
/* Apply unitary matrix used in reduction to tridiagonal
form to eigenvectors returned by CSTEIN. */
magma_cunmtr_gpu(MagmaLeft, uplo, MagmaNoTrans, n, *m, dA, ldda, &work[indtau],
dZ, lddz, wA, ldwa, &iinfo);
}
}
/* If matrix was scaled, then rescale eigenvalues appropriately. */
if (iscale == 1) {
if (*info == 0) {
imax = *m;
} else {
imax = *info - 1;
}
d__1 = 1. / sigma;
blasf77_sscal(&imax, &d__1, &w[1], &ione);
}
/* If eigenvalues are not in order, then sort them, along with
eigenvectors. */
if (wantz) {
for (j = 1; j <= *m-1; ++j) {
i = 0;
tmp1 = w[j];
for (jj = j + 1; jj <= *m; ++jj) {
if (w[jj] < tmp1) {
i = jj;
tmp1 = w[jj];
}
}
if (i != 0) {
itmp1 = iwork[indibl + i - 1];
w[i] = w[j];
iwork[indibl + i - 1] = iwork[indibl + j - 1];
w[j] = tmp1;
iwork[indibl + j - 1] = itmp1;
magma_cswap(n, dZ + (i-1)*lddz, ione, dZ + (j-1)*lddz, ione);
if (*info != 0) {
itmp1 = ifail[i];
ifail[i] = ifail[j];
ifail[j] = itmp1;
}
}
}
}
/* Set WORK[0] to optimal complex workspace size. */
work[1] = MAGMA_C_MAKE( lopt, 0 );
return *info;
} /* magma_cheevx_gpu */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgemm
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, dev_perf, dev_time, cpu_perf, cpu_time;
float magma_error, dev_error, Cnorm, work[1];
magma_int_t M, N, K;
magma_int_t Am, An, Bm, Bn;
magma_int_t sizeA, sizeB, sizeC;
magma_int_t lda, ldb, ldc, ldda, lddb, lddc;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magmaFloatComplex *h_A, *h_B, *h_C, *h_Cmagma, *h_Cdev;
magmaFloatComplex_ptr d_A, d_B, d_C;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex alpha = MAGMA_C_MAKE( 0.29, -0.86 );
magmaFloatComplex beta = MAGMA_C_MAKE( -0.48, 0.38 );
magma_opts opts;
opts.parse_opts( argc, argv );
float tol = opts.tolerance * lapackf77_slamch("E");
#ifdef HAVE_CUBLAS
// for CUDA, we can check MAGMA vs. CUBLAS, without running LAPACK
printf("%% If running lapack (option --lapack), MAGMA and %s error are both computed\n"
"%% relative to CPU BLAS result. Else, MAGMA error is computed relative to %s result.\n\n",
g_platform_str, g_platform_str );
printf("%% transA = %s, transB = %s\n",
lapack_trans_const(opts.transA),
lapack_trans_const(opts.transB) );
printf("%% M N K MAGMA Gflop/s (ms) %s Gflop/s (ms) CPU Gflop/s (ms) MAGMA error %s error\n",
g_platform_str, g_platform_str );
#else
// for others, we need LAPACK for check
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
printf("%% transA = %s, transB = %s\n",
lapack_trans_const(opts.transA),
lapack_trans_const(opts.transB) );
printf("%% M N K %s Gflop/s (ms) CPU Gflop/s (ms) %s error\n",
g_platform_str, g_platform_str );
#endif
printf("%%========================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
K = opts.ksize[itest];
gflops = FLOPS_CGEMM( M, N, K ) / 1e9;
if ( opts.transA == MagmaNoTrans ) {
lda = Am = M;
An = K;
} else {
lda = Am = K;
An = M;
}
if ( opts.transB == MagmaNoTrans ) {
ldb = Bm = K;
Bn = N;
} else {
ldb = Bm = N;
Bn = K;
}
ldc = M;
ldda = magma_roundup( lda, opts.align ); // multiple of 32 by default
lddb = magma_roundup( ldb, opts.align ); // multiple of 32 by default
lddc = magma_roundup( ldc, opts.align ); // multiple of 32 by default
sizeA = lda*An;
sizeB = ldb*Bn;
sizeC = ldc*N;
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, lda*An );
TESTING_MALLOC_CPU( h_B, magmaFloatComplex, ldb*Bn );
TESTING_MALLOC_CPU( h_C, magmaFloatComplex, ldc*N );
TESTING_MALLOC_CPU( h_Cmagma, magmaFloatComplex, ldc*N );
TESTING_MALLOC_CPU( h_Cdev, magmaFloatComplex, ldc*N );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*An );
TESTING_MALLOC_DEV( d_B, magmaFloatComplex, lddb*Bn );
TESTING_MALLOC_DEV( d_C, magmaFloatComplex, lddc*N );
/* Initialize the matrices */
lapackf77_clarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_clarnv( &ione, ISEED, &sizeB, h_B );
lapackf77_clarnv( &ione, ISEED, &sizeC, h_C );
magma_csetmatrix( Am, An, h_A, lda, d_A, ldda, opts.queue );
magma_csetmatrix( Bm, Bn, h_B, ldb, d_B, lddb, opts.queue );
/* =====================================================================
Performs operation using MAGMABLAS (currently only with CUDA)
=================================================================== */
#ifdef HAVE_CUBLAS
magma_csetmatrix( M, N, h_C, ldc, d_C, lddc, opts.queue );
magma_time = magma_sync_wtime( opts.queue );
magmablas_cgemm( opts.transA, opts.transB, M, N, K,
alpha, d_A, ldda,
d_B, lddb,
beta, d_C, lddc,
opts.queue );
magma_time = magma_sync_wtime( opts.queue ) - magma_time;
magma_perf = gflops / magma_time;
magma_cgetmatrix( M, N, d_C, lddc, h_Cmagma, ldc, opts.queue );
#endif
/* =====================================================================
Performs operation using CUBLAS / clBLAS / Xeon Phi MKL
=================================================================== */
magma_csetmatrix( M, N, h_C, ldc, d_C, lddc, opts.queue );
dev_time = magma_sync_wtime( opts.queue );
#ifdef HAVE_CUBLAS
// opts.handle also uses opts.queue
cublasCgemm( opts.handle, cublas_trans_const(opts.transA), cublas_trans_const(opts.transB), M, N, K,
&alpha, d_A, ldda,
d_B, lddb,
&beta, d_C, lddc );
#else
magma_cgemm( opts.transA, opts.transB, M, N, K,
alpha, d_A, 0, ldda,
d_B, 0, lddb,
beta, d_C, 0, lddc, opts.queue );
#endif
dev_time = magma_sync_wtime( opts.queue ) - dev_time;
dev_perf = gflops / dev_time;
magma_cgetmatrix( M, N, d_C, lddc, h_Cdev, ldc, opts.queue );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
blasf77_cgemm( lapack_trans_const(opts.transA), lapack_trans_const(opts.transB), &M, &N, &K,
&alpha, h_A, &lda,
h_B, &ldb,
&beta, h_C, &ldc );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
}
/* =====================================================================
Check the result
=================================================================== */
if ( opts.lapack ) {
// compute relative error for both magma & dev, relative to lapack,
// |C_magma - C_lapack| / |C_lapack|
Cnorm = lapackf77_clange( "F", &M, &N, h_C, &ldc, work );
blasf77_caxpy( &sizeC, &c_neg_one, h_C, &ione, h_Cdev, &ione );
dev_error = lapackf77_clange( "F", &M, &N, h_Cdev, &ldc, work ) / Cnorm;
#ifdef HAVE_CUBLAS
blasf77_caxpy( &sizeC, &c_neg_one, h_C, &ione, h_Cmagma, &ione );
magma_error = lapackf77_clange( "F", &M, &N, h_Cmagma, &ldc, work ) / Cnorm;
printf("%5d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %s\n",
(int) M, (int) N, (int) K,
magma_perf, 1000.*magma_time,
dev_perf, 1000.*dev_time,
cpu_perf, 1000.*cpu_time,
magma_error, dev_error,
(magma_error < tol && dev_error < tol ? "ok" : "failed"));
status += ! (magma_error < tol && dev_error < tol);
#else
printf("%5d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) M, (int) N, (int) K,
dev_perf, 1000.*dev_time,
cpu_perf, 1000.*cpu_time,
dev_error,
(dev_error < tol ? "ok" : "failed"));
status += ! (dev_error < tol);
#endif
}
else {
#ifdef HAVE_CUBLAS
// compute relative error for magma, relative to dev (currently only with CUDA)
Cnorm = lapackf77_clange( "F", &M, &N, h_Cdev, &ldc, work );
blasf77_caxpy( &sizeC, &c_neg_one, h_Cdev, &ione, h_Cmagma, &ione );
magma_error = lapackf77_clange( "F", &M, &N, h_Cmagma, &ldc, work ) / Cnorm;
printf("%5d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) --- ( --- ) %8.2e --- %s\n",
(int) M, (int) N, (int) K,
magma_perf, 1000.*magma_time,
dev_perf, 1000.*dev_time,
magma_error,
(magma_error < tol ? "ok" : "failed"));
status += ! (magma_error < tol);
#else
printf("%5d %5d %5d %7.2f (%7.2f) --- ( --- ) ---\n",
(int) M, (int) N, (int) K,
dev_perf, 1000.*dev_time );
#endif
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_C );
TESTING_FREE_CPU( h_Cmagma );
TESTING_FREE_CPU( h_Cdev );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
TESTING_FREE_DEV( d_C );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/**
Purpose
-------
CHEEVDX_GPU computes selected eigenvalues and, optionally, eigenvectors
of a complex Hermitian matrix A. Eigenvalues and eigenvectors can
be selected by specifying either a range of values or a range of
indices for the desired eigenvalues.
If eigenvectors are desired, it uses a divide and conquer algorithm.
The divide and conquer algorithm makes very mild assumptions about
floating point arithmetic. It will work on machines with a guard
digit in add/subtract, or on those binary machines without guard
digits which subtract like the Cray X-MP, Cray Y-MP, Cray C-90, or
Cray-2. It could conceivably fail on hexadecimal or decimal machines
without guard digits, but we know of none.
Arguments
---------
@param[in]
jobz magma_vec_t
- = MagmaNoVec: Compute eigenvalues only;
- = MagmaVec: Compute eigenvalues and eigenvectors.
@param[in]
range magma_range_t
- = MagmaRangeAll: all eigenvalues will be found.
- = MagmaRangeV: all eigenvalues in the half-open interval (VL,VU]
will be found.
- = MagmaRangeI: the IL-th through IU-th eigenvalues will be found.
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored;
- = MagmaLower: Lower triangle of A is stored.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
dA COMPLEX array on the GPU,
dimension (LDDA, N).
On entry, the Hermitian matrix A. If UPLO = MagmaUpper, the
leading N-by-N upper triangular part of A contains the
upper triangular part of the matrix A. If UPLO = MagmaLower,
the leading N-by-N lower triangular part of A contains
the lower triangular part of the matrix A.
On exit, if JOBZ = MagmaVec, then if INFO = 0, the first m columns
of A contains the required
orthonormal eigenvectors of the matrix A.
If JOBZ = MagmaNoVec, then on exit the lower triangle (if UPLO=MagmaLower)
or the upper triangle (if UPLO=MagmaUpper) of A, including the
diagonal, is destroyed.
@param[in]
ldda INTEGER
The leading dimension of the array DA. LDDA >= max(1,N).
@param[in]
vl REAL
@param[in]
vu REAL
If RANGE=MagmaRangeV, the lower and upper bounds of the interval to
be searched for eigenvalues. VL < VU.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeI.
@param[in]
il INTEGER
@param[in]
iu INTEGER
If RANGE=MagmaRangeI, the indices (in ascending order) of the
smallest and largest eigenvalues to be returned.
1 <= IL <= IU <= N, if N > 0; IL = 1 and IU = 0 if N = 0.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeV.
@param[out]
m INTEGER
The total number of eigenvalues found. 0 <= M <= N.
If RANGE = MagmaRangeAll, M = N, and if RANGE = MagmaRangeI, M = IU-IL+1.
@param[out]
w REAL array, dimension (N)
If INFO = 0, the required m eigenvalues in ascending order.
@param
wA (workspace) COMPLEX array, dimension (LDWA, N)
@param[in]
ldwa INTEGER
The leading dimension of the array wA. LDWA >= max(1,N).
@param[out]
work (workspace) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The length of the array WORK.
If N <= 1, LWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LWORK >= N + N*NB.
If JOBZ = MagmaVec and N > 1, LWORK >= max( N + N*NB, 2*N + N**2 ).
NB can be obtained through magma_get_chetrd_nb(N).
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal sizes of the WORK, RWORK and
IWORK arrays, returns these values as the first entries of
the WORK, RWORK and IWORK arrays, and no error message
related to LWORK or LRWORK or LIWORK is issued by XERBLA.
@param[out]
rwork (workspace) REAL array, dimension (LRWORK)
On exit, if INFO = 0, RWORK[0] returns the optimal LRWORK.
@param[in]
lrwork INTEGER
The dimension of the array RWORK.
If N <= 1, LRWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LRWORK >= N.
If JOBZ = MagmaVec and N > 1, LRWORK >= 1 + 5*N + 2*N**2.
\n
If LRWORK = -1, then a workspace query is assumed; the
routine only calculates the optimal sizes of the WORK, RWORK
and IWORK arrays, returns these values as the first entries
of the WORK, RWORK and IWORK arrays, and no error message
related to LWORK or LRWORK or LIWORK is issued by XERBLA.
@param[out]
iwork (workspace) INTEGER array, dimension (MAX(1,LIWORK))
On exit, if INFO = 0, IWORK[0] returns the optimal LIWORK.
@param[in]
liwork INTEGER
The dimension of the array IWORK.
If N <= 1, LIWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LIWORK >= 1.
If JOBZ = MagmaVec and N > 1, LIWORK >= 3 + 5*N.
\n
If LIWORK = -1, then a workspace query is assumed; the
routine only calculates the optimal sizes of the WORK, RWORK
and IWORK arrays, returns these values as the first entries
of the WORK, RWORK and IWORK arrays, and no error message
related to LWORK or LRWORK or LIWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: if INFO = i and JOBZ = MagmaNoVec, then the algorithm failed
to converge; i off-diagonal elements of an intermediate
tridiagonal form did not converge to zero;
if INFO = i and JOBZ = MagmaVec, then the algorithm failed
to compute an eigenvalue while working on the submatrix
lying in rows and columns INFO/(N+1) through
mod(INFO,N+1).
Further Details
---------------
Based on contributions by
Jeff Rutter, Computer Science Division, University of California
at Berkeley, USA
Modified description of INFO. Sven, 16 Feb 05.
@ingroup magma_cheev_driver
********************************************************************/
extern "C" magma_int_t
magma_cheevdx_gpu(magma_vec_t jobz, magma_range_t range, magma_uplo_t uplo,
magma_int_t n,
magmaFloatComplex *dA, magma_int_t ldda,
float vl, float vu, magma_int_t il, magma_int_t iu,
magma_int_t *m, float *w,
magmaFloatComplex *wA, magma_int_t ldwa,
magmaFloatComplex *work, magma_int_t lwork,
float *rwork, magma_int_t lrwork,
magma_int_t *iwork, magma_int_t liwork,
magma_int_t *info)
{
const char* uplo_ = lapack_uplo_const( uplo );
const char* jobz_ = lapack_vec_const( jobz );
magma_int_t ione = 1;
float d__1;
float eps;
magma_int_t inde;
float anrm;
magma_int_t imax;
float rmin, rmax;
float sigma;
magma_int_t iinfo, lwmin;
magma_int_t lower;
magma_int_t llrwk;
magma_int_t wantz;
magma_int_t indwk2, llwrk2;
magma_int_t iscale;
float safmin;
float bignum;
magma_int_t indtau;
magma_int_t indrwk, indwrk, liwmin;
magma_int_t lrwmin, llwork;
float smlnum;
magma_int_t lquery;
magma_int_t alleig, valeig, indeig;
float *dwork;
magmaFloatComplex *dC;
magma_int_t lddc = ldda;
wantz = (jobz == MagmaVec);
lower = (uplo == MagmaLower);
alleig = (range == MagmaRangeAll);
valeig = (range == MagmaRangeV);
indeig = (range == MagmaRangeI);
lquery = (lwork == -1 || lrwork == -1 || liwork == -1);
*info = 0;
if (! (wantz || (jobz == MagmaNoVec))) {
*info = -1;
} else if (! (alleig || valeig || indeig)) {
*info = -2;
} else if (! (lower || (uplo == MagmaUpper))) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (ldda < max(1,n)) {
*info = -6;
} else if (ldwa < max(1,n)) {
*info = -14;
} else {
if (valeig) {
if (n > 0 && vu <= vl) {
*info = -8;
}
} else if (indeig) {
if (il < 1 || il > max(1,n)) {
*info = -9;
} else if (iu < min(n,il) || iu > n) {
*info = -10;
}
}
}
magma_int_t nb = magma_get_chetrd_nb( n );
if ( n <= 1 ) {
lwmin = 1;
lrwmin = 1;
liwmin = 1;
}
else if ( wantz ) {
lwmin = max( n + n*nb, 2*n + n*n );
lrwmin = 1 + 5*n + 2*n*n;
liwmin = 3 + 5*n;
}
else {
lwmin = n + n*nb;
lrwmin = n;
liwmin = 1;
}
// multiply by 1+eps (in Double!) to ensure length gets rounded up,
// if it cannot be exactly represented in floating point.
real_Double_t one_eps = 1. + lapackf77_slamch("Epsilon");
work[0] = MAGMA_C_MAKE( lwmin * one_eps, 0.);
rwork[0] = lrwmin * one_eps;
iwork[0] = liwmin;
if ((lwork < lwmin) && !lquery) {
*info = -16;
} else if ((lrwork < lrwmin) && ! lquery) {
*info = -18;
} else if ((liwork < liwmin) && ! lquery) {
*info = -20;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info));
return *info;
}
else if (lquery) {
return *info;
}
/* Check if matrix is very small then just call LAPACK on CPU, no need for GPU */
if (n <= 128) {
#ifdef ENABLE_DEBUG
printf("--------------------------------------------------------------\n");
printf(" warning matrix too small N=%d NB=%d, calling lapack on CPU \n", (int) n, (int) nb);
printf("--------------------------------------------------------------\n");
#endif
magmaFloatComplex *A;
magma_cmalloc_cpu( &A, n*n );
magma_cgetmatrix(n, n, dA, ldda, A, n);
lapackf77_cheevd(jobz_, uplo_,
&n, A, &n,
w, work, &lwork,
rwork, &lrwork,
iwork, &liwork, info);
magma_csetmatrix( n, n, A, n, dA, ldda);
magma_free_cpu(A);
*m=n;
return *info;
}
magma_queue_t stream;
magma_queue_create( &stream );
// dC and dwork are never used together, so use one buffer for both;
// unfortunately they're different types (complex and float).
// (this works better in dsyevd_gpu where they're both float).
// n*lddc for chetrd2_gpu, *2 for complex
// n for clanhe
magma_int_t ldwork = n*lddc*2;
if ( wantz ) {
// need 3n^2/2 for cstedx
ldwork = max( ldwork, 3*n*(n/2 + 1) );
}
if (MAGMA_SUCCESS != magma_smalloc( &dwork, ldwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dC = (magmaFloatComplex*) dwork;
/* Get machine constants. */
safmin = lapackf77_slamch("Safe minimum");
eps = lapackf77_slamch("Precision");
smlnum = safmin / eps;
bignum = 1. / smlnum;
rmin = magma_ssqrt(smlnum);
rmax = magma_ssqrt(bignum);
/* Scale matrix to allowable range, if necessary. */
anrm = magmablas_clanhe(MagmaMaxNorm, uplo, n, dA, ldda, dwork);
iscale = 0;
sigma = 1;
if (anrm > 0. && anrm < rmin) {
iscale = 1;
sigma = rmin / anrm;
} else if (anrm > rmax) {
iscale = 1;
sigma = rmax / anrm;
}
if (iscale == 1) {
magmablas_clascl(uplo, 0, 0, 1., sigma, n, n, dA, ldda, info);
}
/* Call CHETRD to reduce Hermitian matrix to tridiagonal form. */
// chetrd rwork: e (n)
// cstedx rwork: e (n) + llrwk (1 + 4*N + 2*N**2) ==> 1 + 5n + 2n^2
inde = 0;
indrwk = inde + n;
llrwk = lrwork - indrwk;
// chetrd work: tau (n) + llwork (n*nb) ==> n + n*nb
// cstedx work: tau (n) + z (n^2)
// cunmtr work: tau (n) + z (n^2) + llwrk2 (n or n*nb) ==> 2n + n^2, or n + n*nb + n^2
indtau = 0;
indwrk = indtau + n;
indwk2 = indwrk + n*n;
llwork = lwork - indwrk;
llwrk2 = lwork - indwk2;
magma_timer_t time=0;
timer_start( time );
#ifdef FAST_HEMV
magma_chetrd2_gpu(uplo, n, dA, ldda, w, &rwork[inde],
&work[indtau], wA, ldwa, &work[indwrk], llwork,
dC, n*lddc, &iinfo);
#else
magma_chetrd_gpu (uplo, n, dA, ldda, w, &rwork[inde],
&work[indtau], wA, ldwa, &work[indwrk], llwork, &iinfo);
#endif
timer_stop( time );
timer_printf( "time chetrd_gpu = %6.2f\n", time );
/* For eigenvalues only, call SSTERF. For eigenvectors, first call
CSTEDC to generate the eigenvector matrix, WORK(INDWRK), of the
tridiagonal matrix, then call CUNMTR to multiply it to the Householder
transformations represented as Householder vectors in A. */
if (! wantz) {
lapackf77_ssterf(&n, w, &rwork[inde], info);
magma_smove_eig(range, n, w, &il, &iu, vl, vu, m);
}
else {
timer_start( time );
magma_cstedx(range, n, vl, vu, il, iu, w, &rwork[inde],
&work[indwrk], n, &rwork[indrwk],
llrwk, iwork, liwork, dwork, info);
timer_stop( time );
timer_printf( "time cstedx = %6.2f\n", time );
timer_start( time );
magma_smove_eig(range, n, w, &il, &iu, vl, vu, m);
magma_csetmatrix( n, *m, &work[indwrk + n * (il-1) ], n, dC, lddc );
magma_cunmtr_gpu(MagmaLeft, uplo, MagmaNoTrans, n, *m, dA, ldda, &work[indtau],
dC, lddc, wA, ldwa, &iinfo);
magma_ccopymatrix( n, *m, dC, lddc, dA, ldda );
timer_stop( time );
timer_printf( "time cunmtr_gpu + copy = %6.2f\n", time );
}
/* If matrix was scaled, then rescale eigenvalues appropriately. */
if (iscale == 1) {
if (*info == 0) {
imax = n;
} else {
imax = *info - 1;
}
d__1 = 1. / sigma;
blasf77_sscal(&imax, &d__1, w, &ione);
}
work[0] = MAGMA_C_MAKE( lwmin * one_eps, 0.); // round up
rwork[0] = lrwmin * one_eps;
iwork[0] = liwmin;
magma_queue_destroy( stream );
magma_free( dwork );
return *info;
} /* magma_cheevdx_gpu */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing ctranspose
Code is very similar to testing_csymmetrize.cpp
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gbytes, gpu_perf, gpu_time, gpu_perf2=0, gpu_time2=0, cpu_perf, cpu_time;
float error, error2, work[1];
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *h_A, *h_B, *h_R;
magmaFloatComplex *d_A, *d_B;
magma_int_t M, N, size, lda, ldda, ldb, lddb;
magma_int_t ione = 1;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
printf("Inplace transpose requires M==N.\n");
printf(" M N CPU GByte/s (ms) GPU GByte/s (ms) check Inplace GB/s (ms) check\n");
printf("====================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
lda = M;
ldda = ((M+31)/32)*32;
ldb = N;
lddb = ((N+31)/32)*32;
// load entire matrix, save entire matrix
gbytes = sizeof(magmaFloatComplex) * 2.*M*N / 1e9;
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, lda*N ); // input: M x N
TESTING_MALLOC_CPU( h_B, magmaFloatComplex, ldb*M ); // output: N x M
TESTING_MALLOC_CPU( h_R, magmaFloatComplex, ldb*M ); // output: N x M
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N ); // input: M x N
TESTING_MALLOC_DEV( d_B, magmaFloatComplex, lddb*M ); // output: N x M
/* Initialize the matrix */
for( int j = 0; j < N; ++j ) {
for( int i = 0; i < M; ++i ) {
h_A[i + j*lda] = MAGMA_C_MAKE( i + j/10000., j );
}
}
for( int j = 0; j < M; ++j ) {
for( int i = 0; i < N; ++i ) {
h_B[i + j*ldb] = MAGMA_C_MAKE( i + j/10000., j );
}
}
magma_csetmatrix( N, M, h_B, ldb, d_B, lddb );
/* =====================================================================
Performs operation using naive out-of-place algorithm
(LAPACK doesn't implement transpose)
=================================================================== */
cpu_time = magma_wtime();
//for( int j = 1; j < N-1; ++j ) { // inset by 1 row & col
// for( int i = 1; i < M-1; ++i ) { // inset by 1 row & col
for( int j = 0; j < N; ++j ) {
for( int i = 0; i < M; ++i ) {
h_B[j + i*ldb] = h_A[i + j*lda];
}
}
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gbytes / cpu_time;
/* ====================================================================
Performs operation using MAGMA, out-of-place
=================================================================== */
magma_csetmatrix( M, N, h_A, lda, d_A, ldda );
magma_csetmatrix( N, M, h_B, ldb, d_B, lddb );
gpu_time = magma_sync_wtime( 0 );
//magmablas_ctranspose( M-2, N-2, d_A+1+ldda, ldda, d_B+1+lddb, lddb ); // inset by 1 row & col
magmablas_ctranspose( M, N, d_A, ldda, d_B, lddb );
gpu_time = magma_sync_wtime( 0 ) - gpu_time;
gpu_perf = gbytes / gpu_time;
/* ====================================================================
Performs operation using MAGMA, in-place
=================================================================== */
if ( M == N ) {
magma_csetmatrix( M, N, h_A, lda, d_A, ldda );
gpu_time2 = magma_sync_wtime( 0 );
//magmablas_ctranspose_inplace( N-2, d_A+1+ldda, ldda ); // inset by 1 row & col
magmablas_ctranspose_inplace( N, d_A, ldda );
gpu_time2 = magma_sync_wtime( 0 ) - gpu_time2;
gpu_perf2 = gbytes / gpu_time2;
}
/* =====================================================================
Check the result
=================================================================== */
// check out-of-place transpose (d_B)
size = ldb*M;
magma_cgetmatrix( N, M, d_B, lddb, h_R, ldb );
blasf77_caxpy( &size, &c_neg_one, h_B, &ione, h_R, &ione );
error = lapackf77_clange("f", &N, &M, h_R, &ldb, work );
if ( M == N ) {
// also check in-place tranpose (d_A)
magma_cgetmatrix( N, M, d_A, ldda, h_R, ldb );
blasf77_caxpy( &size, &c_neg_one, h_B, &ione, h_R, &ione );
error2 = lapackf77_clange("f", &N, &M, h_R, &ldb, work );
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %6s %7.2f (%7.2f) %s\n",
(int) M, (int) N,
cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000.,
(error == 0. ? "ok" : "failed"),
gpu_perf2, gpu_time2,
(error2 == 0. ? "ok" : "failed") );
status += ! (error == 0. && error2 == 0.);
}
else {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %6s --- ( --- )\n",
(int) M, (int) N,
cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000.,
(error == 0. ? "ok" : "failed") );
status += ! (error == 0.);
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
