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_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_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 */
magma_int_t
magma_dbicgstab_merge2( magma_d_sparse_matrix A, magma_d_vector b,
magma_d_vector *x, magma_d_solver_par *solver_par ){
// prepare solver feedback
solver_par->solver = Magma_BICGSTABMERGE2;
solver_par->numiter = 0;
solver_par->info = 0;
// some useful variables
double c_zero = MAGMA_D_ZERO, c_one = MAGMA_D_ONE;
magma_int_t dofs = A.num_rows;
// GPU stream
magma_queue_t stream[2];
magma_event_t event[1];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
magma_event_create( &event[0] );
// workspace
magma_d_vector q, r,rr,p,v,s,t;
double *d1, *d2, *skp;
magma_dmalloc( &d1, dofs*(2) );
magma_dmalloc( &d2, dofs*(2) );
// array for the parameters
magma_dmalloc( &skp, 8 );
// skp = [alpha|beta|omega|rho_old|rho|nom|tmp1|tmp2]
magma_d_vinit( &q, Magma_DEV, dofs*6, c_zero );
// q = rr|r|p|v|s|t
rr.memory_location = Magma_DEV; rr.val = NULL; rr.num_rows = rr.nnz = dofs;
r.memory_location = Magma_DEV; r.val = NULL; r.num_rows = r.nnz = dofs;
p.memory_location = Magma_DEV; p.val = NULL; p.num_rows = p.nnz = dofs;
v.memory_location = Magma_DEV; v.val = NULL; v.num_rows = v.nnz = dofs;
s.memory_location = Magma_DEV; s.val = NULL; s.num_rows = s.nnz = dofs;
t.memory_location = Magma_DEV; t.val = NULL; t.num_rows = t.nnz = dofs;
rr.val = q(0);
r.val = q(1);
p.val = q(2);
v.val = q(3);
s.val = q(4);
t.val = q(5);
// solver variables
double alpha, beta, omega, rho_old, rho_new, *skp_h;
double nom, nom0, betanom, r0, den;
// solver setup
magma_dscal( dofs, c_zero, x->val, 1) ; // x = 0
magma_dcopy( dofs, b.val, 1, q(0), 1 ); // rr = b
magma_dcopy( dofs, b.val, 1, q(1), 1 ); // r = b
rho_new = magma_ddot( dofs, r.val, 1, r.val, 1 ); // rho=<rr,r>
nom = MAGMA_D_REAL(magma_ddot( dofs, r.val, 1, r.val, 1 ));
nom0 = betanom = sqrt(nom); // nom = || r ||
rho_old = omega = alpha = MAGMA_D_MAKE( 1.0, 0. );
beta = rho_new;
solver_par->init_res = nom0;
// array on host for the parameters
magma_dmalloc_cpu( &skp_h, 8 );
skp_h[0]=alpha;
skp_h[1]=beta;
skp_h[2]=omega;
skp_h[3]=rho_old;
skp_h[4]=rho_new;
skp_h[5]=MAGMA_D_MAKE(nom, 0.0);
magma_dsetvector( 8, skp_h, 1, skp, 1 );
magma_d_spmv( c_one, A, r, c_zero, v ); // z = A r
den = MAGMA_D_REAL( magma_ddot(dofs, v.val, 1, r.val, 1) );// den = z dot r
if ( (r0 = nom * solver_par->epsilon) < ATOLERANCE )
r0 = ATOLERANCE;
if ( nom < r0 )
return MAGMA_SUCCESS;
// check positive definite
if (den <= 0.0) {
printf("Operator A is not postive definite. (Ar,r) = %f\n", den);
return -100;
}
//Chronometry
real_Double_t tempo1, tempo2;
magma_device_sync(); tempo1=magma_wtime();
if( solver_par->verbose > 0 ){
solver_par->res_vec[0] = nom0;
solver_par->timing[0] = 0.0;
}
// start iteration
for( solver_par->numiter= 1; solver_par->numiter<solver_par->maxiter;
solver_par->numiter++ ){
magmablasSetKernelStream(stream[0]);
// computes p=r+beta*(p-omega*v)
magma_dbicgmerge1( dofs, skp, v.val, r.val, p.val );
magma_dbicgmerge_spmv1( A, d1, d2, q(2), q(0), q(3), skp );
magma_dbicgmerge2( dofs, skp, r.val, v.val, s.val ); // s=r-alpha*v
magma_dbicgmerge_spmv2( A, d1, d2, q(4), q(5), skp);
magma_dbicgmerge_xrbeta( dofs, d1, d2, q(0), q(1), q(2),
q(4), q(5), x->val, skp);
// check stopping criterion (asynchronous copy)
magma_dgetvector_async( 1 , skp+5, 1,
skp_h+5, 1, stream[1] );
betanom = sqrt(MAGMA_D_REAL(skp_h[5]));
if( solver_par->verbose > 0 ){
magma_device_sync(); tempo2=magma_wtime();
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;
}
}
magma_device_sync(); tempo2=magma_wtime();
solver_par->runtime = (real_Double_t) tempo2-tempo1;
double residual;
magma_dresidual( A, b, *x, &residual );
solver_par->iter_res = betanom;
solver_par->final_res = residual;
if( solver_par->numiter < solver_par->maxiter){
solver_par->info = 0;
}else if( solver_par->init_res > solver_par->final_res ){
if( solver_par->verbose > 0 ){
if( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
solver_par->info = -2;
}
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 = -1;
}
magma_d_vfree(&q); // frees all vectors
magma_free(d1);
magma_free(d2);
magma_free( skp );
magma_free_cpu( skp_h );
return MAGMA_SUCCESS;
} /* dbicgstab_merge2 */
magma_int_t
magma_dpgmres( magma_d_sparse_matrix A, magma_d_vector b, magma_d_vector *x,
magma_d_solver_par *solver_par,
magma_d_preconditioner *precond_par ){
// prepare solver feedback
solver_par->solver = Magma_PGMRES;
solver_par->numiter = 0;
solver_par->info = 0;
// local 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;
magma_int_t i, j, k, m = 0;
magma_int_t restart = min( dofs-1, solver_par->restart );
magma_int_t ldh = restart+1;
double nom, rNorm, RNorm, nom0, betanom, r0 = 0.;
// CPU workspace
magma_setdevice(0);
double *H, *HH, *y, *h1;
magma_dmalloc_pinned( &H, (ldh+1)*ldh );
magma_dmalloc_pinned( &y, ldh );
magma_dmalloc_pinned( &HH, ldh*ldh );
magma_dmalloc_pinned( &h1, ldh );
// GPU workspace
magma_d_vector r, q, q_t, z, z_t, t;
magma_d_vinit( &t, Magma_DEV, dofs, c_zero );
magma_d_vinit( &r, Magma_DEV, dofs, c_zero );
magma_d_vinit( &q, Magma_DEV, dofs*(ldh+1), c_zero );
magma_d_vinit( &z, Magma_DEV, dofs*(ldh+1), c_zero );
magma_d_vinit( &z_t, Magma_DEV, dofs, c_zero );
q_t.memory_location = Magma_DEV;
q_t.val = NULL;
q_t.num_rows = q_t.nnz = dofs;
double *dy, *dH = NULL;
if (MAGMA_SUCCESS != magma_dmalloc( &dy, ldh ))
return MAGMA_ERR_DEVICE_ALLOC;
if (MAGMA_SUCCESS != magma_dmalloc( &dH, (ldh+1)*ldh ))
return MAGMA_ERR_DEVICE_ALLOC;
// GPU stream
magma_queue_t stream[2];
magma_event_t event[1];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
magma_event_create( &event[0] );
magmablasSetKernelStream(stream[0]);
magma_dscal( dofs, c_zero, x->val, 1 ); // x = 0
magma_dcopy( dofs, b.val, 1, r.val, 1 ); // r = b
nom0 = betanom = magma_dnrm2( dofs, r.val, 1 ); // nom0= || r||
nom = nom0 * nom0;
solver_par->init_res = nom0;
H(1,0) = MAGMA_D_MAKE( nom0, 0. );
magma_dsetvector(1, &H(1,0), 1, &dH(1,0), 1);
if ( (r0 = nom * solver_par->epsilon) < ATOLERANCE )
r0 = ATOLERANCE;
if ( nom < r0 )
return MAGMA_SUCCESS;
//Chronometry
real_Double_t tempo1, tempo2;
magma_device_sync(); tempo1=magma_wtime();
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_dcopy(dofs, r.val, 1, q(0), 1); // q[0] = 1.0/H(1,0) r
magma_dscal(dofs, 1./H(1,0), q(0), 1); // (to be fused)
for(k=1; k<=restart; k++) {
q_t.val = q(k-1);
magmablasSetKernelStream(stream[0]);
// preconditioner
// z[k] = M^(-1) q(k)
magma_d_applyprecond_left( A, q_t, &t, precond_par );
magma_d_applyprecond_right( A, t, &z_t, precond_par );
magma_dcopy(dofs, z_t.val, 1, z(k-1), 1);
// r = A q[k]
magma_d_spmv( c_one, A, z_t, c_zero, r );
if (solver_par->ortho == Magma_MGS ) {
// modified Gram-Schmidt
magmablasSetKernelStream(stream[0]);
for (i=1; i<=k; i++) {
H(i,k) =magma_ddot(dofs, q(i-1), 1, r.val, 1);
// H(i,k) = q[i] . r
magma_daxpy(dofs,-H(i,k), q(i-1), 1, r.val, 1);
// r = r - H(i,k) q[i]
}
H(k+1,k) = MAGMA_D_MAKE( magma_dnrm2(dofs, r.val, 1), 0. );
// H(k+1,k) = sqrt(r . r)
if (k < restart) {
magma_dcopy(dofs, r.val, 1, q(k), 1);
// q[k] = 1.0/H[k][k-1] r
magma_dscal(dofs, 1./H(k+1,k), q(k), 1);
// (to be fused)
}
} else if (solver_par->ortho == Magma_FUSED_CGS ) {
// fusing dgemv with dnrm2 in classical Gram-Schmidt
magmablasSetKernelStream(stream[0]);
magma_dcopy(dofs, r.val, 1, q(k), 1);
// dH(1:k+1,k) = q[0:k] . r
magmablas_dgemv(MagmaTrans, dofs, k+1, c_one, q(0),
dofs, r.val, 1, c_zero, &dH(1,k), 1);
// r = r - q[0:k-1] dH(1:k,k)
magmablas_dgemv(MagmaNoTrans, dofs, k, c_mone, q(0),
dofs, &dH(1,k), 1, c_one, r.val, 1);
// 1) dH(k+1,k) = sqrt( dH(k+1,k) - dH(1:k,k) )
magma_dcopyscale( dofs, k, r.val, q(k), &dH(1,k) );
// 2) q[k] = q[k] / dH(k+1,k)
magma_event_record( event[0], stream[0] );
magma_queue_wait_event( stream[1], event[0] );
magma_dgetvector_async(k+1, &dH(1,k), 1, &H(1,k), 1, stream[1]);
// asynch copy dH(1:(k+1),k) to H(1:(k+1),k)
} else {
// classical Gram-Schmidt (default)
// > explicitly calling magmabls
magmablasSetKernelStream(stream[0]);
magmablas_dgemv(MagmaTrans, dofs, k, c_one, q(0),
dofs, r.val, 1, c_zero, &dH(1,k), 1);
// dH(1:k,k) = q[0:k-1] . r
#ifndef DNRM2SCALE
// start copying dH(1:k,k) to H(1:k,k)
magma_event_record( event[0], stream[0] );
magma_queue_wait_event( stream[1], event[0] );
magma_dgetvector_async(k, &dH(1,k), 1, &H(1,k),
1, stream[1]);
#endif
// r = r - q[0:k-1] dH(1:k,k)
magmablas_dgemv(MagmaNoTrans, dofs, k, c_mone, q(0),
dofs, &dH(1,k), 1, c_one, r.val, 1);
#ifdef DNRM2SCALE
magma_dcopy(dofs, r.val, 1, q(k), 1);
// q[k] = r / H(k,k-1)
magma_dnrm2scale(dofs, q(k), dofs, &dH(k+1,k) );
// dH(k+1,k) = sqrt(r . r) and r = r / dH(k+1,k)
magma_event_record( event[0], stream[0] );
// start sending dH(1:k,k) to H(1:k,k)
magma_queue_wait_event( stream[1], event[0] );
// can we keep H(k+1,k) on GPU and combine?
magma_dgetvector_async(k+1, &dH(1,k), 1, &H(1,k), 1, stream[1]);
#else
H(k+1,k) = MAGMA_D_MAKE( magma_dnrm2(dofs, r.val, 1), 0. );
// H(k+1,k) = sqrt(r . r)
if( k<solver_par->restart ){
magmablasSetKernelStream(stream[0]);
magma_dcopy(dofs, r.val, 1, q(k), 1);
// q[k] = 1.0/H[k][k-1] r
magma_dscal(dofs, 1./H(k+1,k), q(k), 1);
// (to be fused)
}
#endif
}
}
magma_queue_sync( stream[1] );
for( k=1; k<=restart; k++ ){
/* Minimization of || b-Ax || in H_k */
for (i=1; i<=k; i++) {
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_ddot_sub( i+1, &H(1,k), 1, &H(1,i), 1, &HH(k,i) );
#else
HH(k,i) = cblas_ddot(i+1, &H(1,k), 1, &H(1,i), 1);
#endif
}
h1[k] = H(1,k)*H(1,0);
if (k != 1)
for (i=1; i<k; i++) {
for (m=i+1; m<k; m++){
HH(k,m) -= HH(k,i) * HH(m,i);
}
HH(k,k) -= HH(k,i) * HH(k,i) / HH(i,i);
HH(k,i) = HH(k,i)/HH(i,i);
h1[k] -= h1[i] * HH(k,i);
}
y[k] = h1[k]/HH(k,k);
if (k != 1)
for (i=k-1; i>=1; i--) {
y[i] = h1[i]/HH(i,i);
for (j=i+1; j<=k; j++)
y[i] -= y[j] * HH(j,i);
}
m = k;
rNorm = fabs(MAGMA_D_REAL(H(k+1,k)));
}
magma_dsetmatrix_async(m, 1, y+1, m, dy, m, stream[0]);
magmablasSetKernelStream(stream[0]);
magma_dgemv(MagmaNoTrans, dofs, m, c_one, z(0), dofs, dy, 1,
c_one, x->val, 1);
magma_d_spmv( c_mone, A, *x, c_zero, r ); // r = - A * x
magma_daxpy(dofs, c_one, b.val, 1, r.val, 1); // r = r + b
H(1,0) = MAGMA_D_MAKE( magma_dnrm2(dofs, r.val, 1), 0. );
// RNorm = H[1][0] = || r ||
RNorm = MAGMA_D_REAL( H(1,0) );
betanom = fabs(RNorm);
if( solver_par->verbose > 0 ){
magma_device_sync(); tempo2=magma_wtime();
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;
}
}
magma_device_sync(); tempo2=magma_wtime();
solver_par->runtime = (real_Double_t) tempo2-tempo1;
double residual;
magma_dresidual( A, b, *x, &residual );
solver_par->iter_res = betanom;
solver_par->final_res = residual;
if( solver_par->numiter < solver_par->maxiter){
solver_par->info = 0;
}else if( solver_par->init_res > solver_par->final_res ){
if( solver_par->verbose > 0 ){
if( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
solver_par->info = -2;
}
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 = -1;
}
// free pinned memory
magma_free_pinned( H );
magma_free_pinned( y );
magma_free_pinned( HH );
magma_free_pinned( h1 );
// free GPU memory
magma_free(dy);
if (dH != NULL ) magma_free(dH);
magma_d_vfree(&t);
magma_d_vfree(&r);
magma_d_vfree(&q);
magma_d_vfree(&z);
magma_d_vfree(&z_t);
// free GPU streams and events
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_event_destroy( event[0] );
magmablasSetKernelStream(NULL);
return MAGMA_SUCCESS;
} /* magma_dgmres */
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_dpbicgstab(
magma_d_sparse_matrix A, magma_d_vector b, magma_d_vector *x,
magma_d_solver_par *solver_par,
magma_d_preconditioner *precond_par,
magma_queue_t queue )
{
// set queue for old dense routines
magma_queue_t orig_queue;
magmablasGetKernelStream( &orig_queue );
// prepare solver feedback
solver_par->solver = Magma_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;
// workspace
magma_d_vector r,rr,p,v,s,t,ms,mt,y,z;
magma_d_vinit( &r, Magma_DEV, dofs, c_zero, queue );
magma_d_vinit( &rr, Magma_DEV, dofs, c_zero, queue );
magma_d_vinit( &p, Magma_DEV, dofs, c_zero, queue );
magma_d_vinit( &v, Magma_DEV, dofs, c_zero, queue );
magma_d_vinit( &s, Magma_DEV, dofs, c_zero, queue );
magma_d_vinit( &t, Magma_DEV, dofs, c_zero, queue );
magma_d_vinit( &ms, Magma_DEV, dofs, c_zero, queue );
magma_d_vinit( &mt, Magma_DEV, dofs, c_zero, queue );
magma_d_vinit( &y, Magma_DEV, dofs, c_zero, queue );
magma_d_vinit( &z, Magma_DEV, dofs, c_zero, queue );
// solver variables
double alpha, beta, omega, rho_old, rho_new;
double nom, betanom, nom0, r0, den, res;
// solver setup
magma_dscal( dofs, c_zero, x->dval, 1) ; // x = 0
magma_dcopy( dofs, b.dval, 1, r.dval, 1 ); // r = b
magma_dcopy( dofs, b.dval, 1, rr.dval, 1 ); // rr = b
nom0 = betanom = magma_dnrm2( dofs, r.dval, 1 ); // nom = || r ||
nom = nom0*nom0;
rho_new = omega = alpha = MAGMA_D_MAKE( 1.0, 0. );
solver_par->init_res = nom0;
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 ) {
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
}
// check positive definite
if (den <= 0.0) {
printf("Operator A is not postive definite. (Ar,r) = %f\n", den);
magmablasSetKernelStream( orig_queue );
return MAGMA_NONSPD;
}
//Chronometry
real_Double_t tempo1, tempo2;
tempo1 = magma_sync_wtime( queue );
if ( solver_par->verbose > 0 ) {
solver_par->res_vec[0] = nom0;
solver_par->timing[0] = 0.0;
}
// start iteration
for( solver_par->numiter= 1; solver_par->numiter<solver_par->maxiter;
solver_par->numiter++ ) {
rho_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
magma_d_applyprecond_left( A, p, &mt, precond_par, queue );
magma_d_applyprecond_right( A, mt, &y, precond_par, queue );
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
magma_d_applyprecond_left( A, s, &ms, precond_par, queue );
magma_d_applyprecond_right( A, ms, &z, precond_par, queue );
magma_d_spmv( c_one, A, z, c_zero, t, queue ); // t=As
// preconditioner
magma_d_applyprecond_left( A, s, &ms, precond_par, queue );
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;
}
}
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t) tempo2-tempo1;
double residual;
magma_dresidual( A, b, *x, &residual, queue );
solver_par->final_res = residual;
solver_par->iter_res = res;
if ( solver_par->numiter < solver_par->maxiter) {
solver_par->info = MAGMA_SUCCESS;
} else if ( solver_par->init_res > solver_par->final_res ) {
if ( solver_par->verbose > 0 ) {
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
solver_par->info = MAGMA_SLOW_CONVERGENCE;
}
else {
if ( solver_par->verbose > 0 ) {
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
solver_par->info = MAGMA_DIVERGENCE;
}
magma_d_vfree(&r, queue );
magma_d_vfree(&rr, queue );
magma_d_vfree(&p, queue );
magma_d_vfree(&v, queue );
magma_d_vfree(&s, queue );
magma_d_vfree(&t, queue );
magma_d_vfree(&ms, queue );
magma_d_vfree(&mt, queue );
magma_d_vfree(&y, queue );
magma_d_vfree(&z, queue );
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
} /* magma_dbicgstab */
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 */
magma_int_t
magma_dpcg( magma_d_sparse_matrix A, magma_d_vector b, magma_d_vector *x,
magma_d_solver_par *solver_par,
magma_d_preconditioner *precond_par ){
// prepare solver feedback
solver_par->solver = Magma_PCG;
solver_par->numiter = 0;
solver_par->info = 0;
// local variables
double c_zero = MAGMA_D_ZERO, c_one = MAGMA_D_ONE;
magma_int_t dofs = A.num_rows;
// GPU workspace
magma_d_vector r, rt, p, q, h;
magma_d_vinit( &r, Magma_DEV, dofs, c_zero );
magma_d_vinit( &rt, Magma_DEV, dofs, c_zero );
magma_d_vinit( &p, Magma_DEV, dofs, c_zero );
magma_d_vinit( &q, Magma_DEV, dofs, c_zero );
magma_d_vinit( &h, Magma_DEV, dofs, c_zero );
// solver variables
double alpha, beta;
double nom, nom0, r0, gammaold, gammanew, den, res;
// solver setup
magma_dscal( dofs, c_zero, x->val, 1) ; // x = 0
magma_dcopy( dofs, b.val, 1, r.val, 1 ); // r = b
// preconditioner
magma_d_applyprecond_left( A, r, &rt, precond_par );
magma_d_applyprecond_right( A, rt, &h, precond_par );
magma_dcopy( dofs, h.val, 1, p.val, 1 ); // p = h
nom = MAGMA_D_REAL( magma_ddot(dofs, r.val, 1, h.val, 1) );
nom0 = magma_dnrm2( dofs, r.val, 1 );
magma_d_spmv( c_one, A, p, c_zero, q ); // q = A p
den = MAGMA_D_REAL( magma_ddot(dofs, p.val, 1, q.val, 1) );// den = p dot q
solver_par->init_res = nom0;
if ( (r0 = nom * solver_par->epsilon) < ATOLERANCE )
r0 = ATOLERANCE;
if ( nom < r0 )
return MAGMA_SUCCESS;
// check positive definite
if (den <= 0.0) {
printf("Operator A is not postive definite. (Ar,r) = %f\n", den);
return -100;
}
//Chronometry
real_Double_t tempo1, tempo2;
magma_device_sync(); tempo1=magma_wtime();
if( solver_par->verbose > 0 ){
solver_par->res_vec[0] = (real_Double_t)nom0;
solver_par->timing[0] = 0.0;
}
// start iteration
for( solver_par->numiter= 1; solver_par->numiter<solver_par->maxiter;
solver_par->numiter++ ){
// preconditioner
magma_d_applyprecond_left( A, r, &rt, precond_par );
magma_d_applyprecond_right( A, rt, &h, precond_par );
gammanew = MAGMA_D_REAL( magma_ddot(dofs, r.val, 1, h.val, 1) );
// gn = < r,h>
if( solver_par->numiter==1 ){
magma_dcopy( dofs, h.val, 1, p.val, 1 ); // p = h
}else{
beta = MAGMA_D_MAKE(gammanew/gammaold, 0.); // beta = gn/go
magma_dscal(dofs, beta, p.val, 1); // p = beta*p
magma_daxpy(dofs, c_one, h.val, 1, p.val, 1); // p = p + h
}
magma_d_spmv( c_one, A, p, c_zero, q ); // q = A p
den = MAGMA_D_REAL(magma_ddot(dofs, p.val, 1, q.val, 1));
// den = p dot q
alpha = MAGMA_D_MAKE(gammanew/den, 0.);
magma_daxpy(dofs, alpha, p.val, 1, x->val, 1); // x = x + alpha p
magma_daxpy(dofs, -alpha, q.val, 1, r.val, 1); // r = r - alpha q
gammaold = gammanew;
res = magma_dnrm2( dofs, r.val, 1 );
if( solver_par->verbose > 0 ){
magma_device_sync(); tempo2=magma_wtime();
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;
}
}
magma_device_sync(); tempo2=magma_wtime();
solver_par->runtime = (real_Double_t) tempo2-tempo1;
double residual;
magma_dresidual( A, b, *x, &residual );
solver_par->iter_res = res;
solver_par->final_res = residual;
if( solver_par->numiter < solver_par->maxiter){
solver_par->info = 0;
}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;
}
}
solver_par->info = -2;
}
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;
}
}
solver_par->info = -1;
}
magma_d_vfree(&r);
magma_d_vfree(&rt);
magma_d_vfree(&p);
magma_d_vfree(&q);
magma_d_vfree(&h);
return MAGMA_SUCCESS;
} /* magma_dcg */
/* ////////////////////////////////////////////////////////////////////////////
-- 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_dbpcg(
magma_d_sparse_matrix A, magma_d_vector b, magma_d_vector *x,
magma_d_solver_par *solver_par,
magma_d_preconditioner *precond_par,
magma_queue_t queue )
{
// set queue for old dense routines
magma_queue_t orig_queue;
magmablasGetKernelStream( &orig_queue );
magma_int_t stat_dev = 0, stat_cpu = 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->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_vector r, rt, p, q, h;
magma_d_vinit( &r, Magma_DEV, dofs*num_vecs, c_zero, queue );
magma_d_vinit( &rt, Magma_DEV, dofs*num_vecs, c_zero, queue );
magma_d_vinit( &p, Magma_DEV, dofs*num_vecs, c_zero, queue );
magma_d_vinit( &q, Magma_DEV, dofs*num_vecs, c_zero, queue );
magma_d_vinit( &h, Magma_DEV, dofs*num_vecs, c_zero, queue );
// solver variables
double *alpha, *beta;
alpha = NULL;
beta = NULL;
stat_cpu += magma_dmalloc_cpu(&alpha, num_vecs);
stat_cpu += magma_dmalloc_cpu(&beta, num_vecs);
double *nom, *nom0, *r0, *gammaold, *gammanew, *den, *res, *residual;
nom = NULL;
nom0 = NULL;
r0 = NULL;
gammaold = NULL;
gammanew = NULL;
den = NULL;
res = NULL;
residual = NULL;
stat_cpu += magma_dmalloc_cpu(&residual, num_vecs);
stat_cpu += magma_dmalloc_cpu(&nom, num_vecs);
stat_cpu += magma_dmalloc_cpu(&nom0, num_vecs);
stat_cpu += magma_dmalloc_cpu(&r0, num_vecs);
stat_cpu += magma_dmalloc_cpu(&gammaold, num_vecs);
stat_cpu += magma_dmalloc_cpu(&gammanew, num_vecs);
stat_cpu += magma_dmalloc_cpu(&den, num_vecs);
stat_cpu += magma_dmalloc_cpu(&res, num_vecs);
stat_cpu += magma_dmalloc_cpu(&residual, num_vecs);
if( stat_cpu != 0 ){
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 );
magma_free_cpu( alpha );
magma_free_cpu( beta );
magma_free_cpu( residual );
magmablasSetKernelStream( orig_queue );
printf("error: memory allocation.\n");
return MAGMA_ERR_HOST_ALLOC;
}
// solver setup
magma_dscal( dofs*num_vecs, c_zero, x->dval, 1) ; // x = 0
magma_dcopy( dofs*num_vecs, b.dval, 1, r.dval, 1 ); // r = b
// preconditioner
magma_d_applyprecond_left( A, r, &rt, precond_par, queue );
magma_d_applyprecond_right( A, rt, &h, precond_par, queue );
magma_dcopy( dofs*num_vecs, h.dval, 1, p.dval, 1 ); // p = h
for( i=0; i<num_vecs; i++) {
nom[i] = MAGMA_D_REAL( magma_ddot(dofs, r(i), 1, h(i), 1) );
nom0[i] = magma_dnrm2( dofs, r(i), 1 );
}
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) ); // den = p dot q
solver_par->init_res = nom0[0];
if ( (r0[0] = nom[0] * solver_par->epsilon) < 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]);
magmablasSetKernelStream( orig_queue );
return MAGMA_NONSPD;
solver_par->info = MAGMA_NONSPD;;
}
if ( nom[0] < r0[0] ) {
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
}
//Chronometry
real_Double_t tempo1, tempo2;
tempo1 = magma_sync_wtime( queue );
if ( solver_par->verbose > 0 ) {
solver_par->res_vec[0] = (real_Double_t)nom0[0];
solver_par->timing[0] = 0.0;
}
// start iteration
for( solver_par->numiter= 1; solver_par->numiter<solver_par->maxiter;
solver_par->numiter++ ) {
// preconditioner
magma_d_applyprecond_left( A, r, &rt, precond_par, queue );
magma_d_applyprecond_right( 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) ); // gn = < r,h>
if ( solver_par->numiter==1 ) {
magma_dcopy( dofs*num_vecs, h.dval, 1, p.dval, 1 ); // 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); // p = beta*p
magma_daxpy(dofs, c_one, h(i), 1, p(i), 1); // p = p + h
}
}
magma_d_spmv( c_one, A, p, c_zero, q, queue ); // q = A p
// 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));
// 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); // x = x + alpha p
magma_daxpy(dofs, -alpha[i], q(i), 1, r(i), 1); // r = r - alpha q
gammaold[i] = gammanew[i];
res[i] = magma_dnrm2( dofs, r(i), 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) res[0];
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
if ( res[0]/nom0[0] < solver_par->epsilon ) {
break;
}
}
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t) tempo2-tempo1;
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;
}
}
solver_par->info = MAGMA_SLOW_CONVERGENCE;
}
else {
if ( solver_par->verbose > 0 ) {
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) res[0];
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
solver_par->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");
magma_d_vfree(&r, queue );
magma_d_vfree(&rt, queue );
magma_d_vfree(&p, queue );
magma_d_vfree(&q, queue );
magma_d_vfree(&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);
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
} /* magma_dbpcg */
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 */
magma_int_t
magma_dcg_merge( magma_d_sparse_matrix A, magma_d_vector b, magma_d_vector *x,
magma_d_solver_par *solver_par ){
// prepare solver feedback
solver_par->solver = Magma_CGMERGE;
solver_par->numiter = 0;
solver_par->info = 0;
// some useful variables
double c_zero = MAGMA_D_ZERO, c_one = MAGMA_D_ONE;
magma_int_t dofs = A.num_rows;
// GPU stream
magma_queue_t stream[2];
magma_event_t event[1];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
magma_event_create( &event[0] );
// GPU workspace
magma_d_vector r, d, z;
magma_d_vinit( &r, Magma_DEV, dofs, c_zero );
magma_d_vinit( &d, Magma_DEV, dofs, c_zero );
magma_d_vinit( &z, Magma_DEV, dofs, c_zero );
double *d1, *d2, *skp;
magma_dmalloc( &d1, dofs*(1) );
magma_dmalloc( &d2, dofs*(1) );
// array for the parameters
magma_dmalloc( &skp, 6 ); // skp = [alpha|beta|gamma|rho|tmp1|tmp2]
// solver variables
double alpha, beta, gamma, rho, tmp1, *skp_h;
double nom, nom0, r0, betanom, den;
// solver setup
magma_dscal( dofs, c_zero, x->val, 1) ; // x = 0
magma_dcopy( dofs, b.val, 1, r.val, 1 ); // r = b
magma_dcopy( dofs, b.val, 1, d.val, 1 ); // d = b
nom0 = betanom = magma_dnrm2( dofs, r.val, 1 );
nom = nom0 * nom0; // nom = r' * r
magma_d_spmv( c_one, A, d, c_zero, z ); // z = A d
den = MAGMA_D_REAL( magma_ddot(dofs, d.val, 1, z.val, 1) ); // den = d'* z
solver_par->init_res = nom0;
// array on host for the parameters
magma_dmalloc_cpu( &skp_h, 6 );
alpha = rho = gamma = tmp1 = c_one;
beta = magma_ddot(dofs, r.val, 1, r.val, 1);
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 );
if ( (r0 = nom * solver_par->epsilon) < ATOLERANCE )
r0 = ATOLERANCE;
if ( nom < r0 )
return MAGMA_SUCCESS;
// check positive definite
if (den <= 0.0) {
printf("Operator A is not postive definite. (Ar,r) = %f\n", den);
return -100;
}
//Chronometry
real_Double_t tempo1, tempo2;
magma_device_sync(); tempo1=magma_wtime();
if( solver_par->verbose > 0 ){
solver_par->res_vec[0] = (real_Double_t) nom0;
solver_par->timing[0] = 0.0;
}
// start iteration
for( solver_par->numiter= 1; solver_par->numiter<solver_par->maxiter;
solver_par->numiter++ ){
magmablasSetKernelStream(stream[0]);
// computes SpMV and dot product
magma_dcgmerge_spmv1( A, d1, d2, d.val, z.val, skp );
// updates x, r, computes scalars and updates d
magma_dcgmerge_xrbeta( dofs, d1, d2, x->val, r.val, d.val, z.val, skp );
// check stopping criterion (asynchronous copy)
magma_dgetvector_async( 1 , skp+1, 1,
skp_h+1, 1, stream[1] );
betanom = sqrt(MAGMA_D_REAL(skp_h[1]));
if( solver_par->verbose > 0 ){
magma_device_sync(); tempo2=magma_wtime();
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;
}
}
magma_device_sync(); tempo2=magma_wtime();
solver_par->runtime = (real_Double_t) tempo2-tempo1;
double residual;
magma_dresidual( A, b, *x, &residual );
solver_par->iter_res = betanom;
solver_par->final_res = residual;
if( solver_par->numiter < solver_par->maxiter){
solver_par->info = 0;
}else if( solver_par->init_res > solver_par->final_res ){
if( solver_par->verbose > 0 ){
if( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
solver_par->info = -2;
}
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 = -1;
}
magma_d_vfree(&r);
magma_d_vfree(&z);
magma_d_vfree(&d);
magma_free( d1 );
magma_free( d2 );
magma_free( skp );
magma_free_cpu( skp_h );
return MAGMA_SUCCESS;
} /* magma_dcg_merge */
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_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_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 */