/* ////////////////////////////////////////////////////////////////////////////
-- Testing ssygvdx
*/
int main( int argc, char** argv)
{
TESTING_INIT_MGPU();
real_Double_t mgpu_time;
float *h_A, *h_Ainit, *h_B, *h_Binit, *h_work;
#if defined(PRECISION_z) || defined(PRECISION_c)
float *rwork;
magma_int_t lrwork;
#endif
float *w1, result=0;
magma_int_t *iwork;
magma_int_t N, n2, info, lwork, liwork;
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
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");
magma_range_t range = MagmaRangeAll;
if (opts.fraction != 1)
range = MagmaRangeI;
if ( opts.check && opts.jobz == MagmaNoVec ) {
fprintf( stderr, "checking results requires vectors; setting jobz=V (option -JV)\n" );
opts.jobz = MagmaVec;
}
printf("using: ngpu = %d, itype = %d, jobz = %s, range = %s, uplo = %s, opts.check = %d, fraction = %6.4f\n",
(int) opts.ngpu, (int) opts.itype,
lapack_vec_const(opts.jobz), lapack_range_const(range), lapack_uplo_const(opts.uplo),
(int) opts.check, opts.fraction);
printf(" N M ngpu MGPU Time (sec)\n");
printf("====================================\n");
magma_int_t threads = magma_get_parallel_numthreads();
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
n2 = N*N;
#if defined(PRECISION_z) || defined(PRECISION_c)
lwork = magma_sbulge_get_lq2(N, threads) + 2*N + N*N;
lrwork = 1 + 5*N +2*N*N;
#else
lwork = magma_sbulge_get_lq2(N, threads) + 1 + 6*N + 2*N*N;
#endif
liwork = 3 + 5*N;
//magma_int_t NB = 96;//magma_bulge_get_nb(N);
//magma_int_t sizvblg = magma_sbulge_get_lq2(N, threads);
//magma_int_t siz = max(sizvblg,n2)+2*(N*NB+N)+24*N;
/* Allocate host memory for the matrix */
TESTING_MALLOC_PIN( h_A, float, n2 );
TESTING_MALLOC_PIN( h_B, float, n2 );
TESTING_MALLOC_PIN( h_work, float, lwork );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_MALLOC_PIN( rwork, float, lrwork);
#endif
TESTING_MALLOC_CPU( w1, float, N );
TESTING_MALLOC_CPU( iwork, magma_int_t, liwork);
/* Initialize the matrix */
lapackf77_slarnv( &ione, ISEED, &n2, h_A );
lapackf77_slarnv( &ione, ISEED, &n2, h_B );
magma_smake_hpd( N, h_B, N );
magma_smake_symmetric( N, h_A, N );
if ( opts.warmup || opts.check ) {
TESTING_MALLOC_CPU( h_Ainit, float, n2 );
TESTING_MALLOC_CPU( h_Binit, float, n2 );
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_Ainit, &N );
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_Binit, &N );
}
magma_int_t m1 = 0;
float vl = 0;
float vu = 0;
magma_int_t il = 0;
magma_int_t iu = 0;
if (range == MagmaRangeI) {
il = 1;
iu = (int) (opts.fraction*N);
}
if ( opts.warmup ) {
// ==================================================================
// Warmup using MAGMA. I prefer to use smalltest to warmup A-
// ==================================================================
magma_ssygvdx_2stage_m(opts.ngpu, opts.itype, opts.jobz, range, opts.uplo,
N, h_A, N, h_B, N, vl, vu, il, iu, &m1, w1,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info);
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_Ainit, &N, h_A, &N );
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_Binit, &N, h_B, &N );
}
// ===================================================================
// Performs operation using MAGMA
// ===================================================================
mgpu_time = magma_wtime();
magma_ssygvdx_2stage_m(opts.ngpu, opts.itype, opts.jobz, range, opts.uplo,
N, h_A, N, h_B, N, vl, vu, il, iu, &m1, w1,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info);
mgpu_time = magma_wtime() - mgpu_time;
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)
// ===================================================================
#if defined(PRECISION_d) || defined(PRECISION_s)
float *rwork = h_work + N*N;
#endif
result = 1.;
result /= lapackf77_slansy("1", lapack_uplo_const(opts.uplo), &N, h_Ainit, &N, rwork);
result /= lapackf77_slange("1", &N , &m1, h_A, &N, rwork);
if (opts.itype == 1) {
blasf77_ssymm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_Ainit, &N, h_A, &N, &c_zero, h_work, &N);
for(int i=0; i<m1; ++i)
blasf77_sscal(&N, &w1[i], &h_A[i*N], &ione);
blasf77_ssymm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_neg_one, h_Binit, &N, h_A, &N, &c_one, h_work, &N);
result *= lapackf77_slange("1", &N, &m1, h_work, &N, rwork)/N;
}
else if (opts.itype == 2) {
blasf77_ssymm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_Binit, &N, h_A, &N, &c_zero, h_work, &N);
for(int i=0; i<m1; ++i)
blasf77_sscal(&N, &w1[i], &h_A[i*N], &ione);
blasf77_ssymm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_Ainit, &N, h_work, &N, &c_neg_one, h_A, &N);
result *= lapackf77_slange("1", &N, &m1, h_A, &N, rwork)/N;
}
else if (opts.itype == 3) {
blasf77_ssymm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_Ainit, &N, h_A, &N, &c_zero, h_work, &N);
for(int i=0; i<m1; ++i)
blasf77_sscal(&N, &w1[i], &h_A[i*N], &ione);
blasf77_ssymm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_Binit, &N, h_work, &N, &c_neg_one, h_A, &N);
result *= lapackf77_slange("1", &N, &m1, h_A, &N, rwork)/N;
}
}
// ===================================================================
// Print execution time
// ===================================================================
printf("%5d %5d %4d %7.2f\n",
(int) N, (int) m1, (int) opts.ngpu, mgpu_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, (result < tol ? "ok" : "failed") );
}
else if (opts.itype==2) {
printf("(1) | A B Z - Z D | / (|A| |Z| N) = %8.2e %s\n", result, (result < tol ? "ok" : "failed") );
}
else if (opts.itype==3) {
printf("(1) | B A Z - Z D | / (|A| |Z| N) = %8.2e %s\n", result, (result < tol ? "ok" : "failed") );
}
printf("\n");
status += ! (result < tol);
}
TESTING_FREE_PIN( h_A );
TESTING_FREE_PIN( h_B );
TESTING_FREE_PIN( h_work );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_FREE_PIN( rwork );
#endif
TESTING_FREE_CPU( w1 );
TESTING_FREE_CPU( iwork );
if ( opts.warmup || opts.check ) {
TESTING_FREE_CPU( h_Ainit );
TESTING_FREE_CPU( h_Binit );
}
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
/* Shutdown */
TESTING_FINALIZE_MGPU();
return status;
}
int main(int argc, char **argv)
{
TESTING_INIT();
const float c_neg_one = MAGMA_S_NEG_ONE;
const magma_int_t ione = 1;
real_Double_t atomics_perf, atomics_time;
real_Double_t gflops, magma_perf, magma_time, cublas_perf, cublas_time, cpu_perf, cpu_time;
float magma_error, atomics_error, cublas_error, work[1];
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t N, lda, ldda, sizeA, sizeX, sizeY, blocks, ldwork;
magma_int_t incx = 1;
magma_int_t incy = 1;
magma_int_t nb = 64;
float alpha = MAGMA_S_MAKE( 1.5, -2.3 );
float beta = MAGMA_S_MAKE( -0.6, 0.8 );
float *A, *X, *Y, *Yatomics, *Ycublas, *Ymagma;
magmaFloat_ptr dA, dX, dY, dwork;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
printf("uplo = %s\n", lapack_uplo_const(opts.uplo) );
printf(" N MAGMA Gflop/s (ms) Atomics Gflop/s CUBLAS Gflop/s CPU Gflop/s MAGMA error Atomics CUBLAS\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;
ldda = ((N + 31)/32)*32;
sizeA = N*lda;
sizeX = N*incx;
sizeY = N*incy;
gflops = FLOPS_SSYMV( N ) / 1e9;
TESTING_MALLOC_CPU( A, float, sizeA );
TESTING_MALLOC_CPU( X, float, sizeX );
TESTING_MALLOC_CPU( Y, float, sizeY );
TESTING_MALLOC_CPU( Yatomics, float, sizeY );
TESTING_MALLOC_CPU( Ycublas, float, sizeY );
TESTING_MALLOC_CPU( Ymagma, float, sizeY );
TESTING_MALLOC_DEV( dA, float, ldda*N );
TESTING_MALLOC_DEV( dX, float, sizeX );
TESTING_MALLOC_DEV( dY, float, sizeY );
blocks = (N + nb - 1) / nb;
ldwork = ldda*blocks;
TESTING_MALLOC_DEV( dwork, float, ldwork );
magmablas_slaset( MagmaFull, ldwork, 1, MAGMA_S_NAN, MAGMA_S_NAN, dwork, ldwork );
magmablas_slaset( MagmaFull, ldda, N, MAGMA_S_NAN, MAGMA_S_NAN, dA, ldda );
/* Initialize the matrix */
lapackf77_slarnv( &ione, ISEED, &sizeA, A );
magma_smake_symmetric( N, A, lda );
// should not use data from the opposite triangle -- fill with NAN to check
magma_int_t N1 = N-1;
if ( opts.uplo == MagmaUpper ) {
lapackf77_slaset( "Lower", &N1, &N1, &MAGMA_S_NAN, &MAGMA_S_NAN, &A[1], &lda );
}
else {
lapackf77_slaset( "Upper", &N1, &N1, &MAGMA_S_NAN, &MAGMA_S_NAN, &A[lda], &lda );
}
lapackf77_slarnv( &ione, ISEED, &sizeX, X );
lapackf77_slarnv( &ione, ISEED, &sizeY, Y );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_ssetmatrix( N, N, A, lda, dA, ldda );
magma_ssetvector( N, X, incx, dX, incx );
magma_ssetvector( N, Y, incy, dY, incy );
cublas_time = magma_sync_wtime( 0 );
cublasSsymv( opts.handle, cublas_uplo_const(opts.uplo),
N, &alpha, dA, ldda, dX, incx, &beta, dY, incy );
cublas_time = magma_sync_wtime( 0 ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_sgetvector( N, dY, incy, Ycublas, incy );
/* =====================================================================
Performs operation using CUBLAS - using atomics
=================================================================== */
cublasSetAtomicsMode( opts.handle, CUBLAS_ATOMICS_ALLOWED );
magma_ssetvector( N, Y, incy, dY, incy );
atomics_time = magma_sync_wtime( 0 );
cublasSsymv( opts.handle, cublas_uplo_const(opts.uplo),
N, &alpha, dA, ldda, dX, incx, &beta, dY, incy );
atomics_time = magma_sync_wtime( 0 ) - atomics_time;
atomics_perf = gflops / atomics_time;
magma_sgetvector( N, dY, incy, Yatomics, incy );
cublasSetAtomicsMode( opts.handle, CUBLAS_ATOMICS_NOT_ALLOWED );
/* =====================================================================
Performs operation using MAGMABLAS
=================================================================== */
magma_ssetvector( N, Y, incy, dY, incy );
magma_time = magma_sync_wtime( 0 );
if ( opts.version == 1 ) {
magmablas_ssymv_work( opts.uplo, N, alpha, dA, ldda, dX, incx, beta, dY, incy, dwork, ldwork, opts.queue );
}
else {
// non-work interface (has added overhead)
magmablas_ssymv( opts.uplo, N, alpha, dA, ldda, dX, incx, beta, dY, incy );
}
magma_time = magma_sync_wtime( 0 ) - magma_time;
magma_perf = gflops / magma_time;
magma_sgetvector( N, dY, incy, Ymagma, incy );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
cpu_time = magma_wtime();
blasf77_ssymv( lapack_uplo_const(opts.uplo), &N, &alpha, A, &lda, X, &incx, &beta, Y, &incy );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/* =====================================================================
Check the result
=================================================================== */
blasf77_saxpy( &N, &c_neg_one, Y, &incy, Ymagma, &incy );
magma_error = lapackf77_slange( "M", &N, &ione, Ymagma, &N, work ) / N;
blasf77_saxpy( &N, &c_neg_one, Y, &incy, Ycublas, &incy );
cublas_error = lapackf77_slange( "M", &N, &ione, Ycublas, &N, work ) / N;
blasf77_saxpy( &N, &c_neg_one, Y, &incy, Yatomics, &incy );
atomics_error = lapackf77_slange( "M", &N, &ione, Yatomics, &N, work ) / N;
bool ok = (magma_error < tol && cublas_error < tol && atomics_error < tol);
status += ! ok;
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %8.2e %s\n",
(int) N,
magma_perf, 1000.*magma_time,
atomics_perf, 1000.*atomics_time,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
magma_error, cublas_error, atomics_error,
(ok ? "ok" : "failed"));
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( X );
TESTING_FREE_CPU( Y );
TESTING_FREE_CPU( Ycublas );
TESTING_FREE_CPU( Yatomics );
TESTING_FREE_CPU( Ymagma );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dX );
TESTING_FREE_DEV( dY );
TESTING_FREE_DEV( dwork );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing ssygvd
*/
int main( int argc, char** argv)
{
TESTING_INIT();
float *h_A, *h_Ainit, *h_B, *h_Binit, *h_work;
#ifdef COMPLEX
float *rwork;
#endif
float *w1, *w2, result[2]={0, 0};
magma_int_t *iwork;
real_Double_t mgpu_time, gpu_time, cpu_time;
/* Matrix size */
magma_int_t N, n2, nb;
magma_int_t info;
magma_int_t ione = 1;
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
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");
float tolulp = opts.tolerance * lapackf77_slamch("P");
// checking NoVec requires LAPACK
opts.lapack |= (opts.check && opts.jobz == MagmaNoVec);
printf("using: ngpu = %d, itype = %d, jobz = %s, uplo = %s, check = %d\n",
(int) opts.ngpu, (int) opts.itype,
lapack_vec_const(opts.jobz), lapack_uplo_const(opts.uplo), (int) opts.check);
printf(" N CPU Time (sec) GPU Time (sec) MGPU Time (sec)\n");
printf("=========================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
// TODO define lda
N = opts.nsize[itest];
n2 = N*N;
nb = magma_get_ssytrd_nb(N);
#ifdef COMPLEX
magma_int_t lwork = max( N + N*nb, 2*N + N*N );
magma_int_t lrwork = 1 + 5*N +2*N*N;
#else
magma_int_t lwork = max( 2*N + N*nb, 1 + 6*N + 2*N*N );
#endif
magma_int_t liwork = 3 + 5*N;
TESTING_MALLOC_PIN( h_A, float, n2 );
TESTING_MALLOC_PIN( h_B, float, n2 );
TESTING_MALLOC_PIN( h_work, float, lwork );
#ifdef COMPLEX
TESTING_MALLOC_PIN( rwork, float, lrwork );
#endif
TESTING_MALLOC_CPU( w1, float, N );
TESTING_MALLOC_CPU( w2, float, N );
TESTING_MALLOC_CPU( iwork, magma_int_t, liwork );
/* Initialize the matrix */
lapackf77_slarnv( &ione, ISEED, &n2, h_A );
lapackf77_slarnv( &ione, ISEED, &n2, h_B );
magma_smake_hpd( N, h_B, N );
magma_smake_symmetric( N, h_A, N );
if ( opts.warmup || opts.check ) {
TESTING_MALLOC_CPU( h_Ainit, float, n2 );
TESTING_MALLOC_CPU( h_Binit, float, n2 );
lapackf77_slacpy( MagmaFullStr, &N, &N, h_A, &N, h_Ainit, &N );
lapackf77_slacpy( MagmaFullStr, &N, &N, h_B, &N, h_Binit, &N );
}
if (opts.warmup) {
// ==================================================================
// Warmup using MAGMA.
// ==================================================================
magma_ssygvd_m( opts.ngpu, opts.itype, opts.jobz, opts.uplo,
N, h_A, N, h_B, N, w1,
h_work, lwork,
#ifdef COMPLEX
rwork, lrwork,
#endif
iwork, liwork,
&info);
lapackf77_slacpy( MagmaFullStr, &N, &N, h_Ainit, &N, h_A, &N );
lapackf77_slacpy( MagmaFullStr, &N, &N, h_Binit, &N, h_B, &N );
}
// ===================================================================
// Performs operation using MAGMA
// ===================================================================
mgpu_time = magma_wtime();
magma_ssygvd_m( opts.ngpu, opts.itype, opts.jobz, opts.uplo,
N, h_A, N, h_B, N, w1,
h_work, lwork,
#ifdef COMPLEX
rwork, lrwork,
#endif
iwork, liwork,
&info);
mgpu_time = magma_wtime() - mgpu_time;
if (info != 0)
printf("magma_ssygvd_m returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.check && opts.jobz != MagmaNoVec ) {
/* =====================================================================
Check the results following the LAPACK's [zc]hegvd 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)
=================================================================== */
#ifdef REAL
float *rwork = h_work + N*N;
#endif
result[0] = 1.;
result[0] /= lapackf77_slansy("1", lapack_uplo_const(opts.uplo), &N, h_Ainit, &N, rwork);
result[0] /= lapackf77_slange("1", &N, &N, h_A, &N, rwork);
if (opts.itype == 1) {
blasf77_ssymm("L", lapack_uplo_const(opts.uplo), &N, &N, &c_one, h_Ainit, &N, h_A, &N, &c_zero, h_work, &N);
for(int i=0; i < N; ++i)
blasf77_sscal(&N, &w1[i], &h_A[i*N], &ione);
blasf77_ssymm("L", lapack_uplo_const(opts.uplo), &N, &N, &c_neg_one, h_Binit, &N, h_A, &N, &c_one, h_work, &N);
result[0] *= lapackf77_slange("1", &N, &N, h_work, &N, rwork)/N;
}
else if (opts.itype == 2) {
blasf77_ssymm("L", lapack_uplo_const(opts.uplo), &N, &N, &c_one, h_Binit, &N, h_A, &N, &c_zero, h_work, &N);
for(int i=0; i < N; ++i)
blasf77_sscal(&N, &w1[i], &h_A[i*N], &ione);
blasf77_ssymm("L", lapack_uplo_const(opts.uplo), &N, &N, &c_one, h_Ainit, &N, h_work, &N, &c_neg_one, h_A, &N);
result[0] *= lapackf77_slange("1", &N, &N, h_A, &N, rwork)/N;
}
else if (opts.itype == 3) {
blasf77_ssymm("L", lapack_uplo_const(opts.uplo), &N, &N, &c_one, h_Ainit, &N, h_A, &N, &c_zero, h_work, &N);
for(int i=0; i < N; ++i)
blasf77_sscal(&N, &w1[i], &h_A[i*N], &ione);
blasf77_ssymm("L", lapack_uplo_const(opts.uplo), &N, &N, &c_one, h_Binit, &N, h_work, &N, &c_neg_one, h_A, &N);
result[0] *= lapackf77_slange("1", &N, &N, h_A, &N, rwork)/N;
}
}
if ( opts.lapack ) {
lapackf77_slacpy( MagmaFullStr, &N, &N, h_Ainit, &N, h_A, &N );
lapackf77_slacpy( MagmaFullStr, &N, &N, h_Binit, &N, h_B, &N );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_ssygvd(opts.itype, opts.jobz, opts.uplo,
N, h_A, N, h_B, N, w2,
h_work, lwork,
#ifdef COMPLEX
rwork, lrwork,
#endif
iwork, liwork,
&info);
gpu_time = magma_wtime() - gpu_time;
if (info != 0)
printf("magma_ssygvd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_ssygvd(&opts.itype, lapack_vec_const(opts.jobz), lapack_uplo_const(opts.uplo),
&N, h_Ainit, &N, h_Binit, &N, w2,
h_work, &lwork,
#ifdef COMPLEX
rwork, &lrwork,
#endif
iwork, &liwork,
&info);
cpu_time = magma_wtime() - cpu_time;
if (info != 0)
printf("lapackf77_ssygvd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
float maxw=0, diff=0;
for(int j=0; j < N; j++) {
maxw = max(maxw, fabs(w1[j]));
maxw = max(maxw, fabs(w2[j]));
diff = max(diff, fabs(w1[j] - w2[j]));
}
result[1] = diff / (N*maxw);
/* =====================================================================
Print execution time
=================================================================== */
printf("%5d %7.2f %7.2f %7.2f\n",
(int) N, cpu_time, gpu_time, mgpu_time);
}
else {
printf("%5d --- --- %7.2f\n",
(int) N, mgpu_time);
}
if ( opts.check && opts.jobz != MagmaNoVec ) {
printf("Testing the eigenvalues and eigenvectors for correctness:\n");
if (opts.itype == 1) {
printf(" | A Z - B Z D | / (|A| |Z| N) = %8.2e %s\n", result[0], (result[0] < tol ? "ok" : "failed") );
}
else if (opts.itype == 2) {
printf(" | A B Z - Z D | / (|A| |Z| N) = %8.2e %s\n", result[0], (result[0] < tol ? "ok" : "failed") );
}
else if (opts.itype == 3) {
printf(" | B A Z - Z D | / (|A| |Z| N) = %8.2e %s\n", result[0], (result[0] < tol ? "ok" : "failed") );
}
status += ! (result[0] < tol);
}
if ( opts.lapack ) {
printf( " | D_mgpu - D_lapack | / |D| = %8.2e %s\n\n", result[1], (result[1] < tolulp ? "ok" : "failed") );
status += ! (result[1] < tolulp);
}
/* Memory clean up */
TESTING_FREE_PIN( h_A );
TESTING_FREE_PIN( h_B );
TESTING_FREE_PIN( h_work );
#ifdef COMPLEX
TESTING_FREE_PIN( rwork );
#endif
TESTING_FREE_CPU( w1 );
TESTING_FREE_CPU( w2 );
TESTING_FREE_CPU( iwork );
if ( opts.warmup || opts.check ) {
TESTING_FREE_CPU( h_Ainit );
TESTING_FREE_CPU( h_Binit );
}
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- 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_CPU( h_A, float, n2 );
TESTING_MALLOC_CPU( h_B, float, n2 );
TESTING_MALLOC_CPU( w1, float, N );
TESTING_MALLOC_CPU( w2, float, N );
TESTING_MALLOC_CPU( iwork, magma_int_t, liwork );
TESTING_MALLOC_PIN( h_R, float, n2 );
TESTING_MALLOC_PIN( h_S, float, n2 );
TESTING_MALLOC_PIN( h_work, float, lwork );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_MALLOC_PIN( rwork, float, lrwork);
#endif
/* Initialize the matrix */
lapackf77_slarnv( &ione, ISEED, &n2, h_A );
lapackf77_slarnv( &ione, ISEED, &n2, h_B );
magma_smake_hpd( N, h_B, N );
magma_smake_symmetric( N, h_A, N );
// ==================================================================
// 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_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( w1 );
TESTING_FREE_CPU( w2 );
TESTING_FREE_CPU( iwork );
TESTING_FREE_PIN( h_R );
TESTING_FREE_PIN( h_S );
TESTING_FREE_PIN( h_work );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_FREE_PIN( rwork );
#endif
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing ssygvdx
*/
int main( int argc, char** argv)
{
TESTING_INIT();
/* Constants */
const float c_zero = MAGMA_S_ZERO;
const float c_one = MAGMA_S_ONE;
const float c_neg_one = MAGMA_S_NEG_ONE;
const magma_int_t ione = 1;
/* Local variables */
real_Double_t gpu_time;
float *h_A, *h_R, *h_B, *h_S, *h_work;
#ifdef COMPLEX
float *rwork;
magma_int_t lrwork;
#endif
float *w1, *w2, result[2]={0,0};
magma_int_t *iwork;
magma_int_t N, n2, info, lda, lwork, liwork;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
float tol = opts.tolerance * lapackf77_slamch("E");
float tolulp = opts.tolerance * lapackf77_slamch("P");
magma_range_t range = MagmaRangeAll;
if (opts.fraction != 1)
range = MagmaRangeI;
// pass ngpu = -1 to test multi-GPU code using 1 gpu
magma_int_t abs_ngpu = abs( opts.ngpu );
printf("%% itype = %d, jobz = %s, range = %s, uplo = %s, fraction = %6.4f, ngpu = %d\n",
int(opts.itype), lapack_vec_const(opts.jobz), lapack_range_const(range), lapack_uplo_const(opts.uplo),
opts.fraction, int(abs_ngpu) );
if (opts.itype == 1) {
printf("%% N M GPU Time (sec) |AZ-BZD| |D - D_magma|\n");
}
else if (opts.itype == 2) {
printf("%% N M GPU Time (sec) |ABZ-ZD| |D - D_magma|\n");
}
else if (opts.itype == 3) {
printf("%% N M GPU Time (sec) |BAZ-ZD| |D - D_magma|\n");
}
printf("%%======================================================\n");
magma_int_t threads = magma_get_parallel_numthreads();
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;
// TODO: test vl-vu range
magma_int_t m1 = 0;
float vl = 0;
float vu = 0;
magma_int_t il = 0;
magma_int_t iu = 0;
if (opts.fraction == 0) {
il = max( 1, magma_int_t(0.1*N) );
iu = max( 1, magma_int_t(0.3*N) );
}
else {
il = 1;
iu = max( 1, magma_int_t(opts.fraction*N) );
}
magma_ssyevdx_getworksize(N, threads, (opts.jobz == MagmaVec),
&lwork,
#ifdef COMPLEX
&lrwork,
#endif
&liwork);
/* Allocate host memory for the matrix */
TESTING_MALLOC_CPU( h_A, float, n2 );
TESTING_MALLOC_CPU( h_B, float, n2 );
TESTING_MALLOC_CPU( w1, float, N );
TESTING_MALLOC_CPU( w2, float, N );
TESTING_MALLOC_CPU( iwork, magma_int_t, liwork );
TESTING_MALLOC_PIN( h_R, float, n2 );
TESTING_MALLOC_PIN( h_S, float, n2 );
TESTING_MALLOC_PIN( h_work, float, max( lwork, N*N )); // check needs N*N
#ifdef COMPLEX
TESTING_MALLOC_PIN( rwork, float, lrwork);
#endif
/* Initialize the matrix */
lapackf77_slarnv( &ione, ISEED, &n2, h_A );
lapackf77_slarnv( &ione, ISEED, &n2, h_B );
magma_smake_hpd( N, h_B, lda );
magma_smake_symmetric( N, h_A, lda );
lapackf77_slacpy( MagmaFullStr, &N, &N, h_A, &lda, h_R, &lda );
lapackf77_slacpy( MagmaFullStr, &N, &N, h_B, &lda, h_S, &lda );
// ===================================================================
// Performs operation using MAGMA
// ===================================================================
gpu_time = magma_wtime();
if (opts.ngpu == 1) {
magma_ssygvdx_2stage( opts.itype, opts.jobz, range, opts.uplo,
N, h_R, lda, h_S, lda, vl, vu, il, iu, &m1, w1,
h_work, lwork,
#ifdef COMPLEX
rwork, lrwork,
#endif
iwork, liwork,
&info );
}
else {
magma_ssygvdx_2stage_m( abs_ngpu, opts.itype, opts.jobz, range, opts.uplo,
N, h_R, lda, h_S, lda, vl, vu, il, iu, &m1, w1,
h_work, lwork,
#ifdef COMPLEX
rwork, lrwork,
#endif
iwork, liwork,
&info );
}
gpu_time = magma_wtime() - gpu_time;
if (info != 0) {
printf("magma_ssygvdx_2stage 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) | D(with V, magma) - D(w/o V, lapack) | / | D |
=================================================================== */
#ifdef REAL
float *rwork = h_work + N*N;
#endif
if ( opts.jobz != MagmaNoVec ) {
result[0] = 1.;
result[0] /= safe_lapackf77_slansy("1", lapack_uplo_const(opts.uplo), &N, h_A, &lda, rwork);
result[0] /= lapackf77_slange("1", &N, &m1, h_R, &lda, rwork);
if (opts.itype == 1) {
blasf77_ssymm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_A, &lda, h_R, &lda, &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", lapack_uplo_const(opts.uplo), &N, &m1, &c_neg_one, h_B, &lda, h_R, &lda, &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", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_B, &lda, h_R, &lda, &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", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_A, &lda, h_work, &N, &c_neg_one, h_R, &lda);
result[0] *= lapackf77_slange("1", &N, &m1, h_R, &lda, rwork)/N;
}
else if (opts.itype == 3) {
blasf77_ssymm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_A, &lda, h_R, &lda, &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", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_B, &lda, h_work, &N, &c_neg_one, h_R, &lda);
result[0] *= lapackf77_slange("1", &N, &m1, h_R, &lda, rwork)/N;
}
}
lapackf77_slacpy( MagmaFullStr, &N, &N, h_A, &lda, h_R, &lda );
lapackf77_slacpy( MagmaFullStr, &N, &N, h_B, &lda, h_S, &lda );
lapackf77_ssygvd( &opts.itype, "N", lapack_uplo_const(opts.uplo), &N,
h_R, &lda, h_S, &lda, w2,
h_work, &lwork,
#ifdef COMPLEX
rwork, &lrwork,
#endif
iwork, &liwork,
&info );
if (info != 0) {
printf("lapackf77_ssygvd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
float maxw=0, diff=0;
for (int j=0; j < m1; j++) {
maxw = max(maxw, fabs(w1[j]));
maxw = max(maxw, fabs(w2[j]));
diff = max(diff, fabs(w1[j] - w2[j]));
}
result[1] = diff / (m1*maxw);
}
/* =====================================================================
Print execution time
=================================================================== */
printf("%5d %5d %9.4f ",
(int) N, (int) m1, gpu_time);
if ( opts.check ) {
bool okay = (result[1] < tolulp);
if ( opts.jobz != MagmaNoVec ) {
okay = okay && (result[0] < tol);
printf(" %8.2e", result[0] );
}
else {
printf(" --- ");
}
printf(" %8.2e %s\n", result[1], (okay ? "ok" : "failed"));
status += ! okay;
}
else {
printf(" ---\n");
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( w1 );
TESTING_FREE_CPU( w2 );
TESTING_FREE_CPU( iwork );
TESTING_FREE_PIN( h_R );
TESTING_FREE_PIN( h_S );
TESTING_FREE_PIN( h_work );
#ifdef COMPLEX
TESTING_FREE_PIN( rwork );
#endif
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing ssygst
*/
int main( int argc, char** argv)
{
TESTING_INIT();
// Constants
const float c_neg_one = MAGMA_S_NEG_ONE;
const magma_int_t ione = 1;
// Local variables
real_Double_t gpu_time, cpu_time;
float *h_A, *h_B, *h_R;
magmaFloat_ptr d_A, d_B;
float Anorm, error, work[1];
magma_int_t N, n2, lda, ldda, info;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
float tol = opts.tolerance * lapackf77_slamch("E");
printf("%% uplo = %s\n", lapack_uplo_const(opts.uplo) );
printf("%% itype N CPU time (sec) GPU time (sec) |R| \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;
ldda = magma_roundup( lda, opts.align );
n2 = N*lda;
TESTING_MALLOC_CPU( h_A, float, lda*N );
TESTING_MALLOC_CPU( h_B, float, lda*N );
TESTING_MALLOC_PIN( h_R, float, lda*N );
TESTING_MALLOC_DEV( d_A, float, ldda*N );
TESTING_MALLOC_DEV( d_B, float, ldda*N );
/* ====================================================================
Initialize the matrix
=================================================================== */
lapackf77_slarnv( &ione, ISEED, &n2, h_A );
lapackf77_slarnv( &ione, ISEED, &n2, h_B );
magma_smake_symmetric( N, h_A, lda );
magma_smake_hpd( N, h_B, lda );
magma_spotrf( opts.uplo, N, h_B, lda, &info );
if (info != 0) {
printf("magma_spotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
magma_ssetmatrix( N, N, h_A, lda, d_A, ldda, opts.queue );
magma_ssetmatrix( N, N, h_B, lda, d_B, ldda, opts.queue );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_ssygst_gpu( opts.itype, opts.uplo, N, d_A, ldda, d_B, ldda, &info );
gpu_time = magma_wtime() - gpu_time;
if (info != 0) {
printf("magma_ssygst_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_ssygst( &opts.itype, lapack_uplo_const(opts.uplo),
&N, h_A, &lda, h_B, &lda, &info );
cpu_time = magma_wtime() - cpu_time;
if (info != 0) {
printf("lapackf77_ssygst returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
magma_sgetmatrix( N, N, d_A, ldda, h_R, lda, opts.queue );
blasf77_saxpy( &n2, &c_neg_one, h_A, &ione, h_R, &ione );
Anorm = safe_lapackf77_slansy("f", lapack_uplo_const(opts.uplo), &N, h_A, &lda, work );
error = safe_lapackf77_slansy("f", lapack_uplo_const(opts.uplo), &N, h_R, &lda, work )
/ Anorm;
bool okay = (error < tol);
status += ! okay;
printf("%3d %5d %7.2f %7.2f %8.2e %s\n",
(int) opts.itype, (int) N, cpu_time, gpu_time,
error, (okay ? "ok" : "failed"));
}
else {
printf("%3d %5d --- %7.2f\n",
(int) opts.itype, (int) N, gpu_time );
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
int main(int argc, char **argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, cublas_perf, cublas_time, cpu_perf, cpu_time;
float magma_error, cublas_error, work[1];
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t N, lda, sizeA, sizeX, sizeY, blocks, ldwork;
magma_int_t incx = 1;
magma_int_t incy = 1;
magma_int_t nb = 64;
float c_neg_one = MAGMA_S_NEG_ONE;
float alpha = MAGMA_S_MAKE( 1.5, -2.3 );
float beta = MAGMA_S_MAKE( -0.6, 0.8 );
float *A, *X, *Y, *Ycublas, *Ymagma;
float *dA, *dX, *dY, *dwork;
magma_opts opts;
parse_opts( argc, argv, &opts );
printf(" N MAGMA Gflop/s (ms) CUBLAS Gflop/s (ms) CPU Gflop/s (ms) MAGMA error CUBLAS error\n");
printf("=============================================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[i];
lda = ((N + 31)/32)*32;
sizeA = N*lda;
sizeX = N*incx;
sizeY = N*incy;
gflops = FLOPS_SSYMV( N ) / 1e9;
TESTING_MALLOC_CPU( A, float, sizeA );
TESTING_MALLOC_CPU( X, float, sizeX );
TESTING_MALLOC_CPU( Y, float, sizeY );
TESTING_MALLOC_CPU( Ycublas, float, sizeY );
TESTING_MALLOC_CPU( Ymagma, float, sizeY );
TESTING_MALLOC_DEV( dA, float, sizeA );
TESTING_MALLOC_DEV( dX, float, sizeX );
TESTING_MALLOC_DEV( dY, float, sizeY );
blocks = (N + nb - 1) / nb;
ldwork = lda * (blocks + 1);
TESTING_MALLOC_DEV( dwork, float, ldwork );
/* Initialize the matrix */
lapackf77_slarnv( &ione, ISEED, &sizeA, A );
magma_smake_symmetric( N, A, lda );
lapackf77_slarnv( &ione, ISEED, &sizeX, X );
lapackf77_slarnv( &ione, ISEED, &sizeY, Y );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_ssetmatrix( N, N, A, lda, dA, lda );
magma_ssetvector( N, X, incx, dX, incx );
magma_ssetvector( N, Y, incy, dY, incy );
cublas_time = magma_sync_wtime( 0 );
cublasSsymv( opts.uplo, N, alpha, dA, lda, dX, incx, beta, dY, incy );
cublas_time = magma_sync_wtime( 0 ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_sgetvector( N, dY, incy, Ycublas, incy );
/* =====================================================================
Performs operation using MAGMA BLAS
=================================================================== */
magma_ssetvector( N, Y, incy, dY, incy );
magma_time = magma_sync_wtime( 0 );
magmablas_ssymv_work( opts.uplo, N, alpha, dA, lda, dX, incx, beta, dY, incy, dwork, ldwork );
// TODO provide option to test non-work interface
//magmablas_ssymv( opts.uplo, N, alpha, dA, lda, dX, incx, beta, dY, incy );
magma_time = magma_sync_wtime( 0 ) - magma_time;
magma_perf = gflops / magma_time;
magma_sgetvector( N, dY, incy, Ymagma, incy );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
cpu_time = magma_wtime();
blasf77_ssymv( &opts.uplo, &N, &alpha, A, &lda, X, &incx, &beta, Y, &incy );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/* =====================================================================
Check the result
=================================================================== */
blasf77_saxpy( &N, &c_neg_one, Y, &incy, Ymagma, &incy );
magma_error = lapackf77_slange( "M", &N, &ione, Ymagma, &N, work ) / N;
blasf77_saxpy( &N, &c_neg_one, Y, &incy, Ycublas, &incy );
cublas_error = lapackf77_slange( "M", &N, &ione, Ycublas, &N, work ) / N;
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e\n",
(int) N,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
magma_error, cublas_error );
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( X );
TESTING_FREE_CPU( Y );
TESTING_FREE_CPU( Ycublas );
TESTING_FREE_CPU( Ymagma );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dX );
TESTING_FREE_DEV( dY );
TESTING_FREE_DEV( dwork );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing ssygvd
*/
int main( int argc, char** argv)
{
TESTING_INIT_MGPU();
float *h_A, *h_Ainit, *h_B, *h_Binit, *h_work;
#if defined(PRECISION_z) || defined(PRECISION_c)
float *rwork;
#endif
float *w1, *w2, result;
magma_int_t *iwork;
float mgpu_time, gpu_time, cpu_time;
/* Matrix size */
magma_int_t N=0, n2;
magma_int_t info;
magma_int_t ione = 1;
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
magma_int_t ISEED[4] = {0,0,0,1};
magma_timestr_t start, end;
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
float tolulp = opts.tolerance * lapackf77_slamch("P");
char jobz = opts.jobz;
int checkres = opts.check;
char uplo = opts.uplo;
magma_int_t itype = opts.itype;
if ( checkres && jobz == MagmaNoVec ) {
fprintf( stderr, "checking results requires vectors; setting jobz=V (option -JV)\n" );
jobz = MagmaVec;
}
printf("using: nrgpu = %d, itype = %d, jobz = %c, uplo = %c, checkres = %d\n",
(int) opts.ngpu, (int) itype, jobz, uplo, (int) checkres);
printf(" N M nr GPU MGPU Time(s) \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;
#if defined(PRECISION_z) || defined(PRECISION_c)
magma_int_t lwork = 2*N + N*N;
magma_int_t lrwork = 1 + 5*N +2*N*N;
// MKL's ssygvd has a bug for small N - it looks like what is returned by a
// query (consistent with LAPACK's number above) is different from a the memory
// requirement ckeck (that returns info -11). The lwork increase below is needed
// to pass this check.
if (N<32)
lwork = 34*32;
#else
magma_int_t lwork = 1 + 6*N + 2*N*N;
#endif
magma_int_t liwork = 3 + 5*N;
TESTING_MALLOC_PIN( h_A, float, n2 );
TESTING_MALLOC_PIN( h_B, float, n2 );
TESTING_MALLOC_PIN( h_work, float, lwork );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_MALLOC_PIN( rwork, float, lrwork );
#endif
TESTING_MALLOC_CPU( w1, float, N );
TESTING_MALLOC_CPU( w2, float, N );
TESTING_MALLOC_CPU( iwork, magma_int_t, liwork );
printf(" N CPU Time(s) GPU Time(s) MGPU Time(s) \n");
printf("==================================================\n");
/* Initialize the matrix */
lapackf77_slarnv( &ione, ISEED, &n2, h_A );
lapackf77_slarnv( &ione, ISEED, &n2, h_B );
magma_smake_hpd( N, h_B, N );
magma_smake_symmetric( N, h_A, N );
if((opts.warmup)||( checkres )){
TESTING_MALLOC_CPU( h_Ainit, float, n2 );
TESTING_MALLOC_CPU( h_Binit, float, n2 );
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_Ainit, &N );
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_Binit, &N );
}
if(opts.warmup){
// ==================================================================
// Warmup using MAGMA.
// ==================================================================
magma_ssygvd_m( opts.ngpu, itype, jobz, uplo,
N, h_A, N, h_B, N, w1,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info);
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_Ainit, &N, h_A, &N );
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_Binit, &N, h_B, &N );
}
// ===================================================================
// Performs operation using MAGMA
// ===================================================================
start = get_current_time();
magma_ssygvd_m( opts.ngpu, itype, jobz, uplo,
N, h_A, N, h_B, N, w1,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info);
end = get_current_time();
if(info != 0)
printf("magma_ssygvd_m returned error %d: %s.\n",
(int) info, magma_strerror( info ));
mgpu_time = GetTimerValue(start,end)/1000.;
if ( checkres ) {
/* =====================================================================
Check the results following the LAPACK's [zc]hegvd 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)
=================================================================== */
#if defined(PRECISION_d) || defined(PRECISION_s)
float *rwork = h_work + N*N;
#endif
result = 1.;
result /= lapackf77_slansy("1",&uplo, &N, h_Ainit, &N, rwork);
result /= lapackf77_slange("1",&N , &N, h_A, &N, rwork);
if (itype == 1){
blasf77_ssymm("L", &uplo, &N, &N, &c_one, h_Ainit, &N, h_A, &N, &c_zero, h_work, &N);
for(int i=0; i<N; ++i)
blasf77_sscal(&N, &w1[i], &h_A[i*N], &ione);
blasf77_ssymm("L", &uplo, &N, &N, &c_neg_one, h_Binit, &N, h_A, &N, &c_one, h_work, &N);
result *= lapackf77_slange("1", &N, &N, h_work, &N, rwork)/N;
}
else if (itype == 2){
blasf77_ssymm("L", &uplo, &N, &N, &c_one, h_Binit, &N, h_A, &N, &c_zero, h_work, &N);
for(int i=0; i<N; ++i)
blasf77_sscal(&N, &w1[i], &h_A[i*N], &ione);
blasf77_ssymm("L", &uplo, &N, &N, &c_one, h_Ainit, &N, h_work, &N, &c_neg_one, h_A, &N);
result *= lapackf77_slange("1", &N, &N, h_A, &N, rwork)/N;
}
else if (itype == 3){
blasf77_ssymm("L", &uplo, &N, &N, &c_one, h_Ainit, &N, h_A, &N, &c_zero, h_work, &N);
for(int i=0; i<N; ++i)
blasf77_sscal(&N, &w1[i], &h_A[i*N], &ione);
blasf77_ssymm("L", &uplo, &N, &N, &c_one, h_Binit, &N, h_work, &N, &c_neg_one, h_A, &N);
result *= lapackf77_slange("1", &N, &N, h_A, &N, rwork)/N;
}
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_Ainit, &N, h_A, &N );
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_Binit, &N, h_B, &N );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
start = get_current_time();
magma_ssygvd(itype, jobz, uplo,
N, h_A, N, h_B, N, w2,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info);
end = get_current_time();
if(info != 0)
printf("magma_ssygvd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
gpu_time = GetTimerValue(start,end)/1000.;
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
start = get_current_time();
lapackf77_ssygvd(&itype, &jobz, &uplo,
&N, h_Ainit, &N, h_Binit, &N, w2,
h_work, &lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, &lrwork,
#endif
iwork, &liwork,
&info);
end = get_current_time();
if (info != 0)
printf("lapackf77_ssygvd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
cpu_time = GetTimerValue(start,end)/1000.;
float temp1 = 0;
float temp2 = 0;
for(int j=0; j<N; j++){
temp1 = max(temp1, absv(w1[j]));
temp1 = max(temp1, absv(w2[j]));
temp2 = max(temp2, absv(w1[j]-w2[j]));
}
float result2 = temp2 / (((float)N)*temp1);
/* =====================================================================
Print execution time
=================================================================== */
printf("%5d %6.2f %6.2f %6.2f\n",
(int) N, cpu_time, gpu_time, mgpu_time);
printf("Testing the eigenvalues and eigenvectors for correctness:\n");
if(itype==1)
printf("(1) | A Z - B Z D | / (|A| |Z| N) = %8.2e%s\n", result, (result < tol ? "" : " failed") );
else if(itype==2)
printf("(1) | A B Z - Z D | / (|A| |Z| N) = %8.2e%s\n", result, (result < tol ? "" : " failed") );
else if(itype==3)
printf("(1) | B A Z - Z D | / (|A| |Z| N) = %8.2e%s\n", result, (result < tol ? "" : " failed") );
printf( "(3) | D(MGPU)-D(LAPACK) |/ |D| = %8.2e%s\n\n", result2, (result2 < tolulp ? "" : " failed") );
}
else {
printf("%5d ------ ------ %6.2f\n",
(int) N, mgpu_time);
}
/* Memory clean up */
TESTING_FREE_PIN( h_A );
TESTING_FREE_PIN( h_B );
TESTING_FREE_PIN( h_work );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_FREE_PIN( rwork );
#endif
TESTING_FREE_CPU( w1 );
TESTING_FREE_CPU( w2 );
TESTING_FREE_CPU( iwork );
if((opts.warmup)||( checkres )){
TESTING_FREE_CPU( h_Ainit );
TESTING_FREE_CPU( h_Binit );
}
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
/* Shutdown */
TESTING_FINALIZE_MGPU();
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing magma_ssymm_mgpu
*/
int main( int argc, char** argv)
{
TESTING_INIT();
float c_neg_one = MAGMA_S_NEG_ONE;
float calpha = MAGMA_S_MAKE( 3.456, 5.678 );
float cbeta = MAGMA_S_MAKE( 1.234, 2.456 );
real_Double_t gflops, gpu_perf=0., cpu_perf=0., gpu_time=0., cpu_time=0.;
real_Double_t gpu_perf2=0., gpu_time2=0.;
float error=0., errorbis=0., work[1];
float *hA, *hX, *hB, *hR;
float *dA[MagmaMaxGPUs], *dX[MagmaMaxGPUs], *dB[MagmaMaxGPUs], *dwork[MagmaMaxGPUs], *hwork[MagmaMaxGPUs+1];
float *dA2;
magma_int_t M, N, size, lda, ldda, msize, nb, nstream;
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");
// default values
nb = (opts.nb > 0 ? opts.nb : 64);
nstream = (opts.nstream > 0 ? opts.nstream : 2);
magma_int_t gnode[MagmaMaxGPUs][MagmaMaxGPUs+2];
magma_int_t nbcmplx = 0;
magma_buildconnection_mgpu(gnode, &nbcmplx, opts.ngpu);
printf("Initializing communication pattern... GPU-ncmplx %d\n\n", (int) nbcmplx);
for (int i=0; i < nbcmplx; ++i) {
int myngpu = gnode[i][MagmaMaxGPUs];
printf("cmplx %d has %d gpu ", i, myngpu);
for(int j=0; j < myngpu; ++j)
printf(" %d", (int) gnode[i][j]);
printf("\n");
}
magma_int_t nbevents = 2;
magma_queue_t streams[MagmaMaxGPUs][20];
magma_event_t redevents[MagmaMaxGPUs][20];
magma_event_t redevents2[MagmaMaxGPUs][MagmaMaxGPUs*MagmaMaxGPUs+10];
for( int d = 0; d < opts.ngpu; ++d ) {
for( magma_int_t i = 0; i < nstream; ++i ) {
magma_queue_create( &streams[d][i] );
}
for( magma_int_t i = 0; i < nbevents; ++i ) {
cudaEventCreateWithFlags(&redevents[d][i], cudaEventDisableTiming);
cudaEventCreateWithFlags(&redevents2[d][i], cudaEventDisableTiming);
}
}
printf( "nb %d, ngpu %d, nstream %d version %d\n", (int) nb, (int) opts.ngpu, (int) nstream, (int) opts.version );
printf(" M N nb offset CPU GFlop/s (sec) GPU GFlop/s (sec) CUBLAS hemm (sec) ||R|| / ||A||*||X||\n");
printf("=========================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
M = opts.msize[itest];
N = opts.nsize[itest];
for( int offset = 0; offset < N; offset += min(N,nb) ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
msize = M - offset;
lda = M;
ldda = ((M + 31)/32)*32;
size = lda*M;
gflops = FLOPS_SSYMM( MagmaLeft, (float)msize, (float)N ) / 1e9;
magma_int_t dworksiz = ldda*N*3;
magma_int_t hworksiz = lda*N;
TESTING_MALLOC_CPU( hA, float, lda*M );
TESTING_MALLOC_CPU( hX, float, lda*N );
TESTING_MALLOC_CPU( hB, float, lda*N );
TESTING_MALLOC_PIN( hR, float, lda*N );
for( int d = 0; d < opts.ngpu; ++d ) {
magma_int_t mlocal = ((M / nb) / opts.ngpu + 1) * nb;
magma_setdevice( d );
TESTING_MALLOC_DEV( dA[d], float, ldda*mlocal );
TESTING_MALLOC_DEV( dX[d], float, ldda*N );
TESTING_MALLOC_DEV( dB[d], float, ldda*N );
TESTING_MALLOC_DEV( dwork[d], float, dworksiz );
TESTING_MALLOC_PIN( hwork[d], float, hworksiz );
}
TESTING_MALLOC_PIN( hwork[opts.ngpu], float, lda*N );
if ( opts.check ) {
magma_setdevice( 0 );
TESTING_MALLOC_DEV( dA2, float, ldda*M );
}
lapackf77_slarnv( &ione, iseed, &size, hA );
magma_smake_symmetric( M, hA, lda );
size = lda*N;
lapackf77_slarnv( &ione, iseed, &size, hX );
lapackf77_slarnv( &ione, iseed, &size, hB );
lapackf77_slacpy( "Full", &M, &N, hB, &lda, hR, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_ssetmatrix_1D_col_bcyclic( M, M, hA, lda, dA, ldda, opts.ngpu, nb );
for( int d = 0; d < opts.ngpu; ++d ) {
magma_setdevice( d );
//magmablasSetKernelStream( streams[ d ][ 0 ] );
magma_ssetmatrix( M, N, hX, lda, dX[d], ldda );
//if (d == 0) magma_ssetmatrix( M, N, hB, lda, dB[d], ldda ); // this is wrong coz when offset != 0 the gpu who do the beta*C may be not 0 so this should be related to stdev(starting device who own i=0 first col)
magma_ssetmatrix( M, N, hB, lda, dB[d], ldda );
}
//memset(hR, 0, lda*N*sizeof(float));
//trace_init( 1, opts.ngpu, nstream, (magma_queue_t*) streams );
//magma_int_t offset = 0; //nb;
gpu_time = magma_sync_wtime(0);
magmablas_ssymm_mgpu_com(
MagmaLeft, MagmaLower, msize, N,
calpha, dA, ldda, offset,
dX, ldda,
cbeta, dB, ldda, dwork, dworksiz, hR, lda, hwork, hworksiz,
opts.ngpu, nb, streams, nstream, redevents2, nbevents, gnode, nbcmplx);
gpu_time = magma_sync_wtime(0) - gpu_time;
gpu_perf = gflops / gpu_time;
#ifdef TRACING
char buf[80];
snprintf( buf, sizeof(buf), "ssymm-m%d-n%d-nb%d-stream%d-ngpu%d-run%d.svg",
(int) M, (int) N, (int) nb, (int) nstream, (int) opts.ngpu, (int) j );
trace_finalize( buf, "trace.css" );
#endif
/* ====================================================================
Performs operation using CUBLAS
=================================================================== */
if ( opts.check && iter == 0 ) {
magma_setdevice( 0 );
magmablasSetKernelStream( 0 );
magma_ssetmatrix( M, M, hA, lda, dA2, ldda );
magma_ssetmatrix( M, N, hX, lda, dX[0], ldda );
magma_ssetmatrix( M, N, hB, lda, dwork[0], ldda );
gpu_time2 = magma_sync_wtime(0);
magma_ssymm(
MagmaLeft, MagmaLower, msize, N,
calpha, dA2+offset*ldda+offset, ldda,
dX[0], ldda,
cbeta, dwork[0], ldda );
gpu_time2 = magma_sync_wtime(0) - gpu_time2;
gpu_perf2 = gflops / gpu_time2;
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.check ) {
// store ||A||*||X||
errorbis = lapackf77_slange("fro", &msize, &msize, hA+offset*lda+offset, &lda, work );
errorbis *= lapackf77_slange("fro", &msize, &N, hX, &lda, work );
//printf( "A =" ); magma_sprint( M, M, hA, lda );
//printf( "X =" ); magma_sprint( M, N, hX, lda );
//printf( "B =" ); magma_sprint( M, N, hB, lda );
cpu_time = magma_wtime();
blasf77_ssymm( "Left", "Lower", &msize, &N,
&calpha, hA+offset*lda+offset, &lda,
hX, &lda,
&cbeta, hB, &lda );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/*
trace_file = fopen("AJETE/C", "w");
for (int j = 0; j < N; j++)
for (int i = 0; i < siz; i++)
fprintf(trace_file, "%10d%10d%40.30e\n", i+1, j+1, hB[j*lda+i]);
fclose(trace_file);
*/
magma_int_t firstprint=0;
for(magma_int_t dev=0; dev < opts.ngpu; ++dev) {
magma_setdevice( dev );
magma_sgetmatrix( M, N, dB[dev], ldda, hR, lda );
// compute relative error ||R||/||A||*||X||, where R := B_magma - B_lapack = R - B
size = lda*N;
blasf77_saxpy( &size, &c_neg_one, hB, &ione, hR, &ione );
error = lapackf77_slange("fro", &msize, &N, hR, &lda, work) / errorbis;
//printf( "R =" ); magma_sprint( M, N, hR, lda );
if (firstprint == 0) {
printf( "%5d %5d %5d %5d %7.1f (%7.4f) %7.1f (%7.4f) %7.1f (%7.4f) %8.2e %s\n",
(int) M, (int) N, (int) nb, (int) offset,
cpu_perf, cpu_time,
gpu_perf, gpu_time,
gpu_perf2, gpu_time2,
error, (error < tol ? "ok" : "failed") );
}
else {
printf( "%89s %8.2e %s\n", " ",
error, (error < tol ? "ok" : "failed") );
}
status += ! (error < tol);
firstprint =1;
}
} else {
printf( "%5d %5d %5d %5d --- ( --- ) %7.1f (%7.4f) --- ( --- ) ---\n",
(int) M, (int) N, (int) nb, (int) offset,
gpu_perf, gpu_time );
}
TESTING_FREE_CPU( hA );
TESTING_FREE_CPU( hX );
TESTING_FREE_CPU( hB );
TESTING_FREE_PIN( hR );
for( int d = 0; d < opts.ngpu; ++d ) {
magma_setdevice( d );
TESTING_FREE_DEV( dA[d] );
TESTING_FREE_DEV( dX[d] );
TESTING_FREE_DEV( dB[d] );
TESTING_FREE_DEV( dwork[d] );
TESTING_FREE_PIN( hwork[d] );
}
TESTING_FREE_PIN( hwork[opts.ngpu] );
if ( opts.check ) {
magma_setdevice( 0 );
TESTING_FREE_DEV( dA2 );
}
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
} // offset
printf( "\n" );
}
for( int d = 0; d < opts.ngpu; ++d ) {
magma_setdevice( d );
for( magma_int_t i = 0; i < nstream; ++i ) {
magma_queue_destroy( streams[d][i] );
}
for( magma_int_t i = 0; i < nbevents; ++i ) {
magma_event_destroy( redevents[d][i] );
magma_event_destroy( redevents2[d][i] );
}
}
TESTING_FINALIZE();
return status;
}
