extern "C" magma_int_t
magma_sbicg(
magma_s_matrix A, magma_s_matrix b, magma_s_matrix *x,
magma_s_solver_par *solver_par,
magma_queue_t queue )
{
magma_int_t info = MAGMA_NOTCONVERGED;
// prepare solver feedback
solver_par->solver = Magma_BICG;
solver_par->numiter = 0;
solver_par->spmv_count = 0;
// some useful variables
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
magma_int_t dofs = A.num_rows * b.num_cols;
// workspace
magma_s_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_s_matrix AT={Magma_CSR}, Ah1={Magma_CSR}, Ah2={Magma_CSR};
CHECK( magma_svinit( &r, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &rt,Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &p, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &pt,Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &q, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &qt,Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &y, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &yt,Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &z, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &zt,Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
// solver variables
float alpha, rho, beta, rho_new, ptq;
float res, nomb, nom0, r0;
// transpose the matrix
magma_smtransfer( A, &Ah1, Magma_DEV, Magma_CPU, queue );
magma_smconvert( Ah1, &Ah2, A.storage_type, Magma_CSR, queue );
magma_smfree(&Ah1, queue );
magma_smtransposeconjugate( Ah2, &Ah1, queue );
magma_smfree(&Ah2, queue );
Ah2.blocksize = A.blocksize;
Ah2.alignment = A.alignment;
magma_smconvert( Ah1, &Ah2, Magma_CSR, A.storage_type, queue );
magma_smfree(&Ah1, queue );
magma_smtransfer( Ah2, &AT, Magma_CPU, Magma_DEV, queue );
magma_smfree(&Ah2, queue );
// solver setup
CHECK( magma_sresidualvec( A, b, *x, &r, &nom0, queue));
res = nom0;
solver_par->init_res = nom0;
magma_scopy( dofs, r.dval, 1, rt.dval, 1, queue ); // rr = r
rho_new = magma_sdot( dofs, rt.dval, 1, r.dval, 1, queue ); // rho=<rr,r>
rho = alpha = MAGMA_S_MAKE( 1.0, 0. );
nomb = magma_snrm2( 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++;
magma_scopy( dofs, r.dval, 1 , y.dval, 1, queue ); // y=r
magma_scopy( dofs, y.dval, 1 , z.dval, 1, queue ); // z=y
magma_scopy( dofs, rt.dval, 1 , yt.dval, 1, queue ); // yt=rt
magma_scopy( dofs, yt.dval, 1 , zt.dval, 1, queue ); // zt=yt
rho= rho_new;
rho_new = magma_sdot( dofs, rt.dval, 1, z.dval, 1, queue ); // rho=<rt,z>
if( magma_s_isnan_inf( rho_new ) ){
info = MAGMA_DIVERGENCE;
break;
}
if( solver_par->numiter==1 ){
magma_scopy( dofs, z.dval, 1 , p.dval, 1, queue ); // yt=rt
magma_scopy( dofs, zt.dval, 1 , pt.dval, 1, queue ); // zt=yt
} else {
beta = rho_new/rho;
magma_sscal( dofs, beta, p.dval, 1, queue ); // p = beta*p
magma_saxpy( dofs, c_one , z.dval, 1 , p.dval, 1, queue ); // p = z+beta*p
magma_sscal( dofs, MAGMA_S_CONJ(beta), pt.dval, 1, queue ); // pt = beta*pt
magma_saxpy( dofs, c_one , zt.dval, 1 , pt.dval, 1, queue ); // pt = zt+beta*pt
}
CHECK( magma_s_spmv( c_one, A, p, c_zero, q, queue )); // v = Ap
CHECK( magma_s_spmv( c_one, AT, pt, c_zero, qt, queue )); // v = Ap
solver_par->spmv_count++;
solver_par->spmv_count++;
ptq = magma_sdot( dofs, pt.dval, 1, q.dval, 1, queue );
alpha = rho_new /ptq;
magma_saxpy( dofs, alpha, p.dval, 1 , x->dval, 1, queue ); // x=x+alpha*p
magma_saxpy( dofs, c_neg_one * alpha, q.dval, 1 , r.dval, 1, queue ); // r=r+alpha*q
magma_saxpy( dofs, c_neg_one * MAGMA_S_CONJ(alpha), qt.dval, 1 , rt.dval, 1, queue ); // r=r+alpha*q
res = magma_snrm2( 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;
float residual;
CHECK( magma_sresidualvec( 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_smfree(&r, queue );
magma_smfree(&rt, queue );
magma_smfree(&p, queue );
magma_smfree(&pt, queue );
magma_smfree(&q, queue );
magma_smfree(&qt, queue );
magma_smfree(&y, queue );
magma_smfree(&yt, queue );
magma_smfree(&z, queue );
magma_smfree(&zt, queue );
magma_smfree(&AT, queue );
magma_smfree(&Ah1, queue );
magma_smfree(&Ah2, queue );
solver_par->info = info;
return info;
} /* magma_sbicg */
magma_int_t
magma_sbicgstab( magma_s_sparse_matrix A, magma_s_vector b, magma_s_vector *x,
magma_s_solver_par *solver_par ){
// prepare solver feedback
solver_par->solver = Magma_BICGSTAB;
solver_par->numiter = 0;
solver_par->info = 0;
// some useful variables
float c_zero = MAGMA_S_ZERO, c_one = MAGMA_S_ONE,
c_mone = MAGMA_S_NEG_ONE;
magma_int_t dofs = A.num_rows;
// workspace
magma_s_vector r,rr,p,v,s,t;
magma_s_vinit( &r, Magma_DEV, dofs, c_zero );
magma_s_vinit( &rr, Magma_DEV, dofs, c_zero );
magma_s_vinit( &p, Magma_DEV, dofs, c_zero );
magma_s_vinit( &v, Magma_DEV, dofs, c_zero );
magma_s_vinit( &s, Magma_DEV, dofs, c_zero );
magma_s_vinit( &t, Magma_DEV, dofs, c_zero );
// solver variables
float alpha, beta, omega, rho_old, rho_new;
float nom, betanom, nom0, r0, den, res;
// solver setup
magma_sscal( dofs, c_zero, x->val, 1) ; // x = 0
magma_scopy( dofs, b.val, 1, r.val, 1 ); // r = b
magma_scopy( dofs, b.val, 1, rr.val, 1 ); // rr = b
nom0 = betanom = magma_snrm2( dofs, r.val, 1 ); // nom = || r ||
nom = nom0*nom0;
rho_old = omega = alpha = MAGMA_S_MAKE( 1.0, 0. );
solver_par->init_res = nom0;
magma_s_spmv( c_one, A, r, c_zero, v ); // z = A r
den = MAGMA_S_REAL( magma_sdot(dofs, v.val, 1, r.val, 1) ); // den = z' * 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++ ){
rho_new = magma_sdot( dofs, rr.val, 1, r.val, 1 ); // rho=<rr,r>
beta = rho_new/rho_old * alpha/omega; // beta=rho/rho_old *alpha/omega
magma_sscal( dofs, beta, p.val, 1 ); // p = beta*p
magma_saxpy( dofs, c_mone * omega * beta, v.val, 1 , p.val, 1 );
// p = p-omega*beta*v
magma_saxpy( dofs, c_one, r.val, 1, p.val, 1 ); // p = p+r
magma_s_spmv( c_one, A, p, c_zero, v ); // v = Ap
alpha = rho_new / magma_sdot( dofs, rr.val, 1, v.val, 1 );
magma_scopy( dofs, r.val, 1 , s.val, 1 ); // s=r
magma_saxpy( dofs, c_mone * alpha, v.val, 1 , s.val, 1 ); // s=s-alpha*v
magma_s_spmv( c_one, A, s, c_zero, t ); // t=As
omega = magma_sdot( dofs, t.val, 1, s.val, 1 ) // omega = <s,t>/<t,t>
/ magma_sdot( dofs, t.val, 1, t.val, 1 );
magma_saxpy( dofs, alpha, p.val, 1 , x->val, 1 ); // x=x+alpha*p
magma_saxpy( dofs, omega, s.val, 1 , x->val, 1 ); // x=x+omega*s
magma_scopy( dofs, s.val, 1 , r.val, 1 ); // r=s
magma_saxpy( dofs, c_mone * omega, t.val, 1 , r.val, 1 ); // r=r-omega*t
res = betanom = magma_snrm2( dofs, r.val, 1 );
nom = betanom*betanom;
rho_old = rho_new; // rho_old=rho
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;
float residual;
magma_sresidual( 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) 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_s_vfree(&r);
magma_s_vfree(&rr);
magma_s_vfree(&p);
magma_s_vfree(&v);
magma_s_vfree(&s);
magma_s_vfree(&t);
return MAGMA_SUCCESS;
} /* magma_sbicgstab */
magma_int_t
magma_sbicgstab_merge( magma_s_sparse_matrix A, magma_s_vector b,
magma_s_vector *x, magma_s_solver_par *solver_par ){
// prepare solver feedback
solver_par->solver = Magma_BICGSTABMERGE;
solver_par->numiter = 0;
solver_par->info = 0;
// some useful variables
float c_zero = MAGMA_S_ZERO, c_one = MAGMA_S_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_s_vector q, r,rr,p,v,s,t;
float *d1, *d2, *skp;
magma_smalloc( &d1, dofs*(2) );
magma_smalloc( &d2, dofs*(2) );
// array for the parameters
magma_smalloc( &skp, 8 );
// skp = [alpha|beta|omega|rho_old|rho|nom|tmp1|tmp2]
magma_s_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
float alpha, beta, omega, rho_old, rho_new, *skp_h;
float nom, nom0, betanom, r0, den;
// solver setup
magma_sscal( dofs, c_zero, x->val, 1) ; // x = 0
magma_scopy( dofs, b.val, 1, q(0), 1 ); // rr = b
magma_scopy( dofs, b.val, 1, q(1), 1 ); // r = b
rho_new = magma_sdot( dofs, r.val, 1, r.val, 1 ); // rho=<rr,r>
nom = MAGMA_S_REAL(magma_sdot( dofs, r.val, 1, r.val, 1 ));
nom0 = betanom = sqrt(nom); // nom = || r ||
rho_old = omega = alpha = MAGMA_S_MAKE( 1.0, 0. );
beta = rho_new;
solver_par->init_res = nom0;
// array on host for the parameters
magma_smalloc_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_S_MAKE(nom, 0.0);
magma_ssetvector( 8, skp_h, 1, skp, 1 );
magma_s_spmv( c_one, A, r, c_zero, v ); // z = A r
den = MAGMA_S_REAL( magma_sdot(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_sbicgmerge1( dofs, skp, v.val, r.val, p.val );
magma_s_spmv( c_one, A, p, c_zero, v ); // v = Ap
magma_smdotc( dofs, 1, q.val, v.val, d1, d2, skp );
magma_sbicgmerge4( 1, skp );
magma_sbicgmerge2( dofs, skp, r.val, v.val, s.val ); // s=r-alpha*v
magma_s_spmv( c_one, A, s, c_zero, t ); // t=As
magma_smdotc( dofs, 2, q.val+4*dofs, t.val, d1, d2, skp+6 );
magma_sbicgmerge4( 2, skp );
magma_sbicgmerge3( dofs, skp, p.val, s.val, // x=x+alpha*p+omega*s
t.val, x->val, r.val ); // r=s-omega*t
magma_smdotc( dofs, 2, q.val, r.val, d1, d2, skp+4);
magma_sbicgmerge4( 3, skp );
// check stopping criterion (asynchronous copy)
magma_sgetvector_async( 1 , skp+5, 1,
skp_h+5, 1, stream[1] );
betanom = sqrt(MAGMA_S_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;
float residual;
magma_sresidual( 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_s_vfree(&q); // frees all vectors
magma_free(d1);
magma_free(d2);
magma_free( skp );
magma_free_cpu( skp_h );
return MAGMA_SUCCESS;
} /* sbicgstab_merge */
extern "C" magma_int_t
magma_scg_merge(
magma_s_sparse_matrix A, magma_s_vector b, magma_s_vector *x,
magma_s_solver_par *solver_par,
magma_queue_t queue )
{
// set queue for old dense routines
magma_queue_t orig_queue;
magmablasGetKernelStream( &orig_queue );
// prepare solver feedback
solver_par->solver = Magma_CGMERGE;
solver_par->numiter = 0;
solver_par->info = MAGMA_SUCCESS;
// some useful variables
float c_zero = MAGMA_S_ZERO, c_one = MAGMA_S_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_s_vector r, d, z;
magma_s_vinit( &r, Magma_DEV, dofs, c_zero, queue );
magma_s_vinit( &d, Magma_DEV, dofs, c_zero, queue );
magma_s_vinit( &z, Magma_DEV, dofs, c_zero, queue );
float *d1, *d2, *skp;
d1 = NULL;
d2 = NULL;
skp = NULL;
magma_int_t stat_dev = 0, stat_cpu = 0;
stat_dev += magma_smalloc( &d1, dofs*(1) );
stat_dev += magma_smalloc( &d2, dofs*(1) );
// array for the parameters
stat_dev += magma_smalloc( &skp, 6 );
// skp = [alpha|beta|gamma|rho|tmp1|tmp2]
if( stat_dev != 0 ){
magma_free( d1 );
magma_free( d2 );
magma_free( skp );
printf("error: memory allocation.\n");
return MAGMA_ERR_DEVICE_ALLOC;
}
// solver variables
float alpha, beta, gamma, rho, tmp1, *skp_h;
float nom, nom0, r0, betanom, den;
// solver setup
magma_sscal( dofs, c_zero, x->dval, 1) ; // x = 0
magma_scopy( dofs, b.dval, 1, r.dval, 1 ); // r = b
magma_scopy( dofs, b.dval, 1, d.dval, 1 ); // d = b
nom0 = betanom = magma_snrm2( dofs, r.dval, 1 );
nom = nom0 * nom0; // nom = r' * r
magma_s_spmv( c_one, A, d, c_zero, z, queue ); // z = A d
den = MAGMA_S_REAL( magma_sdot(dofs, d.dval, 1, z.dval, 1) ); // den = d'* z
solver_par->init_res = nom0;
// array on host for the parameters
stat_cpu += magma_smalloc_cpu( &skp_h, 6 );
if( stat_cpu != 0 ){
magma_free( d1 );
magma_free( d2 );
magma_free( skp );
magma_free_cpu( skp_h );
printf("error: memory allocation.\n");
return MAGMA_ERR_HOST_ALLOC;
}
alpha = rho = gamma = tmp1 = c_one;
beta = magma_sdot(dofs, r.dval, 1, r.dval, 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_S_MAKE(nom, 0.0);
magma_ssetvector( 6, skp_h, 1, skp, 1 );
if ( (r0 = nom * solver_par->epsilon) < ATOLERANCE )
r0 = ATOLERANCE;
if ( nom < r0 ) {
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
}
// check positive definite
if (den <= 0.0) {
printf("Operator A is not postive definite. (Ar,r) = %f\n", den);
magmablasSetKernelStream( orig_queue );
return MAGMA_NONSPD;
solver_par->info = MAGMA_NONSPD;;
}
//Chronometry
real_Double_t tempo1, tempo2;
tempo1 = magma_sync_wtime( queue );
if ( solver_par->verbose > 0 ) {
solver_par->res_vec[0] = (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_scgmerge_spmv1( A, d1, d2, d.dval, z.dval, skp, queue );
// updates x, r, computes scalars and updates d
magma_scgmerge_xrbeta( dofs, d1, d2, x->dval, r.dval, d.dval, z.dval, skp, queue );
// check stopping criterion (asynchronous copy)
magma_sgetvector_async( 1 , skp+1, 1,
skp_h+1, 1, stream[1] );
betanom = sqrt(MAGMA_S_REAL(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 < r0 ) {
break;
}
}
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t) tempo2-tempo1;
float residual;
magma_sresidual( A, b, *x, &residual, queue );
solver_par->iter_res = betanom;
solver_par->final_res = residual;
if ( solver_par->numiter < solver_par->maxiter) {
solver_par->info = MAGMA_SUCCESS;
} else if ( solver_par->init_res > solver_par->final_res ) {
if ( solver_par->verbose > 0 ) {
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
solver_par->info = MAGMA_SLOW_CONVERGENCE;
}
else {
if ( solver_par->verbose > 0 ) {
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
solver_par->info = MAGMA_DIVERGENCE;
}
magma_s_vfree(&r, queue );
magma_s_vfree(&z, queue );
magma_s_vfree(&d, queue );
magma_free( d1 );
magma_free( d2 );
magma_free( skp );
magma_free_cpu( skp_h );
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
} /* magma_scg_merge */
magma_int_t
magma_spcg( magma_s_sparse_matrix A, magma_s_vector b, magma_s_vector *x,
magma_s_solver_par *solver_par,
magma_s_preconditioner *precond_par ){
// prepare solver feedback
solver_par->solver = Magma_PCG;
solver_par->numiter = 0;
solver_par->info = 0;
// local variables
float c_zero = MAGMA_S_ZERO, c_one = MAGMA_S_ONE;
magma_int_t dofs = A.num_rows;
// GPU workspace
magma_s_vector r, rt, p, q, h;
magma_s_vinit( &r, Magma_DEV, dofs, c_zero );
magma_s_vinit( &rt, Magma_DEV, dofs, c_zero );
magma_s_vinit( &p, Magma_DEV, dofs, c_zero );
magma_s_vinit( &q, Magma_DEV, dofs, c_zero );
magma_s_vinit( &h, Magma_DEV, dofs, c_zero );
// solver variables
float alpha, beta;
float nom, nom0, r0, gammaold, gammanew, den, res;
// solver setup
magma_sscal( dofs, c_zero, x->val, 1) ; // x = 0
magma_scopy( dofs, b.val, 1, r.val, 1 ); // r = b
// preconditioner
magma_s_applyprecond_left( A, r, &rt, precond_par );
magma_s_applyprecond_right( A, rt, &h, precond_par );
magma_scopy( dofs, h.val, 1, p.val, 1 ); // p = h
nom = MAGMA_S_REAL( magma_sdot(dofs, r.val, 1, h.val, 1) );
nom0 = magma_snrm2( dofs, r.val, 1 );
magma_s_spmv( c_one, A, p, c_zero, q ); // q = A p
den = MAGMA_S_REAL( magma_sdot(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_s_applyprecond_left( A, r, &rt, precond_par );
magma_s_applyprecond_right( A, rt, &h, precond_par );
gammanew = MAGMA_S_REAL( magma_sdot(dofs, r.val, 1, h.val, 1) );
// gn = < r,h>
if( solver_par->numiter==1 ){
magma_scopy( dofs, h.val, 1, p.val, 1 ); // p = h
}else{
beta = MAGMA_S_MAKE(gammanew/gammaold, 0.); // beta = gn/go
magma_sscal(dofs, beta, p.val, 1); // p = beta*p
magma_saxpy(dofs, c_one, h.val, 1, p.val, 1); // p = p + h
}
magma_s_spmv( c_one, A, p, c_zero, q ); // q = A p
den = MAGMA_S_REAL(magma_sdot(dofs, p.val, 1, q.val, 1));
// den = p dot q
alpha = MAGMA_S_MAKE(gammanew/den, 0.);
magma_saxpy(dofs, alpha, p.val, 1, x->val, 1); // x = x + alpha p
magma_saxpy(dofs, -alpha, q.val, 1, r.val, 1); // r = r - alpha q
gammaold = gammanew;
res = magma_snrm2( 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;
float residual;
magma_sresidual( 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_s_vfree(&r);
magma_s_vfree(&rt);
magma_s_vfree(&p);
magma_s_vfree(&q);
magma_s_vfree(&h);
return MAGMA_SUCCESS;
} /* magma_scg */
extern "C" magma_int_t
magma_spcg_merge(
magma_s_matrix A, magma_s_matrix b, magma_s_matrix *x,
magma_s_solver_par *solver_par,
magma_s_preconditioner *precond_par,
magma_queue_t queue )
{
magma_int_t info = MAGMA_NOTCONVERGED;
// prepare solver feedback
solver_par->solver = Magma_PCGMERGE;
solver_par->numiter = 0;
solver_par->spmv_count = 0;
// solver variables
float alpha, beta, gamma, rho, tmp1, *skp_h={0};
float nom, nom0, r0, res, nomb;
float den;
// some useful variables
float c_zero = MAGMA_S_ZERO, c_one = MAGMA_S_ONE;
magma_int_t dofs = A.num_rows*b.num_cols;
magma_s_matrix r={Magma_CSR}, d={Magma_CSR}, z={Magma_CSR}, h={Magma_CSR},
rt={Magma_CSR};
float *d1=NULL, *d2=NULL, *skp=NULL;
// GPU workspace
CHECK( magma_svinit( &r, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &d, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &z, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &rt, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &h, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_smalloc( &d1, dofs*(2) ));
CHECK( magma_smalloc( &d2, dofs*(2) ));
// array for the parameters
CHECK( magma_smalloc( &skp, 7 ));
// skp = [alpha|beta|gamma|rho|tmp1|tmp2|res]
// solver setup
CHECK( magma_sresidualvec( A, b, *x, &r, &nom0, queue));
// preconditioner
CHECK( magma_s_applyprecond_left( MagmaNoTrans, A, r, &rt, precond_par, queue ));
CHECK( magma_s_applyprecond_right( MagmaNoTrans, A, rt, &h, precond_par, queue ));
magma_scopy( dofs, h.dval, 1, d.dval, 1, queue );
nom = MAGMA_S_ABS( magma_sdot( dofs, r.dval, 1, h.dval, 1, queue ));
CHECK( magma_s_spmv( c_one, A, d, c_zero, z, queue )); // z = A d
den = magma_sdot( dofs, d.dval, 1, z.dval, 1, queue ); // den = d'* z
solver_par->init_res = nom0;
nomb = magma_snrm2( 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 ( nom < r0 ) {
info = MAGMA_SUCCESS;
goto cleanup;
}
// check positive definite
if ( MAGMA_S_ABS(den) <= 0.0 ) {
info = MAGMA_NONSPD;
goto cleanup;
}
// array on host for the parameters
CHECK( magma_smalloc_cpu( &skp_h, 7 ));
alpha = rho = gamma = tmp1 = c_one;
beta = magma_sdot( dofs, h.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_S_MAKE(nom, 0.0);
skp_h[6]=MAGMA_S_MAKE(nom, 0.0);
magma_ssetvector( 7, skp_h, 1, skp, 1, queue );
//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++;
// computes SpMV and dot product
CHECK( magma_scgmerge_spmv1( A, d1, d2, d.dval, z.dval, skp, queue ));
solver_par->spmv_count++;
if( precond_par->solver == Magma_JACOBI ){
CHECK( magma_sjcgmerge_xrbeta( dofs, d1, d2, precond_par->d.dval, x->dval, r.dval, d.dval, z.dval, h.dval, skp, queue ));
}
else if( precond_par->solver == Magma_NONE ){
// updates x, r
CHECK( magma_spcgmerge_xrbeta1( dofs, x->dval, r.dval, d.dval, z.dval, skp, queue ));
// computes scalars and updates d
CHECK( magma_spcgmerge_xrbeta2( dofs, d1, d2, r.dval, r.dval, d.dval, skp, queue ));
} else {
// updates x, r
CHECK( magma_spcgmerge_xrbeta1( dofs, x->dval, r.dval, d.dval, z.dval, skp, queue ));
// preconditioner in between
tempop1 = magma_sync_wtime( queue );
CHECK( magma_s_applyprecond_left( MagmaNoTrans, A, r, &rt, precond_par, queue ));
CHECK( magma_s_applyprecond_right( MagmaNoTrans, A, rt, &h, precond_par, queue ));
// magma_scopy( dofs, r.dval, 1, h.dval, 1 );
tempop2 = magma_sync_wtime( queue );
precond_par->runtime += tempop2-tempop1;
// computes scalars and updates d
CHECK( magma_spcgmerge_xrbeta2( dofs, d1, d2, h.dval, r.dval, d.dval, skp, queue ));
}
//if( solver_par->numiter==1){
// magma_scopy( dofs, h.dval, 1, d.dval, 1 );
//}
// updates x, r, computes scalars and updates d
//CHECK( magma_scgmerge_xrbeta( dofs, d1, d2, x->dval, r.dval, d.dval, z.dval, skp, queue ));
// check stopping criterion (asynchronous copy)
magma_sgetvector( 1 , skp+6, 1, skp_h+6, 1, queue );
res = sqrt(MAGMA_S_ABS(skp_h[6]));
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;
float residual;
CHECK( magma_sresidualvec( 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->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) res;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
solver_par->info = MAGMA_DIVERGENCE;
}
cleanup:
magma_smfree(&r, queue );
magma_smfree(&z, queue );
magma_smfree(&d, queue );
magma_smfree(&rt, queue );
magma_smfree(&h, queue );
magma_free( d1 );
magma_free( d2 );
magma_free( skp );
magma_free_cpu( skp_h );
solver_par->info = info;
return info;
} /* magma_spcg_merge */
extern "C" magma_int_t
magma_scgs(
magma_s_matrix A, magma_s_matrix b, magma_s_matrix *x,
magma_s_solver_par *solver_par,
magma_queue_t queue )
{
magma_int_t info = MAGMA_NOTCONVERGED;
// prepare solver feedback
solver_par->solver = Magma_CGS;
solver_par->numiter = 0;
solver_par->spmv_count = 0;
// constants
const float c_zero = MAGMA_S_ZERO;
const float c_one = MAGMA_S_ONE;
const float c_neg_one = MAGMA_S_NEG_ONE;
// solver variables
float nom0, r0, res=0, nomb;
float rho, rho_l = c_one, alpha, beta;
magma_int_t dofs = A.num_rows* b.num_cols;
// GPU workspace
magma_s_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_svinit( &r, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &rt,Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &r_tld,Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &p, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &p_hat, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &q, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &q_hat, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &u, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &u_hat, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &v, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &v_hat, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &t, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
// solver setup
CHECK( magma_sresidualvec( A, b, *x, &r, &nom0, queue));
magma_scopy( dofs, r.dval, 1, r_tld.dval, 1, queue );
solver_par->init_res = nom0;
nomb = magma_snrm2( 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;
tempo1 = magma_sync_wtime( queue );
solver_par->numiter = 0;
solver_par->spmv_count = 0;
// start iteration
do
{
solver_par->numiter++;
rho = magma_sdot( dofs, r.dval, 1, r_tld.dval, 1, queue );
// rho = < r,r_tld>
if( magma_s_isnan_inf( rho ) ){
info = MAGMA_DIVERGENCE;
break;
}
if ( solver_par->numiter > 1 ) { // direction vectors
beta = rho / rho_l;
magma_scopy( dofs, r.dval, 1, u.dval, 1, queue ); // u = r
magma_saxpy( dofs, beta, q.dval, 1, u.dval, 1, queue ); // u = u + beta q
magma_sscal( dofs, beta, p.dval, 1, queue ); // p = beta*p
magma_saxpy( dofs, c_one, q.dval, 1, p.dval, 1, queue ); // p = q + beta*p
magma_sscal( dofs, beta, p.dval, 1, queue ); // p = beta*(q + beta*p)
magma_saxpy( 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_scopy( dofs, r.dval, 1, u.dval, 1, queue ); // u = r
magma_scopy( dofs, r.dval, 1, p.dval, 1, queue ); // p = r
}
CHECK( magma_s_spmv( c_one, A, p, c_zero, v_hat, queue )); // v = A p
solver_par->spmv_count++;
alpha = rho / magma_sdot( dofs, r_tld.dval, 1, v_hat.dval, 1, queue );
magma_scopy( dofs, u.dval, 1, q.dval, 1, queue ); // q = u
magma_saxpy( dofs, -alpha, v_hat.dval, 1, q.dval, 1, queue ); // q = u - alpha v_hat
magma_scopy( dofs, u.dval, 1, t.dval, 1, queue ); // t = q
magma_saxpy( dofs, c_one, q.dval, 1, t.dval, 1, queue ); // t = u + q
CHECK( magma_s_spmv( c_one, A, t, c_zero, rt, queue )); // t = A u_hat
solver_par->spmv_count++;
magma_saxpy( dofs, c_neg_one*alpha, rt.dval, 1, r.dval, 1, queue ); // r = r -alpha*A u_hat
magma_saxpy( dofs, alpha, t.dval, 1, x->dval, 1, queue ); // x = x + alpha u_hat
rho_l = rho;
res = magma_snrm2( 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;
float residual;
CHECK( magma_sresidualvec( 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_smfree(&r, queue );
magma_smfree(&rt, queue );
magma_smfree(&r_tld, queue );
magma_smfree(&p, queue );
magma_smfree(&q, queue );
magma_smfree(&u, queue );
magma_smfree(&v, queue );
magma_smfree(&t, queue );
magma_smfree(&p_hat, queue );
magma_smfree(&q_hat, queue );
magma_smfree(&u_hat, queue );
magma_smfree(&v_hat, queue );
solver_par->info = info;
return info;
} /* magma_scgs */
magma_int_t
magma_spgmres( magma_s_sparse_matrix A, magma_s_vector b, magma_s_vector *x,
magma_s_solver_par *solver_par,
magma_s_preconditioner *precond_par ){
// prepare solver feedback
solver_par->solver = Magma_PGMRES;
solver_par->numiter = 0;
solver_par->info = 0;
// local variables
float c_zero = MAGMA_S_ZERO, c_one = MAGMA_S_ONE,
c_mone = MAGMA_S_NEG_ONE;
magma_int_t dofs = A.num_rows;
magma_int_t i, j, k, m = 0;
magma_int_t restart = min( dofs-1, solver_par->restart );
magma_int_t ldh = restart+1;
float nom, rNorm, RNorm, nom0, betanom, r0 = 0.;
// CPU workspace
//magma_setdevice(0);
float *H, *HH, *y, *h1;
magma_smalloc_pinned( &H, (ldh+1)*ldh );
magma_smalloc_pinned( &y, ldh );
magma_smalloc_pinned( &HH, ldh*ldh );
magma_smalloc_pinned( &h1, ldh );
// GPU workspace
magma_s_vector r, q, q_t, z, z_t, t;
magma_s_vinit( &t, Magma_DEV, dofs, c_zero );
magma_s_vinit( &r, Magma_DEV, dofs, c_zero );
magma_s_vinit( &q, Magma_DEV, dofs*(ldh+1), c_zero );
magma_s_vinit( &z, Magma_DEV, dofs*(ldh+1), c_zero );
magma_s_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;
float *dy, *dH = NULL;
if (MAGMA_SUCCESS != magma_smalloc( &dy, ldh ))
return MAGMA_ERR_DEVICE_ALLOC;
if (MAGMA_SUCCESS != magma_smalloc( &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_sscal( dofs, c_zero, x->val, 1 ); // x = 0
magma_scopy( dofs, b.val, 1, r.val, 1 ); // r = b
nom0 = betanom = magma_snrm2( dofs, r.val, 1 ); // nom0= || r||
nom = nom0 * nom0;
solver_par->init_res = nom0;
H(1,0) = MAGMA_S_MAKE( nom0, 0. );
magma_ssetvector(1, &H(1,0), 1, &dH(1,0), 1);
if ( (r0 = nom0 * RTOLERANCE ) < ATOLERANCE )
r0 = solver_par->epsilon;
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++ ){
for(k=1; k<=restart; k++) {
magma_scopy(dofs, r.val, 1, q(k-1), 1); // q[0] = 1.0/||r||
magma_sscal(dofs, 1./H(k,k-1), q(k-1), 1); // (to be fused)
q_t.val = q(k-1);
magmablasSetKernelStream(stream[0]);
// preconditioner
// z[k] = M^(-1) q(k)
magma_s_applyprecond_left( A, q_t, &t, precond_par );
magma_s_applyprecond_right( A, t, &z_t, precond_par );
magma_scopy(dofs, z_t.val, 1, z(k-1), 1);
// r = A q[k]
magma_s_spmv( c_one, A, z_t, c_zero, r );
// if (solver_par->ortho == Magma_MGS ) {
// modified Gram-Schmidt
for (i=1; i<=k; i++) {
H(i,k) =magma_sdot(dofs, q(i-1), 1, r.val, 1);
// H(i,k) = q[i] . r
magma_saxpy(dofs,-H(i,k), q(i-1), 1, r.val, 1);
// r = r - H(i,k) q[i]
}
H(k+1,k) = MAGMA_S_MAKE( magma_snrm2(dofs, r.val, 1), 0. ); // H(k+1,k) = ||r||
/*}else if (solver_par->ortho == Magma_FUSED_CGS ) {
// fusing sgemv with snrm2 in classical Gram-Schmidt
magmablasSetKernelStream(stream[0]);
magma_scopy(dofs, r.val, 1, q(k), 1);
// dH(1:k+1,k) = q[0:k] . r
magmablas_sgemv(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_sgemv(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_scopyscale( 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_sgetvector_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_sgemv(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 SNRM2SCALE
// 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_sgetvector_async(k, &dH(1,k), 1, &H(1,k),
1, stream[1]);
#endif
// r = r - q[0:k-1] dH(1:k,k)
magmablas_sgemv(MagmaNoTrans, dofs, k, c_mone, q(0),
dofs, &dH(1,k), 1, c_one, r.val, 1);
#ifdef SNRM2SCALE
magma_scopy(dofs, r.val, 1, q(k), 1);
// q[k] = r / H(k,k-1)
magma_snrm2scale(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_sgetvector_async(k+1, &dH(1,k), 1, &H(1,k), 1, stream[1]);
#else
H(k+1,k) = MAGMA_S_MAKE( magma_snrm2(dofs, r.val, 1), 0. );
// H(k+1,k) = sqrt(r . r)
if( k<solver_par->restart ){
magmablasSetKernelStream(stream[0]);
magma_scopy(dofs, r.val, 1, q(k), 1);
// q[k] = 1.0/H[k][k-1] r
magma_sscal(dofs, 1./H(k+1,k), q(k), 1);
// (to be fused)
}
#endif
}*/
/* Minimization of || b-Ax || in H_k */
for (i=1; i<=k; i++) {
HH(k,i) = magma_cblas_sdot( i+1, &H(1,k), 1, &H(1,i), 1 );
}
h1[k] = H(1,k)*H(1,0);
if (k != 1){
for (i=1; i<k; i++) {
HH(k,i) = HH(k,i)/HH(i,i);//
for (m=i+1; m<=k; m++){
HH(k,m) -= HH(k,i) * HH(m,i) * HH(i,i);
}
h1[k] -= h1[i] * HH(k,i);
}
}
y[k] = h1[k]/HH(k,k);
if (k != 1)
for (i=k-1; i>=1; i--) {
y[i] = h1[i]/HH(i,i);
for (j=i+1; j<=k; j++)
y[i] -= y[j] * HH(j,i);
}
m = k;
rNorm = fabs(MAGMA_S_REAL(H(k+1,k)));
}/* Minimization done */
// compute solution approximation
magma_ssetmatrix(m, 1, y+1, m, dy, m );
magma_sgemv(MagmaNoTrans, dofs, m, c_one, z(0), dofs, dy, 1,
c_one, x->val, 1);
// compute residual
magma_s_spmv( c_mone, A, *x, c_zero, r ); // r = - A * x
magma_saxpy(dofs, c_one, b.val, 1, r.val, 1); // r = r + b
H(1,0) = MAGMA_S_MAKE( magma_snrm2(dofs, r.val, 1), 0. );
// RNorm = H[1][0] = || r ||
RNorm = MAGMA_S_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;
float residual;
magma_sresidual( 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_s_vfree(&t);
magma_s_vfree(&r);
magma_s_vfree(&q);
magma_s_vfree(&z);
magma_s_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_spgmres */
extern "C" magma_int_t
magma_sbicgstab_merge(
magma_s_matrix A, magma_s_matrix b, magma_s_matrix *x,
magma_s_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
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
magma_int_t dofs = A.num_rows * b.num_cols;
// workspace
magma_s_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_svinit( &r, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &rr,Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &p, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &v, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &s, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &t, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &d1, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &d2, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
// solver variables
float alpha, beta, omega, rho_old, rho_new;
float betanom, nom0, r0, res, nomb;
res=0;
//float nom;
//float den;
// solver setup
CHECK( magma_sresidualvec( A, b, *x, &r, &nom0, queue));
magma_scopy( dofs, r.dval, 1, rr.dval, 1, queue ); // rr = r
betanom = nom0;
//nom = nom0*nom0;
rho_new = magma_sdot( dofs, r.dval, 1, r.dval, 1, queue ); // rho=<rr,r>
rho_old = omega = alpha = MAGMA_S_MAKE( 1.0, 0. );
solver_par->init_res = nom0;
CHECK( magma_s_spmv( c_one, A, r, c_zero, v, queue )); // z = A r
//den = MAGMA_S_REAL( magma_sdot( dofs, v.dval, 1, r.dval, 1), queue ); // den = z' * r
nomb = magma_snrm2( 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 ( nomb < 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_sdot( 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_s_isnan_inf( beta ) ){
info = MAGMA_DIVERGENCE;
break;
}
// p = r + beta * ( p - omega * v )
magma_sbicgstab_1(
r.num_rows,
r.num_cols,
beta,
omega,
r.dval,
v.dval,
p.dval,
queue );
CHECK( magma_s_spmv( c_one, A, p, c_zero, v, queue )); // v = Ap
solver_par->spmv_count++;
//alpha = rho_new / tmpval;
alpha = rho_new /magma_sdot( dofs, rr.dval, 1, v.dval, 1, queue );
if( magma_s_isnan_inf( alpha ) ){
info = MAGMA_DIVERGENCE;
break;
}
// s = r - alpha v
magma_sbicgstab_2(
r.num_rows,
r.num_cols,
alpha,
r.dval,
v.dval,
s.dval,
queue );
CHECK( magma_s_spmv( c_one, A, s, c_zero, t, queue )); // t=As
solver_par->spmv_count++;
omega = magma_sdot( dofs, t.dval, 1, s.dval, 1, queue ) // omega = <s,t>/<t,t>
/ magma_sdot( dofs, t.dval, 1, t.dval, 1, queue );
// x = x + alpha * p + omega * s
// r = s - omega * t
magma_sbicgstab_3(
r.num_rows,
r.num_cols,
alpha,
omega,
p.dval,
s.dval,
t.dval,
x->dval,
r.dval,
queue );
res = betanom = magma_snrm2( 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;
float residual;
CHECK( magma_sresidualvec( 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_smfree(&r, queue );
magma_smfree(&rr, queue );
magma_smfree(&p, queue );
magma_smfree(&v, queue );
magma_smfree(&s, queue );
magma_smfree(&t, queue );
magma_smfree(&d1, queue );
magma_smfree(&d2, queue );
solver_par->info = info;
return info;
} /* magma_sbicgstab_merge */
extern "C" magma_int_t
magma_scg_res(
magma_s_matrix A, magma_s_matrix b, magma_s_matrix *x,
magma_s_solver_par *solver_par,
magma_queue_t queue )
{
magma_int_t info = MAGMA_NOTCONVERGED;
// prepare solver feedback
solver_par->solver = Magma_CG;
solver_par->numiter = 0;
solver_par->spmv_count = 0;
// solver variables
float alpha, beta;
float nom0, r0, res, nomb;
float den, gammanew, gammaold = MAGMA_S_MAKE(1.0,0.0);
// local variables
float c_zero = MAGMA_S_ZERO, c_one = MAGMA_S_ONE;
magma_int_t dofs = A.num_rows* b.num_cols;
// GPU workspace
magma_s_matrix r={Magma_CSR}, p={Magma_CSR}, q={Magma_CSR};
CHECK( magma_svinit( &r, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &p, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &q, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
// solver setup
CHECK( magma_sresidualvec( A, b, *x, &r, &nom0, queue));
magma_scopy( dofs, r.dval, 1, p.dval, 1, queue ); // p = h
CHECK( magma_s_spmv( c_one, A, p, c_zero, q, queue )); // q = A p
solver_par->spmv_count++;
den = magma_sdot( dofs, p.dval, 1, q.dval, 1, queue ); // den = p dot q
solver_par->init_res = nom0;
nomb = magma_snrm2( 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 ( nomb < r0 ) {
info = MAGMA_SUCCESS;
goto cleanup;
}
// check positive definite
if ( MAGMA_S_ABS(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++;
gammanew = magma_sdot( dofs, r.dval, 1, r.dval, 1, queue );
// gn = < r,r>
if ( solver_par->numiter == 1 ) {
magma_scopy( dofs, r.dval, 1, p.dval, 1, queue ); // p = r
} else {
beta = (gammanew/gammaold); // beta = gn/go
magma_sscal( dofs, beta, p.dval, 1, queue ); // p = beta*p
magma_saxpy( dofs, c_one, r.dval, 1, p.dval, 1, queue ); // p = p + r
}
CHECK( magma_s_spmv( c_one, A, p, c_zero, q, queue )); // q = A p
solver_par->spmv_count++;
den = magma_sdot( dofs, p.dval, 1, q.dval, 1, queue );
// den = p dot q
alpha = gammanew / den;
magma_saxpy( dofs, alpha, p.dval, 1, x->dval, 1, queue ); // x = x + alpha p
magma_saxpy( dofs, -alpha, q.dval, 1, r.dval, 1, queue ); // r = r - alpha q
gammaold = gammanew;
res = magma_snrm2( 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;
float residual;
CHECK( magma_sresidualvec( 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_smfree(&r, queue );
magma_smfree(&p, queue );
magma_smfree(&q, queue );
solver_par->info = info;
return info;
} /* magma_scg */
extern "C" magma_int_t
magma_scg(
magma_s_sparse_matrix A, magma_s_vector b, magma_s_vector *x,
magma_s_solver_par *solver_par,
magma_queue_t queue )
{
// set queue for old dense routines
magma_queue_t orig_queue;
magmablasGetKernelStream( &orig_queue );
// prepare solver feedback
solver_par->solver = Magma_CG;
solver_par->numiter = 0;
solver_par->info = MAGMA_SUCCESS;
// local variables
float c_zero = MAGMA_S_ZERO, c_one = MAGMA_S_ONE;
magma_int_t dofs = A.num_rows;
// GPU workspace
magma_s_vector r, p, q;
magma_s_vinit( &r, Magma_DEV, dofs, c_zero, queue );
magma_s_vinit( &p, Magma_DEV, dofs, c_zero, queue );
magma_s_vinit( &q, Magma_DEV, dofs, c_zero, queue );
// solver variables
float alpha, beta;
float nom, nom0, r0, betanom, betanomsq, den;
// solver setup
magma_sscal( dofs, c_zero, x->dval, 1) ; // x = 0
magma_scopy( dofs, b.dval, 1, r.dval, 1 ); // r = b
magma_scopy( dofs, b.dval, 1, p.dval, 1 ); // p = b
nom0 = betanom = magma_snrm2( dofs, r.dval, 1 );
nom = nom0 * nom0; // nom = r' * r
magma_s_spmv( c_one, A, p, c_zero, q, queue ); // q = A p
den = MAGMA_S_REAL( magma_sdot(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 ) {
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
}
// check positive definite
if (den <= 0.0) {
printf("Operator A is not postive definite. (Ar,r) = %f\n", den);
magmablasSetKernelStream( orig_queue );
return MAGMA_NONSPD;
solver_par->info = MAGMA_NONSPD;
}
//Chronometry
real_Double_t tempo1, tempo2;
tempo1 = magma_sync_wtime( queue );
if ( solver_par->verbose > 0 ) {
solver_par->res_vec[0] = (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++ ) {
alpha = MAGMA_S_MAKE(nom/den, 0.);
magma_saxpy(dofs, alpha, p.dval, 1, x->dval, 1); // x = x + alpha p
magma_saxpy(dofs, -alpha, q.dval, 1, r.dval, 1); // r = r - alpha q
betanom = magma_snrm2(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_S_MAKE(betanomsq/nom, 0.); // beta = betanoms/nom
magma_sscal(dofs, beta, p.dval, 1); // p = beta*p
magma_saxpy(dofs, c_one, r.dval, 1, p.dval, 1); // p = p + r
magma_s_spmv( c_one, A, p, c_zero, q, queue ); // q = A p
den = MAGMA_S_REAL(magma_sdot(dofs, p.dval, 1, q.dval, 1));
// den = p dot q
nom = betanomsq;
}
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t) tempo2-tempo1;
float residual;
magma_sresidual( A, b, *x, &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;
}
}
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_s_vfree(&r, queue );
magma_s_vfree(&p, queue );
magma_s_vfree(&q, queue );
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
} /* magma_scg */
extern "C" magma_int_t
magma_sidr_strms(
magma_s_matrix A, magma_s_matrix b, magma_s_matrix *x,
magma_s_solver_par *solver_par,
magma_queue_t queue )
{
magma_int_t info = MAGMA_NOTCONVERGED;
// prepare solver feedback
solver_par->solver = Magma_IDRMERGE;
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 float c_zero = MAGMA_S_ZERO;
const float c_one = MAGMA_S_ONE;
const float c_n_one = MAGMA_S_NEG_ONE;
// internal user options
const magma_int_t smoothing = 1; // 0 = disable, 1 = enable
const float 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;
magma_int_t ldd;
magma_int_t q;
float residual;
float nrm;
float nrmb;
float nrmr;
float nrmt;
float rho;
float om;
float gamma;
// matrices and vectors
magma_s_matrix dxs = {Magma_CSR};
magma_s_matrix dr = {Magma_CSR}, drs = {Magma_CSR};
magma_s_matrix dP = {Magma_CSR}, dP1 = {Magma_CSR};
magma_s_matrix dG = {Magma_CSR}, dGcol = {Magma_CSR};
magma_s_matrix dU = {Magma_CSR};
magma_s_matrix dM = {Magma_CSR};
magma_s_matrix df = {Magma_CSR};
magma_s_matrix dt = {Magma_CSR}, dtt = {Magma_CSR};
magma_s_matrix dc = {Magma_CSR};
magma_s_matrix dv = {Magma_CSR};
magma_s_matrix dskp = {Magma_CSR};
magma_s_matrix dalpha = {Magma_CSR};
magma_s_matrix dbeta = {Magma_CSR};
float *hMdiag = NULL;
float *hskp = NULL;
float *halpha = NULL;
float *hbeta = NULL;
float *d1 = NULL, *d2 = NULL;
// queue variables
const magma_int_t nqueues = 3; // number of queues
magma_queue_t queues[nqueues];
// chronometry
real_Double_t tempo1, tempo2;
// create additional queues
queues[0] = queue;
for ( q = 1; q < nqueues; q++ ) {
magma_queue_create( queue->device(), &(queues[q]) );
}
// 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_snrm2( b.num_rows, b.dval, 1, queue );
if ( nrmb == 0.0 ) {
magma_sscal( x->num_rows, MAGMA_S_ZERO, x->dval, 1, queue );
info = MAGMA_SUCCESS;
goto cleanup;
}
// t = 0
// make t twice as large to contain both, dt and dr
ldd = magma_roundup( b.num_rows, 32 );
CHECK( magma_svinit( &dt, Magma_DEV, ldd, 2, c_zero, queue ));
dt.num_rows = b.num_rows;
dt.num_cols = 1;
dt.nnz = dt.num_rows;
// redirect the dr.dval to the second part of dt
CHECK( magma_svinit( &dr, Magma_DEV, b.num_rows, 1, c_zero, queue ));
magma_free( dr.dval );
dr.dval = dt.dval + ldd;
// r = b - A x
CHECK( magma_sresidualvec( 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_svinit( &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_slarnv( &distr, iseed, &dof, dP.val );
// transfer P to device
CHECK( magma_smtransfer( dP, &dP1, Magma_CPU, Magma_DEV, queue ));
magma_smfree( &dP, queue );
// P = ortho(P1)
if ( dP1.num_cols > 1 ) {
// P = magma_sqr(P1), QR factorization
CHECK( magma_sqr( dP1.num_rows, dP1.num_cols, dP1, dP1.ld, &dP, NULL, queue ));
} else {
// P = P1 / |P1|
nrm = magma_snrm2( dof, dP1.dval, 1, queue );
nrm = 1.0 / nrm;
magma_sscal( dof, nrm, dP1.dval, 1, queue );
CHECK( magma_smtransfer( dP1, &dP, Magma_DEV, Magma_DEV, queue ));
}
magma_smfree( &dP1, queue );
//---------------------------------------
// allocate memory for the scalar products
CHECK( magma_smalloc_pinned( &hskp, 5 ));
CHECK( magma_svinit( &dskp, Magma_DEV, 4, 1, c_zero, queue ));
CHECK( magma_smalloc_pinned( &halpha, s ));
CHECK( magma_svinit( &dalpha, Magma_DEV, s, 1, c_zero, queue ));
CHECK( magma_smalloc_pinned( &hbeta, s ));
CHECK( magma_svinit( &dbeta, Magma_DEV, s, 1, c_zero, queue ));
// workspace for merged dot product
CHECK( magma_smalloc( &d1, max(2, s) * b.num_rows ));
CHECK( magma_smalloc( &d2, max(2, s) * b.num_rows ));
// smoothing enabled
if ( smoothing > 0 ) {
// set smoothing solution vector
CHECK( magma_smtransfer( *x, &dxs, Magma_DEV, Magma_DEV, queue ));
// tt = 0
// make tt twice as large to contain both, dtt and drs
ldd = magma_roundup( b.num_rows, 32 );
CHECK( magma_svinit( &dtt, Magma_DEV, ldd, 2, c_zero, queue ));
dtt.num_rows = dr.num_rows;
dtt.num_cols = 1;
dtt.nnz = dtt.num_rows;
// redirect the drs.dval to the second part of dtt
CHECK( magma_svinit( &drs, Magma_DEV, dr.num_rows, 1, c_zero, queue ));
magma_free( drs.dval );
drs.dval = dtt.dval + ldd;
// set smoothing residual vector
magma_scopyvector( dr.num_rows, dr.dval, 1, drs.dval, 1, queue );
}
// G(n,s) = 0
if ( s > 1 ) {
ldd = magma_roundup( A.num_rows, 32 );
CHECK( magma_svinit( &dG, Magma_DEV, ldd, s, c_zero, queue ));
dG.num_rows = A.num_rows;
} else {
CHECK( magma_svinit( &dG, Magma_DEV, A.num_rows, s, c_zero, queue ));
}
// dGcol represents a single column of dG, array pointer is set inside loop
CHECK( magma_svinit( &dGcol, Magma_DEV, dG.num_rows, 1, c_zero, queue ));
magma_free( dGcol.dval );
// U(n,s) = 0
if ( s > 1 ) {
ldd = magma_roundup( A.num_cols, 32 );
CHECK( magma_svinit( &dU, Magma_DEV, ldd, s, c_zero, queue ));
dU.num_rows = A.num_cols;
} else {
CHECK( magma_svinit( &dU, Magma_DEV, A.num_cols, s, c_zero, queue ));
}
// M(s,s) = I
CHECK( magma_svinit( &dM, Magma_DEV, s, s, c_zero, queue ));
CHECK( magma_smalloc_pinned( &hMdiag, s ));
magmablas_slaset( MagmaFull, dM.num_rows, dM.num_cols, c_zero, c_one, dM.dval, dM.ld, queue );
// f = 0
CHECK( magma_svinit( &df, Magma_DEV, dP.num_cols, 1, c_zero, queue ));
// c = 0
CHECK( magma_svinit( &dc, Magma_DEV, dM.num_cols, 1, c_zero, queue ));
// v = r
CHECK( magma_smtransfer( dr, &dv, Magma_DEV, Magma_DEV, queue ));
//--------------START TIME---------------
// chronometry
tempo1 = magma_sync_wtime( queue );
if ( solver_par->verbose > 0 ) {
solver_par->timing[0] = 0.0;
}
cudaProfilerStart();
om = MAGMA_S_ONE;
gamma = MAGMA_S_ZERO;
innerflag = 0;
// new RHS for small systems
// f = P' r
// Q1
magma_sgemvmdot_shfl( dP.num_rows, dP.num_cols, dP.dval, dr.dval, d1, d2, df.dval, queues[1] );
// skp[4] = f(k)
// Q1
magma_sgetvector_async( 1, df.dval, 1, &hskp[4], 1, queues[1] );
// c(k:s) = f(k:s)
// Q1
magma_scopyvector_async( s, df.dval, 1, dc.dval, 1, queues[1] );
// c(k:s) = M(k:s,k:s) \ f(k:s)
// Q1
magma_strsv( MagmaLower, MagmaNoTrans, MagmaNonUnit, s, dM.dval, dM.ld, dc.dval, 1, queues[1] );
// start iteration
do
{
solver_par->numiter++;
// shadow space loop
for ( k = 0; k < s; ++k ) {
sk = s - k;
dGcol.dval = dG.dval + k * dG.ld;
// v = r - G(:,k:s) c(k:s)
// Q1
magmablas_sgemv( MagmaNoTrans, dG.num_rows, sk, c_n_one, dGcol.dval, dG.ld, &dc.dval[k], 1, c_one, dv.dval, 1, queues[1] );
// U(:,k) = om * v + U(:,k:s) c(k:s)
// Q1
magmablas_sgemv( MagmaNoTrans, dU.num_rows, sk, c_one, &dU.dval[k*dU.ld], dU.ld, &dc.dval[k], 1, om, dv.dval, 1, queues[1] );
// G(:,k) = A U(:,k)
// Q1
CHECK( magma_s_spmv( c_one, A, dv, c_zero, dGcol, queues[1] ));
solver_par->spmv_count++;
// bi-orthogonalize the new basis vectors
for ( i = 0; i < k; ++i ) {
// alpha = P(:,i)' G(:,k)
// Q1
halpha[i] = magma_sdot( dP.num_rows, &dP.dval[i*dP.ld], 1, dGcol.dval, 1, queues[1] );
// implicit sync Q1 --> alpha = P(:,i)' G(:,k)
// alpha = alpha / M(i,i)
halpha[i] = halpha[i] / hMdiag[i];
// G(:,k) = G(:,k) - alpha * G(:,i)
// Q1
magma_saxpy( dG.num_rows, -halpha[i], &dG.dval[i*dG.ld], 1, dGcol.dval, 1, queues[1] );
}
// sync Q1 --> G(:,k) = G(:,k) - alpha * G(:,i), skp[4] = f(k)
magma_queue_sync( queues[1] );
// new column of M = P'G, first k-1 entries are zero
// M(k:s,k) = P(:,k:s)' G(:,k)
// Q2
magma_sgemvmdot_shfl( dP.num_rows, sk, &dP.dval[k*dP.ld], dGcol.dval, d1, d2, &dM.dval[k*dM.ld+k], queues[2] );
// non-first s iteration
if ( k > 0 ) {
// alpha = dalpha
// Q0
magma_ssetvector_async( k, halpha, 1, dalpha.dval, 1, queues[0] );
// U update outside of loop using GEMV
// U(:,k) = U(:,k) - U(:,1:k) * alpha(1:k)
// Q0
magmablas_sgemv( MagmaNoTrans, dU.num_rows, k, c_n_one, dU.dval, dU.ld, dalpha.dval, 1, c_one, dv.dval, 1, queues[0] );
}
// Mdiag(k) = M(k,k)
// Q2
magma_sgetvector( 1, &dM.dval[k*dM.ld+k], 1, &hMdiag[k], 1, queues[2] );
// implicit sync Q2 --> Mdiag(k) = M(k,k)
// U(:,k) = v
// Q0
magma_scopyvector_async( dU.num_rows, dv.dval, 1, &dU.dval[k*dU.ld], 1, queues[0] );
// check M(k,k) == 0
if ( MAGMA_S_EQUAL(hMdiag[k], MAGMA_S_ZERO) ) {
innerflag = 1;
info = MAGMA_DIVERGENCE;
break;
}
// beta = f(k) / M(k,k)
hbeta[k] = hskp[4] / hMdiag[k];
// check for nan
if ( magma_s_isnan( hbeta[k] ) || magma_s_isinf( hbeta[k] )) {
innerflag = 1;
info = MAGMA_DIVERGENCE;
break;
}
// r = r - beta * G(:,k)
// Q2
magma_saxpy( dr.num_rows, -hbeta[k], dGcol.dval, 1, dr.dval, 1, queues[2] );
// non-last s iteration
if ( (k + 1) < s ) {
// f(k+1:s) = f(k+1:s) - beta * M(k+1:s,k)
// Q1
magma_saxpy( sk-1, -hbeta[k], &dM.dval[k*dM.ld+(k+1)], 1, &df.dval[k+1], 1, queues[1] );
// c(k+1:s) = f(k+1:s)
// Q1
magma_scopyvector_async( sk-1, &df.dval[k+1], 1, &dc.dval[k+1], 1, queues[1] );
// c(k+1:s) = M(k+1:s,k+1:s) \ f(k+1:s)
// Q1
magma_strsv( MagmaLower, MagmaNoTrans, MagmaNonUnit, sk-1, &dM.dval[(k+1)*dM.ld+(k+1)], dM.ld, &dc.dval[k+1], 1, queues[1] );
// skp[4] = f(k+1)
// Q1
magma_sgetvector_async( 1, &df.dval[k+1], 1, &hskp[4], 1, queues[1] );
}
// smoothing disabled
if ( smoothing <= 0 ) {
// |r|
// Q2
nrmr = magma_snrm2( dr.num_rows, dr.dval, 1, queues[2] );
// implicit sync Q2 --> |r|
// smoothing enabled
} else {
// smoothing operation
//---------------------------------------
// t = rs - r
// Q2
magma_sidr_smoothing_1( drs.num_rows, drs.num_cols, drs.dval, dr.dval, dtt.dval, queues[2] );
// x = x + beta * U(:,k)
// Q0
magma_saxpy( x->num_rows, hbeta[k], &dU.dval[k*dU.ld], 1, x->dval, 1, queues[0] );
// t't
// t'rs
// Q2
CHECK( magma_sgemvmdot_shfl( dt.ld, 2, dtt.dval, dtt.dval, d1, d2, &dskp.dval[2], queues[2] ));
// skp[2-3] = dskp[2-3]
// Q2
magma_sgetvector( 2, &dskp.dval[2], 1, &hskp[2], 1, queues[2] );
// implicit sync Q2 --> skp = dskp
// gamma = (t' * rs) / (t' * t)
gamma = hskp[3] / hskp[2];
// rs = rs - gamma * t
// Q1
magma_saxpy( drs.num_rows, -gamma, dtt.dval, 1, drs.dval, 1, queues[1] );
// xs = xs - gamma * (xs - x)
// Q0
magma_sidr_smoothing_2( dxs.num_rows, dxs.num_cols, -gamma, x->dval, dxs.dval, queues[0] );
// |rs|
// Q1
nrmr = magma_snrm2( drs.num_rows, drs.dval, 1, queues[1] );
// implicit sync Q0 --> |r|
//---------------------------------------
}
// v = r
// Q1
magma_scopyvector_async( dr.num_rows, dr.dval, 1, dv.dval, 1, queues[1] );
// last s iteration
if ( (k + 1) == s ) {
// t = A r
// Q2
CHECK( magma_s_spmv( c_one, A, dr, c_zero, dt, queues[2] ));
solver_par->spmv_count++;
// t't
// t'r
// Q2
CHECK( magma_sgemvmdot_shfl( dt.ld, 2, dt.dval, dt.dval, d1, d2, dskp.dval, queues[2] ));
}
// 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 or iteration limit
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;
}
}
// smoothing disabled
if ( smoothing <= 0 && innerflag != 1 ) {
// dbeta(1:s) = beta(1:s)
// Q0
magma_ssetvector_async( s, hbeta, 1, dbeta.dval, 1, queues[0] );
// x = x + U(:,1:s) * beta(1:s)
// Q0
magmablas_sgemv( MagmaNoTrans, dU.num_rows, s, c_one, dU.dval, dU.ld, dbeta.dval, 1, c_one, x->dval, 1, queues[0] );
}
// check convergence or iteration limit or invalid result of inner loop
if ( innerflag > 0 ) {
break;
}
// computation of a new omega
//---------------------------------------
// skp[0-2] = dskp[0-2]
// Q2
magma_sgetvector( 2, dskp.dval, 1, hskp, 1, queues[2] );
// implicit sync Q2 --> skp = dskp
// |t|
nrmt = magma_ssqrt( MAGMA_S_REAL(hskp[0]) );
// rho = abs((t' * r) / (|t| * |r|))
rho = MAGMA_D_ABS( MAGMA_S_REAL(hskp[1]) / (nrmt * nrmr) );
// om = (t' * r) / (|t| * |t|)
om = hskp[1] / hskp[0];
if ( rho < angle ) {
om = (om * angle) / rho;
}
//---------------------------------------
if ( MAGMA_S_EQUAL(om, MAGMA_S_ZERO) ) {
info = MAGMA_DIVERGENCE;
break;
}
// sync Q1 --> v = r
magma_queue_sync( queues[1] );
// r = r - om * t
// Q2
magma_saxpy( dr.num_rows, -om, dt.dval, 1, dr.dval, 1, queues[2] );
// x = x + om * v
// Q0
magma_saxpy( x->num_rows, om, dv.dval, 1, x->dval, 1, queues[0] );
// smoothing disabled
if ( smoothing <= 0 ) {
// |r|
// Q2
nrmr = magma_snrm2( dr.num_rows, dr.dval, 1, queues[2] );
// implicit sync Q2 --> |r|
// v = r
// Q0
magma_scopyvector_async( dr.num_rows, dr.dval, 1, dv.dval, 1, queues[0] );
// new RHS for small systems
// f = P' r
// Q1
magma_sgemvmdot_shfl( dP.num_rows, dP.num_cols, dP.dval, dr.dval, d1, d2, df.dval, queues[1] );
// skp[4] = f(k)
// Q1
magma_sgetvector_async( 1, df.dval, 1, &hskp[4], 1, queues[1] );
// c(k:s) = f(k:s)
// Q1
magma_scopyvector_async( s, df.dval, 1, dc.dval, 1, queues[1] );
// c(k:s) = M(k:s,k:s) \ f(k:s)
// Q1
magma_strsv( MagmaLower, MagmaNoTrans, MagmaNonUnit, s, dM.dval, dM.ld, dc.dval, 1, queues[1] );
// smoothing enabled
} else {
// smoothing operation
//---------------------------------------
// t = rs - r
// Q2
magma_sidr_smoothing_1( drs.num_rows, drs.num_cols, drs.dval, dr.dval, dtt.dval, queues[2] );
// t't
// t'rs
// Q2
CHECK( magma_sgemvmdot_shfl( dt.ld, 2, dtt.dval, dtt.dval, d1, d2, &dskp.dval[2], queues[2] ));
// skp[2-3] = dskp[2-3]
// Q2
magma_sgetvector( 2, &dskp.dval[2], 1, &hskp[2], 1, queues[2] );
// implicit sync Q2 --> skp = dskp
// gamma = (t' * rs) / (t' * t)
gamma = hskp[3] / hskp[2];
// rs = rs - gamma * (rs - r)
// Q2
magma_saxpy( drs.num_rows, -gamma, dtt.dval, 1, drs.dval, 1, queues[2] );
// xs = xs - gamma * (xs - x)
// Q0
magma_sidr_smoothing_2( dxs.num_rows, dxs.num_cols, -gamma, x->dval, dxs.dval, queues[0] );
// v = r
// Q0
magma_scopyvector_async( dr.num_rows, dr.dval, 1, dv.dval, 1, queues[0] );
// new RHS for small systems
// f = P' r
// Q1
magma_sgemvmdot_shfl( dP.num_rows, dP.num_cols, dP.dval, dr.dval, d1, d2, df.dval, queues[1] );
// skp[4] = f(k)
// Q1
magma_sgetvector_async( 1, df.dval, 1, &hskp[4], 1, queues[1] );
// c(k:s) = f(k:s)
// Q1
magma_scopyvector_async( s, df.dval, 1, dc.dval, 1, queues[1] );
// |rs|
// Q2
nrmr = magma_snrm2( drs.num_rows, drs.dval, 1, queues[2] );
// implicit sync Q2 --> |r|
// c(k:s) = M(k:s,k:s) \ f(k:s)
// Q1
magma_strsv( MagmaLower, MagmaNoTrans, MagmaNonUnit, s, dM.dval, dM.ld, dc.dval, 1, queues[1] );
//---------------------------------------
}
// store current timing and residual
if ( solver_par->verbose > 0 ) {
tempo2 = magma_sync_wtime( queue );
magma_queue_sync( 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 or iteration limit
if ( nrmr <= solver_par->atol ||
nrmr/nrmb <= solver_par->rtol ) {
info = MAGMA_SUCCESS;
break;
}
// sync Q0 --> v = r
magma_queue_sync( queues[0] );
}
while ( solver_par->numiter + 1 <= solver_par->maxiter );
// sync all queues
for ( q = 0; q < nqueues; q++ ) {
magma_queue_sync( queues[q] );
}
// smoothing enabled
if ( smoothing > 0 ) {
// x = xs
magma_scopyvector_async( x->num_rows, dxs.dval, 1, x->dval, 1, queue );
// r = rs
magma_scopyvector_async( dr.num_rows, drs.dval, 1, dr.dval, 1, queue );
}
cudaProfilerStop();
// get last iteration timing
tempo2 = magma_sync_wtime( queue );
magma_queue_sync( queue );
solver_par->runtime = (real_Double_t)tempo2 - tempo1;
//--------------STOP TIME----------------
// get final stats
solver_par->iter_res = nrmr;
CHECK( magma_sresidualvec( 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
// sync all queues, destory additional queues
magma_queue_sync( queues[0] );
for ( q = 1; q < nqueues; q++ ) {
magma_queue_sync( queues[q] );
magma_queue_destroy( queues[q] );
}
// smoothing enabled
if ( smoothing > 0 ) {
drs.dval = NULL; // needed because its pointer is redirected to dtt
magma_smfree( &dxs, queue );
magma_smfree( &drs, queue );
magma_smfree( &dtt, queue );
}
dr.dval = NULL; // needed because its pointer is redirected to dt
dGcol.dval = NULL; // needed because its pointer is redirected to dG
magma_smfree( &dr, queue );
magma_smfree( &dP, queue );
magma_smfree( &dP1, queue );
magma_smfree( &dG, queue );
magma_smfree( &dGcol, queue );
magma_smfree( &dU, queue );
magma_smfree( &dM, queue );
magma_smfree( &df, queue );
magma_smfree( &dt, queue );
magma_smfree( &dc, queue );
magma_smfree( &dv, queue );
magma_smfree( &dskp, queue );
magma_smfree( &dalpha, queue );
magma_smfree( &dbeta, queue );
magma_free_pinned( hMdiag );
magma_free_pinned( hskp );
magma_free_pinned( halpha );
magma_free_pinned( hbeta );
magma_free( d1 );
magma_free( d2 );
solver_par->info = info;
return info;
/* magma_sidr_strms */
}
extern "C" magma_int_t
magma_scg_merge(
magma_s_matrix A, magma_s_matrix b, magma_s_matrix *x,
magma_s_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
float alpha, beta, gamma, rho, tmp1, *skp_h={0};
float nom, nom0, betanom, den, nomb;
// some useful variables
float c_zero = MAGMA_S_ZERO, c_one = MAGMA_S_ONE;
magma_int_t dofs = A.num_rows*b.num_cols;
magma_s_matrix r={Magma_CSR}, d={Magma_CSR}, z={Magma_CSR}, B={Magma_CSR}, C={Magma_CSR};
float *d1=NULL, *d2=NULL, *skp=NULL;
// GPU workspace
CHECK( magma_svinit( &r, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &d, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_svinit( &z, Magma_DEV, A.num_rows, b.num_cols, c_zero, queue ));
CHECK( magma_smalloc( &d1, dofs*(1) ));
CHECK( magma_smalloc( &d2, dofs*(1) ));
// array for the parameters
CHECK( magma_smalloc( &skp, 6 ));
// skp = [alpha|beta|gamma|rho|tmp1|tmp2]
// solver setup
magma_sscal( dofs, c_zero, x->dval, 1, queue ); // x = 0
//CHECK( magma_sresidualvec( A, b, *x, &r, nom0, queue));
magma_scopy( dofs, b.dval, 1, r.dval, 1, queue ); // r = b
magma_scopy( dofs, r.dval, 1, d.dval, 1, queue ); // d = r
nom0 = betanom = magma_snrm2( dofs, r.dval, 1, queue );
nom = nom0 * nom0; // nom = r' * r
CHECK( magma_s_spmv( c_one, A, d, c_zero, z, queue )); // z = A d
den = MAGMA_S_ABS( magma_sdot( dofs, d.dval, 1, z.dval, 1, queue ) ); // den = d'* z
solver_par->init_res = nom0;
nomb = magma_snrm2( dofs, b.dval, 1, queue );
if ( nomb == 0.0 ){
nomb=1.0;
}
// array on host for the parameters
CHECK( magma_smalloc_cpu( &skp_h, 6 ));
alpha = rho = gamma = tmp1 = c_one;
beta = magma_sdot( 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_S_MAKE(nom, 0.0);
magma_ssetvector( 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_scgmerge_spmv1( A, d1, d2, d.dval, z.dval, skp, queue ));
solver_par->spmv_count++;
// updates x, r, computes scalars and updates d
CHECK( magma_scgmerge_xrbeta( dofs, d1, d2, x->dval, r.dval, d.dval, z.dval, skp, queue ));
// check stopping criterion (asynchronous copy)
magma_sgetvector( 1 , skp+1, 1, skp_h+1, 1, queue );
betanom = sqrt(MAGMA_S_ABS(skp_h[1]));
if ( solver_par->verbose > 0 ) {
tempo2 = magma_sync_wtime( queue );
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
if ( betanom < solver_par->atol ||
betanom/nomb < solver_par->rtol ) {
break;
}
}
while ( solver_par->numiter+1 <= solver_par->maxiter );
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t) tempo2-tempo1;
float residual;
CHECK( magma_sresidualvec( 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_smfree(&r, queue );
magma_smfree(&z, queue );
magma_smfree(&d, queue );
magma_smfree(&B, queue );
magma_smfree(&C, queue );
magma_free( d1 );
magma_free( d2 );
magma_free( skp );
magma_free_cpu( skp_h );
solver_par->info = info;
return info;
} /* magma_scg_merge */
/**
Purpose
-------
SGEGQR orthogonalizes the N vectors given by a real M-by-N matrix A:
A = Q * R.
On exit, if successful, the orthogonal vectors Q overwrite A
and R is given in work (on the CPU memory).
The routine is designed for tall-and-skinny matrices: M >> N, N <= 128.
This version uses normal equations and SVD in an iterative process that
makes the computation numerically accurate.
Arguments
---------
@param[in]
ikind INTEGER
Several versions are implemented indiceted by the ikind value:
1: This version uses normal equations and SVD in an iterative process
that makes the computation numerically accurate.
2: This version uses a standard LAPACK-based orthogonalization through
MAGMA's QR panel factorization (magma_sgeqr2x3_gpu) and magma_sorgqr
3: MGS
4. Cholesky QR [ Note: this method uses the normal equations which
squares the condition number of A, therefore
||I - Q'Q|| < O(eps cond(A)^2) ]
@param[in]
m INTEGER
The number of rows of the matrix A. m >= n >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. 128 >= n >= 0.
@param[in,out]
dA REAL array on the GPU, dimension (ldda,n)
On entry, the m-by-n matrix A.
On exit, the m-by-n matrix Q with orthogonal columns.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,m).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
@param
dwork (GPU workspace) REAL array, dimension:
n^2 for ikind = 1
3 n^2 + min(m, n) + 2 for ikind = 2
0 (not used) for ikind = 3
n^2 for ikind = 4
@param[out]
work (CPU workspace) REAL array, dimension 3 n^2.
On exit, work(1:n^2) holds the rectangular matrix R.
Preferably, for higher performance, work should be in pinned memory.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
@ingroup magma_sgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_sgegqr_gpu( magma_int_t ikind, magma_int_t m, magma_int_t n,
float *dA, magma_int_t ldda,
float *dwork, float *work,
magma_int_t *info )
{
#define work(i_,j_) (work + (i_) + (j_)*n)
#define dA(i_,j_) (dA + (i_) + (j_)*ldda)
magma_int_t i = 0, j, k, n2 = n*n;
magma_int_t ione = 1;
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float cn = 200., mins, maxs;
/* check arguments */
*info = 0;
if (ikind < 1 || ikind > 4) {
*info = -1;
} else if (m < 0 || m < n) {
*info = -2;
} else if (n < 0 || n > 128) {
*info = -3;
} else if (ldda < max(1,m)) {
*info = -5;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (ikind == 1) {
// === Iterative, based on SVD ============================================================
float *U, *VT, *vt, *R, *G, *hwork, *tau;
float *S;
R = work; // Size n * n
G = R + n*n; // Size n * n
VT = G + n*n; // Size n * n
magma_smalloc_cpu( &hwork, 32 + 2*n*n + 2*n);
if ( hwork == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_int_t lwork=n*n+32; // First part f hwork; used as workspace in svd
U = hwork + n*n + 32; // Size n*n
S = (float *)(U+n*n); // Size n
tau = U + n*n + n; // Size n
#if defined(PRECISION_c) || defined(PRECISION_z)
float *rwork;
magma_smalloc_cpu( &rwork, 5*n);
if ( rwork == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
#endif
do {
i++;
magma_sgemm(MagmaConjTrans, MagmaNoTrans, n, n, m, c_one, dA, ldda, dA, ldda, c_zero, dwork, n );
magma_sgetmatrix(n, n, dwork, n, G, n);
#if defined(PRECISION_s) || defined(PRECISION_d)
lapackf77_sgesvd("n", "a", &n, &n, G, &n, S, U, &n, VT, &n,
hwork, &lwork, info);
#else
lapackf77_sgesvd("n", "a", &n, &n, G, &n, S, U, &n, VT, &n,
hwork, &lwork, rwork, info);
#endif
mins = 100.f, maxs = 0.f;
for (k=0; k < n; k++) {
S[k] = magma_ssqrt( S[k] );
if (S[k] < mins) mins = S[k];
if (S[k] > maxs) maxs = S[k];
}
for (k=0; k < n; k++) {
vt = VT + k*n;
for (j=0; j < n; j++)
vt[j] *= S[j];
}
lapackf77_sgeqrf(&n, &n, VT, &n, tau, hwork, &lwork, info);
if (i == 1)
blasf77_scopy(&n2, VT, &ione, R, &ione);
else
blasf77_strmm("l", "u", "n", "n", &n, &n, &c_one, VT, &n, R, &n);
magma_ssetmatrix(n, n, VT, n, dwork, n);
magma_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit, m, n, c_one, dwork, n, dA, ldda);
if (mins > 0.00001f)
cn = maxs/mins;
//fprintf(stderr, "Iteration %d, cond num = %f \n", i, cn);
} while (cn > 10.f);
magma_free_cpu( hwork );
#if defined(PRECISION_c) || defined(PRECISION_z)
magma_free_cpu( rwork );
#endif
// ================== end of ikind == 1 ===================================================
}
else if (ikind == 2) {
// ================== LAPACK based ===================================================
magma_int_t min_mn = min(m, n);
magma_int_t nb = n;
float *dtau = dwork + 2*n*n, *d_T = dwork, *ddA = dwork + n*n;
float *tau = work+n*n;
magmablas_slaset( MagmaFull, n, n, c_zero, c_zero, d_T, n );
magma_sgeqr2x3_gpu(m, n, dA, ldda, dtau, d_T, ddA,
(float *)(dwork+min_mn+2*n*n), info);
magma_sgetmatrix( min_mn, 1, dtau, min_mn, tau, min_mn);
magma_sgetmatrix( n, n, ddA, n, work, n);
magma_sorgqr_gpu( m, n, n, dA, ldda, tau, d_T, nb, info );
// ================== end of ikind == 2 ===================================================
}
else if (ikind == 3) {
// ================== MGS ===================================================
for(magma_int_t j = 0; j<n; j++){
for(magma_int_t i = 0; i<j; i++){
*work(i, j) = magma_sdot(m, dA(0,i), 1, dA(0,j), 1);
magma_saxpy(m, -(*work(i,j)), dA(0,i), 1, dA(0,j), 1);
}
for(magma_int_t i = j; i<n; i++)
*work(i, j) = MAGMA_S_ZERO;
//*work(j,j) = MAGMA_S_MAKE( magma_snrm2(m, dA(0,j), 1), 0. );
*work(j,j) = magma_sdot(m, dA(0,j), 1, dA(0,j), 1);
*work(j,j) = MAGMA_S_MAKE( sqrt(MAGMA_S_REAL( *work(j,j) )), 0.);
magma_sscal(m, 1./ *work(j,j), dA(0,j), 1);
}
// ================== end of ikind == 3 ===================================================
}
else if (ikind == 4) {
// ================== Cholesky QR ===================================================
magma_sgemm(MagmaConjTrans, MagmaNoTrans, n, n, m, c_one, dA, ldda, dA, ldda, c_zero, dwork, n );
magma_sgetmatrix(n, n, dwork, n, work, n);
lapackf77_spotrf("u", &n, work, &n, info);
magma_ssetmatrix(n, n, work, n, dwork, n);
magma_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit, m, n, c_one, dwork, n, dA, ldda);
// ================== end of ikind == 4 ===================================================
}
return *info;
} /* magma_sgegqr_gpu */
extern "C" magma_int_t
magma_sbpcg(
magma_s_matrix A, magma_s_matrix b, magma_s_matrix *x,
magma_s_solver_par *solver_par,
magma_s_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
float c_zero = MAGMA_S_ZERO, c_one = MAGMA_S_ONE;
magma_int_t dofs = A.num_rows;
// GPU workspace
magma_s_matrix r={Magma_CSR}, rt={Magma_CSR}, p={Magma_CSR}, q={Magma_CSR}, h={Magma_CSR};
// solver variables
float *alpha={0}, *beta={0};
alpha = NULL;
beta = NULL;
float *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_smalloc_cpu(&alpha, num_vecs));
CHECK( magma_smalloc_cpu(&beta, num_vecs));
CHECK( magma_smalloc_cpu(&residual, num_vecs));
CHECK( magma_smalloc_cpu(&nom, num_vecs));
CHECK( magma_smalloc_cpu(&nom0, num_vecs));
CHECK( magma_smalloc_cpu(&r0, num_vecs));
CHECK( magma_smalloc_cpu(&gammaold, num_vecs));
CHECK( magma_smalloc_cpu(&gammanew, num_vecs));
CHECK( magma_smalloc_cpu(&den, num_vecs));
CHECK( magma_smalloc_cpu(&res, num_vecs));
CHECK( magma_smalloc_cpu(&residual, num_vecs));
CHECK( magma_svinit( &r, Magma_DEV, dofs*num_vecs, 1, c_zero, queue ));
CHECK( magma_svinit( &rt, Magma_DEV, dofs*num_vecs, 1, c_zero, queue ));
CHECK( magma_svinit( &p, Magma_DEV, dofs*num_vecs, 1, c_zero, queue ));
CHECK( magma_svinit( &q, Magma_DEV, dofs*num_vecs, 1, c_zero, queue ));
CHECK( magma_svinit( &h, Magma_DEV, dofs*num_vecs, 1, c_zero, queue ));
// solver setup
CHECK( magma_sresidualvec( A, b, *x, &r, nom0, queue));
// preconditioner
CHECK( magma_s_applyprecond_left( MagmaNoTrans, A, r, &rt, precond_par, queue ));
CHECK( magma_s_applyprecond_right( MagmaNoTrans, A, rt, &h, precond_par, queue ));
magma_scopy( dofs*num_vecs, h.dval, 1, p.dval, 1, queue ); // p = h
for( i=0; i<num_vecs; i++) {
nom[i] = MAGMA_S_REAL( magma_sdot( dofs, r(i), 1, h(i), 1, queue ) );
nom0[i] = magma_snrm2( dofs, r(i), 1, queue );
}
CHECK( magma_s_spmv( c_one, A, p, c_zero, q, queue )); // q = A p
for( i=0; i<num_vecs; i++)
den[i] = MAGMA_S_REAL( magma_sdot( 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_s_applyprecond_left( MagmaNoTrans, A, r, &rt, precond_par, queue ));
CHECK( magma_s_applyprecond_right( MagmaNoTrans, A, rt, &h, precond_par, queue ));
for( i=0; i<num_vecs; i++)
gammanew[i] = MAGMA_S_REAL( magma_sdot( dofs, r(i), 1, h(i), 1, queue ) ); // gn = < r,h>
if ( solver_par->numiter==1 ) {
magma_scopy( dofs*num_vecs, h.dval, 1, p.dval, 1, queue ); // p = h
} else {
for( i=0; i<num_vecs; i++) {
beta[i] = MAGMA_S_MAKE(gammanew[i]/gammaold[i], 0.); // beta = gn/go
magma_sscal( dofs, beta[i], p(i), 1, queue ); // p = beta*p
magma_saxpy( dofs, c_one, h(i), 1, p(i), 1, queue ); // p = p + h
}
}
CHECK( magma_s_spmv( c_one, A, p, c_zero, q, queue )); // q = A p
solver_par->spmv_count++;
// magma_s_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_S_REAL(magma_sdot( dofs, p(i), 1, q(i), 1, queue) );
// den = p dot q
alpha[i] = MAGMA_S_MAKE(gammanew[i]/den[i], 0.);
magma_saxpy( dofs, alpha[i], p(i), 1, x->dval+dofs*i, 1, queue ); // x = x + alpha p
magma_saxpy( dofs, -alpha[i], q(i), 1, r(i), 1, queue ); // r = r - alpha q
gammaold[i] = gammanew[i];
res[i] = magma_snrm2( 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_sresidual( 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_smfree(&r, queue );
magma_smfree(&rt, queue );
magma_smfree(&p, queue );
magma_smfree(&q, queue );
magma_smfree(&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_sbpcg */
extern "C" magma_int_t
magma_spidr(
magma_s_matrix A, magma_s_matrix b, magma_s_matrix *x,
magma_s_solver_par *solver_par,
magma_s_preconditioner *precond_par,
magma_queue_t queue )
{
magma_int_t info = MAGMA_NOTCONVERGED;
// prepare solver feedback
solver_par->solver = Magma_PIDR;
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 float c_zero = MAGMA_S_ZERO;
const float c_one = MAGMA_S_ONE;
const float c_n_one = MAGMA_S_NEG_ONE;
// internal user parameters
const magma_int_t smoothing = 1; // 0 = disable, 1 = enable
const float 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;
float residual;
float nrm;
float nrmb;
float nrmr;
float nrmt;
float rho;
float om;
float tt;
float tr;
float gamma;
float alpha;
float mkk;
float fk;
// matrices and vectors
magma_s_matrix dxs = {Magma_CSR};
magma_s_matrix dr = {Magma_CSR}, drs = {Magma_CSR};
magma_s_matrix dP = {Magma_CSR}, dP1 = {Magma_CSR};
magma_s_matrix dG = {Magma_CSR};
magma_s_matrix dU = {Magma_CSR};
magma_s_matrix dM = {Magma_CSR};
magma_s_matrix df = {Magma_CSR};
magma_s_matrix dt = {Magma_CSR};
magma_s_matrix dc = {Magma_CSR};
magma_s_matrix dv = {Magma_CSR};
magma_s_matrix dbeta = {Magma_CSR}, hbeta = {Magma_CSR};
magma_s_matrix dlu = {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_snrm2( b.num_rows, b.dval, 1, queue );
if ( nrmb == 0.0 ) {
magma_sscal( x->num_rows, MAGMA_S_ZERO, x->dval, 1, queue );
info = MAGMA_SUCCESS;
goto cleanup;
}
// r = b - A x
CHECK( magma_svinit( &dr, Magma_DEV, b.num_rows, 1, c_zero, queue ));
CHECK( magma_sresidualvec( 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_svinit( &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_slarnv( &distr, iseed, &dof, dP.val );
// transfer P to device
CHECK( magma_smtransfer( dP, &dP1, Magma_CPU, Magma_DEV, queue ));
magma_smfree( &dP, queue );
// P = ortho(P1)
if ( dP1.num_cols > 1 ) {
// P = magma_sqr(P1), QR factorization
CHECK( magma_sqr( dP1.num_rows, dP1.num_cols, dP1, dP1.ld, &dP, NULL, queue ));
} else {
// P = P1 / |P1|
nrm = magma_snrm2( dof, dP1.dval, 1, queue );
nrm = 1.0 / nrm;
magma_sscal( dof, nrm, dP1.dval, 1, queue );
CHECK( magma_smtransfer( dP1, &dP, Magma_DEV, Magma_DEV, queue ));
}
magma_smfree( &dP1, queue );
//---------------------------------------
// allocate memory for the scalar products
CHECK( magma_svinit( &hbeta, Magma_CPU, s, 1, c_zero, queue ));
CHECK( magma_svinit( &dbeta, Magma_DEV, s, 1, c_zero, queue ));
// smoothing enabled
if ( smoothing > 0 ) {
// set smoothing solution vector
CHECK( magma_smtransfer( *x, &dxs, Magma_DEV, Magma_DEV, queue ));
// set smoothing residual vector
CHECK( magma_smtransfer( dr, &drs, Magma_DEV, Magma_DEV, queue ));
}
// G(n,s) = 0
CHECK( magma_svinit( &dG, Magma_DEV, A.num_cols, s, c_zero, queue ));
// U(n,s) = 0
CHECK( magma_svinit( &dU, Magma_DEV, A.num_cols, s, c_zero, queue ));
// M(s,s) = I
CHECK( magma_svinit( &dM, Magma_DEV, s, s, c_zero, queue ));
magmablas_slaset( MagmaFull, s, s, c_zero, c_one, dM.dval, s, queue );
// f = 0
CHECK( magma_svinit( &df, Magma_DEV, dP.num_cols, 1, c_zero, queue ));
// t = 0
CHECK( magma_svinit( &dt, Magma_DEV, dr.num_rows, 1, c_zero, queue ));
// c = 0
CHECK( magma_svinit( &dc, Magma_DEV, dM.num_cols, 1, c_zero, queue ));
// v = 0
CHECK( magma_svinit( &dv, Magma_DEV, dr.num_rows, 1, c_zero, queue ));
// lu = 0
CHECK( magma_svinit( &dlu, Magma_DEV, A.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_S_ONE;
innerflag = 0;
// start iteration
do
{
solver_par->numiter++;
// new RHS for small systems
// f = P' r
magmablas_sgemv( 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_scopyvector( sk, &df.dval[k], 1, &dc.dval[k], 1, queue );
magma_strsv( 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_scopyvector( dr.num_rows, dr.dval, 1, dv.dval, 1, queue );
magmablas_sgemv( 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 );
// preconditioning operation
// v = L \ v;
// v = U \ v;
CHECK( magma_s_applyprecond_left( MagmaNoTrans, A, dv, &dlu, precond_par, queue ));
CHECK( magma_s_applyprecond_right( MagmaNoTrans, A, dlu, &dv, precond_par, queue ));
// U(:,k) = om * v + U(:,k:s) c(k:s)
magmablas_sgemv( MagmaNoTrans, dU.num_rows, sk, c_one, &dU.dval[k*dU.ld], dU.ld, &dc.dval[k], 1, om, dv.dval, 1, queue );
magma_scopyvector( dU.num_rows, dv.dval, 1, &dU.dval[k*dU.ld], 1, queue );
// G(:,k) = A U(:,k)
CHECK( magma_s_spmv( c_one, A, dv, c_zero, dv, queue ));
solver_par->spmv_count++;
magma_scopyvector( 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_sdot( dP.num_rows, &dP.dval[i*dP.ld], 1, &dG.dval[k*dG.ld], 1, queue );
// alpha = alpha / M(i,i)
magma_sgetvector( 1, &dM.dval[i*dM.ld+i], 1, &mkk, 1, queue );
alpha = alpha / mkk;
// G(:,k) = G(:,k) - alpha * G(:,i)
magma_saxpy( 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_saxpy( 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_sgemv( 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_sgetvector( 1, &dM.dval[k*dM.ld+k], 1, &mkk, 1, queue );
if ( MAGMA_S_EQUAL(mkk, MAGMA_S_ZERO) ) {
innerflag = 1;
info = MAGMA_DIVERGENCE;
break;
}
// beta = f(k) / M(k,k)
magma_sgetvector( 1, &df.dval[k], 1, &fk, 1, queue );
hbeta.val[k] = fk / mkk;
// check for nan
if ( magma_s_isnan( hbeta.val[k] ) || magma_s_isinf( hbeta.val[k] )) {
innerflag = 1;
info = MAGMA_DIVERGENCE;
break;
}
// r = r - beta * G(:,k)
magma_saxpy( dr.num_rows, -hbeta.val[k], &dG.dval[k*dG.ld], 1, dr.dval, 1, queue );
// smoothing disabled
if ( smoothing <= 0 ) {
// |r|
nrmr = magma_snrm2( dr.num_rows, dr.dval, 1, queue );
// smoothing enabled
} else {
// x = x + beta * U(:,k)
magma_saxpy( x->num_rows, hbeta.val[k], &dU.dval[k*dU.ld], 1, x->dval, 1, queue );
// smoothing operation
//---------------------------------------
// t = rs - r
magma_scopyvector( drs.num_rows, drs.dval, 1, dt.dval, 1, queue );
magma_saxpy( dt.num_rows, c_n_one, dr.dval, 1, dt.dval, 1, queue );
// t't
// t'rs
tt = magma_sdot( dt.num_rows, dt.dval, 1, dt.dval, 1, queue );
tr = magma_sdot( dt.num_rows, dt.dval, 1, drs.dval, 1, queue );
// gamma = (t' * rs) / (t' * t)
gamma = tr / tt;
// rs = rs - gamma * (rs - r)
magma_saxpy( drs.num_rows, -gamma, dt.dval, 1, drs.dval, 1, queue );
// xs = xs - gamma * (xs - x)
magma_scopyvector( dxs.num_rows, dxs.dval, 1, dt.dval, 1, queue );
magma_saxpy( dt.num_rows, c_n_one, x->dval, 1, dt.dval, 1, queue );
magma_saxpy( dxs.num_rows, -gamma, dt.dval, 1, dxs.dval, 1, queue );
// |rs|
nrmr = magma_snrm2( 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_saxpy( 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_ssetvector( s, hbeta.val, 1, dbeta.dval, 1, queue );
magmablas_sgemv( 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;
}
// v = r
magma_scopyvector( dr.num_rows, dr.dval, 1, dv.dval, 1, queue );
// preconditioning operation
// v = L \ v;
// v = U \ v;
CHECK( magma_s_applyprecond_left( MagmaNoTrans, A, dv, &dlu, precond_par, queue ));
CHECK( magma_s_applyprecond_right( MagmaNoTrans, A, dlu, &dv, precond_par, queue ));
// t = A v
CHECK( magma_s_spmv( c_one, A, dv, c_zero, dt, queue ));
solver_par->spmv_count++;
// computation of a new omega
//---------------------------------------
// |t|
nrmt = magma_snrm2( dt.num_rows, dt.dval, 1, queue );
// t'r
tr = magma_sdot( dt.num_rows, dt.dval, 1, dr.dval, 1, queue );
// rho = abs(t' * r) / (|t| * |r|))
rho = MAGMA_D_ABS( MAGMA_S_REAL(tr) / (nrmt * nrmr) );
// om = (t' * r) / (|t| * |t|)
om = tr / (nrmt * nrmt);
if ( rho < angle ) {
om = (om * angle) / rho;
}
//---------------------------------------
if ( MAGMA_S_EQUAL(om, MAGMA_S_ZERO) ) {
info = MAGMA_DIVERGENCE;
break;
}
// update approximation vector
// x = x + om * v
magma_saxpy( x->num_rows, om, dv.dval, 1, x->dval, 1, queue );
// update residual vector
// r = r - om * t
magma_saxpy( dr.num_rows, -om, dt.dval, 1, dr.dval, 1, queue );
// smoothing disabled
if ( smoothing <= 0 ) {
// residual norm
nrmr = magma_snrm2( b.num_rows, dr.dval, 1, queue );
// smoothing enabled
} else {
// smoothing operation
//---------------------------------------
// t = rs - r
magma_scopyvector( drs.num_rows, drs.dval, 1, dt.dval, 1, queue );
magma_saxpy( dt.num_rows, c_n_one, dr.dval, 1, dt.dval, 1, queue );
// t't
// t'rs
tt = magma_sdot( dt.num_rows, dt.dval, 1, dt.dval, 1, queue );
tr = magma_sdot( dt.num_rows, dt.dval, 1, drs.dval, 1, queue );
// gamma = (t' * rs) / (|t| * |t|)
gamma = tr / tt;
// rs = rs - gamma * (rs - r)
magma_saxpy( drs.num_rows, -gamma, dt.dval, 1, drs.dval, 1, queue );
// xs = xs - gamma * (xs - x)
magma_scopyvector( dxs.num_rows, dxs.dval, 1, dt.dval, 1, queue );
magma_saxpy( dt.num_rows, c_n_one, x->dval, 1, dt.dval, 1, queue );
magma_saxpy( dxs.num_rows, -gamma, dt.dval, 1, dxs.dval, 1, queue );
// |rs|
nrmr = magma_snrm2( 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_scopyvector( x->num_rows, dxs.dval, 1, x->dval, 1, queue );
// r = rs
magma_scopyvector( 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_sresidualvec( 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_smfree( &dxs, queue );
magma_smfree( &drs, queue );
}
magma_smfree( &dr, queue );
magma_smfree( &dP, queue );
magma_smfree( &dP1, queue );
magma_smfree( &dG, queue );
magma_smfree( &dU, queue );
magma_smfree( &dM, queue );
magma_smfree( &df, queue );
magma_smfree( &dt, queue );
magma_smfree( &dc, queue );
magma_smfree( &dv, queue );
magma_smfree(&dlu, queue);
magma_smfree( &dbeta, queue );
magma_smfree( &hbeta, queue );
solver_par->info = info;
return info;
/* magma_spidr */
}
extern "C" magma_int_t
magma_sbicgstab_merge3(
magma_s_matrix A, magma_s_matrix b,
magma_s_matrix *x, magma_s_solver_par *solver_par,
magma_queue_t queue )
{
magma_int_t info = MAGMA_NOTCONVERGED;
// prepare solver feedback
solver_par->solver = Magma_BICGSTABMERGE;
solver_par->numiter = 0;
solver_par->spmv_count = 0;
// solver variables
float alpha, beta, omega, rho_old, rho_new, *skp_h={0};
float nom, nom0, betanom, nomb;
// some useful variables
float c_zero = MAGMA_S_ZERO, c_one = MAGMA_S_ONE;
magma_int_t dofs = A.num_rows;
// workspace
magma_s_matrix q={Magma_CSR}, r={Magma_CSR}, rr={Magma_CSR}, p={Magma_CSR}, v={Magma_CSR}, s={Magma_CSR}, t={Magma_CSR};
float *d1=NULL, *d2=NULL, *skp=NULL;
d1 = NULL;
d2 = NULL;
skp = NULL;
CHECK( magma_smalloc( &d1, dofs*(2) ));
CHECK( magma_smalloc( &d2, dofs*(2) ));
// array for the parameters
CHECK( magma_smalloc( &skp, 8 ));
// skp = [alpha|beta|omega|rho_old|rho|nom|tmp1|tmp2]
CHECK( magma_svinit( &q, Magma_DEV, dofs*6, 1, c_zero, queue ));
// q = rr|r|p|v|s|t
rr.memory_location = Magma_DEV; rr.dval = NULL; rr.num_rows = rr.nnz = dofs; rr.num_cols = 1; rr.storage_type = Magma_DENSE;
r.memory_location = Magma_DEV; r.dval = NULL; r.num_rows = r.nnz = dofs; r.num_cols = 1; r.storage_type = Magma_DENSE;
p.memory_location = Magma_DEV; p.dval = NULL; p.num_rows = p.nnz = dofs; p.num_cols = 1; p.storage_type = Magma_DENSE;
v.memory_location = Magma_DEV; v.dval = NULL; v.num_rows = v.nnz = dofs; v.num_cols = 1; v.storage_type = Magma_DENSE;
s.memory_location = Magma_DEV; s.dval = NULL; s.num_rows = s.nnz = dofs; s.num_cols = 1; s.storage_type = Magma_DENSE;
t.memory_location = Magma_DEV; t.dval = NULL; t.num_rows = t.nnz = dofs; t.num_cols = 1; t.storage_type = Magma_DENSE;
rr.dval = q(0);
r.dval = q(1);
p.dval = q(2);
v.dval = q(3);
s.dval = q(4);
t.dval = q(5);
// solver setup
CHECK( magma_sresidualvec( A, b, *x, &r, &nom0, queue));
magma_scopy( dofs, r.dval, 1, q(0), 1, queue ); // rr = r
magma_scopy( dofs, r.dval, 1, q(1), 1, queue ); // q = r
betanom = nom0;
nom = nom0*nom0;
rho_new = magma_sdot( dofs, r.dval, 1, r.dval, 1, queue ); // rho=<rr,r>
rho_old = omega = alpha = MAGMA_S_MAKE( 1.0, 0. );
beta = rho_new;
solver_par->init_res = nom0;
// array on host for the parameters
CHECK( magma_smalloc_cpu( &skp_h, 8 ));
nomb = magma_snrm2( dofs, b.dval, 1, queue );
if ( nomb == 0.0 ){
nomb=1.0;
}
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;
}
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_S_MAKE(nom, 0.0);
magma_ssetvector( 8, skp_h, 1, skp, 1, queue );
CHECK( magma_s_spmv( c_one, A, r, c_zero, v, queue )); // z = A r
nomb = magma_snrm2( dofs, b.dval, 1, queue );
if( nom0 < solver_par->atol ||
nom0/nomb < solver_par->rtol ){
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++;
// computes p=r+beta*(p-omega*v)
CHECK( magma_sbicgmerge1( dofs, skp, v.dval, r.dval, p.dval, queue ));
CHECK( magma_s_spmv( c_one, A, p, c_zero, v, queue )); // v = Ap
solver_par->spmv_count++;
CHECK( magma_smdotc( dofs, 1, q.dval, v.dval, d1, d2, skp, queue ));
CHECK( magma_sbicgmerge4( 1, skp, queue ));
CHECK( magma_sbicgmerge2( dofs, skp, r.dval, v.dval, s.dval, queue )); // s=r-alpha*v
CHECK( magma_s_spmv( c_one, A, s, c_zero, t, queue )); // t=As
solver_par->spmv_count++;
CHECK( magma_smdotc( dofs, 2, q.dval+4*dofs, t.dval, d1, d2, skp+6, queue ));
CHECK( magma_sbicgmerge4( 2, skp, queue ));
CHECK( magma_sbicgmerge_xrbeta( dofs, d1, d2, q.dval, r.dval, p.dval,
s.dval, t.dval, x->dval, skp, queue ));
// check stopping criterion
magma_sgetvector_async( 1 , skp+5, 1, skp_h+5, 1, queue );
betanom = sqrt(MAGMA_S_REAL(skp_h[5]));
if ( solver_par->verbose > 0 ) {
tempo2 = magma_sync_wtime( queue );
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
if ( betanom < solver_par->atol ||
betanom/nomb < solver_par->rtol ) {
break;
}
}
while ( solver_par->numiter+1 <= solver_par->maxiter );
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t) tempo2-tempo1;
float residual;
CHECK( magma_sresidualvec( 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;
}
}
info = MAGMA_DIVERGENCE;
}
cleanup:
magma_smfree(&q, queue ); // frees all vectors
magma_free(d1);
magma_free(d2);
magma_free( skp );
magma_free_cpu( skp_h );
solver_par->info = info;
return info;
} /* sbicgstab_merge */
/* ////////////////////////////////////////////////////////////////////////////
-- testing any solver
*/
int main( int argc, char** argv )
{
magma_int_t info = 0;
/* Initialize */
TESTING_INIT();
magma_queue_t queue=NULL;
magma_queue_create( &queue );
magmablasSetKernelStream( queue );
magma_int_t j, n=1000000, FLOPS;
float one = MAGMA_S_MAKE( 1.0, 0.0 );
float two = MAGMA_S_MAKE( 2.0, 0.0 );
magma_s_matrix a={Magma_CSR}, ad={Magma_CSR}, bd={Magma_CSR}, cd={Magma_CSR};
CHECK( magma_svinit( &a, Magma_CPU, n, 1, one, queue ));
CHECK( magma_svinit( &bd, Magma_DEV, n, 1, two, queue ));
CHECK( magma_svinit( &cd, Magma_DEV, n, 1, one, queue ));
CHECK( magma_smtransfer( a, &ad, Magma_CPU, Magma_DEV, queue ));
real_Double_t start, end, res;
FLOPS = 2*n;
start = magma_sync_wtime( queue );
for (j=0; j<100; j++)
res = magma_snrm2(n, ad.dval, 1);
end = magma_sync_wtime( queue );
printf( " > MAGMA nrm2: %.2e seconds %.2e GFLOP/s\n",
(end-start)/100, FLOPS*100/1e9/(end-start) );
FLOPS = n;
start = magma_sync_wtime( queue );
for (j=0; j<100; j++)
magma_sscal( n, two, ad.dval, 1 );
end = magma_sync_wtime( queue );
printf( " > MAGMA scal: %.2e seconds %.2e GFLOP/s\n",
(end-start)/100, FLOPS*100/1e9/(end-start) );
FLOPS = 2*n;
start = magma_sync_wtime( queue );
for (j=0; j<100; j++)
magma_saxpy( n, one, ad.dval, 1, bd.dval, 1 );
end = magma_sync_wtime( queue );
printf( " > MAGMA axpy: %.2e seconds %.2e GFLOP/s\n",
(end-start)/100, FLOPS*100/1e9/(end-start) );
FLOPS = n;
start = magma_sync_wtime( queue );
for (j=0; j<100; j++)
magma_scopy( n, bd.dval, 1, ad.dval, 1 );
end = magma_sync_wtime( queue );
printf( " > MAGMA copy: %.2e seconds %.2e GFLOP/s\n",
(end-start)/100, FLOPS*100/1e9/(end-start) );
FLOPS = 2*n;
start = magma_sync_wtime( queue );
for (j=0; j<100; j++)
res = MAGMA_S_REAL( magma_sdot(n, ad.dval, 1, bd.dval, 1) );
end = magma_sync_wtime( queue );
printf( " > MAGMA dotc: %.2e seconds %.2e GFLOP/s\n",
(end-start)/100, FLOPS*100/1e9/(end-start) );
printf("# tester BLAS: ok\n");
magma_smfree( &a, queue);
magma_smfree(&ad, queue);
magma_smfree(&bd, queue);
magma_smfree(&cd, queue);
cleanup:
magma_smfree( &a, queue);
magma_smfree(&ad, queue);
magma_smfree(&bd, queue);
magma_smfree(&cd, queue);
magmablasSetKernelStream( NULL );
magma_queue_destroy( queue );
magma_finalize();
return info;
}
/* ////////////////////////////////////////////////////////////////////////////
-- testing zdot
*/
int main( int argc, char** argv )
{
magma_int_t info = 0;
magma_queue_t queue=NULL;
magma_queue_create( 0, &queue );
const float one = MAGMA_S_MAKE(1.0, 0.0);
const float zero = MAGMA_S_MAKE(0.0, 0.0);
float alpha;
TESTING_INIT();
magma_s_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;
float 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_svinit( &a, Magma_DEV, n, num_vecs, one, queue ));
CHECK( magma_svinit( &b, Magma_DEV, n, 1, one, queue ));
CHECK( magma_svinit( &x, Magma_DEV, n, 8, one, queue ));
CHECK( magma_svinit( &y, Magma_DEV, n, 8, one, queue ));
CHECK( magma_svinit( &skp, Magma_DEV, 1, num_vecs, zero, queue ));
// warm up
CHECK( magma_sgemvmdot( 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_sdot( 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_sgemv( 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_sgemv( 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_smdotc( n, 2, a.dval, b.dval, x.dval, y.dval, skp.dval, queue ));
}
for( int c = 0; c<num_vecs%2; c++ ){
CHECK( magma_smdotc( 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_sgemvmdot( 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_sgemvmdot_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_sprint_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_smfree(&a, queue );
magma_smfree(&b, queue );
magma_smfree(&x, queue );
magma_smfree(&y, queue );
magma_smfree(&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;
}
