void magma_sprint(
magma_int_t m, magma_int_t n,
const float *A, magma_int_t lda )
{
#define A(i,j) (A + (i) + (j)*lda)
magma_int_t info = 0;
if ( m < 0 )
info = -1;
else if ( n < 0 )
info = -2;
else if ( lda < max(1,m) )
info = -4;
if (info != 0) {
magma_xerbla( __func__, -(info) );
return; //info;
}
float c_zero = MAGMA_S_ZERO;
if ( m == 1 ) {
printf( "[ " );
}
else {
printf( "[\n" );
}
for( int i = 0; i < m; ++i ) {
for( int j = 0; j < n; ++j ) {
if ( MAGMA_S_EQUAL( *A(i,j), c_zero )) {
#ifdef COMPLEX
printf( " 0. " );
#else
printf( " 0. " );
#endif
}
else {
#ifdef COMPLEX
printf( " %8.4f+%8.4fi", MAGMA_S_REAL( *A(i,j) ), MAGMA_S_IMAG( *A(i,j) ));
#else
printf( " %8.4f", MAGMA_S_REAL( *A(i,j) ));
#endif
}
}
if ( m > 1 ) {
printf( "\n" );
}
else {
printf( " " );
}
}
printf( "];\n" );
}
/** @return true if either real(x) or imag(x) is INF. */
inline bool magma_s_isinf( float x )
{
#ifdef COMPLEX
return isinf( MAGMA_S_REAL( x )) ||
isinf( MAGMA_S_IMAG( x ));
#else
return isinf( x );
#endif
}
void magma_smake_symmetric( magma_int_t N, float* A, magma_int_t lda )
{
magma_int_t i, j;
for( i=0; i < N; ++i ) {
A(i,i) = MAGMA_S_MAKE( MAGMA_S_REAL( A(i,i) ), 0. );
for( j=0; j < i; ++j ) {
A(j,i) = MAGMA_S_CONJ( A(i,j) );
}
}
}
void magma_smake_hpd( magma_int_t N, float* A, magma_int_t lda )
{
magma_int_t i, j;
for( i=0; i<N; ++i ) {
A(i,i) = MAGMA_S_MAKE( MAGMA_S_REAL( A(i,i) ) + N, 0. );
for( j=0; j<i; ++j ) {
A(j,i) = MAGMA_S_CNJG( A(i,j) );
}
}
}
void init_matrix( magma_int_t N, float *h_A, magma_int_t lda )
{
magma_int_t ione = 1, n2 = N*lda;
magma_int_t ISEED[4] = {0,0,0,1};
lapackf77_slarnv( &ione, ISEED, &n2, h_A );
/* Symmetrize and increase the diagonal */
for (magma_int_t i = 0; i < N; ++i) {
h_A(i,i) = MAGMA_S_MAKE( MAGMA_S_REAL(h_A(i,i)) + N, 0 );
for (magma_int_t j = 0; j < i; ++j)
h_A(i, j) = MAGMA_S_CNJG( h_A(j, i) );
}
}
// Initialize matrix to random.
// Having this in separate function ensures the same ISEED is always used,
// so we can re-generate the identical matrix.
void init_matrix(
magma_opts &opts,
magma_int_t m, magma_int_t n,
float *A, magma_int_t lda )
{
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t n2 = lda*n;
lapackf77_slarnv( &ione, ISEED, &n2, A );
if ( opts.version == 2 ) {
for (magma_int_t i=0; i < min(m,n); ++i ) {
A[ i + i*lda ] = MAGMA_S_MAKE( MAGMA_S_REAL( A[ i + i*lda ] ) + max(m,n), 0 );
}
}
}
void magma_sprint( int m, int n, float *A, int lda )
{
float c_zero = MAGMA_S_ZERO;
printf( "[\n" );
for( int i = 0; i < m; ++i ) {
for( int j = 0; j < n; ++j ) {
if ( MAGMA_S_EQUAL( *A(i,j), c_zero )) {
printf( " 0. " );
}
else {
printf( " %8.4f", MAGMA_S_REAL( *A(i,j) ));
}
}
printf( "\n" );
}
printf( "];\n" );
}
/* ////////////////////////////////////////////////////////////////////////////
-- 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;
}
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_ssytrd_sb2st(magma_int_t threads, char uplo, magma_int_t n, magma_int_t nb, magma_int_t Vblksiz,
float *A, magma_int_t lda, float *D, float *E,
float *V, magma_int_t ldv, float *TAU, magma_int_t compT, float *T, magma_int_t ldt)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
Arguments
=========
THREADS (input) INTEGER
Specifies the number of pthreads used.
THREADS > 0
UPLO (input) CHARACTER*1
= 'U': Upper triangles of A is stored;
= 'L': Lower triangles of A is stored.
N (input) INTEGER
The order of the matrix A. N >= 0.
NB (input) INTEGER
The order of the band matrix A. N >= NB >= 0.
VBLKSIZ (input) INTEGER
The size of the block of householder vectors applied at once.
A (input/workspace) REAL array, dimension (LDA, N)
On entry the band matrix stored in the following way:
LDA (input) INTEGER
The leading dimension of the array A. LDA >= 2*NB.
D (output) DOUBLE array, dimension (N)
The diagonal elements of the tridiagonal matrix T:
D(i) = A(i,i).
E (output) DOUBLE array, dimension (N-1)
The off-diagonal elements of the tridiagonal matrix T:
E(i) = A(i,i+1) if UPLO = 'U', E(i) = A(i+1,i) if UPLO = 'L'.
V (output) REAL array, dimension (BLKCNT, LDV, VBLKSIZ)
On exit it contains the blocks of householder reflectors
BLKCNT is the number of block and it is returned by the funtion MAGMA_BULGE_GET_BLKCNT.
LDV (input) INTEGER
The leading dimension of V.
LDV > NB + VBLKSIZ + 1
TAU (output) REAL dimension(BLKCNT, VBLKSIZ)
???
COMPT (input) INTEGER
if COMPT = 0 T is not computed
if COMPT = 1 T is computed
T (output) REAL dimension(LDT *)
if COMPT = 1 on exit contains the matrices T needed for Q2
if COMPT = 0 T is not referenced
LDT (input) INTEGER
The leading dimension of T.
LDT > Vblksiz
INFO (output) INTEGER ????????????????????????????????????????????????????????????????????????????????????
= 0: successful exit
===================================================================== */
#ifdef ENABLE_TIMER
real_Double_t timeblg=0.0;
#endif
//char uplo_[2] = {uplo, 0};
magma_int_t mklth = threads;
magma_int_t INgrsiz=1;
magma_int_t blkcnt = magma_bulge_get_blkcnt(n, nb, Vblksiz);
magma_int_t nbtiles = magma_ceildiv(n, nb);
memset(T, 0, blkcnt*ldt*Vblksiz*sizeof(float));
memset(TAU, 0, blkcnt*Vblksiz*sizeof(float));
memset(V, 0, blkcnt*ldv*Vblksiz*sizeof(float));
magma_int_t* prog;
magma_malloc_cpu((void**) &prog, (2*nbtiles+threads+10)*sizeof(magma_int_t));
memset(prog, 0, (2*nbtiles+threads+10)*sizeof(magma_int_t));
magma_sbulge_id_data* arg;
magma_malloc_cpu((void**) &arg, threads*sizeof(magma_sbulge_id_data));
pthread_t* thread_id;
magma_malloc_cpu((void**) &thread_id, threads*sizeof(pthread_t));
pthread_attr_t thread_attr;
magma_setlapack_numthreads(1);
magma_sbulge_data data_bulge(threads, n, nb, nbtiles, INgrsiz, Vblksiz, compT,
A, lda, V, ldv, TAU, T, ldt, prog);
// Set one thread per core
pthread_attr_init(&thread_attr);
pthread_attr_setscope(&thread_attr, PTHREAD_SCOPE_SYSTEM);
pthread_setconcurrency(threads);
//timing
#ifdef ENABLE_TIMER
timeblg = magma_wtime();
#endif
// Launch threads
for (magma_int_t thread = 1; thread < threads; thread++)
{
arg[thread] = magma_sbulge_id_data(thread, &data_bulge);
pthread_create(&thread_id[thread], &thread_attr, magma_ssytrd_sb2st_parallel_section, &arg[thread]);
}
arg[0] = magma_sbulge_id_data(0, &data_bulge);
magma_ssytrd_sb2st_parallel_section(&arg[0]);
// Wait for completion
for (magma_int_t thread = 1; thread < threads; thread++)
{
void *exitcodep;
pthread_join(thread_id[thread], &exitcodep);
}
// timing
#ifdef ENABLE_TIMER
timeblg = magma_wtime()-timeblg;
printf(" time BULGE+T = %f \n" ,timeblg);
#endif
magma_free_cpu(thread_id);
magma_free_cpu(arg);
magma_free_cpu(prog);
magma_setlapack_numthreads(mklth);
/*================================================
* store resulting diag and lower diag D and E
* note that D and E are always real
*================================================*/
/* Make diagonal and superdiagonal elements real,
* storing them in D and E
*/
/* In real case, the off diagonal element are
* not necessary real. we have to make off-diagonal
* elements real and copy them to E.
* When using HouseHolder elimination,
* the SLARFG give us a real as output so, all the
* diagonal/off-diagonal element except the last one are already
* real and thus we need only to take the abs of the last
* one.
* */
#if defined(PRECISION_z) || defined(PRECISION_c)
if(uplo==MagmaLower){
for (magma_int_t i=0; i < n-1 ; i++)
{
D[i] = MAGMA_S_REAL(A[i*lda ]);
E[i] = MAGMA_S_REAL(A[i*lda+1]);
}
D[n-1] = MAGMA_S_REAL(A[(n-1)*lda]);
} else { /* MagmaUpper not tested yet */
for (magma_int_t i=0; i<n-1; i++)
{
D[i] = MAGMA_S_REAL(A[i*lda+nb]);
E[i] = MAGMA_S_REAL(A[i*lda+nb-1]);
}
D[n-1] = MAGMA_S_REAL(A[(n-1)*lda+nb]);
} /* end MagmaUpper */
#else
if( uplo == MagmaLower ){
for (magma_int_t i=0; i < n-1; i++) {
D[i] = A[i*lda]; // diag
E[i] = A[i*lda+1]; //lower diag
}
D[n-1] = A[(n-1)*lda];
} else {
for (magma_int_t i=0; i < n-1; i++) {
D[i] = A[i*lda+nb]; // diag
E[i] = A[i*lda+nb-1]; //lower diag
}
D[n-1] = A[(n-1)*lda+nb];
}
#endif
return MAGMA_SUCCESS;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing sgeqrf
*/
int main( int argc, char** argv)
{
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
float error, work[1];
float c_neg_one = MAGMA_S_NEG_ONE;
float *h_A, *h_R, *tau, *h_work, tmp[1];
magma_int_t M, N, n2, lda, lwork, info, min_mn, nb;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1}, ISEED2[4];
/* Initialize */
magma_queue_t queue[2];
magma_device_t devices[MagmaMaxGPUs];
int num = 0;
magma_err_t err;
magma_init();
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_int_t status = 0;
float tol, eps = lapackf77_slamch("E");
tol = opts.tolerance * eps;
opts.lapack |= ( opts.check == 2 ); // check (-c2) implies lapack (-l)
err = magma_get_devices( devices, MagmaMaxGPUs, &num );
if ( err != 0 || num < 1 ) {
fprintf( stderr, "magma_get_devices failed: %d\n", err );
exit(-1);
}
// Create two queues on device opts.device
err = magma_queue_create( devices[opts.device], &queue[0] );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", err );
exit(-1);
}
err = magma_queue_create( devices[opts.device], &queue[1] );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", err );
exit(-1);
}
printf("ngpu %d\n", (int) opts.ngpu );
if ( opts.check == 1 ) {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R-Q'A||_1 / (M*||A||_1) ||I-Q'Q||_1 / M\n");
printf("===============================================================================================\n");
} else {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / ||A||_F\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];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
nb = magma_get_sgeqrf_nb(M);
gflops = FLOPS_SGEQRF( M, N ) / 1e9;
lwork = -1;
lapackf77_sgeqrf(&M, &N, h_A, &M, tau, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_S_REAL( tmp[0] );
lwork = max( lwork, max( N*nb, 2*nb*nb ));
TESTING_MALLOC_CPU( tau, float, min_mn );
TESTING_MALLOC_CPU( h_A, float, n2 );
TESTING_MALLOC_PIN( h_R, float, n2 );
TESTING_MALLOC_CPU( h_work, float, lwork );
/* Initialize the matrix */
for ( int j=0; j<4; j++ ) ISEED2[j] = ISEED[j]; // saving seeds
lapackf77_slarnv( &ione, ISEED, &n2, h_A );
lapackf77_slacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_sgeqrf(M, N, h_R, lda, tau, h_work, lwork, &info, queue);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_sgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
float *tau;
TESTING_MALLOC_CPU( tau, float, min_mn );
cpu_time = magma_wtime();
lapackf77_sgeqrf(&M, &N, h_A, &lda, tau, h_work, &lwork, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_sgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
TESTING_FREE_CPU( tau );
}
if ( opts.check == 1 ) {
/* =====================================================================
Check the result
=================================================================== */
magma_int_t lwork = n2+N;
float *h_W1, *h_W2, *h_W3;
float *h_RW, results[2];
TESTING_MALLOC_CPU( h_W1, float, n2 ); // Q
TESTING_MALLOC_CPU( h_W2, float, n2 ); // R
TESTING_MALLOC_CPU( h_W3, float, lwork ); // WORK
TESTING_MALLOC_CPU( h_RW, float, M ); // RWORK
lapackf77_slarnv( &ione, ISEED2, &n2, h_A );
lapackf77_sqrt02( &M, &N, &min_mn, h_A, h_R, h_W1, h_W2, &lda, tau, h_W3, &lwork,
h_RW, results );
results[0] *= eps;
results[1] *= eps;
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time, results[0],results[1] );
printf("%s\n", (results[0] < tol ? "" : " failed"));
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) %8.2e %8.2e",
(int) M, (int) N, gpu_perf, gpu_time, results[0],results[1] );
printf("%s\n", (results[0] < tol ? "" : " failed"));
}
status |= ! (results[0] < tol);
TESTING_FREE_CPU( h_W1 );
TESTING_FREE_CPU( h_W2 );
TESTING_FREE_CPU( h_W3 );
TESTING_FREE_CPU( h_RW );
} else if ( opts.check == 2 ) {
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
error = lapackf77_slange("f", &M, &N, h_A, &lda, work);
blasf77_saxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_slange("f", &M, &N, h_R, &lda, work) / error;
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time, error );
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) %8.2e",
(int) M, (int) N, gpu_perf, gpu_time, error );
}
printf("%s\n", (error < tol ? "" : " failed"));
status |= ! (error < tol);
}
else {
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) ---\n",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time );
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) M, (int) N, gpu_perf, gpu_time);
}
}
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_work );
TESTING_FREE_PIN( h_R );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
magma_queue_destroy( queue[0] );
magma_queue_destroy( queue[1] );
magma_finalize();
return status;
}
int main( int argc, char** argv)
{
real_Double_t gflops, gpu_perf, cpu_perf, gpu_time, cpu_time;
float *h_A, *h_R;
magmaFloat_ptr d_lA[MagmaMaxGPUs];
magma_int_t N = 0, n2, lda, ldda;
magma_int_t size[10] =
{ 1000, 2000, 3000, 4000, 5000, 6000, 7000, 8000, 9000, 10000 };
magma_int_t i, j, k, info;
float mz_one = MAGMA_S_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
float work[1], matnorm, diffnorm;
magma_int_t num_gpus0 = 1, num_gpus, flag = 0;
int nb, mb, n_local, nk;
magma_uplo_t uplo = MagmaLower;
if (argc != 1){
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0){
N = atoi(argv[++i]);
if (N>0) {
size[0] = size[9] = N;
flag = 1;
}else exit(1);
}
if(strcmp("-NGPU", argv[i])==0)
num_gpus0 = atoi(argv[++i]);
if(strcmp("-UPLO", argv[i])==0){
if(strcmp("L", argv[++i])==0){
uplo = MagmaLower;
}else{
uplo = MagmaUpper;
}
}
}
}
else {
printf("\nUsage: \n");
printf(" testing_spotrf_mgpu -N %d -NGPU %d -UPLO -L\n\n", 1024, num_gpus0);
}
/* looking for max. ldda */
ldda = 0;
n2 = 0;
for(i=0;i<10;i++){
N = size[i];
nb = magma_get_spotrf_nb(N);
mb = nb;
if(num_gpus0 > N/nb){
num_gpus = N/nb;
if(N%nb != 0) num_gpus ++;
}else{
num_gpus = num_gpus0;
}
n_local = nb*(1+N/(nb*num_gpus))*mb*((N+mb-1)/mb);
if(n_local > ldda) ldda = n_local;
if(n2 < N*N) n2 = N*N;
if(flag != 0) break;
}
/* Allocate host memory for the matrix */
TESTING_MALLOC_PIN( h_A, float, n2 );
TESTING_MALLOC_PIN( h_R, float, n2 );
/* Initialize */
magma_queue_t queues[MagmaMaxGPUs * 2];
//magma_queue_t queues[MagmaMaxGPUs];
magma_device_t devices[ MagmaMaxGPUs ];
int num = 0;
magma_err_t err;
magma_init();
err = magma_get_devices( devices, MagmaMaxGPUs, &num );
if ( err != 0 || num < 1 ) {
fprintf( stderr, "magma_get_devices failed: %d\n", err );
exit(-1);
}
for(i=0;i<num_gpus;i++){
err = magma_queue_create( devices[i], &queues[2*i] );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", err );
exit(-1);
}
err = magma_queue_create( devices[i], &queues[2*i+1] );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", err );
exit(-1);
}
}
printf("each buffer size: %d\n", ldda);
/* allocate local matrix on Buffers */
for(i=0; i<num_gpus0; i++){
TESTING_MALLOC_DEV( d_lA[i], float, ldda );
}
printf("\n\n");
printf("Using GPUs: %d\n", num_gpus0);
if(uplo == MagmaUpper){
printf("\n testing_spotrf_mgpu -N %d -NGPU %d -UPLO U\n\n", N, num_gpus0);
}else{
printf("\n testing_spotrf_mgpu -N %d -NGPU %d -UPLO L\n\n", N, num_gpus0);
}
printf(" N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R_magma-R_lapack||_F / ||R_lapack||_F\n");
printf("========================================================================================\n");
for(i=0; i<10; i++){
N = size[i];
lda = N;
n2 = lda*N;
ldda = ((N+31)/32)*32;
gflops = FLOPS( (float)N ) * 1e-9;
/* Initialize the matrix */
lapackf77_slarnv( &ione, ISEED, &n2, h_A );
/* Symmetrize and increase the diagonal */
for( int i = 0; i < N; ++i ) {
MAGMA_S_SET2REAL( h_A(i,i), MAGMA_S_REAL(h_A(i,i)) + N );
for( int j = 0; j < i; ++j ) {
h_A(i, j) = MAGMA_S_CNJG( h_A(j,i) );
}
}
lapackf77_slacpy( MagmaFullStr, &N, &N, h_A, &lda, h_R, &lda );
/* Warm up to measure the performance */
nb = magma_get_spotrf_nb(N);
if(num_gpus0 > N/nb){
num_gpus = N/nb;
if(N%nb != 0) num_gpus ++;
printf("too many GPUs for the matrix size, using %d GPUs\n", (int)num_gpus);
}else{
num_gpus = num_gpus0;
}
/* distribute matrix to gpus */
if(uplo == MagmaUpper){
// Upper
ldda = ((N+mb-1)/mb)*mb;
for(j=0;j<N;j+=nb){
k = (j/nb)%num_gpus;
nk = min(nb, N-j);
magma_ssetmatrix(N, nk,
&h_A[j*lda], 0, lda,
d_lA[k], j/(nb*num_gpus)*nb*ldda, ldda,
queues[2*k]);
}
}else{
// Lower
ldda = (1+N/(nb*num_gpus))*nb;
for(j=0;j<N;j+=nb){
k = (j/nb)%num_gpus;
nk = min(nb, N-j);
magma_ssetmatrix(nk, N, &h_A[j], 0, lda,
d_lA[k], (j/(nb*num_gpus)*nb), ldda,
queues[2*k]);
}
}
magma_spotrf_mgpu( num_gpus, uplo, N, d_lA, 0, ldda, &info, queues );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
/* distribute matrix to gpus */
if(uplo == MagmaUpper){
// Upper
ldda = ((N+mb-1)/mb)*mb;
for(j=0;j<N;j+=nb){
k = (j/nb)%num_gpus;
nk = min(nb, N-j);
magma_ssetmatrix(N, nk,
&h_A[j*lda], 0, lda,
d_lA[k], j/(nb*num_gpus)*nb*ldda, ldda,
queues[2*k]);
}
}else{
// Lower
ldda = (1+N/(nb*num_gpus))*nb;
for(j=0;j<N;j+=nb){
k = (j/nb)%num_gpus;
nk = min(nb, N-j);
magma_ssetmatrix(nk, N, &h_A[j], 0, lda,
d_lA[k], (j/(nb*num_gpus)*nb), ldda,
queues[2*k]);
}
}
gpu_time = magma_wtime();
magma_spotrf_mgpu( num_gpus, uplo, N, d_lA, 0, ldda, &info, queues );
gpu_time = magma_wtime() - gpu_time;
if (info != 0)
printf( "magma_spotrf had error %d.\n", info );
gpu_perf = gflops / gpu_time;
/* gather matrix from gpus */
if(uplo==MagmaUpper){
// Upper
for(j=0;j<N;j+=nb){
k = (j/nb)%num_gpus;
nk = min(nb, N-j);
magma_sgetmatrix(N, nk,
d_lA[k], j/(nb*num_gpus)*nb*ldda, ldda,
&h_R[j*lda], 0, lda, queues[2*k]);
}
}else{
// Lower
for(j=0; j<N; j+=nb){
k = (j/nb)%num_gpus;
nk = min(nb, N-j);
magma_sgetmatrix( nk, N,
d_lA[k], (j/(nb*num_gpus)*nb), ldda,
&h_R[j], 0, lda, queues[2*k] );
}
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
if(uplo == MagmaLower){
lapackf77_spotrf( MagmaLowerStr, &N, h_A, &lda, &info );
}else{
lapackf77_spotrf( MagmaUpperStr, &N, h_A, &lda, &info );
}
cpu_time = magma_wtime() - cpu_time;
if (info != 0)
printf( "lapackf77_spotrf had error %d.\n", info );
cpu_perf = gflops / cpu_time;
/* =====================================================================
Check the result compared to LAPACK
|R_magma - R_lapack| / |R_lapack|
=================================================================== */
matnorm = lapackf77_slange("f", &N, &N, h_A, &lda, work);
blasf77_saxpy(&n2, &mz_one, h_A, &ione, h_R, &ione);
diffnorm = lapackf77_slange("f", &N, &N, h_R, &lda, work);
printf( "%5d %6.2f (%6.2f) %6.2f (%6.2f) %e\n",
N, cpu_perf, cpu_time, gpu_perf, gpu_time, diffnorm / matnorm );
if (flag != 0)
break;
}
/* clean up */
TESTING_FREE_PIN( h_A );
TESTING_FREE_PIN( h_R );
for(i=0;i<num_gpus;i++){
TESTING_FREE_DEV( d_lA[i] );
magma_queue_destroy( queues[2*i] );
magma_queue_destroy( queues[2*i+1] );
}
magma_finalize();
}
/* ////////////////////////////////////////////////////////////////////////////
-- testing sparse matrix vector product
*/
int main( int argc, char** argv )
{
magma_int_t info = 0;
TESTING_CHECK( magma_init() );
magma_print_environment();
magma_queue_t queue=NULL;
magma_queue_create( 0, &queue );
magma_s_matrix hA={Magma_CSR}, hA_SELLP={Magma_CSR},
dA={Magma_CSR}, dA_SELLP={Magma_CSR};
magma_s_matrix hx={Magma_CSR}, hy={Magma_CSR}, dx={Magma_CSR},
dy={Magma_CSR}, hrefvec={Magma_CSR}, hcheck={Magma_CSR};
hA_SELLP.blocksize = 8;
hA_SELLP.alignment = 8;
real_Double_t start, end, res;
#ifdef MAGMA_WITH_MKL
magma_int_t *pntre=NULL;
#endif
cusparseHandle_t cusparseHandle = NULL;
cusparseMatDescr_t descr = NULL;
float c_one = MAGMA_S_MAKE(1.0, 0.0);
float c_zero = MAGMA_S_MAKE(0.0, 0.0);
float accuracy = 1e-10;
#define PRECISION_s
#if defined(PRECISION_c)
accuracy = 1e-4;
#endif
#if defined(PRECISION_s)
accuracy = 1e-4;
#endif
magma_int_t i, j;
for( i = 1; i < argc; ++i ) {
if ( strcmp("--blocksize", argv[i]) == 0 ) {
hA_SELLP.blocksize = atoi( argv[++i] );
} else if ( strcmp("--alignment", argv[i]) == 0 ) {
hA_SELLP.alignment = atoi( argv[++i] );
} else
break;
}
printf("\n# usage: ./run_sspmm"
" [ --blocksize %lld --alignment %lld (for SELLP) ] matrices\n\n",
(long long) hA_SELLP.blocksize, (long long) hA_SELLP.alignment );
while( i < argc ) {
if ( strcmp("LAPLACE2D", argv[i]) == 0 && i+1 < argc ) { // Laplace test
i++;
magma_int_t laplace_size = atoi( argv[i] );
TESTING_CHECK( magma_sm_5stencil( laplace_size, &hA, queue ));
} else { // file-matrix test
TESTING_CHECK( magma_s_csr_mtx( &hA, argv[i], queue ));
}
printf("%% matrix info: %lld-by-%lld with %lld nonzeros\n",
(long long) hA.num_rows, (long long) hA.num_cols, (long long) hA.nnz );
real_Double_t FLOPS = 2.0*hA.nnz/1e9;
// m - number of rows for the sparse matrix
// n - number of vectors to be multiplied in the SpMM product
magma_int_t m, n;
m = hA.num_rows;
n = 48;
// init CPU vectors
TESTING_CHECK( magma_svinit( &hx, Magma_CPU, m, n, c_one, queue ));
TESTING_CHECK( magma_svinit( &hy, Magma_CPU, m, n, c_zero, queue ));
// init DEV vectors
TESTING_CHECK( magma_svinit( &dx, Magma_DEV, m, n, c_one, queue ));
TESTING_CHECK( magma_svinit( &dy, Magma_DEV, m, n, c_zero, queue ));
// calling MKL with CSR
#ifdef MAGMA_WITH_MKL
TESTING_CHECK( magma_imalloc_cpu( &pntre, m + 1 ) );
pntre[0] = 0;
for (j=0; j < m; j++ ) {
pntre[j] = hA.row[j+1];
}
MKL_INT num_rows = hA.num_rows;
MKL_INT num_cols = hA.num_cols;
MKL_INT nnz = hA.nnz;
MKL_INT num_vecs = n;
MKL_INT *col;
TESTING_CHECK( magma_malloc_cpu( (void**) &col, nnz * sizeof(MKL_INT) ));
for( magma_int_t t=0; t < hA.nnz; ++t ) {
col[ t ] = hA.col[ t ];
}
MKL_INT *row;
TESTING_CHECK( magma_malloc_cpu( (void**) &row, num_rows * sizeof(MKL_INT) ));
for( magma_int_t t=0; t < hA.num_rows; ++t ) {
row[ t ] = hA.col[ t ];
}
// === Call MKL with consecutive SpMVs, using mkl_scsrmv ===
// warmp up
mkl_scsrmv( "N", &num_rows, &num_cols,
MKL_ADDR(&c_one), "GFNC", MKL_ADDR(hA.val), col, row, pntre,
MKL_ADDR(hx.val),
MKL_ADDR(&c_zero), MKL_ADDR(hy.val) );
start = magma_wtime();
for (j=0; j < 10; j++ ) {
mkl_scsrmv( "N", &num_rows, &num_cols,
MKL_ADDR(&c_one), "GFNC", MKL_ADDR(hA.val), col, row, pntre,
MKL_ADDR(hx.val),
MKL_ADDR(&c_zero), MKL_ADDR(hy.val) );
}
end = magma_wtime();
printf( "\n > MKL SpMVs : %.2e seconds %.2e GFLOP/s (CSR).\n",
(end-start)/10, FLOPS*10/(end-start) );
// === Call MKL with blocked SpMVs, using mkl_scsrmm ===
char transa = 'n';
MKL_INT ldb = n, ldc=n;
char matdescra[6] = {'g', 'l', 'n', 'c', 'x', 'x'};
// warm up
mkl_scsrmm( &transa, &num_rows, &num_vecs, &num_cols, MKL_ADDR(&c_one), matdescra,
MKL_ADDR(hA.val), col, row, pntre,
MKL_ADDR(hx.val), &ldb,
MKL_ADDR(&c_zero),
MKL_ADDR(hy.val), &ldc );
start = magma_wtime();
for (j=0; j < 10; j++ ) {
mkl_scsrmm( &transa, &num_rows, &num_vecs, &num_cols, MKL_ADDR(&c_one), matdescra,
MKL_ADDR(hA.val), col, row, pntre,
MKL_ADDR(hx.val), &ldb,
MKL_ADDR(&c_zero),
MKL_ADDR(hy.val), &ldc );
}
end = magma_wtime();
printf( "\n > MKL SpMM : %.2e seconds %.2e GFLOP/s (CSR).\n",
(end-start)/10, FLOPS*10.*n/(end-start) );
magma_free_cpu( row );
magma_free_cpu( col );
row = NULL;
col = NULL;
#endif // MAGMA_WITH_MKL
// copy matrix to GPU
TESTING_CHECK( magma_smtransfer( hA, &dA, Magma_CPU, Magma_DEV, queue ));
// SpMV on GPU (CSR)
start = magma_sync_wtime( queue );
for (j=0; j < 10; j++) {
TESTING_CHECK( magma_s_spmv( c_one, dA, dx, c_zero, dy, queue ));
}
end = magma_sync_wtime( queue );
printf( " > MAGMA: %.2e seconds %.2e GFLOP/s (standard CSR).\n",
(end-start)/10, FLOPS*10.*n/(end-start) );
TESTING_CHECK( magma_smtransfer( dy, &hrefvec , Magma_DEV, Magma_CPU, queue ));
magma_smfree(&dA, queue );
// convert to SELLP and copy to GPU
TESTING_CHECK( magma_smconvert( hA, &hA_SELLP, Magma_CSR, Magma_SELLP, queue ));
TESTING_CHECK( magma_smtransfer( hA_SELLP, &dA_SELLP, Magma_CPU, Magma_DEV, queue ));
magma_smfree(&hA_SELLP, queue );
magma_smfree( &dy, queue );
TESTING_CHECK( magma_svinit( &dy, Magma_DEV, dx.num_rows, dx.num_cols, c_zero, queue ));
// SpMV on GPU (SELLP)
start = magma_sync_wtime( queue );
for (j=0; j < 10; j++) {
TESTING_CHECK( magma_s_spmv( c_one, dA_SELLP, dx, c_zero, dy, queue ));
}
end = magma_sync_wtime( queue );
printf( " > MAGMA: %.2e seconds %.2e GFLOP/s (SELLP).\n",
(end-start)/10, FLOPS*10.*n/(end-start) );
TESTING_CHECK( magma_smtransfer( dy, &hcheck , Magma_DEV, Magma_CPU, queue ));
res = 0.0;
for(magma_int_t k=0; k < hA.num_rows; k++ ) {
res=res + MAGMA_S_REAL(hcheck.val[k]) - MAGMA_S_REAL(hrefvec.val[k]);
}
printf("%% |x-y|_F = %8.2e\n", res);
if ( res < accuracy )
printf("%% tester spmm SELL-P: ok\n");
else
printf("%% tester spmm SELL-P: failed\n");
magma_smfree( &hcheck, queue );
magma_smfree(&dA_SELLP, queue );
// SpMV on GPU (CUSPARSE - CSR)
// CUSPARSE context //
magma_smfree( &dy, queue );
TESTING_CHECK( magma_svinit( &dy, Magma_DEV, dx.num_rows, dx.num_cols, c_zero, queue ));
//#ifdef PRECISION_d
start = magma_sync_wtime( queue );
TESTING_CHECK( cusparseCreate( &cusparseHandle ));
TESTING_CHECK( cusparseSetStream( cusparseHandle, magma_queue_get_cuda_stream(queue) ));
TESTING_CHECK( cusparseCreateMatDescr( &descr ));
TESTING_CHECK( cusparseSetMatType( descr, CUSPARSE_MATRIX_TYPE_GENERAL ));
TESTING_CHECK( cusparseSetMatIndexBase( descr, CUSPARSE_INDEX_BASE_ZERO ));
float alpha = c_one;
float beta = c_zero;
// copy matrix to GPU
TESTING_CHECK( magma_smtransfer( hA, &dA, Magma_CPU, Magma_DEV, queue) );
for (j=0; j < 10; j++) {
cusparseScsrmm(cusparseHandle,
CUSPARSE_OPERATION_NON_TRANSPOSE,
dA.num_rows, n, dA.num_cols, dA.nnz,
&alpha, descr, dA.dval, dA.drow, dA.dcol,
dx.dval, dA.num_cols, &beta, dy.dval, dA.num_cols);
}
end = magma_sync_wtime( queue );
printf( " > CUSPARSE: %.2e seconds %.2e GFLOP/s (CSR).\n",
(end-start)/10, FLOPS*10*n/(end-start) );
TESTING_CHECK( magma_smtransfer( dy, &hcheck , Magma_DEV, Magma_CPU, queue ));
res = 0.0;
for(magma_int_t k=0; k < hA.num_rows; k++ ) {
res = res + MAGMA_S_REAL(hcheck.val[k]) - MAGMA_S_REAL(hrefvec.val[k]);
}
printf("%% |x-y|_F = %8.2e\n", res);
if ( res < accuracy )
printf("%% tester spmm cuSPARSE: ok\n");
else
printf("%% tester spmm cuSPARSE: failed\n");
magma_smfree( &hcheck, queue );
cusparseDestroyMatDescr( descr );
cusparseDestroy( cusparseHandle );
descr = NULL;
cusparseHandle = NULL;
//#endif
printf("\n\n");
// free CPU memory
magma_smfree( &hA, queue );
magma_smfree( &hx, queue );
magma_smfree( &hy, queue );
magma_smfree( &hrefvec, queue );
// free GPU memory
magma_smfree( &dx, queue );
magma_smfree( &dy, queue );
magma_smfree( &dA, queue);
#ifdef MAGMA_WITH_MKL
magma_free_cpu( pntre );
#endif
i++;
}
magma_queue_destroy( queue );
TESTING_CHECK( magma_finalize() );
return info;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing sgeqrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
const float d_neg_one = MAGMA_D_NEG_ONE;
const float d_one = MAGMA_D_ONE;
const float c_neg_one = MAGMA_S_NEG_ONE;
const float c_one = MAGMA_S_ONE;
const float c_zero = MAGMA_S_ZERO;
const magma_int_t ione = 1;
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
float Anorm, error=0, error2=0;
float *h_A, *h_R, *tau, *h_work, tmp[1];
magmaFloat_ptr d_A, dT;
magma_int_t M, N, n2, lda, ldda, lwork, info, min_mn, nb, size;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
opts.parse_opts( argc, argv );
magma_int_t status = 0;
float tol = opts.tolerance * lapackf77_slamch("E");
// version 3 can do either check
if (opts.check == 1 && opts.version == 1) {
opts.check = 2;
printf( "%% version 1 requires check 2 (solve A*x=b)\n" );
}
if (opts.check == 2 && opts.version == 2) {
opts.check = 1;
printf( "%% version 2 requires check 1 (R - Q^H*A)\n" );
}
printf( "%% version %d\n", (int) opts.version );
if ( opts.check == 1 ) {
printf("%% M N CPU Gflop/s (sec) GPU Gflop/s (sec) |R - Q^H*A| |I - Q^H*Q|\n");
printf("%%==============================================================================\n");
}
else {
printf("%% M N CPU Gflop/s (sec) GPU Gflop/s (sec) |b - A*x|\n");
printf("%%===============================================================\n");
}
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
min_mn = min( M, N );
lda = M;
n2 = lda*N;
ldda = magma_roundup( M, opts.align ); // multiple of 32 by default
nb = magma_get_sgeqrf_nb( M, N );
gflops = FLOPS_SGEQRF( M, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_sgeqrf( &M, &N, NULL, &M, NULL, tmp, &lwork, &info );
lwork = (magma_int_t)MAGMA_S_REAL( tmp[0] );
TESTING_MALLOC_CPU( tau, float, min_mn );
TESTING_MALLOC_CPU( h_A, float, n2 );
TESTING_MALLOC_CPU( h_work, float, lwork );
TESTING_MALLOC_PIN( h_R, float, n2 );
TESTING_MALLOC_DEV( d_A, float, ldda*N );
if ( opts.version == 1 || opts.version == 3 ) {
size = (2*min(M, N) + magma_roundup( N, 32 ) )*nb;
TESTING_MALLOC_DEV( dT, float, size );
magmablas_slaset( MagmaFull, size, 1, c_zero, c_zero, dT, size );
}
/* Initialize the matrix */
lapackf77_slarnv( &ione, ISEED, &n2, h_A );
lapackf77_slacpy( MagmaFullStr, &M, &N, h_A, &lda, h_R, &lda );
magma_ssetmatrix( M, N, h_R, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
nb = magma_get_sgeqrf_nb( M, N );
gpu_time = magma_wtime();
if ( opts.version == 1 ) {
// stores dT, V blocks have zeros, R blocks inverted & stored in dT
magma_sgeqrf_gpu( M, N, d_A, ldda, tau, dT, &info );
}
else if ( opts.version == 2 ) {
// LAPACK complaint arguments
magma_sgeqrf2_gpu( M, N, d_A, ldda, tau, &info );
}
#ifdef HAVE_CUBLAS
else if ( opts.version == 3 ) {
// stores dT, V blocks have zeros, R blocks stored in dT
magma_sgeqrf3_gpu( M, N, d_A, ldda, tau, dT, &info );
}
#endif
else {
printf( "Unknown version %d\n", (int) opts.version );
return -1;
}
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0) {
printf("magma_sgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
if ( opts.check == 1 && (opts.version == 2 || opts.version == 3) ) {
if ( opts.version == 3 ) {
// copy diagonal blocks of R back to A
for( int i=0; i < min_mn-nb; i += nb ) {
magma_int_t ib = min( min_mn-i, nb );
magmablas_slacpy( MagmaUpper, ib, ib, &dT[min_mn*nb + i*nb], nb, &d_A[ i + i*ldda ], ldda );
}
}
/* =====================================================================
Check the result, following zqrt01 except using the reduced Q.
This works for any M,N (square, tall, wide).
Only for version 2, which has LAPACK complaint output.
Or for version 3, after restoring diagonal blocks of A above.
=================================================================== */
magma_sgetmatrix( M, N, d_A, ldda, h_R, lda );
magma_int_t ldq = M;
magma_int_t ldr = min_mn;
float *Q, *R;
float *work;
TESTING_MALLOC_CPU( Q, float, ldq*min_mn ); // M by K
TESTING_MALLOC_CPU( R, float, ldr*N ); // K by N
TESTING_MALLOC_CPU( work, float, min_mn );
// generate M by K matrix Q, where K = min(M,N)
lapackf77_slacpy( "Lower", &M, &min_mn, h_R, &lda, Q, &ldq );
lapackf77_sorgqr( &M, &min_mn, &min_mn, Q, &ldq, tau, h_work, &lwork, &info );
assert( info == 0 );
// copy K by N matrix R
lapackf77_slaset( "Lower", &min_mn, &N, &c_zero, &c_zero, R, &ldr );
lapackf77_slacpy( "Upper", &min_mn, &N, h_R, &lda, R, &ldr );
// error = || R - Q^H*A || / (N * ||A||)
blasf77_sgemm( "Conj", "NoTrans", &min_mn, &N, &M,
&c_neg_one, Q, &ldq, h_A, &lda, &c_one, R, &ldr );
Anorm = lapackf77_slange( "1", &M, &N, h_A, &lda, work );
error = lapackf77_slange( "1", &min_mn, &N, R, &ldr, work );
if ( N > 0 && Anorm > 0 )
error /= (N*Anorm);
// set R = I (K by K identity), then R = I - Q^H*Q
// error = || I - Q^H*Q || / N
lapackf77_slaset( "Upper", &min_mn, &min_mn, &c_zero, &c_one, R, &ldr );
blasf77_ssyrk( "Upper", "Conj", &min_mn, &M, &d_neg_one, Q, &ldq, &d_one, R, &ldr );
error2 = safe_lapackf77_slansy( "1", "Upper", &min_mn, R, &ldr, work );
if ( N > 0 )
error2 /= N;
TESTING_FREE_CPU( Q ); Q = NULL;
TESTING_FREE_CPU( R ); R = NULL;
TESTING_FREE_CPU( work ); work = NULL;
}
else if ( opts.check == 2 && M >= N && (opts.version == 1 || opts.version == 3) ) {
/* =====================================================================
Check the result by solving consistent linear system, A*x = b.
Only for versions 1 & 3 with M >= N.
=================================================================== */
magma_int_t lwork2;
float *x, *b, *hwork;
magmaFloat_ptr d_B;
// initialize RHS, b = A*random
TESTING_MALLOC_CPU( x, float, N );
TESTING_MALLOC_CPU( b, float, M );
lapackf77_slarnv( &ione, ISEED, &N, x );
blasf77_sgemv( "Notrans", &M, &N, &c_one, h_A, &lda, x, &ione, &c_zero, b, &ione );
// copy to GPU
TESTING_MALLOC_DEV( d_B, float, M );
magma_ssetvector( M, b, 1, d_B, 1 );
if ( opts.version == 1 ) {
// allocate hwork
magma_sgeqrs_gpu( M, N, 1,
d_A, ldda, tau, dT,
d_B, M, tmp, -1, &info );
lwork2 = (magma_int_t)MAGMA_S_REAL( tmp[0] );
TESTING_MALLOC_CPU( hwork, float, lwork2 );
// solve linear system
magma_sgeqrs_gpu( M, N, 1,
d_A, ldda, tau, dT,
d_B, M, hwork, lwork2, &info );
if (info != 0) {
printf("magma_sgeqrs returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
TESTING_FREE_CPU( hwork );
}
#ifdef HAVE_CUBLAS
else if ( opts.version == 3 ) {
// allocate hwork
magma_sgeqrs3_gpu( M, N, 1,
d_A, ldda, tau, dT,
d_B, M, tmp, -1, &info );
lwork2 = (magma_int_t)MAGMA_S_REAL( tmp[0] );
TESTING_MALLOC_CPU( hwork, float, lwork2 );
// solve linear system
magma_sgeqrs3_gpu( M, N, 1,
d_A, ldda, tau, dT,
d_B, M, hwork, lwork2, &info );
if (info != 0) {
printf("magma_sgeqrs3 returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
TESTING_FREE_CPU( hwork );
}
#endif
else {
printf( "Unknown version %d\n", (int) opts.version );
return -1;
}
magma_sgetvector( N, d_B, 1, x, 1 );
// compute r = Ax - b, saved in b
blasf77_sgemv( "Notrans", &M, &N, &c_one, h_A, &lda, x, &ione, &c_neg_one, b, &ione );
// compute residual |Ax - b| / (max(m,n)*|A|*|x|)
float norm_x, norm_A, norm_r, work[1];
norm_A = lapackf77_slange( "F", &M, &N, h_A, &lda, work );
norm_r = lapackf77_slange( "F", &M, &ione, b, &M, work );
norm_x = lapackf77_slange( "F", &N, &ione, x, &N, work );
TESTING_FREE_CPU( x );
TESTING_FREE_CPU( b );
TESTING_FREE_DEV( d_B );
error = norm_r / (max(M,N) * norm_A * norm_x);
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_sgeqrf( &M, &N, h_A, &lda, tau, h_work, &lwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0) {
printf("lapackf77_sgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
}
/* =====================================================================
Print performance and error.
=================================================================== */
printf("%5d %5d ", (int) M, (int) N );
if ( opts.lapack ) {
printf( "%7.2f (%7.2f)", cpu_perf, cpu_time );
}
else {
printf(" --- ( --- )" );
}
printf( " %7.2f (%7.2f) ", gpu_perf, gpu_time );
if ( opts.check == 1 ) {
bool okay = (error < tol && error2 < tol);
status += ! okay;
printf( "%11.2e %11.2e %s\n", error, error2, (okay ? "ok" : "failed") );
}
else if ( opts.check == 2 ) {
if ( M >= N ) {
bool okay = (error < tol);
status += ! okay;
printf( "%10.2e %s\n", error, (okay ? "ok" : "failed") );
}
else {
printf( "(error check only for M >= N)\n" );
}
}
else {
printf( " ---\n" );
}
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_work );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
if ( opts.version == 1 || opts.version == 3 ) {
TESTING_FREE_DEV( dT );
}
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
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 */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing sgeqrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float error, work[1];
float c_neg_one = MAGMA_S_NEG_ONE;
float *h_A, *h_R, *tau, *dtau, *h_work, tmp[1];
float *d_A;
float *dwork;
magma_int_t M, N, n2, lda, ldda, lwork, info, min_mn;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
printf(" M N CPU GFlop/s (ms) GPU GFlop/s (ms) ||R||_F / ||A||_F\n");
printf("=======================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = FLOPS_SGEQRF( M, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_sgeqrf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_S_REAL( tmp[0] );
TESTING_MALLOC_CPU( tau, float, min_mn );
TESTING_MALLOC_CPU( h_A, float, n2 );
TESTING_MALLOC_CPU( h_work, float, lwork );
TESTING_MALLOC_PIN( h_R, float, n2 );
TESTING_MALLOC_DEV( d_A, float, ldda*N );
TESTING_MALLOC_DEV( dtau, float, min_mn );
TESTING_MALLOC_DEV( dwork, float, min_mn );
/* Initialize the matrix */
lapackf77_slarnv( &ione, ISEED, &n2, h_A );
lapackf77_slacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
magma_ssetmatrix( M, N, h_R, lda, d_A, ldda );
// warmup
if ( opts.warmup ) {
magma_sgeqr2_gpu( M, N, d_A, ldda, dtau, dwork, &info );
magma_ssetmatrix( M, N, h_R, lda, d_A, ldda );
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_sync_wtime( 0 );
magma_sgeqr2_gpu( M, N, d_A, ldda, dtau, dwork, &info );
gpu_time = magma_sync_wtime( 0 ) - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_sgeqr2_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_sgeqrf(&M, &N, h_A, &lda, tau, h_work, &lwork, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_sgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_sgetmatrix( M, N, d_A, ldda, h_R, M );
error = lapackf77_slange("f", &M, &N, h_A, &lda, work);
blasf77_saxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_slange("f", &M, &N, h_R, &lda, work) / error;
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) M, (int) N, cpu_perf, 1000.*cpu_time, gpu_perf, 1000.*gpu_time,
error, (error < tol ? "ok" : "failed"));
status += ! (error < tol);
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) M, (int) N, gpu_perf, 1000.*gpu_time );
}
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_work );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( dtau );
TESTING_FREE_DEV( dwork );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
extern "C" magma_err_t
magma_ssytrd(char uplo, magma_int_t n,
float *a, magma_int_t lda,
float *d, float *e, float *tau,
float *work, magma_int_t lwork,
magma_int_t *info, magma_queue_t queue)
{
/* -- clMAGMA (version 1.0.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
April 2012
Purpose
=======
SSYTRD reduces a real symmetric matrix A to real symmetric
tridiagonal form T by an orthogonal similarity transformation:
Q**T * A * Q = T.
Arguments
=========
UPLO (input) CHARACTER*1
= 'U': Upper triangle of A is stored;
= 'L': Lower triangle of A is stored.
N (input) INTEGER
The order of the matrix A. N >= 0.
A (input/output) REAL array, dimension (LDA,N)
On entry, the symmetric matrix A. If UPLO = 'U', the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = 'L', the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit, if UPLO = 'U', the diagonal and first superdiagonal
of A are overwritten by the corresponding elements of the
tridiagonal matrix T, and the elements above the first
superdiagonal, with the array TAU, represent the orthogonal
matrix Q as a product of elementary reflectors; if UPLO
= 'L', the diagonal and first subdiagonal of A are over-
written by the corresponding elements of the tridiagonal
matrix T, and the elements below the first subdiagonal, with
the array TAU, represent the orthogonal matrix Q as a product
of elementary reflectors. See Further Details.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
D (output) REAL array, dimension (N)
The diagonal elements of the tridiagonal matrix T:
D(i) = A(i,i).
E (output) REAL array, dimension (N-1)
The off-diagonal elements of the tridiagonal matrix T:
E(i) = A(i,i+1) if UPLO = 'U', E(i) = A(i+1,i) if UPLO = 'L'.
TAU (output) REAL array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details).
WORK (workspace/output) REAL array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The dimension of the array WORK. LWORK >= 1.
For optimum performance LWORK >= N*NB, where NB is the
optimal blocksize.
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
Further Details
===============
If UPLO = 'U', the matrix Q is represented as a product of elementary
reflectors
Q = H(n-1) . . . H(2) H(1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(i+1:n) = 0 and v(i) = 1; v(1:i-1) is stored on exit in
A(1:i-1,i+1), and tau in TAU(i).
If UPLO = 'L', the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(n-1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i) = 0 and v(i+1) = 1; v(i+2:n) is stored on exit in A(i+2:n,i),
and tau in TAU(i).
The contents of A on exit are illustrated by the following examples
with n = 5:
if UPLO = 'U': if UPLO = 'L':
( d e v2 v3 v4 ) ( d )
( d e v3 v4 ) ( e d )
( d e v4 ) ( v1 e d )
( d e ) ( v1 v2 e d )
( d ) ( v1 v2 v3 e d )
where d and e denote diagonal and off-diagonal elements of T, and vi
denotes an element of the vector defining H(i).
===================================================================== */
char uplo_[2] = {uplo, 0};
magma_int_t ldda = lda;
magma_int_t nb = magma_get_ssytrd_nb(n);
float c_neg_one = MAGMA_S_NEG_ONE;
float c_one = MAGMA_S_ONE;
float d_one = MAGMA_D_ONE;
magma_int_t kk, nx;
magma_int_t i, j, i_n;
magma_int_t iinfo;
magma_int_t ldwork, lddwork, lwkopt;
magma_int_t lquery;
*info = 0;
int upper = lapackf77_lsame(uplo_, "U");
lquery = lwork == -1;
if (! upper && ! lapackf77_lsame(uplo_, "L")) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
} else if (lwork < nb*n && ! lquery) {
*info = -9;
}
if (*info == 0) {
/* Determine the block size. */
ldwork = lddwork = n;
lwkopt = n * nb;
// ACD
// MAGMA_S_SET2REAL( work[0], lwkopt );
MAGMA_S_SET2REAL( work[0], (float) lwkopt );
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
/* Quick return if possible */
if (n == 0) {
work[0] = c_one;
return *info;
}
magmaFloat_ptr da;
size_t da_offset = 0;
if (MAGMA_SUCCESS != magma_malloc( &da, (n*ldda + 2*n*nb )*sizeof(float))) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magmaFloat_ptr dwork = da;
size_t dwork_offset = da_offset + (n)*ldda;
if (n < 2048)
nx = n;
else
nx = 512;
if (upper) {
/* Copy the matrix to the GPU */
magma_ssetmatrix( n, n, A(0, 0), 0, lda, dA(0, 0), ldda, queue );
/* Reduce the upper triangle of A.
Columns 1:kk are handled by the unblocked method. */
kk = n - (n - nx + nb - 1) / nb * nb;
for (i = n - nb; i >= kk; i -= nb)
{
/* Reduce columns i:i+nb-1 to tridiagonal form and form the
matrix W which is needed to update the unreduced part of
the matrix */
/* Get the current panel (no need for the 1st iteration) */
if (i!=n-nb)
magma_sgetmatrix( i+nb, nb, dA(0, i), ldda, A(0, i), 0, lda, queue );
magma_slatrd(uplo, i+nb, nb, A(0, 0), lda, e, tau,
work, ldwork, dA(0, 0), ldda, dwork, dwork_offset, lddwork, queue);
/* Update the unreduced submatrix A(0:i-2,0:i-2), using an
update of the form: A := A - V*W' - W*V' */
magma_ssetmatrix( i + nb, nb, work, 0, ldwork, dwork, dwork_offset, lddwork, queue );
magma_ssyr2k(magma_uplo_const(uplo), MagmaNoTrans, i, nb, c_neg_one,
dA(0, i), ldda, dwork, dwork_offset,
lddwork, d_one, dA(0, 0), ldda, queue);
/* Copy superdiagonal elements back into A, and diagonal
elements into D */
for (j = i; j < i+nb; ++j) {
MAGMA_S_SET2REAL( *A(j-1, j), e[j - 1] );
d[j] = MAGMA_S_REAL( *A(j, j) );
}
}
magma_sgetmatrix( kk, kk, dA(0, 0), ldda, A(0, 0), 0, lda, queue );
/* Use unblocked code to reduce the last or only block */
lapackf77_ssytd2(uplo_, &kk, A(0, 0), &lda, d, e, tau, &iinfo);
}
else
{
/* Copy the matrix to the GPU */
if (1<=n-nx)
magma_ssetmatrix( n, n, A(0,0), 0, lda, dA(0,0), ldda, queue );
#ifdef FAST_SYMV
// TODO this leaks memory from da, above
magmaFloat_ptr dwork2;
if (MAGMA_SUCCESS != magma_malloc( &dwork2, (n*n)*sizeof(float) )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
size_t dwork2_offset = 0;
#endif
/* Reduce the lower triangle of A */
for (i = 0; i < n-nx; i += nb)
{
/* Reduce columns i:i+nb-1 to tridiagonal form and form the
matrix W which is needed to update the unreduced part of
the matrix */
/* Get the current panel (no need for the 1st iteration) */
if (i!=0)
magma_sgetmatrix( n-i, nb, dA(i, i), ldda, A(i, i), 0, lda, queue );
#ifdef FAST_SYMV
// unported
magma_slatrd2(uplo, n-i, nb, A(i, i), lda, &e[i],
&tau[i], work, ldwork,
dA(i, i), ldda,
dwork, lddwork, dwork2, n*n);
#else
magma_slatrd(uplo, n-i, nb, A(i, i), lda, &e[i],
&tau[i], work, ldwork,
dA(i, i), ldda,
dwork, dwork_offset, lddwork, queue);
#endif
/* Update the unreduced submatrix A(i+ib:n,i+ib:n), using
an update of the form: A := A - V*W' - W*V' */
magma_ssetmatrix( n-i, nb, work, 0, ldwork, dwork, dwork_offset, lddwork, queue );
magma_ssyr2k(MagmaLower, MagmaNoTrans, n-i-nb, nb, c_neg_one,
dA(i+nb, i), ldda,
dwork, (dwork_offset+nb), lddwork, d_one,
dA(i+nb, i+nb), ldda, queue);
/* Copy subdiagonal elements back into A, and diagonal
elements into D */
for (j = i; j < i+nb; ++j) {
MAGMA_S_SET2REAL( *A(j+1, j), e[j] );
d[j] = MAGMA_S_REAL( *A(j, j) );
}
}
#ifdef FAST_SYMV
magma_free( dwork2 );
#endif
/* Use unblocked code to reduce the last or only block */
if (1<=n-nx)
magma_sgetmatrix( n-i, n-i, dA(i, i), ldda, A(i, i), 0, lda, queue );
i_n = n-i;
lapackf77_ssytrd(uplo_, &i_n, A(i, i), &lda, &d[i], &e[i],
&tau[i], work, &lwork, &iinfo);
}
magma_free( da );
// ACD
// MAGMA_S_SET2REAL( work[0], lwkopt );
MAGMA_S_SET2REAL( work[0], (float) lwkopt );
return *info;
} /* magma_ssytrd */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing sormqr_gpu
*/
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float error, work[1];
float c_neg_one = MAGMA_S_NEG_ONE;
magma_int_t ione = 1;
magma_int_t m, n, k, size, info;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t nb, ldc, lda, lwork, lwork_max, dt_size;
float *C, *R, *A, *W, *tau;
float *dC, *dA, *dT;
magma_opts opts;
parse_opts( argc, argv, &opts );
// test all combinations of input parameters
const char* side[] = { MagmaLeftStr, MagmaRightStr };
const char* trans[] = { MagmaTransStr, MagmaNoTransStr };
printf(" M N K side trans CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / ||QC||_F\n");
printf("===============================================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iside = 0; iside < 2; ++iside ) {
for( int itran = 0; itran < 2; ++itran ) {
m = opts.msize[i];
n = opts.nsize[i];
k = opts.ksize[i];
nb = magma_get_sgeqrf_nb( m );
ldc = ((m + 31)/32)*32;
lda = ((max(m,n) + 31)/32)*32;
gflops = FLOPS_SORMQR( m, n, k, *side[iside] ) / 1e9;
if ( *side[iside] == 'L' && m < k ) {
printf( "%5d %5d %5d %-5s %-9s skipping because side=left and m < k\n",
(int) m, (int) n, (int) k, side[iside], trans[itran] );
continue;
}
if ( *side[iside] == 'R' && n < k ) {
printf( "%5d %5d %5d %-5s %-9s skipping because side=right and n < k\n",
(int) m, (int) n, (int) k, side[iside], trans[itran] );
continue;
}
if ( *side[iside] == 'L' ) {
// side = left
lwork_max = (m - k + nb)*(n + nb) + n*nb;
dt_size = ( 2*min(m,k) + ((k + 31)/32)*32 )*nb;
}
else {
// side = right
lwork_max = (n - k + nb)*(m + nb) + m*nb;
dt_size = ( 2*min(n,k) + ((k + 31)/32)*32 )*nb;
}
TESTING_MALLOC_CPU( C, float, ldc*n );
TESTING_MALLOC_CPU( R, float, ldc*n );
TESTING_MALLOC_CPU( A, float, lda*k );
TESTING_MALLOC_CPU( W, float, lwork_max );
TESTING_MALLOC_CPU( tau, float, k );
TESTING_MALLOC_DEV( dC, float, ldc*n );
TESTING_MALLOC_DEV( dA, float, lda*k );
TESTING_MALLOC_DEV( dT, float, dt_size );
// C is full, m x n
size = ldc*n;
lapackf77_slarnv( &ione, ISEED, &size, C );
magma_ssetmatrix( m, n, C, ldc, dC, ldc );
// A is m x k (left) or n x k (right)
lda = (*side[iside] == 'L' ? m : n);
size = lda*k;
lapackf77_slarnv( &ione, ISEED, &size, A );
// compute QR factorization to get Householder vectors in dA, tau, dT
magma_ssetmatrix( lda, k, A, lda, dA, lda );
magma_sgeqrf_gpu( lda, k, dA, lda, tau, dT, &info );
magma_sgetmatrix( lda, k, dA, lda, A, lda );
if (info != 0)
printf("magma_sgeqrf_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_sormqr( side[iside], trans[itran],
&m, &n, &k,
A, &lda, tau, C, &ldc, W, &lwork_max, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_sormqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
// query for workspace size
lwork = -1;
magma_sormqr_gpu( *side[iside], *trans[itran],
m, n, k,
dA, lda, tau, dC, ldc, W, lwork, dT, nb, &info );
if (info != 0)
printf("magma_sormqr_gpu (lwork query) returned error %d: %s.\n",
(int) info, magma_strerror( info ));
lwork = (magma_int_t) MAGMA_S_REAL( W[0] );
if ( lwork < 0 || lwork > lwork_max )
printf("invalid lwork %d, lwork_max %d\n", (int) lwork, (int) lwork_max );
gpu_time = magma_sync_wtime( 0 ); // sync needed for L,N and R,T cases
magma_sormqr_gpu( *side[iside], *trans[itran],
m, n, k,
dA, lda, tau, dC, ldc, W, lwork, dT, nb, &info );
gpu_time = magma_sync_wtime( 0 ) - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_sormqr_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_sgetmatrix( m, n, dC, ldc, R, ldc );
/* =====================================================================
compute relative error |QC_magma - QC_lapack| / |QC_lapack|
=================================================================== */
error = lapackf77_slange( "Fro", &m, &n, C, &ldc, work );
size = ldc*n;
blasf77_saxpy( &size, &c_neg_one, C, &ione, R, &ione );
error = lapackf77_slange( "Fro", &m, &n, R, &ldc, work ) / error;
printf( "%5d %5d %5d %-5s %-9s %7.2f (%7.2f) %7.2f (%7.2f) %8.2e\n",
(int) m, (int) n, (int) k, side[iside], trans[itran],
cpu_perf, cpu_time, gpu_perf, gpu_time, error );
TESTING_FREE_CPU( C );
TESTING_FREE_CPU( R );
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( W );
TESTING_FREE_CPU( tau );
TESTING_FREE_DEV( dC );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dT );
}} // end iside, itran
printf( "\n" );
}
TESTING_FINALIZE();
return 0;
}
int APPLY_SPECIFIC(magma_svd)(PyGpuArrayObject *A,
PyGpuArrayObject **S,
PyGpuArrayObject **U, // may be NULL
PyGpuArrayObject **VT, // may be NULL
PARAMS_TYPE* params) {
bool compute_uv = (U != NULL);
magma_int_t *iwork = NULL, iunused[1];
magma_int_t M, N, K, ldu, ldv, M_U, N_VT, info;
magma_vec_t jobz;
size_t s_dims[1], u_dims[2], vt_dims[2];
float *a_data = NULL, *s_data = NULL, *u_data = NULL, *vt_data = NULL,
*work = NULL;
float dummy[1];
int res = -1, lwork;
if (A->ga.typecode != GA_FLOAT) {
PyErr_SetString(PyExc_TypeError,
"GpuMagmaMatrixInverse: Unsupported data type");
return -1;
}
// This is early to match the exit() in the fail label.
cuda_enter(params->context->ctx);
magma_init();
if (!GpuArray_IS_C_CONTIGUOUS(&A->ga)) {
PyErr_SetString(PyExc_ValueError,
"GpuMagmaMatrixInverse: requires data to be C-contiguous");
return 1;
}
if (PyGpuArray_NDIM(A) != 2) {
PyErr_SetString(PyExc_ValueError,
"GpuMagmaMatrixInverse: matrix rank error");
goto fail;
}
// magma matrix svd
// reverse dimensions because MAGMA expects column-major matrices:
M = PyGpuArray_DIM(A, 1);
N = PyGpuArray_DIM(A, 0);
K = std::min(M, N);
if (MAGMA_SUCCESS != magma_smalloc_pinned(&a_data, M * N)) {
PyErr_SetString(PyExc_RuntimeError,
"GpuMagmaSVD: failed to allocate memory");
goto fail;
}
cudaMemcpy(a_data, PyGpuArray_DEV_DATA(A), M * N * sizeof(float),
cudaMemcpyDeviceToDevice);
if (MAGMA_SUCCESS != magma_smalloc_pinned(&s_data, K)) {
PyErr_SetString(PyExc_RuntimeError,
"GpuMagmaSVD: failed to allocate memory");
goto fail;
}
if (compute_uv) {
if (params->full_matrices) {
jobz = MagmaAllVec;
} else {
jobz = MagmaSomeVec;
}
M_U = (jobz == MagmaAllVec ? M : K);
N_VT = (jobz == MagmaAllVec ? N : K);
ldu = M;
ldv = N_VT;
if (MAGMA_SUCCESS != magma_smalloc_pinned(&u_data, M_U * M)) {
PyErr_SetString(PyExc_RuntimeError,
"GpuMagmaSVD: failed to allocate memory");
goto fail;
}
if (MAGMA_SUCCESS != magma_smalloc_pinned(&vt_data, N * N_VT)) {
PyErr_SetString(PyExc_RuntimeError,
"GpuMagmaSVD: failed to allocate memory");
goto fail;
}
} else {
jobz = MagmaNoVec;
ldu = M;
ldv = N;
}
// query for workspace size
magma_sgesdd(jobz, M, N, NULL, M, NULL, NULL, ldu, NULL, ldv,
dummy, -1, iunused, &info);
lwork = (magma_int_t) MAGMA_S_REAL(dummy[0]);
if (MAGMA_SUCCESS != magma_smalloc_pinned(&work, lwork)) {
PyErr_SetString(PyExc_RuntimeError,
"GpuMagmaSVD: failed to allocate working memory");
goto fail;
}
if (MAGMA_SUCCESS != magma_imalloc_cpu(&iwork, 8*K)) {
PyErr_SetString(PyExc_RuntimeError,
"GpuMagmaSVD: failed to allocate working memory");
goto fail;
}
// compute svd
magma_sgesdd(jobz, M, N, a_data, M, s_data,
u_data, ldu, vt_data, ldv, work, lwork, iwork, &info);
if (info > 0) {
PyErr_Format(
PyExc_RuntimeError,
"GpuMagmaSVD: the updating process of SBDSDC did not converge (error: %d)",
info);
goto fail;
} else if (info < 0) {
PyErr_Format(
PyExc_RuntimeError,
"GpuMagmaSVD: magma_sgesdd_gpu argument %d has an illegal value", -info);
goto fail;
}
s_dims[0] = K;
if (theano_prep_output(S, 1, s_dims, A->ga.typecode, GA_C_ORDER, params->context) != 0){
PyErr_SetString(PyExc_RuntimeError,
"GpuMagmaSVD: failed to allocate memory");
goto fail;
}
cudaMemcpy(PyGpuArray_DEV_DATA(*S), s_data, K * sizeof(float),
cudaMemcpyDeviceToDevice);
if (compute_uv) {
u_dims[0] = N; u_dims[1] = N_VT;
if (theano_prep_output(U, 2, u_dims, A->ga.typecode, GA_C_ORDER, params->context) != 0){
PyErr_SetString(PyExc_RuntimeError,
"GpuMagmaSVD: failed to allocate memory");
goto fail;
}
// magma expects column-major matrices. Exchange u_data -> VT and vt_data -> U
// to match numpy.linalg.svd output
cudaMemcpy(PyGpuArray_DEV_DATA(*U), vt_data, N * N_VT * sizeof(float),
cudaMemcpyDeviceToDevice);
vt_dims[0] = M_U; vt_dims[1] = M;
if (theano_prep_output(VT, 2, vt_dims, A->ga.typecode, GA_C_ORDER, params->context) != 0){
PyErr_SetString(PyExc_RuntimeError,
"GpuMagmaSVD: failed to allocate memory");
goto fail;
}
// magma expects column-major matrices. Exchange u_data -> VT and vt_data -> U
// to match numpy.linalg.svd output
cudaMemcpy(PyGpuArray_DEV_DATA(*VT), u_data, M_U * M * sizeof(float),
cudaMemcpyDeviceToDevice);
}
res = 0;
fail:
if (a_data != NULL)
magma_free_pinned(a_data);
if (s_data != NULL)
magma_free_pinned(s_data);
if (u_data != NULL)
magma_free_pinned(u_data);
if (vt_data != NULL)
magma_free_pinned(vt_data);
if (work != NULL)
magma_free_pinned(work);
if (iwork != NULL)
magma_free_cpu(iwork);
magma_finalize();
cuda_exit(params->context->ctx);
return res;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing sgegqr
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float e1, e2, e3, e4, e5, *work;
float c_neg_one = MAGMA_S_NEG_ONE;
float c_one = MAGMA_S_ONE;
float c_zero = MAGMA_S_ZERO;
float *h_A, *h_R, *tau, *dtau, *h_work, *h_rwork, tmp[1];
float *d_A, *dwork;
magma_int_t M, N, n2, lda, ldda, lwork, info, min_mn;
magma_int_t ione = 1, ldwork;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
// versions 1...4 are valid
if (opts.version < 1 || opts.version > 4) {
printf("Unknown version %d; exiting\n", opts.version );
return -1;
}
float tol, eps = lapackf77_slamch("E");
tol = 10* opts.tolerance * eps;
printf(" M N CPU GFlop/s (ms) GPU GFlop/s (ms) ||I-Q'Q||_F / M ||I-Q'Q||_I / M ||A-Q R||_I\n");
printf(" MAGMA / LAPACK MAGMA / LAPACK\n");
printf("==========================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
if (N > 128) {
printf("%5d %5d skipping because sgegqr requires N <= 128\n",
(int) M, (int) N);
continue;
}
if (M < N) {
printf("%5d %5d skipping because sgegqr requires M >= N\n",
(int) M, (int) N);
continue;
}
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = FLOPS_SGEQRF( M, N ) / 1e9 + FLOPS_SORGQR( M, N, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_sgeqrf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_S_REAL( tmp[0] );
lwork = max(lwork, 3*N*N);
ldwork = N*N;
if (opts.version == 2) {
ldwork = 3*N*N + min_mn;
}
TESTING_MALLOC_PIN( tau, float, min_mn );
TESTING_MALLOC_PIN( h_work, float, lwork );
TESTING_MALLOC_PIN(h_rwork, float, lwork );
TESTING_MALLOC_CPU( h_A, float, n2 );
TESTING_MALLOC_CPU( h_R, float, n2 );
TESTING_MALLOC_CPU( work, float, M );
TESTING_MALLOC_DEV( d_A, float, ldda*N );
TESTING_MALLOC_DEV( dtau, float, min_mn );
TESTING_MALLOC_DEV( dwork, float, ldwork );
/* Initialize the matrix */
lapackf77_slarnv( &ione, ISEED, &n2, h_A );
lapackf77_slacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
magma_ssetmatrix( M, N, h_R, lda, d_A, ldda );
// warmup
magma_sgegqr_gpu( 1, M, N, d_A, ldda, dwork, h_work, &info );
magma_ssetmatrix( M, N, h_R, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_sync_wtime( 0 );
magma_sgegqr_gpu( opts.version, M, N, d_A, ldda, dwork, h_rwork, &info );
gpu_time = magma_sync_wtime( 0 ) - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_sgegqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_sgetmatrix( M, N, d_A, ldda, h_R, M );
// Regenerate R
// blasf77_sgemm("t", "n", &N, &N, &M, &c_one, h_R, &M, h_A, &M, &c_zero, h_rwork, &N);
// magma_sprint(N, N, h_work, N);
blasf77_strmm("r", "u", "n", "n", &M, &N, &c_one, h_rwork, &N, h_R, &M);
blasf77_saxpy( &n2, &c_neg_one, h_A, &ione, h_R, &ione );
e5 = lapackf77_slange("i", &M, &N, h_R, &M, work) /
lapackf77_slange("i", &M, &N, h_A, &lda, work);
magma_sgetmatrix( M, N, d_A, ldda, h_R, M );
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
/* Orthogonalize on the CPU */
lapackf77_sgeqrf(&M, &N, h_A, &lda, tau, h_work, &lwork, &info);
lapackf77_sorgqr(&M, &N, &N, h_A, &lda, tau, h_work, &lwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_sorgqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
blasf77_sgemm("t", "n", &N, &N, &M, &c_one, h_R, &M, h_R, &M, &c_zero, h_work, &N);
for(int ii = 0; ii < N*N; ii += N+1 ) {
h_work[ii] = MAGMA_S_SUB(h_work[ii], c_one);
}
e1 = lapackf77_slange("f", &N, &N, h_work, &N, work) / N;
e3 = lapackf77_slange("i", &N, &N, h_work, &N, work) / N;
blasf77_sgemm("t", "n", &N, &N, &M, &c_one, h_A, &M, h_A, &M, &c_zero, h_work, &N);
for(int ii = 0; ii < N*N; ii += N+1 ) {
h_work[ii] = MAGMA_S_SUB(h_work[ii], c_one);
}
e2 = lapackf77_slange("f", &N, &N, h_work, &N, work) / N;
e4 = lapackf77_slange("i", &N, &N, h_work, &N, work) / N;
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e / %8.2e %8.2e / %8.2e %8.2e %s\n",
(int) M, (int) N, cpu_perf, 1000.*cpu_time, gpu_perf, 1000.*gpu_time,
e1, e2, e3, e4, e5,
(e1 < tol ? "ok" : "failed"));
status += ! (e1 < tol);
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) M, (int) N, gpu_perf, 1000.*gpu_time );
}
TESTING_FREE_PIN( tau );
TESTING_FREE_PIN( h_work );
TESTING_FREE_PIN( h_rwork );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_R );
TESTING_FREE_CPU( work );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( dtau );
TESTING_FREE_DEV( dwork );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
int main( int argc, char** argv)
{
real_Double_t gflops, gpu_perf, cpu_perf, gpu_time, cpu_time;
float matnorm, work[1];
float mzone = MAGMA_S_NEG_ONE;
float *h_A, *h_R, *tau, *hwork, tmp[1];
magmaFloat_ptr d_A;
/* Matrix size */
magma_int_t M = 0, N = 0, n2, lda, ldda, lhwork;
magma_int_t size[10] = {1024,2048,3072,4032,5184,6016,7040,8064,9088,10176};
magma_int_t i, info, min_mn;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
if (argc != 1){
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0)
N = atoi(argv[++i]);
else if (strcmp("-M", argv[i])==0)
M = atoi(argv[++i]);
}
if ( M == 0 ) {
M = N;
}
if ( N == 0 ) {
N = M;
}
if (M>0 && N>0)
printf(" testing_sgeqrf_gpu -M %d -N %d\n\n", M, N);
else
{
printf("\nUsage: \n");
printf(" testing_sgeqrf_gpu -M %d -N %d\n\n", 1024, 1024);
exit(1);
}
}
else {
printf("\nUsage: \n");
printf(" testing_sgeqrf_gpu -M %d -N %d\n\n", 1024, 1024);
M = N = size[7];
}
/* Initialize */
magma_queue_t queue1, queue2;
magma_device_t device[ MagmaMaxGPUs ];
int num = 0;
magma_err_t err;
magma_init();
err = magma_get_devices( device, MagmaMaxGPUs, &num );
if ( err != 0 || num < 1 ) {
fprintf( stderr, "magma_get_devices failed: %d\n", err );
exit(-1);
}
err = magma_queue_create( device[0], &queue1 );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", err );
exit(-1);
}
err = magma_queue_create( device[0], &queue2 );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", err );
exit(-1);
}
magma_queue_t queues[2] = {queue1, queue2};
ldda = ((M+31)/32)*32;
n2 = M * N;
min_mn = min(M, N);
/* Allocate host memory for the matrix */
TESTING_MALLOC_CPU( tau, float, min_mn );
TESTING_MALLOC_CPU( h_A, float, n2 );
TESTING_MALLOC_PIN( h_R, float, n2 );
TESTING_MALLOC_DEV( d_A, float, ldda*N );
lhwork = -1;
lapackf77_sgeqrf(&M, &N, h_A, &M, tau, tmp, &lhwork, &info);
lhwork = (magma_int_t)MAGMA_S_REAL( tmp[0] );
TESTING_MALLOC_CPU( hwork, float, lhwork );
printf("\n\n");
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / ||A||_F\n");
printf("======================================================================\n");
for(i=0; i<8; i++){
if (argc == 1){
M = N = size[i];
}
min_mn= min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = FLOPS( (float)M, (float)N ) * 1e-9;
/* Initialize the matrix */
lapackf77_slarnv( &ione, ISEED, &n2, h_A );
lapackf77_slacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_sgeqrf(&M, &N, h_A, &M, tau, hwork, &lhwork, &info);
cpu_time = magma_wtime() - cpu_time;
if (info < 0)
printf("Argument %d of lapack_sgeqrf had an illegal value.\n", -info);
cpu_perf = gflops / cpu_time;
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_ssetmatrix( M, N, h_R, 0, lda, d_A, 0, ldda, queue1 );
magma_sgeqrf2_gpu( M, N, d_A, 0, ldda, tau, &info, queues);
magma_ssetmatrix( M, N, h_R, 0, lda, d_A, 0, ldda, queue1 );
clFinish(queue1);
clFinish(queue2);
gpu_time = magma_wtime();
magma_sgeqrf2_gpu( M, N, d_A, 0, ldda, tau, &info, queues);
gpu_time = magma_wtime() - gpu_time;
if (info < 0)
printf("Argument %d of magma_sgeqrf2 had an illegal value.\n", -info);
gpu_perf = gflops / gpu_time;
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_sgetmatrix( M, N, d_A, 0, ldda, h_R, 0, M, queue1 );
matnorm = lapackf77_slange("f", &M, &N, h_A, &M, work);
blasf77_saxpy(&n2, &mzone, h_A, &ione, h_R, &ione);
printf("%5d %5d %6.2f (%6.2f) %6.2f (%6.2f) %e\n",
M, N, cpu_perf, cpu_time, gpu_perf, gpu_time,
lapackf77_slange("f", &M, &N, h_R, &M, work) / matnorm);
if (argc != 1)
break;
}
/* clean up */
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( hwork );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
magma_queue_destroy( queue1 );
magma_queue_destroy( queue2 );
magma_finalize();
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing sgetrf
*/
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float *h_A, *h_R, *work;
magmaFloat_ptr d_A, dwork;
float c_neg_one = MAGMA_S_NEG_ONE;
magma_int_t N, n2, lda, ldda, info, lwork, ldwork;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
float tmp;
float error, rwork[1];
magma_int_t *ipiv;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
// need looser bound (3000*eps instead of 30*eps) for tests
// TODO: should compute ||I - A*A^{-1}|| / (n*||A||*||A^{-1}||)
opts.tolerance = max( 3000., opts.tolerance );
float tol = opts.tolerance * lapackf77_slamch("E");
printf(" N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / (N*||A||_F)\n");
printf("=================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
n2 = lda*N;
ldda = ((N+31)/32)*32;
ldwork = N * magma_get_sgetri_nb( N );
gflops = FLOPS_SGETRI( N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_sgetri( &N, NULL, &lda, NULL, &tmp, &lwork, &info );
if (info != 0)
printf("lapackf77_sgetri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
lwork = int( MAGMA_S_REAL( tmp ));
TESTING_MALLOC_CPU( ipiv, magma_int_t, N );
TESTING_MALLOC_CPU( work, float, lwork );
TESTING_MALLOC_CPU( h_A, float, n2 );
TESTING_MALLOC_PIN( h_R, float, n2 );
TESTING_MALLOC_DEV( d_A, float, ldda*N );
TESTING_MALLOC_DEV( dwork, float, ldwork );
/* Initialize the matrix */
lapackf77_slarnv( &ione, ISEED, &n2, h_A );
error = lapackf77_slange( "f", &N, &N, h_A, &lda, rwork ); // norm(A)
/* Factor the matrix. Both MAGMA and LAPACK will use this factor. */
magma_ssetmatrix( N, N, h_A, lda, d_A, ldda );
magma_sgetrf_gpu( N, N, d_A, ldda, ipiv, &info );
magma_sgetmatrix( N, N, d_A, ldda, h_A, lda );
if ( info != 0 )
printf("magma_sgetrf_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// check for exact singularity
//h_A[ 10 + 10*lda ] = MAGMA_S_MAKE( 0.0, 0.0 );
//magma_ssetmatrix( N, N, h_A, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_sgetri_gpu( N, d_A, ldda, ipiv, dwork, ldwork, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_sgetri_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_sgetmatrix( N, N, d_A, ldda, h_R, lda );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_sgetri( &N, h_A, &lda, ipiv, work, &lwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_sgetri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
blasf77_saxpy( &n2, &c_neg_one, h_A, &ione, h_R, &ione );
error = lapackf77_slange( "f", &N, &N, h_R, &lda, rwork ) / (N*error);
printf( "%5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) N, cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed"));
status += ! (error < tol);
}
else {
printf( "%5d --- ( --- ) %7.2f (%7.2f) ---\n",
(int) N, gpu_perf, gpu_time );
}
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( work );
TESTING_FREE_CPU( h_A );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( dwork );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
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 */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing sormqr
*/
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float error, work[1];
float c_neg_one = MAGMA_S_NEG_ONE;
magma_int_t ione = 1;
magma_int_t mm, m, n, k, size, info;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t nb, ldc, lda, lwork, lwork_max;
float *C, *R, *A, *W, *tau;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
// need slightly looser bound (60*eps instead of 30*eps) for some tests
opts.tolerance = max( 60., opts.tolerance );
float tol = opts.tolerance * lapackf77_slamch("E");
// test all combinations of input parameters
magma_side_t side [] = { MagmaLeft, MagmaRight };
magma_trans_t trans[] = { MagmaTrans, MagmaNoTrans };
printf(" M N K side trans CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / ||QC||_F\n");
printf("===============================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iside = 0; iside < 2; ++iside ) {
for( int itran = 0; itran < 2; ++itran ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
m = opts.msize[itest];
n = opts.nsize[itest];
k = opts.ksize[itest];
nb = magma_get_sgeqrf_nb( m );
ldc = m;
// A is m x k (left) or n x k (right)
mm = (side[iside] == MagmaLeft ? m : n);
lda = mm;
gflops = FLOPS_SORMQR( m, n, k, side[iside] ) / 1e9;
if ( side[iside] == MagmaLeft && m < k ) {
printf( "%5d %5d %5d %4c %5c skipping because side=left and m < k\n",
(int) m, (int) n, (int) k,
lapacke_side_const( side[iside] ),
lapacke_trans_const( trans[itran] ) );
continue;
}
if ( side[iside] == MagmaRight && n < k ) {
printf( "%5d %5d %5d %4c %5c skipping because side=right and n < k\n",
(int) m, (int) n, (int) k,
lapacke_side_const( side[iside] ),
lapacke_trans_const( trans[itran] ) );
continue;
}
// need at least 2*nb*nb for geqrf
lwork_max = max( max( m*nb, n*nb ), 2*nb*nb );
TESTING_MALLOC_CPU( C, float, ldc*n );
TESTING_MALLOC_CPU( R, float, ldc*n );
TESTING_MALLOC_CPU( A, float, lda*k );
TESTING_MALLOC_CPU( W, float, lwork_max );
TESTING_MALLOC_CPU( tau, float, k );
// C is full, m x n
size = ldc*n;
lapackf77_slarnv( &ione, ISEED, &size, C );
lapackf77_slacpy( "Full", &m, &n, C, &ldc, R, &ldc );
size = lda*k;
lapackf77_slarnv( &ione, ISEED, &size, A );
// compute QR factorization to get Householder vectors in A, tau
magma_sgeqrf( mm, k, A, lda, tau, W, lwork_max, &info );
if (info != 0)
printf("magma_sgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_sormqr( lapack_side_const( side[iside] ), lapack_trans_const( trans[itran] ),
&m, &n, &k,
A, &lda, tau, C, &ldc, W, &lwork_max, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_sormqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
// query for workspace size
lwork = -1;
magma_sormqr( side[iside], trans[itran],
m, n, k,
A, lda, tau, R, ldc, W, lwork, &info );
if (info != 0)
printf("magma_sormqr (lwork query) returned error %d: %s.\n",
(int) info, magma_strerror( info ));
lwork = (magma_int_t) MAGMA_S_REAL( W[0] );
if ( lwork < 0 || lwork > lwork_max ) {
printf("optimal lwork %d > lwork_max %d\n", (int) lwork, (int) lwork_max );
lwork = lwork_max;
}
gpu_time = magma_wtime();
magma_sormqr( side[iside], trans[itran],
m, n, k,
A, lda, tau, R, ldc, W, lwork, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_sormqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
compute relative error |QC_magma - QC_lapack| / |QC_lapack|
=================================================================== */
error = lapackf77_slange( "Fro", &m, &n, C, &ldc, work );
size = ldc*n;
blasf77_saxpy( &size, &c_neg_one, C, &ione, R, &ione );
error = lapackf77_slange( "Fro", &m, &n, R, &ldc, work ) / error;
printf( "%5d %5d %5d %4c %5c %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) m, (int) n, (int) k,
lapacke_side_const( side[iside] ),
lapacke_trans_const( trans[itran] ),
cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed") );
status += ! (error < tol);
TESTING_FREE_CPU( C );
TESTING_FREE_CPU( R );
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( W );
TESTING_FREE_CPU( tau );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}} // end iside, itran
printf( "\n" );
}
TESTING_FINALIZE();
return status;
}
/**
Purpose
-------
SSYTRD reduces a real symmetric matrix A to real symmetric
tridiagonal form T by an orthogonal similarity transformation:
Q**H * A * Q = T.
Arguments
---------
@param[in]
num_gpus INTEGER
The number of GPUs. num_gpus > 0.
@param[in]
num_streams INTEGER
The number of GPU streams used for update. 10 >= num_streams > 0.
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored;
- = MagmaLower: Lower triangle of A is stored.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
A REAL array, dimension (LDA,N)
On entry, the symmetric matrix A. If UPLO = MagmaUpper, the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit, if UPLO = MagmaUpper, the diagonal and first superdiagonal
of A are overwritten by the corresponding elements of the
tridiagonal matrix T, and the elements above the first
superdiagonal, with the array TAU, represent the orthogonal
matrix Q as a product of elementary reflectors; if UPLO
= MagmaLower, the diagonal and first subdiagonal of A are over-
written by the corresponding elements of the tridiagonal
matrix T, and the elements below the first subdiagonal, with
the array TAU, represent the orthogonal matrix Q as a product
of elementary reflectors. See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
d REAL array, dimension (N)
The diagonal elements of the tridiagonal matrix T:
D(i) = A(i,i).
@param[out]
e REAL array, dimension (N-1)
The off-diagonal elements of the tridiagonal matrix T:
E(i) = A(i,i+1) if UPLO = MagmaUpper, E(i) = A(i+1,i) if UPLO = MagmaLower.
@param[out]
tau REAL array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details).
@param[out]
work (workspace) REAL array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK. LWORK >= 1.
For optimum performance LWORK >= N*NB, where NB is the
optimal blocksize.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
Further Details
---------------
If UPLO = MagmaUpper, the matrix Q is represented as a product of elementary
reflectors
Q = H(n-1) . . . H(2) H(1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(i+1:n) = 0 and v(i) = 1; v(1:i-1) is stored on exit in
A(1:i-1,i+1), and tau in TAU(i).
If UPLO = MagmaLower, the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(n-1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i) = 0 and v(i+1) = 1; v(i+2:n) is stored on exit in A(i+2:n,i),
and tau in TAU(i).
The contents of A on exit are illustrated by the following examples
with n = 5:
if UPLO = MagmaUpper: if UPLO = MagmaLower:
( d e v2 v3 v4 ) ( d )
( d e v3 v4 ) ( e d )
( d e v4 ) ( v1 e d )
( d e ) ( v1 v2 e d )
( d ) ( v1 v2 v3 e d )
where d and e denote diagonal and off-diagonal elements of T, and vi
denotes an element of the vector defining H(i).
@ingroup magma_ssyev_comp
********************************************************************/
extern "C" magma_int_t
magma_ssytrd_mgpu(
magma_int_t num_gpus, magma_int_t num_streams, magma_uplo_t uplo, magma_int_t n,
float *A, magma_int_t lda,
float *d, float *e, float *tau,
float *work, magma_int_t lwork,
magma_int_t *info)
{
#define A(i, j) (A + (j)*lda + (i))
#define dA(id, i, j) (dA[(id)] + (j)*ldda + (i))
#define dW(id, i, j) (dwork[(id)] + (j)*ldda + (i))
const char* uplo_ = lapack_uplo_const( uplo );
magma_int_t ln, ldda;
magma_int_t nb = magma_get_ssytrd_nb(n), ib;
float c_neg_one = MAGMA_S_NEG_ONE;
float c_one = MAGMA_S_ONE;
float d_one = MAGMA_D_ONE;
//float mv_time = 0.0;
#ifdef PROFILE_SY2RK
float up_time = 0.0;
#endif
magma_int_t kk, nx;
magma_int_t i = 0, ii, iii, j, did, i_n;
magma_int_t iinfo;
magma_int_t ldwork, lddwork, lwkopt, ldwork2;
magma_int_t lquery;
magma_queue_t stream[MagmaMaxGPUs][10];
float *dx[MagmaMaxGPUs], *dy[MagmaMaxGPUs], *hwork;
float *dwork2[MagmaMaxGPUs];
*info = 0;
int upper = (uplo == MagmaUpper);
lquery = (lwork == -1);
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
} else if (lwork < nb*n && ! lquery) {
*info = -9;
} else if ( num_streams > 2 ) {
*info = 2; // TODO fix
}
/* Determine the block size. */
ldwork = lddwork = n;
lwkopt = n * nb;
if (*info == 0) {
work[0] = MAGMA_S_MAKE( lwkopt, 0 );
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
/* Quick return if possible */
if (n == 0) {
work[0] = c_one;
return *info;
}
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
float *dA[MagmaMaxGPUs];
float *dwork[MagmaMaxGPUs];
float times[11];
for( did=0; did < 11; did++ )
times[did] = 0;
//#define PROFILE_SY2RK
#ifdef PROFILE_SY2RK
magma_event_t start, stop;
float etime;
magma_setdevice(0);
magma_event_create( &start );
magma_event_create( &stop );
#endif
ldda = lda;
ln = ((nb*(1+n/(nb*num_gpus))+31)/32)*32;
ldwork2 = (1+ n / nb + (n % nb != 0)) * ldda;
for( did=0; did < num_gpus; did++ ) {
magma_setdevice(did);
// TODO fix memory leak
if ( MAGMA_SUCCESS != magma_smalloc(&dA[did], ln*ldda+3*lddwork*nb) ||
MAGMA_SUCCESS != magma_smalloc(&dx[did], num_streams*n) ||
MAGMA_SUCCESS != magma_smalloc(&dy[did], num_streams*n) ||
MAGMA_SUCCESS != magma_smalloc(&dwork2[did], ldwork2 ) ) {
for( i=0; i < did; i++ ) {
magma_setdevice(i);
magma_free(dA[i]);
magma_free(dx[i]);
magma_free(dy[i]);
}
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dwork[did] = dA[did] + ln*ldda;
for( kk=0; kk < num_streams; kk++ )
magma_queue_create(&stream[did][kk]);
}
magma_setdevice(0);
// TODO fix memory leak dwork2
if ( MAGMA_SUCCESS != magma_smalloc_pinned( &hwork, num_streams*num_gpus*n ) ) {
for( i=0; i < num_gpus; i++ ) {
magma_setdevice(i);
magma_free(dA[i]);
magma_free(dx[i]);
magma_free(dy[i]);
}
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
if (n < 2048)
nx = n;
else
nx = 512;
if (upper) {
/* Copy the matrix to the GPU */
if (1 <= n-nx) {
magma_shtodhe(num_gpus, uplo, n, nb, A, lda, dA, ldda, stream, &iinfo );
}
/* Reduce the upper triangle of A.
Columns 1:kk are handled by the unblocked method. */
for (i = nb*((n-1)/nb); i >= nx; i -= nb) {
ib = min(nb, n-i);
ii = nb*(i/(nb*num_gpus));
did = (i/nb)%num_gpus;
/* wait for the next panel */
if (i != nb*((n-1)/nb)) {
magma_setdevice(did);
magma_queue_sync(stream[did][0]);
}
magma_slatrd_mgpu(num_gpus, uplo, n, i+ib, ib, nb,
A(0, 0), lda, e, tau,
work, ldwork,
dA, ldda, 0,
dwork, i+ib,
dwork2, ldwork2,
1, dx, dy, hwork,
stream, times);
magma_ssyr2k_mgpu(num_gpus, MagmaUpper, MagmaNoTrans, nb, i, ib,
c_neg_one, dwork, i+ib, 0,
d_one, dA, ldda, 0,
num_streams, stream);
/* get the next panel */
if (i-nb >= nx ) {
ib = min(nb, n-(i-nb));
ii = nb*((i-nb)/(nb*num_gpus));
did = ((i-nb)/nb)%num_gpus;
magma_setdevice(did);
magma_sgetmatrix_async( (i-nb)+ib, ib,
dA(did, 0, ii), ldda,
A(0, i-nb), lda,
stream[did][0] );
}
/* Copy superdiagonal elements back into A, and diagonal
elements into D */
for (j = i; j < i+ib; ++j) {
if ( j > 0 ) {
*A(j-1,j) = MAGMA_S_MAKE( e[j - 1], 0 );
}
d[j] = MAGMA_S_REAL( *A(j, j) );
}
} /* end of for i=... */
if ( nx > 0 ) {
if (1 <= n-nx) { /* else A is already on CPU */
for (i=0; i < nx; i += nb) {
ib = min(nb, n-i);
ii = nb*(i/(nb*num_gpus));
did = (i/nb)%num_gpus;
magma_setdevice(did);
magma_sgetmatrix_async( nx, ib,
dA(did, 0, ii), ldda,
A(0, i), lda,
stream[did][0] );
}
}
for( did=0; did < num_gpus; did++ ) {
magma_setdevice(did);
magma_queue_sync(stream[did][0]);
}
/* Use unblocked code to reduce the last or only block */
lapackf77_ssytd2(uplo_, &nx, A(0, 0), &lda, d, e, tau, &iinfo);
}
}
else {
trace_init( 1, num_gpus, num_streams, (CUstream_st**)stream );
/* Copy the matrix to the GPU */
if (1 <= n-nx) {
magma_shtodhe(num_gpus, uplo, n, nb, A, lda, dA, ldda, stream, &iinfo );
}
/* Reduce the lower triangle of A */
for (i = 0; i < n-nx; i += nb) {
ib = min(nb, n-i);
ii = nb*(i/(nb*num_gpus));
did = (i/nb)%num_gpus;
/* Reduce columns i:i+ib-1 to tridiagonal form and form the
matrix W which is needed to update the unreduced part of
the matrix */
/* Get the current panel (no need for the 1st iteration) */
if (i != 0) {
magma_setdevice(did);
trace_gpu_start( did, 0, "comm", "get" );
magma_sgetmatrix_async( n-i, ib,
dA(did, i, ii), ldda,
A(i,i), lda,
stream[did][0] );
trace_gpu_end( did, 0 );
magma_queue_sync(stream[did][0]);
magma_setdevice(0);
}
magma_slatrd_mgpu(num_gpus, uplo, n, n-i, ib, nb,
A(i, i), lda, &e[i],
&tau[i], work, ldwork,
dA, ldda, i,
dwork, (n-i),
dwork2, ldwork2,
1, dx, dy, hwork,
stream, times );
#ifdef PROFILE_SY2RK
magma_setdevice(0);
if ( i > 0 ) {
cudaEventElapsedTime(&etime, start, stop);
up_time += (etime/1000.0);
}
magma_event_record(start, 0);
#endif
magma_ssyr2k_mgpu(num_gpus, MagmaLower, MagmaNoTrans, nb, n-i-ib, ib,
c_neg_one, dwork, n-i, ib,
d_one, dA, ldda, i+ib, num_streams, stream);
#ifdef PROFILE_SY2RK
magma_setdevice(0);
magma_event_record(stop, 0);
#endif
/* Copy subdiagonal elements back into A, and diagonal
elements into D */
for (j = i; j < i+ib; ++j) {
if ( j+1 < n ) {
*A(j+1,j) = MAGMA_S_MAKE( e[j], 0 );
}
d[j] = MAGMA_S_REAL( *A(j, j) );
}
} /* for i=... */
/* Use unblocked code to reduce the last or only block */
if ( i < n ) {
iii = i;
i_n = n-i;
if ( i > 0 ) {
for (; i < n; i += nb) {
ib = min(nb, n-i);
ii = nb*(i/(nb*num_gpus));
did = (i/nb)%num_gpus;
magma_setdevice(did);
magma_sgetmatrix_async( i_n, ib,
dA(did, iii, ii), ldda,
A(iii, i), lda,
stream[did][0] );
}
for( did=0; did < num_gpus; did++ ) {
magma_setdevice(did);
magma_queue_sync(stream[did][0]);
}
}
lapackf77_ssytrd(uplo_, &i_n, A(iii, iii), &lda, &d[iii], &e[iii],
&tau[iii], work, &lwork, &iinfo);
}
}
#ifdef PROFILE_SY2RK
magma_setdevice(0);
if ( n > nx ) {
cudaEventElapsedTime(&etime, start, stop);
up_time += (etime/1000.0);
}
magma_event_destroy( start );
magma_event_destroy( stop );
#endif
trace_finalize( "ssytrd.svg", "trace.css" );
for( did=0; did < num_gpus; did++ ) {
magma_setdevice(did);
for( kk=0; kk < num_streams; kk++ )
magma_queue_sync(stream[did][kk]);
for( kk=0; kk < num_streams; kk++ )
magma_queue_destroy(stream[did][kk]);
magma_free(dA[did]);
magma_free(dx[did]);
magma_free(dy[did]);
magma_free(dwork2[did]);
}
magma_free_pinned(hwork);
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
work[0] = MAGMA_S_MAKE( lwkopt, 0 );
#ifdef PROFILE_SY2RK
printf( " n=%d nb=%d\n", n, nb );
printf( " Time in SLARFG: %.2e seconds\n", times[0] );
//printf( " Time in SSYMV : %.2e seconds\n", mv_time );
printf( " Time in SSYR2K: %.2e seconds\n", up_time );
#endif
return *info;
} /* magma_ssytrd */
/**
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 */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing sormbr
*/
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float Cnorm, error, dwork[1];
float c_neg_one = MAGMA_S_NEG_ONE;
magma_int_t ione = 1;
magma_int_t m, n, k, mi, ni, mm, nn, nq, size, info;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t nb, ldc, lda, lwork, lwork_max;
float *C, *R, *A, *work, *tau, *tauq, *taup;
float *d, *e;
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
// need slightly looser bound (60*eps instead of 30*eps) for some tests
opts.tolerance = max( 60., opts.tolerance );
float tol = opts.tolerance * lapackf77_slamch("E");
// test all combinations of input parameters
magma_vect_t vect [] = { MagmaQ, MagmaP };
magma_side_t side [] = { MagmaLeft, MagmaRight };
magma_trans_t trans[] = { MagmaTrans, MagmaNoTrans };
printf("%% M N K vect side trans CPU Gflop/s (sec) GPU Gflop/s (sec) ||R||_F / ||QC||_F\n");
printf("%%==============================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int ivect = 0; ivect < 2; ++ivect ) {
for( int iside = 0; iside < 2; ++iside ) {
for( int itran = 0; itran < 2; ++itran ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
m = opts.msize[itest];
n = opts.nsize[itest];
k = opts.ksize[itest];
nb = magma_get_sgebrd_nb( m, n );
ldc = m;
// A is nq x k (vect=Q) or k x nq (vect=P)
// where nq=m (left) or nq=n (right)
nq = (side[iside] == MagmaLeft ? m : n );
mm = (vect[ivect] == MagmaQ ? nq : k );
nn = (vect[ivect] == MagmaQ ? k : nq);
lda = mm;
// MBR calls either MQR or MLQ in various ways
if ( vect[ivect] == MagmaQ ) {
if ( nq >= k ) {
gflops = FLOPS_SORMQR( m, n, k, side[iside] ) / 1e9;
}
else {
if ( side[iside] == MagmaLeft ) {
mi = m - 1;
ni = n;
}
else {
mi = m;
ni = n - 1;
}
gflops = FLOPS_SORMQR( mi, ni, nq-1, side[iside] ) / 1e9;
}
}
else {
if ( nq > k ) {
gflops = FLOPS_SORMLQ( m, n, k, side[iside] ) / 1e9;
}
else {
if ( side[iside] == MagmaLeft ) {
mi = m - 1;
ni = n;
}
else {
mi = m;
ni = n - 1;
}
gflops = FLOPS_SORMLQ( mi, ni, nq-1, side[iside] ) / 1e9;
}
}
// workspace for gebrd is (mm + nn)*nb
// workspace for unmbr is m*nb or n*nb, depending on side
lwork_max = max( (mm + nn)*nb, max( m*nb, n*nb ));
// this rounds it up slightly if needed to agree with lwork query below
lwork_max = int( real( magma_smake_lwork( lwork_max )));
TESTING_MALLOC_CPU( C, float, ldc*n );
TESTING_MALLOC_CPU( R, float, ldc*n );
TESTING_MALLOC_CPU( A, float, lda*nn );
TESTING_MALLOC_CPU( work, float, lwork_max );
TESTING_MALLOC_CPU( d, float, min(mm,nn) );
TESTING_MALLOC_CPU( e, float, min(mm,nn) );
TESTING_MALLOC_CPU( tauq, float, min(mm,nn) );
TESTING_MALLOC_CPU( taup, float, min(mm,nn) );
// C is full, m x n
size = ldc*n;
lapackf77_slarnv( &ione, ISEED, &size, C );
lapackf77_slacpy( "Full", &m, &n, C, &ldc, R, &ldc );
size = lda*nn;
lapackf77_slarnv( &ione, ISEED, &size, A );
// compute BRD factorization to get Householder vectors in A, tauq, taup
//lapackf77_sgebrd( &mm, &nn, A, &lda, d, e, tauq, taup, work, &lwork_max, &info );
magma_sgebrd( mm, nn, A, lda, d, e, tauq, taup, work, lwork_max, &info );
if (info != 0) {
printf("magma_sgebrd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
if ( vect[ivect] == MagmaQ ) {
tau = tauq;
} else {
tau = taup;
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_sormbr( lapack_vect_const( vect[ivect] ),
lapack_side_const( side[iside] ),
lapack_trans_const( trans[itran] ),
&m, &n, &k,
A, &lda, tau, C, &ldc, work, &lwork_max, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0) {
printf("lapackf77_sormbr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
// query for workspace size
lwork = -1;
magma_sormbr( vect[ivect], side[iside], trans[itran],
m, n, k,
A, lda, tau, R, ldc, work, lwork, &info );
if (info != 0) {
printf("magma_sormbr (lwork query) returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
lwork = (magma_int_t) MAGMA_S_REAL( work[0] );
if ( lwork < 0 || lwork > lwork_max ) {
printf("Warning: optimal lwork %d > allocated lwork_max %d\n", (int) lwork, (int) lwork_max );
lwork = lwork_max;
}
gpu_time = magma_wtime();
magma_sormbr( vect[ivect], side[iside], trans[itran],
m, n, k,
A, lda, tau, R, ldc, work, lwork, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0) {
printf("magma_sormbr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
compute relative error |QC_magma - QC_lapack| / |QC_lapack|
=================================================================== */
size = ldc*n;
blasf77_saxpy( &size, &c_neg_one, C, &ione, R, &ione );
Cnorm = lapackf77_slange( "Fro", &m, &n, C, &ldc, dwork );
error = lapackf77_slange( "Fro", &m, &n, R, &ldc, dwork ) / (magma_ssqrt(m*n) * Cnorm);
printf( "%5d %5d %5d %c %4c %5c %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) m, (int) n, (int) k,
lapacke_vect_const( vect[ivect] ),
lapacke_side_const( side[iside] ),
lapacke_trans_const( trans[itran] ),
cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed") );
status += ! (error < tol);
TESTING_FREE_CPU( C );
TESTING_FREE_CPU( R );
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( work );
TESTING_FREE_CPU( d );
TESTING_FREE_CPU( e );
TESTING_FREE_CPU( taup );
TESTING_FREE_CPU( tauq );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}}} // end ivect, iside, itran
printf( "\n" );
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
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 */
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing ssygvdx
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gpu_time /*cpu_time*/;
float *h_A, *h_R, *h_B, *h_S, *h_work;
float *w1, *w2, vl=0, vu=0;
float result[2] = {0};
magma_int_t *iwork;
magma_int_t N, n2, info, il, iu, m1, m2, nb, lwork, liwork;
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
#if defined(PRECISION_z) || defined(PRECISION_c)
float *rwork;
magma_int_t lrwork;
#endif
//float d_one = 1.;
//float d_ten = 10.;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
float tolulp = opts.tolerance * lapackf77_slamch("P");
if ( opts.check && opts.jobz == MagmaNoVec ) {
fprintf( stderr, "checking results requires vectors; setting jobz=V (option -JV)\n" );
opts.jobz = MagmaVec;
}
printf(" N M GPU Time (sec)\n");
printf("============================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[i];
n2 = N*N;
nb = magma_get_ssytrd_nb(N);
#if defined(PRECISION_z) || defined(PRECISION_c)
lwork = 2*N*nb + N*N;
lrwork = 1 + 5*N +2*N*N;
#else
lwork = 1 + 6*N*nb + 2* N*N;
#endif
liwork = 3 + 5*N;
if ( opts.fraction == 0 ) {
il = N / 10;
iu = N / 5+il;
}
else {
il = 1;
iu = (int) (opts.fraction*N);
if (iu < 1) iu = 1;
}
TESTING_MALLOC( h_A, float, n2 );
TESTING_MALLOC( h_B, float, n2 );
TESTING_MALLOC( w1, float, N );
TESTING_MALLOC( w2, float, N );
TESTING_MALLOC( iwork, magma_int_t, liwork );
TESTING_HOSTALLOC( h_R, float, n2 );
TESTING_HOSTALLOC( h_S, float, n2 );
TESTING_HOSTALLOC( h_work, float, lwork );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_HOSTALLOC( rwork, float, lrwork);
#endif
/* Initialize the matrix */
lapackf77_slarnv( &ione, ISEED, &n2, h_A );
lapackf77_slarnv( &ione, ISEED, &n2, h_B );
/* increase the diagonal */
for(int i=0; i<N; i++) {
MAGMA_S_SET2REAL( h_B[i*N+i], ( MAGMA_S_REAL(h_B[i*N+i]) + 1.*N ) );
MAGMA_S_SET2REAL( h_A[i*N+i], MAGMA_S_REAL(h_A[i*N+i]) );
}
// ==================================================================
// Warmup using MAGMA
// ==================================================================
if(opts.warmup){
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
magma_ssygvdx( opts.itype, opts.jobz, 'I', opts.uplo,
N, h_R, N, h_S, N, vl, vu, il, iu, &m1, w1,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info );
if (info != 0)
printf("magma_ssygvdx returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
gpu_time = magma_wtime();
magma_ssygvdx( opts.itype, opts.jobz, 'I', opts.uplo,
N, h_R, N, h_S, N, vl, vu, il, iu, &m1, w1,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info );
gpu_time = magma_wtime() - gpu_time;
if (info != 0)
printf("magma_ssygvdx returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.check ) {
/* =====================================================================
Check the results following the LAPACK's [zc]hegvdx routine.
A x = lambda B x is solved
and the following 3 tests computed:
(1) | A Z - B Z D | / ( |A||Z| N ) (itype = 1)
| A B Z - Z D | / ( |A||Z| N ) (itype = 2)
| B A Z - Z D | / ( |A||Z| N ) (itype = 3)
(2) | S(with V) - S(w/o V) | / | S |
=================================================================== */
#if defined(PRECISION_d) || defined(PRECISION_s)
float *rwork = h_work + N*N;
#endif
float temp1, temp2;
result[0] = 1.;
result[0] /= lapackf77_slansy("1", &opts.uplo, &N, h_A, &N, rwork);
result[0] /= lapackf77_slange("1", &N, &m1, h_R, &N, rwork);
if (opts.itype == 1) {
blasf77_ssymm("L", &opts.uplo, &N, &m1, &c_one, h_A, &N, h_R, &N, &c_zero, h_work, &N);
for(int i=0; i < m1; ++i)
blasf77_sscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_ssymm("L", &opts.uplo, &N, &m1, &c_neg_one, h_B, &N, h_R, &N, &c_one, h_work, &N);
result[0] *= lapackf77_slange("1", &N, &m1, h_work, &N, rwork)/N;
}
else if (opts.itype == 2) {
blasf77_ssymm("L", &opts.uplo, &N, &m1, &c_one, h_B, &N, h_R, &N, &c_zero, h_work, &N);
for(int i=0; i < m1; ++i)
blasf77_sscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_ssymm("L", &opts.uplo, &N, &m1, &c_one, h_A, &N, h_work, &N, &c_neg_one, h_R, &N);
result[0] *= lapackf77_slange("1", &N, &m1, h_R, &N, rwork)/N;
}
else if (opts.itype == 3) {
blasf77_ssymm("L", &opts.uplo, &N, &m1, &c_one, h_A, &N, h_R, &N, &c_zero, h_work, &N);
for(int i=0; i < m1; ++i)
blasf77_sscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_ssymm("L", &opts.uplo, &N, &m1, &c_one, h_B, &N, h_work, &N, &c_neg_one, h_R, &N);
result[0] *= lapackf77_slange("1", &N, &m1, h_R, &N, rwork)/N;
}
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
magma_ssygvdx( opts.itype, 'N', 'I', opts.uplo,
N, h_R, N, h_S, N, vl, vu, il, iu, &m2, w2,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info );
if (info != 0)
printf("magma_ssygvdx returned error %d: %s.\n",
(int) info, magma_strerror( info ));
temp1 = temp2 = 0;
for(int j=0; j < m2; j++) {
temp1 = max(temp1, absv(w1[j]));
temp1 = max(temp1, absv(w2[j]));
temp2 = max(temp2, absv(w1[j]-w2[j]));
}
result[1] = temp2 / (((float)m2)*temp1);
}
/* =====================================================================
Print execution time
=================================================================== */
printf("%5d %5d %7.2f\n",
(int) N, (int) m1, gpu_time);
if ( opts.check ) {
printf("Testing the eigenvalues and eigenvectors for correctness:\n");
if (opts.itype==1)
printf("(1) | A Z - B Z D | / (|A| |Z| N) = %8.2e%s\n", result[0], (result[0] < tol ? "" : " failed"));
else if (opts.itype==2)
printf("(1) | A B Z - Z D | / (|A| |Z| N) = %8.2e%s\n", result[0], (result[0] < tol ? "" : " failed"));
else if (opts.itype==3)
printf("(1) | B A Z - Z D | / (|A| |Z| N) = %8.2e%s\n", result[0], (result[0] < tol ? "" : " failed"));
printf( "(2) | D(w/ Z) - D(w/o Z) | / |D| = %8.2e%s\n\n", result[1], (result[1] < tolulp ? "" : " failed"));
}
TESTING_FREE( h_A );
TESTING_FREE( h_B );
TESTING_FREE( w1 );
TESTING_FREE( w2 );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_HOSTFREE( rwork);
#endif
TESTING_FREE( iwork );
TESTING_HOSTFREE( h_work );
TESTING_HOSTFREE( h_R );
TESTING_HOSTFREE( h_S );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
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 */
