magma_int_t
magma_dvspread(
magma_d_matrix *x,
const char * filename,
magma_queue_t queue )
{
magma_int_t info = 0;
magma_d_matrix A={Magma_CSR}, B={Magma_CSR};
magma_int_t entry=0;
// char *vfilename[] = {"/mnt/sparse_matrices/mtx/rail_79841_B.mtx"};
CHECK( magma_d_csr_mtx( &A, filename, queue ));
CHECK( magma_dmconvert( A, &B, Magma_CSR, Magma_DENSE, queue ));
CHECK( magma_dvinit( x, Magma_CPU, A.num_cols, A.num_rows, MAGMA_D_ZERO, queue ));
x->major = MagmaRowMajor;
for(magma_int_t i=0; i<A.num_cols; i++) {
for(magma_int_t j=0; j<A.num_rows; j++) {
x->val[i*A.num_rows+j] = B.val[ i+j*A.num_cols ];
entry++;
}
}
x->num_rows = A.num_rows;
x->num_cols = A.num_cols;
cleanup:
magma_dmfree( &A, queue );
magma_dmfree( &B, queue );
return info;
}
extern "C" magma_int_t
magma_dtfqmr_unrolled(
magma_d_matrix A, magma_d_matrix b, magma_d_matrix *x,
magma_d_solver_par *solver_par,
magma_queue_t queue )
{
magma_int_t info = MAGMA_NOTCONVERGED;
// prepare solver feedback
solver_par->solver = Magma_TFQMR;
solver_par->numiter = 0;
solver_par->spmv_count = 0;
solver_par->spmv_count = 0;
// local variables
double c_zero = MAGMA_D_ZERO, c_one = MAGMA_D_ONE;
// solver variables
double nom0, r0, res, nomb;
double rho = c_one, rho_l = c_one, eta = c_zero , c = c_zero ,
theta = c_zero , tau = c_zero, alpha = c_one, beta = c_zero,
sigma = c_zero;
magma_int_t dofs = A.num_rows* b.num_cols;
// GPU workspace
magma_d_matrix r={Magma_CSR}, r_tld={Magma_CSR},
d={Magma_CSR}, w={Magma_CSR}, v={Magma_CSR},
u_mp1={Magma_CSR}, u_m={Magma_CSR}, Au={Magma_CSR},
Ad={Magma_CSR}, Au_new={Magma_CSR};
CHECK( magma_dvinit( &r, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &u_mp1,Magma_DEV, A.num_rows, b.num_cols, c_one, queue ));
CHECK( magma_dvinit( &r_tld,Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &u_m, Magma_DEV, A.num_rows, b.num_cols, c_one, queue ));
CHECK( magma_dvinit( &v, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &d, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &w, Magma_DEV, A.num_rows, b.num_cols, c_one, queue ));
CHECK( magma_dvinit( &Ad, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &Au_new, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &Au, Magma_DEV, A.num_rows, b.num_cols, c_one, queue ));
// solver setup
CHECK( magma_dresidualvec( A, b, *x, &r, &nom0, queue));
solver_par->init_res = nom0;
magma_dcopy( dofs, r.dval, 1, r_tld.dval, 1, queue );
magma_dcopy( dofs, r.dval, 1, w.dval, 1, queue );
magma_dcopy( dofs, r.dval, 1, u_mp1.dval, 1, queue );
CHECK( magma_d_spmv( c_one, A, u_mp1, c_zero, v, queue )); // v = A u
magma_dcopy( dofs, v.dval, 1, Au.dval, 1, queue );
nomb = magma_dnrm2( dofs, 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;
}
tau = magma_dsqrt( magma_ddot( dofs, r.dval, 1, r_tld.dval, 1, queue ) );
rho = magma_ddot( dofs, r.dval, 1, r_tld.dval, 1, queue );
rho_l = rho;
//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++;
// do this every iteration as unrolled
alpha = rho / magma_ddot( dofs, v.dval, 1, r_tld.dval, 1, queue );
sigma = theta * theta / alpha * eta;
magma_daxpy( dofs, -alpha, v.dval, 1, u_mp1.dval, 1, queue ); // u_mp1 = u_mp_1 - alpha*v;
magma_daxpy( dofs, -alpha, Au.dval, 1, w.dval, 1, queue ); // w = w - alpha*Au;
magma_dscal( dofs, sigma, d.dval, 1, queue );
magma_daxpy( dofs, c_one, u_mp1.dval, 1, d.dval, 1, queue ); // d = u_mp1 + sigma*d;
//magma_dscal( dofs, sigma, Ad.dval, 1, queue );
//magma_daxpy( dofs, c_one, Au.dval, 1, Ad.dval, 1, queue ); // Ad = Au + sigma*Ad;
theta = magma_dsqrt( magma_ddot(dofs, w.dval, 1, w.dval, 1, queue ) ) / tau;
c = c_one / magma_dsqrt( c_one + theta*theta );
tau = tau * theta *c;
eta = c * c * alpha;
sigma = theta * theta / alpha * eta;
printf("sigma: %f+%fi\n", MAGMA_D_REAL(sigma), MAGMA_D_IMAG(sigma) );
CHECK( magma_d_spmv( c_one, A, d, c_zero, Ad, queue )); // Au_new = A u_mp1
solver_par->spmv_count++;
magma_daxpy( dofs, eta, d.dval, 1, x->dval, 1, queue ); // x = x + eta * d
magma_daxpy( dofs, -eta, Ad.dval, 1, r.dval, 1, queue ); // r = r - eta * Ad
// here starts the second part of the loop #################################
magma_daxpy( dofs, -alpha, Au.dval, 1, w.dval, 1, queue ); // w = w - alpha*Au;
magma_dscal( dofs, sigma, d.dval, 1, queue );
magma_daxpy( dofs, c_one, u_mp1.dval, 1, d.dval, 1, queue ); // d = u_mp1 + sigma*d;
magma_dscal( dofs, sigma, Ad.dval, 1, queue );
magma_daxpy( dofs, c_one, Au.dval, 1, Ad.dval, 1, queue ); // Ad = Au + sigma*Ad;
theta = magma_dsqrt( magma_ddot(dofs, w.dval, 1, w.dval, 1, queue ) ) / tau;
c = c_one / magma_dsqrt( c_one + theta*theta );
tau = tau * theta *c;
eta = c * c * alpha;
magma_daxpy( dofs, eta, d.dval, 1, x->dval, 1, queue ); // x = x + eta * d
magma_daxpy( dofs, -eta, Ad.dval, 1, r.dval, 1, queue ); // r = r - eta * Ad
res = magma_dnrm2( dofs, r.dval, 1, queue );
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/nomb <= solver_par->rtol || res <= solver_par->atol ){
break;
}
// do this every loop as unrolled
rho_l = rho;
rho = magma_ddot( dofs, w.dval, 1, r_tld.dval, 1, queue );
beta = rho / rho_l;
magma_dscal( dofs, beta, u_mp1.dval, 1, queue );
magma_daxpy( dofs, c_one, w.dval, 1, u_mp1.dval, 1, queue ); // u_mp1 = w + beta*u_mp1;
CHECK( magma_d_spmv( c_one, A, u_mp1, c_zero, Au_new, queue )); // Au_new = A u_mp1
solver_par->spmv_count++;
// do this every loop as unrolled
magma_dscal( dofs, beta*beta, v.dval, 1, queue );
magma_daxpy( dofs, beta, Au.dval, 1, v.dval, 1, queue );
magma_daxpy( dofs, c_one, Au_new.dval, 1, v.dval, 1, queue ); // v = Au_new + beta*(Au+beta*v);
magma_dcopy( dofs, Au_new.dval, 1, Au.dval, 1, queue );
}
while ( solver_par->numiter+1 <= solver_par->maxiter );
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t) tempo2-tempo1;
double residual;
CHECK( magma_dresidualvec( 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;
} 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) 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 == 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;
}
}
info = MAGMA_DIVERGENCE;
}
cleanup:
magma_dmfree(&r, queue );
magma_dmfree(&r_tld, queue );
magma_dmfree(&d, queue );
magma_dmfree(&w, queue );
magma_dmfree(&v, queue );
magma_dmfree(&u_m, queue );
magma_dmfree(&u_mp1, queue );
magma_dmfree(&d, queue );
magma_dmfree(&Au, queue );
magma_dmfree(&Au_new, queue );
magma_dmfree(&Ad, queue );
solver_par->info = info;
return info;
} /* magma_dfqmr_unrolled */
extern "C" magma_int_t
magma_dpbicgstab(
magma_d_matrix A, magma_d_matrix b, magma_d_matrix *x,
magma_d_solver_par *solver_par,
magma_d_preconditioner *precond_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_PBICGSTAB;
solver_par->numiter = 0;
solver_par->info = MAGMA_SUCCESS;
// some useful variables
double c_zero = MAGMA_D_ZERO, c_one = MAGMA_D_ONE,
c_mone = MAGMA_D_NEG_ONE;
magma_int_t dofs = A.num_rows*b.num_cols;
// workspace
magma_d_matrix r={Magma_CSR}, rr={Magma_CSR}, p={Magma_CSR}, v={Magma_CSR}, s={Magma_CSR}, t={Magma_CSR}, ms={Magma_CSR}, mt={Magma_CSR}, y={Magma_CSR}, z={Magma_CSR};
CHECK( magma_dvinit( &r, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &rr,Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &p, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &v, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &s, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &t, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &ms,Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &mt,Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &y, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &z, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
// solver variables
double alpha, beta, omega, rho_old, rho_new;
double nom, betanom, nom0, r0, den, res;
// solver setup
CHECK( magma_dresidualvec( A, b, *x, &r, &nom0, queue));
magma_dcopy( dofs, r.dval, 1, rr.dval, 1 ); // rr = r
betanom = nom0;
nom = nom0*nom0;
rho_new = omega = alpha = MAGMA_D_MAKE( 1.0, 0. );
solver_par->init_res = nom0;
CHECK( magma_d_spmv( c_one, A, r, c_zero, v, queue )); // z = A r
den = MAGMA_D_REAL( magma_ddot(dofs, v.dval, 1, r.dval, 1) ); // den = z' * r
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;
}
//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;
}
solver_par->numiter = 0;
// start iteration
do
{
solver_par->numiter++;
rho_old = rho_new; // rho_old=rho
rho_new = magma_ddot( dofs, rr.dval, 1, r.dval, 1 ); // rho=<rr,r>
beta = rho_new/rho_old * alpha/omega; // beta=rho/rho_old *alpha/omega
magma_dscal( dofs, beta, p.dval, 1 ); // p = beta*p
magma_daxpy( dofs, c_mone * omega * beta, v.dval, 1 , p.dval, 1 );
// p = p-omega*beta*v
magma_daxpy( dofs, c_one, r.dval, 1, p.dval, 1 ); // p = p+r
// preconditioner
CHECK( magma_d_applyprecond_left( A, p, &mt, precond_par, queue ));
CHECK( magma_d_applyprecond_right( A, mt, &y, precond_par, queue ));
CHECK( magma_d_spmv( c_one, A, y, c_zero, v, queue )); // v = Ap
alpha = rho_new / magma_ddot( dofs, rr.dval, 1, v.dval, 1 );
magma_dcopy( dofs, r.dval, 1 , s.dval, 1 ); // s=r
magma_daxpy( dofs, c_mone * alpha, v.dval, 1 , s.dval, 1 ); // s=s-alpha*v
// preconditioner
CHECK( magma_d_applyprecond_left( A, s, &ms, precond_par, queue ));
CHECK( magma_d_applyprecond_right( A, ms, &z, precond_par, queue ));
CHECK( magma_d_spmv( c_one, A, z, c_zero, t, queue )); // t=As
// preconditioner
CHECK( magma_d_applyprecond_left( A, s, &ms, precond_par, queue ));
CHECK( magma_d_applyprecond_left( A, t, &mt, precond_par, queue ));
// omega = <ms,mt>/<mt,mt>
omega = magma_ddot( dofs, mt.dval, 1, ms.dval, 1 )
/ magma_ddot( dofs, mt.dval, 1, mt.dval, 1 );
magma_daxpy( dofs, alpha, y.dval, 1 , x->dval, 1 ); // x=x+alpha*p
magma_daxpy( dofs, omega, z.dval, 1 , x->dval, 1 ); // x=x+omega*s
magma_dcopy( dofs, s.dval, 1 , r.dval, 1 ); // r=s
magma_daxpy( dofs, c_mone * omega, t.dval, 1 , r.dval, 1 ); // r=r-omega*t
res = betanom = magma_dnrm2( dofs, r.dval, 1 );
nom = betanom*betanom;
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;
}
}
while ( solver_par->numiter+1 <= solver_par->maxiter );
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t) tempo2-tempo1;
double residual;
CHECK( magma_dresidualvec( A, b, *x, &r, &residual, queue));
solver_par->final_res = residual;
solver_par->iter_res = res;
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->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_dmfree(&r, queue );
magma_dmfree(&rr, queue );
magma_dmfree(&p, queue );
magma_dmfree(&v, queue );
magma_dmfree(&s, queue );
magma_dmfree(&t, queue );
magma_dmfree(&ms, queue );
magma_dmfree(&mt, queue );
magma_dmfree(&y, queue );
magma_dmfree(&z, queue );
magmablasSetKernelStream( orig_queue );
solver_par->info = info;
return info;
} /* magma_dbicgstab */
extern "C" magma_int_t
magma_dqmr_merge(
magma_d_matrix A, magma_d_matrix b, magma_d_matrix *x,
magma_d_solver_par *solver_par,
magma_queue_t queue )
{
magma_int_t info = MAGMA_NOTCONVERGED;
// prepare solver feedback
solver_par->solver = Magma_QMRMERGE;
solver_par->numiter = 0;
solver_par->spmv_count = 0;
// local variables
double c_zero = MAGMA_D_ZERO, c_one = MAGMA_D_ONE;
// solver variables
double nom0, r0, res=0, nomb;
double rho = c_one, rho1 = c_one, eta = -c_one , pds = c_one,
thet = c_one, thet1 = c_one, epsilon = c_one,
beta = c_one, delta = c_one, pde = c_one, rde = c_one,
gamm = c_one, gamm1 = c_one, psi = c_one;
magma_int_t dofs = A.num_rows* b.num_cols;
// need to transpose the matrix
magma_d_matrix AT={Magma_CSR}, Ah1={Magma_CSR}, Ah2={Magma_CSR};
// GPU workspace
magma_d_matrix r={Magma_CSR}, r_tld={Magma_CSR},
v={Magma_CSR}, w={Magma_CSR}, wt={Magma_CSR},
d={Magma_CSR}, s={Magma_CSR}, z={Magma_CSR}, q={Magma_CSR},
p={Magma_CSR}, pt={Magma_CSR}, y={Magma_CSR};
CHECK( magma_dvinit( &r, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &r_tld, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &v, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &w, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &wt,Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &d, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &s, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &z, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &q, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &p, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &pt,Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &y, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
// solver setup
CHECK( magma_dresidualvec( A, b, *x, &r, &nom0, queue));
solver_par->init_res = nom0;
magma_dcopy( dofs, r.dval, 1, r_tld.dval, 1, queue );
magma_dcopy( dofs, r.dval, 1, y.dval, 1, queue );
magma_dcopy( dofs, r.dval, 1, v.dval, 1, queue );
magma_dcopy( dofs, r.dval, 1, wt.dval, 1, queue );
magma_dcopy( dofs, r.dval, 1, z.dval, 1, queue );
// transpose the matrix
magma_dmtransfer( A, &Ah1, Magma_DEV, Magma_CPU, queue );
magma_dmconvert( Ah1, &Ah2, A.storage_type, Magma_CSR, queue );
magma_dmfree(&Ah1, queue );
magma_dmtransposeconjugate( Ah2, &Ah1, queue );
magma_dmfree(&Ah2, queue );
Ah2.blocksize = A.blocksize;
Ah2.alignment = A.alignment;
magma_dmconvert( Ah1, &Ah2, Magma_CSR, A.storage_type, queue );
magma_dmfree(&Ah1, queue );
magma_dmtransfer( Ah2, &AT, Magma_CPU, Magma_DEV, queue );
magma_dmfree(&Ah2, queue );
nomb = magma_dnrm2( dofs, 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;
}
psi = magma_dsqrt( magma_ddot( dofs, z.dval, 1, z.dval, 1, queue ));
rho = magma_dsqrt( magma_ddot( dofs, y.dval, 1, y.dval, 1, queue ));
// v = y / rho
// y = y / rho
// w = wt / psi
// z = z / psi
magma_dqmr_1(
r.num_rows,
r.num_cols,
rho,
psi,
y.dval,
z.dval,
v.dval,
w.dval,
queue );
//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++;
if( magma_d_isnan_inf( rho ) || magma_d_isnan_inf( psi ) ){
info = MAGMA_DIVERGENCE;
break;
}
// delta = z' * y;
delta = magma_ddot( dofs, z.dval, 1, y.dval, 1, queue );
if( magma_d_isnan_inf( delta ) ){
info = MAGMA_DIVERGENCE;
break;
}
// no precond: yt = y, zt = z
//magma_dcopy( dofs, y.dval, 1, yt.dval, 1 );
//magma_dcopy( dofs, z.dval, 1, zt.dval, 1 );
if( solver_par->numiter == 1 ){
// p = y;
// q = z;
magma_dcopy( dofs, y.dval, 1, p.dval, 1, queue );
magma_dcopy( dofs, z.dval, 1, q.dval, 1, queue );
}
else{
pde = psi * delta / epsilon;
rde = rho * MAGMA_D_CONJ(delta/epsilon);
// p = y - pde * p
// q = z - rde * q
magma_dqmr_2(
r.num_rows,
r.num_cols,
pde,
rde,
y.dval,
z.dval,
p.dval,
q.dval,
queue );
}
if( magma_d_isnan_inf( rho ) || magma_d_isnan_inf( psi ) ){
info = MAGMA_DIVERGENCE;
break;
}
CHECK( magma_d_spmv( c_one, A, p, c_zero, pt, queue ));
solver_par->spmv_count++;
// epsilon = q' * pt;
epsilon = magma_ddot( dofs, q.dval, 1, pt.dval, 1, queue );
beta = epsilon / delta;
if( magma_d_isnan_inf( epsilon ) || magma_d_isnan_inf( beta ) ){
info = MAGMA_DIVERGENCE;
break;
}
// v = pt - beta * v
// y = v
magma_dqmr_3(
r.num_rows,
r.num_cols,
beta,
pt.dval,
v.dval,
y.dval,
queue );
rho1 = rho;
// rho = norm(y);
rho = magma_dsqrt( magma_ddot( dofs, y.dval, 1, y.dval, 1, queue ));
// wt = A' * q - beta' * w;
CHECK( magma_d_spmv( c_one, AT, q, c_zero, wt, queue ));
solver_par->spmv_count++;
magma_daxpy( dofs, - MAGMA_D_CONJ( beta ), w.dval, 1, wt.dval, 1, queue );
// no precond: z = wt
magma_dcopy( dofs, wt.dval, 1, z.dval, 1, queue );
thet1 = thet;
thet = rho / (gamm * MAGMA_D_MAKE( MAGMA_D_ABS(beta), 0.0 ));
gamm1 = gamm;
gamm = c_one / magma_dsqrt(c_one + thet*thet);
eta = - eta * rho1 * gamm * gamm / (beta * gamm1 * gamm1);
if( magma_d_isnan_inf( thet ) || magma_d_isnan_inf( gamm ) || magma_d_isnan_inf( eta ) ){
info = MAGMA_DIVERGENCE;
break;
}
if( solver_par->numiter == 1 ){
// d = eta * p + pds * d;
// s = eta * pt + pds * d;
// x = x + d;
// r = r - s;
magma_dqmr_4(
r.num_rows,
r.num_cols,
eta,
p.dval,
pt.dval,
d.dval,
s.dval,
x->dval,
r.dval,
queue );
}
else{
pds = (thet1 * gamm) * (thet1 * gamm);
// d = eta * p + pds * d;
// s = eta * pt + pds * d;
// x = x + d;
// r = r - s;
magma_dqmr_5(
r.num_rows,
r.num_cols,
eta,
pds,
p.dval,
pt.dval,
d.dval,
s.dval,
x->dval,
r.dval,
queue );
}
// psi = norm(z);
psi = magma_dsqrt( magma_ddot( dofs, z.dval, 1, z.dval, 1, queue ) );
res = magma_dnrm2( dofs, r.dval, 1, queue );
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;
}
}
// v = y / rho
// y = y / rho
// w = wt / psi
// z = z / psi
magma_dqmr_1(
r.num_rows,
r.num_cols,
rho,
psi,
y.dval,
z.dval,
v.dval,
w.dval,
queue );
if ( res/nomb <= solver_par->rtol || res <= solver_par->atol ){
break;
}
}
while ( solver_par->numiter+1 <= solver_par->maxiter );
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t) tempo2-tempo1;
double residual;
CHECK( magma_dresidualvec( 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_dmfree(&r, queue );
magma_dmfree(&r_tld, queue );
magma_dmfree(&v, queue );
magma_dmfree(&w, queue );
magma_dmfree(&wt, queue );
magma_dmfree(&d, queue );
magma_dmfree(&s, queue );
magma_dmfree(&z, queue );
magma_dmfree(&q, queue );
magma_dmfree(&p, queue );
magma_dmfree(&pt, queue );
magma_dmfree(&y, queue );
magma_dmfree(&AT, queue );
magma_dmfree(&Ah1, queue );
magma_dmfree(&Ah2, queue );
solver_par->info = info;
return info;
} /* magma_dqmr_merge */
extern "C" magma_int_t
magma_dcg(
magma_d_matrix A, magma_d_matrix b, magma_d_matrix *x,
magma_d_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
double c_zero = MAGMA_D_ZERO, c_one = MAGMA_D_ONE;
magma_int_t dofs = A.num_rows * b.num_cols;
// GPU workspace
magma_d_matrix r={Magma_CSR}, p={Magma_CSR}, q={Magma_CSR};
CHECK( magma_dvinit( &r, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &p, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &q, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
// solver variables
double alpha, beta;
double nom, nom0, r0, betanom, betanomsq, den;
// solver setup
CHECK( magma_dresidualvec( A, b, *x, &r, &nom0, queue));
magma_dcopy( dofs, r.dval, 1, p.dval, 1 ); // p = r
betanom = nom0;
nom = nom0 * nom0; // nom = r' * r
CHECK( magma_d_spmv( c_one, A, p, c_zero, q, queue )); // q = A p
den = MAGMA_D_REAL( magma_ddot(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_D_MAKE(nom/den, 0.);
magma_daxpy(dofs, alpha, p.dval, 1, x->dval, 1); // x = x + alpha p
magma_daxpy(dofs, -alpha, q.dval, 1, r.dval, 1); // r = r - alpha q
betanom = magma_dnrm2(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_D_MAKE(betanomsq/nom, 0.); // beta = betanoms/nom
magma_dscal(dofs, beta, p.dval, 1); // p = beta*p
magma_daxpy(dofs, c_one, r.dval, 1, p.dval, 1); // p = p + r
CHECK( magma_d_spmv( c_one, A, p, c_zero, q, queue )); // q = A p
den = MAGMA_D_REAL(magma_ddot(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;
double residual;
CHECK( magma_dresidualvec( 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_dmfree(&r, queue );
magma_dmfree(&p, queue );
magma_dmfree(&q, queue );
magmablasSetKernelStream( orig_queue );
solver_par->info = info;
return info;
} /* magma_dcg */
extern "C" magma_int_t
magma_diterref(
magma_d_matrix A, magma_d_matrix b, magma_d_matrix *x,
magma_d_solver_par *solver_par, magma_d_preconditioner *precond_par,
magma_queue_t queue )
{
magma_int_t info = MAGMA_NOTCONVERGED;
// some useful variables
double c_zero = MAGMA_D_ZERO;
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
// prepare solver feedback
solver_par->solver = Magma_ITERREF;
solver_par->numiter = 0;
solver_par->spmv_count = 0;
magma_int_t dofs = A.num_rows*b.num_cols;
// solver variables
double nom, nom0;
// workspace
magma_d_matrix r={Magma_CSR}, z={Magma_CSR};
CHECK( magma_dvinit( &r, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &z, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
double residual;
CHECK( magma_dresidual( A, b, *x, &residual, queue ));
solver_par->init_res = residual;
// solver setup
magma_dscal( dofs, c_zero, x->dval, 1, queue ); // x = 0
//CHECK( magma_dresidualvec( A, b, *x, &r, nom, queue));
magma_dcopy( dofs, b.dval, 1, r.dval, 1, queue ); // r = b
nom0 = magma_dnrm2( dofs, r.dval, 1, queue ); // nom0 = || r ||
nom = nom0 * nom0;
solver_par->init_res = nom0;
if( nom0 < solver_par->atol ||
nom0/solver_par->init_res < solver_par->rtol ){
solver_par->final_res = solver_par->init_res;
solver_par->iter_res = solver_par->init_res;
info = MAGMA_SUCCESS;
goto cleanup;
}
//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++ ) {
magma_dscal( dofs, MAGMA_D_MAKE(1./nom, 0.), r.dval, 1, queue ); // scale it
CHECK( magma_d_precond( A, r, &z, precond_par, queue )); // inner solver: A * z = r
magma_dscal( dofs, MAGMA_D_MAKE(nom, 0.), z.dval, 1, queue ); // scale it
magma_daxpy( dofs, c_one, z.dval, 1, x->dval, 1, queue ); // x = x + z
CHECK( magma_d_spmv( c_neg_one, A, *x, c_zero, r, queue )); // r = - A x
solver_par->spmv_count++;
magma_daxpy( dofs, c_one, b.dval, 1, r.dval, 1, queue ); // r = r + b
nom = magma_dnrm2( dofs, r.dval, 1, queue ); // nom = || 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) nom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
if( nom < solver_par->atol ||
nom/solver_par->init_res < solver_par->rtol ){
break;
}
}
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t) tempo2-tempo1;
CHECK( magma_dresidualvec( A, b, *x, &r, &residual, queue));
solver_par->final_res = residual;
solver_par->iter_res = nom;
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) nom;
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) nom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
info = MAGMA_DIVERGENCE;
}
cleanup:
magma_dmfree(&r, queue );
magma_dmfree(&z, queue );
solver_par->info = info;
return info;
} /* magma_diterref */
extern "C" magma_int_t
magma_dcg_merge(
magma_d_matrix A, magma_d_matrix b, magma_d_matrix *x,
magma_d_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
double alpha, beta, gamma, rho, tmp1, *skp_h={0};
double nom, nom0, betanom, den, nomb;
// some useful variables
double c_zero = MAGMA_D_ZERO, c_one = MAGMA_D_ONE;
magma_int_t dofs = A.num_rows*b.num_cols;
magma_d_matrix r={Magma_CSR}, d={Magma_CSR}, z={Magma_CSR}, B={Magma_CSR}, C={Magma_CSR};
double *d1=NULL, *d2=NULL, *skp=NULL;
// GPU workspace
CHECK( magma_dvinit( &r, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &d, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &z, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dmalloc( &d1, dofs*(1) ));
CHECK( magma_dmalloc( &d2, dofs*(1) ));
// array for the parameters
CHECK( magma_dmalloc( &skp, 6 ));
// skp = [alpha|beta|gamma|rho|tmp1|tmp2]
// solver setup
magma_dscal( dofs, c_zero, x->dval, 1, queue ); // x = 0
//CHECK( magma_dresidualvec( A, b, *x, &r, nom0, queue));
magma_dcopy( dofs, b.dval, 1, r.dval, 1, queue ); // r = b
magma_dcopy( dofs, r.dval, 1, d.dval, 1, queue ); // d = r
nom0 = betanom = magma_dnrm2( dofs, r.dval, 1, queue );
nom = nom0 * nom0; // nom = r' * r
CHECK( magma_d_spmv( c_one, A, d, c_zero, z, queue )); // z = A d
den = MAGMA_D_ABS( magma_ddot( dofs, d.dval, 1, z.dval, 1, queue ) ); // den = d'* z
solver_par->init_res = nom0;
nomb = magma_dnrm2( dofs, b.dval, 1, queue );
if ( nomb == 0.0 ){
nomb=1.0;
}
// array on host for the parameters
CHECK( magma_dmalloc_cpu( &skp_h, 6 ));
alpha = rho = gamma = tmp1 = c_one;
beta = magma_ddot( 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_D_MAKE(nom, 0.0);
magma_dsetvector( 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_dcgmerge_spmv1( A, d1, d2, d.dval, z.dval, skp, queue ));
solver_par->spmv_count++;
// updates x, r, computes scalars and updates d
CHECK( magma_dcgmerge_xrbeta( dofs, d1, d2, x->dval, r.dval, d.dval, z.dval, skp, queue ));
// check stopping criterion (asynchronous copy)
magma_dgetvector( 1 , skp+1, 1, skp_h+1, 1, queue );
betanom = sqrt(MAGMA_D_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;
double residual;
CHECK( magma_dresidualvec( 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_dmfree(&r, queue );
magma_dmfree(&z, queue );
magma_dmfree(&d, queue );
magma_dmfree(&B, queue );
magma_dmfree(&C, queue );
magma_free( d1 );
magma_free( d2 );
magma_free( skp );
magma_free_cpu( skp_h );
solver_par->info = info;
return info;
} /* magma_dcg_merge */
extern "C" magma_int_t
magma_dbpcg(
magma_d_matrix A, magma_d_matrix b, magma_d_matrix *x,
magma_d_solver_par *solver_par,
magma_d_preconditioner *precond_par,
magma_queue_t queue )
{
magma_int_t info = 0;
magma_int_t i, num_vecs = b.num_rows/A.num_rows;
// prepare solver feedback
solver_par->solver = Magma_PCG;
solver_par->numiter = 0;
solver_par->spmv_count = 0;
solver_par->info = MAGMA_SUCCESS;
// local variables
double c_zero = MAGMA_D_ZERO, c_one = MAGMA_D_ONE;
magma_int_t dofs = A.num_rows;
// GPU workspace
magma_d_matrix r={Magma_CSR}, rt={Magma_CSR}, p={Magma_CSR}, q={Magma_CSR}, h={Magma_CSR};
// solver variables
double *alpha={0}, *beta={0};
alpha = NULL;
beta = NULL;
double *nom={0}, *nom0={0}, *r0={0}, *gammaold={0}, *gammanew={0}, *den={0}, *res={0}, *residual={0};
nom = NULL;
nom0 = NULL;
r0 = NULL;
gammaold = NULL;
gammanew = NULL;
den = NULL;
res = NULL;
residual = NULL;
CHECK( magma_dmalloc_cpu(&alpha, num_vecs));
CHECK( magma_dmalloc_cpu(&beta, num_vecs));
CHECK( magma_dmalloc_cpu(&residual, num_vecs));
CHECK( magma_dmalloc_cpu(&nom, num_vecs));
CHECK( magma_dmalloc_cpu(&nom0, num_vecs));
CHECK( magma_dmalloc_cpu(&r0, num_vecs));
CHECK( magma_dmalloc_cpu(&gammaold, num_vecs));
CHECK( magma_dmalloc_cpu(&gammanew, num_vecs));
CHECK( magma_dmalloc_cpu(&den, num_vecs));
CHECK( magma_dmalloc_cpu(&res, num_vecs));
CHECK( magma_dmalloc_cpu(&residual, num_vecs));
CHECK( magma_dvinit( &r, Magma_DEV, dofs*num_vecs, 1, c_zero, queue ));
CHECK( magma_dvinit( &rt, Magma_DEV, dofs*num_vecs, 1, c_zero, queue ));
CHECK( magma_dvinit( &p, Magma_DEV, dofs*num_vecs, 1, c_zero, queue ));
CHECK( magma_dvinit( &q, Magma_DEV, dofs*num_vecs, 1, c_zero, queue ));
CHECK( magma_dvinit( &h, Magma_DEV, dofs*num_vecs, 1, c_zero, queue ));
// solver setup
CHECK( magma_dresidualvec( A, b, *x, &r, nom0, queue));
// preconditioner
CHECK( magma_d_applyprecond_left( MagmaNoTrans, A, r, &rt, precond_par, queue ));
CHECK( magma_d_applyprecond_right( MagmaNoTrans, A, rt, &h, precond_par, queue ));
magma_dcopy( dofs*num_vecs, h.dval, 1, p.dval, 1, queue ); // p = h
for( i=0; i<num_vecs; i++) {
nom[i] = MAGMA_D_REAL( magma_ddot( dofs, r(i), 1, h(i), 1, queue ) );
nom0[i] = magma_dnrm2( dofs, r(i), 1, queue );
}
CHECK( magma_d_spmv( c_one, A, p, c_zero, q, queue )); // q = A p
for( i=0; i<num_vecs; i++)
den[i] = MAGMA_D_REAL( magma_ddot( dofs, p(i), 1, q(i), 1, queue ) ); // den = p dot q
solver_par->init_res = nom0[0];
if ( (r0[0] = nom[0] * solver_par->rtol) < ATOLERANCE )
r0[0] = ATOLERANCE;
// check positive definite
if (den[0] <= 0.0) {
printf("Operator A is not postive definite. (Ar,r) = %f\n", den[0]);
info = MAGMA_NONSPD;
goto cleanup;
}
if ( nom[0] < r0[0] ) {
solver_par->final_res = solver_par->init_res;
solver_par->iter_res = solver_par->init_res;
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[0];
solver_par->timing[0] = 0.0;
}
solver_par->numiter = 0;
solver_par->spmv_count = 0;
// start iteration
do
{
solver_par->numiter++;
// preconditioner
CHECK( magma_d_applyprecond_left( MagmaNoTrans, A, r, &rt, precond_par, queue ));
CHECK( magma_d_applyprecond_right( MagmaNoTrans, A, rt, &h, precond_par, queue ));
for( i=0; i<num_vecs; i++)
gammanew[i] = MAGMA_D_REAL( magma_ddot( dofs, r(i), 1, h(i), 1, queue ) ); // gn = < r,h>
if ( solver_par->numiter==1 ) {
magma_dcopy( dofs*num_vecs, h.dval, 1, p.dval, 1, queue ); // p = h
} else {
for( i=0; i<num_vecs; i++) {
beta[i] = MAGMA_D_MAKE(gammanew[i]/gammaold[i], 0.); // beta = gn/go
magma_dscal( dofs, beta[i], p(i), 1, queue ); // p = beta*p
magma_daxpy( dofs, c_one, h(i), 1, p(i), 1, queue ); // p = p + h
}
}
CHECK( magma_d_spmv( c_one, A, p, c_zero, q, queue )); // q = A p
solver_par->spmv_count++;
// magma_d_bspmv_tuned( dofs, num_vecs, c_one, A, p.dval, c_zero, q.dval, queue );
for( i=0; i<num_vecs; i++) {
den[i] = MAGMA_D_REAL(magma_ddot( dofs, p(i), 1, q(i), 1, queue) );
// den = p dot q
alpha[i] = MAGMA_D_MAKE(gammanew[i]/den[i], 0.);
magma_daxpy( dofs, alpha[i], p(i), 1, x->dval+dofs*i, 1, queue ); // x = x + alpha p
magma_daxpy( dofs, -alpha[i], q(i), 1, r(i), 1, queue ); // r = r - alpha q
gammaold[i] = gammanew[i];
res[i] = magma_dnrm2( dofs, r(i), 1, queue );
}
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[0];
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
if ( res[0]/nom0[0] < 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;
CHECK( magma_dresidual( A, b, *x, residual, queue ));
solver_par->iter_res = res[0];
solver_par->final_res = residual[0];
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) res[0];
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 ){
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) res[0];
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
info = MAGMA_DIVERGENCE;
}
for( i=0; i<num_vecs; i++) {
printf("%.4e ",res[i]);
}
printf("\n");
for( i=0; i<num_vecs; i++) {
printf("%.4e ",residual[i]);
}
printf("\n");
cleanup:
magma_dmfree(&r, queue );
magma_dmfree(&rt, queue );
magma_dmfree(&p, queue );
magma_dmfree(&q, queue );
magma_dmfree(&h, queue );
magma_free_cpu(alpha);
magma_free_cpu(beta);
magma_free_cpu(nom);
magma_free_cpu(nom0);
magma_free_cpu(r0);
magma_free_cpu(gammaold);
magma_free_cpu(gammanew);
magma_free_cpu(den);
magma_free_cpu(res);
solver_par->info = info;
return info;
} /* magma_dbpcg */
magma_int_t
magma_dcustomilusetup(
magma_d_matrix A,
magma_d_matrix b,
magma_d_preconditioner *precond,
magma_queue_t queue )
{
magma_int_t info = 0;
cusparseHandle_t cusparseHandle=NULL;
cusparseMatDescr_t descrL=NULL;
cusparseMatDescr_t descrU=NULL;
magma_d_matrix hA={Magma_CSR};
char preconditionermatrix[255];
// first L
snprintf( preconditionermatrix, sizeof(preconditionermatrix),
"precondL.mtx" );
CHECK( magma_d_csr_mtx( &hA, preconditionermatrix , queue) );
CHECK( magma_dmtransfer( hA, &precond->L, Magma_CPU, Magma_DEV , queue ));
// extract the diagonal of L into precond->d
CHECK( magma_djacobisetup_diagscal( precond->L, &precond->d, queue ));
CHECK( magma_dvinit( &precond->work1, Magma_DEV, hA.num_rows, 1, MAGMA_D_ZERO, queue ));
magma_dmfree( &hA, queue );
// now U
snprintf( preconditionermatrix, sizeof(preconditionermatrix),
"precondU.mtx" );
CHECK( magma_d_csr_mtx( &hA, preconditionermatrix , queue) );
CHECK( magma_dmtransfer( hA, &precond->U, Magma_CPU, Magma_DEV , queue ));
// extract the diagonal of U into precond->d2
CHECK( magma_djacobisetup_diagscal( precond->U, &precond->d2, queue ));
CHECK( magma_dvinit( &precond->work2, Magma_DEV, hA.num_rows, 1, MAGMA_D_ZERO, queue ));
// CUSPARSE context //
CHECK_CUSPARSE( cusparseCreate( &cusparseHandle ));
CHECK_CUSPARSE( cusparseCreateMatDescr( &descrL ));
CHECK_CUSPARSE( cusparseSetMatType( descrL, CUSPARSE_MATRIX_TYPE_TRIANGULAR ));
CHECK_CUSPARSE( cusparseSetMatDiagType( descrL, CUSPARSE_DIAG_TYPE_UNIT ));
CHECK_CUSPARSE( cusparseSetMatIndexBase( descrL, CUSPARSE_INDEX_BASE_ZERO ));
CHECK_CUSPARSE( cusparseSetMatFillMode( descrL, CUSPARSE_FILL_MODE_LOWER ));
CHECK_CUSPARSE( cusparseCreateSolveAnalysisInfo( &precond->cuinfoL ));
CHECK_CUSPARSE( cusparseDcsrsv_analysis( cusparseHandle,
CUSPARSE_OPERATION_NON_TRANSPOSE, precond->L.num_rows,
precond->L.nnz, descrL,
precond->L.val, precond->L.row, precond->L.col, precond->cuinfoL ));
CHECK_CUSPARSE( cusparseCreateMatDescr( &descrU ));
CHECK_CUSPARSE( cusparseSetMatType( descrU, CUSPARSE_MATRIX_TYPE_TRIANGULAR ));
CHECK_CUSPARSE( cusparseSetMatDiagType( descrU, CUSPARSE_DIAG_TYPE_NON_UNIT ));
CHECK_CUSPARSE( cusparseSetMatIndexBase( descrU, CUSPARSE_INDEX_BASE_ZERO ));
CHECK_CUSPARSE( cusparseSetMatFillMode( descrU, CUSPARSE_FILL_MODE_UPPER ));
CHECK_CUSPARSE( cusparseCreateSolveAnalysisInfo( &precond->cuinfoU ));
CHECK_CUSPARSE( cusparseDcsrsv_analysis( cusparseHandle,
CUSPARSE_OPERATION_NON_TRANSPOSE, precond->U.num_rows,
precond->U.nnz, descrU,
precond->U.val, precond->U.row, precond->U.col, precond->cuinfoU ));
cleanup:
cusparseDestroy( cusparseHandle );
cusparseDestroyMatDescr( descrL );
cusparseDestroyMatDescr( descrU );
cusparseHandle=NULL;
descrL=NULL;
descrU=NULL;
magma_dmfree( &hA, queue );
return info;
}
/* ////////////////////////////////////////////////////////////////////////////
-- testing zdot
*/
int main( int argc, char** argv )
{
magma_int_t info = 0;
magma_queue_t queue=NULL;
magma_queue_create( 0, &queue );
const double one = MAGMA_D_MAKE(1.0, 0.0);
const double zero = MAGMA_D_MAKE(0.0, 0.0);
double alpha;
TESTING_INIT();
magma_d_matrix a={Magma_CSR}, b={Magma_CSR}, x={Magma_CSR}, y={Magma_CSR}, skp={Magma_CSR};
printf("%%=======================================================================================================================================================================\n");
printf("\n");
printf(" | runtime | GFLOPS\n");
printf("%% n num_vecs | CUDOT CUGEMV MAGMAGEMV MDOT MDGM MDGM_SHFL | CUDOT CUGEMV MAGMAGEMV MDOT MDGM MDGM_SHFL\n");
printf("%%------------------------------------------------------------------------------------------------------------------------------------------------------------------------\n");
printf("\n");
for( magma_int_t num_vecs=1; num_vecs <= 32; num_vecs += 1 ) {
for( magma_int_t n=500000; n < 500001; n += 10000 ) {
int iters = 10;
double computations = (2.* n * iters * num_vecs);
#define ENABLE_TIMER
#ifdef ENABLE_TIMER
real_Double_t mdot1, mdot2, mdgm1, mdgm2, magmagemv1, magmagemv2, cugemv1, cugemv2, cudot1, cudot2;
real_Double_t mdot_time, mdgm_time, mdgmshf_time, magmagemv_time, cugemv_time, cudot_time;
#endif
CHECK( magma_dvinit( &a, Magma_DEV, n, num_vecs, one, queue ));
CHECK( magma_dvinit( &b, Magma_DEV, n, 1, one, queue ));
CHECK( magma_dvinit( &x, Magma_DEV, n, 8, one, queue ));
CHECK( magma_dvinit( &y, Magma_DEV, n, 8, one, queue ));
CHECK( magma_dvinit( &skp, Magma_DEV, 1, num_vecs, zero, queue ));
// warm up
CHECK( magma_dgemvmdot( n, num_vecs, a.dval, b.dval, x.dval, y.dval, skp.dval, queue ));
// CUDOT
#ifdef ENABLE_TIMER
cudot1 = magma_sync_wtime( queue );
#endif
for( int h=0; h < iters; h++) {
for( int l=0; l<num_vecs; l++){
alpha = magma_ddot( n, a.dval+l*a.num_rows, 1, b.dval, 1, queue );
//cudaDeviceSynchronize();
}
//cudaDeviceSynchronize();
}
#ifdef ENABLE_TIMER
cudot2 = magma_sync_wtime( queue );
cudot_time=cudot2-cudot1;
#endif
// CUGeMV
#ifdef ENABLE_TIMER
cugemv1 = magma_sync_wtime( queue );
#endif
for( int h=0; h < iters; h++) {
magma_dgemv( MagmaTrans, n, num_vecs, one, a.dval, n, b.dval, 1, zero, skp.dval, 1, queue );
}
#ifdef ENABLE_TIMER
cugemv2 = magma_sync_wtime( queue );
cugemv_time=cugemv2-cugemv1;
#endif
// MAGMAGeMV
#ifdef ENABLE_TIMER
magmagemv1 = magma_sync_wtime( queue );
#endif
for( int h=0; h < iters; h++) {
magmablas_dgemv( MagmaTrans, n, num_vecs, one, a.dval, n, b.dval, 1, zero, skp.dval, 1, queue );
}
#ifdef ENABLE_TIMER
magmagemv2 = magma_sync_wtime( queue );
magmagemv_time=magmagemv2-magmagemv1;
#endif
// MDOT
#ifdef ENABLE_TIMER
mdot1 = magma_sync_wtime( queue );
#endif
for( int h=0; h < iters; h++) {
for( int c = 0; c<num_vecs/2; c++ ){
CHECK( magma_dmdotc( n, 2, a.dval, b.dval, x.dval, y.dval, skp.dval, queue ));
}
for( int c = 0; c<num_vecs%2; c++ ){
CHECK( magma_dmdotc( n, 1, a.dval, b.dval, x.dval, y.dval, skp.dval, queue ));
}
//h++;
}
#ifdef ENABLE_TIMER
mdot2 = magma_sync_wtime( queue );
mdot_time=mdot2-mdot1;
#endif
// MDGM
#ifdef ENABLE_TIMER
mdgm1 = magma_sync_wtime( queue );
#endif
for( int h=0; h < iters; h++) {
CHECK( magma_dgemvmdot( n, num_vecs, a.dval, b.dval, x.dval, y.dval, skp.dval, queue ));
//h++;
}
#ifdef ENABLE_TIMER
mdgm2 = magma_sync_wtime( queue );
mdgm_time=mdgm2-mdgm1;
#endif
// MDGM_shfl
#ifdef ENABLE_TIMER
mdgm1 = magma_sync_wtime( queue );
#endif
for( int h=0; h < iters; h++) {
CHECK( magma_dgemvmdot_shfl( n, num_vecs, a.dval, b.dval, x.dval, y.dval, skp.dval, queue ));
}
#ifdef ENABLE_TIMER
mdgm2 = magma_sync_wtime( queue );
mdgmshf_time=mdgm2-mdgm1;
#endif
//magma_dprint_gpu(num_vecs,1,skp.dval,num_vecs);
//Chronometry
#ifdef ENABLE_TIMER
printf("%d %d %e %e %e %e %e %e || %e %e %e %e %e %e\n",
int(n), int(num_vecs),
cudot_time/iters,
(cugemv_time)/iters,
(magmagemv_time)/iters,
(mdot_time)/iters,
(mdgm_time)/iters,
(mdgmshf_time)/iters,
computations/(cudot_time*1e9),
computations/(cugemv_time*1e9),
computations/(magmagemv_time*1e9),
computations/(mdot_time*1e9),
computations/(mdgm_time*1e9),
computations/(mdgmshf_time*1e9) );
#endif
magma_dmfree(&a, queue );
magma_dmfree(&b, queue );
magma_dmfree(&x, queue );
magma_dmfree(&y, queue );
magma_dmfree(&skp, queue );
}
//printf("%%================================================================================================================================================\n");
//printf("\n");
//printf("\n");
}
// use alpha to silence compiler warnings
if ( isnan( real( alpha ))) {
info = -1;
}
cleanup:
magma_queue_destroy( queue );
TESTING_FINALIZE();
return info;
}
extern "C" magma_int_t
magma_dlsqr(
magma_d_matrix A, magma_d_matrix b, magma_d_matrix *x,
magma_d_solver_par *solver_par,
magma_d_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
double c_zero = MAGMA_D_ZERO, c_one = MAGMA_D_ONE;
// solver variables
double 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_d_matrix AT={Magma_CSR}, Ah1={Magma_CSR}, Ah2={Magma_CSR};
// GPU workspace
magma_d_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_dvinit( &r, Magma_DEV, A.num_cols, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &v, Magma_DEV, A.num_cols, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &z, Magma_DEV, A.num_cols, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &d, Magma_DEV, A.num_cols, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &vt,Magma_DEV, A.num_cols, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &q, Magma_DEV, A.num_cols, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &w, Magma_DEV, A.num_cols, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &u, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &zt,Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
// transpose the matrix
magma_dmtransfer( A, &Ah1, Magma_DEV, Magma_CPU, queue );
magma_dmconvert( Ah1, &Ah2, A.storage_type, Magma_CSR, queue );
magma_dmfree(&Ah1, queue );
magma_dmtransposeconjugate( Ah2, &Ah1, queue );
magma_dmfree(&Ah2, queue );
Ah2.blocksize = A.blocksize;
Ah2.alignment = A.alignment;
magma_dmconvert( Ah1, &Ah2, Magma_CSR, A.storage_type, queue );
magma_dmfree(&Ah1, queue );
magma_dmtransfer( Ah2, &AT, Magma_CPU, Magma_DEV, queue );
magma_dmfree(&Ah2, queue );
// solver setup
CHECK( magma_dresidualvec( A, b, *x, &r, &nom0, queue));
solver_par->init_res = nom0;
nomb = magma_dnrm2( 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_dcopy( m, b.dval, 1, u.dval, 1, queue );
beta = magma_dnrm2( m, u.dval, 1, queue );
magma_dscal( m, MAGMA_D_MAKE(1./beta, 0.0 ), u.dval, 1, queue );
normr = beta;
c = 1.0;
s = 0.0;
phibar = beta;
CHECK( magma_d_spmv( c_one, AT, u, c_zero, v, queue ));
if( precond_par->solver == Magma_NONE ){
;
} else {
CHECK( magma_d_applyprecond_right( MagmaTrans, A, v, &zt, precond_par, queue ));
CHECK( magma_d_applyprecond_left( MagmaTrans, A, zt, &v, precond_par, queue ));
}
alpha = magma_dnrm2( n, v.dval, 1, queue );
magma_dscal( n, MAGMA_D_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_dcopy( n, v.dval, 1 , z.dval, 1, queue );
} else {
CHECK( magma_d_applyprecond_left( MagmaNoTrans, A, v, &zt, precond_par, queue ));
CHECK( magma_d_applyprecond_right( MagmaNoTrans, A, zt, &z, precond_par, queue ));
}
//CHECK( magma_d_spmv( c_one, A, z, MAGMA_D_MAKE(-alpha,0.0), u, queue ));
CHECK( magma_d_spmv( c_one, A, z, c_zero, zt, queue ));
magma_dscal( m, MAGMA_D_MAKE(-alpha, 0.0 ), u.dval, 1, queue );
magma_daxpy( m, c_one, zt.dval, 1, u.dval, 1, queue );
solver_par->spmv_count++;
beta = magma_dnrm2( m, u.dval, 1, queue );
magma_dscal( m, MAGMA_D_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_dscal( n, MAGMA_D_MAKE(-thet, 0.0 ), d.dval, 1, queue );
magma_daxpy( n, c_one, z.dval, 1, d.dval, 1, queue );
magma_dscal( n, MAGMA_D_MAKE(1./rho, 0.0 ), d.dval, 1, queue );
normd = magma_dnrm2( 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_daxpy( n, MAGMA_D_MAKE( phi, 0.0 ), d.dval, 1, x->dval, 1, queue );
normr = fabs(s) * normr;
CHECK( magma_d_spmv( c_one, AT, u, c_zero, vt, queue ));
solver_par->spmv_count++;
if( precond_par->solver == Magma_NONE ){
;
} else {
CHECK( magma_d_applyprecond_right( MagmaTrans, A, vt, &zt, precond_par, queue ));
CHECK( magma_d_applyprecond_left( MagmaTrans, A, zt, &vt, precond_par, queue ));
}
magma_dscal( n, MAGMA_D_MAKE(-beta, 0.0 ), v.dval, 1, queue );
magma_daxpy( n, c_one, vt.dval, 1, v.dval, 1, queue );
alpha = magma_dnrm2( n, v.dval, 1, queue );
magma_dscal( n, MAGMA_D_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;
double residual;
CHECK( magma_dresidualvec( 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_dmfree(&r, queue );
magma_dmfree(&v, queue );
magma_dmfree(&z, queue );
magma_dmfree(&zt, queue );
magma_dmfree(&d, queue );
magma_dmfree(&vt, queue );
magma_dmfree(&q, queue );
magma_dmfree(&u, queue );
magma_dmfree(&w, queue );
magma_dmfree(&AT, queue );
magma_dmfree(&Ah1, queue );
magma_dmfree(&Ah2, queue );
solver_par->info = info;
return info;
} /* magma_dqmr */
extern "C" magma_int_t
magma_didr(
magma_d_matrix A, magma_d_matrix b, magma_d_matrix *x,
magma_d_solver_par *solver_par,
magma_queue_t queue )
{
magma_int_t info = MAGMA_NOTCONVERGED;
// prepare solver feedback
solver_par->solver = Magma_IDR;
solver_par->numiter = 0;
solver_par->spmv_count = 0;
solver_par->init_res = 0.0;
solver_par->final_res = 0.0;
solver_par->iter_res = 0.0;
solver_par->runtime = 0.0;
// constants
const double c_zero = MAGMA_D_ZERO;
const double c_one = MAGMA_D_ONE;
const double c_n_one = MAGMA_D_NEG_ONE;
// internal user parameters
const magma_int_t smoothing = 1; // 0 = disable, 1 = enable
const double angle = 0.7; // [0-1]
// local variables
magma_int_t iseed[4] = {0, 0, 0, 1};
magma_int_t dof;
magma_int_t s;
magma_int_t distr;
magma_int_t k, i, sk;
magma_int_t innerflag;
double residual;
double nrm;
double nrmb;
double nrmr;
double nrmt;
double rho;
double om;
double tt;
double tr;
double gamma;
double alpha;
double mkk;
double fk;
// matrices and vectors
magma_d_matrix dxs = {Magma_CSR};
magma_d_matrix dr = {Magma_CSR}, drs = {Magma_CSR};
magma_d_matrix dP = {Magma_CSR}, dP1 = {Magma_CSR};
magma_d_matrix dG = {Magma_CSR};
magma_d_matrix dU = {Magma_CSR};
magma_d_matrix dM = {Magma_CSR};
magma_d_matrix df = {Magma_CSR};
magma_d_matrix dt = {Magma_CSR};
magma_d_matrix dc = {Magma_CSR};
magma_d_matrix dv = {Magma_CSR};
magma_d_matrix dbeta = {Magma_CSR}, hbeta = {Magma_CSR};
// chronometry
real_Double_t tempo1, tempo2;
// initial s space
// TODO: add option for 's' (shadow space number)
// Hack: uses '--restart' option as the shadow space number.
// This is not a good idea because the default value of restart option is used to detect
// if the user provided a custom restart. This means that if the default restart value
// is changed then the code will think it was the user (unless the default value is
// also updated in the 'if' statement below.
s = 1;
if ( solver_par->restart != 50 ) {
if ( solver_par->restart > A.num_cols ) {
s = A.num_cols;
} else {
s = solver_par->restart;
}
}
solver_par->restart = s;
// set max iterations
solver_par->maxiter = min( 2 * A.num_cols, solver_par->maxiter );
// check if matrix A is square
if ( A.num_rows != A.num_cols ) {
//printf("Matrix A is not square.\n");
info = MAGMA_ERR_NOT_SUPPORTED;
goto cleanup;
}
// |b|
nrmb = magma_dnrm2( b.num_rows, b.dval, 1, queue );
if ( nrmb == 0.0 ) {
magma_dscal( x->num_rows, MAGMA_D_ZERO, x->dval, 1, queue );
info = MAGMA_SUCCESS;
goto cleanup;
}
// r = b - A x
CHECK( magma_dvinit( &dr, Magma_DEV, b.num_rows, 1, c_zero, queue ));
CHECK( magma_dresidualvec( A, b, *x, &dr, &nrmr, queue ));
// |r|
solver_par->init_res = nrmr;
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)nrmr;
}
// check if initial is guess good enough
if ( nrmr <= solver_par->atol ||
nrmr/nrmb <= solver_par->rtol ) {
info = MAGMA_SUCCESS;
goto cleanup;
}
// P = randn(n, s)
// P = ortho(P)
//---------------------------------------
// P = 0.0
CHECK( magma_dvinit( &dP, Magma_CPU, A.num_cols, s, c_zero, queue ));
// P = randn(n, s)
distr = 3; // 1 = unif (0,1), 2 = unif (-1,1), 3 = normal (0,1)
dof = dP.num_rows * dP.num_cols;
lapackf77_dlarnv( &distr, iseed, &dof, dP.val );
// transfer P to device
CHECK( magma_dmtransfer( dP, &dP1, Magma_CPU, Magma_DEV, queue ));
magma_dmfree( &dP, queue );
// P = ortho(P1)
if ( dP1.num_cols > 1 ) {
// P = magma_dqr(P1), QR factorization
CHECK( magma_dqr( dP1.num_rows, dP1.num_cols, dP1, dP1.ld, &dP, NULL, queue ));
} else {
// P = P1 / |P1|
nrm = magma_dnrm2( dof, dP1.dval, 1, queue );
nrm = 1.0 / nrm;
magma_dscal( dof, nrm, dP1.dval, 1, queue );
CHECK( magma_dmtransfer( dP1, &dP, Magma_DEV, Magma_DEV, queue ));
}
magma_dmfree( &dP1, queue );
//---------------------------------------
// allocate memory for the scalar products
CHECK( magma_dvinit( &hbeta, Magma_CPU, s, 1, c_zero, queue ));
CHECK( magma_dvinit( &dbeta, Magma_DEV, s, 1, c_zero, queue ));
// smoothing enabled
if ( smoothing > 0 ) {
// set smoothing solution vector
CHECK( magma_dmtransfer( *x, &dxs, Magma_DEV, Magma_DEV, queue ));
// set smoothing residual vector
CHECK( magma_dmtransfer( dr, &drs, Magma_DEV, Magma_DEV, queue ));
}
// G(n,s) = 0
CHECK( magma_dvinit( &dG, Magma_DEV, A.num_cols, s, c_zero, queue ));
// U(n,s) = 0
CHECK( magma_dvinit( &dU, Magma_DEV, A.num_cols, s, c_zero, queue ));
// M(s,s) = I
CHECK( magma_dvinit( &dM, Magma_DEV, s, s, c_zero, queue ));
magmablas_dlaset( MagmaFull, s, s, c_zero, c_one, dM.dval, s, queue );
// f = 0
CHECK( magma_dvinit( &df, Magma_DEV, dP.num_cols, 1, c_zero, queue ));
// t = 0
CHECK( magma_dvinit( &dt, Magma_DEV, dr.num_rows, 1, c_zero, queue ));
// c = 0
CHECK( magma_dvinit( &dc, Magma_DEV, dM.num_cols, 1, c_zero, queue ));
// v = 0
CHECK( magma_dvinit( &dv, Magma_DEV, dr.num_rows, 1, c_zero, queue ));
//--------------START TIME---------------
// chronometry
tempo1 = magma_sync_wtime( queue );
if ( solver_par->verbose > 0 ) {
solver_par->timing[0] = 0.0;
}
om = MAGMA_D_ONE;
innerflag = 0;
// start iteration
do
{
solver_par->numiter++;
// new RHS for small systems
// f = P' r
magmablas_dgemv( MagmaConjTrans, dP.num_rows, dP.num_cols, c_one, dP.dval, dP.ld, dr.dval, 1, c_zero, df.dval, 1, queue );
// shadow space loop
for ( k = 0; k < s; ++k ) {
sk = s - k;
// f(k:s) = M(k:s,k:s) c(k:s)
magma_dcopyvector( sk, &df.dval[k], 1, &dc.dval[k], 1, queue );
magma_dtrsv( MagmaLower, MagmaNoTrans, MagmaNonUnit, sk, &dM.dval[k*dM.ld+k], dM.ld, &dc.dval[k], 1, queue );
// v = r - G(:,k:s) c(k:s)
magma_dcopyvector( dr.num_rows, dr.dval, 1, dv.dval, 1, queue );
magmablas_dgemv( MagmaNoTrans, dG.num_rows, sk, c_n_one, &dG.dval[k*dG.ld], dG.ld, &dc.dval[k], 1, c_one, dv.dval, 1, queue );
// U(:,k) = om * v + U(:,k:s) c(k:s)
magmablas_dgemv( MagmaNoTrans, dU.num_rows, sk, c_one, &dU.dval[k*dU.ld], dU.ld, &dc.dval[k], 1, om, dv.dval, 1, queue );
magma_dcopyvector( dU.num_rows, dv.dval, 1, &dU.dval[k*dU.ld], 1, queue );
// G(:,k) = A U(:,k)
CHECK( magma_d_spmv( c_one, A, dv, c_zero, dv, queue ));
solver_par->spmv_count++;
magma_dcopyvector( dG.num_rows, dv.dval, 1, &dG.dval[k*dG.ld], 1, queue );
// bi-orthogonalize the new basis vectors
for ( i = 0; i < k; ++i ) {
// alpha = P(:,i)' G(:,k)
alpha = magma_ddot( dP.num_rows, &dP.dval[i*dP.ld], 1, &dG.dval[k*dG.ld], 1, queue );
// alpha = alpha / M(i,i)
magma_dgetvector( 1, &dM.dval[i*dM.ld+i], 1, &mkk, 1, queue );
alpha = alpha / mkk;
// G(:,k) = G(:,k) - alpha * G(:,i)
magma_daxpy( dG.num_rows, -alpha, &dG.dval[i*dG.ld], 1, &dG.dval[k*dG.ld], 1, queue );
// U(:,k) = U(:,k) - alpha * U(:,i)
magma_daxpy( dU.num_rows, -alpha, &dU.dval[i*dU.ld], 1, &dU.dval[k*dU.ld], 1, queue );
}
// new column of M = P'G, first k-1 entries are zero
// M(k:s,k) = P(:,k:s)' G(:,k)
magmablas_dgemv( MagmaConjTrans, dP.num_rows, sk, c_one, &dP.dval[k*dP.ld], dP.ld, &dG.dval[k*dG.ld], 1, c_zero, &dM.dval[k*dM.ld+k], 1, queue );
// check M(k,k) == 0
magma_dgetvector( 1, &dM.dval[k*dM.ld+k], 1, &mkk, 1, queue );
if ( MAGMA_D_EQUAL(mkk, MAGMA_D_ZERO) ) {
innerflag = 1;
info = MAGMA_DIVERGENCE;
break;
}
// beta = f(k) / M(k,k)
magma_dgetvector( 1, &df.dval[k], 1, &fk, 1, queue );
hbeta.val[k] = fk / mkk;
// check for nan
if ( magma_d_isnan( hbeta.val[k] ) || magma_d_isinf( hbeta.val[k] )) {
innerflag = 1;
info = MAGMA_DIVERGENCE;
break;
}
// r = r - beta * G(:,k)
magma_daxpy( dr.num_rows, -hbeta.val[k], &dG.dval[k*dG.ld], 1, dr.dval, 1, queue );
// smoothing disabled
if ( smoothing <= 0 ) {
// |r|
nrmr = magma_dnrm2( dr.num_rows, dr.dval, 1, queue );
// smoothing enabled
} else {
// x = x + beta * U(:,k)
magma_daxpy( x->num_rows, hbeta.val[k], &dU.dval[k*dU.ld], 1, x->dval, 1, queue );
// smoothing operation
//---------------------------------------
// t = rs - r
magma_dcopyvector( drs.num_rows, drs.dval, 1, dt.dval, 1, queue );
magma_daxpy( dt.num_rows, c_n_one, dr.dval, 1, dt.dval, 1, queue );
// t't
// t'rs
tt = magma_ddot( dt.num_rows, dt.dval, 1, dt.dval, 1, queue );
tr = magma_ddot( dt.num_rows, dt.dval, 1, drs.dval, 1, queue );
// gamma = (t' * rs) / (t' * t)
gamma = tr / tt;
// rs = rs - gamma * (rs - r)
magma_daxpy( drs.num_rows, -gamma, dt.dval, 1, drs.dval, 1, queue );
// xs = xs - gamma * (xs - x)
magma_dcopyvector( dxs.num_rows, dxs.dval, 1, dt.dval, 1, queue );
magma_daxpy( dt.num_rows, c_n_one, x->dval, 1, dt.dval, 1, queue );
magma_daxpy( dxs.num_rows, -gamma, dt.dval, 1, dxs.dval, 1, queue );
// |rs|
nrmr = magma_dnrm2( drs.num_rows, drs.dval, 1, queue );
//---------------------------------------
}
// store current timing and residual
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)nrmr;
solver_par->timing[(solver_par->numiter) / solver_par->verbose]
= (real_Double_t)tempo2 - tempo1;
}
}
// check convergence
if ( nrmr <= solver_par->atol ||
nrmr/nrmb <= solver_par->rtol ) {
s = k + 1; // for the x-update outside the loop
innerflag = 2;
info = MAGMA_SUCCESS;
break;
}
// non-last s iteration
if ( (k + 1) < s ) {
// f(k+1:s) = f(k+1:s) - beta * M(k+1:s,k)
magma_daxpy( sk-1, -hbeta.val[k], &dM.dval[k*dM.ld+(k+1)], 1, &df.dval[k+1], 1, queue );
}
}
// smoothing disabled
if ( smoothing <= 0 && innerflag != 1 ) {
// update solution approximation x
// x = x + U(:,1:s) * beta(1:s)
magma_dsetvector( s, hbeta.val, 1, dbeta.dval, 1, queue );
magmablas_dgemv( MagmaNoTrans, dU.num_rows, s, c_one, dU.dval, dU.ld, dbeta.dval, 1, c_one, x->dval, 1, queue );
}
// check convergence or iteration limit or invalid result of inner loop
if ( innerflag > 0 ) {
break;
}
// t = A v
// t = A r
CHECK( magma_d_spmv( c_one, A, dr, c_zero, dt, queue ));
solver_par->spmv_count++;
// computation of a new omega
//---------------------------------------
// |t|
nrmt = magma_dnrm2( dt.num_rows, dt.dval, 1, queue );
// t'r
tr = magma_ddot( dt.num_rows, dt.dval, 1, dr.dval, 1, queue );
// rho = abs(t' * r) / (|t| * |r|))
rho = MAGMA_D_ABS( MAGMA_D_REAL(tr) / (nrmt * nrmr) );
// om = (t' * r) / (|t| * |t|)
om = tr / (nrmt * nrmt);
if ( rho < angle ) {
om = (om * angle) / rho;
}
//---------------------------------------
if ( MAGMA_D_EQUAL(om, MAGMA_D_ZERO) ) {
info = MAGMA_DIVERGENCE;
break;
}
// update approximation vector
// x = x + om * v
// x = x + om * r
magma_daxpy( x->num_rows, om, dr.dval, 1, x->dval, 1, queue );
// update residual vector
// r = r - om * t
magma_daxpy( dr.num_rows, -om, dt.dval, 1, dr.dval, 1, queue );
// smoothing disabled
if ( smoothing <= 0 ) {
// residual norm
nrmr = magma_dnrm2( b.num_rows, dr.dval, 1, queue );
// smoothing enabled
} else {
// smoothing operation
//---------------------------------------
// t = rs - r
magma_dcopyvector( drs.num_rows, drs.dval, 1, dt.dval, 1, queue );
magma_daxpy( dt.num_rows, c_n_one, dr.dval, 1, dt.dval, 1, queue );
// t't
// t'rs
tt = magma_ddot( dt.num_rows, dt.dval, 1, dt.dval, 1, queue );
tr = magma_ddot( dt.num_rows, dt.dval, 1, drs.dval, 1, queue );
// gamma = (t' * rs) / (|t| * |t|)
gamma = tr / tt;
// rs = rs - gamma * (rs - r)
magma_daxpy( drs.num_rows, -gamma, dt.dval, 1, drs.dval, 1, queue );
// xs = xs - gamma * (xs - x)
magma_dcopyvector( dxs.num_rows, dxs.dval, 1, dt.dval, 1, queue );
magma_daxpy( dt.num_rows, c_n_one, x->dval, 1, dt.dval, 1, queue );
magma_daxpy( dxs.num_rows, -gamma, dt.dval, 1, dxs.dval, 1, queue );
// |rs|
nrmr = magma_dnrm2( b.num_rows, drs.dval, 1, queue );
//---------------------------------------
}
// store current timing and residual
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)nrmr;
solver_par->timing[(solver_par->numiter) / solver_par->verbose]
= (real_Double_t)tempo2 - tempo1;
}
}
// check convergence
if ( nrmr <= solver_par->atol ||
nrmr/nrmb <= solver_par->rtol ) {
info = MAGMA_SUCCESS;
break;
}
}
while ( solver_par->numiter + 1 <= solver_par->maxiter );
// smoothing enabled
if ( smoothing > 0 ) {
// x = xs
magma_dcopyvector( x->num_rows, dxs.dval, 1, x->dval, 1, queue );
// r = rs
magma_dcopyvector( dr.num_rows, drs.dval, 1, dr.dval, 1, queue );
}
// get last iteration timing
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t)tempo2 - tempo1;
//--------------STOP TIME----------------
// get final stats
solver_par->iter_res = nrmr;
CHECK( magma_dresidualvec( A, b, *x, &dr, &residual, queue ));
solver_par->final_res = residual;
// set solver conclusion
if ( info != MAGMA_SUCCESS && info != MAGMA_DIVERGENCE ) {
if ( solver_par->init_res > solver_par->final_res ) {
info = MAGMA_SLOW_CONVERGENCE;
}
}
cleanup:
// free resources
// smoothing enabled
if ( smoothing > 0 ) {
magma_dmfree( &dxs, queue );
magma_dmfree( &drs, queue );
}
magma_dmfree( &dr, queue );
magma_dmfree( &dP, queue );
magma_dmfree( &dP1, queue );
magma_dmfree( &dG, queue );
magma_dmfree( &dU, queue );
magma_dmfree( &dM, queue );
magma_dmfree( &df, queue );
magma_dmfree( &dt, queue );
magma_dmfree( &dc, queue );
magma_dmfree( &dv, queue );
magma_dmfree( &dbeta, queue );
magma_dmfree( &hbeta, queue );
solver_par->info = info;
return info;
/* magma_didr */
}
extern "C" magma_int_t
magma_dfgmres(
magma_d_matrix A, magma_d_matrix b, magma_d_matrix *x,
magma_d_solver_par *solver_par,
magma_d_preconditioner *precond_par,
magma_queue_t queue )
{
magma_int_t info = MAGMA_NOTCONVERGED;
magma_int_t dofs = A.num_rows;
// prepare solver feedback
solver_par->solver = Magma_PGMRES;
solver_par->numiter = 0;
solver_par->spmv_count = 0;
//Chronometry
real_Double_t tempo1, tempo2;
magma_int_t dim = solver_par->restart;
magma_int_t m1 = dim+1; // used inside H macro
magma_int_t i, j, k;
double beta;
double rel_resid, resid0=1, r0=0.0, betanom = 0.0, nom;
magma_d_matrix v_t={Magma_CSR}, w_t={Magma_CSR}, t={Magma_CSR}, t2={Magma_CSR}, V={Magma_CSR}, W={Magma_CSR};
v_t.memory_location = Magma_DEV;
v_t.num_rows = dofs;
v_t.num_cols = 1;
v_t.dval = NULL;
v_t.storage_type = Magma_DENSE;
w_t.memory_location = Magma_DEV;
w_t.num_rows = dofs;
w_t.num_cols = 1;
w_t.dval = NULL;
w_t.storage_type = Magma_DENSE;
double temp;
double *H={0}, *s={0}, *cs={0}, *sn={0};
CHECK( magma_dvinit( &t, Magma_DEV, dofs, 1, MAGMA_D_ZERO, queue ));
CHECK( magma_dvinit( &t2, Magma_DEV, dofs, 1, MAGMA_D_ZERO, queue ));
CHECK( magma_dmalloc_pinned( &H, (dim+1)*dim ));
CHECK( magma_dmalloc_pinned( &s, dim+1 ));
CHECK( magma_dmalloc_pinned( &cs, dim ));
CHECK( magma_dmalloc_pinned( &sn, dim ));
CHECK( magma_dvinit( &V, Magma_DEV, dofs*(dim+1), 1, MAGMA_D_ZERO, queue ));
CHECK( magma_dvinit( &W, Magma_DEV, dofs*dim, 1, MAGMA_D_ZERO, queue ));
CHECK( magma_dresidual( A, b, *x, &nom, queue));
solver_par->init_res = nom;
if ( ( nom * solver_par->rtol) < ATOLERANCE )
r0 = ATOLERANCE;
solver_par->numiter = 0;
solver_par->spmv_count = 0;
tempo1 = magma_sync_wtime( queue );
do
{
solver_par->numiter++;
// compute initial residual and its norm
// A.mult(n, 1, x, n, V(0), n); // V(0) = A*x
CHECK( magma_d_spmv( MAGMA_D_ONE, A, *x, MAGMA_D_ZERO, t, queue ));
solver_par->spmv_count++;
magma_dcopy( dofs, t.dval, 1, V(0), 1, queue );
temp = MAGMA_D_MAKE(-1.0, 0.0);
magma_daxpy( dofs,temp, b.dval, 1, V(0), 1, queue ); // V(0) = V(0) - b
beta = MAGMA_D_MAKE( magma_dnrm2( dofs, V(0), 1, queue ), 0.0 ); // beta = norm(V(0))
if( magma_d_isnan_inf( beta ) ){
info = MAGMA_DIVERGENCE;
break;
}
if (solver_par->numiter == 0){
solver_par->init_res = MAGMA_D_REAL( beta );
resid0 = MAGMA_D_REAL( beta );
r0 = resid0 * solver_par->rtol;
if ( r0 < ATOLERANCE )
r0 = ATOLERANCE;
if ( resid0 < r0 ) {
solver_par->final_res = solver_par->init_res;
solver_par->iter_res = solver_par->init_res;
info = MAGMA_SUCCESS;
goto cleanup;
}
}
if ( solver_par->verbose > 0 ) {
solver_par->res_vec[0] = resid0;
solver_par->timing[0] = 0.0;
}
temp = -1.0/beta;
magma_dscal( dofs, temp, V(0), 1, queue ); // V(0) = -V(0)/beta
// save very first residual norm
if (solver_par->numiter == 0)
solver_par->init_res = MAGMA_D_REAL( beta );
for (i = 1; i < dim+1; i++)
s[i] = MAGMA_D_ZERO;
s[0] = beta;
i = -1;
do {
i++;
// M.apply(n, 1, V(i), n, W(i), n);
v_t.dval = V(i);
CHECK( magma_d_applyprecond_left( MagmaNoTrans, A, v_t, &t, precond_par, queue ));
CHECK( magma_d_applyprecond_right( MagmaNoTrans, A, t, &t2, precond_par, queue ));
magma_dcopy( dofs, t2.dval, 1, W(i), 1, queue );
// A.mult(n, 1, W(i), n, V(i+1), n);
w_t.dval = W(i);
CHECK( magma_d_spmv( MAGMA_D_ONE, A, w_t, MAGMA_D_ZERO, t, queue ));
solver_par->spmv_count++;
magma_dcopy( dofs, t.dval, 1, V(i+1), 1, queue );
for (k = 0; k <= i; k++)
{
H(k, i) = magma_ddot( dofs, V(k), 1, V(i+1), 1, queue );
temp = -H(k,i);
// V(i+1) -= H(k, i) * V(k);
magma_daxpy( dofs,-H(k,i), V(k), 1, V(i+1), 1, queue );
}
H(i+1, i) = MAGMA_D_MAKE( magma_dnrm2( dofs, V(i+1), 1, queue), 0. ); // H(i+1,i) = ||r||
temp = 1.0 / H(i+1, i);
// V(i+1) = V(i+1) / H(i+1, i)
magma_dscal( dofs, temp, V(i+1), 1, queue ); // (to be fused)
for (k = 0; k < i; k++)
ApplyPlaneRotation(&H(k,i), &H(k+1,i), cs[k], sn[k]);
GeneratePlaneRotation(H(i,i), H(i+1,i), &cs[i], &sn[i]);
ApplyPlaneRotation(&H(i,i), &H(i+1,i), cs[i], sn[i]);
ApplyPlaneRotation(&s[i], &s[i+1], cs[i], sn[i]);
betanom = MAGMA_D_ABS( s[i+1] );
rel_resid = betanom / resid0;
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 (rel_resid <= solver_par->rtol || betanom <= solver_par->atol ){
info = MAGMA_SUCCESS;
break;
}
}
while (i+1 < dim && solver_par->numiter+1 <= solver_par->maxiter);
// solve upper triangular system in place
for (j = i; j >= 0; j--)
{
s[j] /= H(j,j);
for (k = j-1; k >= 0; k--)
s[k] -= H(k,j) * s[j];
}
// update the solution
for (j = 0; j <= i; j++)
{
// x = x + s[j] * W(j)
magma_daxpy( dofs, s[j], W(j), 1, x->dval, 1, queue );
}
}
while (rel_resid > solver_par->rtol
&& solver_par->numiter+1 <= solver_par->maxiter);
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t) tempo2-tempo1;
double residual;
CHECK( magma_dresidual( A, b, *x, &residual, queue ));
solver_par->iter_res = betanom;
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==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->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==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:
// free pinned memory
magma_free_pinned(s);
magma_free_pinned(cs);
magma_free_pinned(sn);
magma_free_pinned(H);
//free DEV memory
magma_dmfree( &V, queue);
magma_dmfree( &W, queue);
magma_dmfree( &t, queue);
magma_dmfree( &t2, queue);
solver_par->info = info;
return info;
} /* magma_dfgmres */
extern "C" magma_int_t
magma_dpcgs(
magma_d_matrix A, magma_d_matrix b, magma_d_matrix *x,
magma_d_solver_par *solver_par,
magma_d_preconditioner *precond_par,
magma_queue_t queue )
{
magma_int_t info = MAGMA_NOTCONVERGED;
// prepare solver feedback
solver_par->solver = Magma_PCGS;
solver_par->numiter = 0;
solver_par->spmv_count = 0;
// constants
const double c_zero = MAGMA_D_ZERO;
const double c_one = MAGMA_D_ONE;
const double c_neg_one = MAGMA_D_NEG_ONE;
// solver variables
double nom0, r0, res=0, nomb;
double rho, rho_l = c_one, alpha, beta;
magma_int_t dofs = A.num_rows* b.num_cols;
// GPU workspace
magma_d_matrix r={Magma_CSR}, rt={Magma_CSR}, r_tld={Magma_CSR},
p={Magma_CSR}, q={Magma_CSR}, u={Magma_CSR}, v={Magma_CSR}, t={Magma_CSR},
p_hat={Magma_CSR}, q_hat={Magma_CSR}, u_hat={Magma_CSR}, v_hat={Magma_CSR};
CHECK( magma_dvinit( &r, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &rt,Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &r_tld,Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &p, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &p_hat, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &q, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &q_hat, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &u, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &u_hat, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &v, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &v_hat, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &t, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
// solver setup
CHECK( magma_dresidualvec( A, b, *x, &r, &nom0, queue));
magma_dcopy( dofs, r.dval, 1, r_tld.dval, 1, queue );
solver_par->init_res = nom0;
nomb = magma_dnrm2( dofs, 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;
}
//Chronometry
real_Double_t tempo1, tempo2, tempop1, tempop2;
tempo1 = magma_sync_wtime( queue );
solver_par->numiter = 0;
solver_par->spmv_count = 0;
// start iteration
do
{
solver_par->numiter++;
rho = magma_ddot( dofs, r.dval, 1, r_tld.dval, 1, queue );
// rho = < r,r_tld>
if( magma_d_isnan_inf( rho ) ){
info = MAGMA_DIVERGENCE;
break;
}
if ( solver_par->numiter > 1 ) { // direction vectors
beta = rho / rho_l;
magma_dcopy( dofs, r.dval, 1, u.dval, 1, queue ); // u = r
magma_daxpy( dofs, beta, q.dval, 1, u.dval, 1, queue ); // u = r + beta q
magma_dscal( dofs, beta, p.dval, 1, queue ); // p = beta*p
magma_daxpy( dofs, c_one, q.dval, 1, p.dval, 1, queue ); // p = q + beta*p
magma_dscal( dofs, beta, p.dval, 1, queue ); // p = beta*(q + beta*p)
magma_daxpy( dofs, c_one, u.dval, 1, p.dval, 1, queue ); // p = u + beta*(q + beta*p)
//u = r + beta*q;
//p = u + beta*( q + beta*p );
}
else{
magma_dcopy( dofs, r.dval, 1, u.dval, 1, queue ); // u = r
magma_dcopy( dofs, r.dval, 1, p.dval, 1, queue ); // p = r
}
// preconditioner
tempop1 = magma_sync_wtime( queue );
CHECK( magma_d_applyprecond_left( MagmaNoTrans, A, p, &rt, precond_par, queue ));
CHECK( magma_d_applyprecond_right( MagmaNoTrans, A, rt, &p_hat, precond_par, queue ));
tempop2 = magma_sync_wtime( queue );
precond_par->runtime += tempop2-tempop1;
// SpMV
CHECK( magma_d_spmv( c_one, A, p_hat, c_zero, v_hat, queue )); // v = A p
solver_par->spmv_count++;
alpha = rho / magma_ddot( dofs, r_tld.dval, 1, v_hat.dval, 1, queue );
magma_dcopy( dofs, u.dval, 1, q.dval, 1, queue ); // q = u
magma_daxpy( dofs, -alpha, v_hat.dval, 1, q.dval, 1, queue ); // q = u - alpha v_hat
magma_dcopy( dofs, u.dval, 1, t.dval, 1, queue ); // t = q
magma_daxpy( dofs, c_one, q.dval, 1, t.dval, 1, queue ); // t = u + q
// preconditioner
tempop1 = magma_sync_wtime( queue );
CHECK( magma_d_applyprecond_left( MagmaNoTrans, A, t, &rt, precond_par, queue ));
CHECK( magma_d_applyprecond_right( MagmaNoTrans, A, rt, &u_hat, precond_par, queue ));
tempop2 = magma_sync_wtime( queue );
precond_par->runtime += tempop2-tempop1;
// SpMV
CHECK( magma_d_spmv( c_one, A, u_hat, c_zero, t, queue )); // t = A u_hat
solver_par->spmv_count++;
magma_daxpy( dofs, alpha, u_hat.dval, 1, x->dval, 1, queue ); // x = x + alpha u_hat
magma_daxpy( dofs, c_neg_one*alpha, t.dval, 1, r.dval, 1, queue ); // r = r -alpha*A u_hat
res = magma_dnrm2( dofs, r.dval, 1, queue );
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/nomb <= solver_par->rtol || res <= solver_par->atol ){
break;
}
rho_l = rho;
}
while ( solver_par->numiter+1 <= solver_par->maxiter );
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t) tempo2-tempo1;
double residual;
CHECK( magma_dresidualvec( 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 == 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;
}
}
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 == 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;
}
}
info = MAGMA_DIVERGENCE;
}
cleanup:
magma_dmfree(&r, queue );
magma_dmfree(&rt, queue );
magma_dmfree(&r_tld, queue );
magma_dmfree(&p, queue );
magma_dmfree(&q, queue );
magma_dmfree(&u, queue );
magma_dmfree(&v, queue );
magma_dmfree(&t, queue );
magma_dmfree(&p_hat, queue );
magma_dmfree(&q_hat, queue );
magma_dmfree(&u_hat, queue );
magma_dmfree(&v_hat, queue );
solver_par->info = info;
return info;
} /* magma_dpcgs */
extern "C" magma_int_t
magma_dpbicg(
magma_d_matrix A, magma_d_matrix b, magma_d_matrix *x,
magma_d_solver_par *solver_par,
magma_d_preconditioner *precond_par,
magma_queue_t queue )
{
magma_int_t info = MAGMA_NOTCONVERGED;
// prepare solver feedback
solver_par->solver = Magma_PBICG;
solver_par->numiter = 0;
solver_par->spmv_count = 0;
// some useful variables
double c_zero = MAGMA_D_ZERO;
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
magma_int_t dofs = A.num_rows * b.num_cols;
// workspace
magma_d_matrix r={Magma_CSR}, rt={Magma_CSR}, p={Magma_CSR}, pt={Magma_CSR},
z={Magma_CSR}, zt={Magma_CSR}, q={Magma_CSR}, y={Magma_CSR},
yt={Magma_CSR}, qt={Magma_CSR};
// need to transpose the matrix
magma_d_matrix AT={Magma_CSR}, Ah1={Magma_CSR}, Ah2={Magma_CSR};
CHECK( magma_dvinit( &r, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &rt,Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &p, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &pt,Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &q, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &qt,Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &y, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &yt,Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &z, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &zt,Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
// solver variables
double alpha, rho, beta, rho_new, ptq;
double res, nomb, nom0, r0;
// transpose the matrix
magma_dmtransfer( A, &Ah1, Magma_DEV, Magma_CPU, queue );
magma_dmconvert( Ah1, &Ah2, A.storage_type, Magma_CSR, queue );
magma_dmfree(&Ah1, queue );
magma_dmtransposeconjugate( Ah2, &Ah1, queue );
magma_dmfree(&Ah2, queue );
Ah2.blocksize = A.blocksize;
Ah2.alignment = A.alignment;
magma_dmconvert( Ah1, &Ah2, Magma_CSR, A.storage_type, queue );
magma_dmfree(&Ah1, queue );
magma_dmtransfer( Ah2, &AT, Magma_CPU, Magma_DEV, queue );
magma_dmfree(&Ah2, queue );
// solver setup
CHECK( magma_dresidualvec( A, b, *x, &r, &nom0, queue));
res = nom0;
solver_par->init_res = nom0;
magma_dcopy( dofs, r.dval, 1, rt.dval, 1, queue ); // rr = r
rho_new = magma_ddot( dofs, rt.dval, 1, r.dval, 1, queue ); // rho=<rr,r>
rho = alpha = MAGMA_D_MAKE( 1.0, 0. );
nomb = magma_dnrm2( dofs, 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] = nom0;
solver_par->timing[0] = 0.0;
}
if ( nom0 < r0 ) {
info = MAGMA_SUCCESS;
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++;
CHECK( magma_d_applyprecond_left( MagmaNoTrans, A, r, &y, precond_par, queue ));
CHECK( magma_d_applyprecond_right( MagmaNoTrans, A, y, &z, precond_par, queue ));
CHECK( magma_d_applyprecond_right( MagmaTrans, A, rt, &yt, precond_par, queue ));
CHECK( magma_d_applyprecond_left( MagmaTrans, A, yt, &zt, precond_par, queue ));
//magma_dcopy( dofs, r.dval, 1 , y.dval, 1, queue ); // y=r
//magma_dcopy( dofs, y.dval, 1 , z.dval, 1, queue ); // z=y
//magma_dcopy( dofs, rt.dval, 1 , yt.dval, 1, queue ); // yt=rt
//magma_dcopy( dofs, yt.dval, 1 , zt.dval, 1, queue ); // yt=rt
rho= rho_new;
rho_new = magma_ddot( dofs, rt.dval, 1, z.dval, 1, queue ); // rho=<rt,z>
if( magma_d_isnan_inf( rho_new ) ){
info = MAGMA_DIVERGENCE;
break;
}
if( solver_par->numiter==1 ){
magma_dcopy( dofs, z.dval, 1 , p.dval, 1, queue ); // yt=rt
magma_dcopy( dofs, zt.dval, 1 , pt.dval, 1, queue ); // zt=yt
} else {
beta = rho_new/rho;
magma_dscal( dofs, beta, p.dval, 1, queue ); // p = beta*p
magma_daxpy( dofs, c_one , z.dval, 1 , p.dval, 1, queue ); // p = z+beta*p
magma_dscal( dofs, MAGMA_D_CONJ(beta), pt.dval, 1, queue ); // pt = beta*pt
magma_daxpy( dofs, c_one , zt.dval, 1 , pt.dval, 1, queue ); // pt = zt+beta*pt
}
CHECK( magma_d_spmv( c_one, A, p, c_zero, q, queue )); // v = Ap
CHECK( magma_d_spmv( c_one, AT, pt, c_zero, qt, queue )); // v = Ap
solver_par->spmv_count++;
solver_par->spmv_count++;
ptq = magma_ddot( dofs, pt.dval, 1, q.dval, 1, queue );
alpha = rho_new /ptq;
magma_daxpy( dofs, alpha, p.dval, 1 , x->dval, 1, queue ); // x=x+alpha*p
magma_daxpy( dofs, c_neg_one * alpha, q.dval, 1 , r.dval, 1, queue ); // r=r+alpha*q
magma_daxpy( dofs, c_neg_one * MAGMA_D_CONJ(alpha), qt.dval, 1 , rt.dval, 1, queue ); // r=r+alpha*q
res = magma_dnrm2( dofs, r.dval, 1, queue );
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/nomb <= solver_par->rtol || res <= solver_par->atol ){
break;
}
}
while ( solver_par->numiter+1 <= solver_par->maxiter );
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t) tempo2-tempo1;
double residual;
CHECK( magma_dresidualvec( 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;
} 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) 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==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;
}
}
info = MAGMA_DIVERGENCE;
}
cleanup:
magma_dmfree(&r, queue );
magma_dmfree(&rt, queue );
magma_dmfree(&p, queue );
magma_dmfree(&pt, queue );
magma_dmfree(&q, queue );
magma_dmfree(&qt, queue );
magma_dmfree(&y, queue );
magma_dmfree(&yt, queue );
magma_dmfree(&z, queue );
magma_dmfree(&zt, queue );
magma_dmfree(&AT, queue );
magma_dmfree(&Ah1, queue );
magma_dmfree(&Ah2, queue );
solver_par->info = info;
return info;
} /* magma_dpbicg */
extern "C" magma_int_t
magma_dcumilugeneratesolverinfo(
magma_d_preconditioner *precond,
magma_queue_t queue )
{
magma_int_t info = 0;
cusparseHandle_t cusparseHandle=NULL;
cusparseMatDescr_t descrL=NULL;
cusparseMatDescr_t descrU=NULL;
magma_d_matrix hA={Magma_CSR}, hL={Magma_CSR}, hU={Magma_CSR};
if (precond->L.memory_location != Magma_DEV ){
CHECK( magma_dmtransfer( precond->M, &hA,
precond->M.memory_location, Magma_CPU, queue ));
hL.diagorder_type = Magma_UNITY;
CHECK( magma_dmconvert( hA, &hL , Magma_CSR, Magma_CSRL, queue ));
hU.diagorder_type = Magma_VALUE;
CHECK( magma_dmconvert( hA, &hU , Magma_CSR, Magma_CSRU, queue ));
CHECK( magma_dmtransfer( hL, &(precond->L), Magma_CPU, Magma_DEV, queue ));
CHECK( magma_dmtransfer( hU, &(precond->U), Magma_CPU, Magma_DEV, queue ));
magma_dmfree(&hA, queue );
magma_dmfree(&hL, queue );
magma_dmfree(&hU, queue );
}
// CUSPARSE context //
CHECK_CUSPARSE( cusparseCreate( &cusparseHandle ));
CHECK_CUSPARSE( cusparseSetStream( cusparseHandle, queue->cuda_stream() ));
CHECK_CUSPARSE( cusparseCreateMatDescr( &descrL ));
CHECK_CUSPARSE( cusparseSetMatType( descrL, CUSPARSE_MATRIX_TYPE_TRIANGULAR ));
CHECK_CUSPARSE( cusparseSetMatDiagType( descrL, CUSPARSE_DIAG_TYPE_UNIT ));
CHECK_CUSPARSE( cusparseSetMatIndexBase( descrL, CUSPARSE_INDEX_BASE_ZERO ));
CHECK_CUSPARSE( cusparseSetMatFillMode( descrL, CUSPARSE_FILL_MODE_LOWER ));
CHECK_CUSPARSE( cusparseCreateSolveAnalysisInfo( &precond->cuinfoL ));
CHECK_CUSPARSE( cusparseDcsrsm_analysis( cusparseHandle,
CUSPARSE_OPERATION_NON_TRANSPOSE, precond->L.num_rows,
precond->L.nnz, descrL,
precond->L.dval, precond->L.drow, precond->L.dcol, precond->cuinfoL ));
CHECK_CUSPARSE( cusparseCreateMatDescr( &descrU ));
CHECK_CUSPARSE( cusparseSetMatType( descrU, CUSPARSE_MATRIX_TYPE_TRIANGULAR ));
CHECK_CUSPARSE( cusparseSetMatDiagType( descrU, CUSPARSE_DIAG_TYPE_NON_UNIT ));
CHECK_CUSPARSE( cusparseSetMatIndexBase( descrU, CUSPARSE_INDEX_BASE_ZERO ));
CHECK_CUSPARSE( cusparseSetMatFillMode( descrU, CUSPARSE_FILL_MODE_UPPER ));
CHECK_CUSPARSE( cusparseCreateSolveAnalysisInfo( &precond->cuinfoU ));
CHECK_CUSPARSE( cusparseDcsrsm_analysis( cusparseHandle,
CUSPARSE_OPERATION_NON_TRANSPOSE, precond->U.num_rows,
precond->U.nnz, descrU,
precond->U.dval, precond->U.drow, precond->U.dcol, precond->cuinfoU ));
if( precond->maxiter < 50 ){
//prepare for iterative solves
// extract the diagonal of L into precond->d
CHECK( magma_djacobisetup_diagscal( precond->L, &precond->d, queue ));
CHECK( magma_dvinit( &precond->work1, Magma_DEV, precond->U.num_rows, 1, MAGMA_D_ZERO, queue ));
// extract the diagonal of U into precond->d2
CHECK( magma_djacobisetup_diagscal( precond->U, &precond->d2, queue ));
CHECK( magma_dvinit( &precond->work2, Magma_DEV, precond->U.num_rows, 1, MAGMA_D_ZERO, queue ));
}
cleanup:
cusparseDestroyMatDescr( descrL );
cusparseDestroyMatDescr( descrU );
cusparseDestroy( cusparseHandle );
return info;
}
extern "C" magma_int_t
magma_dcumiccsetup(
magma_d_matrix A,
magma_d_preconditioner *precond,
magma_queue_t queue )
{
magma_int_t info = 0;
cusparseHandle_t cusparseHandle=NULL;
cusparseMatDescr_t descrA=NULL;
cusparseMatDescr_t descrL=NULL;
cusparseMatDescr_t descrU=NULL;
#if CUDA_VERSION >= 7000
csric02Info_t info_M=NULL;
void *pBuffer = NULL;
#endif
magma_d_matrix hA={Magma_CSR}, hACSR={Magma_CSR}, U={Magma_CSR};
CHECK( magma_dmtransfer( A, &hA, A.memory_location, Magma_CPU, queue ));
U.diagorder_type = Magma_VALUE;
CHECK( magma_dmconvert( hA, &hACSR, hA.storage_type, Magma_CSR, queue ));
// in case using fill-in
if( precond->levels > 0 ){
magma_d_matrix hAL={Magma_CSR}, hAUt={Magma_CSR};
CHECK( magma_dsymbilu( &hACSR, precond->levels, &hAL, &hAUt, queue ));
magma_dmfree(&hAL, queue);
magma_dmfree(&hAUt, queue);
}
CHECK( magma_dmconvert( hACSR, &U, Magma_CSR, Magma_CSRL, queue ));
magma_dmfree( &hACSR, queue );
CHECK( magma_dmtransfer(U, &(precond->M), Magma_CPU, Magma_DEV, queue ));
// CUSPARSE context //
CHECK_CUSPARSE( cusparseCreate( &cusparseHandle ));
CHECK_CUSPARSE( cusparseSetStream( cusparseHandle, queue->cuda_stream() ));
CHECK_CUSPARSE( cusparseCreateMatDescr( &descrA ));
CHECK_CUSPARSE( cusparseCreateSolveAnalysisInfo( &(precond->cuinfo) ));
// use kernel to manually check for zeros n the diagonal
CHECK( magma_ddiagcheck( precond->M, queue ) );
#if CUDA_VERSION >= 7000
// this version has the bug fixed where a zero on the diagonal causes a crash
CHECK_CUSPARSE( cusparseCreateCsric02Info(&info_M) );
CHECK_CUSPARSE( cusparseSetMatType( descrA, CUSPARSE_MATRIX_TYPE_GENERAL ));
CHECK_CUSPARSE( cusparseSetMatIndexBase( descrA, CUSPARSE_INDEX_BASE_ZERO ));
int buffersize;
int structural_zero;
int numerical_zero;
CHECK_CUSPARSE(
cusparseDcsric02_bufferSize( cusparseHandle,
precond->M.num_rows, precond->M.nnz, descrA,
precond->M.dval, precond->M.drow, precond->M.dcol,
info_M,
&buffersize ) );
CHECK( magma_malloc((void**)&pBuffer, buffersize) );
CHECK_CUSPARSE( cusparseDcsric02_analysis( cusparseHandle,
precond->M.num_rows, precond->M.nnz, descrA,
precond->M.dval, precond->M.drow, precond->M.dcol,
info_M, CUSPARSE_SOLVE_POLICY_NO_LEVEL, pBuffer ));
CHECK_CUSPARSE( cusparseXcsric02_zeroPivot( cusparseHandle, info_M, &numerical_zero ) );
CHECK_CUSPARSE( cusparseXcsric02_zeroPivot( cusparseHandle, info_M, &structural_zero ) );
CHECK_CUSPARSE(
cusparseDcsric02( cusparseHandle,
precond->M.num_rows, precond->M.nnz, descrA,
precond->M.dval, precond->M.drow, precond->M.dcol,
info_M, CUSPARSE_SOLVE_POLICY_NO_LEVEL, pBuffer) );
#else
// this version contains the bug but is needed for backward compability
CHECK_CUSPARSE( cusparseSetMatType( descrA, CUSPARSE_MATRIX_TYPE_SYMMETRIC ));
CHECK_CUSPARSE( cusparseSetMatDiagType( descrA, CUSPARSE_DIAG_TYPE_NON_UNIT ));
CHECK_CUSPARSE( cusparseSetMatIndexBase( descrA, CUSPARSE_INDEX_BASE_ZERO ));
CHECK_CUSPARSE( cusparseSetMatFillMode( descrA, CUSPARSE_FILL_MODE_LOWER ));
CHECK_CUSPARSE( cusparseDcsrsm_analysis( cusparseHandle,
CUSPARSE_OPERATION_NON_TRANSPOSE,
precond->M.num_rows, precond->M.nnz, descrA,
precond->M.dval, precond->M.drow, precond->M.dcol,
precond->cuinfo ));
CHECK_CUSPARSE( cusparseDcsric0( cusparseHandle, CUSPARSE_OPERATION_NON_TRANSPOSE,
precond->M.num_rows, descrA,
precond->M.dval,
precond->M.drow,
precond->M.dcol,
precond->cuinfo ));
#endif
CHECK_CUSPARSE( cusparseCreateMatDescr( &descrL ));
CHECK_CUSPARSE( cusparseSetMatType( descrL, CUSPARSE_MATRIX_TYPE_TRIANGULAR ));
CHECK_CUSPARSE( cusparseSetMatDiagType( descrL, CUSPARSE_DIAG_TYPE_NON_UNIT ));
CHECK_CUSPARSE( cusparseSetMatIndexBase( descrL, CUSPARSE_INDEX_BASE_ZERO ));
CHECK_CUSPARSE( cusparseSetMatFillMode( descrL, CUSPARSE_FILL_MODE_LOWER ));
CHECK_CUSPARSE( cusparseCreateSolveAnalysisInfo( &precond->cuinfoL ));
CHECK_CUSPARSE( cusparseDcsrsm_analysis( cusparseHandle,
CUSPARSE_OPERATION_NON_TRANSPOSE, precond->M.num_rows,
precond->M.nnz, descrL,
precond->M.dval, precond->M.drow, precond->M.dcol, precond->cuinfoL ));
CHECK_CUSPARSE( cusparseCreateMatDescr( &descrU ));
CHECK_CUSPARSE( cusparseSetMatType( descrU, CUSPARSE_MATRIX_TYPE_TRIANGULAR ));
CHECK_CUSPARSE( cusparseSetMatDiagType( descrU, CUSPARSE_DIAG_TYPE_NON_UNIT ));
CHECK_CUSPARSE( cusparseSetMatIndexBase( descrU, CUSPARSE_INDEX_BASE_ZERO ));
CHECK_CUSPARSE( cusparseSetMatFillMode( descrU, CUSPARSE_FILL_MODE_LOWER ));
CHECK_CUSPARSE( cusparseCreateSolveAnalysisInfo( &precond->cuinfoU ));
CHECK_CUSPARSE( cusparseDcsrsm_analysis( cusparseHandle,
CUSPARSE_OPERATION_TRANSPOSE, precond->M.num_rows,
precond->M.nnz, descrU,
precond->M.dval, precond->M.drow, precond->M.dcol, precond->cuinfoU ));
if( precond->maxiter < 50 ){
//prepare for iterative solves
// copy the matrix to precond->L and (transposed) to precond->U
CHECK( magma_dmtransfer(precond->M, &(precond->L), Magma_DEV, Magma_DEV, queue ));
CHECK( magma_dmtranspose( precond->L, &(precond->U), queue ));
// extract the diagonal of L into precond->d
CHECK( magma_djacobisetup_diagscal( precond->L, &precond->d, queue ));
CHECK( magma_dvinit( &precond->work1, Magma_DEV, hA.num_rows, 1, MAGMA_D_ZERO, queue ));
// extract the diagonal of U into precond->d2
CHECK( magma_djacobisetup_diagscal( precond->U, &precond->d2, queue ));
CHECK( magma_dvinit( &precond->work2, Magma_DEV, hA.num_rows, 1, MAGMA_D_ZERO, queue ));
}
/*
// to enable also the block-asynchronous iteration for the triangular solves
CHECK( magma_dmtransfer( precond->M, &hA, Magma_DEV, Magma_CPU, queue ));
hA.storage_type = Magma_CSR;
magma_d_matrix hD, hR, hAt
CHECK( magma_dcsrsplit( 256, hA, &hD, &hR, queue ));
CHECK( magma_dmtransfer( hD, &precond->LD, Magma_CPU, Magma_DEV, queue ));
CHECK( magma_dmtransfer( hR, &precond->L, Magma_CPU, Magma_DEV, queue ));
magma_dmfree(&hD, queue );
magma_dmfree(&hR, queue );
CHECK( magma_d_cucsrtranspose( hA, &hAt, queue ));
CHECK( magma_dcsrsplit( 256, hAt, &hD, &hR, queue ));
CHECK( magma_dmtransfer( hD, &precond->UD, Magma_CPU, Magma_DEV, queue ));
CHECK( magma_dmtransfer( hR, &precond->U, Magma_CPU, Magma_DEV, queue ));
magma_dmfree(&hD, queue );
magma_dmfree(&hR, queue );
magma_dmfree(&hA, queue );
magma_dmfree(&hAt, queue );
*/
cleanup:
#if CUDA_VERSION >= 7000
magma_free( pBuffer );
cusparseDestroyCsric02Info( info_M );
#endif
cusparseDestroySolveAnalysisInfo( precond->cuinfo );
cusparseDestroyMatDescr( descrL );
cusparseDestroyMatDescr( descrU );
cusparseDestroyMatDescr( descrA );
cusparseDestroy( cusparseHandle );
magma_dmfree(&U, queue );
magma_dmfree(&hA, queue );
return info;
}
extern "C" magma_int_t
magma_dbicgstab_merge(
magma_d_matrix A, magma_d_matrix b, magma_d_matrix *x,
magma_d_solver_par *solver_par,
magma_queue_t queue )
{
magma_int_t info = MAGMA_NOTCONVERGED;
// prepare solver feedback
solver_par->solver = Magma_BICGSTAB;
solver_par->numiter = 0;
solver_par->spmv_count = 0;
// some useful variables
double c_zero = MAGMA_D_ZERO;
double c_one = MAGMA_D_ONE;
magma_int_t dofs = A.num_rows * b.num_cols;
// workspace
magma_d_matrix r={Magma_CSR}, rr={Magma_CSR}, p={Magma_CSR}, v={Magma_CSR},
s={Magma_CSR}, t={Magma_CSR}, d1={Magma_CSR}, d2={Magma_CSR};
CHECK( magma_dvinit( &r, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &rr,Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &p, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &v, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &s, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &t, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &d1, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_dvinit( &d2, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
// solver variables
double alpha, beta, omega, rho_old, rho_new;
double nom, betanom, nom0, r0, res, nomb;
res=0;
//double den;
// solver setup
CHECK( magma_dresidualvec( A, b, *x, &r, &nom0, queue));
magma_dcopy( dofs, r.dval, 1, rr.dval, 1, queue ); // rr = r
betanom = nom0;
nom = nom0*nom0;
rho_new = magma_ddot( dofs, r.dval, 1, r.dval, 1, queue ); // rho=<rr,r>
rho_old = omega = alpha = MAGMA_D_MAKE( 1.0, 0. );
solver_par->init_res = nom0;
CHECK( magma_d_spmv( c_one, A, r, c_zero, v, queue )); // z = A r
//den = MAGMA_D_REAL( magma_ddot( dofs, v.dval, 1, r.dval, 1), queue ); // den = z' * r
nomb = magma_dnrm2( dofs, 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] = nom0;
solver_par->timing[0] = 0.0;
}
if ( nom < r0 ) {
info = MAGMA_SUCCESS;
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++;
rho_old = rho_new; // rho_old=rho
rho_new = magma_ddot( dofs, rr.dval, 1, r.dval, 1, queue ); // rho=<rr,r>
beta = rho_new/rho_old * alpha/omega; // beta=rho/rho_old *alpha/omega
if( magma_d_isnan_inf( beta ) ){
info = MAGMA_DIVERGENCE;
break;
}
// p = r + beta * ( p - omega * v )
magma_dbicgstab_1(
r.num_rows,
r.num_cols,
beta,
omega,
r.dval,
v.dval,
p.dval,
queue );
CHECK( magma_d_spmv( c_one, A, p, c_zero, v, queue )); // v = Ap
solver_par->spmv_count++;
//alpha = rho_new / tmpval;
alpha = rho_new /magma_ddot( dofs, rr.dval, 1, v.dval, 1, queue );
if( magma_d_isnan_inf( alpha ) ){
info = MAGMA_DIVERGENCE;
break;
}
// s = r - alpha v
magma_dbicgstab_2(
r.num_rows,
r.num_cols,
alpha,
r.dval,
v.dval,
s.dval,
queue );
CHECK( magma_d_spmv( c_one, A, s, c_zero, t, queue )); // t=As
solver_par->spmv_count++;
omega = magma_ddot( dofs, t.dval, 1, s.dval, 1, queue ) // omega = <s,t>/<t,t>
/ magma_ddot( dofs, t.dval, 1, t.dval, 1, queue );
// x = x + alpha * p + omega * s
// r = s - omega * t
magma_dbicgstab_3(
r.num_rows,
r.num_cols,
alpha,
omega,
p.dval,
s.dval,
t.dval,
x->dval,
r.dval,
queue );
res = betanom = magma_dnrm2( dofs, r.dval, 1, queue );
nom = betanom*betanom;
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/nomb <= solver_par->rtol || res <= solver_par->atol ){
break;
}
}
while ( solver_par->numiter+1 <= solver_par->maxiter );
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t) tempo2-tempo1;
double residual;
CHECK( magma_dresidualvec( 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==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->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==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_dmfree(&r, queue );
magma_dmfree(&rr, queue );
magma_dmfree(&p, queue );
magma_dmfree(&v, queue );
magma_dmfree(&s, queue );
magma_dmfree(&t, queue );
magma_dmfree(&d1, queue );
magma_dmfree(&d2, queue );
solver_par->info = info;
return info;
} /* magma_dbicgstab_merge */