/* ////////////////////////////////////////////////////////////////////////////
-- testing zdot
*/
int main( int argc, char** argv )
{
magma_int_t info = 0;
// set queue for old dense routines
magma_queue_t queue=NULL;
magma_queue_create( /*devices[ opts->device ],*/ &queue );
magmablasGetKernelStream( &queue );
TESTING_INIT();
magma_c_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 | CUDOT CUGEMV MAGMAGEMV MDOT MDGM \n");
printf("#------------------------------------------------------------------------------------------------------------------------------------------------\n");
printf("\n");
for( magma_int_t num_vecs=5; num_vecs<6; num_vecs+=1 ) {
for( magma_int_t n=10000; n<100000001; n=n+10000 ) {
int iters = 10;
float computations = (2.* n * iters * num_vecs);
magmaFloatComplex one = MAGMA_C_MAKE(1.0, 0.0);
magmaFloatComplex zero = MAGMA_C_MAKE(0.0, 0.0);
magmaFloatComplex alpha;
#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, magmagemv_time, cugemv_time, cudot_time;
#endif
CHECK( magma_cvinit( &a, Magma_DEV, n, num_vecs, one, queue ));
CHECK( magma_cvinit( &b, Magma_DEV, num_vecs, 1, one, queue ));
int min_ten = min(num_vecs, 15);
CHECK( magma_cvinit( &x, Magma_DEV, min_ten, n, one, queue ));
CHECK( magma_cvinit( &y, Magma_DEV, min_ten, n, one, queue ));
CHECK( magma_cvinit( &skp, Magma_DEV, num_vecs, 1, zero, queue ));
// warm up
CHECK( magma_cgemvmdot( 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_cdotc(n, a.dval, 1, b.dval, 1);
}
#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_cgemv(MagmaTrans, n, num_vecs, one, a.dval, n, b.dval, 1, zero, skp.dval, 1);
//h++;
}
#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_cgemv(MagmaTrans, n, num_vecs, one, a.dval, n, b.dval, 1, zero, skp.dval, 1);
//h++;
}
#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++) {
//magma_cmdotc( n, num_vecs, a.dval, b.dval, x.dval, y.dval, skp.dval, queue );
CHECK( magma_cmdotc( n, 2, a.dval, b.dval, x.dval, y.dval, skp.dval, queue ));
CHECK( magma_cmdotc( n, 2, a.dval, b.dval, x.dval, y.dval, skp.dval, queue ));
CHECK( magma_cmdotc( 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_cgemvmdot( 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
//magma_cprint_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\n",
n, num_vecs,
cudot_time/iters,
(cugemv_time)/iters,
(magmagemv_time)/iters,
(mdot_time)/iters,
(mdgm_time)/iters,
(float)(computations)/(cudot_time*(1.e+09)),
(float)(computations)/(cugemv_time*(1.e+09)),
(float)(computations)/(magmagemv_time*(1.e+09)),
(float)(computations)/(mdot_time*(1.e+09)),
(float)(computations)/(mdgm_time*(1.e+09)) );
#endif
magma_cmfree(&a, queue );
magma_cmfree(&b, queue );
magma_cmfree(&x, queue );
magma_cmfree(&y, queue );
magma_cmfree(&skp, queue );
}
printf("#================================================================================================================================================\n");
printf("\n");
printf("\n");
}
cleanup:
magma_cmfree(&a, queue );
magma_cmfree(&b, queue );
magma_cmfree(&x, queue );
magma_cmfree(&y, queue );
magma_cmfree(&skp, queue );
magma_queue_destroy( queue );
TESTING_FINALIZE();
return info;
}
int main(int argc, char **argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, cublas_perf, cublas_time, cpu_perf, cpu_time;
float magma_error, cublas_error, work[1];
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t M, N, Xm, Ym, lda, sizeA, sizeX, sizeY;
magma_int_t incx = 1;
magma_int_t incy = 1;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex alpha = MAGMA_C_MAKE( 1.5, -2.3 );
magmaFloatComplex beta = MAGMA_C_MAKE( -0.6, 0.8 );
magmaFloatComplex *A, *X, *Y, *Ycublas, *Ymagma;
magmaFloatComplex *dA, *dX, *dY;
magma_opts opts;
parse_opts( argc, argv, &opts );
printf(" M N MAGMA Gflop/s (ms) CUBLAS Gflop/s (ms) CPU Gflop/s (ms) MAGMA error CUBLAS error\n");
printf("===================================================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[i];
N = opts.nsize[i];
lda = ((M+31)/32)*32;
gflops = FLOPS_CGEMV( M, N ) / 1e9;
if ( opts.transA == MagmaNoTrans ) {
Xm = N;
Ym = M;
} else {
Xm = M;
Ym = N;
}
sizeA = lda*N;
sizeX = incx*Xm;
sizeY = incy*Ym;
TESTING_MALLOC_CPU( A, magmaFloatComplex, sizeA );
TESTING_MALLOC_CPU( X, magmaFloatComplex, sizeX );
TESTING_MALLOC_CPU( Y, magmaFloatComplex, sizeY );
TESTING_MALLOC_CPU( Ycublas, magmaFloatComplex, sizeY );
TESTING_MALLOC_CPU( Ymagma, magmaFloatComplex, sizeY );
TESTING_MALLOC_DEV( dA, magmaFloatComplex, sizeA );
TESTING_MALLOC_DEV( dX, magmaFloatComplex, sizeX );
TESTING_MALLOC_DEV( dY, magmaFloatComplex, sizeY );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &sizeA, A );
lapackf77_clarnv( &ione, ISEED, &sizeX, X );
lapackf77_clarnv( &ione, ISEED, &sizeY, Y );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_csetmatrix( M, N, A, lda, dA, lda );
magma_csetvector( Xm, X, incx, dX, incx );
magma_csetvector( Ym, Y, incy, dY, incy );
cublas_time = magma_sync_wtime( 0 );
cublasCgemv( opts.transA, M, N, alpha, dA, lda, dX, incx, beta, dY, incy );
cublas_time = magma_sync_wtime( 0 ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_cgetvector( Ym, dY, incy, Ycublas, incy );
/* =====================================================================
Performs operation using MAGMABLAS
=================================================================== */
magma_csetvector( Ym, Y, incy, dY, incy );
magma_time = magma_sync_wtime( 0 );
magmablas_cgemv( opts.transA, M, N, alpha, dA, lda, dX, incx, beta, dY, incy );
magma_time = magma_sync_wtime( 0 ) - magma_time;
magma_perf = gflops / magma_time;
magma_cgetvector( Ym, dY, incx, Ymagma, incx );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
cpu_time = magma_wtime();
blasf77_cgemv( &opts.transA, &M, &N,
&alpha, A, &lda,
X, &incx,
&beta, Y, &incy );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/* =====================================================================
Check the result
=================================================================== */
blasf77_caxpy( &Ym, &c_neg_one, Y, &incy, Ymagma, &incy );
magma_error = lapackf77_clange( "M", &Ym, &ione, Ymagma, &Ym, work ) / Ym;
blasf77_caxpy( &Ym, &c_neg_one, Y, &incy, Ycublas, &incy );
cublas_error = lapackf77_clange( "M", &Ym, &ione, Ycublas, &Ym, work ) / Ym;
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e\n",
(int) M, (int) N,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
magma_error, cublas_error );
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( X );
TESTING_FREE_CPU( Y );
TESTING_FREE_CPU( Ycublas );
TESTING_FREE_CPU( Ymagma );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dX );
TESTING_FREE_DEV( dY );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
extern "C" magma_int_t
magma_cidr(
magma_c_matrix A, magma_c_matrix b, magma_c_matrix *x,
magma_c_solver_par *solver_par,
magma_queue_t queue )
{
magma_int_t info = MAGMA_NOTCONVERGED;
// prepare solver feedback
solver_par->solver = Magma_IDR;
solver_par->numiter = 0;
solver_par->spmv_count = 0;
solver_par->init_res = 0.0;
solver_par->final_res = 0.0;
solver_par->iter_res = 0.0;
solver_par->runtime = 0.0;
// constants
const magmaFloatComplex c_zero = MAGMA_C_ZERO;
const magmaFloatComplex c_one = MAGMA_C_ONE;
const magmaFloatComplex c_n_one = MAGMA_C_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;
magmaFloatComplex om;
magmaFloatComplex tt;
magmaFloatComplex tr;
magmaFloatComplex gamma;
magmaFloatComplex alpha;
magmaFloatComplex mkk;
magmaFloatComplex fk;
// matrices and vectors
magma_c_matrix dxs = {Magma_CSR};
magma_c_matrix dr = {Magma_CSR}, drs = {Magma_CSR};
magma_c_matrix dP = {Magma_CSR}, dP1 = {Magma_CSR};
magma_c_matrix dG = {Magma_CSR};
magma_c_matrix dU = {Magma_CSR};
magma_c_matrix dM = {Magma_CSR};
magma_c_matrix df = {Magma_CSR};
magma_c_matrix dt = {Magma_CSR};
magma_c_matrix dc = {Magma_CSR};
magma_c_matrix dv = {Magma_CSR};
magma_c_matrix dbeta = {Magma_CSR}, hbeta = {Magma_CSR};
// chronometry
real_Double_t tempo1, tempo2;
// initial s space
// TODO: add option for 's' (shadow space number)
// Hack: uses '--restart' option as the shadow space number.
// This is not a good idea because the default value of restart option is used to detect
// if the user provided a custom restart. This means that if the default restart value
// is changed then the code will think it was the user (unless the default value is
// also updated in the 'if' statement below.
s = 1;
if ( solver_par->restart != 50 ) {
if ( solver_par->restart > A.num_cols ) {
s = A.num_cols;
} else {
s = solver_par->restart;
}
}
solver_par->restart = s;
// set max iterations
solver_par->maxiter = min( 2 * A.num_cols, solver_par->maxiter );
// check if matrix A is square
if ( A.num_rows != A.num_cols ) {
//printf("Matrix A is not square.\n");
info = MAGMA_ERR_NOT_SUPPORTED;
goto cleanup;
}
// |b|
nrmb = magma_scnrm2( b.num_rows, b.dval, 1, queue );
if ( nrmb == 0.0 ) {
magma_cscal( x->num_rows, MAGMA_C_ZERO, x->dval, 1, queue );
info = MAGMA_SUCCESS;
goto cleanup;
}
// r = b - A x
CHECK( magma_cvinit( &dr, Magma_DEV, b.num_rows, 1, c_zero, queue ));
CHECK( magma_cresidualvec( 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_cvinit( &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_clarnv( &distr, iseed, &dof, dP.val );
// transfer P to device
CHECK( magma_cmtransfer( dP, &dP1, Magma_CPU, Magma_DEV, queue ));
magma_cmfree( &dP, queue );
// P = ortho(P1)
if ( dP1.num_cols > 1 ) {
// P = magma_cqr(P1), QR factorization
CHECK( magma_cqr( dP1.num_rows, dP1.num_cols, dP1, dP1.ld, &dP, NULL, queue ));
} else {
// P = P1 / |P1|
nrm = magma_scnrm2( dof, dP1.dval, 1, queue );
nrm = 1.0 / nrm;
magma_csscal( dof, nrm, dP1.dval, 1, queue );
CHECK( magma_cmtransfer( dP1, &dP, Magma_DEV, Magma_DEV, queue ));
}
magma_cmfree( &dP1, queue );
//---------------------------------------
// allocate memory for the scalar products
CHECK( magma_cvinit( &hbeta, Magma_CPU, s, 1, c_zero, queue ));
CHECK( magma_cvinit( &dbeta, Magma_DEV, s, 1, c_zero, queue ));
// smoothing enabled
if ( smoothing > 0 ) {
// set smoothing solution vector
CHECK( magma_cmtransfer( *x, &dxs, Magma_DEV, Magma_DEV, queue ));
// set smoothing residual vector
CHECK( magma_cmtransfer( dr, &drs, Magma_DEV, Magma_DEV, queue ));
}
// G(n,s) = 0
CHECK( magma_cvinit( &dG, Magma_DEV, A.num_cols, s, c_zero, queue ));
// U(n,s) = 0
CHECK( magma_cvinit( &dU, Magma_DEV, A.num_cols, s, c_zero, queue ));
// M(s,s) = I
CHECK( magma_cvinit( &dM, Magma_DEV, s, s, c_zero, queue ));
magmablas_claset( MagmaFull, s, s, c_zero, c_one, dM.dval, s, queue );
// f = 0
CHECK( magma_cvinit( &df, Magma_DEV, dP.num_cols, 1, c_zero, queue ));
// t = 0
CHECK( magma_cvinit( &dt, Magma_DEV, dr.num_rows, 1, c_zero, queue ));
// c = 0
CHECK( magma_cvinit( &dc, Magma_DEV, dM.num_cols, 1, c_zero, queue ));
// v = 0
CHECK( magma_cvinit( &dv, Magma_DEV, dr.num_rows, 1, c_zero, queue ));
//--------------START TIME---------------
// chronometry
tempo1 = magma_sync_wtime( queue );
if ( solver_par->verbose > 0 ) {
solver_par->timing[0] = 0.0;
}
om = MAGMA_C_ONE;
innerflag = 0;
// start iteration
do
{
solver_par->numiter++;
// new RHS for small systems
// f = P' r
magmablas_cgemv( 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_ccopyvector( sk, &df.dval[k], 1, &dc.dval[k], 1, queue );
magma_ctrsv( 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_ccopyvector( dr.num_rows, dr.dval, 1, dv.dval, 1, queue );
magmablas_cgemv( MagmaNoTrans, dG.num_rows, sk, c_n_one, &dG.dval[k*dG.ld], dG.ld, &dc.dval[k], 1, c_one, dv.dval, 1, queue );
// U(:,k) = om * v + U(:,k:s) c(k:s)
magmablas_cgemv( MagmaNoTrans, dU.num_rows, sk, c_one, &dU.dval[k*dU.ld], dU.ld, &dc.dval[k], 1, om, dv.dval, 1, queue );
magma_ccopyvector( dU.num_rows, dv.dval, 1, &dU.dval[k*dU.ld], 1, queue );
// G(:,k) = A U(:,k)
CHECK( magma_c_spmv( c_one, A, dv, c_zero, dv, queue ));
solver_par->spmv_count++;
magma_ccopyvector( 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_cdotc( dP.num_rows, &dP.dval[i*dP.ld], 1, &dG.dval[k*dG.ld], 1, queue );
// alpha = alpha / M(i,i)
magma_cgetvector( 1, &dM.dval[i*dM.ld+i], 1, &mkk, 1, queue );
alpha = alpha / mkk;
// G(:,k) = G(:,k) - alpha * G(:,i)
magma_caxpy( 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_caxpy( 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_cgemv( 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_cgetvector( 1, &dM.dval[k*dM.ld+k], 1, &mkk, 1, queue );
if ( MAGMA_C_EQUAL(mkk, MAGMA_C_ZERO) ) {
innerflag = 1;
info = MAGMA_DIVERGENCE;
break;
}
// beta = f(k) / M(k,k)
magma_cgetvector( 1, &df.dval[k], 1, &fk, 1, queue );
hbeta.val[k] = fk / mkk;
// check for nan
if ( magma_c_isnan( hbeta.val[k] ) || magma_c_isinf( hbeta.val[k] )) {
innerflag = 1;
info = MAGMA_DIVERGENCE;
break;
}
// r = r - beta * G(:,k)
magma_caxpy( dr.num_rows, -hbeta.val[k], &dG.dval[k*dG.ld], 1, dr.dval, 1, queue );
// smoothing disabled
if ( smoothing <= 0 ) {
// |r|
nrmr = magma_scnrm2( dr.num_rows, dr.dval, 1, queue );
// smoothing enabled
} else {
// x = x + beta * U(:,k)
magma_caxpy( x->num_rows, hbeta.val[k], &dU.dval[k*dU.ld], 1, x->dval, 1, queue );
// smoothing operation
//---------------------------------------
// t = rs - r
magma_ccopyvector( drs.num_rows, drs.dval, 1, dt.dval, 1, queue );
magma_caxpy( dt.num_rows, c_n_one, dr.dval, 1, dt.dval, 1, queue );
// t't
// t'rs
tt = magma_cdotc( dt.num_rows, dt.dval, 1, dt.dval, 1, queue );
tr = magma_cdotc( dt.num_rows, dt.dval, 1, drs.dval, 1, queue );
// gamma = (t' * rs) / (t' * t)
gamma = tr / tt;
// rs = rs - gamma * (rs - r)
magma_caxpy( drs.num_rows, -gamma, dt.dval, 1, drs.dval, 1, queue );
// xs = xs - gamma * (xs - x)
magma_ccopyvector( dxs.num_rows, dxs.dval, 1, dt.dval, 1, queue );
magma_caxpy( dt.num_rows, c_n_one, x->dval, 1, dt.dval, 1, queue );
magma_caxpy( dxs.num_rows, -gamma, dt.dval, 1, dxs.dval, 1, queue );
// |rs|
nrmr = magma_scnrm2( 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_caxpy( 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_csetvector( s, hbeta.val, 1, dbeta.dval, 1, queue );
magmablas_cgemv( MagmaNoTrans, dU.num_rows, s, c_one, dU.dval, dU.ld, dbeta.dval, 1, c_one, x->dval, 1, queue );
}
// check convergence or iteration limit or invalid result of inner loop
if ( innerflag > 0 ) {
break;
}
// t = A v
// t = A r
CHECK( magma_c_spmv( c_one, A, dr, c_zero, dt, queue ));
solver_par->spmv_count++;
// computation of a new omega
//---------------------------------------
// |t|
nrmt = magma_scnrm2( dt.num_rows, dt.dval, 1, queue );
// t'r
tr = magma_cdotc( dt.num_rows, dt.dval, 1, dr.dval, 1, queue );
// rho = abs(t' * r) / (|t| * |r|))
rho = MAGMA_D_ABS( MAGMA_C_REAL(tr) / (nrmt * nrmr) );
// om = (t' * r) / (|t| * |t|)
om = tr / (nrmt * nrmt);
if ( rho < angle ) {
om = (om * angle) / rho;
}
//---------------------------------------
if ( MAGMA_C_EQUAL(om, MAGMA_C_ZERO) ) {
info = MAGMA_DIVERGENCE;
break;
}
// update approximation vector
// x = x + om * v
// x = x + om * r
magma_caxpy( x->num_rows, om, dr.dval, 1, x->dval, 1, queue );
// update residual vector
// r = r - om * t
magma_caxpy( dr.num_rows, -om, dt.dval, 1, dr.dval, 1, queue );
// smoothing disabled
if ( smoothing <= 0 ) {
// residual norm
nrmr = magma_scnrm2( b.num_rows, dr.dval, 1, queue );
// smoothing enabled
} else {
// smoothing operation
//---------------------------------------
// t = rs - r
magma_ccopyvector( drs.num_rows, drs.dval, 1, dt.dval, 1, queue );
magma_caxpy( dt.num_rows, c_n_one, dr.dval, 1, dt.dval, 1, queue );
// t't
// t'rs
tt = magma_cdotc( dt.num_rows, dt.dval, 1, dt.dval, 1, queue );
tr = magma_cdotc( dt.num_rows, dt.dval, 1, drs.dval, 1, queue );
// gamma = (t' * rs) / (|t| * |t|)
gamma = tr / tt;
// rs = rs - gamma * (rs - r)
magma_caxpy( drs.num_rows, -gamma, dt.dval, 1, drs.dval, 1, queue );
// xs = xs - gamma * (xs - x)
magma_ccopyvector( dxs.num_rows, dxs.dval, 1, dt.dval, 1, queue );
magma_caxpy( dt.num_rows, c_n_one, x->dval, 1, dt.dval, 1, queue );
magma_caxpy( dxs.num_rows, -gamma, dt.dval, 1, dxs.dval, 1, queue );
// |rs|
nrmr = magma_scnrm2( 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_ccopyvector( x->num_rows, dxs.dval, 1, x->dval, 1, queue );
// r = rs
magma_ccopyvector( 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_cresidualvec( 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_cmfree( &dxs, queue );
magma_cmfree( &drs, queue );
}
magma_cmfree( &dr, queue );
magma_cmfree( &dP, queue );
magma_cmfree( &dP1, queue );
magma_cmfree( &dG, queue );
magma_cmfree( &dU, queue );
magma_cmfree( &dM, queue );
magma_cmfree( &df, queue );
magma_cmfree( &dt, queue );
magma_cmfree( &dc, queue );
magma_cmfree( &dv, queue );
magma_cmfree( &dbeta, queue );
magma_cmfree( &hbeta, queue );
solver_par->info = info;
return info;
/* magma_cidr */
}
magma_int_t
magma_c_spmv( magmaFloatComplex alpha, magma_c_sparse_matrix A,
magma_c_vector x, magmaFloatComplex beta, magma_c_vector y )
{
if( A.memory_location != x.memory_location ||
x.memory_location != y.memory_location ){
printf("error: linear algebra objects are not located in same memory!\n");
printf("memory locations are: %d %d %d\n",
A.memory_location, x.memory_location, y.memory_location );
return MAGMA_ERR_INVALID_PTR;
}
// DEV case
if( A.memory_location == Magma_DEV ){
if( A.num_cols == x.num_rows ){
if( A.storage_type == Magma_CSR
|| A.storage_type == Magma_CSRL
|| A.storage_type == Magma_CSRU ){
//printf("using CSR kernel for SpMV: ");
magma_cgecsrmv( MagmaNoTrans, A.num_rows, A.num_cols, alpha,
A.val, A.row, A.col, x.val, beta, y.val );
//printf("done.\n");
return MAGMA_SUCCESS;
}
else if( A.storage_type == Magma_ELLPACK ){
//printf("using ELLPACK kernel for SpMV: ");
magma_cgeellmv( MagmaNoTrans, A.num_rows, A.num_cols,
A.max_nnz_row, alpha, A.val, A.col, x.val, beta, y.val );
//printf("done.\n");
return MAGMA_SUCCESS;
}
else if( A.storage_type == Magma_ELL ){
//printf("using ELL kernel for SpMV: ");
magma_cgeelltmv( MagmaNoTrans, A.num_rows, A.num_cols,
A.max_nnz_row, alpha, A.val, A.col, x.val, beta, y.val );
//printf("done.\n");
return MAGMA_SUCCESS;
}
else if( A.storage_type == Magma_ELLRT ){
//printf("using ELLRT kernel for SpMV: ");
magma_cgeellrtmv( MagmaNoTrans, A.num_rows, A.num_cols,
A.max_nnz_row, alpha, A.val, A.col, A.row, x.val,
beta, y.val, A.alignment, A.blocksize );
//printf("done.\n");
return MAGMA_SUCCESS;
}
else if( A.storage_type == Magma_SELLC ){
//printf("using SELLC kernel for SpMV: ");
magma_cgesellcmv( MagmaNoTrans, A.num_rows, A.num_cols,
A.blocksize, A.numblocks, A.alignment,
alpha, A.val, A.col, A.row, x.val, beta, y.val );
//printf("done.\n");
return MAGMA_SUCCESS;
}
else if( A.storage_type == Magma_SELLP ){
//printf("using SELLP kernel for SpMV: ");
magma_cgesellpmv( MagmaNoTrans, A.num_rows, A.num_cols,
A.blocksize, A.numblocks, A.alignment,
alpha, A.val, A.col, A.row, x.val, beta, y.val );
//printf("done.\n");
return MAGMA_SUCCESS;
}
else if( A.storage_type == Magma_DENSE ){
//printf("using DENSE kernel for SpMV: ");
magmablas_cgemv( MagmaNoTrans, A.num_rows, A.num_cols, alpha,
A.val, A.num_rows, x.val, 1, beta, y.val, 1 );
//printf("done.\n");
return MAGMA_SUCCESS;
}
/* else if( A.storage_type == Magma_BCSR ){
//printf("using CUSPARSE BCSR kernel for SpMV: ");
// CUSPARSE context //
cusparseHandle_t cusparseHandle = 0;
cusparseStatus_t cusparseStatus;
cusparseStatus = cusparseCreate(&cusparseHandle);
cusparseMatDescr_t descr = 0;
cusparseStatus = cusparseCreateMatDescr(&descr);
// end CUSPARSE context //
cusparseDirection_t dirA = CUSPARSE_DIRECTION_ROW;
int mb = (A.num_rows + A.blocksize-1)/A.blocksize;
int nb = (A.num_cols + A.blocksize-1)/A.blocksize;
cusparseCbsrmv( cusparseHandle, dirA,
CUSPARSE_OPERATION_NON_TRANSPOSE, mb, nb, A.numblocks,
&alpha, descr, A.val, A.row, A.col, A.blocksize, x.val,
&beta, y.val );
//printf("done.\n");
return MAGMA_SUCCESS;
}*/
else {
printf("error: format not supported.\n");
return MAGMA_ERR_NOT_SUPPORTED;
}
}
else if( A.num_cols < x.num_rows ){
magma_int_t num_vecs = x.num_rows / A.num_cols;
if( A.storage_type == Magma_CSR ){
//printf("using CSR kernel for SpMV: ");
magma_cmgecsrmv( MagmaNoTrans, A.num_rows, A.num_cols,
num_vecs, alpha, A.val, A.row, A.col, x.val, beta, y.val );
//printf("done.\n");
return MAGMA_SUCCESS;
}
else if( A.storage_type == Magma_ELLPACK ){
//printf("using ELLPACK kernel for SpMV: ");
magma_cmgeellmv( MagmaNoTrans, A.num_rows, A.num_cols,
num_vecs, A.max_nnz_row, alpha, A.val, A.col, x.val, beta, y.val );
//printf("done.\n");
return MAGMA_SUCCESS;
}
else if( A.storage_type == Magma_ELL ){
//printf("using ELL kernel for SpMV: ");
magma_cmgeelltmv( MagmaNoTrans, A.num_rows, A.num_cols,
num_vecs, A.max_nnz_row, alpha, A.val,
A.col, x.val, beta, y.val );
//printf("done.\n");
return MAGMA_SUCCESS;
}else if( A.storage_type == Magma_SELLP ){
//printf("using SELLP kernel for SpMV: ");
magma_cmgesellpmv( MagmaNoTrans, A.num_rows, A.num_cols,
num_vecs, A.blocksize, A.numblocks, A.alignment,
alpha, A.val, A.col, A.row, x.val, beta, y.val );
//printf("done.\n");
return MAGMA_SUCCESS;
}/*
if( A.storage_type == Magma_DENSE ){
//printf("using DENSE kernel for SpMV: ");
magmablas_cmgemv( MagmaNoTrans, A.num_rows, A.num_cols,
num_vecs, alpha, A.val, A.num_rows, x.val, 1,
beta, y.val, 1 );
//printf("done.\n");
return MAGMA_SUCCESS;
}*/
else {
printf("error: format not supported.\n");
return MAGMA_ERR_NOT_SUPPORTED;
}
}
}
// CPU case missing!
else{
printf("error: CPU not yet supported.\n");
return MAGMA_ERR_NOT_SUPPORTED;
}
return MAGMA_SUCCESS;
}
extern "C" magma_int_t
magma_cpidr_strms(
magma_c_matrix A, magma_c_matrix b, magma_c_matrix *x,
magma_c_solver_par *solver_par,
magma_c_preconditioner *precond_par,
magma_queue_t queue )
{
magma_int_t info = MAGMA_NOTCONVERGED;
// prepare solver feedback
solver_par->solver = Magma_PIDRMERGE;
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 magmaFloatComplex c_zero = MAGMA_C_ZERO;
const magmaFloatComplex c_one = MAGMA_C_ONE;
const magmaFloatComplex c_n_one = MAGMA_C_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;
magmaFloatComplex om;
magmaFloatComplex gamma;
// matrices and vectors
magma_c_matrix dxs = {Magma_CSR};
magma_c_matrix dr = {Magma_CSR}, drs = {Magma_CSR};
magma_c_matrix dP = {Magma_CSR}, dP1 = {Magma_CSR};
magma_c_matrix dG = {Magma_CSR}, dGcol = {Magma_CSR};
magma_c_matrix dU = {Magma_CSR};
magma_c_matrix dM = {Magma_CSR};
magma_c_matrix df = {Magma_CSR};
magma_c_matrix dt = {Magma_CSR}, dtt = {Magma_CSR};
magma_c_matrix dc = {Magma_CSR};
magma_c_matrix dv = {Magma_CSR};
magma_c_matrix dlu = {Magma_CSR};
magma_c_matrix dskp = {Magma_CSR};
magma_c_matrix dalpha = {Magma_CSR};
magma_c_matrix dbeta = {Magma_CSR};
magmaFloatComplex *hMdiag = NULL;
magmaFloatComplex *hskp = NULL;
magmaFloatComplex *halpha = NULL;
magmaFloatComplex *hbeta = NULL;
magmaFloatComplex *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_scnrm2( b.num_rows, b.dval, 1, queue );
if ( nrmb == 0.0 ) {
magma_cscal( x->num_rows, MAGMA_C_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_cvinit( &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_cvinit( &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_cresidualvec( 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_cvinit( &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_clarnv( &distr, iseed, &dof, dP.val );
// transfer P to device
CHECK( magma_cmtransfer( dP, &dP1, Magma_CPU, Magma_DEV, queue ));
magma_cmfree( &dP, queue );
// P = ortho(P1)
if ( dP1.num_cols > 1 ) {
// P = magma_cqr(P1), QR factorization
CHECK( magma_cqr( dP1.num_rows, dP1.num_cols, dP1, dP1.ld, &dP, NULL, queue ));
} else {
// P = P1 / |P1|
nrm = magma_scnrm2( dof, dP1.dval, 1, queue );
nrm = 1.0 / nrm;
magma_csscal( dof, nrm, dP1.dval, 1, queue );
CHECK( magma_cmtransfer( dP1, &dP, Magma_DEV, Magma_DEV, queue ));
}
magma_cmfree( &dP1, queue );
//---------------------------------------
// allocate memory for the scalar products
CHECK( magma_cmalloc_pinned( &hskp, 5 ));
CHECK( magma_cvinit( &dskp, Magma_DEV, 4, 1, c_zero, queue ));
CHECK( magma_cmalloc_pinned( &halpha, s ));
CHECK( magma_cvinit( &dalpha, Magma_DEV, s, 1, c_zero, queue ));
CHECK( magma_cmalloc_pinned( &hbeta, s ));
CHECK( magma_cvinit( &dbeta, Magma_DEV, s, 1, c_zero, queue ));
// workspace for merged dot product
CHECK( magma_cmalloc( &d1, max(2, s) * b.num_rows ));
CHECK( magma_cmalloc( &d2, max(2, s) * b.num_rows ));
// smoothing enabled
if ( smoothing > 0 ) {
// set smoothing solution vector
CHECK( magma_cmtransfer( *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_cvinit( &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_cvinit( &drs, Magma_DEV, dr.num_rows, 1, c_zero, queue ));
magma_free( drs.dval );
drs.dval = dtt.dval + ldd;
// set smoothing residual vector
magma_ccopyvector( 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_cvinit( &dG, Magma_DEV, ldd, s, c_zero, queue ));
dG.num_rows = A.num_rows;
} else {
CHECK( magma_cvinit( &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_cvinit( &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_cvinit( &dU, Magma_DEV, ldd, s, c_zero, queue ));
dU.num_rows = A.num_cols;
} else {
CHECK( magma_cvinit( &dU, Magma_DEV, A.num_cols, s, c_zero, queue ));
}
// M(s,s) = I
CHECK( magma_cvinit( &dM, Magma_DEV, s, s, c_zero, queue ));
CHECK( magma_cmalloc_pinned( &hMdiag, s ));
magmablas_claset( MagmaFull, dM.num_rows, dM.num_cols, c_zero, c_one, dM.dval, dM.ld, queue );
// f = 0
CHECK( magma_cvinit( &df, Magma_DEV, dP.num_cols, 1, c_zero, queue ));
// c = 0
CHECK( magma_cvinit( &dc, Magma_DEV, dM.num_cols, 1, c_zero, queue ));
// v = r
CHECK( magma_cmtransfer( dr, &dv, Magma_DEV, Magma_DEV, queue ));
// lu = 0
CHECK( magma_cvinit( &dlu, Magma_DEV, dr.num_rows, 1, c_zero, queue ));
//--------------START TIME---------------
// chronometry
tempo1 = magma_sync_wtime( queue );
if ( solver_par->verbose > 0 ) {
solver_par->timing[0] = 0.0;
}
om = MAGMA_C_ONE;
gamma = MAGMA_C_ZERO;
innerflag = 0;
// start iteration
do
{
solver_par->numiter++;
// new RHS for small systems
// f = P' r
// Q1
magma_cgemvmdot_shfl( dP.num_rows, dP.num_cols, dP.dval, dr.dval, d1, d2, df.dval, queues[1] );
// skp[4] = f(k)
// Q1
magma_cgetvector_async( 1, df.dval, 1, &hskp[4], 1, queues[1] );
// c(k:s) = f(k:s)
// Q1
magma_ccopyvector_async( s, df.dval, 1, dc.dval, 1, queues[1] );
// c(k:s) = M(k:s,k:s) \ f(k:s)
// Q1
magma_ctrsv( MagmaLower, MagmaNoTrans, MagmaNonUnit, s, dM.dval, dM.ld, dc.dval, 1, queues[1] );
// 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_cgemv( MagmaNoTrans, dG.num_rows, sk, c_n_one, dGcol.dval, dG.ld, &dc.dval[k], 1, c_one, dv.dval, 1, queues[1] );
// preconditioning operation
// v = L \ v;
// v = U \ v;
// Q1
CHECK( magma_c_applyprecond_left( MagmaNoTrans, A, dv, &dlu, precond_par, queues[1] ));
CHECK( magma_c_applyprecond_right( MagmaNoTrans, A, dlu, &dv, precond_par, queues[1] ));
// sync Q0 --> U(:,k) = U(:,k) - U(:,1:k) * alpha(1:k)
magma_queue_sync( queues[0] );
// U(:,k) = om * v + U(:,k:s) c(k:s)
// Q1
magmablas_cgemv( 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_c_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_cdotc( 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_caxpy( dG.num_rows, -halpha[i], &dG.dval[i*dG.ld], 1, dGcol.dval, 1, queues[1] );
}
// sync Q1 --> compute new G, 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_cgemvmdot_shfl( dP.num_rows, sk, &dP.dval[k*dP.ld], dGcol.dval, d1, d2, &dM.dval[k*dM.ld+k], queues[2] );
// U(:,k) = v
// Q0
magma_ccopyvector_async( dU.num_rows, dv.dval, 1, &dU.dval[k*dU.ld], 1, queues[0] );
// non-first s iteration
if ( k > 0 ) {
// alpha = dalpha
// Q0
magma_csetvector_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_cgemv( MagmaNoTrans, dU.num_rows, k, c_n_one, dU.dval, dU.ld, dalpha.dval, 1, c_one, &dU.dval[k*dU.ld], 1, queues[0] );
}
// Mdiag(k) = M(k,k)
// Q2
magma_cgetvector( 1, &dM.dval[k*dM.ld+k], 1, &hMdiag[k], 1, queues[2] );
// implicit sync Q2 --> Mdiag(k) = M(k,k)
// check M(k,k) == 0
if ( MAGMA_C_EQUAL(hMdiag[k], MAGMA_C_ZERO) ) {
innerflag = 1;
info = MAGMA_DIVERGENCE;
break;
}
// beta = f(k) / M(k,k)
hbeta[k] = hskp[4] / hMdiag[k];
// check for nan
if ( magma_c_isnan( hbeta[k] ) || magma_c_isinf( hbeta[k] )) {
innerflag = 1;
info = MAGMA_DIVERGENCE;
break;
}
// non-last s iteration
if ( (k + 1) < s ) {
// f(k+1:s) = f(k+1:s) - beta * M(k+1:s,k)
// Q1
magma_caxpy( 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_ccopyvector_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_ctrsv( 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_cgetvector_async( 1, &df.dval[k+1], 1, &hskp[4], 1, queues[1] );
}
// r = r - beta * G(:,k)
// Q2
magma_caxpy( dr.num_rows, -hbeta[k], dGcol.dval, 1, dr.dval, 1, queues[2] );
// smoothing disabled
if ( smoothing <= 0 ) {
// |r|
// Q2
nrmr = magma_scnrm2( dr.num_rows, dr.dval, 1, queues[2] );
// implicit sync Q2 --> |r|
// v = r
// Q1
magma_ccopyvector_async( dr.num_rows, dr.dval, 1, dv.dval, 1, queues[1] );
// smoothing enabled
} else {
// x = x + beta * U(:,k)
// Q0
magma_caxpy( x->num_rows, hbeta[k], &dU.dval[k*dU.ld], 1, x->dval, 1, queues[0] );
// smoothing operation
//---------------------------------------
// t = rs - r
// Q2
magma_cidr_smoothing_1( drs.num_rows, drs.num_cols, drs.dval, dr.dval, dtt.dval, queues[2] );
// t't
// t'rs
// Q2
CHECK( magma_cgemvmdot_shfl( dt.ld, 2, dtt.dval, dtt.dval, d1, d2, &dskp.dval[2], queues[2] ));
// skp[2-3] = dskp[2-3]
// Q2
magma_cgetvector( 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];
// xs = xs - gamma * (xs - x)
// Q0
magma_cidr_smoothing_2( dxs.num_rows, dxs.num_cols, -gamma, x->dval, dxs.dval, queues[0] );
// v = r
// Q1
magma_ccopyvector_async( dr.num_rows, dr.dval, 1, dv.dval, 1, queues[1] );
// rs = rs - gamma * t
// Q2
magma_caxpy( drs.num_rows, -gamma, dtt.dval, 1, drs.dval, 1, queues[2] );
// |rs|
// Q2
nrmr = magma_scnrm2( drs.num_rows, drs.dval, 1, queues[2] );
// implicit sync Q2 --> |r|
//---------------------------------------
}
// 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_csetvector_async( s, hbeta, 1, dbeta.dval, 1, queues[0] );
// x = x + U(:,1:s) * beta(1:s)
// Q0
magmablas_cgemv( 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;
}
// preconditioning operation
// v = L \ v;
// v = U \ v;
// Q2
CHECK( magma_c_applyprecond_left( MagmaNoTrans, A, dv, &dlu, precond_par, queues[2] ));
CHECK( magma_c_applyprecond_right( MagmaNoTrans, A, dlu, &dv, precond_par, queues[2] ));
// t = A v
// Q2
CHECK( magma_c_spmv( c_one, A, dv, c_zero, dt, queues[2] ));
solver_par->spmv_count++;
// computation of a new omega
//---------------------------------------
// t't
// t'r
// Q2
CHECK( magma_cgemvmdot_shfl( dt.ld, 2, dt.dval, dt.dval, d1, d2, dskp.dval, queues[2] ));
// skp[0-2] = dskp[0-2]
// Q2
magma_cgetvector( 2, dskp.dval, 1, hskp, 1, queues[2] );
// implicit sync Q2 --> skp = dskp
// |t|
nrmt = magma_ssqrt( MAGMA_C_REAL(hskp[0]) );
// rho = abs((t' * r) / (|t| * |r|))
rho = MAGMA_D_ABS( MAGMA_C_REAL(hskp[1]) / (nrmt * nrmr) );
// om = (t' * r) / (|t| * |t|)
om = hskp[1] / hskp[0];
if ( rho < angle ) {
om = (om * angle) / rho;
}
//---------------------------------------
if ( MAGMA_C_EQUAL(om, MAGMA_C_ZERO) ) {
info = MAGMA_DIVERGENCE;
break;
}
// sync Q1 --> v = r
magma_queue_sync( queues[1] );
// r = r - om * t
// Q2
magma_caxpy( dr.num_rows, -om, dt.dval, 1, dr.dval, 1, queues[2] );
// x = x + om * v
// Q0
magma_caxpy( x->num_rows, om, dv.dval, 1, x->dval, 1, queues[0] );
// smoothing disabled
if ( smoothing <= 0 ) {
// |r|
// Q2
nrmr = magma_scnrm2( dr.num_rows, dr.dval, 1, queues[2] );
// implicit sync Q2 --> |r|
// v = r
// Q1
magma_ccopyvector_async( dr.num_rows, dr.dval, 1, dv.dval, 1, queues[1] );
// smoothing enabled
} else {
// smoothing operation
//---------------------------------------
// t = rs - r
// Q2
magma_cidr_smoothing_1( drs.num_rows, drs.num_cols, drs.dval, dr.dval, dtt.dval, queues[2] );
// t't
// t'rs
// Q2
CHECK( magma_cgemvmdot_shfl( dt.ld, 2, dtt.dval, dtt.dval, d1, d2, &dskp.dval[2], queues[2] ));
// skp[2-3] = dskp[2-3]
// Q2
magma_cgetvector( 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];
// xs = xs - gamma * (xs - x)
// Q0
magma_cidr_smoothing_2( dxs.num_rows, dxs.num_cols, -gamma, x->dval, dxs.dval, queues[0] );
// v = r
// Q1
magma_ccopyvector_async( dr.num_rows, dr.dval, 1, dv.dval, 1, queues[1] );
// rs = rs - gamma * (rs - r)
// Q2
magma_caxpy( drs.num_rows, -gamma, dtt.dval, 1, drs.dval, 1, queues[2] );
// |rs|
// Q2
nrmr = magma_scnrm2( drs.num_rows, drs.dval, 1, queues[2] );
// implicit sync Q2 --> |r|
//---------------------------------------
}
// 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;
}
}
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_ccopyvector_async( x->num_rows, dxs.dval, 1, x->dval, 1, queue );
// r = rs
magma_ccopyvector_async( dr.num_rows, drs.dval, 1, dr.dval, 1, queue );
}
// 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_cresidualvec( 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_cmfree( &dxs, queue );
magma_cmfree( &drs, queue );
magma_cmfree( &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_cmfree( &dr, queue );
magma_cmfree( &dP, queue );
magma_cmfree( &dP1, queue );
magma_cmfree( &dG, queue );
magma_cmfree( &dGcol, queue );
magma_cmfree( &dU, queue );
magma_cmfree( &dM, queue );
magma_cmfree( &df, queue );
magma_cmfree( &dt, queue );
magma_cmfree( &dc, queue );
magma_cmfree( &dv, queue );
magma_cmfree( &dlu, queue );
magma_cmfree( &dskp, queue );
magma_cmfree( &dalpha, queue );
magma_cmfree( &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_cpidr_strms */
}
magma_int_t
magma_cgmres( magma_c_sparse_matrix A, magma_c_vector b, magma_c_vector *x,
magma_c_solver_par *solver_par ){
// prepare solver feedback
solver_par->solver = Magma_GMRES;
solver_par->numiter = 0;
solver_par->info = 0;
// local variables
magmaFloatComplex c_zero = MAGMA_C_ZERO, c_one = MAGMA_C_ONE,
c_mone = MAGMA_C_NEG_ONE;
magma_int_t dofs = A.num_rows;
magma_int_t i, j, k, m = 0;
magma_int_t restart = min( dofs-1, solver_par->restart );
magma_int_t ldh = restart+1;
float nom, rNorm, RNorm, nom0, betanom, r0 = 0.;
// CPU workspace
magma_setdevice(0);
magmaFloatComplex *H, *HH, *y, *h1;
magma_cmalloc_pinned( &H, (ldh+1)*ldh );
magma_cmalloc_pinned( &y, ldh );
magma_cmalloc_pinned( &HH, ldh*ldh );
magma_cmalloc_pinned( &h1, ldh );
// GPU workspace
magma_c_vector r, q, q_t;
magma_c_vinit( &r, Magma_DEV, dofs, c_zero );
magma_c_vinit( &q, Magma_DEV, dofs*(ldh+1), c_zero );
q_t.memory_location = Magma_DEV;
q_t.val = NULL;
q_t.num_rows = q_t.nnz = dofs;
magmaFloatComplex *dy, *dH = NULL;
if (MAGMA_SUCCESS != magma_cmalloc( &dy, ldh ))
return MAGMA_ERR_DEVICE_ALLOC;
if (MAGMA_SUCCESS != magma_cmalloc( &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_cscal( dofs, c_zero, x->val, 1 ); // x = 0
magma_ccopy( dofs, b.val, 1, r.val, 1 ); // r = b
nom0 = betanom = magma_scnrm2( dofs, r.val, 1 ); // nom0= || r||
nom = nom0 * nom0;
solver_par->init_res = nom0;
H(1,0) = MAGMA_C_MAKE( nom0, 0. );
magma_csetvector(1, &H(1,0), 1, &dH(1,0), 1);
if ( (r0 = nom * solver_par->epsilon) < ATOLERANCE )
r0 = ATOLERANCE;
if ( nom < r0 )
return MAGMA_SUCCESS;
//Chronometry
real_Double_t tempo1, tempo2;
magma_device_sync(); tempo1=magma_wtime();
if( solver_par->verbose > 0 ){
solver_par->res_vec[0] = nom0;
solver_par->timing[0] = 0.0;
}
// start iteration
for( solver_par->numiter= 1; solver_par->numiter<solver_par->maxiter;
solver_par->numiter++ ){
magma_ccopy(dofs, r.val, 1, q(0), 1); // q[0] = 1.0/||r||
magma_cscal(dofs, 1./H(1,0), q(0), 1); // (to be fused)
for(k=1; k<=restart; k++) {
q_t.val = q(k-1);
magmablasSetKernelStream(stream[0]);
magma_c_spmv( c_one, A, q_t, c_zero, r );
// r = A q[k]
if (solver_par->ortho == Magma_MGS ) {
// modified Gram-Schmidt
magmablasSetKernelStream(stream[0]);
for (i=1; i<=k; i++) {
H(i,k) =magma_cdotc(dofs, q(i-1), 1, r.val, 1);
// H(i,k) = q[i] . r
magma_caxpy(dofs,-H(i,k), q(i-1), 1, r.val, 1);
// r = r - H(i,k) q[i]
}
H(k+1,k) = MAGMA_C_MAKE( magma_scnrm2(dofs, r.val, 1), 0. );
// H(k+1,k) = sqrt(r . r)
if (k < restart) {
magma_ccopy(dofs, r.val, 1, q(k), 1);
// q[k] = 1.0/H[k][k-1] r
magma_cscal(dofs, 1./H(k+1,k), q(k), 1);
// (to be fused)
}
} else if (solver_par->ortho == Magma_FUSED_CGS ) {
// fusing cgemv with scnrm2 in classical Gram-Schmidt
magmablasSetKernelStream(stream[0]);
magma_ccopy(dofs, r.val, 1, q(k), 1);
// dH(1:k+1,k) = q[0:k] . r
magmablas_cgemv(MagmaTrans, dofs, k+1, c_one, q(0),
dofs, r.val, 1, c_zero, &dH(1,k), 1);
// r = r - q[0:k-1] dH(1:k,k)
magmablas_cgemv(MagmaNoTrans, dofs, k, c_mone, q(0),
dofs, &dH(1,k), 1, c_one, r.val, 1);
// 1) dH(k+1,k) = sqrt( dH(k+1,k) - dH(1:k,k) )
magma_ccopyscale( 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_cgetvector_async(k+1, &dH(1,k), 1, &H(1,k), 1, stream[1]);
// asynch copy dH(1:(k+1),k) to H(1:(k+1),k)
} else {
// classical Gram-Schmidt (default)
// > explicitly calling magmabls
magmablasSetKernelStream(stream[0]);
magmablas_cgemv(MagmaTrans, dofs, k, c_one, q(0),
dofs, r.val, 1, c_zero, &dH(1,k), 1);
// dH(1:k,k) = q[0:k-1] . r
#ifndef SCNRM2SCALE
// start copying dH(1:k,k) to H(1:k,k)
magma_event_record( event[0], stream[0] );
magma_queue_wait_event( stream[1], event[0] );
magma_cgetvector_async(k, &dH(1,k), 1, &H(1,k),
1, stream[1]);
#endif
// r = r - q[0:k-1] dH(1:k,k)
magmablas_cgemv(MagmaNoTrans, dofs, k, c_mone, q(0),
dofs, &dH(1,k), 1, c_one, r.val, 1);
#ifdef SCNRM2SCALE
magma_ccopy(dofs, r.val, 1, q(k), 1);
// q[k] = r / H(k,k-1)
magma_scnrm2scale(dofs, q(k), dofs, &dH(k+1,k) );
// dH(k+1,k) = sqrt(r . r) and r = r / dH(k+1,k)
magma_event_record( event[0], stream[0] );
// start sending dH(1:k,k) to H(1:k,k)
magma_queue_wait_event( stream[1], event[0] );
// can we keep H(k+1,k) on GPU and combine?
magma_cgetvector_async(k+1, &dH(1,k), 1, &H(1,k), 1, stream[1]);
#else
H(k+1,k) = MAGMA_C_MAKE( magma_scnrm2(dofs, r.val, 1), 0. );
// H(k+1,k) = sqrt(r . r)
if( k<solver_par->restart ){
magmablasSetKernelStream(stream[0]);
magma_ccopy(dofs, r.val, 1, q(k), 1);
// q[k] = 1.0/H[k][k-1] r
magma_cscal(dofs, 1./H(k+1,k), q(k), 1);
// (to be fused)
}
#endif
}
}
magma_queue_sync( stream[1] );
for( k=1; k<=restart; k++ ){
/* Minimization of || b-Ax || in H_k */
for (i=1; i<=k; i++) {
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_cdotc_sub( i+1, &H(1,k), 1, &H(1,i), 1, &HH(k,i) );
#else
HH(k,i) = cblas_cdotc(i+1, &H(1,k), 1, &H(1,i), 1);
#endif
}
h1[k] = H(1,k)*H(1,0);
if (k != 1)
for (i=1; i<k; i++) {
for (m=i+1; m<k; m++){
HH(k,m) -= HH(k,i) * HH(m,i);
}
HH(k,k) -= HH(k,i) * HH(k,i) / HH(i,i);
HH(k,i) = HH(k,i)/HH(i,i);
h1[k] -= h1[i] * HH(k,i);
}
y[k] = h1[k]/HH(k,k);
if (k != 1)
for (i=k-1; i>=1; i--) {
y[i] = h1[i]/HH(i,i);
for (j=i+1; j<=k; j++)
y[i] -= y[j] * HH(j,i);
}
m = k;
rNorm = fabs(MAGMA_C_REAL(H(k+1,k)));
}
magma_csetmatrix_async(m, 1, y+1, m, dy, m, stream[0]);
magmablasSetKernelStream(stream[0]);
magma_cgemv(MagmaNoTrans, dofs, m, c_one, q(0), dofs, dy, 1,
c_one, x->val, 1);
magma_c_spmv( c_mone, A, *x, c_zero, r ); // r = - A * x
magma_caxpy(dofs, c_one, b.val, 1, r.val, 1); // r = r + b
H(1,0) = MAGMA_C_MAKE( magma_scnrm2(dofs, r.val, 1), 0. );
// RNorm = H[1][0] = || r ||
RNorm = MAGMA_C_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_cresidual( 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_c_vfree(&r);
magma_c_vfree(&q);
// 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_cgmres */
int main(int argc, char **argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, dev_perf, dev_time, cpu_perf, cpu_time;
float magma_error, dev_error, work[1];
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t M, N, Xm, Ym, lda, sizeA, sizeX, sizeY;
magma_int_t incx = 1;
magma_int_t incy = 1;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex alpha = MAGMA_C_MAKE( 1.5, -2.3 );
magmaFloatComplex beta = MAGMA_C_MAKE( -0.6, 0.8 );
magmaFloatComplex *A, *X, *Y, *Ydev, *Ymagma;
magmaFloatComplex_ptr dA, dX, dY;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
printf("trans = %s\n", lapack_trans_const(opts.transA) );
#ifdef HAVE_CUBLAS
printf(" M N MAGMA Gflop/s (ms) %s Gflop/s (ms) CPU Gflop/s (ms) MAGMA error %s error\n",
g_platform_str, g_platform_str );
#else
printf(" M N %s Gflop/s (ms) CPU Gflop/s (ms) %s error\n",
g_platform_str, g_platform_str );
#endif
printf("===================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
lda = ((M+31)/32)*32;
gflops = FLOPS_CGEMV( M, N ) / 1e9;
if ( opts.transA == MagmaNoTrans ) {
Xm = N;
Ym = M;
} else {
Xm = M;
Ym = N;
}
sizeA = lda*N;
sizeX = incx*Xm;
sizeY = incy*Ym;
TESTING_MALLOC_CPU( A, magmaFloatComplex, sizeA );
TESTING_MALLOC_CPU( X, magmaFloatComplex, sizeX );
TESTING_MALLOC_CPU( Y, magmaFloatComplex, sizeY );
TESTING_MALLOC_CPU( Ydev, magmaFloatComplex, sizeY );
TESTING_MALLOC_CPU( Ymagma, magmaFloatComplex, sizeY );
TESTING_MALLOC_DEV( dA, magmaFloatComplex, sizeA );
TESTING_MALLOC_DEV( dX, magmaFloatComplex, sizeX );
TESTING_MALLOC_DEV( dY, magmaFloatComplex, sizeY );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &sizeA, A );
lapackf77_clarnv( &ione, ISEED, &sizeX, X );
lapackf77_clarnv( &ione, ISEED, &sizeY, Y );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_csetmatrix( M, N, A, lda, dA, lda );
magma_csetvector( Xm, X, incx, dX, incx );
magma_csetvector( Ym, Y, incy, dY, incy );
#ifdef HAVE_CUBLAS
dev_time = magma_sync_wtime( 0 );
cublasCgemv( opts.handle, cublas_trans_const(opts.transA),
M, N, &alpha, dA, lda, dX, incx, &beta, dY, incy );
dev_time = magma_sync_wtime( 0 ) - dev_time;
#else
dev_time = magma_sync_wtime( opts.queue );
magma_cgemv( opts.transA, M, N,
&alpha, dA, lda,
dX, incx,
&beta, dY, incy );
dev_time = magma_sync_wtime( opts.queue ) - dev_time;
#endif
dev_perf = gflops / dev_time;
magma_cgetvector( Ym, dY, incy, Ydev, incy );
/* =====================================================================
Performs operation using MAGMABLAS (currently only with CUDA)
=================================================================== */
#ifdef HAVE_CUBLAS
magma_csetvector( Ym, Y, incy, dY, incy );
magma_time = magma_sync_wtime( 0 );
magmablas_cgemv( opts.transA, M, N, alpha, dA, lda, dX, incx, beta, dY, incy );
magma_time = magma_sync_wtime( 0 ) - magma_time;
magma_perf = gflops / magma_time;
magma_cgetvector( Ym, dY, incy, Ymagma, incy );
#endif
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
cpu_time = magma_wtime();
blasf77_cgemv( lapack_trans_const(opts.transA), &M, &N,
&alpha, A, &lda,
X, &incx,
&beta, Y, &incy );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/* =====================================================================
Check the result
=================================================================== */
float Anorm = lapackf77_clange( "F", &M, &N, A, &lda, work );
float Xnorm = lapackf77_clange( "F", &Xm, &ione, X, &Xm, work );
blasf77_caxpy( &Ym, &c_neg_one, Y, &incy, Ydev, &incy );
dev_error = lapackf77_clange( "F", &Ym, &ione, Ydev, &Ym, work ) / (Anorm * Xnorm);
#ifdef HAVE_CUBLAS
blasf77_caxpy( &Ym, &c_neg_one, Y, &incy, Ymagma, &incy );
magma_error = lapackf77_clange( "F", &Ym, &ione, Ymagma, &Ym, work ) / (Anorm * Xnorm);
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %s\n",
(int) M, (int) N,
magma_perf, 1000.*magma_time,
dev_perf, 1000.*dev_time,
cpu_perf, 1000.*cpu_time,
magma_error, dev_error,
(magma_error < tol && dev_error < tol ? "ok" : "failed"));
status += ! (magma_error < tol && dev_error < tol);
#else
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) M, (int) N,
dev_perf, 1000.*dev_time,
cpu_perf, 1000.*cpu_time,
dev_error,
(dev_error < tol ? "ok" : "failed"));
status += ! (dev_error < tol);
#endif
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( X );
TESTING_FREE_CPU( Y );
TESTING_FREE_CPU( Ydev );
TESTING_FREE_CPU( Ymagma );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dX );
TESTING_FREE_DEV( dY );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
extern "C" magma_int_t
magma_c_spmv(
magmaFloatComplex alpha,
magma_c_matrix A,
magma_c_matrix x,
magmaFloatComplex beta,
magma_c_matrix y,
magma_queue_t queue )
{
magma_int_t info = 0;
magma_c_matrix x2={Magma_CSR};
cusparseHandle_t cusparseHandle = 0;
cusparseMatDescr_t descr = 0;
// make sure RHS is a dense matrix
if ( x.storage_type != Magma_DENSE ) {
printf("error: only dense vectors are supported for SpMV.\n");
info = MAGMA_ERR_NOT_SUPPORTED;
goto cleanup;
}
if ( A.memory_location != x.memory_location ||
x.memory_location != y.memory_location ) {
printf("error: linear algebra objects are not located in same memory!\n");
printf("memory locations are: %d %d %d\n",
A.memory_location, x.memory_location, y.memory_location );
info = MAGMA_ERR_INVALID_PTR;
goto cleanup;
}
// DEV case
if ( A.memory_location == Magma_DEV ) {
if ( A.num_cols == x.num_rows && x.num_cols == 1 ) {
if ( A.storage_type == Magma_CSR ||
A.storage_type == Magma_CUCSR ||
A.storage_type == Magma_CSRL ||
A.storage_type == Magma_CSRU )
{
CHECK_CUSPARSE( cusparseCreate( &cusparseHandle ));
CHECK_CUSPARSE( cusparseSetStream( cusparseHandle, queue->cuda_stream() ));
CHECK_CUSPARSE( cusparseCreateMatDescr( &descr ));
CHECK_CUSPARSE( cusparseSetMatType( descr, CUSPARSE_MATRIX_TYPE_GENERAL ));
CHECK_CUSPARSE( cusparseSetMatIndexBase( descr, CUSPARSE_INDEX_BASE_ZERO ));
cusparseCcsrmv( cusparseHandle,CUSPARSE_OPERATION_NON_TRANSPOSE,
A.num_rows, A.num_cols, A.nnz, &alpha, descr,
A.dval, A.drow, A.dcol, x.dval, &beta, y.dval );
}
else if ( A.storage_type == Magma_ELL ) {
//printf("using ELLPACKT kernel for SpMV: ");
CHECK( magma_cgeelltmv( MagmaNoTrans, A.num_rows, A.num_cols,
A.max_nnz_row, alpha, A.dval, A.dcol, x.dval, beta,
y.dval, queue ));
//printf("done.\n");
}
else if ( A.storage_type == Magma_ELLPACKT ) {
//printf("using ELL kernel for SpMV: ");
CHECK( magma_cgeellmv( MagmaNoTrans, A.num_rows, A.num_cols,
A.max_nnz_row, alpha, A.dval, A.dcol, x.dval, beta,
y.dval, queue ));
//printf("done.\n");
}
else if ( A.storage_type == Magma_ELLRT ) {
//printf("using ELLRT kernel for SpMV: ");
CHECK( magma_cgeellrtmv( MagmaNoTrans, A.num_rows, A.num_cols,
A.max_nnz_row, alpha, A.dval, A.dcol, A.drow, x.dval,
beta, y.dval, A.alignment, A.blocksize, queue ));
//printf("done.\n");
}
else if ( A.storage_type == Magma_SELLP ) {
//printf("using SELLP kernel for SpMV: ");
CHECK( magma_cgesellpmv( MagmaNoTrans, A.num_rows, A.num_cols,
A.blocksize, A.numblocks, A.alignment,
alpha, A.dval, A.dcol, A.drow, x.dval, beta, y.dval, queue ));
//printf("done.\n");
}
else if ( A.storage_type == Magma_DENSE ) {
//printf("using DENSE kernel for SpMV: ");
magmablas_cgemv( MagmaNoTrans, A.num_rows, A.num_cols, alpha,
A.dval, A.num_rows, x.dval, 1, beta, y.dval,
1, queue );
//printf("done.\n");
}
else if ( A.storage_type == Magma_SPMVFUNCTION ) {
//printf("using DENSE kernel for SpMV: ");
CHECK( magma_ccustomspmv( alpha, x, beta, y, queue ));
//printf("done.\n");
}
else if ( A.storage_type == Magma_BCSR ) {
//printf("using CUSPARSE BCSR kernel for SpMV: ");
// CUSPARSE context //
cusparseDirection_t dirA = CUSPARSE_DIRECTION_ROW;
int mb = magma_ceildiv( A.num_rows, A.blocksize );
int nb = magma_ceildiv( A.num_cols, A.blocksize );
CHECK_CUSPARSE( cusparseCreate( &cusparseHandle ));
CHECK_CUSPARSE( cusparseSetStream( cusparseHandle, queue->cuda_stream() ));
CHECK_CUSPARSE( cusparseCreateMatDescr( &descr ));
cusparseCbsrmv( cusparseHandle, dirA,
CUSPARSE_OPERATION_NON_TRANSPOSE, mb, nb, A.numblocks,
&alpha, descr, A.dval, A.drow, A.dcol, A.blocksize, x.dval,
&beta, y.dval );
}
else {
printf("error: format not supported.\n");
info = MAGMA_ERR_NOT_SUPPORTED;
}
}
else if ( A.num_cols < x.num_rows || x.num_cols > 1 ) {
magma_int_t num_vecs = x.num_rows / A.num_cols * x.num_cols;
if ( A.storage_type == Magma_CSR ) {
CHECK_CUSPARSE( cusparseCreate( &cusparseHandle ));
CHECK_CUSPARSE( cusparseSetStream( cusparseHandle, queue->cuda_stream() ));
CHECK_CUSPARSE( cusparseCreateMatDescr( &descr ));
CHECK_CUSPARSE( cusparseSetMatType( descr, CUSPARSE_MATRIX_TYPE_GENERAL ));
CHECK_CUSPARSE( cusparseSetMatIndexBase( descr, CUSPARSE_INDEX_BASE_ZERO ));
if ( x.major == MagmaColMajor) {
cusparseCcsrmm(cusparseHandle,
CUSPARSE_OPERATION_NON_TRANSPOSE,
A.num_rows, num_vecs, A.num_cols, A.nnz,
&alpha, descr, A.dval, A.drow, A.dcol,
x.dval, A.num_cols, &beta, y.dval, A.num_cols);
} else if ( x.major == MagmaRowMajor) {
/*cusparseCcsrmm2(cusparseHandle,
CUSPARSE_OPERATION_NON_TRANSPOSE,
CUSPARSE_OPERATION_TRANSPOSE,
A.num_rows, num_vecs, A.num_cols, A.nnz,
&alpha, descr, A.dval, A.drow, A.dcol,
x.dval, A.num_cols, &beta, y.dval, A.num_cols);
*/
}
} else if ( A.storage_type == Magma_SELLP ) {
if ( x.major == MagmaRowMajor) {
CHECK( magma_cmgesellpmv( MagmaNoTrans, A.num_rows, A.num_cols,
num_vecs, A.blocksize, A.numblocks, A.alignment,
alpha, A.dval, A.dcol, A.drow, x.dval, beta, y.dval, queue ));
}
else if ( x.major == MagmaColMajor) {
// transpose first to row major
CHECK( magma_cvtranspose( x, &x2, queue ));
CHECK( magma_cmgesellpmv( MagmaNoTrans, A.num_rows, A.num_cols,
num_vecs, A.blocksize, A.numblocks, A.alignment,
alpha, A.dval, A.dcol, A.drow, x2.dval, beta, y.dval, queue ));
}
}
/*if ( A.storage_type == Magma_DENSE ) {
//printf("using DENSE kernel for SpMV: ");
magmablas_cmgemv( MagmaNoTrans, A.num_rows, A.num_cols,
num_vecs, alpha, A.dval, A.num_rows, x.dval, 1,
beta, y.dval, 1 );
//printf("done.\n");
}*/
else {
printf("error: format not supported.\n");
info = MAGMA_ERR_NOT_SUPPORTED;
}
}
}
// CPU case missing!
else {
printf("error: CPU not yet supported.\n");
info = MAGMA_ERR_NOT_SUPPORTED;
}
cleanup:
cusparseDestroyMatDescr( descr );
cusparseDestroy( cusparseHandle );
cusparseHandle = 0;
descr = 0;
magma_cmfree(&x2, queue );
return info;
}
